Warp Stall Reason Breakdown

GPU 의 Warp schduler 를 이해하기 위해 

warp 가 어떤 이유로 stall 에 빠지는지 파악 필요

다음과 같은ㅇ 이유로 warp 가 멈춤

• memory dependency stall ( L1TEX / L2 대기 )
• scoreboard dependency stall ( FMA 체인등 )
• IMC stall ( instruction immediate constant miss )
• branch / fech stall
• structural hazard stall

test 를 통해 각각의 stall 을 직접 유발, Nsight Compute 의 메시지 출력 확인

 

테스트 코드 - Stall 생성 전용 Micro - Kernel, Nsight Compute 결과 분석

(1) mem_stall_kernel - L1TEX / L2 miss stall 유발 (메모리 병목)

__global__
void mem_stall_kernel(const float* __restrict__ in, float* __restrict__ out, int N)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    float acc = 0.f;

    // 랜덤 패턴 → L1/L2 miss 증가 → memory stall 강제
    for (int i = 0; i < 4096; i++) {
        int idx = (tid * 97 + i * 57) & (N - 1);
        acc += in[idx];
    }

    out[tid] = acc;
}


stalled waiting for a scoreboard dependency on L1TEX
IMC miss stall ~ 60~80%
L1/L2 hit rate 낮음


OPT   Memory dependency stalls detected.
      Consider improving coalescing or reuse.

• 랜덤 패턴 때문에 coalescing 실패
• 대부분의 load 가 L1/L2 miss -> warp 가 L1TEX 응답을 기다리며 tall
• warp 가 issue-ready 상태로 오래 존재하지 못함

 

 

(2) dep_stall_kernel -> FMA dependency chain (Scoreboard stall)

__global__
void dep_stall_kernel(const float* __restrict__ in, float* __restrict__ out, int N)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    float acc = in[tid];

    // pipeline dependency 강제 → scoreboard stall
    for (int i = 0; i < 4096; i++) {
        acc = acc * 1.0000001f + 0.0000001f;
    }

    out[tid] = acc;
}


stalled waiting on a scoreboard dependency
dependency chain detected

Compute Workload Analysis 의 특징

• Executed IPC 낮게 나옴
• Issue Slots Busy 낮음
• FMA pipeline stall 비중이 매우 큼

의미

• FMA 가 previous FMA 결과를 기다림 - scoreboard stall
• 메모리 stall 이 아니므로 L1/L2 metricsc 는 정상
• IPC 가 크게 떨어

 

(3) mixed_stall_kernel -> 메모리 + 의존성 혼합 stall

__global__
void mixed_stall_kernel(const float* __restrict__ in, float* __restrict__ out, int N)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    float acc = in[tid];

    for (int i = 0; i < 4096; i++) {
        acc = acc * 1.0000001f + 0.0000001f; // dependency
        int idx = (tid * 17 + i * 31) & (N - 1); // random load
        acc += in[idx];
    }

    out[tid] = acc;
}



Memory stall: L1TEX dependency ~40%
Dependency stall: scoreboard ~45%

• 연산 사이에 메모리 load 가 끼어 있어
  • dependency chain + 메모리 latency 가 서로 역학적 stall 발생
• 실전에서 타일링이 부재시 가장 흔한 패턴

 

실험을 통해 얻는 핵심 결론

1) warp stall은 크게 둘이다

① memory stall

• L1TEX/L2 miss
• coalescing 실패
• DRAM latency 노출

② dependency stall

• 연산 파이프라인이 이전 연산 결과를 기다림
• scoreboard dependency
• ILP 부족

이 두 종류가 섞이면 GPU 성능이 가장 빠르게 붕괴한다.

---

2) GPU 성능 최적화란 = stall을 제거하는 과정

• shared memory tiling → global memory stall 줄임
• L2 reuse 개선 → pass-by-pass 성능 증가
• ILP 증가 (register tiling, loop unroll) → dependency stall 감소
• warp specialization → prefetch warp와 compute warp를 분리
• cp.async → L2→shared load를 overlap하여 latency 제거
• double-buffering → 연산/로드 겹침

---

3) Nsight Compute는 stall의 원인을 ‘정확히’ 알려준다

본 실험에서는 특히 다음 섹션을 통해 명확하게 확인할 수 있다:

• Warp State Statistics
• Scheduler Statistics
• Memory Workload Analysis
• Source Counters
• L1TEX stall breakdown

이 정보만 보면 kernel이 어느 방향으로 최적화해야 하는지 바로 보인다.