gpumode · 강의 아카이브
《GPU Mode》 L076 2025 High priority transcript · failed

BackendBench — fixing the LLM kernel correctness problem

"LLM 으로 GPU 커널을 짠다" 는 표어가 폭증한 2025 년의 진짜 문제는 속도가 아니라 정확성이다 — 그럴듯해 보이는 Triton 커널이 silently 잘못된 결과를 내고, 그것이 학습 loss 의 미세한 drift 로만 드러난다. Mark Saroufim 이 PyTorch 팀의 입장에서 만든 BackendBench 는 그 silent failure 를 잡아내려는 첫 정량 도구. 학습 노트.

BackendBench LLM kernels correctness Triton numerical tolerance PyTorch op coverage eval harness auto-fix
M
Speaker
Mark Saroufim
Meta · PyTorch core · GPU Mode 운영진
강의 번호
L076
학습 우선순위
High
자막
failed
다시 볼 때
repo 직접 돌리기
§ 01강의가 풀려는 문제· why this lecture exists

"빠른 커널" 이 아니라 "맞는 커널" 의 정의 자체가 흐려진 자리

2024-2025 년의 큰 흐름 — Claude / GPT / DeepSeek 같은 LLM 이 Triton 커널을 직접 짠다. KernelBench, GPU-Bench 같은 leaderboard 가 등장하고, 가장 빠른 커널이 누구의 모델이 짰는지 경쟁하는 분위기가 생겼다. 그런데 그 커널들이 실제로 맞는가는 별개의 질문이고, 잘 측정되지 않고 있었다.

BackendBench 는 그 공백을 메운다. 강의의 두 줄 요약.

  1. "빠른 커널" 의 leaderboard 는 correctness 를 너무 약하게 본다. 단일 input 에서 atol 1e-2 통과면 "정답" 으로 친다. 그러나 production 학습 / 추론에서는 numerical drift, edge case (NaN, inf, broadcasting), shape variation, dtype 조합이 모두 깨진다.
  2. PyTorch 의 OpInfo 위에 LLM 커널을 검증하는 표준 harness 가 필요하다. PyTorch 의 aten op 들은 이미 잘 정의된 reference 가 있고 (eager mode), 그것에 대해 LLM 이 짠 커널을 입력 다양성 × tolerance × dtype 조합으로 측정.
강의의 인지적 frame

"LLM 이 GPU 커널을 짜준다" 는 hype 의 다음 자리는 그 커널을 production 에 풀어 둘 수 있는가 의 질문이다. BackendBench 는 그 질문에 답하기 위한 PyTorch-eager 기반 차이 측정 harness — 즉 "PyTorch 의 reference 와 일치하는가" 를 op 별, shape 별, dtype 별로 모두 표로 만든다.

"우리는 빠른 커널 leaderboard 를 만든 게 아니다 — '실제로 맞는' 커널의 분포를 본다."Mark Saroufim · 확인 필요
§ 02LLM 커널의 correctness 문제· silent failure

"빠른 코드" 와 "맞는 코드" 의 분포가 다르다

LLM 이 만든 커널이 깨지는 자리는 보통 사람이 잘 안 보는 자리다 — production 에서만 등장하는 입력의 모양, drift 가 만들 수 있는 dtype 조합, 그리고 silent NaN.

흔히 보는 silent failure

  • broadcasting 시 shape mismatch — atol 통과 input 만 테스트하면 안 보임
  • dtype promotion 누락 — fp16 input 을 받아 fp32 reduction 안 하고 직접 sum
  • backward pass 의 grad 부호 / scale 오류
  • NaN propagation — 0/0 처리 누락
  • causal mask 의 off-by-one
  • reduction axis 가 헷갈림 (mean over dim=-1 vs dim=0)
  • numerical instability — exp(large) overflow

BackendBench 가 잡아주는 것

  • 같은 op 의 OpInfo sample 들을 자동 enumerate — shape 다양성
  • fp16/bf16/fp32 + 작은/큰 input 모두
  • NaN/Inf 보존 비교 (단순 atol 이 아닌 equiv)
  • backward grad 비교 (forward 만 보지 않음)
  • per-op scoreboard — 어디가 깨졌는지 정확히 표시
  • OpInfo 의 reference 와의 ulp 단위 차이까지

