CUDA Warp Divergence의 동작 원리와 Independent Thread Scheduling

1. 배경

 

GPU는 수많은 스레드를 병렬로 실행함으로써 높은 처리량을 얻는다

 

CUDA에서 스레드는 개별적으로 작성되지만 GPU에서는 Warp를 기본 단위로 명령어를 실행한다

NVIDIA GPU의 Warp는 32개의 스레드로 구성되며, SIMT 방식에 따라 Warp에 속한 스레드들이 동일한 명령을 수행한다

 

예를 들어 다음과 같이 각 스레드가 배열의 서로 다른 원소를 계산한다고 해보자

__global__ void add(int* data) {
    int tid = threadIndex.x;
    data[tid] += 1;
}

 

각 스레드가 접근하는 데이터는 다르지만 수행하는 명령어는 동일하기 때문에,

하나의 Warp에 속한 스레드들을 효율적으로 병렬 실행할 수 있다

 

그렇다면 다음과 같이 각 스레드마다 서로 다른 분기로 진입한다면 어떻게 될까?

if (threadIdx.x < 16) {
    A();
} else {
    B();
}

 

하나의 Warp에 포함된 32개의 스레드 중 절반은 A를 실행하고 나머지 절반은 B를 실행해야 한다

하지만 하나의 Warp가 동시에 서로 다른 두 명령어를 실행할 수는 없다

 

이처럼 하나의 Warp에 속한 스레드들이 서로 다른 실행 경로를 가지는 상황을 Warp Divergence라고 한다

 

이번 글에서는 Warp Divergence가 발생했을 때 GPU가 서로 다른 실행 경로를 어떻게 처리하는지 살펴보고,

Volta 아키텍처부터 도입된 Independent Thread Scheduling이 어떤 한계를 해결했는지 알아보고자 한다

 

2. Warp Divergence와 Active Mask

 

Volta 이전 NVIDIA GPU에서 하나의 Warp는 32개 스레드가 공유하는 PC(Program Counter)를 가지고 있었다

 

PC는 다음에 실행할 명령어의 위치를 나타낸다

따라서 하나의 Warp에 속한 스레드들은 기본적으로 동일한 명령어 위치를 따라 실행한다

 

하지만 실제 프로그램에서는  if ,  switch 와 같은 분기를 통해 서로 다른 제어 흐름을 가질 수 있다

if (threadIdx.x < 16) {
    A();
} else {
    B();
}

 

예시 코드와 같은 환경의 경우 스레드 0 ~ 15는 A를 실행해야 하고, 16 ~ 31은 B를 실행해야 한다

이때 GPU는 두 경로를 동시에 실행하는 대신 Active Mask를 이용해 현재 명령어에 참여할 스레드를 선택한다

 

 

A 경로를 실행하는 동안에 Active Mask는 위와 같다

이때 활성 비트로 표시된 스레드를 Active Thread라고 하며, 그렇지 않은 스레드를 Inactive Thread라고 한다

 

따라서 Warp Divergence가 발생하면 스레드 자체가 사라지거나 종료되는 것은 아니지만,

하나의 명령어를 실행할 때 사용할 수 있는 실행 자원의 일부가 활용되지 않을 수 있다

 

결과적으로 동일한 Warp의 스레드들이 같은 실행 경로를 가지는 경우보다 처리 효율이 낮아질 수 있다

 

3. Divergent Path에서 Stall이 발생하면 어떻게 될까?

 

Warp Divergence가 발생하면 하나의 Warp에 속한 스레드들은 서로 다른 Branch Path로 나뉘어 실행된다

 

Volta 이전 GPU에서는 하나의 Warp가 하나의 PC와 Active Mask를 공유했기 때문에,

서로 다른 Branch Path의 스레드를 독립적으로 스케줄링할 수 없었다

 

앞선 예시에서 A 경로의 함수는 빠르게 수행을 마치지만,

B 경로의 함수는 메모리 로드에 지연이 생긴다면 Warp는 다음 명령어를 바로 실행할 수 없게 된다

 

문제는 A의 스레드들이 당장 실행 가능한 연산을 가지고 있더라도,

같은 Warp의 다른 Path를 독립적으로 선택해서 실행할 수 없다는 점이다

 

 

Volta 이전에는 하나의 Warp가 공유된 실행 상태를 기반으로 동작했기 때문에,

