passing and using 2d arrays in kernel 2d arrays in struct

Hey guys,

I’m trying to use a few 2d arrays which I put into a struct, filled with data, and sent over to the kernel.

I have confirmed that the data is being passed into the kernel and is correct.

However once inside the kernel I am having issues passing the arrays from function to function.

Right now I’m trying to use the following code, but getting many odd errors here and there.

obviously my struct is dev_input and the arrays are all 10x100.

double *Isomer1_arr[10][100] = dev_input->Isomer1_arr[][];

double *Isomer2_arr[10][100] = dev_input->Isomer2_arr[][];

double *Shift_arr[10][100] = dev_input->Shift_arr[][];

double *SP_arr[10][100] = dev_input->sp_arr[][];

for the above section of code the error im getting is

“initialization with “{…}” expected for aggregate object”

Trying to pass them into a function I have the following…

dev_E1 = getInteraction(1, 1, dev_x1n, dev_y1n, dev_z1n, dev_thetax1n, dev_thetay1n, dev_thetaz1n, Isomer1_arr, Isomer2_arr, Shift_arr, sp_arr);

For this part I’m getting

“error : argument of type “double ()[100]” is incompatible with parameter of type “double (*)[100]””

I’m sure I’m making a stupid mistake but I want another opinion on this.

Thanks!

Also I realize that the one variable

double *SP_arr[10][100] = dev_input->sp_arr[][];

should reflect the change in the function call.

I have already changed that.

You are trying to make it a lot more complicated than it is. The simplest way to pass a 2D array to a kernel is to just pass a pointer to the first element in the array, and then use a 2D->1D indexing scheme.

Example:

__global__

void my_kernel(double* Isomer1_arr,int xdim,int ydim)

{

   unsigned int idx = threadIdx.x;

   unsigned int idy = threadIdx.y;

   unsigned int gidx = blockIdx.x*blockDim.x+idx;

   unsigned int gidy = blockIdx.y*blockDim.y+idy;

double local_data;

if((gidx < xdim)&&(gidy < ydim))

   {

       local_data = Isomer1_arr[gidx+xdim*gidy];  

   }

}

int xdim = 10;

int ydim = 100;

dim3 cudaBlockSize(8,8,1);

dim3 cudaGridSize((8+xdim-1)/8,(8+ydim-1)/8,1);

double* Isomer1_arr;

cudaMalloc((void**)&Isomer1_arr,sizeof(double)*xdim*ydim);

CUDA_SAFE_CALL(my_kernel<<<cudaGridSize,cudaBlockSize>>>(Isomer1_arr,xdim,ydim));

Basically you never really want to make a “2D” array. All arrays that you pass from the host to the device will become 1D arrays anyway. Even cudaMallocPitch() and cudaMalloc3D() only allocate linear memory. The difference is that cudaMallocPitch() and cudaMalloc3D() give the linear memory a pitch that ensures that the start address of each dimension will conform to the alignment requirements of that datatype.

If you want a simple front end for allocating 3D memory space on the device and accessing that space using 1/2/3 indices then I can give you the code for my ‘cudaMatrix’ class template.

Is there no way to create a 2d array that can be accessed in the normal way?
I’m trying to basically recreate an excel chart to read data in to make it simple to visualize the program.

I have it running but I’m getting wrong answers which would be in line with what you’ve said.

Is your example going to work for trying to access data spots such as Isomer1_arr[2][4]?

Yes, the example would work for that. You can access Isomer1_arr[2][4] through a 1d array.

// First you need a pointer to the very first element in your array

double* Isomer1_arr_1D;

Isomer1_arr_1D = &Isomer1_arr[0][0];

// Now you can access the data with a 1D array. 

result = Isomer1_arr_1D[4+100*2];

// This should be equivalent to:

result = Isomer1_arr[2][4];

I forgot that you were using C arrays, which have a different indexing scheme than Fortran arrays. I use the Fortran scheme because I do a lot of data transfer between C and Fortran code, and multi-dimensional CUDA allocations use the Fortran method of indexing, not the C method.

In C or C++ a multidimensional array is all continuous in memory space anyway. An array such as Isomer1_arr[10][100] is an array of pointers that points to the first element in each 100 element row. There is no gap between the rows unless you have some weird memory pitch.

See this link for a little more insight: Tutorial: Using C/C++ and Fortran together

Awesome thank you so much!
I will work on changing my code over to fit this.

You certainly can use 2D arrays in the kernel, and index them in the “normal” manner. However, this entails making two memory accesses. In contrast, if you flatten the 2D array to a 1D array, you can do it in 1 memory access.

Yup.

It seems like a lot of people don’t quite understand how multidimensional arrays in C work.

For a 3dimensional array “double array[nx][ny][nz];” you actually have 3 different sets of arrays containing different data types.

The first array, ‘array’ is an array of ‘double**’ pointers. You can visualize this as the address of the first element of a yz plane located at a specific x.

The second array ‘array[i]’ is an array of ‘double*’ pointers. The elements of this array contain the address of the first element of a z-row in a yz-plane at x=i;

The final array, ‘array[i][j][k]’ is the data located at index k of the array starting at ‘array[i][j]’.

For most applications all of the data arrays ‘array[][]’ are all aligned linearly in memory. However, there may be cases in which they are not. Such as the following:

int nx;

int ny;

// column1,... columnnx are double*'s

double column1[ny];

double column2[ny];

double column3[ny];

double column4[ny];

....

double columnnx[ny];

// array2d is a double**

double* array2d[nx];

array2d[0] = column1;

array2d[1] = column2;

....

array2d[nx-1] = columnnx;