CUDA纹理内存绑定全局内存的子部分

Posted 2023-03-23

【中文标题】CUDA纹理内存绑定全局内存的子部分

【英文标题】：CUDA texture memory to bind a sub-portion of global memory

【发布时间】：2011-07-26 03:13:40

【问题描述】：

我在将纹理内存绑定到全局设备内存的子部分时遇到问题。

我有一个大型的全局设备数组，其中填充了如下内存：

双 * device_global;

cudaMalloc((void **)&device_global, sizeof(double)*N));

cudaMemcpy(device_global, host, sizeof(double)*N, cudaMemcpyHostToDevice) );

我在一个 for 循环中运行多个内核。

每个内核需要device_global 的一小部分（int offset = 100），我通过以下方式将其绑定到纹理：

cudaBindTexture(0, texRef, device_global, channelDesc, sizeof(double)*10);

但是我面临的问题是我无法使用指针算法仅通过循环偏移绑定device_global 的循环部分。

我想做这样的事情：

cudaBindTexture(0, texRef, device_global+ offsett * i , channelDesc, sizeof(double)*10);

需要注意的是，如果将偏移量设置为0，上述方法确实有效，但不知何故，指针算法不起作用。

任何帮助或其他指南将不胜感激。

【问题讨论】：

IIRC 指针算法是可以的，即使对于设备指针也是如此。您是否在循环结束时取消绑定纹理？什么错误？

【参考方案1】：

将0 或NULL 作为cudaBindTexture 的第一个参数传递是一种不好的做法。 CUDA 纹理绑定要求要绑定的指针必须对齐。对齐要求可以通过cudaDeviceProp::textureAlignment设备属性来确定。

cudaBindTexture 可以将任何设备指针绑定到纹理。如果指针未对齐，它会返回距cudaBindTexture 的第一个参数中最近的对齐地址的字节偏移量。如果第一个参数是NULL，则函数调用失败。

绑定应该是这样的：

size_t texture_offset = 0;
cudaBindTexture(&texture_offset, texRef, device_global+ offsett * i , channelDesc, sizeof(double)*10);

【参考方案2】：

纹理内存的偏移量必须对齐。您不能仅将内存的任何部分绑定到正确对齐的部分，这是因为内部高性能硬件的工作方式。

一种解决方案是使用 Pitched Memory 而不是使用非常小的纹理 有几个大的，每个都从矩阵的对齐行开始。

我在这里猜测，但我认为使用

sizeof(double)*10

作为纹理内存的数据大小，设置内存本身比读取它需要更多的时间。

总矩阵有多大？

【讨论】：

为了说明范围，我们假设每个内核的纹理内存约为 2000 双倍，并且程序需要 100 次内核启动（因此device_global 的总大小为 100 * 2000）。因此，我希望在每次内核启动时绑定和取消绑定。这也回答了@talonmies 的评论。非常感谢任何其他见解或提示。

【参考方案3】：

我不相信有可能做你想做的事。我怀疑有一些幕后地址转换，这意味着如果您传递给绑定调用的指针对于运行时内存管理器来说是未知的并且适当地对齐到页面边界，它将不允许绑定纹理到地址。

最好将整个数组绑定到纹理，然后将索引偏移量传递给每个内核以用于纹理提取。