Hi,
Change the following code to this operation of sorting wit bitonicSort, The kernels to be written should sort the contents of the file datSeq1M.bin in 10 iteration
steps. In iteration 1, 1024 subsequences of 1024 integer values are to be sorted in parallel. In
iteration 2, 512 subsequences of 2048 integer values are to be sorted in parallel based on the
merging of the previously sorted halves. And so on and so forth, until iteration 10 is reached,
where 2 previously sorted subsequences of 512K values are merged, yielding the whole sorted
sequence.
Assume that the values of the whole sequence, when stored in memory, are organized as the
elements of 1024 X 1024 matrix. The threads in a block thread process successive matrix rows

This is the code that has to be changed:

/**
 *   Hugo Leal (93059) and Luísa Amaral (93001), June 2022
 *   Matrices are processed by rows
 */
#include <stdio.h>
#include <stdlib.h>
#include <sys/types.h>
#include <libgen.h>
#include <unistd.h>
#include <stdbool.h>
#include <math.h>
#include <string.h>
#include <time.h>


#include "common.h"
#include <cuda_runtime.h>

/* allusion to internal functions */
static void calculateDeterminantCPU(double *matrix, double *determinant, int size);
__global__ void calculateDeterminantGPUKernel(double *mat, double *determinants, int n);
static double get_delta_time(void);
static void printUsage(char *cmdName);
int parseArgs(int argc, char* argv[], char *filenames[]);

/**
 *   main program
 */
int main(int argc, char **argv)
{
    printf("%s Starting...\n", argv[0]);
    if (sizeof(unsigned int) != (size_t)4)
        return 1; // it fails with prejudice if an integer does not have 4 bytes
    
    /* set up the device */
    int dev = 0;

    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("Using Device %d: %s\n", dev, deviceProp.name);
    CHECK(cudaSetDevice(dev));


    int nArgs = (int)(argc - 1) / 2 + 50; /* size of the array of file names */
    char *fileNames[nArgs];               /* Stores the provided file names  */
    memset(fileNames, 0, nArgs);          /* Initializes the array with zeros */
    parseArgs(argc, argv, fileNames);     /* parse command line arguments */

    int num, size;
    size_t matSize, matAreaSize;
    size_t numDet;

    double *mat;                          /* stores the matrices coeficients on host side */
    double *mat_device;                   /* stores the matrices coeficients on device side */
    double *det;                          /* used to initialize determinants data on host side */
    double *det_device;                   /* stores the determinants on device side */
    double *det_cpu;                      /* stores the determinants calculated with CPU kernel */
    double *new_det;                      /* stores the determinants from device side*/


    /* Open file */
    FILE *f;
    f = fopen(fileNames[0], "rb");

    if (!f)
    {
        printf("Error opening file %s\n", fileNames[0]);
        exit(EXIT_FAILURE);
    }

    printf("File '%s'\n", fileNames[0]); //only the first file is parsed

    if(!fread(&num, sizeof(int), 1, f)) // num of matrices in file
    {
        printf("Error reading from file %s\n", fileNames[0]);
        exit(EXIT_FAILURE);
    } 
    printf("Number of matrices to be read = %d\n", num);

    if(!fread(&size, sizeof(int), 1, f)) // size of each matrix
    {
        printf("Error reading from file %s\n", fileNames[0]);
        exit(EXIT_FAILURE);
    } 
    printf("Matrices order = %d\n\n", size);

    /* create memory areas in host and device memory where the coefficients will be stored */
    matSize = (size_t)size * (size_t)size * sizeof(double);
    matAreaSize = (size_t)num * matSize;
    numDet = (size_t)num * sizeof(double);

    if ((matAreaSize + numDet) > (size_t)1.3e9)
    {
        fprintf(stderr, "The GeForce GTX 1660 Ti cannot handle more than 5GB of memory!\n");
        exit(1);
    }

    mat = (double *)malloc(matAreaSize);
    det = (double *)malloc(numDet);
    CHECK(cudaMalloc((void **)&mat_device, matAreaSize));
    CHECK(cudaMalloc((void **)&det_device, numDet));
    

    
    /* initialize the host data */
    for (int j = 0; j < num * size * size; j++)
    {
        char bytes[8];
        if(!fread(&bytes, sizeof(double), 1, f))
        {
            printf("Error reading from file %s\n", fileNames[0]);
            exit(EXIT_FAILURE);
        }
        double d = *((double *)bytes);
        mat[j] = d;
    }
    fclose(f);

    for(int j= 0; j < num; j++) {
        det[j] = 1;
    }

    
    /* copy the host data to the device memory */
    (void) get_delta_time ();
    CHECK (cudaMemcpy (mat_device, mat, matAreaSize, cudaMemcpyHostToDevice));
    CHECK (cudaMemcpy (det_device, det, numDet, cudaMemcpyHostToDevice));
    printf ("The transfer of %ld bytes from the host to the device took %.3e seconds\n",
          (long) matAreaSize + (long) numDet, get_delta_time ());


    /* run the computational kernel */
    unsigned int gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ;
    blockDimX = size;                                      
    blockDimY = 1;                                      
    blockDimZ = 1;                                      
    gridDimX = num;                                      
    gridDimY = 1;                                       
    gridDimZ = 1;                                       

    dim3 grid (gridDimX, gridDimY, gridDimZ);
    dim3 block (blockDimX, blockDimY, blockDimZ);

    (void) get_delta_time ();
    calculateDeterminantGPUKernel<<<grid, block>>>(mat_device, det_device, size);
    CHECK (cudaDeviceSynchronize ());                            // wait for kernel to finish
    CHECK (cudaGetLastError ());                                 // check for kernel errors
    printf("The CUDA kernel <<<(%d,%d,%d), (%d,%d,%d)>>> took %.3e seconds to run\n",
         gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, get_delta_time ());

    
    
    /* copy kernel result back to host side */
    new_det = (double *) malloc (numDet);
    (void) get_delta_time ();
    CHECK (cudaMemcpy (new_det, det_device, numDet, cudaMemcpyDeviceToHost));
    printf ("The transfer of %ld bytes from the device to the host took %.3e seconds\n",
          (long) numDet, get_delta_time ());

    
    /* free device global memory */
    CHECK(cudaFree(mat_device));
    CHECK(cudaFree(det_device));
    
    /* reset the device */
    CHECK(cudaDeviceReset());
    
    /* compute the determinants on the CPU */
    det_cpu = (double *)malloc(numDet);
    (void) get_delta_time ();
    for(int i = 0; i < num; i++) 
    {
        det_cpu[i] = 1;
        calculateDeterminantCPU(mat + (i*size*size), &det_cpu[i], size);
    }
    printf("The single threaded cpu kernel took %.3e seconds to run \n",get_delta_time ());

    for(int i = 0; i < num; i++) {  
        printf("\nDeterminant of matrix %d   GPU: %.3E   CPU: %.3E\n", i+1, new_det[i], det_cpu[i]);
    }

    free(mat);
    free(det);
    free(new_det);

    exit(EXIT_SUCCESS);

    return 0;
}

