HSAFoundation / CLOC

CL Offline Compiler : Compile OpenCL kernels to HSAIL
Other
50 stars 18 forks source link

How to make pass-by-value struct in the kernel code work in CLOC 1.0.x? #17

Open PXAHyLee opened 8 years ago

PXAHyLee commented 8 years ago

Hi developer,

When I leverage the new CLOC toolchain, I am struggled to make a kernel that has a struct in the argument work.

Consider the following simple kernel code

typedef struct SimpleStruct
{
    int a;
    int b;
} SimpleStruct;

__kernel void struct_copy(SimpleStruct a, global SimpleStruct *b)
{
    int i = get_global_id(0);
    b[i] = a;
}

The host side may have the following kernel argument structure

typedef struct __attribute__((aligned(16))) {
    int a;
    int b;
} SimpleStruct;

typedef struct __attribute__((aligned(16))) _KernelArg {
#ifdef BRIG
    uint64_t global_offset_0;
    uint64_t global_offset_1;
    uint64_t global_offset_2;
    uint64_t printf_buffer;
    uint64_t vqueue_pointer;
    uint64_t aqlwrap_pointer;
#endif
    SimpleStruct *a;
    SimpleStruct *b;
} KernelArg;

In the older hlc 3.2, after doing some survey on the semantic of the hsail code, I put the pass-by-value struct in the global memory and successfully copy the content to b on the GPU. It works.

However, in the new toolchain, it seems that the way that work in the older toolchain now doesn't work. Maybe the assumption that I made in the older toolchain is changed. The content of b isn't correct. How do I set the content of a now? How do I make such cases work in the CLOC toolchain?

Any advice or hint would be appreciated. Thanks.