CUDA 활용 팁: 벡터화된 메모리 접근으로 성능 향상하기

Reading Time: 4 minutes

CUDA 커널의 성능은 종종 메모리 대역폭에 의해 제한되며, 최신 하드웨어일수록 연산 성능 대비 대역폭이 상대적으로 부족합니다. 따라서 코드에서 대역폭 병목 현상을 완화하기 위한 최적화가 매우 중요합니다. 이번 포스트에서는 CUDA C++에서 벡터 단위 로드 및 저장 방식을 활용해 대역폭 활용도를 높이고, 실행 명령 수를 줄이는 방법을 소개합니다.

먼저 아래의 간단한 메모리 복사 커널 예제부터 살펴보겠습니다.

__global__ void device_copy_scalar_kernel(int* d_in, int* d_out, int N) { 
  int idx = blockIdx.x * blockDim.x + threadIdx.x; 
  for (int i = idx; i < N; i += blockDim.x * gridDim.x) { 
    d_out[i] = d_in[i]; 
  } 
} 
 
void device_copy_scalar(int* d_in, int* d_out, int N) 
{ 
  int threads = 256; 
  int blocks = min((N + threads-1) / threads, MAX_BLOCKS);  
  device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N); 
}

이 코드에서는 이전 CUDA 팁 게시물에서 소개한 grid-stride 루프를 사용하고 있습니다. 그림 1은 복사 크기에 따른 커널의 처리량(GB/s)을 보여줍니다

CUDA Toolkit에 포함된 도구인 cuobjdump를 사용하면 해당 커널의 어셈블리 코드를 확인할 수 있습니다.

%> cuobjdump -sass executable

스칼라 복사 커널의 본문에서 생성된 SASS는 다음과 같은 명령어를 포함합니다:

...
LDG.E R3, desc[UR6][R2.64] ;
...
STG.E desc[UR6][R4.64], R3 ; 
...

LDG.E와 STG.E 명령어는 각각 글로벌 메모리에서 32비트를 읽고 쓰는 역할을 합니다.

이 연산의 성능은 LDG.E.{64,128} 및 STG.E.{64,128}와 같은 벡터화된 로드/스토어 명령어를 사용해 향상시킬 수 있습니다. 이 명령어들은 동일한 동작을 수행하지만, 한 번에 64비트 또는 128비트 단위로 데이터를 처리합니다. 벡터화된 로드를 사용하면 명령어 수를 줄이고, 지연 시간을 단축하며, 대역폭 활용률도 높일 수 있습니다.

벡터화된 로드를 가장 쉽게 사용하는 방법은 CUDA C++ 표준 헤더에 정의된 벡터 자료형(int2, int4, float2, float4 등)을 활용하는 것입니다. 이 자료형은 여러 값을 하나의 데이터 단위로 묶어 표현하며, C++에서는 reinterpret_cast<int2*>(d_in)과 같이 포인터 타입을 변환해 사용할 수 있습니다. 이 포인터는 두 개의 int 값을 하나의 단위로 처리합니다. C99에서는 (int2*)(d_in)과 같은 형 변환 연산자를 사용할 수 있습니다.

이러한 포인터를 역참조하면 컴파일러는 자동으로 벡터화된 명령어를 생성합니다.

int2* int2Ptr = reinterpret_cast<int2*>(d_in); 
int2 data = int2Ptr[0]; // Loads the first two int values as one int2

하지만 한 가지 중요한 주의사항이 있습니다. 벡터화된 명령어는 **정렬(alignment)**된 데이터를 요구합니다. 장치 메모리는 데이터 타입 크기의 배수로 자동 정렬되지만, 포인터에 오프셋을 줄 경우 해당 오프셋 역시 정렬 기준을 만족해야 합니다. 예를 들어, reinterpret_cast<int2*>(d_in+1)은 잘못된 예시입니다. d_in+1은 int2의 크기인 8바이트 단위로 정렬되어 있지 않기 때문입니다.

반면, reinterpret_cast<int2*>(d_in+2)처럼 정렬된 오프셋을 사용하는 경우는 안전하게 사용할 수 있습니다.

또한, 구조체의 크기가 2의 거듭제곱인 경우에도 벡터화된 로드를 생성할 수 있습니다.

struct Foo {int a, int b, double c}; // 16 bytes in size
Foo *x, *y;
…
x[i]=y[i];

