4.21. CUDA 与 API 的互操作性

在 CUDA 中直接访问来自 API 的 GPU 数据，可以使用 CUDA 核函数读写这些数据，从而在其他 API 中使用这些数据时提供 CUDA 功能。主要有两个概念：一种是直接方式，即与 OpenGL 和 Direct3D[9-11] 进行图形互操作，它能够将 OpenGL 和 Direct3D 的资源映射到 CUDA 地址空间；另一种是更灵活的外部资源互操作，它允许通过操作系统级句柄进行导入和导出，以访问内存和同步对象。以下 API 支持这种方式：Direct3D[11-12]、Vulkan 和 NVIDIA 软件通信接口互操作。

4.21.1. 图形互操作性

在使用 CUDA 访问 Direct3D 或 OpenGL 资源（例如 VBO，即顶点缓冲区对象）之前，必须先注册和映射该资源。使用相应的 CUDA 函数进行注册（参见下面的示例），会返回一个类型为struct cudaGraphicsResource的 CUDA 图形资源，它持有 CUDA 设备指针或数组。要在核函数中访问设备数据，必须映射该资源。在资源处于已注册状态时，可以根据需要多次映射和取消映射它。核函数使用返回的设备内存地址来访问已映射的资源：对于缓冲区，该地址由cudaGraphicsResourceGetMappedPointer()返回；对于 CUDA 数组，则由cudaGraphicsSubResourceGetMappedArray()返回。一旦 CUDA 不再需要该资源，就可以注销它。

主要步骤如下：

向 CUDA 注册图形缓冲区
映射资源
访问已映射资源的设备指针或数组
在 CUDA 核函数中使用设备指针或数组
取消资源映射
注销资源

注意，注册资源的开销很大，因此理想情况下，每个资源只调用一次注册操作，但是，对于每个打算使用该资源的 CUDA 上下文，都必须单独注册该资源。可以调用cudaGraphicsResourceSetMapFlags()来指定使用提示（只写、只读），CUDA 驱动程序可以利用这些提示优化资源管理。还需注意，在资源处于已映射状态时，通过 OpenGL、Direct3D 或其他 CUDA 上下文访问该资源会产生未定义的结果。

4.21.1.1. OpenGL 互操作性

可以映射到 CUDA 地址空间的 OpenGL 资源包括 OpenGL 缓冲区、纹理和渲染缓冲区对象。缓冲区对象使用cudaGraphicsGLRegisterBuffer()注册，在 CUDA 中表现为普通的设备指针。纹理或渲染缓冲区对象使用cudaGraphicsGLRegisterImage()注册，在 CUDA 中表现为 CUDA 数组。

如果纹理或渲染缓冲区对象已使用cudaGraphicsRegisterFlagsSurfaceLoadStore标志注册，则可以向其写入。cudaGraphicsGLRegisterImage()支持所有具有 1、2 或 4 个分量且内部类型为 float（例如GL_RGBA_FLOAT32）、归一化整数（例如GL_RGBA8, GL_INTENSITY16）和非归一化整数（例如GL_RGBA8UI）的纹理格式。

示例：simpleGL 互操作

以下代码示例使用核函数动态修改存储在顶点缓冲区对象（VBO）中的二维widthxheight顶点网格，并依次执行以下主要步骤：

向 CUDA 注册 VBO
循环：映射 VBO，以便从 CUDA 写入
循环：运行 CUDA 核函数以修改顶点位置
循环：取消 VBO 映射
循环：使用 OpenGL 渲染结果
注销并删除 VBO

本节的完整示例 simpleGL 可以在此处找到：NVIDIA/cuda-samples。

__global__ void simple_vbo_kernel(float4 *pos, unsigned int width, unsigned int height, float time)
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;

    // calculate uv coordinates
    float u = x / (float)width;
    float v = y / (float)height;
    u = u * 2.0f - 1.0f;
    v = v * 2.0f - 1.0f;

    // calculate simple sine wave pattern
    float freq = 4.0f;
    float w = sinf(u * freq + time) * cosf(v * freq + time) * 0.5f;

    // write output vertex
    pos[y * width + x] = make_float4(u, w, v, 1.0f);
}

int main(int argc, char **argv)
{
    char *ref_file = NULL;

    pArgc = &argc;
    pArgv = argv;

#if defined(__linux__)
    setenv("DISPLAY", ":0", 0);
#endif

    printf("%s starting...\n", sSDKsample);

    if (argc > 1) {
        if (checkCmdLineFlag(argc, (const char **)argv, "file")) {
            // In this mode, we are running non-OpenGL and doing a compare of the VBO was generated correctly
            getCmdLineArgumentString(argc, (const char **)argv, "file", (char **)&ref_file);
        }
    }

    printf("\n");

    // First initialize OpenGL context
    if (false == initGL(&argc, argv)) {
        return false;
    }

    // register callbacks
    glutDisplayFunc(display);
    glutKeyboardFunc(keyboard);
    glutMouseFunc(mouse);
    glutMotionFunc(motion);
    glutCloseFunc(cleanup);

    // Create an empty vertex buffer object (VBO)
    // 1. Register the VBO with CUDA
    createVBO(&vbo, &cuda_vbo_resource, cudaGraphicsMapFlagsWriteDiscard);

    // start rendering mainloop
    //  5. Render the results using OpenGL
    glutMainLoop();

    printf("%s completed, returned %s\n", sSDKsample, (g_TotalErrors == 0) ? "OK" : "ERROR!");
    exit(g_TotalErrors == 0 ? EXIT_SUCCESS : EXIT_FAILURE);
    
}

void createVBO(GLuint *vbo, struct cudaGraphicsResource **vbo_res, unsigned int vbo_res_flags)
{
    assert(vbo);

    // create buffer object
    glGenBuffers(1, vbo);
    glBindBuffer(GL_ARRAY_BUFFER, *vbo);

    // initialize buffer object
    unsigned int size = mesh_width * mesh_height * 4 * sizeof(float);
    glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);

    glBindBuffer(GL_ARRAY_BUFFER, 0);

    // register this buffer object with CUDA
    checkCudaErrors(cudaGraphicsGLRegisterBuffer(vbo_res, *vbo, vbo_res_flags));

    SDK_CHECK_ERROR_GL();
}

void display()
{
    float4 *dptr;
    // 2. Map the VBO for writing from CUDA
    checkCudaErrors(cudaGraphicsMapResources(1, &cuda_vbo_resource, 0));
    size_t num_bytes;
    checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr, &num_bytes, cuda_vbo_resource));

    // 3. Run CUDA kernel to modify the vertex positions
    //call the CUDA kernel
    dim3 block(8, 8, 1);
    dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
    simple_vbo_kernel<<<grid, block>>>(dptr, mesh_width, mesh_height, g_fAnim);

    //  4. Unmap the VBO    
    checkCudaErrors(cudaGraphicsUnmapResources(1, &cuda_vbo_resource, 0));

    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);

    // set view matrix
    glMatrixMode(GL_MODELVIEW);
    glLoadIdentity();
    glTranslatef(0.0, 0.0, translate_z);
    glRotatef(rotate_x, 1.0, 0.0, 0.0);
    glRotatef(rotate_y, 0.0, 1.0, 0.0);

    // 5. Render the updated  using OpenGL
    glBindBuffer(GL_ARRAY_BUFFER, vbo);
    glVertexPointer(4, GL_FLOAT, 0, 0);

    glEnableClientState(GL_VERTEX_ARRAY);
    glColor3f(1.0, 0.0, 0.0);
    glDrawArrays(GL_POINTS, 0, mesh_width * mesh_height);
    glDisableClientState(GL_VERTEX_ARRAY);

    glutSwapBuffers();

    g_fAnim += 0.01f;

}

void deleteVBO(GLuint *vbo, struct cudaGraphicsResource *vbo_res)
{
    // 6. Unregister and delete VBO
    checkCudaErrors(cudaGraphicsUnregisterResource(vbo_res));

    glBindBuffer(1, *vbo);
    glDeleteBuffers(1, vbo);

    *vbo = 0;
}

void cleanup()
{

    if (vbo) {
        deleteVBO(&vbo, cuda_vbo_resource);
    }
}

限制与注意事项。

对于进行任何 OpenGL 互操作 API 调用的主机线程，其当前 OpenGL 上下文必须是资源正在被共享的那个上下文。
当 OpenGL 纹理变为无绑定纹理时（例如，使用glGetTextureHandle或glGetImageHandleAPI 请求图像或纹理句柄），便无法向 CUDA 注册该纹理。应用程序需要在请求图像或纹理句柄之前注册该纹理以用于互操作。

4.21.1.2. Direct3D 互操作性

Direct3D 互操作支持 Direct3D9、Direct3D10 和 Direct3D11，但不支持 Direct3D12。这里重点介绍 Direct3D11，有关 Direct3D9 和 Direct3D10，请参阅 CUDA 编程指南 12.9。可以映射到 CUDA 地址空间的 Direct3D 资源包括 Direct3D 缓冲区、纹理和表面。这些资源使用cudaGraphicsD3D11RegisterResource()注册。

CUDA 上下文只能与创建时将DriverType设置为D3D_DRIVER_TYPE_HARDWARE的 Direct3D11 设备互操作。

示例：二维纹理 Direct3D11 互操作

以下代码片段来自NVIDIA/cuda-samples中的 simpleD3D11Texture 示例。完整示例包含大量 DX11 样板代码，这里重点介绍 CUDA 端。

