CUDA 编程手册系列 附录L – CUDA底层驱动API（二）

L.4. Interoperability between Runtime and Driver APIs

应用程序可以将运行时 API 代码与驱动程序 API 代码混合。

如果通过驱动程序 API 创建上下文并使其成为当前上下文，则后续运行时调用将获取此上下文，而不是创建新上下文。

如果运行时已初始化（如 CUDA 运行时中提到的那样）， cuCtxGetCurrent() 可用于检索在初始化期间创建的上下文。 后续驱动程序 API 调用可以使用此上下文。

从运行时隐式创建的上下文称为主上下文（请参阅 初始化 ）。 它可以通过具有 主要上下文管理 功能的驱动程序 API 进行管理。

可以使用任一 API 分配和释放设备内存。 CUdeviceptr 可以转换为常规指针，反之亦然：

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;

特别是，这意味着使用驱动程序 API 编写的应用程序可以调用使用运行时 API 编写的库（例如 cuFFT、cuBLAS...）。

参考手册的设备和版本管理部分的所有功能都可以互换使用。

L.5. Driver Entry Point Access

L.5.1. Introduction

驱动程序入口点访问 API 提供了一种检索 CUDA 驱动程序函数地址的方法。 从 CUDA 11.3 开始，用户可以使用从这些 API 获得的函数指针调用可用的 CUDA 驱动程序 API。

这些 API 提供的功能类似于它们的对应物，POSIX 平台上的 dlsym 和 Windows 上的 GetProcAddress 。 提供的 API 将允许用户：

• 使用 CUDA 驱动程序 API 检索 驱动程序 函数的地址。
• 使用 CUDA 运行时 API 检索驱动程序函数的地址。
• 请求 CUDA 驱动程序函数的每线程默认流版本。 有关更多详细信息，请参阅 检索每个线程的默认流版本
• 使用较新的驱动程序访问旧工具包上的新 CUDA 功能。

L.5.2. Driver Function Typedefs

为了帮助检索 CUDA 驱动程序 API 入口点，CUDA 工具包提供对包含所有 CUDA 驱动程序 API 的函数指针定义的头文件的访问。 这些头文件与 CUDA Toolkit 一起安装，并且在工具包的 include/ 目录中可用。 下表总结了包含每个 CUDA API 头文件的 typedef 的头文件。

Table 17. Typedefs header files for CUDA driver APIs

上面的头文件本身并没有定义实际的函数指针； 他们为函数指针定义了 typedef 。 例如， cudaTypedefs.h 具有驱动 API cuMemAlloc 的以下 typedef ：

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

CUDA 驱动程序符号具有基于版本的命名方案，其名称中带有 _v* 扩展名，但第一个版本除外。 当特定 CUDA 驱动程序 API 的签名或语义发生变化时，我们会增加相应驱动程序符号的版本号。 对于 cuMemAlloc 驱动程序 API，第一个驱动程序符号名称是 cuMemAlloc ，下一个符号名称是 cuMemAlloc_v2 。 CUDA 2.0 (2000) 中引入的第一个版本的 typedef 是 PFN_cuMemAlloc_v2000 。 CUDA 3.2 (3020) 中引入的下一个版本的 typedef 是 PFN_cuMemAlloc_v3020 。

typedef 可用于更轻松地在代码中定义适当类型的函数指针：

PFN_cuMemAlloc_v3020 pfn_cuMemAlloc_v2;
    PFN_cuMemAlloc_v2000 pfn_cuMemAlloc_v1;

如果用户对 API 的特定版本感兴趣，则上述方法更可取。 此外，头文件中包含所有驱动程序符号的最新版本的预定义宏，这些驱动程序符号在安装的 CUDA 工具包发布时可用； 这些 typedef 没有 _v* 后缀。 对于 CUDA 11.3 工具包， cuMemAlloc_v2 是最新版本，所以我们也可以定义它的函数指针如下：

PFN_cuMemAlloc pfn_cuMemAlloc;

L.5.3. Driver Function Retrieval

使用驱动程序入口点访问 API 和适当的 typedef ，我们可以获得指向任何 CUDA 驱动程序 API 的函数指针。

L.5.3.1. Using the driver API

