I'm trying to use a structure to pass a pointer to buffers on the host side, e.g., cl_mem* buffs, as a pointer to pointers in the kernel side. So far, I've tried to see if the kernel code will build ok. So I did the following kernel test code.
//ex ok:
typedef struct struct_buffs{
float **tab; //or __global float **tab;
}struct_buffs;
__kernel void test(struct_buffs buffs)
{
uint gid = get_global_id(0);
float* buff = buffs.tab[gid]; //or __global float* buff = buffs.tab[gid];
buff[0] = gid;
}
//
This code will build fine, no problem (although I haven't tried to run it yet). So, I could be happy, but I've tried the version below which will make the .exe crash when running clBuildProgram(). It's an access violation (test.exe: 0xC0000005: access violation writing location 0x00000008) (it will also make kernel analyser 1.10 not respond, I guess for the same reason).
// ex no ok:
typedef struct struct_buffs{
float **tab;
}struct_buffs;
__kernel void test(__global struct_buffs* buffs)
{
uint gid = get_global_id(0);
float* buff = buffs->tab[gid];
buff[0] = gid;
}
So I have 2 questions:
1- Is the second code supposed to crash??
2- In the case of the first code, will the struct_buffs buffs variable be in the global or private memory space?. I mean, will I be able to pass things from the host as I usually do when manipulating a simple buffer (cl_mem buff on the host side and __global float* buff on the kernel side)?
Solved! Go to Solution.
kernel arguments can't be a pointer to pointer. also struct members can't be a OpenCL objects like cl_mem.
kernel arguments can't be a pointer to pointer. also struct members can't be a OpenCL objects like cl_mem.
even if it does not accept pointer to pointer, it does not explain why it crashes. It should return a build error and report it in the log. Besides, if it crashes with the second kernel, why it does not with the first one, which does essentially the same thing?
nou, about the "struct members can't be a OpenCL objects like cl_mem", can you point me where is that in the spec?
and just to finish with this, I could not, indeed, make it work with structure containing cl_mem, so I had to pass all buffers individually as kernel arguments and use the buffer offset thing (http://devgurus.amd.com/message/1256237#1256237), something like:
__kernel void test(__global float* buff1, __global float* buff2, __global float* buff3)
{
uint gid = get_global_id(0);
ptrdiff_t buffs_offsets[] = {0, buff2-buff1, buff3-buff1};
__global float* buff = buff1 + buffs_offsets[gid];
buff[0] = gid;
}
laobrasuca,
cl_mem is not a valid data type in the OpenCL C language, so it cannot be used in the kernel.
section 6.9 of OpenCL specification a) and p)
Thanks for providing a test case, we will investigate the crash. As nou said, pointer to a pointer is illegal, however, it is only illegal as a kernel argument.
yes, I'm aware about the pointer to pointer thing as kernel arg, I remember an old discussion we had on this quite long ago. I suppose support for pointer to buffers still is something that was not proposed or accepted yet. I still believe that run a kernel only once for several buffers is better than run several times the kernel for a buffer at a time.
It was more like for curiosity that I posted this second kernel, as I suppose crashes should not exist at all.
laobrasuca,
The problem is that the frontend is giving a warning and not an error, and so invalid code is making it to the passes that assumes all code has been validated. A future release will make this an error and so it won't compile.
thx micah! Could u give an insight on what consists the difference between the first kernel code and the second one so that the first one passes and the second one doesn't? (note: if u comment the line with buff[0] = gid; the second kernel will also build without problem).
Both are invalid and therefore the behavior is undefined, so the crash can be random based on invalid data.
ok, then. Thank you for your time
thx for pointing in the spec. The statement p. does not exists in specs 1.0 and 1.1 thats why I could not find it.
it's well know that all the cl_types are not part of the OpenCL C language, but, since a cl_mem in host side is often (not to say always) interpreted as a __global some_type * in the kernel (some_type being (u)int, float or something alike, except some restrictions, namely section 6.9, item k. of spec 1.2), I though that some structure like
struct host_struct{
cl_mem buff1;
cl_mem buff2;
cl_mem buff3;
} struc_buffers;
could be seen in the kernel as
struct kernel_struct{
__global float* buff1;
__global float* buff2;
__global float* buff3;
} struc_buffers;
whose elements can be accessed by some index i like buff_i = *(&struc_buffers.buff1+i). If all this was allowed, it would require lot less kernel arguments.