はじめに

CUDAプログラミングを行う場合、たいていの場合はRuntime APIが利用されると思います。
Runtime APIより細かい制御をしたい場合や、ホストコードとデバイスコードをプロジェクト上分けたいときなど、Driver APIを利用したいときのため、基本的な実行方法を整理しておきます。

実行環境

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

kernelコード

kernelコードはCUDAコンパイラ(nvcc)を使います。
kernelはNVIDIAプロジェクトに自動生成されるサンプル(int型配列同士の加算
)をそのまま利用します。
- 新規作成→プロジェクト→NVIDIA→CUDA 9.0→CUDA 9.0 Runtime

kernel.cu

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でkernelを実行する場合は、cuModuleLoadでptxファイルを外部モジュールとしてロードし、cuModuleGetFunctionで関数名をキーにエントリポイントを取得します。
その後、cuLaunchKernelにgridやthreadblockの設定とともに渡すことで、kernelが実行されます。

ちなみに、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,