CUDA 核函数cuda_kernel_texture_2d在闪烁的蓝色背景上，以移动的红色/绿色阴影线图案绘制二维纹理，它依赖于先前的纹理值。底层数据是二维 CUDA 数组，其中行偏移量由行距定义。

/*
 * Paint a 2D texture with a moving red/green hatch pattern on a
 * strobing blue background.  Note that this kernel reads to and
 * writes from the texture, hence why this texture was not mapped
 * as WriteDiscard.
 */
__global__ void cuda_kernel_texture_2d(unsigned char *surface, int width,
                                       int height, size_t pitch, float t) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;
  float *pixel;

  // in the case where, due to quantization into grids, we have
  // more threads than pixels, skip the threads which don't
  // correspond to valid pixels
  if (x >= width || y >= height) return;

  // get a pointer to the pixel at (x,y)
  pixel = (float *)(surface + y * pitch) + 4 * x;

  // populate it
  float value_x = 0.5f + 0.5f * cos(t + 10.0f * ((2.0f * x) / width - 1.0f));
  float value_y = 0.5f + 0.5f * cos(t + 10.0f * ((2.0f * y) / height - 1.0f));
  pixel[0] = 0.5 * pixel[0] + 0.5 * pow(value_x, 3.0f);  // red
  pixel[1] = 0.5 * pixel[1] + 0.5 * pow(value_y, 3.0f);  // green
  pixel[2] = 0.5f + 0.5f * cos(t);                       // blue
  pixel[3] = 1;                                          // alpha
}

extern "C" void cuda_texture_2d(void *surface, int width, int height,
                                size_t pitch, float t) {
  cudaError_t error = cudaSuccess;

  dim3 Db = dim3(16, 16);  // block dimensions are fixed to be 256 threads
  dim3 Dg = dim3((width + Db.x - 1) / Db.x, (height + Db.y - 1) / Db.y);

  cuda_kernel_texture_2d<<<Dg, Db>>>((unsigned char *)surface, width, height,
                                     pitch, t);

  error = cudaGetLastError();

  if (error != cudaSuccess) {
    printf("cuda_kernel_texture_2d() failed to launch error = %d\n", error);
  }
}

为使指针与数据缓冲区保持对应，使用以下数据结构：

// Data structure for 2D texture shared between DX11 and CUDA
struct {
  ID3D11Texture2D *pTexture;
  ID3D11ShaderResourceView *pSRView;
  cudaGraphicsResource *cudaResource;
  void *cudaLinearMemory;
  size_t pitch;
  int width;
  int height;
  int offsetInShader;
} g_texture_2d;

初始化 Direct3D 设备和纹理后，向 CUDA 注册这些资源一次。为匹配 Direct3D 像素格式，分配 CUDA 数组时采用相同的宽度和高度，以及与 Direct3D 纹理行距匹配的行距。

    // register the Direct3D resources that are used in the CUDA kernel
    // we'll read to and write from g_texture_2d, so don't set any special map flags for it
    cudaGraphicsD3D11RegisterResource(&g_texture_2d.cudaResource,
                                      g_texture_2d.pTexture,
                                      cudaGraphicsRegisterFlagsNone);
    getLastCudaError("cudaGraphicsD3D11RegisterResource (g_texture_2d) failed");
    // CUDA cannot write into the texture directly : the texture is seen as a
    // cudaArray and can only be mapped as a texture
    // Create a buffer so that CUDA can write into it
    // the pixel fmt is DXGI_FORMAT_R32G32B32A32_FLOAT
    cudaMallocPitch(&g_texture_2d.cudaLinearMemory, &g_texture_2d.pitch,
                    g_texture_2d.width * sizeof(float) * 4,
                    g_texture_2d.height);
    getLastCudaError("cudaMallocPitch (g_texture_2d) failed");
    cudaMemset(g_texture_2d.cudaLinearMemory, 1,
               g_texture_2d.pitch * g_texture_2d.height);

在渲染循环中，先映射资源，再启动 CUDA 核函数以更新纹理数据，然后取消资源映射。完成此步骤后，使用 Direct3D 设备在屏幕上绘制更新后的纹理。

    cudaStream_t stream = 0;
    const int nbResources = 3;
    cudaGraphicsResource *ppResources[nbResources] = {
        g_texture_2d.cudaResource, g_texture_3d.cudaResource,
        g_texture_cube.cudaResource,
    };
    cudaGraphicsMapResources(nbResources, ppResources, stream);
    getLastCudaError("cudaGraphicsMapResources(3) failed");

    // run kernels which will populate the contents of those textures
    RunKernels();

    // unmap the resources
    cudaGraphicsUnmapResources(nbResources, ppResources, stream);
    getLastCudaError("cudaGraphicsUnmapResources(3) failed");

最后，当 CUDA 不再需要这些资源时，将注销它们并释放设备数组。

  // unregister the Cuda resources
  cudaGraphicsUnregisterResource(g_texture_2d.cudaResource);
  getLastCudaError("cudaGraphicsUnregisterResource (g_texture_2d) failed");
  cudaFree(g_texture_2d.cudaLinearMemory);
  getLastCudaError("cudaFree (g_texture_2d) failed");

4.21.1.3. 可扩展链接接口（SLI）配置中的互操作性

在具有多个 GPU 的系统中，所有支持 CUDA 的 GPU 都可以通过 CUDA 驱动程序和运行时作为独立设备进行访问。当系统处于 SLI 模式时，情况有所不同。SLI 是一种通过硬件配置的多 GPU 配置，它通过将工作负载分配到多个 GPU 上来提高渲染性能。由驱动程序作出假设的隐式 SLI 模式已不再受支持，但显式 SLI 仍受支持。显式 SLI 意味着应用程序通过 API（例如 Vulkan、DirectX、GL）了解并管理 SLI 组中所有设备的 SLI 状态。

当系统处于 SLI 模式时，需要特别考虑以下事项：

在一个 GPU 上的一个 CUDA 设备中进行的内存分配，会消耗属于该 Direct3D 或 OpenGL 设备的 SLI 配置中的其他 GPU 上的内存。因此，内存分配可能会比通常预期的更早失败。
应用程序应创建多个 CUDA 上下文，为 SLI 配置中的每个 GPU 各创建一个。虽然这不是严格要求，但这样可以避免设备之间不必要的数据传输。应用程序可以使用用于 Direct3D 的cudaD3D[9|10|11]GetDevices()和用于 OpenGL 的cudaGLGetDevices()这两组调用，来确定在当前帧和下一帧中执行渲染的设备所对应的 CUDA 设备句柄。有了这些信息，应用程序通常会选择适当的设备，并将 Direct3D 或 OpenGL 资源映射到当deviceList参数设置为cudaD3D[9|10|11]DeviceListCurrentFrame或cudaGLDeviceListCurrentFrame时，由cudaD3D[9|10|11]GetDevices()或cudaGLGetDevices()返回的 CUDA 设备。
从cudaGraphicsD3D[9|10|11]RegisterResource和cudaGraphicsGLRegister[Buffer|Image]返回的资源只能在进行注册的设备上使用。因此，在 SLI 配置中，当不同帧的数据在不同的 CUDA 设备上计算时，必须为每个设备分别注册资源。

4.21.2. 外部资源互操作性

外部资源互操作性允许 CUDA 导入由 API 显式导出的某些资源。这些对象通常使用操作系统的原生句柄导出，例如 Linux 上的文件描述符或 Windows 上的 NT 句柄。这样可以在其他 API 与 CUDA 之间高效地共享资源，而无须在两者之间进行复制或创建副本。以下 API 支持此功能：Direct3D[11-12]、Vulkan 和 NVIDIA 软件通信接口互操作性。可以导入两种类型的资源：

内存对象
可以使用cudaImportExternalMemory()导入 CUDA。然后，可以在核函数内使用通过cudaExternalMemoryGetMappedBuffer()映射到该内存对象的设备指针，或通过cudaExternalMemoryGetMappedMipmappedArray()映射的 CUDA mipmap 数组来访问已导入的内存对象。根据内存对象的类型，可能可以在单个内存对象上建立多个映射。这些映射必须与导出方 API 建立的映射相匹配。任何不匹配的映射都会导致未定义行为。已导入的内存对象必须使用cudaDestroyExternalMemory()释放。释放内存对象不会释放到该对象的任何映射。因此，映射到该对象的任何设备指针都必须使用cudaFree()显式释放，映射到该对象的任何 CUDA mipmap 数组都必须使用cudaFreeMipmappedArray()显式释放。在对象被销毁后访问到该对象的映射是非法的。
同步对象
可以使用cudaImportExternalSemaphore()导入 CUDA。然后，可以使用cudaSignalExternalSemaphoresAsync()向已导入的同步对象发出信号，并使用cudaWaitExternalSemaphoresAsync()等待该对象。在相应的信号发出之前发起等待是非法的。此外，根据所导入的同步对象的类型，可能还会对向其发出信号和等待它的方式施加额外的约束，如后续各节所述。已导入的信号量对象必须使用cudaDestroyExternalSemaphore()释放。在销毁信号量对象之前，所有尚未完成的信号和等待操作都必须已完成。

4.21.2.1. Vulkan 互操作性

在同一硬件上耦合执行 Vulkan 图形和计算工作负载，可以最大限度地提高 GPU 利用率并避免不必要的复制。注意，这不是 Vulkan 指南，我们仅关注与 CUDA 的互操作性，有关 Vulkan 指南，请参阅https://www.vulkan.org/learn#vulkan-tutorials。

实现 Vulkan-CUDA 互操作性所需的主要步骤包括：

