我有一个每个块的数据数组。我在 cuda Grid 中有 N 个块和一个大小为 N 的常量数据数组“block_data []”。
因此,给定块“X”中的所有线程只访问一次 block_data[X],并使用该值执行某些操作。
我的问题是:这种广播方案是否有效?如果没有,我应该采取什么方法?
评论后编辑:我对常量内存的唯一问题是它的大小有限,因为我可以拥有超过 64K 的块。这意味着超过 64KB
问候
如果您只使用普通的全局内存访问,那么事务的效率就相当低,尽管取决于您的内核所做的工作量,影响可能非常小。
我假设sizeof(block_data)
是一个字节(从你的问题推断“......可能有超过 64K 的块。这意味着超过 64KB”)。
sizeof(block_data)
),如果块中的其他扭曲请求相同的数据,那么它们应该从 L1 获取。负载的效率是 1/128,但你应该只为块支付一次。另一种方法是将数据标记为const __restrict__
向编译器指示数据是 a) 只读的并且 b) 没有被任何其他指针别名。由于编译器可以检测到访问是统一的,因此它可以优化访问以使用只读缓存之一(例如,常量缓存或计算能力 >=3.5 的只读数据缓存,即纹理缓存)。
如果要更改block_data[N]
数组中的值,最好使用共享内存的概念__shared__
。如果您不更改 的值block_data[N]
,请使用__const__
或使用缓存的概念。通过使用 L2 Cache,可以获得 1536KB 的内存(Kepler)。