cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

laobrasuca
Journeyman III

Build program crash, also crashes on kernel analyser

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)?

0 Likes
1 Solution
nou
Exemplar

kernel arguments can't be a pointer to pointer. also struct members can't be a OpenCL objects like cl_mem.

View solution in original post

0 Likes
12 Replies
nou
Exemplar

kernel arguments can't be a pointer to pointer. also struct members can't be a OpenCL objects like cl_mem.

0 Likes

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?

0 Likes

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;

}

0 Likes

laobrasuca,

cl_mem is not a valid data type in the OpenCL C language, so it cannot be used in the kernel.

0 Likes

section 6.9 of OpenCL specification a) and p)

0 Likes

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.

0 Likes

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.

0 Likes

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.

laobrasuca
Journeyman III

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).

0 Likes

Both are invalid and therefore the behavior is undefined, so the crash can be random based on invalid data.

0 Likes

ok, then. Thank you for your time

0 Likes
laobrasuca
Journeyman III

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.

0 Likes