cub BlockRadixSort：如何處理大的瓷磚大小或排序多個瓷磚？

Question

當使用cub :: BlockRadixSort在塊中進行排序時，如果元素數量太大，我們該如何處理？如果我們將圖塊大小設置得太大，臨時存儲的共享內存很快就無法存儲。如果我們將它分成多個圖塊，我們如何在對每個圖塊進行排序後對其進行後處理？

Answer 1

• 警告：我不是一個幼仔專家（遠離它）。
• 您可能想檢查此question/answer，因爲我正在構建我在那裏做的一些工作。
• 當然，如果問題的規模足夠大，那麼device-wide sort似乎是你可能要考慮的事情。但你的問題似乎集中在塊排序。

從我的測試中，幼崽並沒有真正的要求你的原始數據的位置，或你放置臨時存儲的位置。因此，一種可能的解決方案就是將臨時存儲放置在全局內存中。爲了分析這一點，我創建了一個包含3個不同測試用例的代碼：

1. 測試一個帶有臨時存儲在全局內存中的小型塊排序版本。
2. 測試的原始版本的幼崽塊從例如here
3. 測試版本的幼崽塊排序從我以前的答案，在沒有從全局內存，即數據/複製衍生的排序調整。假定數據已經駐留在「片上」即共享存儲器中。這

沒有經過廣泛的測試，但因爲我建立在幼崽積木，並在前兩種情況下的測試我的結果，希望我還沒有做出任何嚴重錯誤。下面是完整的測試代碼，而我會在下面的補充意見：

$ cat t10.cu 
#include <cub/cub.cuh> 
#include <stdio.h> 
#include <stdlib.h> 
#include <thrust/sort.h> 
#define nTPB 512 
#define ELEMS_PER_THREAD 2 
#define RANGE (nTPB*ELEMS_PER_THREAD) 
#define DSIZE (nTPB*ELEMS_PER_THREAD) 



#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

using namespace cub; 
// GLOBAL CUB BLOCK SORT KERNEL 
// Specialize BlockRadixSort collective types 
typedef BlockRadixSort<int, nTPB, ELEMS_PER_THREAD> my_block_sort; 
__device__ int my_val[DSIZE]; 
__device__ typename my_block_sort::TempStorage sort_temp_stg; 

// Block-sorting CUDA kernel (nTPB threads each owning ELEMS_PER THREAD integers) 
__global__ void global_BlockSortKernel() 
{ 
    // 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)))); 

} 

// ORIGINAL CUB BLOCK SORT KERNEL 
template <int BLOCK_THREADS, int ITEMS_PER_THREAD> 
__global__ void BlockSortKernel(int *d_in, int *d_out) 
{ 
// Specialize BlockLoad, BlockStore, and BlockRadixSort collective types 
    typedef cub::BlockLoad<int*, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_LOAD_TRANSPOSE> BlockLoadT; 
    typedef cub::BlockStore<int*, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_STORE_TRANSPOSE> BlockStoreT; 
    typedef cub::BlockRadixSort<int, BLOCK_THREADS, ITEMS_PER_THREAD> BlockRadixSortT; 
// Allocate type-safe, repurposable shared memory for collectives 
    __shared__ union { 
    typename BlockLoadT::TempStorage load; 
    typename BlockStoreT::TempStorage store; 
    typename BlockRadixSortT::TempStorage sort; 
    } temp_storage; 
// Obtain this block's segment of consecutive keys (blocked across threads) 
    int thread_keys[ITEMS_PER_THREAD]; 
    int block_offset = blockIdx.x * (BLOCK_THREADS * ITEMS_PER_THREAD); 
    BlockLoadT(temp_storage.load).Load(d_in + block_offset, thread_keys); 
    __syncthreads(); // Barrier for smem reuse 
// Collectively sort the keys 
    BlockRadixSortT(temp_storage.sort).Sort(thread_keys); 
    __syncthreads(); // Barrier for smem reuse 
// Store the sorted segment 
    BlockStoreT(temp_storage.store).Store(d_out + block_offset, thread_keys); 
} 



// SHARED MEM CUB BLOCK SORT KERNEL 
// Block-sorting CUDA kernel (nTPB threads each owning ELEMS_PER THREAD integers) 
template <int BLOCK_THREADS, int ITEMS_PER_THREAD> 
__global__ void shared_BlockSortKernel(int *d_out) 
{ 
    __shared__ int my_val[BLOCK_THREADS*ITEMS_PER_THREAD]; 
    // Specialize BlockRadixSort collective types 
    typedef BlockRadixSort<int, BLOCK_THREADS, ITEMS_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*ITEMS_PER_THREAD] = (threadIdx.x + 5); // synth data 
    my_val[threadIdx.x*ITEMS_PER_THREAD+1] = (threadIdx.x + BLOCK_THREADS + 5); // synth data 
    __syncthreads(); 
// printf("thread %d data = %d\n", threadIdx.x, my_val[threadIdx.x*ITEMS_PER_THREAD]); 

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

// printf("thread %d sorted data = %d\n", threadIdx.x, my_val[threadIdx.x*ITEMS_PER_THREAD]); 
    if (threadIdx.x == clock()){ // dummy to prevent compiler optimization 
     d_out[threadIdx.x*ITEMS_PER_THREAD] = my_val[threadIdx.x*ITEMS_PER_THREAD]; 
     d_out[threadIdx.x*ITEMS_PER_THREAD+1] = my_val[threadIdx.x*ITEMS_PER_THREAD+1];} 
} 


