cuda - Hacer CUB blockradixsort en chip por completo?

gpu radix-sort

Question

Estoy leyendo la documentación y ejemplos de 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);
...
}

En el ejemplo, cada hilo tiene 4 claves. Parece que ''thread_keys'' se asignará en la memoria local global. Si solo tengo 1 clave por hilo, podría declarar "int thread_key;" y hacer que esta variable solo se registre?

BlockRadixSort (temp_storage) .Sort () toma un puntero a la clave como parámetro. ¿Significa que las claves tienen que estar en la memoria global?

Me gustaría utilizar este código, pero quiero que cada hilo contenga una clave en el registro y la mantenga en el chip en el registro / memoria compartida después de que se hayan ordenado. ¡Gracias por adelantado!

Answer 1

Puede hacerlo utilizando la memoria compartida (que lo mantendrá "en el chip"). No estoy seguro de saber cómo hacerlo utilizando estrictamente registros sin BlockRadixSort objeto BlockRadixSort .

Aquí hay un código de ejemplo que usa memoria compartida para contener los datos iniciales que se ordenarán y los resultados finales ordenados. Esta muestra está configurada principalmente para un elemento de datos por hilo, ya que parece ser lo que está pidiendo. No es difícil extenderlo a múltiples elementos por hilo, y he colocado la mayoría de las tuberías en su lugar para hacerlo, con la excepción de la síntesis de datos y las impresiones de depuración:

#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();

}

Esto parece funcionar correctamente para mí, en este caso yo estaba usando RHEL 5.5 / gcc 4.1.2, CUDA 6.0 RC y CUB v1.2.0 (que es bastante reciente).

Por lo que yo sé, se necesita una transmisión estática extraña / fea, porque CUB Sort espera una referencia a una matriz de longitud igual al parámetro de personalización ITEMS_PER_THREAD (es decir, ELEMS_PER_THREAD ):

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