title: 【CUDA 基础】5.5 常量内存
categories:
- CUDA
- Freshman
tags:
- CUDA常量内存
- CUDA只读缓存
toc: true
date: 2018-06-06 06:20:07
Abstract: 本文介绍另外两种内存——常量内存，只读缓存
Keywords: CUDA常量内存，CUDA只读缓存
开篇废话
本来早上想写数值分析类的博客，但是看CUDA已经拖太久了，还是写一篇CUDA的博客，争取快些结束本系列，本章完成后还有五章。
常量内存
本文介绍常量内存和只读缓存，常量内存是专用内存，他用于只读数据和线程束统一访问某一个数据，常量内存对内核代码而言是只读的，但是主机是可以修改（写）只读内存的，当然也可以读。
注意，常量内存并不是在片上的，而是在DRAM上，而其有在片上对应的缓存，其片上缓存就和一级缓存和共享内存一样， 有较低的延迟，但是容量比较小，合理使用可以提高内和效率，每个SM常量缓存大小限制为64KB。
我们可以发现，所有的片上内存，我们是不能通过主机赋值的，我们只能对DRAM上内存进行赋值。
每种内存访问都有最优与最坏的访问方式，主要原因是内存的硬件结构和底层设计原因，比如全局内存按照连续对去访问最优，交叉访问最差，共享内存无冲突最优，都冲突就会最差，其根本原因在于硬件设计，而我们的常量内存的最优访问模式是线程束所有线程访问一个位置，那么这个访问是最优的。如果要访问不同的位置，就要编程串行了，作为对比，这种情况相当于全局内存完全不连续，共享内存的全部冲突。
数学上，一个常量内存读取成本与线程束中线程读取常量内存地址个数呈线性关系。
常量内存的声明方式：
__constant
常量内存变量的生存周期与应用程序生存周期相同，所有网格对声明的常量内存都是可以访问的，运行时对主机可见，当CUDA独立编译被使用的，常量内存跨文件可见，这个要后面才会介绍。
初始化常量内存使用一下函数完成
cudaError_t cudaMemcpyToSymbol(const void *symbol, const void * src, size_t count, size_t offset, cudaMemcpyKind kind)
和我们之前使用的copy到全局内存的函数类似，参数也类似，包含传输到设备，以及从设备读取，kind的默认参数是传输到设备。
posted on
