OpenCL Backend: Argument Mapping

[m4rs-mt]

OpenCL Backend: Argument Mapping

[m4rs-mt]

Unfortunately, it is not possible to pass (global) pointers to OpenCL buffers within structures, since they require a user-defined driver-specific argument mapping (interested readers should refer to this post). In other words, a C# structure

public struct Foo
{
    public ArrayView<int> Data;
    public long SomeField;
}

cannot simply be passed to a kernel by creating a marshalling OpenCL structure (and a corresponding host structure in managed code) View_Int:

struct View_Int
{
    __global int *ptr;
    int length;
}

struct Foo
{
    struct View Data;
    long SomeField;
}

One solution to this problem is to extract all nested ArrayView<> values from the arguments of a kernel. We should extract all source view objects (i.e. CLBuffer instances) and map their values separately with distinct calls to clSetKernelArg. After bringing the buffer pointers into the OpenCL world, we can reconstruct the actual view structures to simulate the intended behavior. Consider the following ILGPU kernel program:

void Kernel(Index index, ArrayView<int> source, Foo foo)
{
     // ...
}

As previously mentioned, this kernel cannot be represented in OpenCL via:

__kernel void Kernel(struct View_Int source, struct Foo foo)
{
    // ...
}

However, to avoid the problems presented, we can output the following OpenCL code:

__kernel void Kernel(
    __global int *ptr1,         // First nested view pointer
    __global int *ptr2,         // Second nested view pointer
    struct View_Int source, // Source structure without pointer information from host
    struct Foo foo)             // Source structure without pointer information from host
{
    // Setup view pointers in kernel code
    source.ptr = ptr1;
    foo.Data.ptr = ptr2;
    // ...
}

[m4rs-mt]

Unfortunately, the proposed solution does not work properly, since OpenCL does not allow nested pointers within structures to be passed to the kernel. However, some drivers accept such code, while others simply fail. Therefore, passing arguments requires an interop structure to be passed to the kernel, which stores the required offset and length information for each view. Then, each kernel argument is converted into its internal representation, which can contain (and work on) pointers to device-memory addresses.