cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

timchist
Elite

Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

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;

        }

    }

}

Tags (2)
0 Likes
16 Replies
dipak
Staff
Staff

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

Thanks for reporting this and providing the reproducible test case. We'll check and get back to you.

0 Likes
timchist
Elite

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

Hi dipak.

Did you manage to reproduce the issue? Is there any news on the problem?

Thanks, Tim

0 Likes
dipak
Staff
Staff

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

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.

Regards,

0 Likes
timchist
Elite

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

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

0 Likes
timchist
Elite

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

Merry Christmas, Dipak.

Any progress on the issue?

0 Likes
dipak
Staff
Staff

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

Merry Christmas.

I've filed an internal bug report against the issue. If I get any update, I'll share with you.

Regards,

0 Likes
dipak
Staff
Staff

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

Hi Tim,

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.

Regards,

0 Likes
timchist
Elite

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

Thanks Dipak.

Hopefully other problems recently reported by me will be fixed soon too.

0 Likes
timchist
Elite

Re: Stack overflow in clBuildProgram on Catalyst 14.12 with R9 290

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).

0 Likes