2

我正在阅读 CUB 文档和示例： 

#include <cub/cub.cuh>   // or equivalently <cub/block/block_radix_sort.cuh>
__global__ void ExampleKernel(...)
{
    // Specialize BlockRadixSort for 128 threads owning 4 integer items each
typedef cub::BlockRadixSort<int, 128, 4> BlockRadixSort;
    // Allocate shared memory for BlockRadixSort
__shared__ typename BlockRadixSort::TempStorage temp_storage;
    // Obtain a segment of consecutive items that are blocked across threads
int thread_keys[4];
...
    // Collectively sort the keys
BlockRadixSort(temp_storage).Sort(thread_keys);
...
}

在示例中，每个线程有 4 个键。看起来“thread_keys”将被分配到全局本地内存中。如果我每个线程只有 1 个密钥，我可以声明“int thread_key;”吗？并仅在寄存器中创建此变量？

BlockRadixSort(temp_storage).Sort() 将指向键的指针作为参数。这是否意味着密钥必须在全局内存中？

我想使用这段代码，但我希望每个线程在寄存器中保存一个键，并在排序后将其保存在寄存器/共享内存中。提前致谢！

4

1 回答 1

4

您可以使用共享内存来做到这一点（这将保持它“在芯片上”）。我不确定我是否知道如何在不解构BlockRadixSort对象的情况下使用严格的寄存器来做到这一点。

这是一个使用共享内存来保存要排序的初始数据和最终排序结果的示例代码。此示例主要针对每个线程的一个数据元素设置，因为这似乎是您所要求的。将它扩展到每个线程的多个元素并不难，我已经把大部分管道都放在了地方，除了数据合成和调试打印输出：

#include <cub/cub.cuh>
#include <stdio.h>
#define nTPB 32
#define ELEMS_PER_THREAD 1

// Block-sorting CUDA kernel (nTPB threads each owning ELEMS_PER THREAD integers)
__global__ void BlockSortKernel()
{
    __shared__ int my_val[nTPB*ELEMS_PER_THREAD];
    using namespace cub;
    // Specialize BlockRadixSort collective types
    typedef BlockRadixSort<int, nTPB, ELEMS_PER_THREAD> my_block_sort;
    // Allocate shared memory for collectives
    __shared__ typename my_block_sort::TempStorage sort_temp_stg;

    // need to extend synthetic data for ELEMS_PER_THREAD > 1
    my_val[threadIdx.x*ELEMS_PER_THREAD]  = (threadIdx.x + 5)%nTPB; // synth data
    __syncthreads();
    printf("thread %d data = %d\n", threadIdx.x,  my_val[threadIdx.x*ELEMS_PER_THREAD]);

    // Collectively sort the keys
    my_block_sort(sort_temp_stg).Sort(*static_cast<int(*)[ELEMS_PER_THREAD]>(static_cast<void*>(my_val+(threadIdx.x*ELEMS_PER_THREAD))));
    __syncthreads();

    printf("thread %d sorted data = %d\n", threadIdx.x,  my_val[threadIdx.x*ELEMS_PER_THREAD]);
}

int main(){
    BlockSortKernel<<<1,nTPB>>>();
    cudaDeviceSynchronize();

}

这对我来说似乎工作正常，在这种情况下，我碰巧使用的是 RHEL 5.5/gcc 4.1.2、CUDA 6.0 RC 和 CUB v1.2.0（这是最近的）。

据我所知，需要奇怪/丑陋的静态转换Sort，因为 CUB期望引用长度等于自定义参数（即）的数组：ITEMS_PER_THREADELEMS_PER_THREAD

   __device__ __forceinline__ void Sort(
        Key     (&keys)[ITEMS_PER_THREAD],          
        int     begin_bit   = 0,                   
        int     end_bit     = sizeof(Key) * 8)      
   { ...
于 2014-02-27T16:17:51.277 回答