CUDA立方体贴图纹理

15

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

当想要在CUDA kernel中使用OpenGL纹理时,需要做的事情之一是从已注册的图像和映射资源中检索一个CUDA数组,在这种情况下是一个纹理。在驱动程序API中,可以通过cuGraphicsSubResourceGetMappedArray调用来完成,对于2D纹理不是问题。但是当涉及前述的立方体贴图时,这个函数的第三个参数需要一个面的枚举(例如CU_CUBEMAP_FACE_POSITIVE_X)。因此,会有一些问题 - 当传递这样的枚举时,返回的纹理数组将仅包含该特定面的数据,对吗?那么如何作为一个整体使用立方体纹理来执行立方体映射,同样:

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

在CUDA内核中是否可以这样做,还是需要在用户代码中使用2D纹理进行适当的计算和采样?


2
你看过cube map texture sample code了吗? - Robert Crovella
@RobertCrovella 谢谢您的建议,但恐怕这个演示并没有涉及在CUDA内核中使用OpenGL立方体贴图。该演示使用直接创建和填充的CUDA数组,而不是从OpenGL GL_TEXTURE_CUBE_MAP纹理对象获取的。 - Matso
是的,我同意它并没有展示OGL互操作性,这是你问题的主要部分。但我认为它可能有助于回答你问题中似乎在询问如何在内核中访问/使用立方体贴图的部分。该示例演示了内核访问所有6个面的可能性。 - Robert Crovella
请查看以下两个链接:http://www.nvidia.com/object/cube_map_ogl_tutorial.html 和 http://docs.nvidia.com/cuda/index.html#axzz4ibxwC8x2。它们应该作为一个很好的参考。 - Francis Cugler
2个回答

4

好的 - 我自己解决了问题,但解决方案并不像使用另一个CUDA函数那么简单。

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

在任何时候,我们需要绑定引用的CUDA数组。立方体贴图纹理也是如此。但在这种情况下,数组必须是三维的。问题在于,CUDA驱动程序API仅提供上述功能以从此类纹理资源检索单个层,并将其映射到单个二维数组。为了得到我们想要的结果,我们必须自己制作包含所有层的3D数组(或立方体贴图的面)。

首先,我们必须使用上述函数为每个层/面获取数组。下一步是通过调用cuArray3DCreate创建3D数组,并提供正确的参数集(大小/层数、详细级别、数据格式、每个纹素的通道数和一些标志)。然后,我们必须使用一系列调用cuMemcpy3D将层的数组复制到3D数组中,每个层/面数组一个。

