AnsweredAssumed Answered

Implementing 64 bit atomics

Question asked by chersanya on Jul 22, 2012
Latest reply on Oct 29, 2012 by notzed

I know that no ATI cards and drivers support 64 bit atomics, but they are needed for my project. I've tried to implement them this way:

 

long _atomic_xchg(volatile __global int3 *p, long val)
{
    __global long *p_val = (__global long)p;
    __global int *p_lock = (__global int)p + 2;
    while(atomic_cmpxchg(p_lock, 0, 1));
    long old = *p_val;
    *p_val = val;
    atomic_xchg(p_lock, 0);
    *p_lock = 0;
    return old;
}

__kernel void build_hashtable(
    __global int3 *table)
{    
    _atomic_xchg(table, 1);
}

This contains _atomic_xchg for long (actually it gets a pointer to int3, where the first two ints are the value, and the third is used for locking) and simple kernel to test this. However, it hangs up and the driver restarts. It happens even for small global work sizes, like 64, and I don't see the idea how to fix it; atomic_xchg is the only function needed.

Outcomes