CUDA编程之内存

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）‌	主机内存	手动释放	高带宽	主机与设备间快速数据传输（如流处理）	分配开销大，避免过量使用