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

CUDA Dynamically allocated constant or texture memory for array of structs

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

__constant__ SKY_GRID_TYPE const_patch_grid_lat[5];

其中大小已经定义。但是因为我需要的实际大小取决于运行时完成的其他计算,所以我似乎不能使用常量内存。

上面的答案建议改为使用它所说的纹理内存 "can be set dynamically and are cached." 但是,我在内存中需要的数据类型是一个结构数组,根据 Structure in Texture memory on CUDA,它看起来像纹理内存仅支持 CUDA 内置类型。

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

常量内存最大为 64 KB。据我所知,分配所有 64 kbytes 没有任何缺点。

只需为数组分配最大大小(64kbytes/结构的大小)。使用任何你需要的东西。这还假设您将 uniform access 跨越扭曲,对于每次访问。

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

对于较大的常量区域,and/or对于没有统一访问权限的情况,我对 cc3.5 和更新的 GPU 的建议是使用 read-only cache mechanismconst __restrict____ldg()).