🧠 워프 레인과 SIMT 실행

워프 프로그래밍 vs SIMD 멘탈 모델

워프란 무엇인가?

워프는 32개(또는 64개)의 GPU 스레드가 서로 다른 데이터에 대해 동일한 명령을 동시에 실행하는 그룹입니다. 각 스레드가 벡터 프로세서의 “레인” 역할을 하는 동기화된 벡터 유닛이라고 생각하면 됩니다.

간단한 예시:

from gpu.primitives.warp import sum
# 워프 내 32개 스레드가 동시에 실행:
var my_value = input[my_thread_id]     # 각 스레드가 서로 다른 데이터를 가져옴
var warp_total = sum(my_value)         # 모든 스레드가 하나의 합계에 기여

무슨 일이 일어난 걸까요? 32개의 개별 스레드가 복잡한 조율을 하는 대신, 워프가 자동으로 동기화하여 하나의 결과를 만들어냈습니다. 이것이 바로 SIMT(Single Instruction, Multiple Thread) 실행입니다.

SIMT vs SIMD 비교

CPU 벡터 프로그래밍(SIMD)에 익숙하다면, GPU 워프는 비슷하지만 핵심적인 차이가 있습니다:

CPU SIMD 방식 (C++ intrinsics):

// 명시적 벡터 연산 - 8개의 float를 병렬로
__m256 result = _mm256_add_ps(a, b);   // 8쌍을 동시에 덧셈

CPU SIMD 방식 (Mojo):

# Mojo에서 SIMD는 일급 시민 타입이므로 a, b가 SIMD 타입이면
# 덧셈이 병렬로 수행됩니다
var result = a + b # 8쌍을 동시에 덧셈

GPU SIMT 방식 (Mojo):

# 스레드 기반 코드가 벡터 연산으로 변환됩니다
from gpu.primitives.warp import sum

var my_data = input[thread_id]         # 각 스레드가 자기 요소를 가져옴
var partial = my_data * coefficient    # 모든 스레드가 동시에 계산
var total = sum(partial)               # 하드웨어가 합산을 조율

워프를 강력하게 만드는 핵심 개념

1. 레인 식별: 각 스레드는 사실상 비용 없이 접근할 수 있는 “레인 ID” (0~31)를 갖습니다

var my_lane = lane_id()  # 하드웨어 레지스터를 읽을 뿐

2. 암시적 동기화: 워프 내에서 배리어가 필요 없습니다

# 그냥 동작 - 모든 스레드가 자동으로 동기화
var sum = sum(my_contribution)

3. 효율적인 통신: 메모리 없이도 스레드 간 데이터 공유가 가능합니다

# 레인 0의 값을 다른 모든 레인으로 전달
var broadcasted = shuffle_idx(my_value, 0)

핵심 통찰: SIMT를 사용하면 자연스러운 스레드 코드를 작성하면서도 효율적인 벡터 연산으로 실행할 수 있어, 스레드 프로그래밍의 편리함과 벡터 처리의 성능을 모두 얻을 수 있습니다.

GPU 실행 계층 구조에서 워프의 위치

워프가 전체 GPU 실행 모델과 어떻게 연결되는지 자세히 알아보려면 GPU 스레딩 vs SIMD 개념을 참고하세요. 워프의 위치는 다음과 같습니다:

GPU 디바이스
├── 그리드 (전체 문제)
│   ├── 블록 1 (스레드 그룹, 공유 메모리)
│   │   ├── 워프 1 (32 스레드, 록스텝 실행) ← 이 레벨
│   │   │   ├── 스레드 1 → SIMD 연산
│   │   │   ├── 스레드 2 → SIMD 연산
│   │   │   └── ... (총 32개 스레드)
│   │   └── 워프 2 (32 스레드)
│   └── 블록 2 (독립적인 그룹)

워프 프로그래밍은 “워프 레벨“에서 동작합니다 - 단일 워프 내의 32개 스레드를 모두 조율하는 연산을 다루며, 그렇지 않으면 복잡한 공유 메모리 조율이 필요한 sum() 같은 강력한 기본 요소를 사용할 수 있습니다.

이 멘탈 모델은 문제가 워프 연산에 자연스럽게 매핑되는 경우와 기존의 공유 메모리 방식이 필요한 경우를 구분하는 데 도움이 됩니다.

워프 프로그래밍의 하드웨어 기반

Single Instruction, Multiple Thread(SIMT) 실행을 이해하는 것은 효과적인 워프 프로그래밍에 필수적입니다. 이것은 단순한 소프트웨어 추상화가 아니라, GPU 하드웨어가 실리콘 수준에서 실제로 작동하는 방식입니다.

SIMT 실행이란?