初始化 Vulkan，创建并导出外部缓冲区和/或同步对象
通过匹配设备 UUID，将 CUDA 设备设置为 Vulkan 正在运行的设备
获取内存和/或同步句柄
使用这些句柄在 CUDA 中导入内存和/或同步对象
将设备指针或 mipmap 数组映射到内存对象
通过向同步对象发出信号和等待同步对象来定义执行顺序，从而在 CUDA 和 Vulkan 中交替使用已导入的内存对象。

本节借助NVIDIA/cuda-samples中的simpleVulkan示例说明上述步骤。我们逐步讲解该示例，重点介绍实现 CUDA 互操作性所需的部分。对于一些变体，则使用独立的代码片段加以说明。

注意

本节使用的代码示例采用直接内存分配和资源创建方式。由于多种原因，包括可创建的实例数量受到限制，这并不是当前最先进的方式。但是，要理解互操作性，就需要了解底层 Vulkan 代码以及特定的标志。有关使用VulkanMemoryAllocator的更先进的示例，请参阅NVProSamples仓库中的sample_cuda_interop。

整个示例中都使用以下数据结构：

class VulkanCudaSineWave : public VulkanBaseApp {
  typedef struct UniformBufferObject_st {
    mat4x4 modelViewProj;
  } UniformBufferObject;

  VkBuffer m_heightBuffer, m_xyBuffer, m_indexBuffer;
  VkDeviceMemory m_heightMemory, m_xyMemory, m_indexMemory;
  UniformBufferObject m_ubo;
  VkSemaphore m_vkWaitSemaphore, m_vkSignalSemaphore;
  SineWaveSimulation m_sim;
  cudaStream_t m_stream;
  cudaExternalSemaphore_t m_cudaWaitSemaphore, m_cudaSignalSemaphore, m_cudaTimelineSemaphore;
  cudaExternalMemory_t m_cudaVertMem;
  float *m_cudaHeightMap;
  // ...

4.21.2.1.1. 设置 Vulkan 设备

为了导出内存对象，创建 Vulkan 实例时必须启用VK_KHR_external_memory_capabilities扩展，创建设备时必须启用VK_KHR_external_memory。此外，还必须启用平台特定的句柄类型，在 Windows 上为VK_KHR_external_memory_win32，在基于 UNIX 的系统上为VK_KHR_external_memory_fd。

类似地，为了导出同步对象，需要在设备级别启用VK_KHR_external_semaphore_capabilities，在实例级别启用VK_KHR_external_semaphore。还需要启用平台特定的句柄扩展，即 Windows 上的VK_KHR_external_semaphore_win32和基于 Unix 的系统上的VK_KHR_external_semaphore_fd。

在simpleVulkan示例中，使用以下枚举启用这些扩展。

  std::vector<const char *> getRequiredExtensions() const {
    std::vector<const char *> extensions;
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME);
    return extensions;
  }

  std::vector<const char *> getRequiredDeviceExtensions() const {
    std::vector<const char *> extensions;
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME);
    extensions.push_back(VK_KHR_TIMELINE_SEMAPHORE_EXTENSION_NAME);
#ifdef _WIN64
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME);
#else
    extensions.push_back(VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME);
    extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_FD_EXTENSION_NAME);
#endif /* _WIN64 */
    return extensions;
  }

随后，将这些扩展添加到 Vulkan 实例和设备创建信息中，详情请参阅simpleVulkan示例。

4.21.2.1.2. 使用匹配的设备 UUID 初始化 CUDA

导入由 Vulkan 导出的内存和同步对象时，必须在创建这些对象的同一设备上导入并映射它们。通过比较 CUDA 设备的 UUID 与 Vulkan 物理设备的 UUID，可以确定与创建这些对象的 Vulkan 物理设备相对应的 CUDA 设备，如以下来自 simpleVulkan 示例的代码片段所示。其中，vkDeviceUUID是 Vulkan API 结构体vkPhysicalDeviceIDProperties.deviceUUID的成员，定义了当前 Vulkan 实例的物理设备 ID。

// from the CUDA example `simpleVulkan`
int SineWaveSimulation::initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE) {
  int current_device = 0;
  int device_count = 0;
  int devices_prohibited = 0;

  cudaDeviceProp deviceProp;
  checkCudaErrors(cudaGetDeviceCount(&device_count));

  if (device_count == 0) {
    fprintf(stderr, "CUDA error: no devices supporting CUDA.\n");
    exit(EXIT_FAILURE);
  }

  // Find the GPU which is selected by Vulkan
  while (current_device < device_count) {
    cudaGetDeviceProperties(&deviceProp, current_device);

    if ((deviceProp.computeMode != cudaComputeModeProhibited)) {
      // Compare the cuda device UUID with vulkan UUID
      int ret = memcmp((void *)&deviceProp.uuid, vkDeviceUUID, UUID_SIZE);
      if (ret == 0) {
        checkCudaErrors(cudaSetDevice(current_device));
        checkCudaErrors(cudaGetDeviceProperties(&deviceProp, current_device));
        printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n",
               current_device, deviceProp.name, deviceProp.major,
               deviceProp.minor);

        return current_device;
      }

    } else {
      devices_prohibited++;
    }

    current_device++;
  }

  if (devices_prohibited == device_count) {
    fprintf(stderr,
            "CUDA error:"
            " No Vulkan-CUDA Interop capable GPU found.\n");
    exit(EXIT_FAILURE);
  }

  return -1;
}

注意，该 Vulkan 物理设备不应属于包含多个 Vulkan 物理设备的设备组。也就是说，由vkEnumeratePhysicalDeviceGroups返回的、包含给定 Vulkan 物理设备的设备组，其物理设备数量必须为 1。

4.21.2.1.3. 导出 Vulkan 内存对象

为了导出 Vulkan 内存对象，必须创建带有相应导出标志的缓冲区。注意，句柄类型的枚举是平台特定的。

void VulkanBaseApp::createExternalBuffer(
    VkDeviceSize size, VkBufferUsageFlags usage,
    VkMemoryPropertyFlags properties,
    VkExternalMemoryHandleTypeFlagsKHR extMemHandleType, VkBuffer &buffer,
    VkDeviceMemory &bufferMemory) {
  VkBufferCreateInfo bufferInfo = {};
  bufferInfo.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
  bufferInfo.size = size;
  bufferInfo.usage = usage;
  bufferInfo.sharingMode = VK_SHARING_MODE_EXCLUSIVE;

  VkExternalMemoryBufferCreateInfo externalMemoryBufferInfo = {};
  externalMemoryBufferInfo.sType =
      VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_BUFFER_CREATE_INFO;
  externalMemoryBufferInfo.handleTypes = extMemHandleType;
  bufferInfo.pNext = &externalMemoryBufferInfo;

  if (vkCreateBuffer(m_device, &bufferInfo, nullptr, &buffer) != VK_SUCCESS) {
    throw std::runtime_error("failed to create buffer!");
  }

  VkMemoryRequirements memRequirements;
  vkGetBufferMemoryRequirements(m_device, buffer, &memRequirements);

#ifdef _WIN64
  WindowsSecurityAttributes winSecurityAttributes;

  VkExportMemoryWin32HandleInfoKHR vulkanExportMemoryWin32HandleInfoKHR = {};
  vulkanExportMemoryWin32HandleInfoKHR.sType =
      VK_STRUCTURE_TYPE_EXPORT_MEMORY_WIN32_HANDLE_INFO_KHR;
  vulkanExportMemoryWin32HandleInfoKHR.pNext = NULL;
  vulkanExportMemoryWin32HandleInfoKHR.pAttributes = &winSecurityAttributes;
  vulkanExportMemoryWin32HandleInfoKHR.dwAccess =
      DXGI_SHARED_RESOURCE_READ | DXGI_SHARED_RESOURCE_WRITE;
  vulkanExportMemoryWin32HandleInfoKHR.name = (LPCWSTR)NULL;
#endif /* _WIN64 */
  VkExportMemoryAllocateInfoKHR vulkanExportMemoryAllocateInfoKHR = {};
  vulkanExportMemoryAllocateInfoKHR.sType =
      VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO_KHR;
#ifdef _WIN64
  vulkanExportMemoryAllocateInfoKHR.pNext =
      extMemHandleType & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR
          ? &vulkanExportMemoryWin32HandleInfoKHR
          : NULL;
  vulkanExportMemoryAllocateInfoKHR.handleTypes = extMemHandleType;
#else
  vulkanExportMemoryAllocateInfoKHR.pNext = NULL;
  vulkanExportMemoryAllocateInfoKHR.handleTypes =
      VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT;
#endif /* _WIN64 */
  VkMemoryAllocateInfo allocInfo = {};
  allocInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
  allocInfo.pNext = &vulkanExportMemoryAllocateInfoKHR;
  allocInfo.allocationSize = memRequirements.size;
  allocInfo.memoryTypeIndex = findMemoryType(
      m_physicalDevice, memRequirements.memoryTypeBits, properties);

  if (vkAllocateMemory(m_device, &allocInfo, nullptr, &bufferMemory) !=
      VK_SUCCESS) {
    throw std::runtime_error("failed to allocate external buffer memory!");
  }

  vkBindBufferMemory(m_device, buffer, bufferMemory, 0);
}

4.21.2.1.4. 导出 Vulkan 同步对象

