当前位置： 首页 > news >正文

CUDA编程之内存

news 来源：原创 2025/5/13 20:17:18

        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;  //col
  int y = threadIdx.y + blockDim.y * blockIdx.y;  //row


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

相关文章：

  • axios的二次封装
  • nginx配置转发到另一个网站或另一台服务器的服务
  • HOT100——栈篇Leetcode739. 每日温度
  • 简单的bug+1
  • 前沿计组知识入门（三）
  • React 和 Vue 框架设计原理对比分析
  • springboot集成flink实现DM数据库同步到ES
  • 反射（第三篇）、代理模式、静态代理和动态代理、InvocationHandler实际应用
  • 淘宝API实时监控系统开发：商品价格波动预警与竞品分析实战
  • 基于python+django+vue.js开发的医院门诊管理系统/医疗管理系统源码+运行
  • 大型语言模型（LLM）：解码人工智能的“语言基因“
  • 数据结构（C\C++）——算法复杂度
  • 影刀RPA结合Pandas的优势
  • @Autowired 注解在构造器上的使用规则（字段注入也挺好的）
  • DeepSeek在医学领域的应用
  • Go语言对于MySQL的基本操作
  • .NET 9 中 OpenAPI 替代 Swagger 文档生成
  • Python精进系列：解包（Unpacking）用法之 *args 和 **kwargs
  • 使用py-ffmpeg批量合成视频的脚本
  • HarmonyOS NEXT开发进阶（十二）：build-profile.json5 文件解析
  • 图讯丨习近平出席中国－拉美和加勒比国家共同体论坛第四届部长级会议开幕式
  • 2025年上海好护士揭晓，上海护士五年增近两成达12.31万人
  • 湖南湘西州副州长刘冬生主动交代问题，接受审查调查
  • 巫蛊：文化的历史暗流
  • 石家庄推动城市能级与民生福祉并进
  • 梅花奖在上海｜穿上初演时的服装，“鹮仙”朱洁静再起飞