CUDA：纹理内存的访问时间是否类似于合并的全局内存？

Question

3

我的内核线程以合并的方式访问一个线性字符数组。如果我将该数组映射到纹理中，我并没有看到任何加速。运行时间几乎相同。我正在使用计算能力为2.0的Tesla C2050，读到某个地方说全局访问具有缓存。这是真的吗？也许这就是我没有看到运行时间差异的原因。

主程序中的数组为

char *dev_database = NULL;
cudaMalloc( (void**) &dev_database, JOBS * FRAGMENTSIZE * sizeof(char) );

我将其绑定到纹理texture<char> texdatabase上。

cudaBindTexture(NULL, texdatabase, dev_database, JOBS * FRAGMENTSIZE * sizeof(char) );

每个线程都会读取一个字符ch = tex1Dfetch(texdatabase, p + id)，其中id是threadIdx.x + blockIdx.x * blockDim.x，而p是一个偏移量。

我只绑定一次，而dev_database是一个很大的数组。事实上，我发现如果大小太大，绑定会失败。是否有绑定数组大小的限制？非常感谢。

- Ross

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- harrism · Accepted Answer

有几种可能导致您在性能方面看不到任何差异，但最有可能的是这个内存访问不是您的瓶颈。如果它不是您的瓶颈，使它更快将对性能没有影响。

关于缓存：对于这种情况，由于您只读取字节，因此每个warp将读取32字节，这意味着每组4个warps将映射到每个缓存行。因此，假设很少有缓存冲突，您将从缓存中获得高达4倍的重复使用率。因此，如果这个内存访问是瓶颈，纹理缓存可能不比通用缓存更有用。

您应该首先确定是否受带宽限制，以及这个数据访问是否是罪魁祸首。一旦您完成了这项工作，然后优化您的内存访问。另一个要考虑的策略是每次加载访问4到16个字符（使用char4或int4结构与字节打包/解包），而不是每个线程访问一个字符，以增加同时进行的内存事务数 - 这可以帮助饱和全局内存总线。

有来自GTC 2010的Paulius Micikevicius的一份好演示文稿，您可能想观看。它涵盖了基于分析的优化和内存事务在飞行中的特定概念。