PDA

View Full Version : Assigning a number to a __global int *p and resolving it



poonaatsoc
07-03-2009, 10:02 PM
Thanks for taking time to read this. I have 2 kernels. To both these kernels I send the same 2 memory buffers as arguments, where each of the memory buffers have the size -> sizeof(int) * 10.

I call init_mpm( ) first and I set m[i] to the subsequent address locations in bmasks. Basically I am holding an array of addresses.

I next call resolve_address( ) and I assign the value m[i], which actually now holds the address values from bmasks from the previous call to init_mpm( ) to a pointer p. But now when I try to resolve the pointer as p[0], I get clBuildProgram error of -42. Can't I resolve the address like this? Isn't p now holding the address and can't it be resolved as an integer using p[0]? I am actually trying to hold an array of address from a memory object across multiple kernel invocations and later resolve these addresses in another kernel.

__kernel void init_mpm(__global int *m, __global int *bmasks)
{
unsigned int i = 0;

for (i = 0; i < 10; i++) {
bmasks[i] = i;
m[i] = bmasks + i;
}
}

__kernel void resolve_address(__global int *m, __global int *bmasks)
{
unsigned int i = 0;
__global int *p;

for (i = 0; i < 10; i++) {
p = m[i];
m[i] = p[0]; <-- this line gives me an error. If I comment this line I get no errors during compilation
}

}

bgaster
07-06-2009, 09:04 AM
There seems to be few issues with the code as presented. Firstly, in the first kernel you are assigning a value (bmasks + i) of type __global int * to a value (m[i]) of type int:

t1.cl(7): warning: a value of type "__global int *" cannot be assigned to an
entity of type "int"
m[i] = bmasks + i;
^
This can be fixed by assigning the correct type to the argument m, with something like:

typedef __global int * globalInt;

__kernel void init_mpm(__global globalInt *m, __global int *bmasks)
{
unsigned int i = 0;

for (i = 0; i < 10; i++) {
bmasks[i] = i;
m[i] = bmasks + i;
}
}


The bigger problem is with the 2nd kernel as it does not seem to be
valid. This is made clearer if we change it to also use the typedef
globalInt:

__kernel void resolve_address(__global globalInt *m, __global int *bmasks)
{
unsigned int i = 0;
__global int *p;

for (i = 0; i < 10; i++) {
p = (globalInt)m[i];
m[i] = p[0];
}
}

Compiling this we get:

t1.cl(20): warning: a value of type "int" cannot be assigned to an entity of
type "globalInt"
m[i] = p[0];

The problem is that you are assigning m[i] (__global * int) to p (also
a __global * int) and then on the following line dereferencing p[0]
and assigning back to the the same m[i] on the previous line. Of
course, this is a type error and while you could probalby get it to
statically compile it is not clear this is the expected behaviour. On the other hand it is completly valid to do something like:

int foo = p[0];

A useful OpenCL API call is clGetProgramBuildInfo in conjunction with
the flag CL_PROGRAM_BUILD_LOG that can be used to obtain the log for a
particular call to clBuildProgram.