OpenCL Backend: Argument Mapping

[m4rs-mt]

The new CLArgumentMapper maps managed values from the .Net world to the OpenCL driver world. Since we have to map all intrinsic view types (which are an abstract concept) to actual kernel types, the CLTypeGenerator class should do the job for us. However, it is not currently known whether OpenCL drivers accept structures with nested pointer declarations (they will not be passed as pointers but as constants - which should be acceptable...). We need to investigate this issue and change the classes CLArgumentMapper/CLTypeGenerator accordingly (if necessary).

[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.