cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

Deluxe
Adept I

Possible OpenCL kernel (CPU) compiler bug

I have following (quite simple) function my kernel code. Problem is that this function compiles without any problems on GPU (RV770), but on CPU (AMD FX-8350) call to clBuildProgram causes "access violation exception" in amdocl.dll.

float4 RandNormalFloat4(float4 mi, float2 sigma, mwc64xvec4_state_t *s)

{

    float4 out;

    float4 randSeed = RandFloat4((float4)(0.0f), (float4)(1.0f), s);

    // TODO: Try to use "native_" versions

    out = (sigma * sqrt((float2)(-2.0f) * log((float2)(1.0f) - randSeed.xz))).xxyy;

    out.xz = out.xz * sincos(2.0f*M_PI_F * randSeed.yw, &out.yw);   

    return out + mi;

}

Problem seems to be call of "sincos" function, if I replace this line with following separate calls of "sin" and "cos" kernel compiles without problems on both GPU and CPU devices.

out.xz = out.xz * sin(2.0f * M_PI_F * randSeed.yw);

out.yw = out.xz * cos(2.0f * M_PI_F * randSeed.yw);

I'm using Catalyst 13.1 legacy driver (OpenCL 1.0 AMD-APP 937.2) and AMD-APP 2.8 (tried 2.6 RC3 and 2.7 RC too). Actually AMD-APP 2.7 and AMD APP CPU Runtime from AMD-APP 2.8.

"clinfo.exe"output is in attachment.

0 Likes
1 Solution

Hi Deluxe,

As per OpenCL Spec section 6.1.7, page 208, it is illegal to take address of a vector element. This should have actually resulted in a compilation error, but it crashed. The crash will be fixed in future releases.

View solution in original post

0 Likes
4 Replies
himanshu_gautam
Grandmaster

Can you please attach a small kernel code, so that i can verify, the cause of the issue. The above function has many dependencies (mwc64xvec struct, Randfloat4 etc). I tried commenting out some of the code, and made the above function to compile for both GPU & CPU.

0 Likes

Thanks for reply. Here is minimal code, which still crashes (for me) same way as I mentioned in first post.

float4 RandNormalFloat4(float4 mi)
{
    float4 out = (float4)(1.0f);
   
    out.xz = out.xz * sincos(2.0f, &out.yw);
   
    return out + mi;
}

__kernel void Test(__global float4 *out)
{
    size_t id = get_global_id(0);
   
    float4 random = RandNormalFloat4((float4)(1.0f));
    out[id] = random;
}

0 Likes

Thanks for the testcase. I could reproduce the issue.

I have forwarded it to relevent AMD Engg Team. It will be fixed in upcoming releases.

I guess this bug does not stop you severely as sine and cosine functions are anyways working fine.

0 Likes

Hi Deluxe,

As per OpenCL Spec section 6.1.7, page 208, it is illegal to take address of a vector element. This should have actually resulted in a compilation error, but it crashed. The crash will be fixed in future releases.

0 Likes