정렬: 라딕스와 머지소트

13.1 배경: 비교 정렬과 분배 정렬

정렬은 컴퓨터 과학에서 가장 오래된 문제이자, 동시에 GPU에서 가장 흥미로운 케이스다. 직렬에서야 퀵소트 한 줄로 끝낼 일이지만, 수만 개 스레드가 동시에 같은 배열을 다투면 이야기가 달라진다. 우선 정렬을 두 부류로 나눠 보자.

첫째, 비교 기반 정렬(comparison sort)이다. 퀵소트, 힙소트, 머지소트가 여기 속한다. 원소들을 두 개씩 비교해서 순서를 결정한다는 공통점이 있다. 정보 이론적으로 N개의 원소를 비교만으로 구별하려면 log₂(N!) ≈ N log N 번의 비교가 필요하다는 하한이 알려져 있다. 즉 비교 정렬은 절대 N log N보다 빨라질 수 없다.

둘째, 분배 기반 정렬(distribution sort)이다. 카운팅 소트(counting sort), 라딕스 소트(radix sort), 버킷 소트(bucket sort)가 대표적이다. 이 친구들은 비교를 하지 않는다. 대신 키의 값 자체를 보고 어디로 갈지 결정한다. 그래서 N log N 하한을 비껴가서 O(N) 또는 O(N · k)로 끝낼 수 있다. 단, 키가 정수처럼 분배 가능한 형태여야 한다는 제약이 따른다.

GPU 친화성이라는 잣대로 보면 둘은 또 다른 풍경을 보여 준다. 비교 정렬 중에서도 머지소트(merge sort)는 12장에서 본 병렬 머지를 그대로 활용할 수 있어 GPU에 잘 어울린다. 반면 퀵소트는 피벗을 잡는 순간 병렬성이 잘게 쪼개져서 그다지 매력적이지 않다. 분배 정렬 쪽에서는 라딕스 소트가 단연 챔피언이다. 비트 단위로 동일한 작업을 반복하기에 SIMD에 잘 맞고, 11장에서 익힌 스캔을 그대로 끼워 넣어 쓰면 된다.

13.2 라딕스 소트: 한 비트씩 줄 세우기

라딕스 소트의 아이디어는 우체국 분류대를 떠올리면 가깝다. 편지를 우편번호 마지막 자리부터 한 자리씩 보면서 0~9 칸에 나눠 담고, 다음 자리, 그 다음 자리로 넘어가며 같은 작업을 반복한다. 모든 자리를 처리하면 자연스럽게 정렬이 끝난다. 핵심은 매 패스가 안정 정렬(stable sort)이어야 한다는 점이다. 같은 자릿값을 가진 두 원소의 상대적 순서가 보존되어야 다음 자리 패스가 망가지지 않는다.

이진 표현으로 옮겨 보자. 키가 32비트 정수라면, 가장 단순한 라딕스 소트는 1비트 라딕스다. 즉 한 패스에서 한 비트만 본다. 0인 원소는 앞쪽으로, 1인 원소는 뒤쪽으로. 32비트면 패스를 32번 돌면 끝난다. 좀 더 일반화하자면 b-bit 라딕스는 한 패스에서 b 비트를 보고 2^b개의 빈(bin)으로 분배한다. 이 경우 패스 수는 N_bits/b가 된다. 32비트 키에 4-bit 라딕스를 쓰면 패스가 8번이다.

그림 13.1 — 1비트 라딕스 한 패스의 모습. 안정 정렬이라 0묶음 안에서의 원래 순서(6,2,4,0)가 그대로 보존된다.

13.3 병렬 라딕스 소트: 마스크, 스캔, 스캐터

병렬 환경에서 한 패스를 어떻게 수행할까. 단순하게 "0이면 왼쪽, 1이면 오른쪽" 한 줄을 GPU로 옮기려면 의외로 까다롭다. 각 스레드가 자기 원소를 0묶음의 어느 위치, 또는 1묶음의 어느 위치에 써야 할지 먼저 알아야 한다. 이 자리 계산이 바로 스캔(scan)의 임무다.

한 패스를 두 단계로 쪼갠다.

1. 마스크 생성: 각 원소의 해당 비트가 0인지 1인지 본다. 0이면 1, 1이면 0인 비트맵을 만든다(0묶음의 카운트를 위해). 또는 0/1 자체를 저장한다.
2. 익스클루시브 스캔(exclusive scan): 0묶음 위치를 알기 위해 0-마스크에 대해 exclusive scan을 돌린다. 그러면 각 원소의 출력 인덱스가 즉시 나온다. 1묶음은 전체 0의 개수만큼 오프셋을 더해 주면 된다.
3. 스캐터(scatter): 계산된 인덱스로 출력 배열에 원소를 배치한다.