驱动程序 API 需要 CUDA 版本作为参数来获取请求的驱动程序符号的 ABI 兼容版本。 CUDA 驱动程序 API 有一个以 _v* 扩展名表示的按功能 ABI。 例如，考虑 cudaTypedefs.h 中 cuStreamBeginCapture 的版本及其对应的 typedef ：

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

从上述代码片段中的typedefs，版本后缀 _v10000 和 _v10010 表示上述API分别在 CUDA 10.0 和 CUDA 10.1 中引入。

#include <cudaTypedefs.h>
    // 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);

参考上面的代码片段，要检索到驱动程序 API cuStreamBeginCapture 的 _v1 版本的地址，CUDA 版本参数应该正好是 10.0 (10000) 。同样，用于检索 _v2 版本 API 的地址的 CUDA 版本应该是 10.1 (10010 )。为检索特定版本的驱动程序 API 指定更高的 CUDA 版本可能并不总是可移植的。例如，在此处使用 11030 仍会返回 _v2 符号，但如果在 CUDA 11.3 中发布假设的 _v3 版本，则当与 CUDA 11.3 驱动程序配对时， cuGetProcAddress API 将开始返回较新的 _v3 符号。由于 _v2 和 _v3 符号的 ABI 和函数签名可能不同，使用用于 _v2 符号的 _v10010 typedef 调用 _v3 函数将表现出未定义的行为。

要检索给定 CUDA 工具包的驱动程序 API 的最新版本，我们还可以指定 CUDA_VERSION 作为版本参数，并使用未版本化的 typedef 来定义函数指针。由于 _v2 是 CUDA 11.3 中驱动程序 API cuStreamBeginCapture 的最新版本，因此下面的代码片段显示了检索它的不同方法。

// Assuming we are using CUDA 11.3 Toolkit
    #include <cudaTypedefs.h>
    // 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);

请注意，请求具有无效 CUDA 版本的驱动程序 API 将返回错误 CUDA_ERROR_NOT_FOUND。 在上面的代码示例中，传入小于 10000 (CUDA 10.0) 的版本将是无效的。

L.5.3.2. Using the runtime API

运行时 API 使用 CUDA 运行时版本来获取请求的驱动程序符号的 ABI 兼容版本。 在下面的代码片段中，所需的最低 CUDA 运行时版本将是 CUDA 11.2，因为当时引入了 cuMemAllocAsync 。

#include <cudaTypedefs.h>
    // 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(...);

L.5.3.3. Retrieve per-thread default stream versions

一些 CUDA 驱动程序 API 可以配置为具有默认流或每线程默认流语义。具有每个线程默认流语义的驱动程序 API 在其名称中以 _ptsz 或 _ptds 为后缀。例如，cuLaunchKernel 有一个名为 cuLaunchKernel_ptsz 的每线程默认流变体。使用驱动程序入口点访问 API，用户可以请求驱动程序 API cuLaunchKernel 的每线程默认流版本，而不是默认流版本。为默认流或每线程默认流语义配置 CUDA 驱动程序 API 会影响同步行为。更多详细信息可以在这里找到。

驱动API的默认流或每线程默认流版本可以通过以下方式之一获得：

• 使用编译标志 --default-stream per-thread 或定义宏 CUDA_API_PER_THREAD_DEFAULT_STREAM 以获取每个线程的默认流行为。
• 分别使用标志 CU_GET_PROC_ADDRESS_LEGACY_STREAM/cudaEnableLegacyStream 或 CU_GET_PROC_ADDRESS_PER_THREAD_DEFAULT_STREAM/cudaEnablePerThreadDefaultStream 强制默认流或每个线程的默认流行为。

L.5.3.4. Access new CUDA features

始终建议安装最新的 CUDA 工具包以访问新的 CUDA 驱动程序功能，但如果出于某种原因，用户不想更新或无法访问最新的工具包，则可以使用 API 来访问新的 CUDA 功能 只有更新的 CUDA 驱动程序。 为了讨论，让我们假设用户使用 CUDA 11.3，并希望使用 CUDA 12.0 驱动程序中提供的新驱动程序 API cuFoo 。 下面的代码片段说明了这个用例：

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