15

如何在 CUDA 中处理 OpenGL 立方体贴图纹理？

当想要在 CUDA 内核中使用 OpenGL 纹理时，要做的一件事是从注册的图像和映射资源中检索一个 CUDA 数组，在这种情况下是一个纹理。在驱动程序 API 中，它是通过cuGraphicsSubResourceGetMappedArray调用完成的，这在 2D 纹理的情况下不是问题。但是在谈到上述立方体贴图时，该函数的第三个参数需要一个面枚举（如CU_CUBEMAP_FACE_POSITIVE_X）。因此出现了一些问题——当一个人通过这样一个枚举时，返回的纹理数组将只包含那个特定人脸的数据，对吧？那么如何将立方体纹理作为一个整体来进行立方体映射，同样：

color = texCube(cubeMap, x, y, z);

还是在 CUDA 内核中不可能这样做，并且需要在用户代码中使用具有适当计算和采样的 2D 纹理？

4

2 回答 2

4

好的 - 我自己设法解决了这个问题，虽然解决方案不像使用另一个 CUDA 函数那么简单。

要将 CUDA 纹理引用与任何纹理绑定，无论是从 OpenGL 还是 D3D 获得的纹理，都必须提供一个映射到资源的 CUDA 数组，cuGraphicsSubResourceGetMappedArray用于检索它。正如我在问题中提到的，在一维或二维纹理的情况下很简单。但是对于其他可用类型，它会更加复杂。

在任何时候，我们都需要一个引用绑定到的 CUDA 数组。立方体贴图纹理也是如此。但在这种情况下，阵列必须是 3D 阵列。问题在于 CUDA 驱动程序 API 仅提供上述函数来从此类纹理资源中检索单个图层，并将其映射到单个二维数组。为了得到我们想要的，我们必须让自己成为一个包含所有层（或立方体贴图的面）的 3D 数组。

首先，我们必须使用上述函数为每个层/面获取数组。下一步是通过调用来创建 3D 数组 cuArray3DCreate，并输入适当的参数集（大小/层数、细节级别、数据格式、每个纹素的通道数和一些标志）。然后我们必须通过一系列调用将图层的数组复制cuMemcpy3D到 3D 数组，每个图层/面数组调用一个。

cuTexRefSetArray最后，我们使用我们创建并复制到的 3D 数组来设置目标 CUDA 纹理参考。在设备代码内部，我们创建了一个具有适当纹理类型和模式（float4 和立方体贴图）的引用，并使用texCubemap.

下面我放了一个函数片段，它可以在CIRT 存储库（cirt_server.c 文件，函数cirtTexImage3D）中完整地获得所有这些功能。

//...
if (result)
{
    // Create a 3D array...
    CUDA_ARRAY3D_DESCRIPTOR layeredTextureDescr;
    layeredTextureDescr.Width = w;
    layeredTextureDescr.Height = h;
    layeredTextureDescr.Depth = d;
    layeredTextureDescr.Format = map_type_to_format(type);
    layeredTextureDescr.NumChannels = format == CIRT_RGB ? CIRT_RGBA : format;
    layeredTextureDescr.Flags = map_target_to_flags(target);

    if (result) result = LogCUDADriverCall(cuArray3DCreate(&hTexRefArray, &layeredTextureDescr),
        FUN_NAME(": cuArray3DCreate_tex3D"), __FILE_LINE__);

    // Copy the acquired layer/face arrays into the collective 3D one...
    CUDA_MEMCPY3D layerCopyDescr;
    layerCopyDescr.srcMemoryType = CU_MEMORYTYPE_ARRAY;
    layerCopyDescr.srcXInBytes = 0;
    layerCopyDescr.srcZ = 0;
    layerCopyDescr.srcY = 0;
    layerCopyDescr.srcLOD = 0;

    layerCopyDescr.dstMemoryType = CU_MEMORYTYPE_ARRAY;
    layerCopyDescr.dstLOD = 0;

    layerCopyDescr.WidthInBytes = layeredTextureDescr.NumChannels * w;
    layerCopyDescr.Height = h;
    layerCopyDescr.Depth = target == CIRT_TEXTURE_CUBE_MAP ? 1 : d;
    layerCopyDescr.dstArray = hTexRefArray;

    for (i = 0; i < num_layers; ++i)
    {
        layer = ((num_layers == 6) ? CU_CUBEMAP_FACE_POSITIVE_X + i : i);
        layerCopyDescr.dstXInBytes = 0;
        layerCopyDescr.dstY = 0;
        layerCopyDescr.dstZ = i;
        layerCopyDescr.srcArray = hLayres[i];

        if (result) result = LogCUDADriverCall(cuMemcpy3D(&layerCopyDescr), 
            FUN_NAME(": cuMemcpy3D_tex3D"), __FILE_LINE__);
    }

    // Finally bind the 3D array with texture reference...
    if (result) LogCUDADriverCall(cuTexRefSetArray(hTexRef, hTexRefArray, CU_TRSA_OVERRIDE_FORMAT),
        FUN_NAME(": cuTexRefSetArray_tex3D"), __FILE_LINE__);

    if (hLayres)
        free(hLayres);

    if (result)
        current->m_oTextureManager.m_cuTextureRes[current->m_oTextureManager.m_nTexCount++] = hTexResource;
}
//...

