배치 루프 (여러 번 런치) vs 배치 결합 (한 번에 런치)

위는 CUDA 에서 자주 쓰이는 선택지

 

• 배치 루프 (멀티 런치)

각 배치 b 마다 같은 커널을 별도로 호출

for (int b = 0; b < B; ++b) {
  kernel<<<grid, block>>>(in + b*stride, out + b*stride, rows, cols);
}

 

• 배치 결합(싱글 런치)

배치 차원을 행으로 합쳐 한 번에 호출

kernel<<<gridB, block>>>(in /*배치까지 연속*/, out, B*rows, cols);

 

오버헤드 / 유연성 / 확장성에서 차이가 발생

 

성능·동작 관점의 차이

1) 커널 런치 오버헤드

• 멀티 런치: 런치 오버헤드가 배치 수만큼 누적. 작은 텐서일수록 상대 비용이 큼.
• 싱글 런치: 한 번만 호출 → 보통 더 빠름.
  • 특히 짧은 커널(activation, elementwise)은 차이가 눈에 띕니다.
  • CUDA Graph 캡처/재생과도 궁합이 좋음(런치 수 최소화).

2) 스트림·파이프라이닝(겹치기)

• 멀티 런치: 배치별로 다른 스트림에 던져서 H2D/D2H 복사와 오버랩 가능.
  • 예: b는 계산, b+1은 H2D 복사 동시 수행 → 파이프라인 구성 유리.
• 싱글 런치: 한 덩어리라 파이프라이닝 자유도가 낮음.
  • 대신 한 번의 큰 커널로 SM 점유가 잘 이뤄지면 순수 계산 위주에서는 빠름.

3) 그리드 차원·안전성

• 싱글 런치: rows' = B*rows가 아주 크면 grid.y 상한(구형 65,535 등)에 닿을 수 있음.
  • 해결: block.y 키우거나, 1D 인덱싱으로 분해, 또는 두 번에 나눠 호출.
  • 일반적인 DL 배치/크기에선 거의 문제 없음.
• 멀티 런치: 각 호출의 그리드가 작아 상한 걱정 적음.

4) 메모리 연속성/어드레싱

• 싱글 런치는 in/out이 배치까지 연속 메모리여야 인덱싱이 간단하고 coalesced 유지.
• 멀티 런치는 배치 간 stride가 달라도 각 호출에서 포인터만 맞추면 OK.

5) 디버깅·로깅

• 멀티 런치: 문제난 배치를 쉽게 특정(루프 인덱스).
• 싱글 런치: (global_row / rows)로 역산해야 해서 로그가 약간 불편.

6) 스케줄링·레지스터/SMEM 압박

• 싱글 런치: 커널 하나가 길게 돌면(다른 이유로) Watchdog(TDR) 우려(Windows 디스플레이 GPU).
  • Activation 같은 경량 커널은 안전. 초장시간 커널(거대 합성)은 유의.
• 멀티 런치: 커널이 짧아 TDR 위험 낮고, 작업 나눔으로 스케줄러 유연성↑.

7) 라이브러리/그래프 연동

• 싱글 런치: CUDA Graph나 “한 step을 캡처해서 재생”하려면 런치 수가 적을수록 유리.
• 멀티 런치: 그래프 캡처 시 노드 수가 증가 → 초기 캡처/검증 비용↑.

---

메모리 접근·캐시 관점

• 두 방식 모두 x축을 col에 매핑하고, warp가 col 연속을 잡으면 coalescing 동일.
• 캐시 관점에서도 큰 차이는 없음. 다만 싱글 런치는 전체 데이터를 한 번에 훑어 스케줄러가 더 큰 워크풀을 가진다는 장점이 있고, 멀티 런치는 배치 경계마다 스케줄·캐시가 리셋되는 형태.

---

유지보수·유연성 관점

• 멀티 런치
  • 배치마다 다른 파라미터/shape(ragged batch)도 처리 가능.
  • per-batch 조건 분기·메트릭 수집이 쉽다.
• 싱글 런치
  • 모든 샘플이 동일 shape & 동일 파라미터 가정이 편함.
  • 코드가 더 단순(한 번 인덱싱, 한 번 호출).

---

언제 무엇을 쓰나 (의사결정 가이드)

• 짧은 커널(activation, elementwise), 동형 배치:
  → 싱글 런치 권장(오버헤드 최소화, Graph-friendly).
• 입출력 복사와 겹치고 싶다 / 배치별 비동기 파이프라인:
  → 멀티 런치(+다중 스트림).
• 매우 큰 rows’ = B*rows로 그리드 한계가 걱정**:**
  → 멀티 런치 또는 싱글 런치지만 블록/그리드 재구성.
• 디버깅·로그 가독성이 최우선:
  → 멀티 런치.

---

실전 팁

• 싱글 런치로 바꿀 때:
  • 인덱싱만 rowsB = B*rows; colsB = cols로 바꾸면 끝.
  • 인덱스는 32비트로 충분한지 확인(아주 큰 텐서면 size_t/64-bit 연산 사용).
• 멀티 런치에서 병렬화:
  • 스트림 풀 만들고 b % num_streams로 라운드로빈.
  • H2D(b+1) ↔ Compute(b) ↔ D2H(b-1) 파이프라인.
• CUDA Graph:
  • 싱글 런치가 노드 수를 줄이므로 캡처/재생 성능이 유리.

---

결론 (한 줄 버전)

• 기본값: 싱글 런치(배치 결합)가 빠르고, 그래프/오버헤드 면에서 유리.
• 필요할 때만: 멀티 런치(배치 루프)를 써서 스트림 파이프라인, 디버깅, 특수 배치 형태를 처리.