CUDA 컨텍스트 비의존 모듈 로딩

Reading Time: 3 minutes

대부분의 CUDA 개발자는 CUDA 컨텍스트에 디바이스 코드를 포함하는 모듈을 로드하기 위해 cuModuleLoad API와 이에 상응하는 API를 사용하는 것에 익숙합니다. 대부분의 상황에서는 모든 디바이스에 동일한 디바이스 코드를 로드하게 됩니다. 이를 위해서는 각 CUDA 컨텍스트에 장치 코드를 명시적으로 로드해야 합니다. 또한 컨텍스트 생성과 파괴를 제어하지 않는 라이브러리와 프레임워크는 모듈을 명시적으로 로드하고 언로드하기 위해 추적해야 합니다.

이 게시물에서는 이러한 문제를 해결하기 위해 CUDA 12.0에 도입된 컨텍스트에 구애받지 않는 로딩에 대해 설명합니다.

컨텍스트 의존 로딩(Context-dependent loading)

통상적으로 모듈 로딩은 항상 CUDA 컨텍스트와 연관이 있었습니다. 다음 코드 예제에서는 동일한 디바이스 코드를 두 디바이스에 로드한 다음 해당 디바이스에서 커널을 실행하는 전통적인 방법을 보여줍니다.


// Device 0
cuDeviceGet(&device0, 0);
cuDevicePrimaryCtxRetain(&ctx0, device0);
cuModuleLoad(&module0, “myModule.cubin”);
// Device 1
cuDeviceGet(&device1, 1);
cuDevicePrimaryCtxRetain(&ctx1, device1);
cuModuleLoad(&module1, “myModule.cubin”);

각 디바이스에서 커널을 실행하려면 다음 코드 예제에서 볼 수 있듯 모듈별 CUfunction을 검색해야 합니다:

// Device 0
cuModuleGetFuntion(&function0, module0, “myKernel”);
cuLaunchKernel(function0, …);
// Device 1
cuModuleGetFuntion(&function1, module1, “myKernel”);
cuLaunchKernel(function1, …);

이 경우 컨텍스트별/모듈별 유형을 검색하고 추적해야 하므로 애플리케이션의 코드 복잡성이 증가합니다. 게다가 cuModuleUnload API를 사용하여 각 모듈을 명시적으로 언로드해야 하죠.

이 문제는 라이브러리나 프레임워크에서 자체 모듈을 불러오기 위해 주로 CUDA 드라이버 API를 사용하는 경우 더욱 악화됩니다. 애플리케이션에서 소유한 컨텍스트의 수명에 대한 완전한 제어권이 없을 수도 있습니다.

// Application code

libraryInitialize();
cuDevicePrimaryCtxRetain(&ctx0, device0);
libraryFunc();
cuDevicePrimaryCtxRetain(&ctx0, device1);
libraryFunc();
libraryDeinitialize();

// Library code

libraryInitialize() {
  map<CUcontext, CUmodule> moduleContextMap;
}

libraryFunc() {
  cuCtxGetCurrent(&ctx);
  if (!moduleContextMap.contains(ctx)){
    cuModuleLoad(&module, “myModule.cubin”);
    moduleContextMap[ctx] = module;
  }
  else {
    module = moduleContextMap[ctx];
  }

  cuModuleGetFuntion(&function, module, “myKernel”);
  cuLaunchKernel(function, …);
}

libraryDeinitialize() {
  moduleContextMap.clear();
}

코드 예제에서 라이브러리는 새 컨텍스트를 확인하고, 컨텍스트에 모듈을 명시적으로 로드해야 합니다. 또한 모듈이 이미 컨텍스트에 로드되어 있는지 확인하기 위해 상태를 유지해야 합니다.

이상적인 상황에서라면 상태는 컨텍스트가 소멸된 후에야 해제할 수 있습니다. 하지만 라이브러리에 컨텍스트 수명에 대한 제어권이 없다면 이는 불가능합니다.

즉, 라이브러리 초기화가 해제될 때까지 리소스 해제가 지연되어야 한다는 뜻입니다. 이는 코드 복잡성을 증가시킬 뿐만 아니라, 라이브러리가 리소스를 필요 이상으로 길게 사용함으로써 애플리케이션의 다른 부분에서 해당 메모리를 사용하지 못하게 될 수 있습니다.

다른 대안 하나는 라이브러리와 프레임워크가 리소스 할당 및 정리를 위한 충분한 제어력을 갖출 수 있도록 사용자에게 추가적인 제약 조건을 강제하는 것입니다.

컨텍스트 비의존 로딩