int main(){ 
    int *h_data, *h_result; 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    h_data=(int *)malloc(DSIZE*sizeof(int)); 
    h_result=(int *)malloc(DSIZE*sizeof(int)); 
    if (h_data == 0) {printf("malloc fail\n"); return 1;} 
    if (h_result == 0) {printf("malloc fail\n"); return 1;} 
    for (int i = 0 ; i < DSIZE; i++) h_data[i] = rand()%RANGE; 
    // first test sorting directly out of global memory 
    global_BlockSortKernel<<<1,nTPB>>>(); //warm up run 
    cudaDeviceSynchronize(); 
    cudaMemcpyToSymbol(my_val, h_data, DSIZE*sizeof(int)); 
    cudaCheckErrors("memcpy to symbol fail"); 
    cudaEventRecord(start); 
    global_BlockSortKernel<<<1,nTPB>>>(); //timing run 
    cudaEventRecord(stop); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("cub 1 fail"); 
    cudaEventSynchronize(stop); 
    float et; 
    cudaEventElapsedTime(&et, start, stop); 
    cudaMemcpyFromSymbol(h_result, my_val, DSIZE*sizeof(int)); 
    cudaCheckErrors("memcpy from symbol fail"); 
    if(!thrust::is_sorted(h_result, h_result+DSIZE)) { printf("sort 1 fail!\n"); return 1;} 
    printf("global Elapsed time: %fms\n", et); 
    printf("global Kkeys/s: %d\n", (int)(DSIZE/et)); 
    // now test original CUB block sort copying global to shared 
    int *d_in, *d_out; 
    cudaMalloc((void **)&d_in, DSIZE*sizeof(int)); 
    cudaMalloc((void **)&d_out, DSIZE*sizeof(int)); 
    cudaCheckErrors("cudaMalloc fail"); 
    BlockSortKernel<nTPB, ELEMS_PER_THREAD><<<1, nTPB>>>(d_in, d_out); // warm up run 
    cudaMemcpy(d_in, h_data, DSIZE*sizeof(int), cudaMemcpyHostToDevice); 
    cudaEventRecord(start); 
    BlockSortKernel<nTPB, ELEMS_PER_THREAD><<<1, nTPB>>>(d_in, d_out); // timing run 
    cudaEventRecord(stop); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("cub 2 fail"); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&et, start, stop); 
    cudaMemcpy(h_result, d_out, DSIZE*sizeof(int), cudaMemcpyDeviceToHost); 
    cudaCheckErrors("cudaMemcpy D to H fail"); 
    if(!thrust::is_sorted(h_result, h_result+DSIZE)) { printf("sort 2 fail!\n"); return 1;} 
    printf("CUB Elapsed time: %fms\n", et); 
    printf("CUB Kkeys/s: %d\n", (int)(DSIZE/et)); 
    // now test shared memory-only version of block sort 
    shared_BlockSortKernel<nTPB, ELEMS_PER_THREAD><<<1, nTPB>>>(d_out); // warm-up run 
    cudaEventRecord(start); 
    shared_BlockSortKernel<nTPB, ELEMS_PER_THREAD><<<1, nTPB>>>(d_out); // timing run 
    cudaEventRecord(stop); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("cub 3 fail"); 
    cudaEventSynchronize(stop); 
    cudaEventElapsedTime(&et, start, stop); 
    printf("shared Elapsed time: %fms\n", et); 
    printf("shared Kkeys/s: %d\n", (int)(DSIZE/et)); 
    return 0; 
} 
$ nvcc -O3 -arch=sm_20 -o t10 t10.cu 
$ ./t10 
global Elapsed time: 0.236960ms 
global Kkeys/s: 4321 
CUB Elapsed time: 0.042816ms 
CUB Kkeys/s: 23916 
shared Elapsed time: 0.040192ms 
shared Kkeys/s: 25477 
$ 

對於這個測試，我使用CUDA 6.0RC，幼獸V1.2.0（這是相當近期的），RHEL5.5/gcc4.1.2和Quadro5000 GPU（cc2.0,11SMs，比GTX480慢大約40％）。這裏有一些意見，即發生給我：

1. 原始幼獸排序（2）向所述全局存儲器排序的變速比（1）爲約6：1，這大約是共享存儲器的帶寬比（ 〜1TB/s）到全局存儲器（〜150GB/s）。
2. 原始Cub類（2）具有吞吐量，當SMs（11）的數量縮放時，產生263MKeys/s，是我在此設備上看到的最佳設備範圍排序的很大一部分（thrust sort ，產生〜480MKeys/s）
3. 只有共享內存的排序並不比從原始的Cub文件排序快得多，它將輸入/輸出從/複製到全局內存，表明從全局內存複製到Cub文件臨時存儲不是整個處理時間的很大一部分。

6：1的罰款是一個很大的支付。所以我的建議是，如果可能的話，在問題的大小上使用設備範圍的排序大於小塊排序容易處理的問題。這使您可以利用一些最優秀的GPU代碼編寫器的專業知識進行分類，並實現更接近整個設備的吞吐量。

請注意，所以我可以在類似條件下進行測試，此處的問題大小（512個線程，每個線程2個元素）不會超過您在CUB塊排序中可以執行的操作。但是，將數據集大小擴展到更大的值（例如，每個線程有1024個元素）並不難，只能使用第一種方法處理（在這種情況下，這些選擇之間）。如果我這樣做了更大的問題，那麼在我的GPU上，我的cc2.0設備上的全局內存塊排序爲的吞吐量大約爲6Mkeys/s。