Archives Discussions

[settle]

In the OpenCL Static C++ Kernel Language Extension, Rev. 04, Sec. 4.1 (page 10) the example class code on the host and kernel are exactly the same.  However, shouldn't the host side use cl_int to ensure matching data types with the kernel side?  If so, every class needs two code versions: host code (using OpenCL API) and kernel code (using OpenCL Langauge).  Is this correct?

Kernel code:

Class Test {

   setX (int value); private: int x; 

}

__kernel foo (__global Test* InClass, ...) {

If (get_global_id(0) == 0) InClass->setX(5);

}

Host Code:

Class Test {

   setX (int value); private: int x; 

}

MyFunc () {

tempClass = new(Test);
... // Some OpenCL startup code – create context, queue, etc. cl_mem classObj = clCreateBuffer(context, CL_USE_HOST_PTR,

sizeof(Test),&tempClass, event); clEnqueueMapBuffer(...,classObj,...);

tempClass.setX(10);
clEnqueueUnmapBuffer(...,classObj,...); //class is passed to the Device clEnqueueNDRange(..., fooKernel, ...); clEnqueueMapBuffer(...,classObj,...); //class is passed back to the Host

}

[settle]

bump.

This post over a month old and hasn't received any replies, but I'm still hoping to get some clarification on the correct usage of the OpenCL Static C++ Kernel Language Extension in the above example.  If someone at AMD could comment I'd greatly appreciate it.  Thanks!

[LeeHowes]

Yes, I think you're right. However, I'm usually too lazy to use the cl_x types. Additionally, I'm not 100% sure what the alignment guarantees are for structs defined in host code matching with structs in OpenCL "C" code. I'm not sure that using cl_x is really enough, you have to follow more general alignment rules.

[brg · Accepted solution]

Lee is correct that just using cl_XXX types is not enough. The problem is that Visual Studio will not guaranteeing alignments beyond 64-bit. and so the following would not be correctly aligned:

struct Foo

{

   cl_int x;

   cl_float4 y;

};

You need to add the explicit padding to match what OpenCL is expecting. So

struct Foo

{

  cl_int x;

   cl_int;

   cl_int;

   cl_int;

  cl_float4 y;

};

Of course, this is still valid on the host. While it is true that you could argue that you needed to versions, i.e. one for the host and one for the device, you can safely just use the cl_XXX version in both cases. The only different you will find in practice is with respect to:

  cl_long

and

  long

which are 8 and 4 bytes, respectively.

Alternatively, you can just define you own header that abstracts the typedefs and hides the choice.