/**
 * \brief Method calculateDeterminantGPUKernel.
 *
 * Its role is to calculate the determinants of the matrix using a GPU
 *
 *
 * \param mat pointer to the matrix coefficients
 * \param determinants pointer to struct that holds the calculated determinants
 * \param n order of matrix
 */
__global__ void calculateDeterminantGPUKernel(double *mat, double *determinants, int n)
{

    int bkx = blockIdx.x + gridDim.x * blockIdx.y + gridDim.x * gridDim.y * blockIdx.z;
    int pos = bkx * n * n; // point to the first coefficient of the current matrix
    mat += pos;            
    bool reverseSignal = false;
    double aux;

    int idx;
    idx = threadIdx.x + blockDim.x * threadIdx.y + blockDim.x * blockDim.y * threadIdx.z;
    double det;

    for (int i = 0; i <= idx; i++) // each loop has a smaller matrix to process (submatrix)
    {
        // in case of a 0 in a diagonal
        if (mat[(i * n + i)] == 0) // test diagonal cell in current row
        {
            bool allNull = true;
            for (int j = i + 1; j < n; j++)
            {
                if (mat[(i + j*n)] != 0)
                {
                    // swap rows
                    aux = mat[i*n+idx];
                    mat[i*n+idx] = mat[j*n + idx];
                    mat[j*n + idx] = aux;
                    
                    
                    __syncthreads();
                    reverseSignal = !reverseSignal;
                    allNull = false;
                }
            }
            // if the matrix contains a full column of zeros, anywhere in the matrix, the determinant will be 0
            if (allNull)
            {
                determinants[bkx] = 0;
                break;
            }
        }

        double mul;
        for (int k = i + 1; k < n; k++)
        {
            mul = (mat[(k + n*i)] / mat[(i + i * n)]);
            __syncthreads();
            mat[(k + n*idx)] -= mul * mat[i + n*idx];
        }
        

        if(i == idx) {
            // calculate the determinant by multiplying the diagonal elements
            det = mat[idx*n + idx];
            determinants[bkx] *= det;
            if (reverseSignal)
            {
                determinants[bkx] *= -1;
                reverseSignal =  false;
            }
        }
    }    
}

