完全在片上制作CUB blockradixsort？

Question

我读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'将被分配在全局本地内存中。如果我每个线程只有一个键，我可以声明“int thread_key;”并只在这个变量中注册？

BlockRadixSort（temp_storage）.Sort（）将指向该键的指针作为参数。这是否意味着密钥必须位于全局内存中？

我想使用此代码，但我希望每个线程在寄存器中保存一个密钥，并在排序后将它保存在寄存器/共享内存中。 在此先感谢！

Answer 1

您可以使用共享内存（这将保持它“片上”）。我不知道如何使用严格的寄存器来完成这个操作，而不需要去构造这个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（这是相当新的）。

奇怪/丑static casting是，据我可以告诉需要，因为古巴Sort的长度为expecting a reference to an array等于自定义参数ITEMS_PER_THREAD（即ELEMS_PER_THREAD）：

__device__ __forceinline__ void Sort(
     Key  (&keys)[ITEMS_PER_THREAD],   
     int  begin_bit = 0,     
     int  end_bit  = sizeof(Key) * 8)  
    { ...