RDC + Device LTO(-dlto) 란??

여러 .cu/.cpp 파일로 쪼개서 빌드하더라도, 커널 내부의 __device__ 템플릿/인라인 코드가 TU(번역단위) 경계를 넘어 진짜로 인라인, 최적화되게 만들고 싶다면 CUDA 의 분리 컴파일(RDC) + 디바이스 LTO(-dlto) 를 같이 켜야 한다.

 

기본 빌드 = 단일 TU 에서만 인라인

• nvcc 는 같은 번역단위 (.cu) 안에 있는 __device__ __forceinline__ 함수/템플릿은 인라인한다.
• 하지만 코드를 여러 .cu 로 나누면, TU 경계 너머의 __device__ 함수는 보통 인라인이 안 됨, 
  • 호출 경계(콜) 이 남아 성능, 레지스터 배치가 최적이 아니게 될 수 있음

 

RDC (Relocatable Device Code, 분리 컴파일)

• 플래그 : -rdc=true 또는 CMake CMAKE_CUDA_SEPARABLE_COMPILATION ON
• 효과 : .cu 를 각각 개별 오브젝트로 컴파일하고 디바이스 링킹에서 서로의 __device__ 심볼을 연결할 수 있게 해줌
• 하지만 RDC 만 켜면 "연결" 은 되지만 인라인/IPO(번역단위 간 최적화) 는 제한적, 즉 함수 호출 경계는 남기 쉬움

 

Device LTO (디바이스 링크 타임 최적화)

• 플래그: -dlto (컴파일/링크 둘 다)
• 효과: 디바이스 코드도 LTO 를 적용해서 TU 경계를 넘어 인라인, DCE, 상수 전파 같은 최적화를 수행
• 즉, 여러 모듈로 분리해도, 마치 한 파일에 있는 것철머 고급 최적화가 가능

 

현재 Epilogeu, stBias, StAct 같은 정책/스테이지를 헤더로 두고, 여러 커널 TU에서 이걸 호출, 

모듈화를 하되 성능은 그대로 가져가려면, 

RDC 로 모듈을 나누고

Device LTo 로 TU 경계까지 최적화해야 동일 수준이 나온다