驅動程序 API 在 cuda 動態庫（cuda.dll或cuda.so）中實現，該庫在安裝設備驅動程序期間復制到系統上。 它的所有入口點都以 cu 為前綴。

它是一個基于句柄的命令式 API：大多數對象都由不透明的句柄引用，這些句柄可以指定給函數來操作對象。

驅動程序 API 中可用的對象匯總在下表中。Table 16. Objects Available in the CUDA Driver API

在調用驅動程序 API 的任何函數之前，必須使用cuInit()初始化驅動程序 API。 然后必須創建一個附加到特定設備的 CUDA 上下文，并使其成為當前調用主機線程，如上下文中所述。

在 CUDA 上下文中，內核作為 PTX 或二進制對象由主機代碼顯式加載，如模塊中所述。 因此，用 C++ 編寫的內核必須單獨編譯成 PTX 或二進制對象。 內核使用 API 入口點啟動，如內核執行中所述。

任何想要在未來設備架構上運行的應用程序都必須加載 PTX，而不是二進制代碼。 這是因為二進制代碼是特定于體系結構的，因此與未來的體系結構不兼容，而 PTX 代碼在加載時由設備驅動程序編譯為二進制代碼。

以下是使用驅動程序 API 編寫的內核示例的主機代碼：

int main()
{
    int N = ...;
    size_t size = N * sizeof(float);

    // Allocate input vectors h_A and h_B in host memory
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);

    // Initialize input vectors
    ...

    // Initialize
    cuInit(0);

    // Get number of devices supporting CUDA
    int deviceCount = 0;
    cuDeviceGetCount(&deviceCount);
    if (deviceCount == 0) {
        printf("There is no device supporting CUDA.\n");
        exit (0);
    }

    // Get handle for device 0
    CUdevice cuDevice;
    cuDeviceGet(&cuDevice, 0);

    // Create context
    CUcontext cuContext;
    cuCtxCreate(&cuContext, 0, cuDevice);

    // Create module from binary file
    CUmodule cuModule;
    cuModuleLoad(&cuModule, "VecAdd.ptx");

    // Allocate vectors in device memory
    CUdeviceptr d_A;
    cuMemAlloc(&d_A, size);
    CUdeviceptr d_B;
    cuMemAlloc(&d_B, size);
    CUdeviceptr d_C;
    cuMemAlloc(&d_C, size);

    // Copy vectors from host memory to device memory
    cuMemcpyHtoD(d_A, h_A, size);
    cuMemcpyHtoD(d_B, h_B, size);

    // Get function handle from module
    CUfunction vecAdd;
    cuModuleGetFunction(&vecAdd, cuModule, "VecAdd");

    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid =
            (N + threadsPerBlock - 1) / threadsPerBlock;
    void* args[] = { &d_A, &d_B, &d_C, &N };
    cuLaunchKernel(vecAdd,
                   blocksPerGrid, 1, 1, threadsPerBlock, 1, 1,
                   0, 0, args, 0);

    ...
}

完整的代碼可以在 vectorAddDrv CUDA 示例中找到。

L.1. Context

CUDA 上下文類似于 CPU 進程。驅動 API 中執行的所有資源和操作都封裝在 CUDA 上下文中，當上下文被銷毀時，系統會自動清理這些資源。除了模塊和紋理或表面引用等對象外，每個上下文都有自己獨特的地址空間。因此，來自不同上下文的 CUdeviceptr 值引用不同的內存位置。

主機線程一次可能只有一個設備上下文當前。當使用 cuCtxCreate（） 創建上下文時，它對調用主機線程是當前的。如果有效上下文不是線程當前的，則在上下文中操作的 CUDA 函數（大多數不涉及設備枚舉或上下文管理的函數）將返回 CUDA_ERROR_INVALID_CONTEXT。

