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.
Solved! Go to 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.
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.
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;
}
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.
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.