Divergence 된 스레드들은 Reconvergence 이전까지 서로 독립적인 실행 흐름을 가지기 어려웠다

 

단순히 일부 스레드의 실행 자원을 충분히 활용하지 못하는 경우라면 이는 성능 저하에 그칠 수 있다

하지만 현재 실행 중인 스레드가 다른 Branch Path에 속한 스레드의 실행 결과를 기다린다면 어떻게 될까?

 

4. Divergent Thread 간 의존성과 Deadlock

 

Warp 내부의 스레드가 항상 서로 독립적인 연산만 수행하는 것은 아니다

 

경우에 따라 동일한 Warp에 속한 스레드가 다른 스레드에서 생성한 데이터를 사용하거나,

다른 스레드의 작업이 완료되기를 기다려야 할 수도 있다

 

예를 들어 다음과 같이 서로 다른 Branch Path 사이에 데이터 의존성이 존재한다고 가정해보자

if (condition) {
    waitForData(); // 다른 Thread의 결과를 기다리는 spin-wait
    consume();
} else {
    produce();
}

 

여기서  consume() 을 수행하는 스레드는  produce() 를 수행하는 다른 스레드가 데이터를 생성해야 다음 연산을 진행할 수 있다

 

만약  produce() 가 실행되는 Path가 먼저 처리된다면 데이터가 준비된 이후이므로  consume()  Path도 정상 실행이 가능하다

하지만 반대로  consume()  Path가 먼저 실행되고, 데이터가 생성될 때 까지 실행을 종료하지 않는다면 문제가 발생할 수 있다

 

Volta 이전 GPU에서는 Divergence된 스레드가 하나의 Warp 단위 실행 상태를 공유했기 때문에,

현재 Path가 다른 Path의 실행을 기다리는 동안 필요한 스레드만 독립적으로 진행시킬 수 없었다

 

즉 현재 실행 중인 Path는 다른 Path가 실행되기를 기다리고 있지만,

해당 Path의 스레드들은 현재 실행 흐름과 독립적으로 진행할 수 없기에 Warp 전체가 진행할 수 없는 상황이 발생할 수 있다

 

---

Threads from the same warp in divergent regions or different states of execution cannot signal each other or exchange data, and algorithms requiring fine-grained sharing of data guarded by locks or mutexes can easily lead to deadlock, depending on which warp the contending threads come from. (출처)

---

 

NVIDIA 공식 문서에서도 Volta 이전 실행 모델에 대해 Divergence 된 스레드들이 서로 신호를 보내거나 데이터 교환이 불가하며,

Lock이나 Mutex를 통해 세밀하게 데이터를 공유하는 알고리즘에서는 이러한 특성이 Deadlock과 이어질 수 있다고 설명한다

 

이러한 Warp 단위 실행의 제약을 완화하기 위해 Volta 아키텍처부터 Independent Thread Scheduling이 도입되었다

 

5. Volta와 Independent Thread Scheduling

 

기존에는 하나의 Warp가 하나의 PC를 공유하며 실행되었지만,

Volta부터는 Warp에 속한 각 스레드가 개별적인 PC와 Call Stack 등의 실행 상태를 가질 수 있게 되었다

 

이를 통해 동일한 Warp에 속한 스레드라도 서로 다른 실행 위치에 머무를 수 있으며,

현재 진행 가능한 스레드들을 다시 묶어 실행하는 보다 유연한 Scheduling이 가능해졌다

 

앞선 예시를 다시 생각해보자

if (condition) {
    waitForData(); // 다른 Thread의 결과를 기다리는 spin-wait
    consume();
} else {
    produce();
}

 

기존 방식에서는  waitForData() 를 수행하는 스레드가 다른 Branch Path의 실행을 기다리더라도,

같은 Warp의 다른 스레드를 독립적인 실행 흐름으로 진행시키기 어려웠다

 

하지만 Independent Thread Scheduling에서는 대기 중인 스레드와 별개로

실행 가능한 스레드가 다른 위치의 명령어를 진행할 수 있게 되었다

 

다만 Independent Thread Scheduling이 Warp Divergence 자체의 비용을 제거하는 것은 아니다

 

서로 다른 Branch Path를 실행하는 도중에는 여전히 각 명령어에 참여하지 않는 스레드가 존재할 수 있으며,

따라서 Divergence로 인한 실행 효율 저하는 그대로 발생할 수 있다

 

