Выполнение полной блокировки блока CUB blockixixsort?

Question

Я читаю детёныш документацию и примеры:

#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() принимает указатель на ключ в качестве параметра. Означает ли это, что ключи должны находиться в глобальной памяти?

Я хотел был бы использовать этот код, но я хочу, чтобы каждый поток удерживал один ключ в регистре и сохранял его на чипе в регистровой/разделяемой памяти после их сортировки. Спасибо заранее!

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 нужен, насколько я могу судить, потому что CUB 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)  
    { ...