CUDA内存类型memory
2013-09-30 10:34
330 查看
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
![](http://images.cnitblog.com/blog/491395/201304/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一样快.
4、 全局存储器 global memory 特点:所有线程都可以访问;没有缓存
5、 常数存储器constant memory 用于存储访问频繁的只读参数 特点:只读;有缓存;空间小(64KB) 注:定义常数存储器时,需要将其定义在所有函数之外,作用于整个文件
CUDA存储器类型:每个线程拥有自己的register and loacal memory;每个线程块拥有一块shared memory;所有线程都可以访问global memory;还有,可以被所有线程访问的只读存储器:constant memory and texture memory
![](http://images.cnitblog.com/blog/491395/201304/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一样快.
//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
相关文章推荐
- 自适应网页前端设计相关
- Delphi XE2 之 FireMonkey 入门(23) - 数据绑定: TBindingsList: TBindExpression
- 【JAVA正则表达式使用方法】
- [R] subset a list
- 怎样看待360 提示可以卸载手机预装软件商店引发的事件?
- 1
- 网站被泛解析后的解决方法
- java写的ftp客户端源码
- BFS+思维-poj-3182-The Grove
- Hibernate 的两种继承映射关系,单表映射
- 教你怎么鉴别发霉大米?
- Delphi XE2 之 FireMonkey 入门(22) - 数据绑定: BindingSource、BindingName、FindBinding()、Binding[]
- 若何塑造行业品牌的几种门径
- Android 之 内存管理
- 58同城提交赴美IPO申请:图解招股书要点
- java面试基础
- 0930
- Delphi XE2 之 FireMonkey 入门(21) - 和 FMX 相关的类(表)
- iOS制作Static Library(静态库),实现多工程的连编
- Geoserver 标准乱码、属性查询源码分析