完全在片上製作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)  
    { ...