1

我有一个不支持 cl_khr_int64_base_atomics 的 ATI Firepro V4800 显卡。我正在尝试将 RadixSort 算法调整为长整数。该算法使用 atomic_inc，它的 64 位是 atom_inc，我不能在内核中使用它。所以，我的问题是，是否有一段代码可以执行与 atomic_inc 相同的功能？内核代码如下：

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

有什么建议么？谢谢你。

编辑： 这个博客网站给出了另一种方法：http: //suhorukov.blogspot.in/2011/12/opencl-11-atomic-operations-on-floating.html。这给出了 Atomic Inc. 的一个非常通用的实现。

4

1 回答 1

1

你可以尝试这样的事情：

void atomInc64 (__local uint *counter)
   {
   uint old, carry;

   old = atomic_inc (&counter [0]);
   carry = old == 0xFFFFFFFF;
   atomic_add (&counter [1], carry);
   }

其中 counter 是两个 32 位整数的数组。虽然两半不会同时增加，但程序完成时总数应该是正确的。

于 2013-05-24T00:21:25.403 回答