通过表面写入 CUDA 中的浮点 OpenGL 纹理

Posted 2023-03-23

【中文标题】通过表面写入 CUDA 中的浮点 OpenGL 纹理

【英文标题】：Writing to a floating point OpenGL texture in CUDA via a surface

【发布时间】：2015-02-28 18:27:17

【问题描述】：

我正在编写一个 OpenGL/CUDA (6.5) 互操作应用程序。尝试通过我的 CUDA 内核中的表面引用将浮点值写入 OpenGL 纹理时出现编译时错误。

在这里，我对如何设置互操作进行了高级描述，但是我成功地从我的 CUDA 内核中的纹理读取，所以我相信这是正确的。我有一个用

声明的 OpenGL 纹理

glTexImage2D(GL_TEXTURE_RECTANGLE_ARB, 0, GL_RGB32F_ARB, 512, 512, 0, GL_RGB, GL_FLOAT, NULL);

创建纹理后，我调用cudaGraphicsGLRegisterImage 并设置cudaGraphicsRegisterFlagsSurfaceLoadStore。在运行我的 CUDA 内核之前，我取消绑定纹理并在从cudaGraphicsGLRegisterImage 获得的cudaGraphicsResource 指针上调用cudaGraphicsMapResources。然后我从cudaGraphicsSubResourceGetMappedArray 得到一个cudaArray，为该数组创建一个适当的资源描述符，并调用cudaCreateSurfaceObject 来获取一个指向cudaSurfaceObject_t 的指针。然后我用cudaMemcpyHostToDevice 调用cudaMemcpy 将cudaSurfaceObject_t 复制到cudaMalloc 分配的设备上的缓冲区。

在我的 CUDA 内核中，我可以使用类似这样的内容从表面参考中读取数据，并且我已经验证这可以按预期工作。

__global__ void cudaKernel(cudaSurfaceObject_t tex) 
    int x = blockIdx.x*blockDim.x + threadIdx.x;
    int y = blockIdx.y*blockDim.y + threadIdx.y;
    float4 sample = surf2Dread<float4>(tex, (int)sizeof(float4)*x, y, cudaBoundaryModeClamp);

在内核中，我想修改样本并将其写回纹理。 GPU 具有 5.0 的计算能力，所以这应该是可能的。我正在尝试这个

surf2Dwrite<float4>(sample, tex, (int)sizeof(float4)*x, y, cudaBoundaryModeClamp);

但我得到了错误：

error: no instance of overloaded function "surf2Dwrite" matches the argument list
argument types are: (float4, cudaSurfaceObject_t, int, int, cudaSurfaceBoundaryMode)

我可以看到

cuda-6.5/include/surface_functions.h

只有surf2Dwrite 的完整版本的原型接受void * 作为第二个参数。我确实看到了surf2Dwrite 的原型，它接受带有模板化surface 对象的float4，但是，我不确定如何使用OpenGL 互操作声明模板化surface 对象。我无法找到有关如何执行此操作的任何其他内容。任何帮助表示赞赏。谢谢。

【问题讨论】：

这不是错字，它是这样初始化的。也许您认为这是内部格式。内部格式为GL_RGB32F_ARB。 （请参阅glTexImage2D 手册页。）GL_RGB 的格式要求 opengl 将第一个浮点数视为 R、第二个 G 和第三个 B。据我了解，从 OpenGL 渲染到纹理时忽略第四个浮点数.

文档说 GL_RGB32F 是不允许的。 docs.nvidia.com/cuda/cuda-runtime-api/…

有趣。好吧，实际上它似乎只是忽略了 alpha 通道。

也许他们在 6.5 中修复了它，但在 5.5 中使用 RGB 会产生 Seg-Fault。 ;)

【参考方案1】：

事实证明，答案很简单，虽然我不知道为什么会这样。而不是调用

surf2Dwrite<float4>(sample, tex, (int)sizeof(float4)*x, y, cudaBoundaryModeClamp);

我需要打电话

surf2Dwrite(sample, tex, (int)sizeof(float4)*x, y, cudaBoundaryModeClamp);

说实话，我不确定我是否完全理解 CUDA 在 c++ 中对模板的使用。有人解释一下吗？

【参考方案2】：

有关 CUDA 写入链接到 OpenGL 纹理的表面的完整示例，请参阅此项目：

https://github.com/nvpro-samples/gl_cuda_interop_pingpong_st

【讨论】：

对于这个问题，这不是一个非常有用的答案，因为 1）您的示例使用 CUDA 驱动程序 API，而我询问的是运行时 API，2）您的示例使用较旧的表面引用而不是较新的表面我的问题中的对象。并不是只有一种公认的方法可以做到这一点。由于结论只是“不要将模板符号与surf2Dwrite 一起使用”，因此要求人们挖掘整个 git 项目以得出与五个月前在接受的答案中发布的相同结论也不是很有帮助...

【参考方案3】：

来自CUDA Documentation，这里是表面模板函数的定义：

template<class T>
T surf2Dread(cudaSurfaceObject_t surfObj,
              int x, int y,
              boundaryMode = cudaBoundaryModeTrap);

template<class T>
void surf2Dread(T* data,
                 cudaSurfaceObject_t surfObj,
                 int x, int y,
                 boundaryMode = cudaBoundaryModeTrap);