CUDA Vision Studio Trouble with Kernel

Question

I have a project where I have to create a program in Visual Studio/CUDA using GPU Threads and Blocks that contain 2 random Arrays A and B used for computation before storing the result in a third Array C the values based on this equation:

C(1,j)=max(A(:,j)+max(B(:,j)) (Note: the ":" operator means for all the lines)

Here is my Kernel function

           __global__ void mykernel(int **a, int **b,int **c,const int width)
             {
int col= threadIdx.x;
int tempa=0;
int tempb=0;

for(int k=0;k<width;k++){

    int maxA=a[k][col];

    if  (maxA>tempa){
    tempa=maxA;
    }

    int maxB=b[k][col];

    if  (maxB>tempb){
    tempb=maxB;
    }

    }


c[0][col] =tempa+tempb;
}

And my main

      int main()
       {
const int  dim= 5;
const int rows=5;
const int columns=5;
size_t dsize = rows*columns*sizeof(int);
//Αντίγραφα πινάκων δεδομένων της CPU
int *A[dim][dim];
int *B[dim][dim];
int *C[1][dim];

//Αντίγραφα πινάκων δεδομένων της GPU
int *d_A[dim][dim],*d_B[dim][dim],*d_C[1][dim]; 

//Εξασφάλιση μνήμης για τα αντίγραφα δεδομένων της CPU 
A[dim][dim]= (int *)malloc(dsize); 
B[dim][dim] = (int *)malloc(dsize); 
C[1][dim]= (int *)malloc(dsize);

//Γέμισμα των πινάκων με τυχαίες τιμές μεταξυ 
 for (int i=0;i<rows;i++)
   for (int j=0;j<columns;j++){


  *A[i][j]=rand() %5+1;
  *B[i][j]=rand() %5+1;

 }

//Εξασφάλιση μνήμης για τα αντίγραφα δεδομένων της GPU και αντιγραφή δεδομένων CPU προς GPU  
cudaMalloc((void **)&d_A, dsize);
cudaMemcpy(d_A, A, dsize, cudaMemcpyHostToDevice);

cudaMalloc((void **)&d_B, dsize);
cudaMemcpy(d_B, B, dsize, cudaMemcpyHostToDevice);

cudaMalloc((void **)&d_C, dsize);

//Κλήση Kernel συνάρτησης στην GPU με χρήση 5 Threads σε 1 Block

mykernel<<<1,5>>>(d_A,d_B,d_C,dim);

//Αντιγραφή αποτελέσματος στην μνήμη της CPU
cudaMemcpy(C, d_C, dsize, cudaMemcpyDeviceToHost);

//Εκκαθάριση Μνήμης για CPU και GPU
free(A);
free(B);
free(C);

cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

while(1){};
return 0;
  }

I think I got the algorithm right but in this line I get the following error:

Line

mykernel<<<1,5>>>(d_A,d_B,d_C,dim);

Error

argument of type "int *(*)[5]" is incompatible with parameter of type "int **"

Any suggestions on what I should do?

Answer 1

1. First of all, any time you are having trouble with a CUDA code, I recommend using proper CUDA error checking and running your code with cuda-memcheck (e.g. from the command line).

2. This isn't doing what you think:

   int *A[dim][dim];
   

   this is creating a 2-dimensional array of pointers. What you want is a pointer to a 2-dimensional array (of int).

3. In C, when you define an array, you cannot assign anything to elements at the array dimension:

   A[dim][dim]= (int *)malloc(dsize); 
   

   arrays only go up to size of dim- 1. A[dim][dim] does not exist and is out-of-range for your array definition. This appears to be related to your general confusion around using 2D arrays in C or C++.

4. This line is similarly broken (assigning a numerical value to a region associated with an unallocated pointer) and is further evidence of your confusion around 2D arrays in C:

   *A[i][j]=rand() %5+1;
   

5. Your handling of d_A, d_B and d_C are similarly broken.

It's evident that you are wanting to use a 2D array in your CUDA kernel, so the right approach here is to probably pick a method from a set of canonical examples. Since you appear to know the array dimensions at compile time, we'll leverage that. Here is a fully worked example showing the mods:

$ cat t1344.cu
#include <iostream>
const int dim = 5;
typedef int my_arr[dim];

__global__ void mykernel(my_arr *a, my_arr *b, my_arr *c,const int width)
{
int col= threadIdx.x;
int tempa=0;
int tempb=0;

for(int k=0;k<width;k++){

    int maxA=a[k][col];

    if  (maxA>tempa){
    tempa=maxA;
    }

    int maxB=b[k][col];

    if  (maxB>tempb){
    tempb=maxB;
    }

    }


c[0][col] =tempa+tempb;
}

int main()
       {
const int rows=dim;
const int columns=dim;
size_t dsize = rows*columns*sizeof(int);
//Αντίγραφα πινάκων δεδομένων της CPU;
my_arr *A, *B, *C;

//Αντίγραφα πινάκων δεδομένων της GPU
my_arr *d_A,*d_B,*d_C;

//Εξασφάλιση μνήμης για τα αντίγραφα δεδομένων της CPU
A = (my_arr *)malloc(dsize);
B = (my_arr *)malloc(dsize);
C = (my_arr *)malloc(dsize);

//Γέμισμα των πινάκων με τυχαίες τιμές μεταξυ
 for (int i=0;i<rows;i++)
   for (int j=0;j<columns;j++){


  A[i][j]=rand() %5+1;
  B[i][j]=rand() %5+1;

 }

//Εξασφάλιση μνήμης για τα αντίγραφα δεδομένων της GPU και αντιγραφή δεδομένων CPU προς GPU
cudaMalloc((void **)&d_A, dsize);
cudaMemcpy(d_A, A, dsize, cudaMemcpyHostToDevice);

cudaMalloc((void **)&d_B, dsize);
cudaMemcpy(d_B, B, dsize, cudaMemcpyHostToDevice);

cudaMalloc((void **)&d_C, dsize);

//Κλήση Kernel συνάρτησης στην GPU με χρήση 5 Threads σε 1 Block

mykernel<<<1,5>>>(d_A,d_B,d_C,dim);

//Αντιγραφή αποτελέσματος στην μνήμη της CPU
cudaMemcpy(C, d_C, dsize, cudaMemcpyDeviceToHost);
for (int i = 0; i < dim; i++)  std::cout << C[0][i] << std::endl;
//Εκκαθάριση Μνήμης για CPU και GPU
free(A);
free(B);
free(C);

cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);

return 0;
  }
$ nvcc -arch=sm_35 -o t1344 t1344.cu
$ cuda-memcheck ./t1344
========= CUDA-MEMCHECK
9
8
8
9
8
========= ERROR SUMMARY: 0 errors
$

I haven't fully verified the results, but they seem plausible to me.