소개

CUDA 프로그래밍을 할 때 대부분의 경우 Runtime API가 사용된다고 생각합니다.
Runtime API보다 세밀한 제어를 하고 싶은 경우나, 호스트 코드와 디바이스 코드를 프로젝트 상으로 나누고 싶을 때 등, Driver API를 이용하고 싶을 때를 위해, 기본적인 실행 방법을 정리해 둡니다.

실행 환경

Windows 10 Home
Microsoft Visual Studio Community 2017
CUDA 9.0
NVIDIA GeForce GPU

커널 코드

커널 코드는 CUDA 컴파일러 (nvcc)를 사용합니다.
커널은 NVIDIA 프로젝트에 자동 생성되는 샘플 (int형 배열끼리의 가산
)을 그대로 이용합니다.
- 신규작성→프로젝트→NVIDIA→CUDA 9.0→CUDA 9.0 Runtime

털 r 값 l. 구

extern "C"
{
    __global__ void addKernel(int *c, const int *a, const int *b)
    {
        int i = threadIdx.x;
        c[i] = a[i] + b[i];
    }
}

프로젝트 속성에서 ptx 파일이 출력되도록 합니다.
컴파일하면 kernel.ptx가 출력됩니다.

호스트 코드

호스트 코드는 VC 컴파일러를 사용합니다.
- 신규 작성 → 프로젝트 → Visual C++ → Win32 → Win32 콘솔 애플리케이션

Driver API에서 커널을 실행하는 경우 cuModuleLoad 에서 ptx 파일을 외부 모듈로 로드하고 cuModuleGetFunction 에서 함수 이름을 키로 엔트리 포인트를 가져옵니다.
그런 다음 cuLaunchKernel에 grid 및 threadblock 설정과 함께 전달하여 커널을 실행합니다.

덧붙여서, CUDA 9로부터 CUDA_LAUNCH_PARAMS 구조체가 추가되고 있어, cuLaunchKernel 의 인수 상당의 값을 구조체를 해 정리해 둘 수 있습니다만, cuLaunchKernel 에는 건네줄 수 없습니다.
현재는 멀티 디바이스용 cuLaunchCooperativeKernelMultiDevice 밖에 사용할 수 없는 것 같습니다.

main.cpp

#include "cuda.h"
#include <string>
#include <iostream>

//! CUresultエラーコード→エラーメッセージ
std::string GetErrorName(const CUresult& cuResult)
{
    const char* pStr;
    if (cuGetErrorName(cuResult, &pStr) != CUDA_ERROR_INVALID_VALUE)
    {
        return std::string(pStr);
    }

    return std::string();
}

int main()
{
    CUresult result = CUDA_SUCCESS;

    result = cuInit(0);
    if (result != CUDA_SUCCESS)
    {
        std::cout << GetErrorName(result) << std::endl;
        return 1;
    }

    int deviceNum = 0;
    result = cuDeviceGetCount(&deviceNum);
    if (result != CUDA_SUCCESS)
    {
        std::cout << GetErrorName(result) << std::endl;
        return 1;
    }
    if (deviceNum < 1)
    {
        std::cout << "No Device."<< std::endl;
        return 1;
    }

    CUdevice deviceID = 0;
    CUcontext context;
    result = cuCtxCreate(&context, CU_CTX_SCHED_AUTO, deviceID);
    if (result != CUDA_SUCCESS) 
    {
        std::cout << GetErrorName(result) << std::endl;
        return 1;
    }

    cuCtxPushCurrent(context);

    //! kernel用入出力データ作成
    constexpr size_t length = 10;
    CUdeviceptr arrayAd, arrayBd, arrayCd;
    cuMemAlloc(&arrayAd, sizeof(int)*length);
    cuMemAlloc(&arrayBd, sizeof(int)*length);
    cuMemAlloc(&arrayCd, sizeof(int)*length);

    constexpr int arrayAh[length] = { 100, 100, 100, 100, 100, 100, 100, 100, 100, 100 };
    constexpr int arrayBh[length] = {   0,   1,   2,   3,   4,   5,   6,   7,   8,   9 };
    int arrayCh[length] = {};

    cuMemcpyHtoD(arrayAd, arrayAh, sizeof(int)*length);
    cuMemcpyHtoD(arrayBd, arrayBh, sizeof(int)*length);

    //! ptxファイルから関数のエントリポイントを取得
    CUmodule module;
    result = cuModuleLoad(&module, "kernel.ptx"); // 環境に合わせてパスを設定
    if (result != CUDA_SUCCESS)
    {
        std::cout << GetErrorName(result) << std::endl;
        return 1;
    }

    CUfunction addFunc;
    result = cuModuleGetFunction(&addFunc, module, "addKernel");
    if (result != CUDA_SUCCESS)
    {
        std::cout << GetErrorName(result) << std::endl;
        return 1;
    }

    //! kernel実行
    void* kernelArgs[] = { &arrayCd, &arrayAd, &arrayBd };
    result = cuLaunchKernel(addFunc, 1, 1, 1, length, 1, 1, 0, NULL, kernelArgs, 0);
    if (result != CUDA_SUCCESS)
    {
        std::cout << GetErrorName(result) << std::endl;
        return 1;
    }

    //! 結果確認
    cuMemcpyDtoH(arrayCh, arrayCd, sizeof(int)*length);

    for (int i = 0; i < length; i++)
    {
        std::cout << arrayCh[i] << ",";
    }
    std::cout << std::endl;

    cuModuleUnload(module);
    cuMemFree(arrayAd);
    cuMemFree(arrayBd);
    cuMemFree(arrayCd);
    cuCtxPopCurrent(&context);
    cuCtxDestroy(context);

    return 0;
}

실행 결과

100,101,102,103,104,105,106,107,108,109,