2011-03-30 74 views
2

我有问题绑定到纹理内存全局设备内存的一个子部分。CUDA纹理内存绑定全局内存的子部分

我有充满存储器的大全球装置阵列如下:

双* device_global; ((** **)& device_global,sizeof(double)* N));

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

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

每个内核所需其我绑定到纹理通过device_global一小部分(INT偏移量= 100):

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

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

我想这样做:

cudaBindTexture(0,将texRef,device_global + offsett * 1,channelDesc,的sizeof(双)* 10);

应该指出的是,如果偏移量设置为0,上述方法可以工作,但某种程度上指针算术不起作用。

任何帮助或其他指导方针将不胜感激。

+0

IIRC指针算法是确定的,即使是设备指针。您是否在循环结束时解除了纹理?什么是错误? – LumpN 2011-03-30 17:07:52

回答

2

纹理存储器的偏移必须对齐。您无法将内存中的任何部分绑定到只有正确对齐的部分,这是因为内部高性能硬件的工作原理。

一个解决方案是使用倾斜内存,而不是有非常小的纹理 有几个大的开始在矩阵对齐的行。

我猜在这里,但我觉得用

sizeof(double)*10 

作为纹理显存是命令datasize,需要更多的设置记忆本身,而不是阅读。

总矩阵有多大?

+0

给出了范围的指示,假设每个内核的纹理内存大约是2000倍,程序需要启动100次内核(因此'device_global'的总大小为100 * 2000)。因此我希望在每次启动内核时都绑定并取消绑定。这也回答了@talonmies的评论。任何其他见解或提示非常赞赏。 – user683994 2011-03-30 18:25:57

2

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

将整个数组绑定到纹理,然后将索引偏移量传递到要用于纹理获取的每个内核中可能会更好。

2

将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);