2의 거듭제곱이 아닌 크기를 가진 구조체는 메모리 정렬이 비효율적으로 이루어질 수 있으며, 이로 인해 컴파일러가 자동으로 패딩을 추가해 정렬을 맞추는 경우가 발생할 수 있습니다.

이제 벡터화된 명령어를 생성하는 방법을 살펴봤으니, 메모리 복사 커널을 벡터 로드를 활용하는 방식으로 수정해보겠습니다.

__global__ void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
    reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
  }
 
  // in only one thread, process final element (if there is one)
  if (idx==N/2 && N%2==1)
    d_out[N-1] = d_in[N-1];
}
 
void device_copy_vector2(int* d_in, int* d_out, int n) {
  threads = 256; 
  blocks = min((N/2 + threads-1) / threads, MAX_BLOCKS); 
 
  device_copy_vector2_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

이 커널은 몇 가지 부분만 수정되었습니다. 첫째, 한 번의 반복에서 두 개의 요소를 처리하므로 루프는 이제 전체 N의 절반 횟수만 실행됩니다. 둘째, 복사 과정에서 앞서 설명한 형 변환 기법을 사용합니다. 셋째, N이 2로 나누어떨어지지 않는 경우를 대비해 남은 요소를 별도로 처리합니다. 마지막으로, 스칼라 방식의 커널보다 절반의 스레드만 실행합니다.

SASS를 확인해 보면 다음과 같은 변화가 나타납니다:

...
LDG.E.64 R2, desc[UR4][R2.64] ; 
...
STG.E.64 desc[UR4][R4.64], R2 ; 
...

이제 컴파일러가 LDG.E.64 및 STG.E.64 명령어를 생성하는 것을 확인할 수 있습니다. 그 외의 명령어들은 이전과 동일하지만, 루프가 N/2번만 실행되므로 전체 명령어 실행 수는 절반으로 줄어듭니다. 명령어 수가 2배 줄어드는 이러한 최적화는 명령어 수에 의해 병목이 생기거나, 지연 시간에 민감한 커널에서 특히 중요합니다.

또한, 벡터 크기를 4로 확장한 vector4 버전의 복사 커널도 작성할 수 있습니다.

__global__ void device_copy_vector4_kernel(int* d_in, int* d_out, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  for(int i = idx; i < N/4; i += blockDim.x * gridDim.x) {
    reinterpret_cast<int4*>(d_out)[i] = reinterpret_cast<int4*>(d_in)[i];
  }
 
  // in only one thread, process final elements (if there are any)
  int remainder = N%4;
  if (idx==N/4 && remainder!=0) {
    while(remainder) {
      int idx = N - remainder--;
      d_out[idx] = d_in[idx];
    }
  }
}
 
void device_copy_vector4(int* d_in, int* d_out, int N) {
  int threads = 256;
  int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);
 
  device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

이에 따라 생성되는 SASS에서도 다음과 같은 변화가 나타납니다:

...
LDG.E.128 R4, desc[UR4][R4.64] ;  
...
STG.E.128 desc[UR4][R8.64], R4 ;         
...

여기서는 LDG.E.128과 STG.E.128 명령어가 생성된 것을 확인할 수 있습니다. 이 버전의 코드는 명령어 수를 기존보다 4분의 1로 줄였습니다. 세 가지 커널의 전체 성능 비교는 그림 2에서 확인할 수 있습니다.

거의 대부분의 경우, 벡터화된 로드는 스칼라 로드보다 더 나은 선택입니다. 하지만 주의할 점도 있습니다. 벡터화된 로드를 사용하면 레지스터 사용량이 증가하고, 그로 인해 전체 병렬성이 감소할 수 있습니다. 따라서 이미 레지스터 사용에 제한이 있거나 병렬성이 낮은 커널이라면 스칼라 로드를 유지하는 것이 더 나을 수 있습니다. 또한 앞서 설명했듯이, 포인터가 정렬되어 있지 않거나 데이터 타입의 크기가 2의 거듭제곱이 아닐 경우 벡터화된 로드를 사용할 수 없습니다.

벡터화된 로드는 대역폭 활용을 높이고, 명령어 수와 지연 시간을 줄여주는 핵심 CUDA 최적화 기법입니다. 이번 포스트에서는 기존 커널에 벡터 로드를 간단히 적용하는 방법을 살펴봤으며, 적용 가능하다면 꼭 고려해볼 만한 최적화 전략입니다.

이 블로그는 원래 2013년 12월 4일에 게시되었으며, 현재 GPU의 동작 방식에 맞춰 업데이트되었습니다.