SIMT란 워프 내에서 모든 스레드가 서로 다른 데이터에 대해 같은 명령을 동시에 실행한다는 뜻입니다. 이는 완전히 다른 명령을 독립적으로 실행할 수 있는 CPU 스레드와 근본적으로 다릅니다.

CPU vs GPU 실행 모델

GPU 워프 실행 패턴:

• 명령: 32개 레인 모두 동일: add r1, r2
• 레인 0: Data0에 연산 → Result0
• 레인 1: Data1에 연산 → Result1
• 레인 2: Data2에 연산 → Result2
• … (모든 레인이 동시에 실행)
• 레인 31: Data31에 연산 → Result31

핵심 통찰: 모든 레인이 서로 다른 데이터에 대해 같은 명령을 동시에 실행합니다.

SIMT가 GPU에 적합한 이유

GPU는 지연 시간이 아닌 처리량에 최적화되어 있습니다. SIMT가 가능하게 하는 것들:

• 하드웨어 단순화: 하나의 명령 디코더가 32개 또는 64개 스레드를 처리
• 실행 효율성: 워프 내 스레드 간 복잡한 스케줄링 불필요
• 메모리 대역폭: 병합된 메모리 접근 패턴
• 전력 효율성: 레인 전체에 걸쳐 제어 로직 공유

워프 실행 메커니즘

레인 번호와 식별

워프 내 각 스레드는 0부터 WARP_SIZE-1까지의 레인 ID를 갖습니다:

from gpu import lane_id
from gpu.primitives.warp import WARP_SIZE

# 커널 함수 내에서:
my_lane = lane_id()  # 0-31 (NVIDIA/RDNA) 또는 0-63 (CDNA) 반환

핵심 통찰: lane_id()는 비용이 없습니다 - 값을 계산하는 것이 아니라 하드웨어 레지스터를 읽을 뿐입니다.

워프 내 동기화

SIMT의 가장 강력한 측면: 암시적 동기화.

# thread_idx.x < WARP_SIZE인 경우의 예시

# 1. 기존 공유 메모리 방식:
shared[thread_idx.x] = partial_result
barrier()  # 명시적 동기화 필요
var total = shared[0] + shared[1] + ... + shared[WARP_SIZE] # 합산 리덕션

# 2. 워프 방식:
from gpu.primitives.warp import sum

var total = sum(partial_result)  # 암시적 동기화!

왜 배리어가 필요 없을까요? 모든 레인이 각 명령을 정확히 같은 시점에 실행하기 때문입니다. sum()이 시작될 때, 모든 레인은 이미 partial_result 계산을 마친 상태입니다.

워프 분기와 수렴

조건 코드에서 무슨 일이 일어날까?

if lane_id() % 2 == 0:
    # 짝수 레인이 이 경로를 실행
    result = compute_even()
else:
    # 홀수 레인이 이 경로를 실행
    result = compute_odd()
# 모든 레인이 여기서 수렴

하드웨어 동작 단계:

예시 분석:

• 2단계: 짝수 레인만 compute_even()을 실행하고 홀수 레인은 대기
• 3단계: 홀수 레인만 compute_odd()를 실행하고 짝수 레인은 대기
• 총 소요 시간: time(compute_even) + time(compute_odd) (순차 실행)
• 분기 없는 경우: max(time(compute_even), time(compute_odd)) (병렬 실행)

성능 영향:

1. 분기: 워프가 실행을 분리 - 일부 레인은 활성, 나머지는 대기
2. 순차 실행: 서로 다른 경로가 병렬이 아닌 순차적으로 실행
3. 수렴: 모든 레인이 다시 합류하여 함께 진행
4. 비용: 분기가 있는 워프는 통합 실행 대비 2배 이상의 시간 소요

워프 효율을 위한 모범 사례

워프 효율 패턴

✅ 우수: 균일 실행 (100% 효율)

# 모든 레인이 같은 작업 수행 - 분기 없음
var partial = a[global_i] * b[global_i]
var total = sum(partial)

성능: 32개 레인 모두 동시 활성

⚠️ 허용: 예측 가능한 분기 (~95% 효율)

# lane_id() 기반 분기 - 하드웨어 최적화됨
if lane_id() == 0:
    output[block_idx] = sum(partial)

성능: 단일 레인의 짧은 연산, 예측 가능한 패턴

🔶 주의: 구조화된 분기 (~50-75% 효율)

# 규칙적인 패턴은 컴파일러가 최적화 가능
if (global_i / 4) % 2 == 0:
    result = method_a()
else:
    result = method_b()

성능: 예측 가능한 그룹, 일부 최적화 가능

❌ 회피: 데이터 의존적 분기 (~25-50% 효율)

