首页 文章

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

提问于
浏览
2

我有问题绑定到纹理内存全局设备内存的子部分 .

我有一个充满内存的大型全局设备数组,如下所示:

double * 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,则上述方法确实有效,某种方式指针算法不起作用 .

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

3 回答

  • 2

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

    一种解决方案是使用Pitchhed Memory而不是具有非常小的纹理,而不是具有几个大的纹理,每个大的纹理从矩阵的对齐行开始 .

    我猜这里,但我认为使用

    sizeof(double)*10
    

    作为纹理存储器的数据集,设置存储器本身比读取它需要更多 .

    总矩阵有多大?

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

相关问题