我现在只用立方体贴图检查过它，但它应该也适用于 3D 纹理。

于 2017-06-06T07:42:01.530 回答
0

我并不真正熟悉 CUDA，但我确实在 OpenGL 和 DirectX 方面有一些经验，而且我也熟悉 3D 图形渲染 API、库和管道，并且能够设置和使用这些 API。

当我看到你的问题时：

如何在 CUDA 中处理 OpenGL 立方体贴图纹理？

然后你继续解释它：

当想要在 CUDA 内核中使用 OpenGL 纹理时，要做的一件事是从注册的图像和映射资源中检索一个 CUDA 数组，在这种情况下是一个纹理。在驱动程序 API 中，它由 cuGraphicsSubResourceGetMappedArray 调用完成，在 2D 纹理的情况下这不是问题。但是在谈到上述立方体贴图时，此函数的第三个参数需要一个面枚举（如 CU_CUBEMAP_FACE_POSITIVE_X）。因此出现了一些问题 - 当一个人通过这样的枚举时，返回的纹理数组将只包含该特定面部的数据，对吗？那么如何将立方体纹理作为一个整体来进行立方体映射，同样：

color = texCube(cubeMap, x, y, z);

或者在 CUDA 内核中不可能这样做，并且需要使用 2D 纹理并在用户代码中进行适当的计算和采样？

我访问了 CUDA 的网站以获取他们的 API SDK 和编程文档。并找到了有问题的函数cuGraphicsSubResourceGetMappedArray()

CUresult cuGraphicsSubResourceGetMappedArray ( CUarray* pArray,
                                               CUgraphicsResource resource, 
                                               unsigned int arrayIndex,
                                               unsigned int mipLevel )

获取一个数组，通过该数组访问映射的图形资源的子资源。

参数

  • pArray - 返回的数组，通过该数组可以访问资源的子资源
  • resource - 要访问的映射资源
  • arrayIndex - 数组纹理的数组索引或立方体贴图面索引，由 CUarray_cubemap_face 为要访问的子资源的立方体贴图纹理定义
  • mipLevel - 要访问的子资源的 Mipmap 级别

退货

  • CUDA_SUCCESS, CUDA_ERROR_DEINITIALIZED, CUDA_ERROR_NOT_INITIALIZED,
  • CUDA_ERROR_INVALID_CONTEXT, CUDA_ERROR_INVALID_VALUE,
  • CUDA_ERROR_INVALID_HANDLE, CUDA_ERROR_NOT_MAPPED,
  • CUDA_ERROR_NOT_MAPPED_AS_ARRAY

描述

在*pArray 中返回一个数组，通过该数组可以访问对应于数组索引arrayIndex 和mipmap 级别mipLevel 的映射图形资源资源的子资源。每次映射资源时，*pArray 中设置的值可能会发生变化。

如果resource不是 atexture则无法通过 an 访问array并CUDA_ERROR_NOT_MAPPED_AS_ARRAY返回。如果arrayIndex不是有效array index的，resource则CUDA_ERROR_INVALID_VALUE返回。如果mipLevel不是有效mipmap level的，resource则CUDA_ERROR_INVALID_VALUE返回。如果资源不是mapped，则CUDA_ERROR_NOT_MAPPED返回。

注意：请注意，此函数也可能从以前的异步启动返回错误代码。

也可以看看：

cuGraphicsResourceGetMappedPointer

阅读更多：http ://docs.nvidia.com/cuda/cuda-driver-api/index.html#ixzz4ic22V4Dz 关注我们：@GPUComputing on Twitter | 英伟达在 Facebook

这个函数方法是在 NVidia CUDA 中找到的DriverAPI，而不是在他们的RuntimeAPI. 当了解具有 CUDA 功能的硬件时，可以在此处找到可编程管道Host和可编程管道之间的区别：http: //docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#axzz4ic6tFjXRDevice

2. 异构计算

CUDA 编程涉及同时在两个不同的平台上运行代码：具有一个或多个 CPU 的主机系统和一个或多个支持 CUDA 的 NVIDIA GPU 设备。

虽然 NVIDIA GPU 经常与图形相关联，但它们也是强大的算术引擎，能够并行运行数千个轻量级线程。这种能力使它们非常适合可以利用并行执行的计算。

