How to pass multidimensional arrays as function arguments

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I’m trying to pass a multidimensional array as function argument. But I am unable to find the correct way. The following code is not working:

#define MAX_NO_Q 4

__device__ void kernel(float* y, int q)

{

	float A_b[MAX_NO_Q][MAX_NO_Q + 1];

	calculate(A_b, q);

}

__device__ void calculate(float** A, int q)

{

}

I’m getting the error: argument of type “float (*)[5]” is incompatible with parameter of type “float **”

Hopefully somebody can tell me what the correct way is. Thanks.

Kind regards,

Nick

2

In fact, float** differs from float[5]. In order do pass static arrays as arguments, you need to specify at least the number of rows of your array, like this

#define MAX_NO_Q 4

__device__ void calculate(float A[][MAX_NO_Q + 1], int q)

{

}

__device__ void kernel(float* y, int q)

{

        float A_b[MAX_NO_Q][MAX_NO_Q + 1];

calculate(A_b, q);

}

But remember that the size of array must be determined in compiling time or it won’t work.

Hope this can help

3

Your answer solved my problem. Thanks.

4

This is not necessarily true.

If you are using global memory (the compiler will automatically assume global memory if you don’t specify cause it won’t know otherwise), then you can do it like you were before. You will need to specify the width and height of the 2D array, however. I wrote a memory class to help with allocating 2D arrays in CUDA, and can pass those 2D arrays to kernels. If you want to check it out it’s here on sourceforge.

Example…

// Kernel that executes on the CUDA device

__global__ void square_array2d(int** a, int width, int height)

{

    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    int idy = blockIdx.y * blockDim.y + threadIdx.y;

if (idx < width && idy < height) {

        a[idx][idy] = a[idx][idy] * a[idx][idy];

    }

}

extern "C" void square2d(int** a, int width, int height)

{

    dim3 threads(width, height);

    square_array2d<<< 1, threads >>>(a, width, height);

}

Using the 2D memory…

// allocate the memory on the device

    DeviceMemory2D<float> device2d_1;

    device2d_1.allocate(10, 10);

// call the kernel that squares each element of a 2d array

    square2d(device2d_1.ptr(), device2d_1.width(), device2d_1.height());
5

Pardon my ignorance, but how does any of that address the original question, which was about passing a statically declared, 2D shared memory array to a device function inside a kernel?

6

The OP didn’t say anything about static or shared memory. The only hint he really gave is from the error he posted. Giving information about how it can be done in a general sense, and clarifying the passing of float** to a kernel (since his original kernel had that) doesn’t seem too far from topic to me. Not trying to start anything, only justifying.