// 1-bit radix sort, 한 패스 (단순화)
__global__ void radixPassNaive(
    const unsigned int* in,    // 입력 키
    unsigned int* out,           // 출력 키
    int* mask,                     // 0이면 1, 1이면 0
    int* scan,                     // exclusive scan 결과
    int totalZeros,                // 전체 0의 수
    int N, int bit) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i >= N) return;

  unsigned int key = in[i];
  int b = (key >> bit) & 1;

  int dst;
  if (b == 0) {
    // 0묶음: scan 결과가 곧 출력 위치
    dst = scan[i];
  } else {
    // 1묶음: 자기 앞에 있는 1의 개수 = i - scan[i]
    dst = totalZeros + (i - scan[i]);
  }
  out[dst] = key;
}

이 코드에서 scan[]은 미리 별도 커널로 11장의 work-efficient scan을 돌려서 만들어 둔 결과라고 가정했다. 실제 라이브러리 구현은 마스크와 스캔을 같은 블록에서 합쳐서 처리해 메모리 트래픽을 줄인다.

이걸 32번 반복하면 32비트 정수 정렬이 끝난다. 한 패스가 O(N), 패스가 32번이니 총 비용은 O(32N) = O(N). 비교 정렬의 N log N 하한을 비껴간 셈이다.

13.4 코얼레싱을 살리는 트릭

위 단순 구현에는 큰 함정이 있다. 마지막 out[dst] = key가 글로벌 메모리에 쓰는 위치다. 인접 스레드들의 dst가 들쭉날쭉하면 코얼레싱이 무너진다. 0묶음 영역과 1묶음 영역으로 분리되니, 같은 워프 내에서 어떤 스레드는 0묶음 앞쪽으로, 어떤 스레드는 1묶음 뒷쪽으로 흩어져 쓴다. 이는 한 트랜잭션이 둘로 갈라지는 결과를 낳고, 메모리 효율을 절반으로 떨어뜨린다.

해법은 간단하다. 먼저 셰어드 메모리에 정리한 뒤 한꺼번에 글로벌에 쓰는 것이다. 블록 내에서 자체 라딕스 패스를 돌려, 블록 단위로는 이미 0묶음/1묶음이 깔끔하게 정렬된 상태를 만든다. 그 다음 블록 단위로 글로벌 위치를 받아서 연속적으로 쓰면, 워프 내 스레드들의 글로벌 주소가 인접해져 코얼레싱이 회복된다.

// 블록 단위 라딕스 패스 (write coalescing 회복)
__global__ void radixPassBlock(
    const unsigned int* in,
    unsigned int* out,
    int* blockOffsets,  // 블록별 0/1묶음 글로벌 시작 위치
    int N, int bit) {
  __shared__ unsigned int sKey[BLOCK];
  __shared__ int sFlag[BLOCK];

  int tid = threadIdx.x;
  int i = blockIdx.x * BLOCK + tid;
  sKey[tid] = (i < N) ? in[i] : 0xFFFFFFFF;
  sFlag[tid] = (i < N) ? 1 - ((sKey[tid] >> bit) & 1) : 0;
  __syncthreads();

  // 블록 내 exclusive scan (생략, 11장 참고)
  blockExclusiveScan(sFlag);

  // 블록 내 재배치 (셰어드 메모리에서)
  int zerosInBlock = sFlag[BLOCK-1] + (((sKey[BLOCK-1] >> bit) & 1) == 0);
  int dst = (((sKey[tid] >> bit) & 1) == 0)
           ? sFlag[tid]
           : zerosInBlock + (tid - sFlag[tid]);
  __shared__ unsigned int sSorted[BLOCK];
  sSorted[dst] = sKey[tid];
  __syncthreads();

  // 글로벌에 연속 쓰기 → coalesced
  if (i < N) {
    int globalDst = blockOffsets[blockIdx.x] + tid;
    out[globalDst] = sSorted[tid];
  }
}

실제 구현은 한 단계 더 나아가서 블록별 카운트를 또 한번 글로벌 스캔해서 블록 간 위치까지 한꺼번에 결정한다. 이렇게 두 단계 스캔으로 패스를 구성하는 것이 표준이다.

13.5 라딕스 비트 폭의 trade-off

이제 b를 얼마로 잡을지 결정해야 한다. b가 크면 패스 수가 줄어든다. 하지만 한 패스에서 처리해야 할 빈의 개수가 2^b로 늘어난다.

실험적으로 보면 4-bit 라딕스가 가장 자주 채택된다. 패스가 8번으로 많이 줄어들면서도 빈 16개의 스캔과 히스토그램이 셰어드 메모리에 무리 없이 들어간다. CUB 라이브러리도 기본값이 4 또는 8 비트 근방이다. 8-bit는 패스를 4번까지 줄여 주지만 256개의 빈 카운트와 256-way scan이 부담이다. 1-bit는 단순하지만 패스가 32번이라 글로벌 메모리 트래픽이 32번 왕복한다.

13.6 코얼레싱을 위한 코어스닝

11장과 12장에서 봤던 스레드 코어스닝(thread coarsening)이 여기서도 약을 바른다. 한 스레드가 키 한 개만 다루면 매번 새로 데이터를 읽어 오는 오버헤드가 누적된다. 한 스레드가 키를 4개 또는 8개씩 처리하도록 묶으면 셰어드 메모리 재사용이 늘고, 레지스터 위에서 작은 정렬을 할 수 있게 되어 셰어드/글로벌 트래픽도 함께 줄어든다.