但是，该设备基于与主机系统截然不同的设计，因此了解这些差异以及它们如何确定 CUDA 应用程序的性能以有效地使用 CUDA 非常重要。

  • 2.1。主机和设备之间 的区别主要区别在于线程模型和单独的物理内存：
    • 线程资源——主机系统上的执行管道可以支持有限数量的并发线程。如今拥有四个六核处理器的服务器只能同时运行 24 个线程（如果 CPU 支持超线程，则可以运行 48 个线程。）相比之下，CUDA 设备上最小的并行执行单元包含 32 个线程（称为线程束） . 现代 NVIDIA GPU 最多可支持每个多处理器并发 1536 个活动线程（请参阅 CUDA C 编程指南的特性和规范）在具有 16 个多处理器的 GPU 上，这会导致超过 24,000 个并发活动线程。
    • 线程- CPU 上的线程通常是重量级实体。操作系统必须在 CPU 执行通道上交换线程以提供多线程能力。上下文切换（当交换两个线程时）因此缓慢且昂贵。相比之下，GPU 上的线程非常轻量级。在一个典型的系统中，数千个线程排队等待工作（每个线程有 32 个线程）。如果 GPU 必须等待一个线程束，它只会开始在另一个线程上执行工作。因为单独的寄存器分配给所有活动线程，所以在 GPU 线程之间切换时不需要发生寄存器交换或其他状态。资源一直分配给每个线程，直到它完成执行。简而言之，CPU 内核旨在最大程度地减少每个线程一次或两个线程的延迟，
    • RAM - 主机系统和设备都有各自不同的附加物理内存。由于主机和设备内存由 PCI Express (PCIe) 总线分隔，因此主机内存中的项目有时必须通过总线与设备内存进行通信，反之亦然，如在启用 CUDA 的设备上运行什么？

这些是 CPU 主机和 GPU 设备在并行编程方面的主要硬件差异。其他差异将在本文档其他地方出现时进行讨论。考虑到这些差异的应用程序可以将主机和设备一起视为一个有凝聚力的异构系统，其中每个处理单元都被利用来完成它最擅长的工作：主机上的顺序工作和设备上的并行工作。

阅读更多内容：http ://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#ixzz4ic8ch2fq 关注我们：@GPUComputing on Twitter | 英伟达在 Facebook

现在知道 CUDA API 库有两种不同的 API，我们必须了解两者之间的区别：驱动程序和运行时 API 之间的区别

1.驱动和运行时API的区别

驱动程序和运行时 API 非常相似，并且大部分可以互换使用。但是，两者之间存在一些值得注意的关键差异。

复杂性与控制

运行时 API 通过提供隐式初始化、上下文管理和模块管理来简化设备代码管理。这导致代码更简单，但它也缺乏驱动程序 API 所具有的控制级别。

相比之下，驱动 API 提供了更细粒度的控制，尤其是对上下文和模块加载的控制。内核启动实现起来要复杂得多，因为必须使用显式函数调用指定执行配置和内核参数。然而，与运行时不同，所有内核在初始化期间自动加载并在程序运行期间一直保持加载状态，使用驱动程序 API 可以只保持当前需要加载的模块，甚至动态地重新加载模块。驱动程序 API 也与语言无关，因为它只处理 cubin 对象。

上下文管理

上下文管理可以通过驱动 API 来完成，但不会暴露在运行时 API 中。相反，运行时 API 自行决定为线程使用哪个上下文：如果已通过驱动程序 API 使调用线程成为当前上下文，则运行时将使用该上下文，但如果没有这样的上下文，它使用“主语境。” 主要上下文根据需要创建，每个设备每个进程一个，被引用计数，然后在不再引用它们时被销毁。在一个进程中，运行时 API 的所有用户都将共享主上下文，除非每个线程都有一个当前上下文。运行时使用的上下文，即当前上下文或主上下文，可以用 cudaDeviceSynchronize() 同步，并用 cudaDeviceReset() 销毁。

然而，将运行时 API 与主要上下文一起使用有其权衡。例如，如果所有插件在同一个进程中运行，它们将共享一个上下文，但可能无法相互通信，这可能会给用户编写较大软件包的插件带来麻烦。因此，如果其中一个插件在完成其所有 CUDA 工作后调用 cudaDeviceReset()，其他插件将失败，因为它们正在使用的上下文在他们不知情的情况下被破坏。为避免此问题，CUDA 客户端可以使用驱动程序 API 创建和设置当前上下文，然后使用运行时 API 来处理它。但是，上下文可能会消耗大量资源，例如设备内存、额外的主机线程以及设备上上下文切换的性能成本。

