Get_global_id(0) and other - how

  • laterafter

    laterafter - 2011-05-18

    I need help with … hmm very basic stuff…

    I wrote kernel but I have some problems running it. So I just came up with an idea to check some basic values in my kernel. I want to check get_global_id(0), get_local_id(0) and so one. All are of type size_t (unsigned 32/64 bit values).

    I have such code:
    ComputeBuffer<uint> iBuffer = new ComputeBuffer<uint>(_context, ComputeMemoryFlags.WriteOnly | ComputeMemoryFlags.AllocateHostPointer, 3);
    _kernel.SetMemoryArgument(9, iBuffer);
    uint iarr = new uint;
    GCHandle iHandle = GCHandle.Alloc(iarr, GCHandleType.Pinned);
    queue.Read(iBuffer, true, 0, 3, iHandle.AddrOfPinnedObject(), null);

    in kernel code I put into iarr values like get_global_id(0) or get_global_size(0);

    but after executing kernel (with some fixed values of workGroupSize and globalOffset=0) iarr _arry has very strange values (very big). I expect that in  one of _iarr _array cell I will have workGroupSize (_get_global_size(0)).

    Can you please help me with this problem? How can I get values like get_global_size() out of kernel code (to c# code) ?

  • laterafter

    laterafter - 2011-05-18

    ** Does this forum have a "edit post" option ???

  • nythrix

    nythrix - 2011-05-22

    1. There are several bugs with that code. Replace with this:

    ComputeBuffer<uint> iBuffer = new ComputeBuffer<uint>(_context, ComputeMemoryFlags.WriteOnly, 3);
    _kernel.SetMemoryArgument(9, iBuffer);
    //This is not necessary but much nicer:
    uint[] iarr = new uint[3];
    queue.ReadFromBuffer(iBuffer, ref iarr, true, null);

    2. You're assuming that get_global_id or get_global_size return something (size_t in OpenCL C) that can match the C# uint. Which might not be true. Your app and the drivers can be 32 or 64 bits. This can change the size of size_t. Check that one first by returning sizeof(size_t) (or whatever is used in OpenCL C) from inside your kernel.
    3. I'm not  sure if size_t is signed or not. Check that too.
    4. No, it isn't possible to edit posts here. I can delete my posts and reupload them, if they contain serious bugs. Otherwise I don't bother.

  • laterafter

    laterafter - 2011-05-23

    Thank you for the answer. Indeed your idea with ReadFromBuffer is really much more readable and elegant.
    So I ran a small kernel with parameter_ __global int *i which does only _i = sizeof(size_t).

    In iarr there is value '4'. So it's 4*8=32 bits. ( I was aware that size_t can be 32 or 64 bits long but now it's clear).
    I'wll try now to correct my previous code and read finally get_global_id(0).


  • nythrix

    nythrix - 2011-05-23

    No prob. By the way, avoid using size_t as a kernel argument type. Getting that right can be REALLY tricky. This is not a Cloo specific problem but a general issue arising from different combinations of OpenCL drivers, managed code and OS bitness.


Log in to post a comment.