왜 이게 KernelBench 같은 다른 도구와 다른가 — KernelBench 는 한 종류의 op (matmul, attention) 의 best-of-N 속도를 본다. BackendBench 는 PyTorch 의 op 전체를 cover 하는 평면적 harness 다. 한 LLM 의 커널이 attention 에서 빛난다 해도 dropout 에서 fail 하면 production 에서는 동일하게 사용 불가능.

왜 silent 가 silent 인가

학습 loss 곡선은 평균 metric. 한 batch 에서 일부 element 만 잘못된 값을 받아도 loss 의 차이는 작다. 며칠 지난 뒤에 evaluation accuracy 에서만 1-2% 격차로 드러난다. 이 시점이면 어떤 커널이 원인인지 추적이 거의 불가능.

§ 03benchmark 설계· how BackendBench works

OpInfo 위에 얹는 표준 harness — 한 op 당 수천 sample

PyTorch 의 testing 인프라는 이미 거의 모든 aten op 에 대해 OpInfo entry 를 가진다 — input sample generator, output dtype rule, autograd support 정보. BackendBench 는 그 위에 얹힌다.

FIG · BackendBench 의 evaluation pipeline한 op 당 한 줄
01target op 선택aten::softmax 같은 PyTorch op 한 개
02LLM 으로 Triton 커널 생성prompt: "write Triton kernel for aten::softmax with same semantics"
03OpInfo sample iterateshape × dtype × device. 보통 100~수백 sample
04eager reference 실행PyTorch 의 표준 구현으로 ground truth 생성
05LLM kernel 실행생성된 코드를 안전한 sandbox 에서 import / 호출
06numerical 비교atol/rtol pass · NaN equiv · backward grad · 모두 표로
07per-op aggregate통과율, 실패 sample list, 첫 실패 input
08leaderboard 갱신모델 × op 의 매트릭스
"sandbox" 가 핵심 — LLM 이 생성한 임의 코드를 그냥 import 하면 보안 위험. BackendBench 는 별도 process 로 격리하고 timeout 을 강제 (확인 필요).

이 설계의 진짜 미덕은 분산 가능하다는 점. op 단위로 독립이라 1000 개 op 를 1000 개 worker 에 뿌릴 수 있다. PyTorch 가 OpInfo 를 잘 관리해온 부산물이 LLM 커널 검증 인프라가 된다.

§ 04채점 기준· tolerance & coverage

"맞다" 의 정의를 직접 구체화한다

numerical kernel 의 correctness 는 binary 가 아니다. 어디까지를 "통과" 로 볼지 명확하지 않으면 같은 커널이 누구는 통과 누구는 실패가 된다. BackendBench 의 채점 기준 — 강의에서 Mark 가 의도적으로 conservative 하게 잡았다고 강조한 부분.

  • fp32: atol=1e-5, rtol=1e-5 — eager 와 거의 비트 동일까지 요구.
  • fp16: atol=1e-3, rtol=1e-3 — fp16 의 정상 noise 범위.
  • bf16: atol=1e-2, rtol=1e-2 — bf16 의 mantissa 7-bit 한계.
  • NaN/Inf 보존 — input 에 NaN 이 있으면 output 에도 같은 위치에 NaN 이 있어야 한다. 단순 numerical close 가 아닌 equiv 비교.
  • backward grad — forward 통과 후 random gradient 로 backward, eager 와 비교. autograd 미지원 op 는 skip.
  • shape coverage — OpInfo sample 의 모든 shape 을 통과해야 op-level pass. 80% 통과면 "partial" 로 표시.
왜 conservative 인가

production 에서 학습 / 추론에 풀어두려면 "거의 비트 동일" 이 사실상 기준. 학습 길이 1M step 누적이면 1e-3 의 차이가 final loss 에서 정상 모델과 분기한다.

# BackendBench 의 비교 함수 (개념)
def compare(out_kernel, out_ref, dtype):
    # NaN 위치 일치
    if torch.isnan(out_ref).any():
        if not torch.equal(
            torch.isnan(out_kernel),
            torch.isnan(out_ref),
        ): return Fail("nan_mask")

    # Inf 위치 일치
    if torch.isinf(out_ref).any():
        if not torch.equal(
            torch.isinf(out_kernel),
            torch.isinf(out_ref),
        ): return Fail("inf_mask")

    atol, rtol = TOLERANCES[dtype]
    if not torch.allclose(
        out_kernel, out_ref,
        atol=atol, rtol=rtol,
        equal_nan=True,
    ): return Fail("value_mismatch")

    return Pass()
