Reputation: 33
Got a question about struct handling in OpenCL that I didn't find on here. I've gathered the all of the data I use in a struct, which itself consists of several structs. I want to do the following:
typedef struct tag_OwnStruct
{
float a;
float b;
float c;
int d;
float e;
int f;
}OwnStruct;
typedef struct tag_DataStruct
{
OwnStruct g;
//+ Alot of other structs... not written for simplicity
}DataStruct;
void PrintOwnStruct(OwnStruct* g)
{
printf("Current lane id : %f\n",g->a);
}
__kernel void Test(__global DataStruct *data)
{
PrintOwnStruct(&data->g);
}
So I want, from the given data I get sent in from the host side to the device, to send the reference to a struct inside of it. That doesn't work somehow, and I don't know why. I've tried the same thing in plain C code and it worked..
If I change PrintOwnStruct to :
void PrintOwnStruct(OwnStruct g)
{
printf("Current lane id : %f\n",g.a);
}
and call the function as : PrintOwnStruct(data->g)
the code will run on the device side. Is there any other way to do this? Since I'm not sending the reference to the function, is it being passed by value? And shouldn't that be slower than passing function parameters by reference?
Upvotes: 0
Views: 494
Reputation: 129524
So the problem appears (from comments) to be the confusion between __private
and __global
address spaces, and possibly that the compiler/runtime is not very helpful in informing about the mix of pointers.
void PrintOwnStruct(OwnStruct* g)
{
printf("Current lane id : %f\n",g->a);
}
__kernel void Test(__global DataStruct *data)
{
PrintOwnStruct(&data->g);
}
The __global DataStruct *data
is a pointer to something in __global
address space [in other words having the same address for all CL threads], the argument to void PrintOwnStruct
OwnStruct* g)
declares an argument that is pointing to OwnStruct
in the default __private
address space [in other words on the stack of this thread].
The correct thing to do is to maintain the address space for both pointers to __global
by declaring the function PrintOwnStruct(__global OwnStruct* g)
.
I'm pretty sure SOME OpenCL compilers will give an error for this, but apparently not this one. I expect that true syntax errors, such as adding %-&6
to the code will in fact give you a kernel that doesn't run at all, so when you call clCreateKernel
or clBuildProgram
, you'll get an error - which can be displayed by clGetProgramBuildInfo
. But if the compiler isn't detecting different address spaces, then it's a bug/feature of the compiler.
[In fact, if your compiler is based on Clang, you may want to have a look at this bug: https://llvm.org/bugs/show_bug.cgi?id=19957 - half an hour of googling gives a result of some sort! :) ]
In the newer CL2.0 the default address-space is generic
, which allows "any" address space to be used.
Upvotes: 1