# 데이터에 따라 레인마다 다른 경로를 탈 수 있음
if input[global_i] > threshold:  # 예측 불가능한 분기
    result = expensive_computation()
else:
    result = simple_computation()

성능: 무작위 분기가 워프 효율을 떨어뜨림

💀 최악: 중첩된 데이터 의존적 분기 (~10-25% 효율)

# 예측 불가능한 분기의 다단계 중첩
if input[global_i] > threshold1:
    if input[global_i] > threshold2:
        result = very_expensive()
    else:
        result = expensive()
else:
    result = simple()

성능: 워프 효율이 사실상 무너짐

크로스 아키텍처 호환성

NVIDIA vs AMD 워프 크기

from gpu.primitives.warp import WARP_SIZE

# NVIDIA GPUs:     WARP_SIZE = 32
# AMD RDNA GPUs:   WARP_SIZE = 32 (wavefront32 모드)
# AMD CDNA GPUs:   WARP_SIZE = 64 (전통적인 wavefront64)

왜 중요할까요:

• 메모리 패턴: 병합된 접근이 워프 크기에 의존
• 알고리즘 설계: 리덕션 트리가 워프 크기를 고려해야 함
• 성능 확장: AMD에서 워프당 레인이 2배

이식 가능한 워프 코드 작성

아키텍처 적응 전략

✅ 이식 가능: 항상 WARP_SIZE 사용

comptime THREADS_PER_BLOCK = (WARP_SIZE, 1)  # 자동으로 적응
comptime ELEMENTS_PER_WARP = WARP_SIZE       # 하드웨어에 맞게 확장

결과: NVIDIA/AMD (32)와 AMD (64) 모두에서 최적으로 동작

❌ 잘못된 방식: 워프 크기를 하드코딩하지 마세요

comptime THREADS_PER_BLOCK = (32, 1)  # AMD GPU에서 동작 안 함!
comptime REDUCTION_SIZE = 32          # AMD에서 잘못된 값!

결과: AMD에서 성능 저하, 정확성 문제 가능

실제 하드웨어 영향

64 vs 32의 성능 차이:

• CDNA 장점: 워프당 2배의 메모리 대역폭
• CDNA 장점: 워프당 2배의 연산량
• NVIDIA/RDNA 장점: 블록당 더 많은 워프 (더 높은 점유율)
• 코드 이식성: 같은 소스 코드로 양쪽 모두 최적 성능

워프와 메모리 접근 패턴

병합된 메모리 접근 패턴

✅ 완벽: 병합된 접근 (100% 대역폭 활용)

# 인접 레인 → 인접 메모리 주소
var value = input[global_i]  # 레인 0→input[0], 레인 1→input[1], 등

메모리 접근 패턴:

주소 예시:

• 병합: 레인 0→0, 레인 1→4, 레인 2→8, 레인 3→12, …
• 분산: 레인 0→1000, 레인 1→52, 레인 2→997, 레인 3→8, …

공유 메모리 뱅크 충돌

뱅크 충돌이란?

GPU 공유 메모리가 동시 접근이 가능한 32개의 독립적인 뱅크로 나뉘어 있다고 가정합니다. 뱅크 충돌은 워프 내 여러 스레드가 같은 뱅크의 서로 다른 주소에 동시에 접근하려 할 때 발생합니다. 이 경우 하드웨어가 접근을 직렬화해야 하므로, 단일 사이클이어야 할 연산이 여러 사이클로 늘어납니다.

핵심 개념:

• 충돌 없음: 각 스레드가 서로 다른 뱅크에 접근 → 모든 접근이 동시에 발생 (1 사이클)
• 뱅크 충돌: 여러 스레드가 같은 뱅크에 접근 → 접근이 순차적으로 발생 (N개 스레드에 N 사이클)
• 브로드캐스트: 모든 스레드가 같은 주소에 접근 → 하드웨어가 1 사이클로 최적화

공유 메모리 뱅크 구성:

뱅크 충돌 예시:

워프 프로그래밍의 실전 활용

워프 연산이 가장 효과적인 경우

1. 리덕션 연산: sum(), max() 등
2. 브로드캐스트 연산: shuffle_idx()로 값 공유
3. 이웃 통신: shuffle_down()으로 슬라이딩 윈도우
4. 누적 합 연산: prefix_sum()으로 scan 알고리즘

성능 특성

다음 단계

SIMT의 기반을 이해했으니, 이 개념이 어떻게 강력한 워프 연산을 가능하게 하는지 알아볼 차례입니다. 다음 섹션에서는 sum()이 복잡한 리덕션 패턴을 간단하고 효율적인 함수 호출로 어떻게 변환하는지 보여줍니다.

→ 다음: warp.sum()의 핵심