CUDA 显卡 GPU memory

﹏ヽ暗。殇╰゛Y 2022-05-08 01:34 402阅读 0赞

[http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015482.html][http_www.cnblogs.com_traceorigin_archive_2013_04_11_3015482.html]

CUDA存储器类型：

每个线程拥有自己的register and loacal memory;

每个线程块拥有一块shared memory;

所有线程都可以访问global memory;

还有，可以被所有线程访问的只读存储器：constant memory and texture memory

![11210324-3d471997b48843338346966adad096da.jpg][]

**1、 寄存器Register**

寄存器是GPU上的高速缓存器，其基本单元是寄存器文件，每个寄存器文件大小为32bit.

Kernel中的局部(简单类型)变量第一选择是被分配到Register中。

特点：每个线程私有，速度快。

**2、 局部存储器 local memory**

当register耗尽时，数据将被存储到local memory。如果每个线程中使用了过多的寄存器，或声明了大型结构体或数组，或编译器无法确定数组大小，线程的私有数据就会被分配到local 　　memory中。

特点：每个线程私有；没有缓存，慢。

注：在声明局部变量时，尽量使变量可以分配到register。如：

unsigned int mt\[3\];

改为：　unsigned int mt0, mt1, mt2;

**3、 共享存储器 shared memory**

可以被同一block中的所有线程读写

特点：block中的线程共有；访问共享存储器几乎与register一样快.

1.  `//u(i)= u(i)^2 + u(i-1)`
2.  `//Static`
3.  `__global__ example(float* u) {`
4.  `int i=threadIdx.x;`
5.  `__shared__ int tmp[4];`
6.  `tmp[i]=u[i];`
7.  `u[i]=tmp[i]*tmp[i]+tmp[3-i];`
8.  `}`
9.  
10. `int main() {`
11. `float hostU[4] = {1, 2, 3, 4};`
12. `float* devU;`
13. `size_t size = sizeof(float)*4;`
14. `cudaMalloc(&devU, size);`
15. `cudaMemcpy(devU, hostU, size,`
16. `cudaMemcpyHostToDevice);`
17. `example<<<1,4>>>(devU, devV);`
18. `cudaMemcpy(hostU, devU, size,`
19. `cudaMemcpyDeviceToHost);`
20. `cudaFree(devU);`
21. `return 0;`
22. `}`
23. 
24. `//Dynamic`
25. `extern __shared__ int tmp[];`
26. 
27. `__global__ example(float* u) {`
28. `int i=threadIdx.x;`
29. `tmp[i]=u[i];`
30. `u[i]=tmp[i]*tmp[i]+tmp[3-i];`
31. `}`
32. 
33. `int main() {`
34. `float hostU[4] = {1, 2, 3, 4};`
35. `float* devU;`
36. `size_t size = sizeof(float)*4;`
37. `cudaMalloc(&devU, size);`
38. `cudaMemcpy(devU, hostU, size, cudaMemcpyHostToDevice);`
39. `example<<<1,4,size>>>(devU, devV);`
40. `cudaMemcpy(hostU, devU, size, cudaMemcpyDeviceToHost);`
41. `cudaFree(devU);`
42. `return 0;`
43. `}`

**4、 全局存储器 global memory**

特点：所有线程都可以访问；没有缓存

1.  `//Dynamic`
2.  `__global__ add4f(float* u, float* v) {`
3.  `int i=threadIdx.x;`
4.  `u[i]+=v[i];`
5.  `}`
6.  `int main() {`
7.  `float hostU[4] = {1, 2, 3, 4};`
8.  `float hostV[4] = {1, 2, 3, 4};`
9.  `float* devU, devV;`
10. `size_t size = sizeof(float)*4;`
11. `cudaMalloc(&devU, size);`
12. `cudaMalloc(&devV, size);`
13. `cudaMemcpy(devU, hostU, size,`
14. `cudaMemcpyHostToDevice);`
15. `cudaMemcpy(devV, hostV, size,`
16. `cudaMemcpyHostToDevice);`
17. `add4f<<<1,4>>>(devU, devV);`
18. `cudaMemcpy(hostU, devU, size,`
19. `cudaMemcpyDeviceToHost);`
20. `cudaFree(devV);`
21. `cudaFree(devU);`
22. `return 0;`
23. `}`
24. 
25. `//static`
26. `__device__ float devU[4];`
27. `__device__ float devV[4];`
28. 
29. `__global__ addUV() {`
30. `int i=threadIdx.x;`
31. `devU[i]+=devV[i];`
32. `}`
33. 
34. `int main() {`
35. `float hostU[4] = {1, 2, 3, 4};`
36. `float hostV[4] = {1, 2, 3, 4};`
37. `size_t size = sizeof(float)*4;`
38. `cudaMemcpyToSymbol(devU, hostU, size, 0, cudaMemcpyHostToDevice);`
39. `cudaMemcpyToSymbol(devV, hostV, size, 0, cudaMemcpyHostToDevice);`
40. `addUV<<<1,4>>>();`
41. `cudaMemcpyFromSymbol(hostU, devU, size, 0, cudaMemcpyDeviceToHost);`
42. `return 0;`
43. `}`

**5、 常数存储器constant memory**

用于存储访问频繁的只读参数

特点：只读；有缓存；空间小(64KB)

注：定义常数存储器时，需要将其定义在所有函数之外，作用于整个文件

1 __constant__ int devVar;
    2 cudaMemcpyToSymbol(devVar, hostVar, sizeof(int), 0, cudaMemcpyHostToDevice)
    3 cudaMemcpyFromSymbol(hostVar, devVar, sizeof(int), 0, cudaMemcpyDeviceToHost)

