#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 + (isizesize), &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 + jn)] != 0)
{
// swap rows
aux = mat[in+idx];
mat[in+idx] = mat[jn + idx];
mat[jn + 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 + ni)] / mat[(i + i * n)]);
__syncthreads();
mat[(k + nidx)] -= mul * mat[i + nidx];
}


if(i == idx) {
// calculate the determinant by multiplying the diagonal elements
det = mat[idxn + 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);
}


change this code to do bitonic sort with this discription:
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

Give me the full code solution