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
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一样快.
-
//u(i)= u(i)^2 + u(i-1)
-
//Static
-
__global__ example(float* u) {
-
int i=threadIdx.x;
-
__shared__ int tmp[4];
-
tmp[i]=u[i];
-
u[i]=tmp[i]*tmp[i]+tmp[3-i];
-
}
-
int main() {
-
float hostU[4] = {1, 2, 3, 4};
-
float* devU;
-
size_t size = sizeof(float)*4;
-
cudaMalloc(&devU, size);
-
cudaMemcpy(devU, hostU, size,
-
cudaMemcpyHostToDevice);
-
example<<<1,4>>>(devU, devV);
-
cudaMemcpy(hostU, devU, size,
-
cudaMemcpyDeviceToHost);
-
cudaFree(devU);
-
return 0;
-
}
-
//Dynamic
-
extern __shared__ int tmp[];
-
__global__ example(float* u) {
-
int i=threadIdx.x;
-
tmp[i]=u[i];
-
u[i]=tmp[i]*tmp[i]+tmp[3-i];
-
}
-
int main() {
-
float hostU[4] = {1, 2, 3, 4};
-
float* devU;
-
size_t size = sizeof(float)*4;
-
cudaMalloc(&devU, size);
-
cudaMemcpy(devU, hostU, size, cudaMemcpyHostToDevice);
-
example<<<1,4,size>>>(devU, devV);
-
cudaMemcpy(hostU, devU, size, cudaMemcpyDeviceToHost);
-
cudaFree(devU);
-
return 0;
-
}
4、 全局存储器 global memory
特点:所有线程都可以访问;没有缓存
-
//Dynamic
-
__global__ add4f(float* u, float* v) {
-
int i=threadIdx.x;
-
u[i]+=v[i];
-
}
-
int main() {
-
float hostU[4] = {1, 2, 3, 4};
-
float hostV[4] = {1, 2, 3, 4};
-
float* devU, devV;
-
size_t size = sizeof(float)*4;
-
cudaMalloc(&devU, size);
-
cudaMalloc(&devV, size);
-
cudaMemcpy(devU, hostU, size,
-
cudaMemcpyHostToDevice);
-
cudaMemcpy(devV, hostV, size,
-
cudaMemcpyHostToDevice);
-
add4f<<<1,4>>>(devU, devV);
-
cudaMemcpy(hostU, devU, size,
-
cudaMemcpyDeviceToHost);
-
cudaFree(devV);
-
cudaFree(devU);
-
return 0;
-
}
-
//static
-
__device__ float devU[4];
-
__device__ float devV[4];
-
__global__ addUV() {
-
int i=threadIdx.x;
-
devU[i]+=devV[i];
-
}
-
int main() {
-
float hostU[4] = {1, 2, 3, 4};
-
float hostV[4] = {1, 2, 3, 4};
-
size_t size = sizeof(float)*4;
-
cudaMemcpyToSymbol(devU, hostU, size, 0, cudaMemcpyHostToDevice);
-
cudaMemcpyToSymbol(devV, hostV, size, 0, cudaMemcpyHostToDevice);
-
addUV<<<1,4>>>();
-
cudaMemcpyFromSymbol(hostU, devU, size, 0, cudaMemcpyDeviceToHost);
-
return 0;
-
}
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: 访问延迟极低;
基本单元: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