Делать 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 будет размещен в глобальной локальной памяти. Если бы у меня был только 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 (что довольно недавно).

Насколько я могу судить, странное / безобразное статическое приведение в порядок необходимо, потому что CUB Sort ожидает ссылку на массив длины, равный параметру настройки 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)      
   { ...