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

Question

我的内核线程以合并的方式访问线性字符数组。如果我映射 数组到纹理我没有看到任何加速。运行时间是 差不多一样。我正在研究具有计算能力 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 是一个大数组。其实我发现 如果大小太大，则绑定失败。数组的大小是否有限制 绑定？非常感谢。

Answer 1

您没有看到任何性能差异的原因有多种可能性，但最有可能的是这种内存访问不是您的瓶颈。如果它不是您的瓶颈，那么加快它对性能没有影响。

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

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

您可能想观看a good presentation by Paulius Micikevicius from GTC 2010。它涵盖了分析驱动的优化和动态内存事务的具体概念。