驅(qū)動程序 API 在 cuda 動態(tài)庫（cuda.dll或cuda.so）中實現(xiàn)，該庫在安裝設(shè)備驅(qū)動程序期間復(fù)制到系統(tǒng)上。 它的所有入口點都以 cu 為前綴。

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

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

在調(diào)用驅(qū)動程序 API 的任何函數(shù)之前，必須使用cuInit()初始化驅(qū)動程序 API。 然后必須創(chuàng)建一個附加到特定設(shè)備的 CUDA 上下文，并使其成為當(dāng)前調(diào)用主機(jī)線程，如上下文中所述。

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

任何想要在未來設(shè)備架構(gòu)上運行的應(yīng)用程序都必須加載 PTX，而不是二進(jìn)制代碼。 這是因為二進(jìn)制代碼是特定于體系結(jié)構(gòu)的，因此與未來的體系結(jié)構(gòu)不兼容，而 PTX 代碼在加載時由設(shè)備驅(qū)動程序編譯為二進(jìn)制代碼。

以下是使用驅(qū)動程序 API 編寫的內(nèi)核示例的主機(jī)代碼：

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 進(jìn)程。驅(qū)動 API 中執(zhí)行的所有資源和操作都封裝在 CUDA 上下文中，當(dāng)上下文被銷毀時，系統(tǒng)會自動清理這些資源。除了模塊和紋理或表面引用等對象外，每個上下文都有自己獨特的地址空間。因此，來自不同上下文的 CUdeviceptr 值引用不同的內(nèi)存位置。

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

每個主機(jī)線程都有一堆當(dāng)前上下文。 cuCtxCreate（） 將新上下文推送到堆棧頂部。可以調(diào)用 cuCtxPopCurrent（） 將上下文與主機(jī)線程分離。然后上下文是“浮動的”，并且可以作為任何主機(jī)線程的當(dāng)前上下文推送。 cuCtxPopCurrent（） 還會恢復(fù)先前的當(dāng)前上下文（如果有）。

還為每個上下文維護(hù)使用計數(shù)。 cuCtxCreate（） 創(chuàng)建使用計數(shù)為 1 的上下文。cuCtxAttach（） 增加使用計數(shù)，而 cuCtxDetach（） 減少使用計數(shù)。當(dāng)調(diào)用 cuCtxDetach（） 或 cuCtxDestroy（） 時使用計數(shù)變?yōu)?0，上下文將被銷毀。

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

使用計數(shù)有助于在相同上下文中運行的第三方編寫的代碼之間的互操作性。例如，如果加載三個庫以使用相同的上下文，則每個庫將調(diào)用 cuCtxAttach（） 來增加使用計數(shù)，并在庫使用上下文完成時調(diào)用 cuCtxDetach（） 來減少使用計數(shù)。對于大多數(shù)庫，預(yù)計應(yīng)用程序會在加載或初始化庫之前創(chuàng)建上下文；這樣，應(yīng)用程序可以使用自己的啟發(fā)式方法創(chuàng)建上下文，并且?guī)熘恍鑼鬟f給它的上下文進(jìn)行操作。希望創(chuàng)建自己的上下文的庫（可能會或可能沒有創(chuàng)建自己的上下文的 API 客戶端不知道）將使用 cuCtxPushCurrent（） 和 cuCtxPopCurrent（），如下圖所示。

L.2. Module

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

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

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（） 啟動具有給定執(zhí)行配置的內(nèi)核。

參數(shù)可以作為指針數(shù)組（在 cuLaunchKernel（） 的最后一個參數(shù)旁邊）傳遞，其中第 n 個指針對應(yīng)于第 n 個參數(shù)并指向從中復(fù)制參數(shù)的內(nèi)存區(qū)域，或者作為額外選項之一（ cuLaunchKernel（）） 的最后一個參數(shù)。

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

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

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

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

#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);

結(jié)構(gòu)的對齊要求等于其字段的對齊要求的最大值。 因此，包含內(nèi)置向量類型 CUdeviceptr 或未對齊的 double 和 long long 的結(jié)構(gòu)的對齊要求可能在設(shè)備代碼和主機(jī)代碼之間有所不同。 這種結(jié)構(gòu)也可以用不同的方式填充。 例如，以下結(jié)構(gòu)在主機(jī)代碼中根本不填充，但在設(shè)備代碼中填充了字段 f 之后的 12 個字節(jié)，因為字段 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
        
    

關(guān)于作者

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

審核編輯：郭婷