Bump.

Bump.

Do you have an example of such a kernel?

Here are two kernels that fail:
__kernel void uint_div(__global const uint *a,  __global const uint *b, __global uint *c){
    c[0]=a[0]/b[0];
}
__kernel void uint_add16(__global const uint *a,  __global const uint *b, __global uint *c){
    for(uint i=0;i<16;i++) c[i]=a[i]+b[i];
}

This one works:
__kernel void uint_add(__global const uint *a,  __global const uint *b, __global uint *c){
    c[0]=a[0]+b[0];
}

Sadly I don't have access to the hardware anymore (it's probably in a dump somewhere).

Created attachment 121742 [details]
dmesg after running the uint_add16() kernel

(In reply to Hristo Venev from comment #4)
> Here are two kernels that fail:
> __kernel void uint_div(__global const uint *a,  __global const uint *b,
> __global uint *c){
>     c[0]=a[0]/b[0];
> }
> __kernel void uint_add16(__global const uint *a,  __global const uint *b,
> __global uint *c){
>     for(uint i=0;i<16;i++) c[i]=a[i]+b[i];
> }

I actually can confirm that the 2nd kernel does cause a GPU stall in the radeon driver (r600/Barts). It causes no problem with amdgpu (radeonsi/Tonga).

I am using kernel 4.5.0-rc3, current llvm 3.8 branch and current mesa git.

As I am not overly good with OpenCL, is this kernel somehow problematic?

I think this is probably a duplicate of bug 60879 and that it was basically luck that "small kernels" didn't hang on Hristo's card.

(In reply to EoD from comment #5)
> I actually can confirm that the 2nd kernel does cause a GPU stall in the
> radeon driver (r600/Barts). It causes no problem with amdgpu
> (radeonsi/Tonga).
> 
> I am using kernel 4.5.0-rc3, current llvm 3.8 branch and current mesa git.
> 
> As I am not overly good with OpenCL, is this kernel somehow problematic?

EoD, if these stalls on r600 are still reproducible, please open a separate bug.

*** This bug has been marked as a duplicate of bug 60879 ***