4.1.1 内存模型
CUDA内存模型提出了多种可编程内存的类型:
·寄存器 – 最快的内存空间,有限的资源,核函数中声明的没有其他修饰符的自变量,核函数使用超出限制,则用本地内存。
·共享内存 – 线程块(block)中所有线程都可见,类似CPU一级缓存,但可编程。一个块内的线程可以通过共享内存合作,访问共享内存必须同步,命令__syncthreads();
·本地内存 – 线程私有,本质上与全局内存在同一块存储区域,所以高延迟低带宽。
·常量内存 – 只读,修饰符是__constant__。通过cudaMemcpyToSymbol 来初始化。
·纹理内存 – 只读
·全局内存 – GPU中最大、延迟最高并且最常使用的内存。静态声明用修饰符__device__, 动态分配用cudaMalloc , 释放用cudaFree
如图所示是内存空间的层次结构
除了可编程的内存,GPU上还有不可编程的内存(缓存):
·一级缓存
·二级缓存
·只读常量缓存
·只读纹理缓存
下表总结了CUDA变量声明和它们相应的存储位置、作用域、生命周期和修饰符。
下表总结了各类存储器的主要特征
下面代码说明如何静态声明一个全局变量
#include <cuda_runtime.h>
#include <stdio.h>
#include "../common/common.h"
__device_