[실습] CUDA MergeSort 구현 (Naive 버전)

 

CUDA에서 정렬을 구현한다고 할때, 가장 먼저 드는 생각은 우리가 알고 있는 알고리즘들(Quick Sort, Merge Sort, Buble Sort 등등.. )을 병렬로 바꿔볼까 하는생각들을 하게 될것 입니다. 그러면 이게 생각보다 쉽지 않다는 것을 느낄 수 있고 여기에 성능까지 고려하게 된다면 CUDA를 사용해 병렬로 정렬하다는 것이 의미 없어보이기도 합니다. 
(왜냐하면, std::sort에 비해 성능이 너무 안나오기 때문이죠.)

단순히 우리가 이번에 이야기 해볼 Merge Sort만 놓고 고민해 본다고 하더라도, 분활하여 비교하는 부분은 병렬로 처리할 수 있다고 생각하지만, 이후에 Merge 하는 과정은 특정 Thread만 수행하게 되어 CUDA를 사용하는 의미를 많이 잃게 됩니다.
(일반적인 Merge Sort의 알고리즘을 따라간다고 했을 때)

 

따라서 이번 글은 Merge Sort에 대해서 간단하게 이야기 하고, Merge Sort를 CUDA로 구현하는 과정에 대해서 이야기 해보려고 합니다.

---

일반 적으로 우리가 생각하는 Merge Sort는 아래와 같은 방식으로 진행됩니다.

Merge Sort - wikipedia

주어진 배열을 잘게 쪼개고 이를 정렬하여 새로운 배열를 만들어가는 방식입니다.

이때 새로운 배열를 만들어 가는 과정에서 입력으로 들어온 배열의 크기만큼의 추가 메모리가 필요합니다.

그리고 위 알고리즘을 그대로 CUDA로 구현한다면 아래 처럼 구현해 볼 수 있습니다.

// 두 개의 sub배열을 정렬된 하나의 배열로 합치는 함수
__device__ void merge(TARGET_INPUT_TYPE* input, TARGET_INPUT_TYPE* temp, unsigned int start, unsigned int idx, unsigned int stride, unsigned int size) {
    unsigned int left = start;
    unsigned int right = start + (stride / 2);
    const unsigned int mid = right < SIZE ? right : SIZE;
    const unsigned int end = start + stride < SIZE ? start + stride : SIZE;

    while (left < mid || right < end) {
        if (left < mid && right < end) {
            if (input[left] < input[right]) {
                temp[idx++] = input[left++];
            } else {
                temp[idx++] = input[right++];
            }
        } else if (left >= mid && right < end) {
            temp[idx++] = input[right++];
        } else if (right >= end && left < mid) {
            temp[idx++] = input[left++];
        }
    }
}

// block 사이즈만큼의 정렬된 sub 배열을 만드는 함수
__global__ void mergeSortStep1(TARGET_INPUT_TYPE* input)
{
    unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;
    __shared__ TARGET_INPUT_TYPE temp[THREADS];
    unsigned int stride = 2;

    while (stride < THREADS) {
        if (idx % stride == 0) {
            merge(input, temp, idx, threadIdx.x, stride, THREADS);
        }
        __syncthreads();
        if (idx < SIZE) {
            input[idx] = temp[threadIdx.x];
        }
        __syncthreads();
        stride <<= 1;
    }
}

//  block 사이즈 이상의 배열들을 하나라 합치는 함수입니다.
__global__ void mergeSortStep2(TARGET_INPUT_TYPE* input, TARGET_OUTPUT_TYPE* output, unsigned int stride)
{
    bool flag = true;
    unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;
    TARGET_INPUT_TYPE* buffer1 = input;
    TARGET_INPUT_TYPE* buffer2 = output;

    while (stride / 2 < SIZE) {
        if (idx % stride == 0) {
            if (flag) {
                merge(buffer1, buffer2, idx, idx, stride, SIZE);
            } else {
                merge(buffer2, buffer1, idx, idx, stride, SIZE);
            }
        }
        __syncthreads();
        flag = !flag;
        stride <<= 1;
    }
    
    if (flag && idx < SIZE) {
        output[idx] = input[idx];
    }
}

...

dim3 gridDim(divideUp(SIZE, THREADS));
dim3 blockDim(THREADS);

mergeSortStep1<<<gridDim, blockDim>>>(inputs[DEVICE_INPUT]);
mergeSortStep2<<<gridDim, blockDim>>>(inputs[DEVICE_INPUT], outputs[DEVICE_OUTPUT], THREADS);

mergeSortStep1부터 확인하면 CPU에서 존재하던 배열을 분해하는 단계가 필요 없습니다. 왜냐하면 각각의 Thread가 자기가 맡고 있는 구역을 가지고 바로 판단하면 되기 때문입니다. 이렇게 mergeSortStep1에서는 Bottom-Up 방식으로 Shared Memory 크기만큼 배열을 정렬시키고 있습니다. 그런데 한눈에 보기에도 굉장히 비효율적인 부분을 많이 찾아 볼 수 있습니다.

 

우선 모든 Thread가 merge에 가담하지 않고 있습니다. 이 말인 즉 노는 Thread가 많다는 의미입니다. 그리고 이렇게 노는 Thread는 stride가 커질 수록 지수함수 크기로 증가하게 됩니다. 그리고 결국 단 한개의 Thread가 두 개의 sub배열에서 값을 읽어오며 정렬된 배열을 만들게 됩니다.
(물론 shared memory에서 데이터를 복사하는 과정에서 모든 스레드를 활용한다고 볼 수 있습니다.)

 

두 번째로 Coalesced Memory Access하지 않고 있습니다. 특히 이 역시 stride가 커질 수록 Global Memory에 접근하는 스레드 자체가 줄어들게 되고, 최후에는 warp내 Thread 들이 2개의 데이터에만 접근하게 되어 bandwidth가 낭비되게 됩니다.

 

그러면 어느정도 성능이 나오는지 판단해 보기 위해서 1,000,000개의 랜덤한 데이터를 가지를 배열을 정렬해 봤습니다.

CUDA 구현체의 경우에는 총 216.094ms가 걸린데 반해, std::qsort는 124.551ms가 걸렸습니다.

100만개 정렬

데이터를 1억개로 늘려도 std::qsort가 더 빠르게 측정됩니다.

1억개 정렬

단순히 CPU에서 동작하는 방식을 CUDA로 구현 하는것만으론 실제로 활용할 수 있는 sort 알고리즘을 구현하기 어려워 보입니다. CUDA에 맞는 새로운 알고리즘을 찾아 봐야 할거 같습니다.

 

전체적인 소스코드는 아래의 링크에서 참고해보실 수 있습니다.

Cuda_Study/0_1_MergeSort.cu at main · hotstone1993/Cuda_Study (github.com)

GitHub - hotstone1993/Cuda_Study