CUDA内存类型memory
2013-09-30 10:34
260 查看
CUDA内存类型memory
2013-09-30 10:344224人阅读 评论(1)
收藏
举报
![](http://static.blog.csdn.net/images/category_icon.jpg)
分类:
CUDA(106)
![](http://static.blog.csdn.net/images/arrow_triangle%20_down.jpg)
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一样快.
[cpp]
view plain
copy
//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
特点:所有线程都可以访问;没有缓存
[cpp]
view plain
copy
//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
相关文章推荐
- DB2的JDBC驱动
- PE文件结构详解(三)PE导出表
- winform 模仿 百度搜索 时 自动匹配 效果
- CUDA内存类型memory
- 自适应网页前端设计相关
- 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面试基础