AnsweredAssumed Answered

Atomic operations not working on HD7970 Tahiti

Question asked by smistad on Oct 4, 2012
Latest reply on Oct 5, 2012 by smistad

The small kernel below which use an atomic_min operation fails on HD7970 (meaning the buffer C is not changed at all, can test with C = (0,1,2,3,4,5,6,7,8, ... 99) and global size 99). But if I unquote the printf(a); line, it works (meaning C is changed correctly)!? But using printf offcourse slows down the computation. I am currently using Ubuntu 12.04, AMD driver 12.9 beta and AMD APP 2.7

 

Note that this works on AMD 5870 and 5850, and NVIDIA cards.

 

Does anyone know a solution or workaround for this problem? Thanks.

 

__kernel void test(
    volatile __global int * C 
    ) { 
    int id = get_global_id(0);
    atomic_min(C+id*2+1, C[id*2]);
    //printf("a");
}

 

I also discovered that some compiler optimizations doesn't work on the HD7970 Tahiti device which results in the compiler seg faulting. Using the option "-cl-opt-disable" avoids this problem.

Outcomes