Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Open
PXAHyLee opened this issue May 18, 2016 · 0 comments
Open

Comments

@PXAHyLee
Copy link

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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

1 participant