cub BlockRadixSort: как бороться с большим размером плитки или сортировать несколько плиток?

Question

При использовании cub::BlockRadixSort для сортировки внутри блока, если количество элементов слишком велико, как мы с этим справимся? Если мы установим слишком большой размер тайла, общая память для временного хранилища скоро не сможет его удержать. Если мы разделим его на несколько плиток, как мы постобработаем его после сортировки каждой плитки?

Answer 1

• Предостережение: я не эксперт по детёнышу (далеко не так).
• Возможно, вы захотите просмотреть этот вопрос / ответ, так как я строю некоторые работы, которые я там делал.
• Конечно, если размер проблемы достаточно велик, то сортировка по всему устройству может показаться вам интересной. Но ваш вопрос кажется сосредоточенным на сортировке блоков.

Из моего тестирования у cub нет никаких требований относительно того, где находятся ваши исходные данные или где вы размещаете временное хранилище. Поэтому одним из возможных решений было бы просто разместить ваше временное хранилище в глобальной памяти. Чтобы проанализировать это, я создал код, который имеет 3 разных тестовых случая:

1. Протестируйте версию сортировки кубов с временным хранилищем в глобальной памяти.
2. Протестируйте оригинальную версию сортировки кубов, адаптированную из примера, здесь
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, cub v1.2.0 (что довольно недавно), RHEL5.5/gcc4.1.2 и графический процессор Quadro5000 (cc2.0, 11SM, примерно на 40% медленнее, чем GTX480). Вот некоторые наблюдения, которые происходят со мной:

1. Отношение скорости исходной сортировки куба (2) к сортировке глобальной памяти (1) составляет примерно 6:1, что примерно равно отношению пропускной способности разделяемой памяти (~1 ТБ / с) к глобальной памяти (~150 ГБ / с).
2. Исходная сортировка куба (2) обладает пропускной способностью, которая при масштабировании на количество SM (11), приводящее к 263MKeys/s, представляет собой значительную долю лучшей сортировки по всему устройству, которую я видел на этом устройстве ( упорная сортировка, уступающая ~480MKeys/ с)
3. Сортировка только с разделяемой памятью не намного быстрее, чем исходная сортировка куба, которая копирует ввод / вывод из / в глобальную память, что указывает на то, что копирование из глобальной памяти во временную память куба не является большой долей общего времени обработки.

Штраф 6: 1 - большой штраф. Поэтому я бы порекомендовал, если возможно, использовать сортировку по всему устройству для задач с размерами, большими, чем те, которые могут быть легко обработаны сортировкой кубов. Это позволяет вам использовать опыт некоторых из лучших разработчиков кода GPU для вашей сортировки и получать пропускную способность, намного более близкую к возможностям устройства в целом.

Обратите внимание, что, чтобы я мог протестировать в аналогичных условиях, размер проблемы здесь (512 потоков, 2 элемента на поток) не превышает того, что вы можете сделать при сортировке блоков CUB. Но нетрудно расширить размер набора данных до больших значений (скажем, 1024 элемента на поток), которые могут обрабатываться только (в данном контексте, среди этих вариантов) с использованием первого подхода. Если я выполняю такие большие задачи, то на моем GPU я наблюдаю пропускную способность около 6Mkeys/s для сортировки блоков глобальной памяти на моем устройстве cc2.0.