%%cu
#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 "/content/common.h"
#include <cuda_runtime.h>

#define ROWS 1024
#define COLS 1024
#define ITERATIONS 10

__global__ void bitonicSortKernel(int *mat, int iteration, bool dir)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int offset = idx * COLS;
    int i, j, k;
    int temp;

    for (i = 1 << (iteration - 1); i > 0; i = i >> 1)
    {
        for (j = i; j > 0; j = j >> 1)
        {
            for (k = 0; k < (COLS >> 1); ++k)
            {
                int ixj = k ^ j;
                if ((k & i) == 0 && (ixj > k) && (dir == (mat[offset + k] > mat[offset + ixj])))
                {
                    // Swap elements
                    temp = mat[offset + k];
                    mat[offset + k] = mat[offset + ixj];
                    mat[offset + ixj] = temp;
                }
            }
            __syncthreads();
        }
    }
}

bool isSorted(int* data, int size) {
  for (int i = 0; i < size - 1; i++) {
    if (data[i] > data[i + 1])
      return false;
  }
  return true;
}

int main(int argc, char **argv)
{
    printf("%s Starting…\n", argv[0]);

    int *mat;

    mat = (int *)malloc(sizeof(int) * ROWS * COLS);

    // Read the input data from file
    FILE* inputFile = fopen("/content/datSeq1M.bin", "rb");
    fseek(inputFile, sizeof(int), SEEK_SET); // Skip the first integer
    fread(mat, sizeof(int), ROWS * COLS, inputFile);
    fclose(inputFile);

    int *mat_device;
    CHECK(cudaMalloc((void**)&mat_device, sizeof(int) * ROWS * COLS));
    CHECK(cudaMemcpy(mat_device, mat, sizeof(int) * ROWS * COLS, cudaMemcpyHostToDevice));

    dim3 grid(ROWS / COLS, 1, 1);
    dim3 block(COLS, 1, 1);
    int numPairs = ROWS;

    for (int i = 0; i < ITERATIONS; ++i)
    {
        for (int j = 0; j < i; ++j)
        {
            bool dir = (numPairs & 0x01) == 0;
            bitonicSortKernel<<<grid, block>>>(mat_device, i - j + 1, dir);
            CHECK (cudaDeviceSynchronize ());                            // wait for kernel to finish
            CHECK (cudaGetLastError ());                                 // check for kernel errors
        }
        numPairs = numPairs >> 1;
    }

    CHECK(cudaMemcpy(mat, mat_device, sizeof(int) * ROWS * COLS, cudaMemcpyDeviceToHost));
    CHECK(cudaFree(mat_device));
    // Print the sorted data
    for (int i = 0; i < 20; i++) {
      printf("%d ", mat[i]);
      if( i == 10 ) break;
    }
    printf("\n");
    // Check if the array is sorted
    bool isOrdered = isSorted(mat, ROWS * COLS);
 
    // Print the result of the ordering check
    if (isOrdered) {
      printf("The array is ordered.\n");
    } else {
      printf("The array is not ordered.\n");
    }
 
    free(mat);
    return 0;
}

This isnt ordering the array, please fix it by keeping on mind this:
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