CUDA 全局内存拷贝

Question

CUDA C 编程指南 (p.70) 说，

全局内存驻留在设备内存中，设备内存被访问 通过 32、64 或 128 字节的内存事务。这些记忆 事务必须自然对齐：仅 32、64 或 128 字节 与其大小对齐的设备内存段（即其 第一个地址是它们大小的倍数）可以读取或写入 内存事务。

所以，如果我想在 device 函数中一次访问 32、64 或 128 个连续字节，（例如复制到共享内存）什么是最合适的函数（或分配）对于这个操作？

传统 C 的 memcpy 函数似乎不能一次访问 32 个字节（非常慢）。而且因为这不是向量数据，所以我想要一个线程一次读取这些数据。

---

致 dbaupp

memcpy 运行良好，但我说的是速度。 例如，假设我有设备内存指针 p 并在设备函数中运行以下代码。

a) 字符 c[8]; memcpy(c, p, 8);

b) 字符 c[8]; * (双 * )c = * (双 * )p;

对于以上两种情况，结果是相同的，但是情况 b 比情况 a 快了近 8 倍（我在我的代码中测试并确认了）。

仅供参考，cudaMemcpy 功能在设备功能中不起作用。

所以，我想知道是否有任何方法可以从单个操作中复制 16 个字节。 （希望比 memcpy(c, p, 16); 快 16 倍；）

Answer 1

目前还不是 100% 清楚您要做什么。如果您尝试将数据从全局复制到共享内存，那么它可能具有某种结构，例如chars 或 floats 或其他东西的数组。以下答案将假设您正在处理一个 chars 数组（您可以将 char 替换为任何数据类型）。

总结：不要考虑一次显式访问 32/64/128 字节，只需编写代码以便合并内存访问即可。

---

您可以使用 CUDA 随心所欲地访问数据，就像在普通的 C/C++/其他语言中一样。您甚至可以深入到单个字节。编程指南所说的是，每当访问数据时，都必须读取 32/64/128 字节的块。例如。如果您有char a[128] 并且想要获得a[17]，那么GPU 必须从a[0] 读取到a[31] 才能获得a[17] 中的数据。这是透明地发生的，因为您无需编写任何不同的代码即可访问各个字节。

主要考虑的是内存访问速度：如果必须为每个信息字节读取 31 个垃圾字节，那么您的有效内存带宽将减少 32 倍（这也意味着您必须进行更多的全局内存访问，这是 sloowww)！

但是，GPU 上的内存访问可以在一个块中跨线程“合并”（this question 为优化合并提供了一个合理的起点。）。简而言之，合并允许一个块中多个线程同时发生的内存访问可以“批处理”在一起，这样只需要发生一次读取。

这里的重点是合并发生在块内的线程之间（而不是在单个线程内），因此对于复制到共享内存中可以做到（array 是 chars 的数组全局内存）：

__shared__ char shrd[SIZE];

shrd[threadIdx.x] = array[blockDim.x * blockIdx.x  + threadIdx.x];
__syncthreads();

这将使每个线程将一个字节复制到共享数组中。这种 memcpy 操作本质上是并行发生的，并且数据访问是合并的，因此不会浪费带宽（或时间）。

上述策略比让单个线程逐个字节地迭代和复制要好得多。

还可以将数组的每个 n 字节块视为单个 n 字节数据类型，并让每个线程复制它。例如对于 n==16，对uint4进行一些强制转换

__shared__ char shrd[SIZE];

((uint4*)shrd)[threadIdx.x] = ((uint4*)array)[blockDim.x * blockIdx.x  + threadIdx.x];
__syncthreads();

这将允许每个线程一次复制 16 个字节。关于那段代码的注释：

• 我尚未对其进行测试或基准测试
• 我不知道这是否是一种好的做法（我强烈希望它不是）。）
• 索引按 16 倍缩放（例如，threadIdx.x == 1 对应于写入 shrd[16],shrd[17],...,shrd[31]）

---

附带说明：根据您的具体用例，built-in cudaMemcpy functions 可能有用。