2D thread blocks accessing a 3D array

Hey everybody, this is the kernel I am working with:

global void seedKernel(float * target, float * source)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;

target[x+1][y][1] = source[x+1][y+1][0];
target[x+1][y][6] = source[x+1][y+1][1];
target[x+1][y][11] = source[x+1][y+1][2];
target[x+1][y][16] = source[x+1][y+1][3];

target[x+2][y+1][2] = source[x+1][y+1][0];
target[x+2][y+1][7] = source[x+1][y+1][1];
target[x+2][y+1][12] = source[x+1][y+1][2];
target[x+2][y+1][17] = source[x+1][y+1][3];

target[x+1][y+2][3] = source[x+1][y+1][0];
target[x+1][y+2][8] = source[x+1][y+1][1];
target[x+1][y+2][13] = source[x+1][y+1][2];
target[x+1][y+2][18] = source[x+1][y+1][3];

target[x][y+1][4] = source[x+1][y+1][0];
target[x][y+1][9] = source[x+1][y+1][1];
target[x][y+1][14] = source[x+1][y+1][2];
target[x][y+1][19] = source[x+1][y+1][3];

}

I am using a 2D array of threads (22 X 22) to access a 3 dimensional array. I got the errors below. I know sort of what I am doing wrong, but not how to fix it…

kernel.cu(9): error: expression must have pointer-to-object type

kernel.cu(9): error: expression must have pointer-to-object type

kernel.cu(10): error: expression must have pointer-to-object type

kernel.cu(10): error: expression must have pointer-to-object type

kernel.cu(11): error: expression must have pointer-to-object type

kernel.cu(11): error: expression must have pointer-to-object type

kernel.cu(12): error: expression must have pointer-to-object type

kernel.cu(12): error: expression must have pointer-to-object type

kernel.cu(14): error: expression must have pointer-to-object type

kernel.cu(14): error: expression must have pointer-to-object type

kernel.cu(15): error: expression must have pointer-to-object type

kernel.cu(15): error: expression must have pointer-to-object type

kernel.cu(16): error: expression must have pointer-to-object type

kernel.cu(16): error: expression must have pointer-to-object type

kernel.cu(17): error: expression must have pointer-to-object type

kernel.cu(17): error: expression must have pointer-to-object type

kernel.cu(19): error: expression must have pointer-to-object type

kernel.cu(19): error: expression must have pointer-to-object type

kernel.cu(20): error: expression must have pointer-to-object type

kernel.cu(20): error: expression must have pointer-to-object type

kernel.cu(21): error: expression must have pointer-to-object type

kernel.cu(21): error: expression must have pointer-to-object type

kernel.cu(22): error: expression must have pointer-to-object type

kernel.cu(22): error: expression must have pointer-to-object type

kernel.cu(24): error: expression must have pointer-to-object type

kernel.cu(24): error: expression must have pointer-to-object type

kernel.cu(25): error: expression must have pointer-to-object type

kernel.cu(25): error: expression must have pointer-to-object type

kernel.cu(26): error: expression must have pointer-to-object type

kernel.cu(26): error: expression must have pointer-to-object type

kernel.cu(27): error: expression must have pointer-to-object type

kernel.cu(27): error: expression must have pointer-to-object type

Thank you for your time!

>>>Chris

You’re tring to access a flat array (source,target) like it was a “real” 3D array.

you should use only one set of to access source and target.

eyal

I am sorry, but how exactly would I do that?

By computing indices… something like
target[i][j][k] <=> target[k+(jnk)+(inj*nk)]

Thank you for your help.

Chris: Can you tell me how exactly you solved this problem?

You have to manually calculate element index in your array. To do so you need to multiply each index by the “step size”:

target[i][j][k] <=> target[k+(j*(max_index_of_k+1))+(i*(max_index_of_k+1)*(max_index_of_j+1))]

Think about it in 2D first. You’ve gotta a table of size N X M, how do you calculate absolute index of the element (x,y) in it, considering element at 0th colume 0th row has absolute index of “0”, while element at 0th column 1st row has index N.