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

Question

我的内核线程以合并方式访问线性字符数组。如果我映射纹理数组我没有看到任何加速。运行时间是差不多一样。我正在研究具有计算能力 2.0 的 Tesla C2050 并阅读全局访问被缓存的地方。真的吗？也许这就是为什么我我没有看到运行时间的差异。

主程序中的数组为

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

然后我将它绑定(bind)到纹理 texture<char> texdatabase与

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

然后每个线程读取一个字符 ch = tex1Dfetch(texdatabase, p + id)在哪里是threadIdx.x + blockIdx.x * blockDim.x和 p是一个偏移量。

我只绑定(bind)一次 dev_database是一个大数组。其实我发现如果大小太大，则绑定(bind)失败。数组的大小是否有限制绑定(bind)？非常感谢。

Answer 1

对于为什么您看不到任何性能差异，有多种可能性，但最有可能的是这种内存访问不是您的瓶颈。如果它不是你的瓶颈，让它更快对性能没有影响。

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

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

有a good presentation by Paulius Micikevicius from GTC 2010你可能想看。它涵盖了分析驱动的优化和正在运行的内存事务的具体概念。