(실습) Nvidia의 MergeSort 예제 분석 - 1

(실습) CUDA MergeSort 구현 (순진한 버전) (tistory.com)

(연습) CUDA MergeSort 구현(순진한 버전)

CUDA에서 정렬을 구현할 때 가장 먼저 떠오르는 생각은 우리가 알고 있는 알고리즘(퀵 정렬, 병합 정렬, 버블 정렬 등)을 병렬로 전환하는 것입니다. 그런 다음 생각

hotstone.tistory.com

'추천 관련글,

과거에는 병합 정렬이 CPU에서와 동일하게 CUDA에서 구현되었습니다. 하지만 당연히 std::qsort보다 훨씬 느린 성능을 보였기 때문에 CUDA에서 정렬을 구현하는 방법에 대한 Nvidia의 예제를 참고했습니다. 오늘은 제가 분석한 Nvidia MergeSort에 대해 이야기 해보려고 합니다.

이전의 순진한 버전의 병합 정렬을 구현할 때 가장 성능에 민감한 부분은 위 이미지와 같이 정렬된 두 개의 하위 배열을 가져와 병합하는 부분이었습니다. 정렬된 2개의 배열에서 값을 하나씩 빼서 값을 비교하는 방식으로 값을 정렬해야 했기 때문이다. 단계 크기가 증가함에 따라 병합에 포함되지 않은 스레드 수가 2의 기하급수적으로 증가하여 성능이 저하되었습니다. 그리고 이러한 값을 비교해야 했던 이유는 정렬된 두 배열의 값이 머지할 배열의 어느 위치에 놓일지 모르기 때문이었습니다. 따라서 병합할 배열에서 두 개의 정렬된 배열의 각 값을 배치할 위치를 알고 있으면 병렬로 처리할 수 있습니다.

그리고 이것을 이미지로 표현하면 이렇게 됩니다.

(위의 배열은 인덱스를 나타냅니다.)

위의 그림을 보면 각각의 요소는 합쳐질 배열의 인덱스를 가지고 있고 그 값을 기준으로 정렬된 배열이 형성됩니다. 그리고 다음 코드는 전역 메모리 인덱스를 조정하고 공유 메모리를 선언하고 초기화합니다.

참고로 이때 d_SrcKey는 정렬할 값을 담고 있고 d_SrcVal은 인덱스 값을 담고 있다.

위 이미지의 예를 사용하여 상위 배열은 d_SrcVal, 하위 배열은 d_SrcKey라고 할 수 있습니다.

__global__ void mergeSortSharedKernel(uint *d_DstKey, uint *d_DstVal,
                                      uint *d_SrcKey, uint *d_SrcVal,
                                      uint arrayLength) {
  __shared__ uint s_key(SHARED_SIZE_LIMIT); // 1024, thread 개수 = 512
  __shared__ uint s_val(SHARED_SIZE_LIMIT);

  // global memory의 인덱스를 편하게 다루기 위해서 현재 인덱스 기준으로 바꿉니다.
  d_SrcKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
  d_SrcVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
  d_DstKey += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
  d_DstVal += blockIdx.x * SHARED_SIZE_LIMIT + threadIdx.x;
  
  // shared memory를 초기화 합니다.
  s_key(threadIdx.x + 0) = d_SrcKey(0);
  s_val(threadIdx.x + 0) = d_SrcVal(0);
  s_key(threadIdx.x + (SHARED_SIZE_LIMIT / 2)) =
      d_SrcKey((SHARED_SIZE_LIMIT / 2));
  s_val(threadIdx.x + (SHARED_SIZE_LIMIT / 2)) =
      d_SrcVal((SHARED_SIZE_LIMIT / 2));
      
      ...
}

이제부터 공유 메모리 초기화가 완료되면 stride를 증가시켜 하위 배열을 정렬합니다.

하지만 이 시점에서 앞서 말했듯이 병합할 배열의 인덱스를 찾아야 합니다. 그리고 두 개의 배열이 정렬된 상태이므로 성능을 위해 다른 배열에서 해당 값의 위치 값을 찾기 위해 이진 검색을 사용합니다.

