Support for Local Params

  • Jeff Heaton
    Jeff Heaton

    First of all, great job on CLOO.  We've been using it on the Encog Neural Network project and CLOO is working out great for us.

    I do have one question.  Do you have support for "local" parameters, for example:

    __kernel void fft1D_1024 (__global float2 *in, __global float2 *out,
                              __local float *sMemx, __local float *sMemy) {

    How would I add the last two parameters.  No matter what I do I end up with an InvalidArgument.  I know at the C level in OpenCL you have to specify a NULL memory location for a local param, but I can't see where that would ever happen in CLOO.

  • nythrix

    that's nice to hear!

    You can use the lowest level method: ComputeKernel.SetArgument(argIndex, sizeOfArg, IntPtr.Zero).
    It offers the exact same functionality as the native clSetKernelArg(…) with the addition of error checking.

  • Jeff Heaton
    Jeff Heaton

    Worked great! Thanks!

  • Hi!
    I'm taking my first steps of OpenCL with CLOO and I guess that I've run into the same problem as OP allthough the solution given didn't do it for me :/

    I have a kernel to part an image into different segment,

    kernel void MyKernel(constant uchar* origSegments, local uchar* varSegments, private uint pixelCount, global uchar* src_img, global uchar* dest_img, local uint* totalColor, local uint* totalCount, local uint* pxlAff)

    So far I've learned how to pass privates (SetValueArgument<>) and constants/globals (by ComputeBuffers and SetMemoryArgument()), but when I try to pass locals I've run into the wall.

    From what I understand I should be able to use either the solution given here (I actually tried SetArgument() - even before I saw this solution) or just use SetLocalArgument() which by a look at the sourcecode seems to just pass parameters to a call to SetArgument, but neither of the works :(


    MyKernel.SetArgument(1, new IntPtr(varSegments.Length * sizeof(cl_uchar4)), IntPtr.Zero); 
    results in 
    Cloo.MemoryObjectAllocationFailureComputeException: OpenCL error code detected: MemoryObjectAllocationFailure when i call Execute() from my ComputeCommandQueue.

    True is, however, that I don't know if it is this particular local or another that screws up. But since I do the same kind of pass with the other locals (but with different indices and sizes of course) it feels like it should either be none or all of them.

    Same error goes for when I try

    MyKernel.SetLocalArgument(1, varSegments.Length * sizeof(cl_uchar4));

    Now I'm not sure, but I think I should also be able to send data into the local parameter/argument of the kernel, right? So I also tried

    MyKernel.SetArgument(1, new IntPtr(varSegments.Length * sizeof(cl_uchar4)), varSegmentsHandle.AddrOfPinnedObject());
    Where varSegmentsHandle is of type GCHandle. But this gave me another error, namely
    Cloo.InvalidArgumentValueComputeException: OpenCL error code detected: InvalidArgumentValue.

    Well. What I'd really want to know is how to properly pass data to local arguments of kernels, or how to pass the sizes of them at least. It could also help with some clearing out what kind of wrongdoings I've been up to :)

  • Oh big thanks!
    The thread kind of does explain why the errors appear. But then, if I get the picture right, if I somehow want to work with an array in local memory (since it's faster) how would I be able to declare it with the correct size? If I understand correctly I can't declare and pass an array in local memory space as an argument to a kernel and because of the inability to declare an array inside a opencl function without a constant value I can't really declare arrays which size I don't know at compile time, right?

    At the moment I'm thinking of solutions like having a row in the top of my kernel source where I have some #define FOO, where I can change my source string value so the the #define matches my array size and then compile the ComputeProgram afterwards… or something. What is the usual way of doing this?


  • Anonymous

    Just take a look at ComputeKernel.SetLocalArgument(int index, long dataSize). WIth this function u could set the size of the local buffer in bytes. I think u use local memory as a kind of a cache. so the first step would be filling the cache and than work with it. As douglas125 said u can't pass any data to local memory. it's a kind of private memory, but shared in a workgroup and not accessible from host. If u declare an __local short * array u can use it as short array but need to pass SHORT_SIZE*ArraySize to SetLocalArgument(…)

    If u could determine the array_size at compile time, the #define would be ok and sometimes to only solution, because openCL doesn't supports dynamic memory. If the size is only avaliable at runtime and u want to avoid recompiling, u should use buffers.