OpenCL pass local memory

Question

I've got the following OpenCL function:

__kernel void mergeSort(__global const float * in,__global float * out, __local float *aux)

When I try to call it, I only receive an internal error

OpenCLFunction::invprop: OpenCLLink encountered an invalid property. 

I've allocated the memory with

OpenCLMemoryAllocate["Float", 8]

When I remove the "__local", the kernel can be executed. The kernel works fine in a standalone C++ application.

Thx

Answer 1

The OpenCL specification has this to say about clSetKernelArg which is called internally when the arguments are passed in:

If the argument is declared with the __local qualifier, the arg_value entry must be NULL.

When you allocate using OpenCLMemoryAllocate, you are setting that pointer to something non-null, and you are allocating global memory.

Do not use OpenCLMemoryAllocate for this. Instead, remove the aux argument and handle this inside the kernel:

__kernel void mergeSort(__global const float * in,__global float * out)
{
    __local float aux[8];    
    ...
}

I have dug into the Mathematica packages and found this in SystemFiles/Links/GPUTools/WGLPrivate.m

SetKernelLocalMemoryArgument[size_Integer] :=
    Module[{res},
        res = cSetLocalMemoryArgument[size];
        If[SuccessQ[res],
            res,
            Throw[$Failed]
        ]
    ]

... which appears to do what we want. However, I see no way to call this hidden function.

---

The Mathematica documentation also says:

Use of memory modifiers such as __constant is not supported by OpenCLLink. Memory passed into an OpenCLFunction must be __global.