阅读更多：http ://docs.nvidia.com/cuda/cuda-driver-api/index.html#ixzz4icCoAXb7 关注我们：@GPUComputing 在 Twitter | 英伟达在 Facebook

由于这恰好出现在DriverAPI它对程序员的控制方面具有更大的灵活性，但也需要更多的责任来管理RuntimeAPI库在哪里做的事情更自动化但给你的控制更少。

这很明显，因为您提到您正在使用他们Kernels，但从他们对功能实现的描述中

 CUresult cuGraphicsSubResourceGetMappedArray ( CUarray* pArray,
                                                CUgraphicsResource resource, 
                                                unsigned int arrayIndex,
                                                unsigned int mipLevel )

文档告诉我，此函数采用的第一个参数是一个返回的数组，通过该数组可以访问资源的子资源。该函数的第二个参数是映射的图形资源本身。我相信的第三个参数是你有问题的参数，它是一个面的枚举类型，然后你问：当一个人通过这样一个枚举时，返回的纹理数组将只包含那个特定面的数据，正确的？根据我从文档中收集和理解的信息，这是array您的立方体贴图资源的索引值。

从他们的文档中可以看出：

arrayIndex - 数组纹理的数组索引或立方体贴图面索引，由 CUarray_cubemap_face 为要访问的子资源的立方体贴图纹理定义

阅读更多内容：http ://docs.nvidia.com/cuda/cuda-driver-api/index.html#ixzz4icHnwe9v 关注我们：@GPUComputing on Twitter | 英伟达在 Facebook

它恰好是构成典型立方体贴图unsigned int的纹理中的一个或索引位置，或者如果立方体的内部和外部都被映射，则最多。因此，如果我们查看立方体贴图以及纹理及其与伪代码的关系，我们可以看到：cube map6 faces12

// Texture
struct Texture {
    unsigned pixelsWidth;
    unsigned pixelsHeight;        
    // Other Texture member variables or fields here.
};

// Only interested in the actual size of the texture `width by height`
// where these would be used to map this texture to one of the 6 faces
// of a cube:

struct CubeMap {
    Texture face[6];
    // face[0] = frontFace
    // face[1] = backFace
    // face[2] = leftFace
    // face[3] = rightFace
    // face[4] = topFace
    // face[5] = bottomFace
};

立方体贴图对象有一个构成其面的纹理数组，根据文档，您使用其第三个参数的函数要求您对该纹理数组的索引，整个函数将返回：

在*pArray 中返回一个数组，通过该数组可以访问对应于数组索引arrayIndex 和mipmap 级别mipLevel 的映射图形资源资源的子资源。每次映射资源时，*pArray 中设置的值可能会发生变化。

阅读更多：http ://docs.nvidia.com/cuda/cuda-driver-api/index.html#ixzz4icKF1c00 关注我们：@GPUComputing 在 Twitter | 英伟达在 Facebook

我希望这有助于回答您关于将第三个参数用于您尝试从他们的 API 使用的函数的问题。

编辑

OP 曾询问，当将此枚举传递CU_CUBEMAP_FACE_POSITIVE_X给上述函数调用的第三个参数时，它是否只返回立方体贴图的那个面，它恰好是一个纹理。在查看有关此处找到的枚举值或类型的文档时：enum CUarray_cubemap_face

enum CUarray_cubemap_face - 立方体面的数组索引

价值观

  • CU_CUBEMAP_FACE_POSITIVE_X = 0x00
    • 立方体贴图的正 X 面
  • CU_CUBEMAP_FACE_NEGATIVE_X = 0x01
    • 立方体贴图的负 X 面
  • CU_CUBEMAP_FACE_POSITIVE_Y = 0x02
    • 立方体贴图的正 Y 面
  • CU_CUBEMAP_FACE_NEGATIVE_Y = 0x03
    • 立方体贴图的负 Y 面
  • CU_CUBEMAP_FACE_POSITIVE_Z = 0x04
    • 立方体贴图的正 Z 面
  • CU_CUBEMAP_FACE_NEGATIVE_Z = 0x05
    • 立方体贴图的负 Z 面

阅读更多内容：http ://docs.nvidia.com/cuda/cuda-driver-api/index.html#ixzz4idOT67US 在 Twitter 上关注我们：@GPUComputing | 英伟达在 Facebook

在我看来，当使用这种方法查询或获取存储到立方体贴图数组中的纹理信息时，第三个参数的要求就是这个枚举值；无非就是0-index进入那个数组。因此，CU_CUBEMAP_FACE_POSITIVE_X作为第三个参数传递给我并不一定意味着您只会取回该特定面部的纹理。在我看来，因为这是0th index它将返回整个纹理数组。C像指针一样传递数组的旧样式。

于 2017-05-31T01:55:59.720 回答