대표적인 코어스닝 패턴은 다음과 같다. 한 스레드가 4개를 들고 와서 레지스터에서 정렬한 뒤, 셰어드 메모리에 모아 블록 단위 라딕스 패스를 수행한다. 이때 한 블록이 처리하는 키의 양은 (스레드 수) × 4가 되어 처리량이 극대화된다.

// 한 스레드가 K개 키를 들고 시작
unsigned int reg[K];
// vectorized load: float4/uint4로 한꺼번에
uint4 v = reinterpret_cast<const uint4*>(in)[i];
reg[0] = v.x; reg[1] = v.y; reg[2] = v.z; reg[3] = v.w;

// 레지스터 안에서 작은 라딕스 정렬 (K = 4)
sort4InRegisters(reg, bit);

// 셰어드에 적재 → 블록 단위 패스
// ...

13.7 병렬 머지소트: 12장을 다시 꺼내 들기

비교 기반 쪽으로 옮겨 보자. 머지소트는 분할 정복으로 작은 부분을 정렬해 가며 병합한다. 직렬에서는 재귀적으로 반을 나누지만, 병렬에서는 거꾸로 가는 bottom-up 구성이 자연스럽다.

1. 전체 배열을 작은 타일(예: 1024 원소)로 나누고, 각 타일을 한 블록 안에서 정렬한다. 작은 타일이라 셰어드 메모리에 들어가니 블록 내 머지소트나 비토닉 정렬을 쓰면 된다.
2. 크기 T의 정렬된 두 조각을 병합해 크기 2T의 정렬된 조각을 만든다. 이 단계에서 12장의 병렬 머지(parallel merge)를 그대로 호출한다. co-rank 함수로 두 입력의 분할점을 찾아 여러 블록이 협력해 한 큰 머지를 처리한다.
3. 크기를 두 배씩 키워 가며 step ∈ {T, 2T, 4T, …, N/2}로 반복한다. 단계 수는 log₂(N/T)이다.

그림 13.2 — bottom-up 병렬 머지소트의 단계 구조. 각 단계는 12장의 병렬 merge를 활용한다.

// 호스트 측 의사코드
void parallelMergesort(int* d_data, int N) {
  const int T = 1024;
  // 1단계: 타일별 블록 내 정렬
  blockSortKernel<<<(N+T-1)/T, T>>>(d_data, N);

  // 2단계: bottom-up 병합
  int step = T;
  while (step < N) {
    parallelMergeKernel<<<..., ...>>>(d_in, d_out, N, step);
    std::swap(d_in, d_out);
    step *= 2;
  }
}

비교 정렬이라 N log N 하한을 그대로 따른다. 그래서 라딕스에 비해 보통 1.5~2배 느리다. 하지만 부동소수, 사용자 정의 키, 정렬과 동시에 다른 키-값 페어 정렬 등 유연성이 필요할 때 머지소트가 빛난다.

13.8 그 밖의 병렬 정렬

두 챔피언 외에도 GPU 정렬에는 흥미로운 친구들이 있다.

비토닉 소트(bitonic sort)는 정렬 네트워크(sorting network)의 대표주자다. N 원소를 정렬하기 위해 (log₂ N)²개의 비교-교환 단계를 정해진 순서로 적용한다. 모든 비교가 데이터에 무관해서 분기가 없고, 같은 워프 안에서 실행 흐름이 완벽히 동기화된다는 점이 매력이다. 작은 N(블록 안에 들어오는 크기)에서 매우 빠르다. 하지만 큰 N에서는 단계 수가 (log N)²로 늘어 머지소트보다 느려진다. 그래서 실무에서는 블록 내 정렬의 부속품으로 자주 쓰인다.

샘플 소트(sample sort)는 분할 정복의 또 다른 방향이다. 입력에서 작은 표본을 뽑아 분위수(splitter)를 정하고, 모든 원소를 그 분위수가 만드는 버킷으로 분배한 뒤 각 버킷을 독립적으로 정렬한다. 한번에 여러 갈래로 갈라지므로 거대한 데이터를 메모리 분산 환경에 펼치기 좋다. multi-GPU나 외부 정렬에서 자주 채용된다. 단, splitter 선택이 잘못되면 부하 불균형으로 망할 수 있다.

13.9 정리

정렬은 GPU의 모든 무기를 모아 쓰는 종합 격투기다. 비트 마스크, 스캔, 스캐터, 코얼레싱 트릭, 코어스닝, 그리고 12장의 머지까지. 정수 키라면 라딕스가 정답에 가깝고, 부동소수나 임의 키라면 머지소트가 안전한 선택이다. 어느 쪽이든 잘 짠 GPU 정렬은 CPU 대비 10~30배 처리량을 보여주며, 데이터베이스, 그래프 알고리즘, 충돌 검출, 머신러닝 전처리 등 거의 모든 후속 워크로드의 시드가 된다.