2 Replies Latest reply on May 26, 2013 9:00 AM by shunyo

    Substitute for atomic_inc

    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];
      }