在 GPU 上执行的 Vulkan API 调用是异步的。为了定义执行顺序，Vulkan 提供了可与 CUDA 共享的信号量和栅栏。与内存对象类似，信号量可以由 Vulkan 导出，创建信号量时需要根据其类型指定导出标志。信号量分为二进制信号量和时间线信号量。二进制信号量只有一个 1 位计数器，表示已发出信号或未发出信号。时间线信号量有一个 64 位计数器，可以使用同一个信号量来定义执行顺序。simpleVulkan示例中包含时间线信号量和二进制信号量两者的代码路径。

void VulkanBaseApp::createExternalSemaphore(
    VkSemaphore &semaphore, VkExternalSemaphoreHandleTypeFlagBits handleType) {
  VkSemaphoreCreateInfo semaphoreInfo = {};
  semaphoreInfo.sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO;
  VkExportSemaphoreCreateInfoKHR exportSemaphoreCreateInfo = {};
  exportSemaphoreCreateInfo.sType =
      VK_STRUCTURE_TYPE_EXPORT_SEMAPHORE_CREATE_INFO_KHR;

#ifdef _VK_TIMELINE_SEMAPHORE
  VkSemaphoreTypeCreateInfo timelineCreateInfo;
  timelineCreateInfo.sType = VK_STRUCTURE_TYPE_SEMAPHORE_TYPE_CREATE_INFO;
  timelineCreateInfo.pNext = NULL;
  timelineCreateInfo.semaphoreType = VK_SEMAPHORE_TYPE_TIMELINE;
  timelineCreateInfo.initialValue = 0;
  exportSemaphoreCreateInfo.pNext = &timelineCreateInfo;
#else
  exportSemaphoreCreateInfo.pNext = NULL;
#endif /* _VK_TIMELINE_SEMAPHORE */
  exportSemaphoreCreateInfo.handleTypes = handleType;
  semaphoreInfo.pNext = &exportSemaphoreCreateInfo;

  if (vkCreateSemaphore(m_device, &semaphoreInfo, nullptr, &semaphore) !=
      VK_SUCCESS) {
    throw std::runtime_error(
        "failed to create synchronization objects for a CUDA-Vulkan!");
  }
}

4.21.2.1.5. 导入内存对象

由 Vulkan 导出的专用和非专用内存对象都可以导入 CUDA。导入 Vulkan 专用内存对象时，必须设置标志cudaExternalMemoryDedicated。

在 Windows 中，可以使用与对象关联的 NT 句柄，将使用VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT导出的 Vulkan 内存对象导入 CUDA，如下所示。注意，CUDA 不会接管 NT 句柄的所有权，应用程序有责任在不再需要该句柄时将其关闭。NT 句柄持有对该资源的引用，因此必须先显式释放该句柄，才能释放底层内存。

在 Linux 中，可以使用与对象关联的文件描述符，将使用VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT导出的 Vulkan 内存对象导入 CUDA，如下所示。注意，一旦导入，CUDA 就会接管该文件描述符的所有权。在成功导入后使用该文件描述符会导致未定义行为。

  // from the CUDA example `simpleVulkan`
  void importCudaExternalMemory(void **cudaPtr, cudaExternalMemory_t &cudaMem,
                                VkDeviceMemory &vkMem, VkDeviceSize size,
                                VkExternalMemoryHandleTypeFlagBits handleType) {
    cudaExternalMemoryHandleDesc externalMemoryHandleDesc = {};

    if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) {
      externalMemoryHandleDesc.type = cudaExternalMemoryHandleTypeOpaqueWin32;
    } else if (handleType &
               VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) {
      externalMemoryHandleDesc.type =
          cudaExternalMemoryHandleTypeOpaqueWin32Kmt;
    } else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) {
      externalMemoryHandleDesc.type = cudaExternalMemoryHandleTypeOpaqueFd;
    } else {
      throw std::runtime_error("Unknown handle type requested!");
    }

    externalMemoryHandleDesc.size = size;

#ifdef _WIN64
    externalMemoryHandleDesc.handle.win32.handle =
        (HANDLE)getMemHandle(vkMem, handleType);
#else
    externalMemoryHandleDesc.handle.fd =
        (int)(uintptr_t)getMemHandle(vkMem, handleType);
#endif

    checkCudaErrors(
        cudaImportExternalMemory(&cudaMem, &externalMemoryHandleDesc));

对于使用VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT导出的 Vulkan 内存对象，如果存在命名句柄，也可以使用该命名句柄进行导入，如下面的独立代码片段所示。

cudaExternalMemory_t importVulkanMemoryObjectFromNamedNTHandle(LPCWSTR name, unsigned long long size, bool isDedicated) {
   cudaExternalMemory_t extMem = NULL;
   cudaExternalMemoryHandleDesc desc = {};

   memset(&desc, 0, sizeof(desc));

   desc.type = cudaExternalMemoryHandleTypeOpaqueWin32;
   desc.handle.win32.name = (void *)name;
   desc.size = size;
   if (isDedicated) {
       desc.flags |= cudaExternalMemoryDedicated;
   }

   cudaImportExternalMemory(&extMem, &desc);

   return extMem;
}

4.21.2.1.6. 将缓冲区映射到已导入的内存对象

导入内存对象后，必须先对其进行映射才能使用。可以将设备指针映射到已导入的内存对象，如下所示。映射的偏移量和大小必须与使用相应 Vulkan API 创建映射时指定的值相匹配。所有已映射的设备指针都必须使用cudaFree()释放。

    // from the CUDA example `simpleVulkan`, continuation of function `importCudaExternalMemory`
    cudaExternalMemoryBufferDesc externalMemBufferDesc = {};
    externalMemBufferDesc.offset = 0;
    externalMemBufferDesc.size = size;
    externalMemBufferDesc.flags = 0;

    checkCudaErrors(cudaExternalMemoryGetMappedBuffer(cudaPtr, cudaMem,
                                                      &externalMemBufferDesc));
  }

4.21.2.1.7. 将 mipmap 数组映射到已导入的内存对象

可以将 CUDA mipmap 数组映射到已导入的内存对象，如下所示。偏移量、各维尺寸、格式和 mip 级别数必须与使用相应 Vulkan API 创建映射时指定的值相匹配。此外，如果该 mipmap 数组在 Vulkan 中被绑定为颜色目标，则必须设置标志cudaArrayColorAttachment。所有已映射的 mipmap 数组都必须使用cudaFreeMipmappedArray()释放。以下独立代码片段展示了在将 mipmap 数组映射到已导入的内存对象时，如何将 Vulkan 参数转换为相应的 CUDA 参数。

cudaMipmappedArray_t mapMipmappedArrayOntoExternalMemory(cudaExternalMemory_t extMem, unsigned long long offset, cudaChannelFormatDesc *formatDesc, cudaExtent *extent, unsigned int flags, unsigned int numLevels) {
    cudaMipmappedArray_t mipmap = NULL;
    cudaExternalMemoryMipmappedArrayDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.offset = offset;
    desc.formatDesc = *formatDesc;
    desc.extent = *extent;
    desc.flags = flags;
    desc.numLevels = numLevels;

    // Note: 'mipmap' must eventually be freed using cudaFreeMipmappedArray()
    cudaExternalMemoryGetMappedMipmappedArray(&mipmap, extMem, &desc);

    return mipmap;
}
//end mapMipmappedArrayOntoExternalMemory

//begin getCudaChannelFormatDescForVulkanFormat
cudaChannelFormatDesc getCudaChannelFormatDescForVulkanFormat(VkFormat format)
{
    cudaChannelFormatDesc d;

    memset(&d, 0, sizeof(d));
 
    switch (format) {
       case VK_FORMAT_R8_UINT:             d.x = 8;  d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R8_SINT:             d.x = 8;  d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R8G8_UINT:           d.x = 8;  d.y = 8;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R8G8_SINT:           d.x = 8;  d.y = 8;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R8G8B8A8_UINT:       d.x = 8;  d.y = 8;  d.z = 8;  d.w = 8;  d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R8G8B8A8_SINT:       d.x = 8;  d.y = 8;  d.z = 8;  d.w = 8;  d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R16_UINT:            d.x = 16; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R16_SINT:            d.x = 16; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R16G16_UINT:         d.x = 16; d.y = 16; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R16G16_SINT:         d.x = 16; d.y = 16; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R16G16B16A16_UINT:   d.x = 16; d.y = 16; d.z = 16; d.w = 16; d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R16G16B16A16_SINT:   d.x = 16; d.y = 16; d.z = 16; d.w = 16; d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R32_UINT:            d.x = 32; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R32_SINT:            d.x = 32; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R32_SFLOAT:          d.x = 32; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindFloat;    break;
       case VK_FORMAT_R32G32_UINT:         d.x = 32; d.y = 32; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R32G32_SINT:         d.x = 32; d.y = 32; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R32G32_SFLOAT:       d.x = 32; d.y = 32; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindFloat;    break;
       case VK_FORMAT_R32G32B32A32_UINT:   d.x = 32; d.y = 32; d.z = 32; d.w = 32; d.f = cudaChannelFormatKindUnsigned; break;
       case VK_FORMAT_R32G32B32A32_SINT:   d.x = 32; d.y = 32; d.z = 32; d.w = 32; d.f = cudaChannelFormatKindSigned;   break;
       case VK_FORMAT_R32G32B32A32_SFLOAT: d.x = 32; d.y = 32; d.z = 32; d.w = 32; d.f = cudaChannelFormatKindFloat;    break;
       default: assert(0);
    }
    return d;
}
//end getCudaChannelFormatDescForVulkanFormat

