开发者社区> night李> 正文
阿里云
为了无法计算的价值
打开APP
阿里云APP内打开

CUDA实践指南(十四)

简介:
+关注继续查看

简单的访问模式:
第一个也是最简单的合并案例可以通过任何支持CUDA的设备来实现:第k个线程访问缓存行中的第k个字。 并非所有线程都需要参与。
例如,如果warp访问的线程相邻4字节字(例如,相邻浮点值),单个128B L1高速缓存线并因此单个合并事务将服务该存储器访问。 图3显示了这种模式。
1

如果线的某些字未被任何线程请求(例如,如果多个线程已经访问了相同的字或者某些线程没有参与该访问),则高速缓存线中的所有数据都被获取。 此外,如果在该段内已经对warp线程进行了访问,则只有一个128字节的L1事务将由具有计算能力2.x的设备执行。
顺序但未对齐的访问模式:
如果顺序线程顺序但未与高速缓存线对齐,则会请求两个128字节的L1高速缓存,如图4所示。
2

对于非高速缓存事务(即绕过L1并仅使用L2高速缓存的事务),除32字节L2段的级别外,可以看到类似的效果。 在图5中,我们看到一个例子:使用与图4相同的访问模式,但现在禁用了L1缓存,因此现在需要五个32字节的L2段来满足请求。
3

通过CUDA Runtime API分配的内存(如通过cudaMalloc())保证至少对齐至少256个字节。 因此,选择明显的线程块大小,例如翘曲大小的倍数(即,当前GPU上的32),便于通过与高速缓存线对齐的弯曲进行存储器访问。 (例如,如果线程块大小不是warp大小的倍数,请考虑第二个,第三个和后续线程块访问的内存地址会发生什么情况。)
错位访问的影响:
使用简单的复制内核来探究未对齐的访问的后果是容易的和信息性的,例如复制内核中的一个,它说明未对齐的访问。
说明错位访问的复制内核:

__global__ void offsetCopy(float *odata, float* idata, int offset)
{
    int xid = blockIdx.x * blockDim.x + threadIdx.x + offset;
    odata[xid] = idata[xid];
}

在说明未对齐访问的副本内核中,将数据从输入数组idata复制到输出数组,这两者都存在于全局内存中。 内核在主机代码的循环中执行,该代码将参数偏移量从0改为32.(图4和图4分别对应于缓存和非缓存内存访问情况下的错位)。副本的有效带宽 在NVIDIA Tesla M2090上具有各种偏移量(计算能力2.0,开启了ECC,因为它是默认设置)如图6所示。
对于NVIDIA Tesla M2090,全局内存访问没有偏移或偏移量为32个字的倍数,导致单个L1高速缓存线事务或4个L2高速缓存段负载(用于非L1高速缓存负载)。 实现的带宽约为130GB / s。 否则,每个warp都会加载两个L1高速缓存行(高速缓存模式)或者四个到五个L2高速缓存段(非高速缓存模式),从而实现了无偏移量的约四分之一的内存吞吐量。
有趣的一点是,我们可能预期缓存情况比此示例的非缓存情况更差,因为缓存情况下的每个翘曲都会按需要获取两倍的字节数,而在非缓存情况下, 5/4根据需要获取的字节数是按每个warp获取的。 然而,在这个特殊的例子中,这种效果并不明显,因为相邻的经线会重复使用它们的邻居所获取的缓存线。 因此,虽然缓存负载的影响仍然很明显,但并不像我们预期的那么好。 如果相邻的经线没有表现出超高速缓存线的高度重用,情况会更加如此。

版权声明:本文内容由阿里云实名注册用户自发贡献,版权归原作者所有,阿里云开发者社区不拥有其著作权,亦不承担相应法律责任。具体规则请查看《阿里云开发者社区用户服务协议》和《阿里云开发者社区知识产权保护指引》。如果您发现本社区中有涉嫌抄袭的内容,填写侵权投诉表单进行举报,一经查实,本社区将立刻删除涉嫌侵权内容。

相关文章
+关注
night李
对机器视觉 图像处理有一定的兴趣 要做一条有梦想的咸鱼
文章
问答
文章排行榜
最热
最新
相关电子书
更多
从零到一:IOS平台TensorFlow入门及应用详解
立即下载
从零到一:IOS平台TensorFlow入门及应用详解(附源
立即下载
Serverless 开发速查手册
立即下载