AnsweredAssumed Answered

Substitute for atomic_inc

Question asked by shunyo on May 23, 2013
Latest reply on May 26, 2013 by shunyo

Hi,

I have an ATI Firepro V4800 graphics card which does not support cl_khr_int64_base_atomics. I am trying to adapt the RadixSort algo given in the SDK Examples for long integers. The algo uses atomic_inc, the 64-bit of which is atom_inc, which I cannot use in the kernel. So, my question is, is there a piece of code which performs the same function as atomic_inc which can be used? The piece of kernel code is given below:

 

__kernel void histogram(__global uint* unsortedData,
           __global uint* buckets,
           uint shiftCount,
            __local uint* sharedArray)
{
    size_t localId = get_local_id(0);
    size_t globalId = get_global_id(0);
    size_t groupId = get_group_id(0);
    size_t groupSize = get_local_size(0);

    uint numGroups = get_global_size(0) / get_local_size(0);


    // Initialize shared array to zero //

    sharedArray[localId] = 0;

    barrier(CLK_LOCAL_MEM_FENCE);

   // Calculate thread-histograms //
    uint value = unsortedData[globalId];
  value =  value >> shiftCount & 0xFFU;   
    atomic_inc(sharedArray+value);


    barrier(CLK_LOCAL_MEM_FENCE);
    // Copy calculated histogram bin to global memory //

    uint bucketPos = groupId  * groupSize + localId ;
    //uint bucketPos = localId * numGroups + groupId ;
    buckets[bucketPos] = sharedArray[localId];
}

Outcomes