//begin getCudaExtentForVulkanExtent
cudaExtent getCudaExtentForVulkanExtent(VkExtent3D vkExt, uint32_t arrayLayers, VkImageViewType vkImageViewType) {
    cudaExtent e = { 0, 0, 0 };

    switch (vkImageViewType) {
        case VK_IMAGE_VIEW_TYPE_1D:         e.width = vkExt.width; e.height = 0;            e.depth = 0;           break;
        case VK_IMAGE_VIEW_TYPE_2D:         e.width = vkExt.width; e.height = vkExt.height; e.depth = 0;           break;
        case VK_IMAGE_VIEW_TYPE_3D:         e.width = vkExt.width; e.height = vkExt.height; e.depth = vkExt.depth; break;
        case VK_IMAGE_VIEW_TYPE_CUBE:       e.width = vkExt.width; e.height = vkExt.height; e.depth = arrayLayers; break;
        case VK_IMAGE_VIEW_TYPE_1D_ARRAY:   e.width = vkExt.width; e.height = 0;            e.depth = arrayLayers; break;
        case VK_IMAGE_VIEW_TYPE_2D_ARRAY:   e.width = vkExt.width; e.height = vkExt.height; e.depth = arrayLayers; break;
        case VK_IMAGE_VIEW_TYPE_CUBE_ARRAY: e.width = vkExt.width; e.height = vkExt.height; e.depth = arrayLayers; break;
        default: assert(0);
    }

    return e;
}
//end getCudaExtentForVulkanExtent

//begin getCudaMipmappedArrayFlagsForVulkanImage
unsigned int getCudaMipmappedArrayFlagsForVulkanImage(VkImageViewType vkImageViewType,
                                                      VkImageUsageFlags vkImageUsageFlags,
                                                      bool allowSurfaceLoadStore) {
    unsigned int flags = 0;

    switch (vkImageViewType) {
        case VK_IMAGE_VIEW_TYPE_CUBE:       flags |= cudaArrayCubemap;                    break;
        case VK_IMAGE_VIEW_TYPE_CUBE_ARRAY: flags |= cudaArrayCubemap | cudaArrayLayered; break;
        case VK_IMAGE_VIEW_TYPE_1D_ARRAY:   flags |= cudaArrayLayered;                    break;
        case VK_IMAGE_VIEW_TYPE_2D_ARRAY:   flags |= cudaArrayLayered;                    break;
        default: break;
    }
    if (vkImageUsageFlags & VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT) {
        flags |= cudaArrayColorAttachment;
    }

    if (allowSurfaceLoadStore) {
        flags |= cudaArraySurfaceLoadStore;
    }
    
    return flags;
}

4.21.2.1.8. 导入同步对象

可以使用与对象关联的文件描述符，将使用VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT导出的 Vulkan 信号量对象导入 CUDA，如下所示。注意，一旦导入，CUDA 就会接管该文件描述符的所有权。在成功导入后使用该文件描述符会导致未定义行为。

而对于使用VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT导出的 Vulkan 信号量对象，可以使用与该对象关联的 NT 句柄将其导入 CUDA，如下所示。注意，CUDA 不会接管 NT 句柄的所有权，应用程序有责任在不再需要该句柄时将其关闭。NT 句柄持有对该资源的引用，因此必须先显式释放该句柄，才能释放底层信号量。

此外，可以使用与对象关联的全局共享 D3DKMT 句柄，将使用VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT导出的 Vulkan 信号量对象导入 CUDA，如下所示。由于全局共享的 D3DKMT 句柄不持有对底层信号量的引用，因此在对该资源的所有其他引用被销毁时，它会自动销毁。

  void importCudaExternalSemaphore(
      cudaExternalSemaphore_t &cudaSem, VkSemaphore &vkSem,
      VkExternalSemaphoreHandleTypeFlagBits handleType) {
    cudaExternalSemaphoreHandleDesc externalSemaphoreHandleDesc = {};

#ifdef _VK_TIMELINE_SEMAPHORE
    if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) {
      externalSemaphoreHandleDesc.type =
          cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32;
    } else if (handleType &
               VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) {
      externalSemaphoreHandleDesc.type =
          cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32;
    } else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) {
      externalSemaphoreHandleDesc.type =
          cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd;
    }
#else
    if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) {
      externalSemaphoreHandleDesc.type =
          cudaExternalSemaphoreHandleTypeOpaqueWin32;
    } else if (handleType &
               VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) {
      externalSemaphoreHandleDesc.type =
          cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt;
    } else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) {
      externalSemaphoreHandleDesc.type =
          cudaExternalSemaphoreHandleTypeOpaqueFd;
    }
#endif /* _VK_TIMELINE_SEMAPHORE */
    else {
      throw std::runtime_error("Unknown handle type requested!");
    }

#ifdef _WIN64
    externalSemaphoreHandleDesc.handle.win32.handle =
        (HANDLE)getSemaphoreHandle(vkSem, handleType);
#else
    externalSemaphoreHandleDesc.handle.fd =
        (int)(uintptr_t)getSemaphoreHandle(vkSem, handleType);
#endif

    externalSemaphoreHandleDesc.flags = 0;

    checkCudaErrors(
        cudaImportExternalSemaphore(&cudaSem, &externalSemaphoreHandleDesc));
  }

4.21.2.1.9. 向已导入的同步对象发出信号/等待已导入的同步对象

可以向已导入的 Vulkan 信号量发出信号并等待该信号量，如下所示。向信号量发出信号会将其设置为已发出信号状态；对于时间线信号量，还会将计数器设置为信号调用中指定的值。等待此信号的相应等待操作必须在 Vulkan 中发起。此外，对于二进制信号量，等待此信号的等待操作必须在此信号发出之后才能发起。

等待信号量会一直等待，直到其达到已发出信号状态或指定的等待值。随后，已发出信号的二进制信号量会将其重置回未发出信号状态。此等待操作所等待的相应信号必须在 Vulkan 中发出。此外，对于二进制信号量，必须先发出信号，才能发起此等待操作。

在以下摘自simpleVulkan示例的代码中，只有在 Vulkan 向用于顶点缓冲区的信号量发出信号后，才会调用模拟步骤/CUDA 核函数。在模拟步骤之后，CUDA 会向另一个信号量发出信号，或者在使用时间线信号量的情况下递增同一个信号量，从而使等待此信号量的 Vulkan 部分能够使用更新后的顶点缓冲区继续渲染。

#ifdef _VK_TIMELINE_SEMAPHORE
    static uint64_t waitValue = 1;
    static uint64_t signalValue = 2;

    cudaExternalSemaphoreWaitParams waitParams = {};
    waitParams.flags = 0;
    waitParams.params.fence.value = waitValue;

    cudaExternalSemaphoreSignalParams signalParams = {};
    signalParams.flags = 0;
    signalParams.params.fence.value = signalValue;
    // Wait for vulkan to complete it's work
    checkCudaErrors(cudaWaitExternalSemaphoresAsync(&m_cudaTimelineSemaphore,
                                                    &waitParams, 1, m_stream));
    // Now step the simulation, call CUDA kernel
    m_sim.stepSimulation(time, m_stream);
    // Signal vulkan to continue with the updated buffers
    checkCudaErrors(cudaSignalExternalSemaphoresAsync(
        &m_cudaTimelineSemaphore, &signalParams, 1, m_stream));

    waitValue += 2;
    signalValue += 2;
#else
    cudaExternalSemaphoreWaitParams waitParams = {};
    waitParams.flags = 0;
    waitParams.params.fence.value = 0;

    cudaExternalSemaphoreSignalParams signalParams = {};
    signalParams.flags = 0;
    signalParams.params.fence.value = 0;

    // Wait for vulkan to complete it's work
    checkCudaErrors(cudaWaitExternalSemaphoresAsync(&m_cudaWaitSemaphore,
                                                    &waitParams, 1, m_stream));
    // Now step the simulation, call CUDA kernel
    m_sim.stepSimulation(time, m_stream);
    // Signal vulkan to continue with the updated buffers
    checkCudaErrors(cudaSignalExternalSemaphoresAsync(
        &m_cudaSignalSemaphore, &signalParams, 1, m_stream));
#endif /* _VK_TIMELINE_SEMAPHORE */

4.21.2.1.10. OpenGL 互操作性

OpenGL 互操作性中概述的传统 OpenGL-CUDA 互操作方式，是由 CUDA 直接使用在 OpenGL 中创建的句柄。但是，由于 OpenGL 也可以使用在 Vulkan 中创建的内存和同步对象，因此还存在另一种实现 OpenGL-CUDA 互操作的方式。其基本思想是，可以将由 Vulkan 导出的内存和同步对象同时导入 OpenGL 和 CUDA，然后用它们来协调 OpenGL 与 CUDA 之间的内存访问。有关如何导入由 Vulkan 导出的内存和同步对象的更多详细信息，请参阅以下 OpenGL 扩展：

GL_EXT_memory_object
GL_EXT_memory_object_fd
GL_EXT_memory_object_win32
GL_EXT_semaphore
GL_EXT_semaphore_fd
GL_EXT_semaphore_win32

4.21.2.2. Direct3D 互操作性

Direct3D11 和 Direct3D12 均支持将 Direct3D[11|12] 资源导入 CUDA。我们仅讨论 Direct3D12，有关 Direct3D11，请参阅 CUDA 编程指南 12.9。

4.21.2.2.1. 匹配设备 LUID

