Битонная сортировка
В этом примере будет реализована битонная сортировка. Битонная сортировка – это параллельный алгоритм сортировки, основанный на битонной последовательности. Такой последовательностью называют последовательность, которая монотонно не убывает, а затем монотонно не возрастает. Пример такой последовательности – {8, 10, 16, 12, 4, -2, -8}.
Ее можно эффективно отсортировать следующим образом: разобьем массив на две части, создадим два массива, в первый добавим все элементы, равные минимуму из соответственных элементов каждой из двух частей, а во второй – равные максимуму. Утверждается, что получатся две битонные последовательности, каждую из которых можно рекурсивно отсортировать тем же образом, после чего можно склеить два массива (так как любой элемент первого меньше или равен любого элемента второго). Для того, чтобы преобразовать исходный массив в битонную последовательность, сделаем следующее: если массив состоит из двух элементов, можно просто завершиться, иначе разделим массив пополам, рекурсивно вызовем от половинок алгоритм, после чего отсортируем первую часть по порядку, вторую в обратном порядке и склеим.
С помощью CUDA будет реализован данный алгоритм.
#include <stdlib.h> #include <stdio.h> #include <time.h> /* Каждый поток получает ровно одно значение в несортированном массиве */ #define THREADS 512 // 2^9 #define BLOCKS 32768 // 2^15 #define NUM_VALS THREADS*BLOCKS /* Функция печати результата выполнения программы*/ void print_elapsed(clock_t start, clock_t stop) { double elapsed = ((double) (stop - start)) / CLOCKS_PER_SEC; printf("Elapsed time: %.3fs\n", elapsed); } /* Генерация случайных чисел массива */ float random_float() { return (float)rand()/(float)RAND_MAX; } /* Печать массива */ void array_print(float *arr [], int length) { int i; for (i = 0; i < length; ++i) { printf("%1.3f ", arr[i]); } printf("\n"); } /* Заполнение массива */ void array_fill(float *arr, int length) { srand(time(NULL)); int i; for (i = 0; i < length; ++i) { arr[i] = random_float(); } } __global__ void bitonic_sort_step(float *dev_values, int j, int k) { unsigned int i, ixj; /* Сортировка i и ixj */ i = threadIdx.x + blockDim.x * blockIdx.x; ixj = i^j; /* Нити с наименьшими идентификаторами сортируют массив. */ if ((ixj)>i) { if ((i&k)==0) { /* Сортировка по возрастанию */ if (dev_values[i]>dev_values[ixj]) { /* обмен(i,ixj) */ float temp = dev_values[i]; dev_values[i] = dev_values[ixj]; dev_values[ixj] = temp; } } if ((i&k)!=0) { /* Сортировка по убыванию */ if (dev_values[i]<dev_values[ixj]) { /* обмен(i,ixj); */ float temp = dev_values[i]; dev_values[i] = dev_values[ixj]; dev_values[ixj] = temp; } } } } /** * Битонная сортировка на CUDA. */ void bitonic_sort(float *values) { float *dev_values; size_t size = NUM_VALS * sizeof(float); cudaMalloc((void**) &dev_values, size); cudaMemcpy(dev_values, values, size, cudaMemcpyHostToDevice); dim3 blocks(BLOCKS,1); /* Количество блоков */ dim3 threads(THREADS,1); /* Количество тредов */ int j, k; /* Основной шаг выполнения сортировки */ for (k = 2; k <= NUM_VALS; k <<= 1) { /* Второстепенный шаг выполнения сортировки */ for (j=k>>1; j>0; j=j>>1) { bitonic_sort_step<<<blocks, threads>>>(dev_values, j, k); } } cudaMemcpy(values, dev_values, size, cudaMemcpyDeviceToHost); cudaFree(dev_values); } int main(void) { clock_t start, stop; float *values = (float*) malloc( NUM_VALS * sizeof(float)); array_fill(values, NUM_VALS); start = clock(); bitonic_sort(values); stop = clock(); print_elapsed(start, stop); }
Результатом выполнения программы будет затраченное время на сортировку случайных значений несортированного массива по 512 тредам и 32768 блокам. Количество значений будет равно их произведению – NUM_VALS THREADS*BLOCKS. Затраченное время на исполнение:
