CUDA 스레드 계층

스레드

CUDA 스레드 계층에서 가장 작은 단위는 스레드이다. 따라서 CUDA 연산을 수행하거나, 코어를 사용하는 기본 단위이다. 커널 호출 시, CUDA 커널 코드는 모든 스레드에 공유된다. 각 스레드는 커널을 독립적으로 실행한다.

워프

CUDA 스레드 계층의 두 번째 계층은 워프이다. 워프는 32개 스레드를 하나로 묶은 단위이다. 중요한 점은 워프는 디바이스에서 하나의 제어 장치에 의해 제어된다. GPU의 SIMT 구조에서 멀티 스레드 단위가 바로 워프이다. 이 말은 1개의 명령어에 의해 32개 스레드가 동시에 움직이는 것을 의미한다.

블록

블록은 워프보다 상위 개념으로, 워프들의 집합이다. 하나의 블록에 포함된 각 스레드는 고유한 스레드 인덱스를 가진다. 이는 동일한 블록 안에서는 동일한 인덱스의 스레드가 없다는 것을 의미한다. 반면 서로 다른 블록 간 스레드들은 같은 인덱스를 가질 수 있다. 따라서 특정 스레드를 가리키려면, 블록 번호까지 알아야 한다. 이때 블록 내 스레드는 1차원, 2차원, 3차원 형태로 배치 가능하다.

그리드

여러 개의 블록을 하나로 묶은 것은 그리드이다. 하나의 그리드에 포함된 블록은 자신만의 고유한 블록 인덱스를 가진다. 블록과 마찬가지로 그리드 내 블록 또한 1차원, 2차원, 3차원 형태로 배치 가능하다. 하나의 커널은 하나의 그리드와 1:1 매핑된다. 따라서 커널이 호출되면 그리드가 생성되고, 이는 커널을 수행하는 스레드를 생성한 것과 같다.

CUDA 스레드 계층을 위한 내장 변수

내장 변수

  • 그리드 내 블록은 1~3차원 형태이다.
  • 블록 내 스레드는 1~3차원 형태이다.

따라서 스레드들은 각자 처리한 데이터가 무엇인지, 어떤 블록인지 등 블록 내 인덱스를 알아야 한다. CUDA는 현재 그리드 및 블록의 형태와 스레드가 속한 블록 번호 및 스레드 번호를 확인하는 내장 변수를 제공한다. 내장 변수 값은 커널 실행 시, 결정되며 수정할 수 없다.

gridDim

그리드의 형태 정보를 담고 있는 구조체형 내장 변수이다. x, y, z 멤버 변수로 1, 2, 3차원 정보를 표현한다.

gridDim.x -> 1차원
gridDim.y -> 2차원
gridDim.z -> 3차원

blockIdx

현재 스레드의 블록 정보를 담고 있는 구조체형 내장 변수이다. 각 차원은 girdDim 방식과 같다.

threadIdx.x -> 1차원
threadIdx.y -> 2차원
threadIdx.z -> 3차원

blockDim

블록의 형태 정보를 담고 있는 구조체형 내장 변수이다. 커널이 실행될 때, 그리드 및 블록 형태가 결정된다. 한 그리드 내, 블록은 모두 동일한 형태이므로 blockDim은 그리드 내 모든 스레드가 공유한다.

threadIdx

블록 내에서 현재 스레드가 부여받은 인덱스를 담고 있는 구조체형 내장 변수이다. 한 블록 내 스레드들은 모두 고유한 인덱스를 갖지만, 블록 형태에 따라 달라진다.

스레드 인덱스와 워프의 구성

  • 워프는 연속된 32개의 스레드로 구성된다.
  • 커널 성능에 영향을 미치는 요소인 메모리 접근 패턴을 이해해야 한다.

그리드 및 블록의 최대 크기 제한

그리드의 크기

그리드 크기에서 y, z 차원은 65,535로 제한된다. (x는 사실상 제한 없다.)

블록의 크기

블록 크기에서 x, y 차원은 1,024가 최대 크기이며, z 차원은 64이다.
또한 블록 하나가 가질 수 있는 최대 스레드 수는 1,024이다.

CUDA 스레드 구조와 커널 호출

스레드 레이아웃과 커널 호출

# 스레드의 배치 형태를 지정하는 방법

__global__ Kernel {
    ...
}

Kernel <<<그리드의 형태, 블록의 형태>>>()

실제로 Kerenl«<1, n»> 은 그리드 형태가 (1, 1, 1)이며 블록 형태는 (n, 1, 1)이다. 따라서 좀 더 정확하게 기술하기 위해 (x, y, z) 정보를 담는 구조체를 사용할 수 있다.

// 1. dim3 구조체로 (x, y, z)를 선언한다.
// 2. dimGrid, dimBlock을 커널 레이아웃 형태로 전달한다.
// 3. dimGrid는 블록의 레이아웃을 결정한다.
// 4. dimBlock은 스레드의 레이아웃을 결정한다.

struct dim3(int x, int y, int z) {
    int x = x;
    int y = y;
    int z = z;
}

dim3 dimGrid(4, 1, 1);
dim3 dimBlock(8, 1, 1);

Kernel <<<dimGrid, dimBlock>>>();

위 코드에서 블록 인덱스와 스레드 인덱스는 다음과 같다.
blockIdx.x: 0, 1, 2, 3; // dimGrd(4, 1, 1)
threadIdx.x: 0, 1, 2, …, 7 // dimBlock(8, 1, 1)

특징

  • (x, y, z) 블록에서 스레드 수는 threadIdx.x * threadIdx.y * threadIdx.z 이다.
  • 블록 내 스레드는 고유한 인덱스이다. 따라서 블록 수 만큼 같은 스레드 인덱스들이 있다. 동일한 스레드 인덱스들을 구분하기 위해서는 블록 인덱스가 필요하다.

큰 벡터의 합을 연산하는 CUDA 커널 (1)

이전 벡터 연산 코드에서 무작정 벡터 차원을 늘리면, 커널 연산이 느려진다. 그렇다고 (예를 들어 백터 차원이 1024보다 큰 경우처럼) 벡터 차원에 맞게 스레드 수를 늘리면 블록 하나에서 가능한 최대 스레드 수가 1024개 이므로 아래 코드는 커널을 호출할 수 없다.

NUM_DATA = 4224;
addVec <<<1, NUM_DATA>>>();

떠올릴 수 있는 방법은?

  1. 블록 형태를 늘린다. (스레드 인덱스의 y, z를 추가하기)
  2. 블록을 여러개 사용한다.

방법 1.은 블록이 가지는 스레드 수의 제한으로 똑같이 사용할 수 없다 따라서 블록 형태를 다시 조정해야 한다 방법 2.로 블록 수를 늘려 addVec 커널을 호출하면 다음과 같다.

NUM_DATA = 4224;
addVec <<<4224/1024, 1024>>>();

또 발생하는 문제는 블록, 스레드의 인덱스는 정수만 가질 수 있다. 그래서 4224/1024 처럼 소수점 결과가 나오는 인덱스를 커널에 전달할 수 없다. 따라서 올림 처리하여 수정한다.

NUM_DATA = 4224;
addVec <<<ceil(4224/1024), 1024>>>();

위처럼 블록 레이아웃을 전달하면, 4224개 벡터 차원에 대해서 5개 블록의 커널을 호출할 수 있다. 이때 인덱스 0부터 4까지 블록은 Full utilization으로 인덱스마다 데이터가 있지만, 5번 블록은 그렇지 않다. 이의 처리 방법은 스레드 인덱스를 처리하는 방법을 고민해야 한다.

하지만 블록 수를 늘려 커널을 호출하여도 실행은 되지만 올바른 결과가 나오지 않는다. 1024번째 원소 이후로는 잘못된 결과가 나온다. 기존 코드는 단일 블록의 1024개 스레드에서만 데이터 연산을 고려하였기 때문이다. (즉, 각 스레드가 속한 다른 블록은 무시)

// 단일 블록의 스레드에 대해서만 벡터 합을 연산

__global__ void vecAdd(int* _a, int* _b, int* _c) {
    int tID = threadIdx.x;
    _c[tID] = _a[tID] + _b[tID];
}

구체적으로 기존 코드의 커널은 threadIdx 내장 변수를 통해 x 차원 스레드 인덱스를 갖고 온다. 이 코드에서는 스레드 인덱스와 동일한 벡터 인덱스의 원소간 합을 연산하지만, 다른 블록의 동일한 인덱스를 가진 스레드는 고려하지 않은 것을 알 수 있다.