정책화된 커널 내부 에필로그 호출 체계 policy-based epilogue invocation

Epilogue 가 하나의 연산이 아니라 정책 (policy), trait, functor, 커널 호출자 (launcher), API layer 로 층위화된 구조

 

전반적 호출 체계 개요

flowchart LR
    A[API Layer<br>epilogue.h / dtype.h] --> B[Policy Interface<br>ep_policy.cuh / ep_traits.cuh]
    B --> C[Functor Layer<br>ep_functors.cuh]
    C --> D[Kernel Policy<br>ep_kernel_policy.cuh]
    D --> E[Kernel Impl<br>epilogue_kernels_policy.cu]
    E --> F[Launcher<br>epilogue_launcher_policy.cu]
    F --> G[Python Binding<br>epilogue_pybind.cpp]

 

1. api/

dtype.h

• 템플릿 타입 매핑
• 내부 functor, policy 계층이 공통적으로 사용하는 정적 타입 추론 유틸

epilogue.h

• Epilogue API entry - 커널 호출자들이 직접 include 하는 상위 헤더
• 내부적으로 EpilogueParams 를 생성 후 ep_policy 기반으로 호출 경로를 결정
• template<typename Policy, typename Params> 구조로 각 정책을 주입받음

 

2. kernels/policy/

ep_policy.cuh

• 에필로그의 정책 클래스 policy class 핵심
• 이 정책은 kernel 내부에서 에필로그 연산 조합을 결정하는 blueprint 역할

ep_traits.cuh

• 정책별 trait 정의
• 커널 compile-time specialization 을 유도

ep_functors.cuh

• 실제 수학적 연산 functor 집합
• Policy 에서 이 functor 들을 조합해 호출한다.

ep_apply.cuh

• functor 들을 순차/병렬 적용하는 장치
• 즉, Epilogue 의 실제 호출 순서를 담당

ep_kernel_policy.cuh

• 커널 수준에서의 정책 연결부

 

3. kernels/

epilogue_kernels_policy.cu

• 위 정책/trait/functor 들을 실제 instantiation 해 등록하는 파일
• __global__ 커널 정의 포함
• NVTX 범위나 arch 별 specialization 도 여기서

epilogue_params.cuh

• 런타임 파라미터 구조체 정의
• policy.apply 시 이 구조를 전달

philox.cuh

• dropout 등 stochastic 연산을 위한 philox RNG 구현

 

4. launcher/epilogue_launcher_policy.cu

• 런타임 진입점
• dispatch_policy(params) 로 정책을 분기시키고, kernel launch configuration 을 설정

 

5. pybind/epilogue_pybind.cpp

• Python binding
• launch_epilogue() 를 PyBind11 로 export 하여 Python 에서 접근 가능