CODEWITHSUNDEEP
structopencl
Batko
Batko

Reputation: 33

OpenCL: Struct within struct sending to internal function on device side

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

Answers (1)

Mats Petersson
Mats Petersson

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

Related Questions