/**
 * \brief Method calculateDeterminantCPU.
 *
 * Its role is to calculate the determinant of a matrix
 * 
 *
 * \param matrix pointer to the memory region with the matrix coefficients
 * \param determinant pointer that holds the calculated determinant
 * \param size order of the matrix
 */
static void calculateDeterminantCPU(double *matrix, double *determinant, int size)
{
    bool reverseSignal = false;

    for (int i = 0; i < size - 1; i++)
    {
        if (matrix[i * size + i] == 0)
        {
            bool allNull = true;
            for (int j = i + 1; j < size; j++)
            {
                if (matrix[j * size + i] != 0)
                {
                    // swap rows
                    int aux;
                    for (int k = 0; k < size; k++)
                    {
                        aux = matrix[(i-1) * size + k];
                        matrix[(i-1) * size + k] = matrix[(j-1) * size + k];
                        matrix[(j-1) * size + k] = aux;
                    }
                    reverseSignal = !reverseSignal;
                    allNull = false;
                }
            }
            // if the matrix contains a full column of zeros, anywhere in the matrix, the determinant will be 0
            if (allNull) 
            {
                
                *determinant = 0;
                return;
            }
        }
        double aux;
        for (int k = i + 1; k < size; k++)
        {
            aux = (matrix[i * size + k] / matrix[i * size + i]);
            for (int j = 0; j < size; j++)
            {
                matrix[j * size + k] -= aux * matrix[j * size + i];
            }
        }
    }

    double prod = 1;
    for (int i = 0; i < size; i++)
    {
        prod *= matrix[i * size + i];
    }

    if(reverseSignal) {
        prod *= -1;
    }
    
    *determinant = prod;
    
}


/**
 *  \brief parseArgs function.
 *
 *  \param argc number of words of the command line
 *  \param argv list of words of the command line
 *  \param filenames pointer to store names of files
 *
 *  \return status of operation
 */
int parseArgs(int argc, char* argv[], char *filenames[])
{
    /* process command line options */

    int opt;                 /* selected option */
    // char *fName = "no name"; /* file name (initialized to "no name" by default) */
    int cnt = 0;

    do
    {
        switch ((opt = getopt(argc, argv, "f:n:h")))
        {
        case 'f': /* file name */
            if (optarg[0] == '-')
            {
                fprintf(stderr, "%s: file name is missing\n", basename(argv[0]));
                printUsage(basename(argv[0]));
                return EXIT_FAILURE;
            }
            filenames[cnt] = optarg;
            cnt++;
            break;
        case 'h': /* help mode */
            printUsage(basename(argv[0]));
            return EXIT_SUCCESS;
        case '?': /* invalid option */
            fprintf(stderr, "%s: invalid option\n", basename(argv[0]));
            printUsage(basename(argv[0]));
            return EXIT_FAILURE;
        case -1:
            break;
        }
    } while (opt != -1);
    
    if (argc == 1)
    {
        fprintf(stderr, "%s: invalid format\n", basename(argv[0]));
        printUsage(basename(argv[0]));
        return EXIT_FAILURE;
    }

    /* that's all */

    return EXIT_SUCCESS;

} /* end of parseArgs */


/**
 *  \brief Print command usage.
 *
 *  A message specifying how the program should be called is printed.
 *
 *  \param cmdName string with the name of the command
 */

void printUsage(char *cmdName)
{
    fprintf(stderr, "\nSynopsis: %s OPTIONS [filename / positive number]\n"
                    "  OPTIONS:\n"
                    "  -h      --- print this help\n"
                    "  -f      --- filename\n",
            cmdName);
}


static double get_delta_time(void)
{
  static struct timespec t0,t1;

  t0 = t1;
  if(clock_gettime(CLOCK_MONOTONIC,&t1) != 0)
  {
    perror("clock_gettime");
    exit(1);
  }
  return (double)(t1.tv_sec - t0.tv_sec) + 1.0e-9 * (double)(t1.tv_nsec - t0.tv_nsec);
}