그리고 현재 배열에 위치값(lPos)을 더하면 최종 정렬된 배열의 위치값을 얻을 수 있다. 마지막으로 최종 위치 값을 가져오고 해당 위치에서 키와 값을 업데이트합니다.

  // stride 값을 2배씩 증가해 가며, 인덱스를 갱신합니다.
  for (uint stride = 1; stride < arrayLength; stride <<= 1) {
    // 기준이 되는 인덱스 계산합니다.
    // e.g: stride가 4일 경우 threadIdx.x에 따라서 lPos는 0, 1, 2, 3, 0, 1, 2, 3, ... 으로 구성됩니다.
    // 따라서 threadIdx.x - lPos를 함으로써 stride별 기준이 되는 인덱스를 구할 수 있습니다.
    // e.g: 0, 0, 0, 0, 8, 8, 8, 8, 16, 16, ...
    uint lPos = threadIdx.x & (stride - 1);
    uint *baseKey = s_key + 2 * (threadIdx.x - lPos);
    uint *baseVal = s_val + 2 * (threadIdx.x - lPos);

    cg::sync(cta);
    // 현재의 key와 value를 저장합니다.
    uint keyA = baseKey(lPos + 0);
    uint valA = baseVal(lPos + 0);
    uint keyB = baseKey(lPos + stride);
    uint valB = baseVal(lPos + stride);
    
    // binary Search를 활용하여 값(key)에 대한 인덱스를 구합니다.
    uint posA =
        binarySearchExclusive<sortDir>(keyA, baseKey + stride, stride, stride) +
        lPos;
    uint posB =
        binarySearchInclusive<sortDir>(keyB, baseKey + 0, stride, stride) +
        lPos;

    cg::sync(cta);
    baseKey(posA) = keyA;
    baseVal(posA) = valA;
    baseKey(posB) = keyB;
    baseVal(posB) = valB;
  }

그러나 주의해야 할 상황이 있습니다. 두 하위 배열에 동일한 값이 있는 경우입니다. 이때 동일한 함수를 사용하여 위치를 찾는 함수를 사용하면 아래와 같이 인덱스와 동일한 위치값을 가지게 됩니다.

따라서 각 하위 배열에서 위치 값을 검색할 때 하나의 하위 배열에 값이 포함된 위치와 다른 하위 배열에 값이 포함되지 않은 위치를 검색하여 인덱스가 중복되는 것을 방지할 수 있습니다.

template <uint sortDir>
static inline __device__ uint binarySearchInclusive(uint val, uint *data,
                                                    uint L, uint stride) {
  if (L == 0) {
    return 0;
  }

  uint pos = 0;

  for (; stride > 0; stride >>= 1) {
    uint newPos = umin(pos + stride, L);

    // 이 함수는 val를 포함한 position값을 리턴합니다.
    if ((sortDir && (data(newPos - 1) <= val)) ||
        (!sortDir && (data(newPos - 1) >= val))) {
      pos = newPos;
    }
  }

  return pos;
}

template <uint sortDir>
static inline __device__ uint binarySearchExclusive(uint val, uint *data,
                                                    uint L, uint stride) {
  if (L == 0) {
    return 0;
  }

  uint pos = 0;

  for (; stride > 0; stride >>= 1) {
    uint newPos = umin(pos + stride, L);

    // 이 함수는 val를 제외한 position값을 리턴합니다.
    if ((sortDir && (data(newPos - 1) < val)) ||
        (!sortDir && (data(newPos - 1) > val))) {
      pos = newPos;
    }
  }

  return pos;
}

마지막으로 다음 코드와 같이 해당 커널 함수는 공유 메모리에 정렬된 값을 전역 메모리로 이동하여 작업을 완료합니다.

  cg::sync(cta);
  d_DstKey(0) = s_key(threadIdx.x + 0);
  d_DstVal(0) = s_val(threadIdx.x + 0);
  d_DstKey((SHARED_SIZE_LIMIT / 2)) =
      s_key(threadIdx.x + (SHARED_SIZE_LIMIT / 2));
  d_DstVal((SHARED_SIZE_LIMIT / 2)) =
      s_val(threadIdx.x + (SHARED_SIZE_LIMIT / 2));

그러나 아직 정렬이 완료되지 않았습니다.

지금까지 블록 크기(512) x 2 = 1024 블록에 대한 정렬은 공유 메모리를 사용하여 수행되었으며 전체 정렬을 위해서는 전역 메모리 수준에서 정렬을 수행해야 합니다.

그리고 전체 배열을 정렬할 때 값을 사용하기 때문에 이진 검색으로 위치 값을 검색할 때 키만 사용하고 값은 업데이트만 하고 사용하지 않았습니다.

그런 다음 다음 기사에서는 전체 어레이를 업데이트하는 프로세스에 대해 설명합니다.

전체 소스 코드는 다음 링크에서 확인할 수 있습니다.

v11.6의 cuda-samples/mergeSort.cu NVIDIA/cuda-samples(github.com)

GitHub – NVIDIA/cuda-samples: CUDA 툴킷의 기능을 시연하는 CUDA 개발자용 샘플

CUDA 툴킷의 기능을 시연하는 CUDA 개발자 샘플 – GitHub – NVIDIA/cuda-samples: CUDA 툴킷의 기능을 시연하는 CUDA 개발자 샘플

github.com