CUDA全局内存拷贝

Question

CUDA C 编程指南(第 70 页)说，

Global memory resides in device memory and device memory is accessed via 32-, 64-, or 128-byte memory transactions. These memory transactions must be naturally aligned: Only the 32-, 64-, or 128-byte segments of device memory that are aligned to their size (i.e. whose first address is a multiple of their size) can be read or written by memory transactions.

所以，如果我想在 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% 清楚您要尝试做什么。如果您尝试将数据从全局复制到共享内存，那么它可能具有某种结构，例如char 或 float 之类的数组。以下答案将假定您正在处理一个 char 数组(您可以将 char 替换为任何数据类型)。

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

---

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

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

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

重点是合并发生在一个 block 内的线程之间(而不是在单个线程内)，因此对于复制到共享内存，我们可以这样做(array 是一个数组全局内存中的 char:

__shared__ char shrd[SIZE];

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

这将使每个线程将一个字节复制到共享数组中。这个 memcpy 操作基本上是并行发生的，数据访问是合并的，所以没有浪费带宽(或时间)。

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

也可以将数组的每个 n 字节 block 视为单个 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 对应于写入 shr​​d[16],shrd[17],...,shrd[31])

---

作为旁注:根据您的具体用例，built-in cudaMemcpy functions可能会有用。