CUDA学习(三十二)

简介:

本地内存:
本地内存访问只对一些自动变量发生,如可变内存空间说明符中所述。 编译器可能在本地内存中放置的自动变量是:

  • 不能确定它们是以固定数量进行索引的数组,
  • 大型结构或阵列会消耗太多的寄存器空间,
  • 任何变量,如果内核使用更多的寄存器比可用(这也被称为寄存器溢出)。

检查PTX汇编代码(通过使用-ptx orkeep选项编译获得)将会告诉在第一个编译阶段变量是否被放置在本地内存中,因为它将使用.local助记符进行声明,并使用ld.local进行访问 和st.local助记符。 即使没有,后续的编译阶段也可能会做出决定,但是如果他们发现它为目标体系结构消耗太多寄存器空间:使用cuobjdump检查cubin对象会告诉是否是这种情况。 此外,使用--ptxas-options = -v选项进行编译时,编译器会报告每个内核的本地内存使用情况(lmem)。 请注意,一些数学函数具有可能访问本地内存的实现路径。
在某些计算能力的设备上,3.x本地内存访问总是以与全局内存访问相同的方式缓存在L1和L2中(请参见计算功能3.x)。
在计算能力5.x和6.x的设备上,本地内存访问总是以与全局内存访问相同的方式缓存在L2中(请参阅计算能力5.x和计算能力6.x)。
共享内存:
由于它是片上的,共享内存比本地或全局内存具有更高的带宽和更低的延迟。
为了实现高带宽,共享内存被分成大小相同的内存模块,称为banks,可以同时访问。 因此,由n个地址落入n个不同存储体的任何存储器读或写请求可以同时得到服务,产生的总带宽是单个模块带宽的n倍。
但是,如果一个内存请求的两个地址落入同一个内存区,则存在一个bank冲突,并且访问必须被序列化。 硬件将存储器请求与银行冲突拆分为许多单独的无冲突请求,并将吞吐量降低一个等于单独存储器请求数量的因子。 如果单独的内存请求的数量是n,则初始的内存请求被认为是导致n路bank冲突。
为了获得最大性能,重要的是要了解内存地址如何映射到内存组以便调度内存请求以最小化组冲突。 这在计算能力3.x,5.x,6.x和7.x的设备中分别描述在计算能力3.x,计算能力5.x,计算能力6.x和计算能力7.x中。
纹理和表面内存:
纹理和表面内存空间驻留在设备内存中,并缓存在纹理缓存中,所以纹理读取或表面读取只需要在缓存未命中时从设备内存读取一个内存,否则只花费一次从纹理缓存读取。 纹理缓存针对2D空间局部性进行了优化,因此读取纹理或表面地址相同的变形的线程可以达到最佳性能。 此外,它设计用于具有恒定延迟的流式抓取; 高速缓存命中减少了DRAM带宽需求,但不能提取等待时间.
通过纹理或表面读取读取设备内存可带来一些好处,可以使其成为从全局或常量内存中读取设备内存的有利替代方案:

  • 如果内存读取不遵循全局或常量内存读取必须遵循的访问模式以获得良好的性能,则可以实现更高的带宽,只要在纹理提取或表面读取中存在局部性;
  • 寻址计算在内核之外由专用单元执行;
  • 打包的数据可以在单个操作中广播以分离变量;
  • 8位和16位整型输入数据可以选择性地转换为[0.0,1.0]或[-1.0,1.0]范围内的32位浮点值(请参阅纹理存储器)
    timg
目录
相关文章
|
并行计算 C语言 存储
|
并行计算 API 调度
CUDA学习(八十八)
3.虽然__syncthreads()一直被记录为同步线程块中的所有线程,但Pascal和以前的体系结构只能在warp级别强制执行同步。 在某些情况下,只要每条经线中至少有一条线达到屏障,就可以在不被每条线执行的情况下成功实现屏障。
1721 0
|
并行计算 C语言 编译器
|
并行计算 API 编译器
CUDA学习(六十五)
很早之前就发现云栖社区的编辑器有一个Bug,往草稿箱存博客,当草稿箱博客数超过十篇时,无法再选择十篇前的博客进行编辑
2380 0
|
并行计算 API
|
存储 并行计算 API
|
并行计算 API
|
并行计算 API 索引