CUDA总结:合并访问coalesced

传输延迟(latency)

在host端和device端之间存在latency,数据通过PCI-E总线从CPU传输给GPU,我们必须避免
频繁的host、device间数据传输,即使是最新的PCIE 3.0 x16接口,其双向带宽也只有32GB/s

在device内部也存在latency,即数据从gpu的存储器到multi-processor(SM)的传输。
设备内存带宽
访问一次全局内存,将耗费400~600个cycle,成本是非常高的,所以必须谨慎对待全局内存的访问

合并(coalesced)

数据从全局内存到SM(stream-multiprocessor)的传输,会进行cache,如果cache命中了,下一次的访问的耗时将大大减少。
每个SM都具有单独的L1 cache,所有的SM共用一个L2 cache。
在计算能力2.x之前的设备,全局内存的访问会在L1\L2 cache上缓存;在计算能力3.x以上的设备,全局内存的访问只在L2 cache上缓存。
对于L1 cache,每次按照128字节进行缓存;对于L2 cache,每次按照32字节进行缓存。
参考:《CUDA_C_Programming_Guide-V8.0》 Appendix G. COMPUTE CAPABILITIES

合并访问是指所有线程访问连续的对齐的内存块,对于L1 cache,内存块大小支持32字节、64字节以及128字节,分别表示线程束中每个线程以一个字节(1*32=32)、16位(2*32=64)、32位(4*32=128)为单位读取数据。前提是,访问必须连续,并且访问的地址是以32字节对齐。(类似于SSE\AVX的向量指令,cuda中的合并访存也是向量指令)

例子,假设每个thread读取一个float变量,那么一个warp(32个thread)将会执行32*4=128字节的合并访存指令,通过一次访存操作完成所有thread的读取请求。
coalesced示意图

对于L2 cache,合并访存的字节减少为32字节,那么L2 cache相对L1 cache的好处?
在非对齐访问、分散访问(非连续访问)的情况下,提高吞吐量(cache的带宽利用率)

非对齐访问(unaligned)

L1 cache的非对齐访问

L2 cache的非对齐访问

以上是L1、L2 cache的非对齐访问的对比,128字节的数据没有进行内存对齐,首地址位于96~128之间,
L1为了访问128之前的数据,必须访问从0~127的整段内存,其cache的有效利用率是128/256=50%,L2则只需要访问96~127的内存,其cache的有效利用率是128/160=80%

分散访问(scattered)
warp请求访问位于不同地址的数据,数据是非连续的,此时warp无法进行合并访问,每个thread访问一个float,一共需要执行32次访存指令。下面观察L1 和 L2 的区别

L1 cache的分散访问
L1 cache,访存请求分布在0~383的内存之间,cache的有效利用率是128/384=33%

L2 cache的分散访问
L2 cache,相比L1在scatterd情况下要好得多,cache的有效利用率达到128/192=67%

关于L2 cache的读写操作
L1 has a cache line size that is fixed at 128 bytes and cannot be changed.
L2 has a cache line size that is fixed at 32 bytes and cannot be changed.
Note that the L1 may be disabled by default on some GPUs, and can be disabled in software.
If the L1 is enabled, a cache line miss will force a load of that cache line, i.e. a 128byte load. This will necessarily result in 4 L2 transactions (4x32=128).
If the L1 is disabled, transactions may attempt to hit in the L2. If they miss in the L2, a DRAM transaction will be generated. The size of this transaction would be 32 bytes.
If the L1 is enabled, and a transaction misses in the L1, it will generate 4 L2 transactions. If all 4 of those L2 transactions also miss, then 4 DRAM read transactions (each of 32 bytes) will be generated. In effect, in this scenario, 128 bytes will be read from DRAM as a result of the miss in L1 and L2.

  • 8
    点赞
  • 26
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
`cuda::memcpy_async` 和 `cooperative_groups::memcpy_async` 都是异步的内存拷贝函数,但是它们的使用场景和特点略有不同。 `cuda::memcpy_async` 是 CUDA Runtime API 中提供的异步内存拷贝函数,它可以在主机和设备之间进行数据拷贝,并且可以在拷贝过程中执行其他的 CUDA 操作。这个函数的使用方法和 `cudaMemcpy` 类似,但是需要额外传入一个 `cudaStream_t` 参数来指定使用的 CUDA 流。与 `cudaMemcpy` 不同的是,`cuda::memcpy_async` 不会阻塞主机线程,而是立即返回并在后台执行数据拷贝。因此,它可以提高程序的并发性能,特别是在数据量较大的情况下。 `cooperative_groups::memcpy_async` 是 CUDA cooperative groups 库中提供的异步内存拷贝函数,它是在协作线程组(cooperative thread groups)中进行的数据拷贝。与 `cuda::memcpy_async` 不同的是,这个函数只能在协作线程组中使用,而且需要传入一个 `cooperative_groups::coalesced_group` 参数来指定线程组。由于协作线程组中的线程可以协同工作,因此这种方式可以进一步提高内存拷贝的效率。 综上所述,`cuda::memcpy_async` 和 `cooperative_groups::memcpy_async` 都是异步内存拷贝函数,但是使用场景和特点略有不同。`cuda::memcpy_async` 可以在主机和设备之间进行数据拷贝,并且可以在拷贝过程中执行其他的 CUDA 操作;而 `cooperative_groups::memcpy_async` 则是在协作线程组中进行的数据拷贝,可以进一步提高内存拷贝的效率。

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值