CUDA的内存类型有全局内存、共享内存、常量内存、纹理内存、本地内存、寄存器等。我们需要分别了解它们的特点和使用场景。在CUDA编程中,合理利用各种内存类型对性能优化至关重要。
1. 全局内存(Global Memory)
- 特点:设备中最大、最慢的内存,所有线程均可访问,需通过合并访问优化带宽。
- 使用方法:
- 分配与传输:
float *d_data; cudaMalloc(&d_data, size); // 分配 cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice); // 数据传输
- 核函数访问:
__global__ void kernel(float *data) {int idx = blockIdx.x * blockDim.x + threadIdx.x;data[idx] *= 2; // 合并访问(连续线程访问连续地址) }
- 分配与传输:
2. 共享内存(Shared Memory)
- 特点:块内线程共享,速度快,需避免Bank Conflict。
- 使用方法:
- 静态声明:
__global__ void reduce0(float *g_in,float *g_out) {//每个线程从全局内存中加载一个对应位置元素到共享内存__shared__ float s_data[256]; //共享内存大小等于线程块的大小int tid = threadIdx.x; //共享内存中的索引,即在线程块中的编号int i = blockIdx.x * blockDim.x + threadIdx.x; //全局内存中的索引s_data[tid] = g_in[i];__syncthreads(); //同步等待共享内存加载完毕//归约操作优化(避免Bank Conflict)//在共享内存做相邻配对归约,线程和数据序号一一对应for(int s = 1; s < blockDim.x; s *= 2) {if(tid % (2 * s) == 0) {s_data[tid] += s_data[tid + s];}__syncthreads();}//把结果写回全局内存if (tid == 0) g_out[blockIdx.x] = s_data[0]; }__global__ void reduce1(float *g_in,float *g_out) {//每个线程从全局内存中加载一个对应位置元素到共享内存__shared__ float s_data[256];//共享内存大小等于线程块的大小int tid = threadIdx.x; //共享内存中的索引,即在线程块中的编号int i = blockIdx.x*blockDim.x + threadIdx.x;//全局内存中的索引s_data[tid] = g_in[i];__syncthreads(); //同步等待共享内存加载完毕//归约操作优化(避免Bank Conflict)//在共享内存做相邻配对归约,线程和数据序号间隔对应for(int s = 1; s < blockDim.x; s *= 2) {int index = 2 * s * tid;if (index < blockDim.x) {s_data[index] += s_data[index + s];}__syncthreads();}//把结果写回全局内存if (tid == 0) g_out[blockIdx.x] = s_data[0]; }__global__ void reduce2(float *g_in,float *g_out) {//每个线程从全局内存中加载一个对应位置元素到共享内存__shared__ float s_data[256];//共享内存大小等于线程块的大小int tid = threadIdx.x; //共享内存中的索引,即在线程块中的编号int i = blockIdx.x * blockDim.x + threadIdx.x;//全局内存中的索引s_data[tid] = g_in[i];__syncthreads(); //同步等待共享内存加载完毕//归约操作优化(避免Bank Conflict)//在共享内存做交错配对归约for(int s = (blockDim.x >> 1); s > 0; s >>= 1) {if (tid < s) {s_data[tid] += s_data[tid + s];}__syncthreads();}//把结果写回全局内存if (tid == 0) g_out[blockIdx.x] = s_data[0]; }
- 动态声明:
extern __shared__ float sdata[]; // 核函数调用时指定大小:<<<grid, block, smem_size>>>
- 静态声明:
3. 常量内存(Constant Memory)
- 特点:只读,适合频繁访问的常量数据,具有缓存优化。
- 使用方法:
- 声明与数据拷贝:
__constant__ float const_data[1024]; cudaMemcpyToSymbol(const_data, h_data, sizeof(float)*1024);
- 核函数中直接访问:
__global__ void kernel() {float value = const_data[threadIdx.x]; }
- 声明与数据拷贝:
4. 纹理内存(Texture Memory)
- 特点:适合具有空间局部性的访问,如图像处理。
-
纹理内存的寻址模式:
寻址模式有几种:cudaAddressModeWrap、cudaAddressModeClamp、cudaAddressModeBorder、cudaAddressModeMirror
cudaAddressModeWrap:循环模式,超出范围的坐标会循环到另一侧。比如,当x坐标超过宽度时,会回到起始位置,类似取模操作。
cudaAddressModeClamp:防止数据外溢,越界坐标取边界值。
cudaAddressModeBorder:严格限制数据范围,越界坐标返回零值。
cudaAddressModeMirror:对称信号处理,越界坐标镜像对称。 - 使用方法:
- 创建纹理对象:
texture<float, 1, cudaReadModeElementType> tex_ref; cudaArray *cuArray; cudaMallocArray(&cuArray, &tex_ref.channelDesc, size); cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice); cudaBindTextureToArray(tex_ref, cuArray);
- 核函数采样:
__global__ void kernel() {float value = tex1Dfetch(tex_ref, threadIdx.x); }
- 创建纹理对象:
硬件插值功能实现代码:
核函数:
//定义纹理内存变量
texture<float, cudaTextureType2D, cudaReadModeElementType> tex_src;__global__ void resize_img_ker(int row, int col, float x_a, float y_a, uchar *out)
{int x = threadIdx.x + blockDim.x * blockIdx.x; //colint y = threadIdx.y + blockDim.y * blockIdx.y; //rowif (x < col && y < row){float xx = x*x_a;float yy = y*y_a;//这里的xx和yy都是浮点数,tex2D函数返回的数值已经过硬件插值了,所以不需要开发者再进行插值啦~out[y*col+x] = (uchar)tex2D(tex_src, xx, yy);}
}
主体函数:
void resize_img_cuda(Mat src, Mat &dst, float row_m, float col_m)
{const int row = (int)(src.rows*row_m);const int col = (int)(src.cols*col_m);const int srcimg_size = src.rows*src.cols*sizeof(float);const int dstimg_size = row*col;const float x_a = 1.0 / col_m;const float y_a = 1.0 / row_m;uchar *dst_cuda;cudaMalloc((void**)&dst_cuda, dstimg_size);Mat src_tmp;src.convertTo(src_tmp, CV_32F); //注意这里要把图像转换为float浮点型,否则线性插值模式无法使用cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();//声明数据类型cudaArray *cuArray_src; //定义CUDA数组cudaMallocArray(&cuArray_src, &channelDesc, src_tmp.cols, src_tmp.rows); //分配大小为col*row的CUDA数组tex_src.addressMode[0] = cudaAddressModeWrap;//寻址方式tex_src.addressMode[1] = cudaAddressModeWrap;//寻址方式 如果是三维数组则设置texRef.addressMode[2]tex_src.normalized = false;//是否对纹理坐标归一化tex_src.filterMode = cudaFilterModeLinear;//硬件插值方式:最邻近插值--cudaFilterModePoint 双线性插值--cudaFilterModeLinearcudaBindTextureToArray(&tex_src, cuArray_src, &channelDesc); //把CUDA数组绑定到纹理内存cudaMemcpyToArray(cuArray_src, 0, 0, src_tmp.data, srcimg_size, cudaMemcpyHostToDevice); //把源图像数据拷贝到CUDA数组dim3 Block_resize(16, 16);dim3 Grid_resize((col + 15) / 16, (row + 15) / 16);//调用核函数resize_img_ker << <Grid_resize, Block_resize >> > (row, col, x_a, y_a, dst_cuda);dst = Mat::zeros(row, col, CV_8UC1);cudaMemcpy(dst.data, dst_cuda, dstimg_size, cudaMemcpyDeviceToHost);cudaFree(dst_cuda);cudaFreeArray(cuArray_src);cudaUnbindTexture(tex_src);
}
5. 本地内存(Local Memory)
- 特点:线程私有,速度慢,由编译器自动分配(如大数组或寄存器不足时)。
- 优化:减少使用,优先使用寄存器或共享内存。
__global__ void kernel() {float local_var; // 寄存器分配// float large_array[100]; // 可能溢出到本地内存 }
6. 寄存器(Registers)
- 特点:最快的内存,但数量有限。优化方法包括减少变量或循环展开。
__global__ void kernel() {int tid = threadIdx.x; // 使用寄存器// 循环展开减少寄存器压力float sum = data + data + data + data; }
7. 固定内存(Pinned Memory)
- 特点:主机内存,加速主机与设备间传输。
- 使用方法:
float *h_pinned; cudaHostAlloc(&h_pinned, size, cudaHostAllocDefault); // 分配固定内存 cudaMemcpy(d_data, h_pinned, size, cudaMemcpyHostToDevice); // 快速传输
八、总结
内存类型 | 作用域 | 生命周期 | 速度 | 使用场景 | 注意事项 |
---|---|---|---|---|---|
全局内存 | 所有线程 | 手动释放 | 最慢 | 大数据传输,跨线程块通信 | 需合并访问(连续地址访问) |
共享内存 | 线程块内 | 线程块执行期间 | 快 | 线程块内协作(如归约、矩阵分块) | 避免Bank Conflict(线程访问不同Bank) |
常量内存 | 所有线程 | 程序结束 | 较快 | 频繁读取的常量数据(如配置参数) | 只读,需提前用cudaMemcpyToSymbol 拷贝 |
纹理内存 | 所有线程 | 手动释放 | 较快 | 具有空间局部性的访问(如图像采样) | 需绑定纹理对象,支持插值和缓存 |
本地内存 | 单个线程 | 线程执行期间 | 慢 | 寄存器溢出时的临时变量(如大数组) | 尽量避免使用,优先用寄存器/共享内存 |
寄存器 | 单个线程 | 线程执行期间 | 最快 | 局部变量和临时计算 | 数量有限(每个线程约255个) |
固定内存(Pinned) | 主机内存 | 手动释放 | 高带宽 | 主机与设备间快速数据传输(如流处理) | 分配开销大,避免过量使用 |