"correctness 는 binary 가 아니지만 'maybe' 같은 답을 받아들이면 production 에서는 silent fail 만 남는다."Mark Saroufim · 확인 필요
§ 05자주 보는 오류· failure taxonomy

같은 모델이 같은 종류의 실수를 반복한다 — 그것이 패턴

BackendBench 의 부산물 — LLM 이 짠 커널이 실패하는 패턴이 표로 정리된다. 강의에서 Mark 가 강조한 5 가지 failure mode.

dtype promotion miss
fp16 input 을 받아 직접 sum 한다. eager 는 fp16 → fp32 promotion 후 sum. 큰 reduction 에서 catastrophic accumulation error.
broadcasting off
shape (B,1,N) 과 (B,M,N) 의 broadcasting 에서 stride 계산을 혼동. 작은 shape 에선 통과하다 큰 shape 에서 깨짐.
numerical instability
softmax 에서 subtract-max 누락 또는 잘못 적용. 큰 logit 입력에서만 silent NaN.
backward 누락 / 부호 오류
forward 만 짜고 backward 는 PyTorch 의 autograd 가 알아서 짜준다고 가정. 사실은 custom op 면 직접 짜야 함.
causal / mask off-by-one
attention 의 causal mask 에서 한 element 어긋남. 짧은 sequence 에선 별 차이 없음, 긴 sequence 에서 학습이 미세하게 fail.
contiguous 가정
non-contiguous tensor (예: transpose 결과) 가 들어왔을 때 stride 무시. 작은 tile 에선 통과, 큰 sequence 에선 wrong stride read.
edge shape
empty tensor (shape=(0,)), shape=1, dim=0 의 reduce 같은 edge case 에서 분기 누락. eager 는 잘 처리.
launch shape mismatch
grid 계산이 ceil_div 를 안 함. shape 의 마지막 element 가 누락 — silent.
왜 LLM 이 이런 실수를 하는가

학습 데이터에 있는 Triton 코드의 대부분이 happy-path 만 있다. github 의 인기 repo 가 가지는 typical input shape, dtype 위에서만 검증된 코드. 그래서 LLM 이 학습한 분포가 OpInfo 의 corner case 분포와 어긋난다. BackendBench 는 그 어긋남을 정량화한다.

§ 06LLM 별 결과· leaderboard

"Claude vs GPT vs DeepSeek 의 정확성 비교" — 거친 그림

강의에서 Mark 가 발표 시점의 leaderboard 한 장면을 공유한다. 정확한 수치는 변동 — 모든 수치는 확인 필요이고 시점에 따라 달라진다. 의미는 모든 모델이 비슷한 수준에서 깨진다 는 frame 자체.

TAB · BackendBench overall pass rate (개념적, 발표 시점)op 단위 strict pass
모델op pass ratestrictpartial
Claude 3.5 Sonnet
62%21%
GPT-4o
54%24%
DeepSeek-V3
51%26%
Llama 3.1 70B
38%22%
Qwen 2.5 Coder 32B
45%25%
Codestral 22B
34%20%
strict = 모든 OpInfo sample pass. partial = ≥80% pass. 최상위 모델도 strict 60% 대 — 즉 LLM 이 짜준 커널 10 개 중 4 개는 어디선가 silent fail. production 에 그대로 풀어둘 수 없다는 뜻.

강의에서 의외였던 관찰 두 가지.

§ 07자동 수정 워크플로· auto-fix loop

"실패 → 에러 메시지 → 재시도" 의 closed loop

BackendBench 의 두 번째 가치 — LLM 이 만든 커널을 자동으로 다시 짜게 만드는 loop의 base. failure mode 를 prompt 에 다시 넣으면 같은 모델이 두 번째 시도에서 더 많이 통과한다.

FIG · auto-fix loop실패 evidence 를 prompt 로 다시
prompt: write Triton kernel LLM draft kernel BackendBench evaluate PASS? FAIL prompt 에 failure evidence 추가 (failed input · expected vs got · stack trace) retry up to N times — typical N = 3-5
강의에서 Mark 가 보여준 핵심 결과 — 3 회 retry 면 strict pass rate 가 ~15-25%p 상승 (확인 필요). 즉 LLM 이 한 번에는 못 짜지만 evidence 를 받으면 빠르게 수렴한다.

