CUDA 立方体贴图纹理

Posted 2023-03-16

【中文标题】CUDA 立方体贴图纹理

【英文标题】：CUDA cube map textures

【发布时间】：2014-05-06 00:53:20

【问题描述】：

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

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

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

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

【问题讨论】：

你看过cube map texture sample code吗？

@RobertCrovella 谢谢你的建议，但我担心这个演示没有提到在 CUDA 内核中使用 OpenGL 立方体贴图。该演示使用直接创建和填充的 CUDA 数组，而不是从 OpenGL GL_TEXTURE_CUBE_MAP 纹理对象获取。

是的，我同意它没有展示 OGL 互操作，这是您问题的主要部分。我认为这可能有助于您的问题中似乎询问如何在内核中访问/使用立方体映射纹理的部分。该示例演示了可以在内核中访问所有 6 个面。

看看这两个链接：nvidia.com/object/cube_map_ogl_tutorial.html & docs.nvidia.com/cuda/index.html#axzz4ibxwC8x2 这些应该可以作为很好的参考。

【参考方案1】：

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

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

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

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

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

下面我放了一个函数片段，它可以完成所有这些，完整的长度在CIRT Repository（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 纹理。

【参考方案2】：

我并不真正熟悉 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 不是texture，则无法通过array 访问它，并返回CUDA_ERROR_NOT_MAPPED_AS_ARRAY。如果arrayIndex 不是resource 的有效array index，则返回CUDA_ERROR_INVALID_VALUE。如果mipLevel 不是resource 的有效mipmap level，则返回CUDA_ERROR_INVALID_VALUE。如果资源不是mapped，则返回CUDA_ERROR_NOT_MAPPED。

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

另见：

cuGraphicsResourceGetMappedPointer

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

---

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

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 内核旨在将每个线程的延迟降至最低，而 GPU 旨在处理大量并发的轻量级线程，以最大限度地提高吞吐量。 RAM - 主机系统和设备都有各自不同的附加物理内存。由于主机和设备内存由 PCI Express (PCIe) 总线分开，因此主机内存中的项目有时必须通过总线与设备内存进行通信，反之亦然，如在启用 CUDA 的设备上运行什么？李>

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

阅读更多：http://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#ixzz4ic8ch2fq 在 Twitter 上关注我们：@GPUComputing | Facebook 上的 NVIDIA

---

现在知道 CUDA API 库有两种不同的 API，我们必须了解两者之间的区别：Difference Between the driver and runtime APIs

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

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

复杂性与控制

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

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

上下文管理

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

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

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

因为这恰好出现在 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 在 Twitter 上关注我们：@GPUComputing | Facebook 上的 NVIDIA

恰好是一个unsigned int 或构成cube map 的纹理的索引位置.因此，如果我们查看立方体贴图以及纹理及其与伪代码的关系，我们可以看到：

// 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 在 Twitter 上关注我们：@GPUComputing | Facebook 上的 NVIDIA

---

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

---

编辑

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

枚举 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 上的 NVIDIA

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

【讨论】：

我看不到其中的任何实际上是如何回答问题的

@talonmies OP 询问他们是否可以执行任务A，前提是他们基于该功能的实现正确理解了所述功能的工作原理。我认为他们的假设是错误的，因此他们需要清楚地了解这个函数作为其参数列表的内容、这些参数是什么以及它们的用途、实际函数的作用以及它的返回值是什么以及任何返回的错误消息。在我看来，他们理解第三个参数是选择单个纹理......

@talonmies ...续 事实并非如此；第三个参数正在寻找或期待一个无符号值，该值恰好是该映射资源（立方体贴图）对象的纹理数组的索引。一旦OP理解了这一点；然后他们可以根据这个功能的性质重新设计他们的实现。我的回答也来自 CUDA 的文档。

@talonmies 我对我的原始答案进行了编辑，提供了有关第三个参数的更多信息，以及它是作为枚举类型或无符号值进入立方体贴图纹理数组的第一个索引值这一事实.

@FrancisCugler 因此，如果我将CU_CUBEMAP_FACE_NEGATIVE_Z 而不是CU_CUBEMAP_FACE_POSITIVE_X 作为arrayIndex 参数，则返回的数组将只包含最后一张脸（并且分别在传递时说CU_CUBEMAP_FACE_NEGATIVE_Y 将只包含从CU_CUBEMAP_FACE_NEGATIVE_Y 到CU_CUBEMAP_FACE_NEGATIVE_Z 的面孔，最后三个是）？