cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

shunyo
Journeyman III

Substitute for atomic_inc

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

0 Likes
1 Solution
himanshu_gautam
Grandmaster

check http://devgurus.amd.com/thread/139081 and http://suhorukov.blogspot.in/2011/12/opencl-11-atomic-operations-on-floating.html

You can probably try something like:

Union mydata { int2 intVal; long longVal;};

Now using atomic_incs for intVal element. Please let us know, if this works out for you.

View solution in original post

0 Likes
2 Replies
himanshu_gautam
Grandmaster

check http://devgurus.amd.com/thread/139081 and http://suhorukov.blogspot.in/2011/12/opencl-11-atomic-operations-on-floating.html

You can probably try something like:

Union mydata { int2 intVal; long longVal;};

Now using atomic_incs for intVal element. Please let us know, if this works out for you.

0 Likes

Thanks Himanshu for pointing me to the blogpost. I have worked it out.

0 Likes