cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

viscocoa
Adept I

local memory conflict?

Just experienced a bizarre problem. I can't past the entire code to confuse you. The program is like this conceptually:

struct st_a{

    ...

    int3 vInA;

    ...

    int num;

};

struct st_b{

    ...

    float3 vInB;

    ...

}

__kernel void MyKernel()

{

    __local struct st_a A;

    __local struct st_b B;

    A.vInA = (int3)0;

    B.vInB = (float3).6f;

    // OMG, A.vInA is now very big numbers!!!

    ...

}

I only launched one work item, so there is no write conflict.

Does anybody have any suggestions?

Thank you in advance!

VisCocoa

0 Likes
1 Solution

viscocoa,

I've found the problem. The reason this is occurring is that the code in question was written a long time ago, before OpenCL 1.1 and vec3 existed. In OpenCL, vec3 types actually store the same amount of memory as a vec4 type and this section of code wasn't taking that into account, so the offsets into the structure were based on the wrong size. The only work-around that is guaranteed to work is to use vec4 types instead of vec3 types. This will provide the correct size calculation and will fix all the issues.

Expect this fix to make it into catalyst in the April or May driver.

View solution in original post

0 Likes
18 Replies