CUDA 12.0에서는 cuLibrary* 및 cuKernel* API를 추가함으로써 컨텍스트 비의존 로딩을 도입하여 이러한 문제를 해결합니다. 컨텍스트 비의존 로딩을 사용할 경우 애플리케이션에서 컨텍스트를 생성하고 파괴하므로, 각 CUDA 컨텍스트에 대한 모듈 로딩 및 언로딩이 CUDA 드라이버에 의해 자동으로 수행됩니다.


// Load library
cuLibraryLoadFromFile(&library,“myModule.cubin”, …);
cuLibraryGetKernel(&kernel, library, “myKernel”);

// Launch kernel on the primary context of device 0
cuDevicePrimaryCtxRetain(&ctx0, device0);
cuLaunchKernel((CUkernel)kernel, …);

// Launch kernel on the primary context of device 1
cuDevicePrimaryCtxRetain(&ctx1, device1);
cuLaunchKernel((CUkernel)kernel, …);

// Unload library
cuLibraryUnload(library);

코드 예제에서 볼 수 있듯이 cuLibraryLoadFromFile API는 컨텍스트가 생성되거나 초기화되면 모듈 로딩을 처리합니다. 예제에서는 cuDevicePrimaryCtxRetain 중에 수행됩니다.

또한 이제 컨텍스트별 CUfunction을 유지할 필요 없이 CUkernel 컨텍스트 비의존 핸들을 사용하여 커널을 실행할 수 있습니다. cuLibraryGetKernel은 디바이스 함수 myKernel에 대한 컨텍스트 독립 핸들을 검색합니다. 그런 다음, CUkernel 컨텍스트 비의존 핸들을 지정하여 cuLaunchKernel을 통해 디바이스 함수를 실행할 수 있습니다. CUDA 드라이버는 해당 시점에서 활성 상태인 컨텍스트에 따라 적절한 컨텍스트에서 디바이스 함수를 실행합니다.

이제 라이브러리와 프레임워크는 초기화 및 초기화 해지 중 모듈을 한 번만 로드하고 언로드하면 됩니다.


// Application code

libraryInitialize();
cuDevicePrimaryCtxRetain(&ctx0, device0);
libraryFunc();
cuDevicePrimaryCtxRetain(&ctx0, device1);
libraryFunc();
libraryDeinitialize();

// Library code

libraryInitialize() {
  cuLibraryLoadFromFile(&library,“myModule.cubin”, …);
  cuLibraryGetKernel(&kernel, library, “myKernel”);
}

libraryFunc() {
  cuLaunchKernel((CUkernel)kernel, …);
}

libraryDeinitialize() {
  cuLibraryUnload(library);
}

라이브러리가 더 이상 컨텍스트별 상태를 유지하고 추적할 필요가 없다는 뜻이죠. 컨텍스트 비의존 로딩의 설계 덕분에, CUDA 드라이버가 모듈과 컨텍스트를 추적하고 모듈 로딩 및 언로딩 작업을 수행할 수 있습니다.

__managed__ 변수에 액세스

관리형 변수는 디바이스와 호스트 코드 모두에서 참조할 수 있습니다. 예를 들어, 관리형 변수의 주소는 쿼리할 수 있고, 디바이스 또는 호스트 함수에서 바로 읽거나 작성할 수 있습니다. 생성되어 속한 CUDA 컨텍스트만큼의 수명을 갖는 __device__ 변수와는 달리, 모듈에 속한 __managed__ 변수는 모든 CUDA 컨텍스트 또는 디바이스에서도 동일한 메모리를 가리킵니다.

CUDA 12.0 이전에는 드라이버 API를 통해 CUDA 컨텍스트에 전반에 걸쳐 고유한 관리형 변수로 핸들을 검색할 방법이 없었습니다. CUDA 12.0에는 새 드라이버 API cuLibraryGeManaged가 도입되어 CUDA 컨텍스트 전반에 걸쳐 고유한 핸들을 확보할 수 있습니다.

컨텍스트 비의존 로딩 시작하기

이 게시물에서는 CUDA 컨텍스트에 의존하지 않고 디바이스 코드를 로드할 수 있는 기능을 제공하는 새로운 CUDA 드라이버 API를 소개해 드렸습니다. 커널을 실행하기 위한 컨텍스트 비의존 핸들에 대해서도 이야기해 보았습니다. 둘을 함께 사용하면 기존의 로딩 메커니즘에 비해 GPU에서 코드를 로드하고 실행하기가 보다 간단해져, 코드 복잡성은 줄고 컨텍스트별 상태를 유지할 필요도 없습니다.

이러한 API 사용을 시작하려면 CUDA 드라이버 및 툴킷 버전 12 이상을 다운로드하세요. cuLibrary* 및 cuKernel* API에 대한 자세한 내용은 CUDA 드라이버 API 문서를 참조하시기 바랍니다.