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