GPU 관련 개념 꽉 잡기

gpu_time_duration.sm [ms]

GPU 전체가 해당 커널을 실행하는 데 사용한 총 시간

여러 SM, 여러 서브 유닛이 관여하면, 그 모든 sub-unit 의 시간을 합산한 값을 나타냄

 

gpu_time_duration_measured_user

Nsight Compute 가 정확하게 수집 가능한 상황일때, 사용하는 기준 시간

즉, Nsight 가 실제 유저 코드 구간을 기준으로 정확히 측정할 수 있을 때 이 값으로 표시됨

 

gpu_time_duratoin_measured_wallclock

위 user 시간 측정이 불가능할 때 fallback 하는 시간

 

sm_throughput.avg.pct_of_peak_sustained_elapsed

현재 커널이SM 이 낼 수 있는 이론적 지속 처리량 대비 얼마나 활용했는지의 지표

즉, 이 커널이 SM 성능을 얼마나 끌어냈는가를 나타내는 핵심점수

 

sm Streaming Multiprocessor

GPU 의 실제 실행 유닛

커널은 32 스레드 단위 warp 로 실행

warp 들은 block (CTA) 로 묶여 있고, 하나의 block 은 반드시 한 SM 위에서만 실행됨

같은 block 내 스레드들은 shared memory 등 여러 리소스를 공유한다.

 

SMSPs

하나의 SM 은 4개의 sub partition 으로나뉜 구조

각 SMSP 는 

독립적인 warp 처리 파이프라인을 갖고 있음, 각자 일정 수의 warps 를 스케줄링함

실제 계산은 대부분 SMSP 단위에서 이뤄지고, Nsight Compute 의 많은 metric 은 sm 전체 - smsp 묶음 기준으로 계산된다.

 

launch_registers_per_thread

커널 실행 시 각 스레드가 할당받는 레지스터 개수

해당 값은 컴파일러가 결정, 레지스터 수가 많을수록 스레드당 성능은 좋아지지만, 한 SM 에 동시에 올릴 수 있는 warp 수는 줄어든다.

 

registers

SM 은 4개의 sub partition 으로 나뉘는데, 각 sub partition 이 자체적인 레지스터 파일을 가짐

레지스터는 HW 에서 정해진 크기 단위로 chunk 할당된다.

한 warp 가 많은 레지스터를 요구하면 레지스터 파일을 많이 차지해서 occupancy 가 줄어듦

 

thread   

GPU 의 기본 실행 단위

각 thread 는 자신만의 독립적인 레지스터 세트를갖는다.

warp ( 32 threads ) 단위로 실제 실행되므로, 사용 레지스터 수 * 32 가 warp 레지스터 사용량이다.

 

레지스터 - occupanc - latency hiding - SM utilization

  

1. 레지스터(registers_per_thread)가 많아지면?

✔ 각 스레드가 더 많은 상태를 보관 가능

• tile fragment
• accumulator fragment
• A/B sub-tile
• loop counter
• pointer offset
  → 단일 thread 의 계산 효율은 증가.

✔ 하지만 대가가 있다

SM 의 레지스터는 한정되어 있고, warp 단위로 나눠쓴다.

예:

• Ampere 기준 SMSP 당 약 ~16K 32-bit 레지스터
• warp 1개 = 32 스레드
• 레지스터 per thread 가 올라가면, warp 1개가 차지하는 레지스터도 커짐
  → 동시에 투입 가능한 warp 수가 줄어듦.

---

2. occupancy 감소로 연결

✔ occupancy = “SM 안에서 동시에 활성화될 수 있는 warp 수”

• 많은 warp 가 동시에 존재할수록 GPU 는 더 많은 latency hiding 가능.
• 레지스터 사용량이 증가하면 occupancy 가 줄어든다.

예시:

• 64 regs/thread → SM 에 warp 48개 가능
• 128 regs/thread → SM 에 warp 24개만 가능
• 192 regs/thread → SM 에 warp 12개 이하

Tensor Core GEMM 은 보통 120~160 regs/thread 를 쉽게 넘어간다.

핵심 결과

• 레지스터 증가 → occupancy 감소
• 이는 GEMM 의 streaming pipeline 구조에 매우 직접적인 영향을 준다.

---

3. occupancy 감소 → latency hiding 악화

GPU 성능의 핵심은 **“메모리/파이프라인 지연을 숨기는 능력”**이다.

SM 이 실제로 처리하는 과정

• warp A: A/B tile 글로벌 로드 (cp.async 또는 ld.global)
• warp B: MMA 수행
• warp C: accumulator write
• warp D: 다음 tile preload

이런 파이프라인을 구성하려면 동시에 충분한 warp 가 필요하다.

occupancy 가 낮으면?

• 글로벌 메모리에서 오는 latency (400~800 cycles) 숨길 warp 가 부족
• shared memory bank conflict 회피 능력 감소
• tensor core pipeline 에 bubble 발생
• issue slot 에 idle cycle 증가

즉, “stall → stall → stall” 이 된다.

---

4. latency hiding 악화 → SM utilization 하락

Nsight Compute 의 페이지에서 자주 보는 값:

• sm__throughput.pct_of_peak_sustained_elapsed
• sm__warps_active.avg.pct_of_peak
• sm__pipe_tensor_active
• sm__sass_average_branch_cost
• sm__inst_executed_pipe_tensor.sum

이런 지표들이 떨어지는 근본 원인 중 하나가 바로
레지스터로 인해 occupancy 가 제한되어 파이프라인이 채워지지 않았기 때문이다.

Tensor Core GEMM 같은 커널에서는 몇 가지 계층이 있다:

✔ 충분한 warp → 충분한 파이프 스테이지 → 높은 SM utilization

✔ warp 부족 → pipeline bubble → 낮은 utilization

✔ 너무 많은 레지스터 사용 → warp 부족

→ 결국 SM utilization 이 바닥으로 떨어진다.

---

⭐ 전체 흐름을 한 줄로 요약하면

레지스터를 많이 쓰면 단일 thread 성능은 좋아지지만, 전체 warp 수는 줄고,
warp 가 줄면 지연을 숨기기 어려워져 pipeline 이 비고, SM utilization 이 떨어진다.

---

✔ GEMM 최적화에서 실제로 어떻게 적용하냐?

① 레지스터를 너무 줄이면?

• fragment reuse가 줄고
• indexing / pointer math 가 반복
  → 성능 저하

② 레지스터를 너무 많이 쓰면?

• occupancy 감소
• latency hiding 불가
  → SM 이 idle 상태 증가

그래서 cuBLAS / CUTLASS / TRT가 하는 방식

• tile size, warp tile shape, mma per warp 수를 조절하여
  ‘적정한 레지스터 사용량’에서 최대 throughput을 뽑는 설계.

대표적으로:

• Ampere FP16 TC → warp tile 64x64x16 / 2-stage pipeline → 120~160 regs/thread
• Hopper FP8 TC → warp tile 64x128 / 2~3 stage → 160~200 regs/thread