2d array (_device_ or _constant_)

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,
In my program I have 2d arrays, like this
float ** host_2d;
and I need to use it on GPU,
My fist questions is How can i do it? Example will be best answer)
Second Is it possible to use constant memory for 2d Array?

2

This kind of 2D array representation (using a vector of row or column vectors) is usually a poor match for the GPU (and, I would claim, for CPUs as well where high performance code is concerned).

Since the individual row or column vectors are not guaranteed contiguous, they cannot be copied between host and device with a single 2D copy operation.

For the device, it is best to map a multi-dimensional array to one-dimensional array which provide contious storage, using either row-major or column-major storage arrangements as needed. For a 2D array, access might look as follows:

#define M ...  // #rows
#define N ...  // #cols

T *arr_d;
cudaMalloc ((void*)&arr_d, M*N*sizeof(T));
kernel<<<blks,thrds>>>(arr_d, ...);

#define arr(row,col) arr[(col)*M+(row)]  // column-major 
__global__ void kernel (T *arr, ...)
{
...
arr (i,j) = ...
... = arr (i,j) 
...
}

If you want to bind a 2D array to a texture