导入由 Direct3D12 导出的内存和同步对象时，必须在创建这些对象的同一设备上导入并映射它们。通过比较 CUDA 设备的 LUID 与 Direct3D12 设备的 LUID，可以确定与创建这些对象的 Direct3D12 设备相对应的 CUDA 设备，如以下代码示例所示。注意，Direct3D12 设备不得在链接节点适配器上创建，即ID3D12Device::GetNodeCount返回的节点数量必须为 1。

int getCudaDeviceForD3D12Device(ID3D12Device *d3d12Device) {
    LUID d3d12Luid = d3d12Device->GetAdapterLuid();

    int cudaDeviceCount;
    cudaGetDeviceCount(&cudaDeviceCount);

    for (int cudaDevice = 0; cudaDevice < cudaDeviceCount; cudaDevice++) {
        cudaDeviceProp deviceProp;
        cudaGetDeviceProperties(&deviceProp, cudaDevice);
        char *cudaLuid = deviceProp.luid;

        if (!memcmp(&d3d12Luid.LowPart, cudaLuid, sizeof(d3d12Luid.LowPart)) &&
            !memcmp(&d3d12Luid.HighPart, cudaLuid + sizeof(d3d12Luid.LowPart), sizeof(d3d12Luid.HighPart))) {
            return cudaDevice;
        }
    }
    return cudaInvalidDeviceId;
}

4.21.2.2.2. 导入内存对象

从 NT 句柄导入内存对象有几种不同的方法。注意，当不再需要 NT 句柄时，应用程序有责任将其关闭。NT 句柄持有对资源的引用，因此必须显式释放该句柄，才能释放底层内存。导入 Direct3D 资源时，必须设置标志cudaExternalMemoryDedicated，如下面的代码片段所示。

在调用ID3D12Device::CreateHeap时通过设置标志D3D12_HEAP_FLAG_SHARED创建的可共享 Direct3D12 堆内存对象，可以使用与该对象关联的 NT 句柄导入 CUDA，如下所示。

cudaExternalMemory_t importD3D12HeapFromNTHandle(HANDLE handle, unsigned long long size) {
    cudaExternalMemory_t extMem = NULL;
    cudaExternalMemoryHandleDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.type = cudaExternalMemoryHandleTypeD3D12Heap;
    desc.handle.win32.handle = (void *)handle;
    desc.size = size;

    cudaImportExternalMemory(&extMem, &desc);

    // Input parameter 'handle' should be closed if it's not needed anymore
    CloseHandle(handle);

    return extMem;
}

如果存在命名句柄，也可以使用该句柄导入可共享的 Direct3D12 堆内存对象：

cudaExternalMemory_t importD3D12HeapFromNamedNTHandle(LPCWSTR name, unsigned long long size) {
    cudaExternalMemory_t extMem = NULL;
    cudaExternalMemoryHandleDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.type = cudaExternalMemoryHandleTypeD3D12Heap;
    desc.handle.win32.name = (void *)name;
    desc.size = size;

    cudaImportExternalMemory(&extMem, &desc);

    return extMem;
}

在调用D3D12Device::CreateCommittedResource时通过设置标志D3D12_HEAP_FLAG_SHARED创建的可共享 Direct3D12 已提交资源，可以使用与该对象关联的 NT 句柄导入 CUDA，如下所示。导入 Direct3D12 已提交资源时，必须设置标志cudaExternalMemoryDedicated。

cudaExternalMemory_t importD3D12CommittedResourceFromNTHandle(HANDLE handle, unsigned long long size) {
    cudaExternalMemory_t extMem = NULL;
    cudaExternalMemoryHandleDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.type = cudaExternalMemoryHandleTypeD3D12Resource;
    desc.handle.win32.handle = (void *)handle;
    desc.size = size;
    desc.flags |= cudaExternalMemoryDedicated;

    cudaImportExternalMemory(&extMem, &desc);

    // Input parameter 'handle' should be closed if it's not needed anymore
    CloseHandle(handle);

    return extMem;
}

如果存在命名句柄，也可以使用该句柄导入可共享的 Direct3D12 已提交资源，如下所示。

cudaExternalMemory_t importD3D12CommittedResourceFromNamedNTHandle(LPCWSTR name, unsigned long long size) {
    cudaExternalMemory_t extMem = NULL;
    cudaExternalMemoryHandleDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.type = cudaExternalMemoryHandleTypeD3D12Resource;
    desc.handle.win32.name = (void *)name;
    desc.size = size;
    desc.flags |= cudaExternalMemoryDedicated;

    cudaImportExternalMemory(&extMem, &desc);

    return extMem;
}

4.21.2.2.3. 将缓冲区映射到导入的内存对象

可以将设备指针映射到导入的内存对象，如下所示。映射的偏移量和大小必须与使用相应的 Direct3D12 API 创建映射时指定的值一致。所有映射的设备指针都必须使用cudaFree()释放。

void * mapBufferOntoExternalMemory(cudaExternalMemory_t extMem, unsigned long long offset, unsigned long long size) {
    void *ptr = NULL;
    cudaExternalMemoryBufferDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.offset = offset;
    desc.size = size;

    cudaExternalMemoryGetMappedBuffer(&ptr, extMem, &desc);

    // Note: 'ptr' must eventually be freed using cudaFree()
    return ptr;
}

4.21.2.2.4. 将多级渐远纹理数组映射到导入的内存对象

可以将 CUDA 多级渐远纹理数组映射到导入的内存对象，如下所示。偏移量、维度、格式和 mip 级别数必须与使用相应的 Direct3D12 API 创建映射时指定的值一致。此外，如果该多级渐远纹理数组可以在 Direct3D12 中绑定为渲染目标，则必须设置标志cudaArrayColorAttachment。所有映射的多级渐远纹理数组都必须使用cudaFreeMipmappedArray()释放。以下代码示例展示了在将多级渐远纹理数组映射到导入的内存对象时，如何将参数转换为相应的 CUDA 参数。

cudaMipmappedArray_t mapMipmappedArrayOntoExternalMemory(cudaExternalMemory_t extMem, unsigned long long offset, cudaChannelFormatDesc *formatDesc, cudaExtent *extent, unsigned int flags, unsigned int numLevels) {
    cudaMipmappedArray_t mipmap = NULL;
    cudaExternalMemoryMipmappedArrayDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.offset = offset;
    desc.formatDesc = *formatDesc;
    desc.extent = *extent;
    desc.flags = flags;
    desc.numLevels = numLevels;

    // Note: 'mipmap' must eventually be freed using cudaFreeMipmappedArray()
    cudaExternalMemoryGetMappedMipmappedArray(&mipmap, extMem, &desc);

    return mipmap;
}

cudaChannelFormatDesc getCudaChannelFormatDescForDxgiFormat(DXGI_FORMAT dxgiFormat)
{
    cudaChannelFormatDesc d;

    memset(&d, 0, sizeof(d));

    switch (dxgiFormat) {
        case DXGI_FORMAT_R8_UINT:            d.x = 8;  d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R8_SINT:            d.x = 8;  d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R8G8_UINT:          d.x = 8;  d.y = 8;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R8G8_SINT:          d.x = 8;  d.y = 8;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R8G8B8A8_UINT:      d.x = 8;  d.y = 8;  d.z = 8;  d.w = 8;  d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R8G8B8A8_SINT:      d.x = 8;  d.y = 8;  d.z = 8;  d.w = 8;  d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R16_UINT:           d.x = 16; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R16_SINT:           d.x = 16; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R16G16_UINT:        d.x = 16; d.y = 16; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R16G16_SINT:        d.x = 16; d.y = 16; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R16G16B16A16_UINT:  d.x = 16; d.y = 16; d.z = 16; d.w = 16; d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R16G16B16A16_SINT:  d.x = 16; d.y = 16; d.z = 16; d.w = 16; d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R32_UINT:           d.x = 32; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R32_SINT:           d.x = 32; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R32_FLOAT:          d.x = 32; d.y = 0;  d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindFloat;    break;
        case DXGI_FORMAT_R32G32_UINT:        d.x = 32; d.y = 32; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R32G32_SINT:        d.x = 32; d.y = 32; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R32G32_FLOAT:       d.x = 32; d.y = 32; d.z = 0;  d.w = 0;  d.f = cudaChannelFormatKindFloat;    break;
        case DXGI_FORMAT_R32G32B32A32_UINT:  d.x = 32; d.y = 32; d.z = 32; d.w = 32; d.f = cudaChannelFormatKindUnsigned; break;
        case DXGI_FORMAT_R32G32B32A32_SINT:  d.x = 32; d.y = 32; d.z = 32; d.w = 32; d.f = cudaChannelFormatKindSigned;   break;
        case DXGI_FORMAT_R32G32B32A32_FLOAT: d.x = 32; d.y = 32; d.z = 32; d.w = 32; d.f = cudaChannelFormatKindFloat;    break;
        default: assert(0);
    }
    return d;
}

cudaExtent getCudaExtentForD3D12Extent(UINT64 width, UINT height, UINT16 depthOrArraySize, D3D12_SRV_DIMENSION d3d12SRVDimension) {
    cudaExtent e = { 0, 0, 0 };

    switch (d3d12SRVDimension) {
        case D3D12_SRV_DIMENSION_TEXTURE1D:        e.width = width; e.height = 0;      e.depth = 0;                break;
        case D3D12_SRV_DIMENSION_TEXTURE2D:        e.width = width; e.height = height; e.depth = 0;                break;
        case D3D12_SRV_DIMENSION_TEXTURE3D:        e.width = width; e.height = height; e.depth = depthOrArraySize; break;
        case D3D12_SRV_DIMENSION_TEXTURECUBE:      e.width = width; e.height = height; e.depth = depthOrArraySize; break;
        case D3D12_SRV_DIMENSION_TEXTURE1DARRAY:   e.width = width; e.height = 0;      e.depth = depthOrArraySize; break;
        case D3D12_SRV_DIMENSION_TEXTURE2DARRAY:   e.width = width; e.height = height; e.depth = depthOrArraySize; break;
        case D3D12_SRV_DIMENSION_TEXTURECUBEARRAY: e.width = width; e.height = height; e.depth = depthOrArraySize; break;
        default: assert(0);
    }

    return e;
}

