2

CUDA C 编程指南（第 70 页）说，

全局内存驻留在设备内存中，设备内存通过 32、64 或 128 字节内存事务访问。这些内存事务必须自然对齐：内存事务只能读取或写入与其大小对齐的设备内存的 32、64 或 128 字节段（即其首地址是其大小的倍数） .

那么，如果我想在设备功能中一次访问 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 倍）

4

1 回答 1

6

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

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

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

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

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

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

__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]）

附带说明：根据您的特定用例，内置的 cudaMemcpy 函数可能很有用。

于 2012-05-31T14:29:02.913 回答