[빌드] CUDA에서 Separate Compilation과 Linking 이슈

최근 PTX 어셈블리를 활용하여 실행 중인 SM 확인하기 (tistory.com) 이라는 글을 보고서, Debug 용으로 실행중인 Kernal의 Warp나 Thread 정보들을 출력할 수 있는 유틸 함수들을 추가하면 좋겠다는 생각을 하게 되었습니다.

그래서 기존에 헤더로만 구성되어 있던 util 코드를 소스코드와 헤더로 분리하고 위에 블로그에서 아래의 코드들을 복사해서 가져왔습니다.

// utils.cu

#include "utils.cuh"

__device__ uint32_t __mysmid() {
    uint32_t smid;
    asm volatile("mov.u32 %0, %%smid;" : "=r"(smid));
    return smid;
}

__device__ uint32_t __mywarpid() {
    uint32_t warpid;
    asm volatile("mov.u32 %0, %%warpid;" : "=r"(warpid));
    return warpid;
}
  
__device__ uint32_t __mylaneid() {
    uint32_t laneid;
    asm volatile("mov.u32 %0, %%laneid;" : "=r"(laneid));
    return laneid;
}

그런데 아래와 같은 compile 오류를 내면서 빌드가 되지 않았습니다. 

혹시나 내가 CMakeLists.txt에 소스 파일을 추가하는 것을 까먹은 것인가 생각했는데 그게 아니었고, 아래의 stackoverflow글에서 확인해 볼 수 있었습니다. 

c++ - CUDA __device__ Unresolved extern function - Stack Overflow

CUDA __device__ Unresolved extern function

원인은 __device__ 함수와 이를 호출하는 __global__함수가 별도의 컴파일 단위에서 정의되었기 때문이었습니다. 따라서 --device-c 혹은 -dc 플래그를 추가하여, 재배치 가능 장치 코드 모드를 명시적으로 활성화하거나 정의를 동일한 단위로 이동해야 했습니다.

하지만 위의 내용에 따라서 CMakeLists.txt에서 CMAKE_CUDA_FLAGS에 -dc 플래그를 추가할 경우 다른 에러가 발생합니다.

이 문제의 원인은 생성된 glue code가 제대로 링킹되지 않아서 발생한 이슈였습니다.

따라서 아래 이미지 처럼 두 가지 옵션을 활성화해야 합니다. 단 컴파일러 버전이 CUDA 5.0 이상이여야 하고 Compute Capability가 2.0 이상이여야 합니다. 그리고 이를 위해 Compute Capability를 설정하는 부분이 옵션을 설정하는 부분보다 먼저와야 합니다.
(단 이때 이전에 설정했던 -dc 플래그는 제거해줘야 합니다.)

위의 이슈에 대한 내용은 이 링크에서 더 자세한 확인할 수 있습니다.

참고로, Separate Compilation을 하기 위해 위의 옵션을 켰을 경우 기존 보다 컴파일 속도가 느려질 수 있습니다.  왜냐하면 호출 부분과 함수가 다른 컴파일 단위에 존재하는 경우 ABI를 완전히 준수해야 하고 이것이 최적화를 방해하기 때문입니다.

 

또한 nvcc 컴파일에 대한 전반적인 내용은 아래의 링크에서 확인할 수 있습니다.

Separate Compilation and Linking of CUDA C++ Device Code | NVIDIA Technical Blog

Separate Compilation and Linking of CUDA C++ Device Code | NVIDIA Technical Blog

1. Introduction — cuda-compiler-driver-nvcc 12.1 documentation (nvidia.com)

NVIDIA CUDA Compiler Driver NVCC

 

마지막으로, 제가 작업한 내용들은 아래의 github 페이지 전체 코드로 확인해 볼 수 있습니다.

hotstone1993/Cuda_Study (github.com)

GitHub - hotstone1993/Cuda_Study