unsigned int getCudaMipmappedArrayFlagsForD3D12Resource(D3D12_SRV_DIMENSION d3d12SRVDimension, D3D12_RESOURCE_FLAGS d3d12ResourceFlags, bool allowSurfaceLoadStore) {
    unsigned int flags = 0;

    switch (d3d12SRVDimension) {
        case D3D12_SRV_DIMENSION_TEXTURECUBE:      flags |= cudaArrayCubemap;                    break;
        case D3D12_SRV_DIMENSION_TEXTURECUBEARRAY: flags |= cudaArrayCubemap | cudaArrayLayered; break;
        case D3D12_SRV_DIMENSION_TEXTURE1DARRAY:   flags |= cudaArrayLayered;                    break;
        case D3D12_SRV_DIMENSION_TEXTURE2DARRAY:   flags |= cudaArrayLayered;                    break;
        default: break;
    }

    if (d3d12ResourceFlags & D3D12_RESOURCE_FLAG_ALLOW_RENDER_TARGET) {
        flags |= cudaArrayColorAttachment;
    }
    if (allowSurfaceLoadStore) {
        flags |= cudaArraySurfaceLoadStore;
    }

    return flags;
}

4.21.2.2.5. 导入同步对象

在调用ID3D12Device::CreateFence时通过设置标志D3D12_FENCE_FLAG_SHARED创建的可共享 Direct3D12 栅栏对象，可以使用与该对象关联的 NT 句柄导入 CUDA，如下所示。注意，当不再需要该句柄时，应用程序有责任将其关闭。NT 句柄持有对资源的引用，因此必须显式释放该句柄，才能释放底层信号量。

cudaExternalSemaphore_t importD3D12FenceFromNTHandle(HANDLE handle) {
    cudaExternalSemaphore_t extSem = NULL;
    cudaExternalSemaphoreHandleDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.type = cudaExternalSemaphoreHandleTypeD3D12Fence;
    desc.handle.win32.handle = handle;

    cudaImportExternalSemaphore(&extSem, &desc);

    // Input parameter 'handle' should be closed if it's not needed anymore
    CloseHandle(handle);

    return extSem;
}

如果存在命名句柄，也可以使用该句柄导入可共享的 Direct3D12 栅栏对象，如下所示。

cudaExternalSemaphore_t importD3D12FenceFromNamedNTHandle(LPCWSTR name) {
    cudaExternalSemaphore_t extSem = NULL;
    cudaExternalSemaphoreHandleDesc desc = {};
 
    memset(&desc, 0, sizeof(desc));

    desc.type = cudaExternalSemaphoreHandleTypeD3D12Fence;
    desc.handle.win32.name = (void *)name;

    cudaImportExternalSemaphore(&extSem, &desc);

    return extSem;
}

4.21.2.2.6. 向导入的同步对象发出信号或等待这些对象

一旦从 Direct3D12 导入了带有栅栏的信号量，就可以向其发出信号并等待这些信号量。

向栅栏对象发出信号会设置其值。等待此信号的相应等待操作必须在 Direct3D12 中发出。注意，等待此信号的等待操作必须在发出此信号之后发出。

void signalExternalSemaphore(cudaExternalSemaphore_t extSem, unsigned long long value, cudaStream_t stream) {
    cudaExternalSemaphoreSignalParams params = {};

    memset(&params, 0, sizeof(params));

    params.params.fence.value = value;

    cudaSignalExternalSemaphoresAsync(&extSem, &params, 1, stream);
}

栅栏对象会等待，直到其值大于或等于指定值。它所等待的相应信号必须在 Direct3D12 中发出。注意，必须先发出该信号，才能发出此等待操作。

void waitExternalSemaphore(cudaExternalSemaphore_t extSem, unsigned long long value, cudaStream_t stream) {
    cudaExternalSemaphoreWaitParams params = {};

    memset(&params, 0, sizeof(params));

    params.params.fence.value = value;

    cudaWaitExternalSemaphoresAsync(&extSem, &params, 1, stream);
}

4.21.2.3. NVIDIA 软件通信接口互操作性（NVSCI）

NvSciBuf 和 NvSciSync 是为满足以下用途而开发的接口：

NvSciBuf：允许应用程序在内存中分配和交换缓冲区
NvSciSync：允许应用程序在操作边界处管理同步对象

有关这些接口的更多详细信息，请参阅：https://docs.nvidia.com/drive。

4.21.2.3.1. 导入内存对象

要分配与给定 CUDA 设备兼容的 NvSciBuf 对象，必须在 NvSciBuf 属性列表中通过NvSciBufGeneralAttrKey_GpuId设置相应的 GPU ID，如下所示。应用程序还可以选择指定以下属性：

NvSciBufGeneralAttrKey_NeedCpuAccess：指定缓冲区是否需要 CPU 访问
NvSciBufRawBufferAttrKey_Align：指定NvSciBufType_RawBuffer的对齐要求
NvSciBufGeneralAttrKey_RequiredPerm：对于每个 NvSciBuf 内存对象实例，可以为不同的 UMD 配置不同的访问权限。例如，要为 GPU 提供对缓冲区的只读访问权限，可使用NvSciBufObjDupWithReducePerm()创建一个 NvSciBuf 对象副本，并将NvSciBufAccessPerm_Readonly作为输入参数。然后，将这个新创建的、权限降低的对象副本导入 CUDA，如下所示
NvSciBufGeneralAttrKey_EnableGpuCache：用于控制 GPU L2 可缓存性
NvSciBufGeneralAttrKey_EnableGpuCompression：用于指定 GPU 压缩

注意

有关这些属性及其有效输入选项的更多详细信息，请参阅 NvSciBuf 文档。

以下代码片段展示了它们的用法示例。

NvSciBufObj createNvSciBufObject() {
   // Raw Buffer Attributes for CUDA
    NvSciBufType bufType = NvSciBufType_RawBuffer;
    uint64_t rawsize = SIZE;
    uint64_t align = 0;
    bool cpuaccess_flag = true;
    NvSciBufAttrValAccessPerm perm = NvSciBufAccessPerm_ReadWrite;

    NvSciRmGpuId gpuid[] ={};
    CUuuid uuid;
    cuDeviceGetUuid(&uuid, dev);

    memcpy(&gpuid[0].bytes, &uuid.bytes, sizeof(uuid.bytes));
    // Disable cache on dev
    NvSciBufAttrValGpuCache gpuCache[] = {{gpuid[0], false}};
    NvSciBufAttrValGpuCompression gpuCompression[] = {{gpuid[0], NvSciBufCompressionType_GenericCompressible}};
    // Fill in values
    NvSciBufAttrKeyValuePair rawbuffattrs[] = {
         { NvSciBufGeneralAttrKey_Types, &bufType, sizeof(bufType) },
         { NvSciBufRawBufferAttrKey_Size, &rawsize, sizeof(rawsize) },
         { NvSciBufRawBufferAttrKey_Align, &align, sizeof(align) },
         { NvSciBufGeneralAttrKey_NeedCpuAccess, &cpuaccess_flag, sizeof(cpuaccess_flag) },
         { NvSciBufGeneralAttrKey_RequiredPerm, &perm, sizeof(perm) },
         { NvSciBufGeneralAttrKey_GpuId, &gpuid, sizeof(gpuid) },
         { NvSciBufGeneralAttrKey_EnableGpuCache, &gpuCache, sizeof(gpuCache) },
         { NvSciBufGeneralAttrKey_EnableGpuCompression, &gpuCompression, sizeof(gpuCompression) }
    };

    // Create list by setting attributes
    err = NvSciBufAttrListSetAttrs(attrListBuffer, rawbuffattrs,
            sizeof(rawbuffattrs)/sizeof(NvSciBufAttrKeyValuePair));

    NvSciBufAttrListCreate(NvSciBufModule, &attrListBuffer);

    // Reconcile And Allocate
    NvSciBufAttrListReconcile(&attrListBuffer, 1, &attrListReconciledBuffer,
                       &attrListConflictBuffer)
    NvSciBufObjAlloc(attrListReconciledBuffer, &bufferObjRaw);
    return bufferObjRaw;
}

NvSciBufObj bufferObjRo; // Readonly NvSciBuf memory obj
// Create a duplicate handle to the same memory buffer with reduced permissions
NvSciBufObjDupWithReducePerm(bufferObjRaw, NvSciBufAccessPerm_Readonly, &bufferObjRo);
return bufferObjRo;

可以使用 NvSciBufObj 句柄将分配的 NvSciBuf 内存对象导入 CUDA，如下所示。应用程序应查询分配的 NvSciBufObj，以获取填写 CUDA 外部内存描述符所需的属性。注意，属性列表和 NvSciBuf 对象应由应用程序维护。如果导入 CUDA 的 NvSciBuf 对象也被其他驱动程序映射，那么应用程序必须根据NvSciBufGeneralAttrKey_GpuSwNeedCacheCoherency输出属性的值，使用 NvSciSync 对象（请参阅导入同步对象）作为适当的屏障，以维持 CUDA 与其他驱动程序之间的相干性。

