CUDA实践指南(二十)-阿里云开发者社区

开发者社区> 开发与运维> 正文

CUDA实践指南(二十)

简介:

常量内存:
设备上总共有64 KB的常量内存。 恒定的内存空间被缓存。 因此,从恒定内存读取仅从缓存未命中的设备内存读取一个内存; 否则,它只需要从常量缓存中读取一次。 序列化由warp中的线程访问不同地址的成本,因此成本与warp中所有线程读取的唯一地址数成线性比例关系。 因此,当相同warp中的线程只访问几个不同的位置时,常量缓存是最好的。 如果一个warp的所有线程访问相同的位置,那么恒定的内存可以像寄存器访问一样快。
寄存器:
通常,访问寄存器消耗每个指令零个额外的时钟周期,但由于寄存器读写后依赖性和寄存器存储体冲突可能会发生延迟。
写后读取依赖关系的延迟大约为24个周期,但这种延迟完全隐藏在多处理器上,每个多处理器具有足够的并发线程warp。对于计算能力2.0的设备(每个多处理器具有32个CUDA核心),可能需要多达768个线程(24个warps)才能完全隐藏延迟等等,以满足更高计算能力的设备。
编译器和硬件线程调度器将尽可能优化地调度指令以避免寄存器存储器组冲突。当每个块的线程数是64的倍数时,它们会达到最佳效果。除了遵循此规则之外,应用程序无法直接控制这些存储库冲突。特别是,没有与寄存器相关的理由将数据打包到float4或int4类型中。
Register Pressure(寄存器压力):
当给定任务没有可用寄存器时发生寄存器压力。
尽管每个多处理器都包含数千个32位寄存器(请参见CUDA C编程指南的功能和技术规格),但它们在并发线程之间进行分区。 要防止编译器分配过多的寄存器,请使用-maxrregcount = N编译器命令行选项(请参见nvcc)或启动界限内核定义限定符(请参阅CUDA C编程指南的执行配置)来控制寄存器的最大数量 分配给每个线程。
分配:
通过cudaMalloc()和cudaFree()进行设备内存分配和解除分配是昂贵的操作,因此应尽可能重新使用和/或分配设备内存,以尽量减少分配对整体性能的影响。

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

分享:
开发与运维
使用钉钉扫一扫加入圈子
+ 订阅

集结各类场景实战经验,助你开发运维畅行无忧

其他文章