본문으로 건너뛰기

[Triton] 소스 위치의 열 번호 수정

들어가며

Triton 컴파일러가 생성하는 IR에서 소스 위치(source location) 정보는 디버깅과 프로파일링에 필수적이다. 기존에는 열 번호에 두 가지 버그가 있었다: (1) 들여쓰기가 반영되지 않아 dedent된 커널 기준의 열 번호가 나왔고, (2) 중첩된 토큰 방문 시 부모 op의 위치가 자식 인자의 위치로 대체되었다.

핵심 코드 분석

Before

# 소스 위치 출력 예시
# 모든 줄이 column 0에서 시작하는 것처럼 보임
test_stdout = 'Hello From First Instruction of GPU Kernel: kernel1\ttest_gpuhello.py:17:4\n'

After

# 실제 들여쓰기를 반영한 열 번호
test_stdout = 'Hello From First Instruction of GPU Kernel: kernel1\ttest_gpuhello.py:17:5\n'

코드 생성기에서 Location 클래스에 get_name 메서드가 추가되어 named location 복원이 가능해졌다:

.def("get_name", [](Location &self) -> std::optional<std::string> {
    if (auto nameLoc = dyn_cast<NameLoc>(self)) {
        return nameLoc.getName().str();
    }
    return std::nullopt;
})

테스트에서 column 번호의 정확성을 검증한다:

# CHECK: #loc1 = loc({{.*}}:535:20)   # x_plus_1 할당의 정확한 열
# CHECK: #loc3 = loc({{.*}}:530:15)   # pid 할당의 정확한 열
# CHECK: #loc4 = loc({{.*}}:531:18)   # offset 할당의 정확한 열

왜 이게 좋은가

  • 정확한 디버깅: GPU 프로파일러(nsight 등)에서 정확한 소스 위치를 가리킬 수 있다.
  • 계측 정확성: GPU instrumentation hook이 올바른 소스 위치를 보고한다.
  • op 위치 보존: tl.load(ptr, mask=mask) 같은 호출에서 load op의 위치가 mask의 위치로 오염되는 문제가 해결되었다.

정리

+153/-63 변경으로, 컴파일러 인프라의 기본 품질을 높이는 중요한 PR이다. 디버깅 정보의 정확성이 개발자 경험에 직접적으로 영향을 미친다.

참고 자료

이 글은 AI(Claude)의 도움을 받아 작성되었습니다. 코드 분석 내용은 실제 PR diff를 기반으로 합니다.

댓글