__global__
void warp_divergence_test(int* out) {
int tid = threadIdx.x;
int warp_id = tid / 32;
int lane_id = tid % 32;
if (lane_id < 16)
out[tid] = 111;
else
out[tid] = 222;
}
////////////////////////////////////////
[0..15] = 111
[16..31] = 222
- warp 는 동일 명령만 동시에 실행함
- 분기하면 warp 는 A-path -> B-path 순차 실행
- 즉, divergence = 순차처리 -> 성능 감소
'GPU-KERNEL' 카테고리의 다른 글
| Warp Schedulig - clock64 기반 스케줄링 관찰 실험 (0) | 2025.11.29 |
|---|---|
| GPU Hierarchy - Block / Warp / Thread 구조 실험 (0) | 2025.11.29 |
| GPU Architecture & Execution Model (0) | 2025.11.29 |
| Optimization Principles 와 Tensor Core Optimization 의 분리 (0) | 2025.11.28 |
| TensorCore K 방향 타일 루프 + cp.async 2 stage pipeline (0) | 2025.11.28 |