CUDA 为结构数组动态分配常量或纹理内存答案

【问题标题】：CUDA Dynamically allocated constant or texture memory for array of structsCUDA 为结构数组动态分配常量或纹理内存
【发布时间】：2020-01-09 00:37:40
【问题描述】：

我需要为我的内核使用常量内存中的结构数组，其中数组的实际大小直到运行时才知道。正如Correct way to use __constant__ memory on CUDA? 中的回答，我意识到在编译期间分配了常量内存，因此需要将数组声明为：

__constant__ SKY_GRID_TYPE const_patch_grid_lat[5];

已经定义了尺寸。但是因为我需要的实际大小取决于运行时完成的其他计算，所以我似乎无法使用常量内存。

上面的答案建议改用纹理内存，它说“可以动态设置并被缓存”。但是，我在内存中需要的数据类型是结构数组，根据Structure in Texture memory on CUDA，纹理内存似乎只支持 CUDA 内置类型。

那么有没有办法解决这个问题？常量内存对于我的结构数组来说是完美的，但大小是动态确定的，所以它不起作用。纹理内存本来可以工作，但它只允许 CUDA 内置类型。还有什么我可以使用或一些聪明的方法来解决这个问题吗？

【问题讨论】：

标签： cuda nvidia

【解决方案1】：

恒定内存最大为 64 KB。据我所知，分配所有 64 KB 没有缺点。

只需为您的数组分配最大大小（64kbytes/您的结构大小）。使用任何你需要的东西。这还假设您将在每次访问时使 uniform access 跨越整个经线。

如果您需要超过 64 KB，那么这当然行不通，但它会质疑您问题的整个前提。

对于较大的常量区域和/或您没有统一访问权限的情况，我对 cc3.5 和更新的 GPU 的建议是使用 read-only cache mechanism（const __restrict__ 或 __ldg()）。

【讨论】：

我无法分配所有 64 KB，因为我还需要常量内存中的其他内容，但我喜欢分配比我需要的更多的东西，然后只使用我需要的东西听起来是个好主意，并且应该管用。常量内存和只读缓存全局内存的速度对比如何？
它们都是由 GPU DRAM 内存支持的片上 GPU 缓存。在性能方面，它们在缓存命中时将相似，并且受制于对恒定内存的统一访问的期望。否则，NVIDIA 通常不会公布缓存大小、行为、性能。