Mergesort using CUDA

Question

Я попытался реализовать свой собственный Mergesort, основанный на алгоритме с нижним/итеративным объединением. Этот алгоритм разбил данные на 2 элемента и отсортировал их. Затем по 4 элементам и отсортированы и так далее, пока все данные не отсортированы. Итак, мой план назначает каждый поток на 2 элемента. Так я это сделать:

__global__ void mergeBU(int *d_a, int *d_aux, int sz, int N) 
{ 
    int idk = blockIdx.x*blockDim.x+threadIdx.x; 
    int lo = 2 * sz * idk; 
    int mid = lo + sz - 1; 
    float hi = fminf(lo + sz + sz - 1, N - 1); 
    merge(d_a, d_aux, lo, mid, hi); 
} 

__device__ void merge(int *d_a, int *d_aux, int lo, int mid, float hi) 
{ 
int i = lo; 
int j = mid + 1; 

    for (int k = lo; k <= hi; k++) 
    { 
     d_aux[k] = d_a[k]; 
    } 

    for (int k = lo; k <= hi; k++) 
    { 
     if (i > mid)     { d_a[k] = d_aux[j]; j++; } 
     else if (j > hi)    { d_a[k] = d_aux[i]; i++; } 
     else if (d_aux[j] < d_aux[i]) { d_a[k] = d_aux[j]; j++; } 
     else        { d_a[k] = d_aux[i]; i++; } 
    } 
} 

Допустим, я призываю мое ядро ​​< < < 2,4 >>> (который 8 потоков), поэтому я могу только сортировать 16 элементов макс. Если я ввожу 32 элемента, то остальная часть индекса данных не затронута (16-31). Как заставить индекс потока продолжать обрабатывать остальную часть индекса данных? По-прежнему я имею в виду, что индекс потоков (0,1,2,3,4,5,6,7) продолжает обрабатывать остальную часть индекса данных, он должен быть похож на threadindex (dataindex, dataindex) -> 0 (16, 17); 1 (18,19); 2 (20,21); и так далее. Любые комментарии приветствуются.

Answer 1

Без учета вашего фактического кода: Сортировка слияний - это многопроходный алгоритм. Поскольку различные блоки обычно не синхронизируются вообще при выполнении ядра (если вы не используете атомарную аппаратуру на уровне устройства), вам, вероятно, следует рассмотреть несколько последующих запусков ядра, по одному для каждого прохода. Например, при первом запуске каждый блок потоков сортирует n_1 элементов; при втором запуске каждая пара блоков объединяет 2 * n_1 элементов и так далее. Конечно, это не так просто, как кажется: как вы можете сказать, какой блок должен делать то, что именно?

Кроме того, вы можете взглянуть на approach used in the ModernGPU library на другие идеи.

Answer 2

Похоже, что ваш подход состоит в том, чтобы разделить массив размером n на n/2 подматрицы, объединить пары этих подматриц, чтобы в итоге были добавлены n/4 вспомогательные массивы и т. Д. Однако этот подход, вероятно, будет ограничен ограниченной пропускной способностью.

Предположим, вы решили использовать 8 "нитей". Разделите массив на 8 под-массивов размером n/8 каждый (последний подребер может быть другого размера), затем используйте 8 потоков для объединения сортировки подматрицы, а затем 4 потока для объединения 4 пар отсортированной подстанции -arrays, затем 2 потока для объединения двух пар объединенных подмассивов, затем 1 поток для объединения последних двух пар.

Основываясь на моем опыте с многопоточным сортированием, вы достигаете предела пропускной способности памяти в 8 потоках для процессора, но если память gpu может использоваться для хранения больших секций массива, тогда более 8 потоков могут быть полезным. Я не знаю, какие операции (сравнивать, перемещать, ...) возможны внутри gpu и его памяти.