개요

워프 스케줄링(Warp Scheduling)은 GPU에서 중요한 역할을 하는 개념이다. GPU는 대량의 쓰레드를 동시에 처리하는 능력을 가지며, 이러한 병렬 처리 능력을 최적화하기 위해 워프라는 단위를 사용한다. 워프는 일반적으로 한 번에 실행되는 32개의 쓰레드로 구성된다. 이 단위는 실행 유닛이 수천 개의 쓰레드를 관리하고 실행하는 데 효율적이도록 설계되었다.

워프의 개념

워프는 일반적으로 32개의 쓰레드를 포함하지만, 이는 GPU 아키텍처에 따라 다를 수 있다. 동일한 프로그램 카운터(PC)를 가지고 있으므로 같은 명령을 실행하게 되며, 동일한 명령어 스트림에서 작동한다. 워프는 SIMD(Single Instruction Multiple Data) 방식으로 동작하여 다수의 데이터를 하나의 명령어로 처리할 수 있다.

워프 스케줄링의 필요성

GPU에서 다수의 워프를 효율적으로 관리하기 위해 스케줄링이 필요하다. 이유는 다음과 같다:

1. 리소스 효율성: 하드웨어 리소스를 효율적으로 사용하기 위해 여러 워프를 동시에 실행함.
2. 메모리 대기 시간 은폐: 메모리 액세스 지연 문제를 숨기기 위해 다른 일을 병행함.
3. 연산 유닛 활용: 연산 유닛의 공백 시간을 최소화함.

워프 스케줄링 알고리즘

워프 스케줄링 알고리즘은 여러 종류가 있으며, 주로 다음 방법들이 사용된다:

1. 라운드 로빈(Round-Robin): 간단한 방법으로, 워프를 순차적으로 실행한다.
2. 우선순위 기반 스케줄링: 특정 조건에 따라 우선순위를 부여하고, 높은 우선순위의 워프를 먼저 실행한다.
3. 연결 스케줄링(Convergence Scheduling): 분기 점에서 워프들을 합쳐서 실행 효율을 높인다.

이러한 전략들이 어떻게 작동하는지 자세히 알아보겠다.

라운드 로빈 스케줄링

라운드 로빈 스케줄링은 각 워프를 순차적으로 실행하는 방식이다. 이 방식은 구현이 간단하고 공정한다. 모든 워프에게 동일한 기회를 제공한다.

예를 들어, 워프 A, B, C가 있을 때: - 첫 번째 사이클: 워프 A 실행 - 두 번째 사이클: 워프 B 실행 - 세 번째 사이클: 워프 C 실행

이 패턴을 반복한다.

장점

• 공평함
• 구현의 용이성

단점

• 특정 워프가 자주 블록될 경우 전체 성능 저하

우선순위 기반 스케줄링

우선순위 기반 스케줄링에서는 각 워프에게 우선순위를 설정하고, 높은 우선순위를 가진 워프를 먼저 실행한다. 우선순위는 여러 기준에 따라 설정될 수 있다: - 워프의 생명 주기 - 남은 명령어 수 - 메모리 액세스 패턴

장점

• 중요한 작업을 신속하게 처리 가능

단점

• 우선순위가 낮은 작업들이 무한히 기다릴 수 있음 (Starvation 문제)

연결 스케줄링

연결 스케줄링은 분기점에서 워프를 합쳐서 실행 효율을 높이는 방법이다. 머신에서 더 많은 워프를 동시에 실행하기 위해 분기점에서 분기된 여러 워프를 가능한 한 많이 합친다.

장점

• 실행 효율 극대화
• 분기 이후 명령어 실행 병목 현상 감소

단점

• 복잡한 구현
• 특정 분기 패턴에 최적화 필요

워프 다이버전스

워프 다이버전스(Warp Divergence)는 워프 내 쓰레드들이 서로 다른 경로를 따라가면서 발생하는 문제이다. GPU는 각 워프 내에서 동일한 명령어를 동시에 실행하도록 설계되었기 때문에, 워프 내 쓰레드들이 분기점에서 다른 경로를 따를 경우, 병목 현상이 발생하게 된다. 이는 GPU 성능에 큰 영향을 미칠 수 있다.

원인

워프 다이버전스는 일반적으로 조건문과 반복문 등에서 발생한다. 예를 들어, 한 워프 내에 있는 32개의 쓰레드 중 16개는 조건문을 만족하고, 16개는 그렇지 않은 경우를 생각해볼 수 있다.

다이버전스 해결 방법

1. 분기 예측: 분기점에서 수행될 명령어를 미리 예측하여 준비함.
2. 조건부 실행: 특정 조건을 만족하는 쓰레드만 명령어를 실행함.

다이버전스 영향을 줄이는 코드 작성법

• 코드 경로 일치: 가능한 한 워프 내 쓰레드들이 동일한 경로를 따르도록 코드 작성.
• 작은 분기: 분기 내의 명령어 수를 최소화하여 다이버전스 수행을 빠르게 마무리.
• 병렬 연산: 다이버전스가 적은 연산을 우선적으로 병렬 처리.

워프 압축(SIMT 모델)

SIMT(Single Instruction, Multiple Threads) 모델은 워프 다이버전스를 최소화하기 위한 접근 방식이다. 이 모델은 다음과 같은 특징을 갖는다:

1. 스칼라 명령어 스트림: 동일한 명령어 스트림을 여러 쓰레드에 적용.
2. 동기화: 쓰레드 간의 동기화를 통해 효율적으로 병렬 연산을 수행.

워프 압축(SIMT) 장점

• 효율성: 동일한 명령어 스트림을 동시에 여러 쓰레드에 적용함으로써 연산 유닛을 최적으로 활용.
• 간소화된 스케줄링: 워프 단위로 스케줄링하여 관리 복잡성 감소.

실전 예제

코드 예제: 다이버전스가 발생하는 경우

__global__ void exampleKernel(int *data) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    if (data[idx] > 0) {
        data[idx] *= 2;
    } else {
        data[idx] -= 1;
    }
}

위의 코드는 워프 다이버전스를 일으킬 수 있다. 왜냐하면 if 조건에 따라 워프 내 쓰레드들이 서로 다른 코드를 실행하기 때문이다.

코드 예제: 다이버전스를 최소화한 경우

__global__ void optimizedKernel(int *data) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int val = data[idx];
    data[idx] = (val > 0) ? val * 2 : val - 1;
}

위의 코드는 삼항 연산자(? :를 사용)를 활용하여 다이버전스를 최소화하였다. 따라서 워프 내 모든 쓰레드가 동일한 명령을 실행한다.

---

워프 스케줄링과 다이버전스 관리는 GPU 성능 최적화에 중요한 요소이다. 올바른 스케줄링 알고리즘 선택 및 워프 다이버전스를 최소화하는 코드 작성은 GPU 프로그램의 효율성을 크게 향상시킬 수 있다.