Independent Thread Scheduling의 핵심은 Divergence를 없애는 것이 아니라,

Divergence된 스레드들의 실행 상태를 보다 독립적으로 관리할 수 있도록 Warp의 Scheduling 모델을 확장한 것에 가깝다

 

6. Warp Divergence 최적화와 동기화

 

Independent Thread Scheduling이 도입되면서 하나의 Warp에 속한 스레드들이 서로 다른 상태를 가질 수 있지만,

CUDA의 기본 실행 모델이 SIMT라는 점은 바뀌지 않았다

 

Warp는 여전히 동일한 명령어를 여러 스레드가 함께 실행하며,

일부 스레드만 특정 Branch Path를 따라가는 경우 참여하지 않는 스레드는 Masking된다

 

---

Utilization of the GPU is maximized when threads within a warp follow the same control flow path. (출처)

---

 

NVIDIA 역시 Warp 내부의 스레드가 동일한 Control Flow Path를 따를 때 GPU 활용도가 가장 높다고 설명한다

따라서 Independent Thread Scheduling이 도입된 이후에도 서로 다른 Branch Path로 나뉘지 않도록 구성하는 것이 유리하다

 

한편 Independent Thread Scheduling은 기존의 Warp Synchronous Programming에도 영향을 준다

예를 들어 Shared Memory를 통해 Warp 내부의 스레드가 값을 전달한다고 해보자

__global__ void example(int *result) {
    __shared__ int shared[32];

    int lane = threadIdx.x % 32;

    if (lane < 16) {
        shared[lane] = lane + 1;
    } else {
        shared[lane] = lane * 2;
    }

    int partner = (lane + 16) % 32;
    result[threadIdx.x] = shared[lane] + shared[partner];
}

 

이 코드에서 스레드 0 ~ 15는 덧셈을 수행하고 스레드 16 ~ 31은 곱셈을 수행하기 때문에 서로 다른 Branch Path가 만들어진다

이후 각 스레드는 자신의 계산 결과 뿐만 아니라 반대쪽 브랜치에 속한 스레드가 계산한 값도 사용한다

 

Volta 이전에는 Divergence된 Warp가 다시 합류하는 지점을 기준으로 실행되는 특성에 기대어,

같은 Warp에 속한 스레드들이 분기 이후에는 함께 진행할 것이라고 가정하는 코드가 사용되기도 했다

 

하지만 Volta 이후에는 서로 다른 Branch Path가 서로 다른 실행 위치를 유지할 수 있으며,

하나의 스레드가 분기 이후의 코드까지 진행했다고 해서 다른 Branch의 스레드 역시 계산을 마쳤다고 가정할 수 없다

 

따라서 위 코드에서는  shared[partner] 에 대한 Write가 완료되기 전에 해당 값을 읽는 문제가 발생할 수 있다

 

Warp 내부 스레드 사이에 이러한 데이터 의존성이 존재한다면  __syncwarp() 를 이용해 명시적으로 동기화해야 한다

__global__ void example(int *result) {
    __shared__ int shared[32];

    int lane = threadIdx.x % 32;

    if (lane < 16) {
        shared[lane] = lane + 1;
    } else {
        shared[lane] = lane * 2;
    }
    
    __syncwarp();

    int partner = (lane + 16) % 32;
    result[threadIdx.x] = shared[lane] + shared[partner];
}

 

 __syncwarp() 는 스레드들이 해당 지점에 도달할 때 까지 기다려 메모리 동기화를 제공한다

따라서  __syncwarp()  이후에는 각 스레드가 다른 Branch에 속한  partner 의 계산 결과까지 안전하게 사용할 수 있다

 

7. 마무리하며

 

CUDA 코드는 스레드 단위로 작성되지만,

실제 동작과 성능을 이해하기 위해서는 해당 코드가 Warp 단위에서 어떻게 실행되는지 함께 바라볼 필요가 있다

 

특히 단순해 보이는 분기문 하나도 SIMT 실행 모델과 Scheduling 방식에 따라 예상과 다른 실행 특성을 가질 수 있다

 

이번 내용을 정리하면서 CUDA에서는 코드를 작성하는 것 뿐만 아니라,

GPU가 해당 코드를 어떤 단위와 방식으로 실행하는지를 이해하는 것이 중요하다는 점을 다시 확인할 수 있었다