CUDA全局内存副本

Question

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

全局存储器驻留在设备存储器中,并且通过32,64或128字节存储器事务访问设备存储器.这些内存事务必须自然对齐:只有与其大小对齐的设备内存的32,64或128字节段(即其第一个地址是其大小的倍数)可以由内存事务读取或写入.

因此,如果我想在设备功能中一次访问32,64或128个连续字节,(例如,复制到共享内存)这个操作最合适的功能(或赋值)是什么？

传统的C memcpy函数似乎一次不能访问32个字节(非常慢).因为这不是矢量数据,所以我希望单个线程一次读取这些数据.

---

到dbaupp

memcpy效果很好,但我说的是速度.例如,假设我有设备内存指针p并在设备功能中运行以下代码.

a)char c [8]; memcpy(c,p,8);

b)char c [8];*(double*)c =*(double*)p;

对于上述两种情况,结果相同,但情况b比情况a(我在我的代码中测试和确认)快近8倍.

而且,仅供参考,cudaMemcpy功能在设备功能中不起作用.

所以,我想知道的是,是否有任何方法可以在单个操作中复制16个字节.(希望比memcpy快16倍(c,p,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上的内存访问可以跨块中的线程"合并"(这个问题为优化合并提供了合理的起点).简而言之,合并允许对块中的多个线程同时发生的存储器访问可以"批处理"在一起,从而只需要进行一次读取.

这一点的重点是在块内的线程(不在单个线程中)之间发生合并,因此对于可以执行的复制到共享内存(array是char全局内存中的s 数组):

__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函数可能很有用.