每個主機線程都有一堆當前上下文。 cuCtxCreate（） 將新上下文推送到堆棧頂部?？梢哉{用 cuCtxPopCurrent（） 將上下文與主機線程分離。然后上下文是“浮動的”，并且可以作為任何主機線程的當前上下文推送。 cuCtxPopCurrent（） 還會恢復先前的當前上下文（如果有）。

還為每個上下文維護使用計數。 cuCtxCreate（） 創建使用計數為 1 的上下文。cuCtxAttach（） 增加使用計數，而 cuCtxDetach（） 減少使用計數。當調用 cuCtxDetach（） 或 cuCtxDestroy（） 時使用計數變為 0，上下文將被銷毀。

驅動程序 API 可與運行時互操作，并且可以通過 cuDevicePrimaryCtxRetain（） 從驅動程序 API 訪問由運行時管理的主上下文（參見初始化）。

使用計數有助于在相同上下文中運行的第三方編寫的代碼之間的互操作性。例如，如果加載三個庫以使用相同的上下文，則每個庫將調用 cuCtxAttach（） 來增加使用計數，并在庫使用上下文完成時調用 cuCtxDetach（） 來減少使用計數。對于大多數庫，預計應用程序會在加載或初始化庫之前創建上下文；這樣，應用程序可以使用自己的啟發式方法創建上下文，并且庫只需對傳遞給它的上下文進行操作。希望創建自己的上下文的庫（可能會或可能沒有創建自己的上下文的 API 客戶端不知道）將使用 cuCtxPushCurrent（） 和 cuCtxPopCurrent（），如下圖所示。

L.2. Module

模塊是設備代碼和數據的動態可加載包，類似于 Windows 中的 DLL，由 nvcc 輸出（請參閱使用 NVCC 編譯）。 所有符號的名稱，包括函數、全局變量和紋理或表面引用，都在模塊范圍內維護，以便獨立第三方編寫的模塊可以在相同的 CUDA 上下文中互操作。

此代碼示例加載一個模塊并檢索某個內核的句柄：

CUmodule cuModule;
cuModuleLoad(&cuModule, "myModule.ptx");
CUfunction myKernel;
cuModuleGetFunction(&myKernel, cuModule, "MyKernel");

此代碼示例從 PTX 代碼編譯和加載新模塊并解析編譯錯誤：

#define BUFFER_SIZE 8192
CUmodule cuModule;
CUjit_option options[3];
void* values[3];
char* PTXCode = "some PTX code";
char error_log[BUFFER_SIZE];
int err;
options[0] = CU_JIT_ERROR_LOG_BUFFER;
values[0]  = (void*)error_log;
options[1] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
values[1]  = (void*)BUFFER_SIZE;
options[2] = CU_JIT_TARGET_FROM_CUCONTEXT;
values[2]  = 0;
err = cuModuleLoadDataEx(&cuModule, PTXCode, 3, options, values);
if (err != CUDA_SUCCESS)
    printf("Link error:\n%s\n", error_log);

此代碼示例從多個 PTX 代碼編譯、鏈接和加載新模塊，并解析鏈接和編譯錯誤：

#define BUFFER_SIZE 8192
CUmodule cuModule;
CUjit_option options[6];
void* values[6];
float walltime;
char error_log[BUFFER_SIZE], info_log[BUFFER_SIZE];
char* PTXCode0 = "some PTX code";
char* PTXCode1 = "some other PTX code";
CUlinkState linkState;
int err;
void* cubin;
size_t cubinSize;
options[0] = CU_JIT_WALL_TIME;
values[0] = (void*)&walltime;
options[1] = CU_JIT_INFO_LOG_BUFFER;
values[1] = (void*)info_log;
options[2] = CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES;
values[2] = (void*)BUFFER_SIZE;
options[3] = CU_JIT_ERROR_LOG_BUFFER;
values[3] = (void*)error_log;
options[4] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES;
values[4] = (void*)BUFFER_SIZE;
options[5] = CU_JIT_LOG_VERBOSE;
values[5] = (void*)1;
cuLinkCreate(6, options, values, &linkState);
err = cuLinkAddData(linkState, CU_JIT_INPUT_PTX,
                    (void*)PTXCode0, strlen(PTXCode0) + 1, 0, 0, 0, 0);
if (err != CUDA_SUCCESS)
    printf("Link error:\n%s\n", error_log);
err = cuLinkAddData(linkState, CU_JIT_INPUT_PTX,
                    (void*)PTXCode1, strlen(PTXCode1) + 1, 0, 0, 0, 0);
if (err != CUDA_SUCCESS)
    printf("Link error:\n%s\n", error_log);
cuLinkComplete(linkState, &cubin, &cubinSize);
printf("Link completed in %fms. Linker Output:\n%s\n", walltime, info_log);
cuModuleLoadData(cuModule, cubin);
cuLinkDestroy(linkState);

完整的代碼可以在 ptxjit CUDA 示例中找到。

L.3. Kernel Execution

cuLaunchKernel（） 啟動具有給定執行配置的內核。

參數可以作為指針數組（在 cuLaunchKernel（） 的最后一個參數旁邊）傳遞，其中第 n 個指針對應于第 n 個參數并指向從中復制參數的內存區域，或者作為額外選項之一（ cuLaunchKernel（）） 的最后一個參數。

當參數作為額外選項（CU_LAUNCH_PARAM_BUFFER_POINTER 選項）傳遞時，它們作為指向單個緩沖區的指針傳遞，在該緩沖區中，通過匹配設備代碼中每個參數類型的對齊要求，參數被假定為彼此正確偏移。

表 4 列出了內置向量類型的設備代碼中的對齊要求。對于所有其他基本類型，設備代碼中的對齊要求與主機代碼中的對齊要求相匹配，因此可以使用 __alignof（） 獲得。唯一的例外是當宿主編譯器在一個字邊界而不是兩個字邊界上對齊 double 和 long long（在 64 位系統上為 long）（例如，使用 gcc 的編譯標志 -mno-align-double ） 因為在設備代碼中，這些類型總是在兩個字的邊界上對齊。

CUdeviceptr是一個整數，但是代表一個指針，所以它的對齊要求是__alignof（void*）。

以下代碼示例使用宏 （ALIGN_UP（）） 調整每個參數的偏移量以滿足其對齊要求，并使用另一個宏 （ADD_TO_PARAM_BUFFER（）） 將每個參數添加到傳遞給 CU_LAUNCH_PARAM_BUFFER_POINTER 選項的參數緩沖區。

#define ALIGN_UP(offset, alignment) \
      (offset) = ((offset) + (alignment) - 1) & ~((alignment) - 1)

char paramBuffer[1024];
size_t paramBufferSize = 0;

#define ADD_TO_PARAM_BUFFER(value, alignment)                   \
    do {                                                        \
        paramBufferSize = ALIGN_UP(paramBufferSize, alignment); \
        memcpy(paramBuffer + paramBufferSize,                   \
               &(value), sizeof(value));                        \
        paramBufferSize += sizeof(value);                       \
    } while (0)

int i;
ADD_TO_PARAM_BUFFER(i, __alignof(i));
float4 f4;
ADD_TO_PARAM_BUFFER(f4, 16); // float4's alignment is 16
char c;
ADD_TO_PARAM_BUFFER(c, __alignof(c));
float f;
ADD_TO_PARAM_BUFFER(f, __alignof(f));
CUdeviceptr devPtr;
ADD_TO_PARAM_BUFFER(devPtr, __alignof(devPtr));
float2 f2;
ADD_TO_PARAM_BUFFER(f2, 8); // float2's alignment is 8

void* extra[] = {
    CU_LAUNCH_PARAM_BUFFER_POINTER, paramBuffer,
    CU_LAUNCH_PARAM_BUFFER_SIZE,    ¶mBufferSize,
    CU_LAUNCH_PARAM_END
};
cuLaunchKernel(cuFunction,
               blockWidth, blockHeight, blockDepth,
               gridWidth, gridHeight, gridDepth,
               0, 0, 0, extra);

結構的對齊要求等于其字段的對齊要求的最大值。 因此，包含內置向量類型 CUdeviceptr 或未對齊的 double 和 long long 的結構的對齊要求可能在設備代碼和主機代碼之間有所不同。 這種結構也可以用不同的方式填充。 例如，以下結構在主機代碼中根本不填充，但在設備代碼中填充了字段 f 之后的 12 個字節，因為字段 f4 的對齊要求是 16。

typedef struct {
    float  f;
    float4 f4;
} myStruct;

CUdeviceptr devPtr;
float* d_data;

// Allocation using driver API
cuMemAlloc(&devPtr, size);
d_data = (float*)devPtr;

// Allocation using runtime API
cudaMalloc(&d_data, size);
devPtr = (CUdeviceptr)d_data;

    typedef CUresult (CUDAAPI *PFN_cuMemAlloc_v3020)(CUdeviceptr_v2 *dptr, size_t bytesize);
    typedef CUresult (CUDAAPI *PFN_cuMemAlloc_v2000)(CUdeviceptr_v1 *dptr, unsigned int bytesize);

    PFN_cuMemAlloc_v3020 pfn_cuMemAlloc_v2;
    PFN_cuMemAlloc_v2000 pfn_cuMemAlloc_v1;

  PFN_cuMemAlloc pfn_cuMemAlloc;

    // cuda.h
    CUresult CUDAAPI cuStreamBeginCapture(CUstream hStream);
    CUresult CUDAAPI cuStreamBeginCapture_v2(CUstream hStream, CUstreamCaptureMode mode);
            
    // cudaTypedefs.h
    typedef CUresult (CUDAAPI *PFN_cuStreamBeginCapture_v10000)(CUstream hStream);
    typedef CUresult (CUDAAPI *PFN_cuStreamBeginCapture_v10010)(CUstream hStream, CUstreamCaptureMode mode);

    #include 
    
    // Declare the entry points for cuStreamBeginCapture
    PFN_cuStreamBeginCapture_v10000 pfn_cuStreamBeginCapture_v1;
    PFN_cuStreamBeginCapture_v10010 pfn_cuStreamBeginCapture_v2;
    
    // Get the function pointer to the cuStreamBeginCapture driver symbol
    cuGetProcAddress("cuStreamBeginCapture", &pfn_cuStreamBeginCapture_v1, 10000, CU_GET_PROC_ADDRESS_DEFAULT);
    // Get the function pointer to the cuStreamBeginCapture_v2 driver symbol
    cuGetProcAddress("cuStreamBeginCapture", &pfn_cuStreamBeginCapture_v2, 10010, CU_GET_PROC_ADDRESS_DEFAULT);

    // Assuming we are using CUDA 11.3 Toolkit

    #include 
    
    // Declare the entry point
    PFN_cuStreamBeginCapture pfn_cuStreamBeginCapture_latest;
    
    // Intialize the entry point. Specifying CUDA_VERSION will give the function pointer to the
    // cuStreamBeginCapture_v2 symbol since it is latest version on CUDA 11.3.
    cuGetProcAddress("cuStreamBeginCapture", &pfn_cuStreamBeginCapture_latest, CUDA_VERSION, CU_GET_PROC_ADDRESS_DEFAULT);
  

    #include 
    
    // Declare the entry point
    PFN_cuMemAllocAsync pfn_cuMemAllocAsync;
    
    // Intialize the entry point. Assuming CUDA runtime version >= 11.2
    cudaGetDriverEntryPoint("cuMemAllocAsync", &pfn_cuMemAllocAsync, cudaEnableDefault);
    
    // Call the entry point
    pfn_cuMemAllocAsync(...);

    int main()
    {
        // Assuming we have CUDA 12.0 driver installed.

        // Manually define the prototype as cudaTypedefs.h in CUDA 11.3 does not have the cuFoo typedef
        typedef CUresult (CUDAAPI *PFN_cuFoo)(...);
        PFN_cuFoo pfn_cuFoo = NULL;
        
        // Get the address for cuFoo API using cuGetProcAddress. Specify CUDA version as
        // 12000 since cuFoo was introduced then or get the driver version dynamically
        // using cuDriverGetVersion 
        int driverVersion;
        cuDriverGetVersion(&driverVersion);
        cuGetProcAddress("cuFoo", &pfn_cuFoo, driverVersion, CU_GET_PROC_ADDRESS_DEFAULT);
        
        if (pfn_cuFoo) {
            pfn_cuFoo(...);
        }
        else {
            printf("Cannot retrieve the address to cuFoo. Check if the latest driver for CUDA 12.0 is installed.\n");
            assert(0);
        }
        
        // rest of code here
        
    

關于作者

Ken He 是 NVIDIA 企業級開發者社區經理 & 高級講師，擁有多年的 GPU 和人工智能開發經驗。自 2017 年加入 NVIDIA 開發者社區以來，完成過上百場培訓，幫助上萬個開發者了解人工智能和 GPU 編程開發。在計算機視覺，高性能計算領域完成過多個獨立項目。并且，在機器人和無人機領域，有過豐富的研發經驗。對于圖像識別，目標的檢測與跟蹤完成過多種解決方案。曾經參與 GPU 版氣象模式GRAPES，是其主要研發者。

審核編輯：郭婷