When I try to compile certain OpenCL code on a machine with R9 290 installed running a Catalyst 14.12 driver, the program crashes with stack overflow in amdocl64.dll when calling clBuildProgram.
The same code compiles just fine on a machine running HD 7970 with the same driver.
Below is the minimum kernel causing this to happen:
__kernel void test(__global float* srcBase, __global float* dstBase, int width, int height)
int x = get_global_id(0);
if(x < width)
int y = get_global_id(1);
if(y < height)
__global float* dst = dstBase + x + y * width;
__global const float* head = srcBase + y * width;
int offsetLeft = x > 0? -1: 0;
float leftFactor = x > 0? 1: 0;
float curAverage = leftFactor * (head[offsetLeft]);
*dst = curAverage;
dst += width;
Yes, I'm able to reproduce the issue using a sample test project. However, when I'm trying to compile the kernel code using CodeXL, its working fine. Another interesting point is the test code is working fine for following cases:
1) Disable the optimization by passing "-O0" or "-cl-opt-disable" during clBuildProgram()
2) Building as OpenCL 2.0 kernel code by passing "-cl-std=
CL2.0" during clBuildProgram() [with/without optimization]
I guess its a compiler bug. I've asked someone for clarification. If needed, I'll file an bug report against it. Meanwhile, you may try the above workarounds and let me know your findings.
Compilation works with either "-O0" or "-cl-std=CL2.0" or both, however, in these cases application still fails with stack overflow in amdocl64.dll when attempting to call clBuildProgram after clCreateProgramWithBinary with the binary produced after compilation. As I have explained in another post, we are using offline compilation and do not ship our apps with OpenCL sources.
Also, when "-fno-bin-llvmir" flag is passed, compilation results in this error:
An error with the ELF object was encountered.
Error while BRIG Codegen phase: compilation error
Good news. Recently I got an update that this issue has been fixed in latest internal driver build. Hope the fix will be available to public version soon.
We have just found out that the same kernel causes compilation to crash with Catalyst 14.9 driver too when GPU_FORCE_64BIT_PTR is set to 1 on both machines (with R9 290 and HD 7970).