注意

有关如何分配和维护 NvSciBuf 对象的更多详细信息，请参阅NvSciBuf API 文档。

cudaExternalMemory_t importNvSciBufObject (NvSciBufObj bufferObjRaw) {

    /*************** Query NvSciBuf Object **************/
    NvSciBufAttrKeyValuePair bufattrs[] = {
                { NvSciBufRawBufferAttrKey_Size, NULL, 0 },
                { NvSciBufGeneralAttrKey_GpuSwNeedCacheCoherency, NULL, 0 },
                { NvSciBufGeneralAttrKey_EnableGpuCompression, NULL, 0 }
    };
    NvSciBufAttrListGetAttrs(retList, bufattrs,
        sizeof(bufattrs)/sizeof(NvSciBufAttrKeyValuePair)));
                ret_size = *(static_cast<const uint64_t*>(bufattrs[0].value));

    // Note cache and compression are per GPU attributes, so read values for specific gpu by comparing UUID
    // Read cacheability granted by NvSciBuf
    int numGpus = bufattrs[1].len / sizeof(NvSciBufAttrValGpuCache);
    NvSciBufAttrValGpuCache[] cacheVal = (NvSciBufAttrValGpuCache *)bufattrs[1].value;
    bool ret_cacheVal;
    for (int i = 0; i < numGpus; i++) {
        if (memcmp(gpuid[0].bytes, cacheVal[i].gpuId.bytes, sizeof(CUuuid)) == 0) {
            ret_cacheVal = cacheVal[i].cacheability);
        }
    }

    // Read compression granted by NvSciBuf
    numGpus = bufattrs[2].len / sizeof(NvSciBufAttrValGpuCompression);
    NvSciBufAttrValGpuCompression[] compVal = (NvSciBufAttrValGpuCompression *)bufattrs[2].value;
    NvSciBufCompressionType ret_compVal;
    for (int i = 0; i < numGpus; i++) {
        if (memcmp(gpuid[0].bytes, compVal[i].gpuId.bytes, sizeof(CUuuid)) == 0) {
            ret_compVal = compVal[i].compressionType);
        }
    }

    /*************** NvSciBuf Registration With CUDA **************/

    // Fill up CUDA_EXTERNAL_MEMORY_HANDLE_DESC
    cudaExternalMemoryHandleDesc memHandleDesc;
    memset(&memHandleDesc, 0, sizeof(memHandleDesc));
    memHandleDesc.type = cudaExternalMemoryHandleTypeNvSciBuf;
    memHandleDesc.handle.nvSciBufObject = bufferObjRaw;
    // Set the NvSciBuf object with required access permissions in this step
    memHandleDesc.handle.nvSciBufObject = bufferObjRo;
    memHandleDesc.size = ret_size;
    cudaImportExternalMemory(&extMemBuffer, &memHandleDesc);
    return extMemBuffer;
 }

4.21.2.3.2. 将缓冲区映射到导入的内存对象

可以将设备指针映射到导入的内存对象，如下所示。可以根据分配的NvSciBufObj的属性填写映射的偏移量和大小。所有映射的设备指针都必须使用cudaFree()释放。

void * mapBufferOntoExternalMemory(cudaExternalMemory_t extMem, unsigned long long offset, unsigned long long size) {
    void *ptr = NULL;
    cudaExternalMemoryBufferDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.offset = offset;
    desc.size = size;

    cudaExternalMemoryGetMappedBuffer(&ptr, extMem, &desc);

    // Note: 'ptr' must eventually be freed using cudaFree()
    return ptr;
}

4.21.2.3.3. 将多级渐远纹理数组映射到导入的内存对象

可以将 CUDA 多级渐远纹理数组映射到导入的内存对象，如下所示。可以根据分配的NvSciBufObj的属性填写偏移量、维度和格式。所有映射的多级渐远纹理数组都必须使用cudaFreeMipmappedArray()释放。以下代码示例展示了在将多级渐远纹理数组映射到导入的内存对象时，如何将 NvSciBuf 属性转换为相应的 CUDA 参数。

注意

mip 级别数必须为 1。

cudaMipmappedArray_t mapMipmappedArrayOntoExternalMemory(cudaExternalMemory_t extMem, unsigned long long offset, cudaChannelFormatDesc *formatDesc, cudaExtent *extent, unsigned int flags, unsigned int numLevels) {
    cudaMipmappedArray_t mipmap = NULL;
    cudaExternalMemoryMipmappedArrayDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.offset = offset;
    desc.formatDesc = *formatDesc;
    desc.extent = *extent;
    desc.flags = flags;
    desc.numLevels = numLevels;

    // Note: 'mipmap' must eventually be freed using cudaFreeMipmappedArray()
    cudaExternalMemoryGetMappedMipmappedArray(&mipmap, extMem, &desc);

    return mipmap;
}

4.21.2.3.4. 导入同步对象

可以使用cudaDeviceGetNvSciSyncAttributes()生成与给定 CUDA 设备兼容的 NvSciSync 属性。返回的属性列表可用于创建一个保证与给定 CUDA 设备兼容的NvSciSyncObj。

NvSciSyncObj createNvSciSyncObject() {
    NvSciSyncObj nvSciSyncObj
    int cudaDev0 = 0;
    int cudaDev1 = 1;
    NvSciSyncAttrList signalerAttrList = NULL;
    NvSciSyncAttrList waiterAttrList = NULL;
    NvSciSyncAttrList reconciledList = NULL;
    NvSciSyncAttrList newConflictList = NULL;

    NvSciSyncAttrListCreate(module, &signalerAttrList);
    NvSciSyncAttrListCreate(module, &waiterAttrList);
    NvSciSyncAttrList unreconciledList[2] = {NULL, NULL};
    unreconciledList[0] = signalerAttrList;
    unreconciledList[1] = waiterAttrList;

    cudaDeviceGetNvSciSyncAttributes(signalerAttrList, cudaDev0, CUDA_NVSCISYNC_ATTR_SIGNAL);
    cudaDeviceGetNvSciSyncAttributes(waiterAttrList, cudaDev1, CUDA_NVSCISYNC_ATTR_WAIT);

    NvSciSyncAttrListReconcile(unreconciledList, 2, &reconciledList, &newConflictList);

    NvSciSyncObjAlloc(reconciledList, &nvSciSyncObj);

    return nvSciSyncObj;
}

可以使用 NvSciSyncObj 句柄将 NvSciSync 对象（按上述方式创建）导入 CUDA，如下所示。注意，即使在导入之后，NvSciSyncObj 句柄的所有权仍归应用程序所有。

cudaExternalSemaphore_t importNvSciSyncObject(void* nvSciSyncObj) {
    cudaExternalSemaphore_t extSem = NULL;
    cudaExternalSemaphoreHandleDesc desc = {};

    memset(&desc, 0, sizeof(desc));

    desc.type = cudaExternalSemaphoreHandleTypeNvSciSync;
    desc.handle.nvSciSyncObj = nvSciSyncObj;

    cudaImportExternalSemaphore(&extSem, &desc);

    // Deleting/Freeing the nvSciSyncObj beyond this point will lead to undefined behavior in CUDA

    return extSem;
}

4.21.2.3.5. 向导入的同步对象发出信号或等待这些对象

可以按下述方式向导入的NvSciSyncObj对象发出信号。向由 NvSciSync 支持的信号量对象发出信号，会初始化作为输入传入的栅栏参数。与上述信号对应的等待操作会等待此栅栏参数。此外，等待此信号的等待操作必须在发出此信号之后发出。如果将标志设置为cudaExternalSemaphoreSignalSkipNvSciBufMemSync，则会跳过默认作为发出信号操作的一部分执行的内存同步操作（针对该进程中所有导入的 NvSciBuf）。当NvsciBufGeneralAttrKey_GpuSwNeedCacheCoherency为 FALSE 时，应设置此标志。

void signalExternalSemaphore(cudaExternalSemaphore_t extSem, cudaStream_t stream, void *fence) {
    cudaExternalSemaphoreSignalParams signalParams = {};

    memset(&signalParams, 0, sizeof(signalParams));

    signalParams.params.nvSciSync.fence = (void*)fence;
    signalParams.flags = 0; //OR cudaExternalSemaphoreSignalSkipNvSciBufMemSync

    cudaSignalExternalSemaphoresAsync(&extSem, &signalParams, 1, stream);

}

可以按下述方式等待导入的NvSciSyncObj对象。等待由 NvSciSync 支持的信号量对象，会一直等待，直到相应的信号发出者向输入的栅栏参数发出信号。此外，必须先发出信号，才能发出等待操作。如果将标志设置为cudaExternalSemaphoreWaitSkipNvSciBufMemSync，则会跳过默认作为发出信号操作的一部分执行的内存同步操作（针对该进程中所有导入的 NvSciBuf）。当NvsciBufGeneralAttrKey_GpuSwNeedCacheCoherency为 FALSE 时，应设置此标志。

void waitExternalSemaphore(cudaExternalSemaphore_t extSem, cudaStream_t stream, void *fence) {
     cudaExternalSemaphoreWaitParams waitParams = {};

    memset(&waitParams, 0, sizeof(waitParams));

    waitParams.params.nvSciSync.fence = (void*)fence;
    waitParams.flags = 0; //OR cudaExternalSemaphoreWaitSkipNvSciBufMemSync

    cudaWaitExternalSemaphoresAsync(&extSem, &waitParams, 1, stream);
}