c - 如何使用 GPGPU 高效地执行加载和位操作？

Question

我需要将长度为 128 到 256 字节的数组加载到 GPU 共享内存。我想在有效执行按位运算的同时最小化全局内存访问。让我们解释一下。

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

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

1. 如果我为每个线程加载 16 字节的数组，那么我只有一个(warp size)*(16bytes) 数据的内存访问或最好的方法是每个线程 4 个字节？
2. 如果我将每个线程 16 字节加载到共享内存中，那么我是否会将这 16 字节与相同的 16 字节变量进行“或”操作以将某个位设置为 1？
3. 如果我为每个线程加载 16 个字节，那么在 16 字节数据类型上或在 4 字节数据类型或更小的数据类型上执行 OR 操作会更快？

例如我想将第 3 位设置为 1

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

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

Answer 1

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

指定 64 位的存储体大小

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

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