如何使用GPGPU高效地执行加载和位运算？

Question

我需要加载一个 128 到 256 字节长度的数组到 GPU 共享内存。

我想在有效执行按位操作的同时最小化全局内存访问。

我有一个 256 字节的数组加载到全局内存，我想在内核代码的开头将所有 256 字节加载到共享内存。

然后，必须对每个字节执行按位运算，而此运算可以应用于两个 1 字节变量、两个 2 字节变量或两个 4 字节变量和更大的变量。

问题是：

1. 如果我为每个线程加载 16 个字节的数组，那么我只有一个内存访问（扭曲大小）*（16 字节）数据，或者最好的方法是每个线程 4 个字节？
2. 如果我将每个线程 16 个字节加载到共享内存中，那么我是否会将这 16 个字节与相同的 16 个字节变量进行 OR 运算以将某个位设置为 1？
3. 如果我为每个线程加载 16 字节，然后对 16 字节数据类型执行 OR 运算，使用 4 字节数据类型或更小的数据类型是否更快？

例如我想将第3位设置为

1

：

__shared__ (which data type?) temp = ((which data type?) *)array[i];
temp |= (a variable with third bit set to 1)

array

在全局内存中，我想通过最小的全局内存访问量将它加载到共享内存中。

Answer 1

将内存事务的大小与银行大小相匹配。大多数时候，bank 大小是 32 位。从 Kepler 开始，您可以通过调用

cudaDeviceSharedMemConfig()

指定 64 位的存储体大小

Acceleware 在这个主题上有很好的资源：

http://acceleware.com/blog/maximizing-shared-memory-bandwidth-nvidia-kepler-gpus