最后,我们使用cuTexRefSetArray设置目标CUDA纹理引用,其中包含我们创建并复制到的3D数组。在设备代码内部,我们使用适当的纹理类型和模式(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纹理也应该能够正常工作。

0

我不太熟悉CUDA,但我有OpenGL和DirectX的经验,并且熟悉3D图形渲染API、库和管线,有能力设置和使用这些API。


当我看到你的问题:

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

然后你继续解释:

When one want to use OpenGL textures in CUDA kernel one of the things to do is to retrieve a CUDA array from registered image and mapped resource, in this case a texture. In driver API it is done by cuGraphicsSubResourceGetMappedArray call, which in a case of 2D texture is not a problem. But when talking about aforementioned cube map, third parameter of this function requires a face enum (like CU_CUBEMAP_FACE_POSITIVE_X). Thus some questions arise - when one passes such an enum, then the returned texture array will contain only data of that particular face, right? Then how to use cube texture as a whole, to perform cube mapping, likewise:

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

Or is it impossible to do so in CUDA kernal and one need to use 2D textures with proper calculations and sampling in user code?


我访问了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中返回一个数组,通过它可以访问映射图形资源resource对应于数组索引arrayIndex和mipmap级别mipLevel的子资源。每次映射资源时,*pArray中设置的值可能会发生变化。

如果resource不是纹理,则无法通过数组访问它,并返回CUDA_ERROR_NOT_MAPPED_AS_ARRAY。如果arrayIndex不是resource的有效数组索引,则返回CUDA_ERROR_INVALID_VALUE。如果mipLevel不是resource的有效mipmap级别,则返回CUDA_ERROR_INVALID_VALUE。如果资源未被映射,则返回CUDA_ERROR_NOT_MAPPED

注意: 请注意,此函数还可以返回先前异步启动的错误代码。

另请参阅:

cuGraphicsResourceGetMappedPointer

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


这个函数方法是在NVidia CUDA的DriverAPI中发现的,而不是在他们的RuntimeAPI中。当理解具有CUDA能力的硬件时,可以在此处找到HostDevice可编程管道之间的区别:http://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#axzz4ic6tFjXR2. 异构计算 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 关注我们:@GPUComputing on Twitter | NVIDIA on Facebook

现在我们知道了CUDAs API库有两个不同的API,我们必须理解这两者之间的区别,可以在此处找到:驱动程序API和运行时API之间的区别

1. 驱动程序API和运行时API之间的区别 驱动程序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 关注我们:@GPUComputing on Twitter | NVIDIA on 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 | NVIDIA on Facebook

这个位置恰好是一个unsigned int或者说是作为一个cube map中所组成的纹理的索引。一个典型的cube map6个面,如果同时内部和外部都映射的话最多也只有12个面。因此,如果我们看一下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
};

cubemap对象具有由纹理组成的面的数组,根据文档,你所质疑的函数与其第三个参数正在要求您进入此纹理数组的索引,整个函数将返回以下内容:

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

阅读更多信息:http://docs.nvidia.com/cuda/cuda-driver-api/index.html#ixzz4icKF1c00 关注我们:@GPUComputing on Twitter | NVIDIA on Facebook


我希望这能解答你关于在API中使用第三个参数的函数的问题。


编辑

原帖中询问当将枚举值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 上的 NVIDIA

在使用此方法查询或获取存储在立方体贴图数组中的纹理信息时,第三个参数要求是枚举值,似乎只是该数组中的0索引。因此,将CU_CUBEMAP_FACE_POSITIVE_X作为第三个参数传递给我并不一定意味着您只会获得该特定面的纹理。看来,由于这是0索引,它将返回整个纹理数组。就像旧的C风格传递数组一样,好像它们是指针。


4
我看不出这些内容中有任何一个能够真正回答这个问题。 - talonmies
@talonmies OP 问是否可以执行任务A,前提是他们能够正确理解该函数的工作原理,基于该函数的实现。我认为他们的假设是错误的,因此他们需要清楚地了解该函数需要哪些参数列表,这些参数是什么,以及它们的目的,实际的功能是什么,以及返回值和任何返回的错误信息。对我来说,他们理解第3个参数是选择单个纹理... - Francis Cugler
@talonmies ...继续说,情况并非如此。第三个参数正在寻找或期望一个无符号值,它恰好是这个映射资源(立方体贴图)对象的纹理数组的索引。一旦OP理解了这一点,他们就可以根据这个函数的性质重新设计他们的实现。我的答案也来自CUDA的文档。 - Francis Cugler
@talonmies 我对我的原始答案进行了编辑,提供了有关第三个参数的更多信息,以及它作为枚举类型或无符号值的第一个索引值进入立方体贴图数组的纹理。 - Francis Cugler
@Matso(...续)然而,我不知道这个CUDA的API函数调用的确切内部工作或实现。我基于我所阅读的文档以及它所接受的值、它所做的事情以及它返回的内容来进行操作。我熟悉使用SDK和API调用的3D图形结构和管道。我在OpenGL(传统和现代)和DirectX 9-11中都有一些经验。我曾经从头开始构建过一个3D渲染引擎,因此我对这个我尚未使用或编程的CUDA API进行了一些研究,并且... - Francis Cugler
显示剩余3条评论

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接