**6、 纹理存储器 texture memory**

是一种只读存储器，其中的数据以一维、二维或者三维数组的形式存储在显存中。在通用计算中，其适合实现图像处理和查找，对大量数据的随机访问和非对齐访问也有良好的加速效果。

特点：具有纹理缓存，只读。

TNE END

\--------------------- 本文来自 Augusdi 的CSDN 博客 ，全文地址请点击：https://blog.csdn.net/augusdi/article/details/12186939?utm\_source=copy

=============================================================================================

CUDA存储器模型：

GPU片内：register，shared memory；

板载显存：local memory,constant memory, texture memory, texture memory,global memory;

host 内存： host memory, pinned memory.

基本单元：register file （32bit/each）

计算能力1.0/1.1版本硬件：8192/SM；

计算能力1.2/1.3版本硬件： 16384/SM;

每个线程占有的register有限，编程时不要为其分配过多私有变量；

local memory：寄存器被使用完毕，数据将被存储在局部存储器中；

大型结构体或者数组；

无法确定大小的数组；

线程的输入和中间变量；

定义线程私有数组的同时进行初始化的数组被分配在寄存器中；

shared memory：访问速度与寄存器相似；

实现线程间通信的延迟最小；

保存公用的计数器或者block的公用结果；

硬件1.0~1.3中，16KByte/SM,被组织为16个bank；

声明关键字 \_shared\_ int sdata\_static\[16\];

global memory:存在于显存中，也称为线性内存（显存可以被定义为线性存储器或者CUDA数组）；

cudaMalloc（）函数分配，cudaFree（）函数释放，cudaMemcpy（）进行主机端与设备端的数据传输；

初始化共享存储器需要调用cudaMemset（）；

二维三维数组：cudaMallocPitch（）和cudaMalloc3D（）分配线性存储空间，可以确保分配满足对齐要求；

cudaMemcpy2D（），cudaMemcpy3D（）与设备端存储器进行拷贝；

host内存：分为pageable memory 和 pinned memory

pageable memory： 通过操作系统API（malloc（），new（））分配的存储器空间；、

pinned memory：始终存在于物理内存中，不会被分配到低速的虚拟内存中，能够通过DMA加速与设备端进行通信；

cudaHostAlloc(), cudaFreeHost()来分配和释放pinned memory；

使用pinned memory优点：主机端-设备端的数据传输带宽高；

某些设备上可以通过zero-copy功能映射到设备地址空间，从GPU直接访问，省掉主存与显存间进行数据拷贝的工作；

pinned memory 不可以分配过多：导致操作系统用于分页的物理内存变， 导致系统整体性能下降；通常由哪个cpu线程分配，就只有这个线程才有访问权限；

cuda2.3版本中，pinned memory功能扩充：

portable memory：让控制不同GPU的主机端线程操作同一块portable memory，实现cpu线程间通信；使用cudaHostAlloc（）分配页锁定内存时，加上cudaHostAllocPortable标志；

write-combined Memory：提高从cpu向GPU单向传输数据的速度；不使用cpu的L1，L2 cache对一块pinned memory中的数据进行缓冲，将cache资源留给其他程序使用；在pci-e总线传输期间不会被来自cpu的监视打断；在调用cudaHostAlloc（）时加上cudaHostAllocWriteCombined标志；cpu从这种存储器上读取的速度很低；

mapped memory：两个地址：主机端地址（内存地址），设备端地址（显存地址）。 可以在kernnel程序中直接访问mapped memory中的数据，不必在内存和显存之间进行数据拷贝，即zero-copy功能；在主机端可以由cudaHostAlloc（）函数获得，在设备端指针可以通过cudaHostGetDevicePointer（）获得；通过cudaGetDeviceProperties（）函数返回的canMapHostMemory属性知道设备是否支持mapped memory；在调用cudaHostAlloc（）时加上cudaHostMapped标志，将pinned memory映射到设备地址空间；必须使用同步来保证cpu和GPu对同一块存储器操作的顺序一致性；显存中的一部分可以既是portable memory又是mapped memory；在执行CUDA操作前，先调用cudaSetDeviceFlags（）（加cudaDeviceMapHost标志）进行页锁定内存映射。

constant memory：只读地址空间；位于显存，有缓存加速；64Kb；用于存储需要频繁访问的只读参数 ；只读；使用\_constant\_ 关键字，定义在所有函数之外；两种常数存储器的使用方法：直接在定义时初始化常数存储器；定义一个constant数组，然后使用函数进行赋值；

texture memory：只读；不是一块专门的存储器，而是牵涉到显存、两级纹理缓存、纹理拾取单元的纹理流水线；数据常以一维、二维或者三维数组的形式存储在显存中；缓存加速；可以声明大小比常数存储器大得多；适合实现图像树立和查找表；对大量数据的随机访问或非对齐访问有良好的加速效果；在kernel中访问纹理存储器的操作成为纹理拾取（texture fetching）；纹理拾取使用的坐标与数据在显存中的位置可以不同，通过纹理参照系约定二者的映射方式；将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定（texture binding）；显存中可以绑定到纹理的数据有：普通线性存储器和cuda数组；存在缓存机制；可以设定滤波模式，寻址模式等；

\--------------------- 本文来自 caoeryingzi 的CSDN 博客 ，全文地址请点击：https://blog.csdn.net/caoeryingzi/article/details/21323475?utm\_source=copy

[http_www.cnblogs.com_traceorigin_archive_2013_04_11_3015482.html]: http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015482.html
[11210324-3d471997b48843338346966adad096da.jpg]: /images/20220508/f2e34c8ffa724a5998961a3aba0ac0d9.png