OpenCL allocating struct on GPU results in garbage

I have defined the following struct in C++ host code :

struct __declspec(align(16)) MyNode {
cl_uchar mData;
cl_int3 mPos;
};


and in OpenCL :

struct __attribute__((aligned(16))) MyNode {
uchar mData;
int3 mPos;
};


now from host code Im calling :

MyNode node= {0};
node.mPos.x = 1;
node.mPos.y = 2;
node.mPos.z = 3;

cl_mem clnode_mem = clCreateBuffer(
mOpenCLctx,
CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
sizeof(MyNode),
&node,
&err);

clSetKernelArg(m_kernel_test, 0, sizeof(cl_mem), &clnode_mem));


then Im calling a test kernel on GPU which is defined as :

__kernel void test(__global MyNode* node)
{
printf("pos = %d,%d,%dn",
node->mPos.x,
node->mPos.y,
node->mPos.z);
}


but what I see as output is pos = 0,0,0, and if I remove the mData member from host and device struct definitions, then it is printing correctly .
what is going on here ?

pos = 0,0,0

mData

I'm building with VS2015 with x64 configuration and running OpenCL 1.2 on an NVIDIA GPU.

1 Answer
1

Seems likely to be an alignment problem. int3 probably has 16-byte alignment, whereas the C type cl_int3 may only have 4-byte alignment. So the kernel would be looking for it at offset 16 whereas the CPU provides it at offset 4.

int3

cl_int3

A good sanity check would be to print the sizeof(MyNode) on host and kernel. If those diverge, I suggest trying explicit amounts of padding or alignment annotations for the struct fields.

sizeof(MyNode)

By clicking "Post Your Answer", you acknowledge that you have read our updated terms of service, privacy policy and cookie policy, and that your continued use of the website is subject to these policies.