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

Question

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

我有一个大的全局设备数组，里面装满了内存，如下所示:

double * device_global;

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

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

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

每个内核都需要 device_global 的一小部分(int offset = 100)，我通过以下方式将其绑定(bind)到纹理:

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

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

我想做这样的事情:

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

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

如有任何帮助或其他指南，我们将不胜感激。

Answer 1

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

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

绑定(bind)应按以下方式完成:

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