이 loop 가 시사하는 더 큰 그림 — LLM + evaluation harness 의 closed loop 가 새로운 컴파일러 형태가 된다. PyTorch op 의 reference 가 spec 이고, LLM 이 그 spec 을 만족하는 커널을 generate-and-test 로 찾는다. BackendBench 는 그 loop 의 testing 부분.

test-time compute 와의 연결

이 워크플로는 OpenAI o1 / Claude 의 reasoning model 과 자연스럽게 결합한다. failure evidence 를 reasoning 안으로 넣고 더 깊이 생각하게 만들면 단일 op 의 strict pass rate 가 더 올라간다 — 강의 시점에서 초기 실험이 진행 중 (확인 필요).

§ 08채택 사례· who's using it

research benchmark 에서 PyTorch 의 CI 로

채택의 큰 그림 — BackendBench 는 leaderboard 라기보다 standard test harness 로 자리잡는 길로 가고 있다. "어떤 모델이 1등" 이 아니라 "이 모델이 production-safe 한가" 의 답.

§ 09다음 단계· future work

강의가 명시적으로 던진 다음 질문들

§ 10기억할 메모· key takeaways

다시 열었을 때 손에 잡혀야 할 것

silent failure
LLM 이 짠 커널이 atol 통과 input 에서는 정답, edge case 에선 wrong. production 학습 며칠 지나야 드러남.
OpInfo 위에 얹는다
PyTorch 가 이미 가진 op 별 sample generator + reference 를 그대로 활용. LLM 커널 검증의 base.
conservative tolerance
fp32 atol 1e-5, fp16 1e-3, bf16 1e-2. NaN/Inf mask 일치까지 강제.
failure taxonomy
dtype promotion · broadcasting · 수치 안정성 · backward 부호 · causal off-by-one · stride · edge shape · launch grid.
leaderboard 의 의미
"누가 1등" 이 아닌 "production safe?" 의 척도. 최상위 모델도 strict 60% 대 — 절반 가까이 silent.
auto-fix loop
실패 evidence 를 prompt 에 추가하고 LLM 에 재요청. 3 회면 strict pass +15-25%p (확인 필요).
vs KernelBench
KernelBench 는 한 op 의 best 속도. BackendBench 는 op 전체 cover 의 정확성. 보완재.
future
fused op spec, perf+correct 통합, multi-GPU collective, end-to-end loss drift 검증.
YouTube강의 영상 (확인 필요)
Repogithub.com/meta-pytorch/backendbench (또는 PyTorch org · 확인 필요)
관련KernelBench · GPU-Bench · TritonBench

손에 새기기 — 실습 시퀀스

  1. BackendBench 직접 돌리기 — repo 를 clone, 단일 op (예: aten::softmax) 만 target 으로 한 번 통과. fail 한 sample 의 input shape 을 직접 본다.
  2. LLM 한 모델 평가 — Claude 또는 GPT-4o 에 같은 prompt 로 5 개 op 의 Triton 커널 짜게 한다. BackendBench 결과 표를 직접 만든다.
  3. auto-fix loop 직접 돌려보기 — failure evidence 를 prompt 에 넣고 retry 3 회. pass rate 변화 그래프.
  4. failure taxonomy 분류 — § 05 의 5 가지 mode 중 자기가 받은 fail 들이 어디로 분포되는지 직접 분류.
  5. tolerance 변화 실험 — atol 를 1e-5 / 1e-3 / 1e-2 로 바꿔가며 같은 커널의 pass rate 변화. 어떤 op 가 tolerance 에 민감한지.
  6. OpInfo 직접 읽기 — PyTorch repo 의 torch/testing/_internal/common_methods_invocations.py 를 열어서 한 op 의 sample generator 가 어떤 shape/dtype 을 만드는지 본다.
§ 11다른 강의로의 연결· connections

이 강의가 시리즈 안에서 어디로 이어지는가

§ 12열린 질문· open questions

다음에 다시 들었을 때 직접 검증해야 할 것들

검증 메모

BackendBench 는 활발히 개발 중인 도구. 강의 시점의 기능과 현재 repo 의 기능이 다를 수 있다. 이 노트의 대부분은 "PyTorch 기반 op 검증 harness" 의 일반론 + 강의 metadata. 정확한 API, 명령행 사용법은 repo README 직접 참조.

← Lecture 075GPU Programming Fundamentals + ThunderKittens Lecture 077 →DSLs for GPU Kernels — Tri Dao