CODEWITHSUNDEEP
c++opencl
stix
stix

Reputation: 1146

What is the proper way to handle OpenCL buffers of structs?

I have the following OpenCL Kernel:

//OPENCL KERNEL
struct MyStruct
{
  float A;
  float B;
  float C;
  float D;
  float E;
};

__kernel void kernelMain(struct MyStruct* outputBuffer)
{
   size_t idx = get_global_id(0);

   //Do some stuff here with the outputBuffer
   outputBuffer[idx].A = 42.0;
}

As you can see, it defines a custom type, called MyStruct.

On the host side, I have the same struct definition (copy-pasta'd):

//HOST SIDE
struct MyStruct
{
  float A;
  float B;
  float C;
  float D;
  float E;
};

And I am trying to create a buffer for the writing of the kernel data, again host side code:

//HOST SIDE
cl::Buffer outputBuffer(clContext, CL_MEM_READ_WRITE, (size_t)numElements * sizeof(MyStruct));

clKernel.setArg(0, outputBuffer);

The problem occurs when I call clKernel.setArg. It fails with an error code -51, which according to OpenCL documentation is a kernel invalid argument size error.

I've tried using openCL datatypes, re-writing the host definition of the struct to be:

struct MyStruct
{
  cl_float A;
  cl_float B;
  cl_float C;
  cl_float D;
  cl_float E;
};

But this also gives an error.

My question is this: What is the proper way to create an OpenCL buffer for working with custom structs?

Upvotes: 2

Views: 1254

Answers (2)

Kyle Lutz
Kyle Lutz

Reputation: 8036

If you're able to use C++, you could use the BOOST_COMPUTE_ADAPT_STRUCT() macro which takes care of wrapping a struct and making it available for use in OpenCL kernels.

After wrapping, you can create a OpenCL memory buffer for a collection of your structs with the boost::compute::vector<T> container class:

// adapt "MyStruct" for OpenCL
BOOST_COMPUTE_ADAPT_STRUCT(MyStruct, MyStruct, (A, B, C, D, E));

// create a OpenCL buffer with 100 "MyStruct" objects
boost::compute::vector<MyStruct> my_structs(100);

// use "my_structs" with an opencl kernel
my_kernel.set_arg(0, my_structs);

Upvotes: 1

maZZZu
maZZZu

Reputation: 3615

If adding __global keyword doesn't help it sounds like you might have a different struct padding between host and device.That would make the size of the struct or data positions within the struct different. Struct padding is platform and compiler dependent.

This is something you should consider when using same structs on host and device with OpenCL. Especially if you are targeting your software to multiple platforms.

One workaround is to use array of floats to move data between host and device.

Upvotes: 1

Related Questions