Cannot use struct inside a kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

hello, i’m having some odd problems using struct inside a kernel. I paste my little code here:

typedef struct {

    float **elementi;

} matrix_t;

__kernel void moltiplicazione(

							  __global float* dim_block,

							  __global float* num_block,

							  __global matrix_t* blocchiA,

							  __global matrix_t* blocchiB,

							  __global matrix_t* blocchiC,

							  __global matrix_t* swap

							  )

{

	int iGID = get_global_id(0);

	if (iGID > num_block[0]) return;

        blocchiC[iGID].elementi[0][0]=1.0f;

}

really simple (and useless in this form), but everytime i assign some value to blocchiC.elementi thread stops and i got error from the host’s clEnqueueReadBuffer function.

If i try blocchiC=blocchiA it works fine and the output is correct.

What i can do to manipulate every single element of blocchiA, B and C?

Last thing, if i do something like “float foo = blocchiC[iGID].elementi[0][0];” i got no error but i didn’t verified the value of that foo.

I read the result from host with this:

clEnqueueReadBuffer(queue, d_matC, CL_TRUE, 0, sizeof(matrix_t)*n_blocchi, result, 0, NULL, NULL);

so i’m reading “blocchiC” from the kernel, but even if only do some write operation in blocchiA or blocchiB i got error.

Thank you all :)

2

I don’t think it’s a struct issue, I think it’s more that OpenCL doesn’t support the use of double pointers (2D arrays). You can use 2D images, but I know for sure you can’t use a 2D buffer, or a struct that contains a 2D buffer.

3

thank you, i finally got it working with a single pointer :D
the odd thing is that it works only with driver 256, with other version i got some assembler error when building the program! i think i’m back to cuda :D

4

It’s a memory space problem.

float * * a;

declares a pointer in private memory space to a pointer in private memory pointing to a private float (as this is the default). Should be something like

__global float  * * a;

as the float itself is in global space.