본문 바로가기

GPU-KERNEL

Warp Divergence - SIMT 분기 실험

__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 = 순차처리 -> 성능 감소