CUDA线性内存分配
2013-09-30 10:35
246 查看
http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015554.html
概述:线性存储器可以通过cudaMalloc()、cudaMallocPitch()和cudaMalloc3D()分配1、1D线性内存分配
概述:线性存储器可以通过cudaMalloc()、cudaMallocPitch()和cudaMalloc3D()分配1、1D线性内存分配
![](http://common.cnblogs.com/images/copycode.gif)
![](http://common.cnblogs.com/images/copycode.gif)
3、3D线性内存
![](http://common.cnblogs.com/images/copycode.gif)
概述:线性存储器可以通过cudaMalloc()、cudaMallocPitch()和cudaMalloc3D()分配1、1D线性内存分配
1 cudaMalloc(void**,int) //在设备端分配内存 2 cudaMemcpy(void* dest,void* source,int size,enum direction) //数据拷贝 3 cudaMemcpyToSymbol //将数据复制到__constant__变量中,或者__device__变量中 4 cudaMemcpyFromSynbol //同上相反 5 cudaFree() //内存释放 6 cudaMemset() //内存初始化注意:主机和设备间的数据交换会自动同步,而设备与设备却不会,需要使用cudaThreadSynchronize()2、2D线性内存分配2.1 分配
1 cudaMallocPitch( void** devPtr,size_t* pitch,size_t widthInBytes,size_t height ) //在线性内存中分配二维数组,width的单位是字节,而height单位是数据类型c语言申请2维内存时,一般是连续存放的。a[y][x]存放在第y*widthofx*sizeof(元素)+x*sizeof(元素)个字节。但在cuda的global memory访问中,从256字节对齐的地址(addr=0, 256, 512, ...)开始的连续访问是最有效率的。这样,为了提高内存访问的效率,有了cudaMallocPitch函数。cudaMallocPitch函数分配的内存中,数组的每一行的第一个元素的开始地址都保证是对齐的。因为每行有多少个数据是不确定的,widthofx*sizeof(元素)不一定是256的倍数。故此,为保证数组的每一行的第一个元素的开始地址对齐,cudaMallocPitch在分配内存时,每行会多分配一些字节,以保证widthofx*sizeof(元素)+多分配的字节是256的倍数(对齐)。这样,上面的y*widthofx*sizeof(元素)+x*sizeof(元素)来计算a[y][x]的地址就不正确了。而应该是y*[widthofx*sizeof(元素)+多分配的字节]+x*sizeof(元素)。而函数中返回的pitch的值就是widthofx*sizeof(元素)+多分配的字节。说明:widthInBytes作为输入参数,应该是widthofx*sizeof(元素);这样的话,复制内容时也要作相应的修改。2.2 访问
1 T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column; //元素访问方式cudaMallocPitch()以*pitch的形式返回间距,即所分配存储器的宽度,以字节为单位。间距用作存储器分配的一个独立参数,用于在2D数组内计算地址。2.3 拷贝
1 cudaMemcpy2D( void* dst,size_t dpitch,const void* src,size_t spitch,size_t width,size_t height,enum cudaMemcpyKind kind )这里需要特别注意width与pitch的区别,width是实际需要拷贝的数据宽度而pitch是2D线性存储空间分配时对齐的行宽,而当数据传递发生在设备与主机之间时,主机端pitch==width.综上我们可以看到,CUDA下对二维线性空间的访问是不提供多下标支持的,访问时依然是通过计算偏移量得到,不同的地方在于使用pitch对齐后非常利于实现coalesce访问例:下面的代码分配了一个尺寸为width*height的二维浮点数组,同时演示了怎样在设备代码中遍历数组元素
1 // Host code 2 int width = 64, height = 64; 3 float* devPtr; 4 int pitch; 5 cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(float), height); 6 MyKernel<<<100, 512>>>(devPtr, pitch, width, height); 7 // Device code 8 __global__ void MyKernel(float* devPtr, int pitch, int width, int height){ 9 for (int r = 0; r < height; ++r) { 10 float* row = (float*)((char*)devPtr + r * pitch); 11 for (int c = 0; c < width; ++c) { 12 float element = row[c]; 13 } 14 } 15 }3、3D线性内存
1 cudaError_t cudaMalloc3D( 2 struct cudaPitchedPtr * pitchedDevPtr, 3 struct cudaExtent extent 4 )例:下面的代码分配了一个尺寸为width*height*depth的三维浮点数组,同时演示了怎样在设备代码中遍历数组元素
1 // Host code 2 cudaPitchedPtr devPitchedPtr; 3 cudaExtent extent = make_cudaExtent(64, 64, 64); 4 cudaMalloc3D(&devPitchedPtr, extent); 5 MyKernel<<<100, 512>>>(devPitchedPtr, extent); 6 // Device code 7 __global__ void MyKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent) { 8 char* devPtr = devPitchedPtr.ptr; 9 size_t pitch = devPitchedPtr.pitch; 10 size_t slicePitch = pitch * extent.height; 11 for (int z = 0; z < extent.depth; ++z) { 12 char* slice = devPtr + z * slicePitch; 13 for (int y = 0; y < extent.height; ++y) { 14 float* row = (float*)(slice + y * pitch); 15 for (int x = 0; x < extent.width; ++x) { float element = row[x]; 16 } 17 } 18 }
概述:线性存储器可以通过cudaMalloc()、cudaMallocPitch()和cudaMalloc3D()分配1、1D线性内存分配
1 cudaMalloc(void**,int) //在设备端分配内存 2 cudaMemcpy(void* dest,void* source,int size,enum direction) //数据拷贝 3 cudaMemcpyToSymbol //将数据复制到__constant__变量中,或者__device__变量中 4 cudaMemcpyFromSynbol //同上相反 5 cudaFree() //内存释放 6 cudaMemset() //内存初始化注意:主机和设备间的数据交换会自动同步,而设备与设备却不会,需要使用cudaThreadSynchronize()2、2D线性内存分配2.1 分配
1 cudaMallocPitch( void** devPtr,size_t* pitch,size_t widthInBytes,size_t height ) //在线性内存中分配二维数组,width的单位是字节,而height单位是数据类型c语言申请2维内存时,一般是连续存放的。a[y][x]存放在第y*widthofx*sizeof(元素)+x*sizeof(元素)个字节。但在cuda的global memory访问中,从256字节对齐的地址(addr=0, 256, 512, ...)开始的连续访问是最有效率的。这样,为了提高内存访问的效率,有了cudaMallocPitch函数。cudaMallocPitch函数分配的内存中,数组的每一行的第一个元素的开始地址都保证是对齐的。因为每行有多少个数据是不确定的,widthofx*sizeof(元素)不一定是256的倍数。故此,为保证数组的每一行的第一个元素的开始地址对齐,cudaMallocPitch在分配内存时,每行会多分配一些字节,以保证widthofx*sizeof(元素)+多分配的字节是256的倍数(对齐)。这样,上面的y*widthofx*sizeof(元素)+x*sizeof(元素)来计算a[y][x]的地址就不正确了。而应该是y*[widthofx*sizeof(元素)+多分配的字节]+x*sizeof(元素)。而函数中返回的pitch的值就是widthofx*sizeof(元素)+多分配的字节。说明:widthInBytes作为输入参数,应该是widthofx*sizeof(元素);这样的话,复制内容时也要作相应的修改。2.2 访问
1 T* pElement = (T*)((char*)BaseAddress + Row * pitch) + Column; //元素访问方式cudaMallocPitch()以*pitch的形式返回间距,即所分配存储器的宽度,以字节为单位。间距用作存储器分配的一个独立参数,用于在2D数组内计算地址。2.3 拷贝
1 cudaMemcpy2D( void* dst,size_t dpitch,const void* src,size_t spitch,size_t width,size_t height,enum cudaMemcpyKind kind )这里需要特别注意width与pitch的区别,width是实际需要拷贝的数据宽度而pitch是2D线性存储空间分配时对齐的行宽,而当数据传递发生在设备与主机之间时,主机端pitch==width.综上我们可以看到,CUDA下对二维线性空间的访问是不提供多下标支持的,访问时依然是通过计算偏移量得到,不同的地方在于使用pitch对齐后非常利于实现coalesce访问例:下面的代码分配了一个尺寸为width*height的二维浮点数组,同时演示了怎样在设备代码中遍历数组元素
![](http://common.cnblogs.com/images/copycode.gif)
1 // Host code 2 int width = 64, height = 64; 3 float* devPtr; 4 int pitch; 5 cudaMallocPitch((void**)&devPtr, &pitch, width * sizeof(float), height); 6 MyKernel<<<100, 512>>>(devPtr, pitch, width, height); 7 // Device code 8 __global__ void MyKernel(float* devPtr, int pitch, int width, int height){ 9 for (int r = 0; r < height; ++r) { 10 float* row = (float*)((char*)devPtr + r * pitch); 11 for (int c = 0; c < width; ++c) { 12 float element = row[c]; 13 } 14 } 15 }
![](http://common.cnblogs.com/images/copycode.gif)
3、3D线性内存
1 cudaError_t cudaMalloc3D( 2 struct cudaPitchedPtr * pitchedDevPtr, 3 struct cudaExtent extent 4 )例:下面的代码分配了一个尺寸为width*height*depth的三维浮点数组,同时演示了怎样在设备代码中遍历数组元素
![](http://common.cnblogs.com/images/copycode.gif)
1 // Host code 2 cudaPitchedPtr devPitchedPtr; 3 cudaExtent extent = make_cudaExtent(64, 64, 64); 4 cudaMalloc3D(&devPitchedPtr, extent); 5 MyKernel<<<100, 512>>>(devPitchedPtr, extent); 6 // Device code 7 __global__ void MyKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent) { 8 char* devPtr = devPitchedPtr.ptr; 9 size_t pitch = devPitchedPtr.pitch; 10 size_t slicePitch = pitch * extent.height; 11 for (int z = 0; z < extent.depth; ++z) { 12 char* slice = devPtr + z * slicePitch; 13 for (int y = 0; y < extent.height; ++y) { 14 float* row = (float*)(slice + y * pitch); 15 for (int x = 0; x < extent.width; ++x) { float element = row[x]; 16 } 17 } 18 }
![](http://common.cnblogs.com/images/copycode.gif)
相关文章推荐
- Delphi XE2 之 FireMonkey 入门(24) - 数据绑定: TBindingsList: TBindExpression.Direction
- 群发邮件要注意什么
- CUDA内存类型memory
- 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[]
- 若何塑造行业品牌的几种门径