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

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.rhkb.cn/news/34258.html

如若内容造成侵权/违法违规/事实不符,请联系长河编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

睡不着运动锻炼贴士

在快节奏的现代生活中&#xff0c;失眠似乎已成为许多人的“夜间伴侣”。夜晚辗转反侧&#xff0c;白天精神不振&#xff0c;这样的恶性循环让许多人苦不堪言。其实&#xff0c;除了调整作息和饮食习惯&#xff0c;适当的运动也是改善睡眠的一剂良药。今天&#xff0c;就让我们…

java数据结构(复杂度)

一.时间复杂度和空间复杂度 1.时间复杂度 衡量一个程序好坏的标准&#xff0c;除了能处理各种异常&#xff0c;还有就是时间效率&#xff0c;当然&#xff0c;对于一些配置好的电脑数据处理起来就是比配置低的高&#xff0c;但从后期发展来看&#xff0c;当数据量足够庞大时&…

NAT和NAPT的介绍

一、NAT的介绍以及作用 二、NAPT的介绍以及作用 三、NAT vs NAPT 一、NAT的介绍以及作用 1.1 NAT的介绍 NAT&#xff08;Network Address Translation&#xff09;是一种广泛应用于互联网的技术&#xff0c;主要用于解决IPv4地址耗尽问题&#xff0c;同时提供网络安全和网络…

VSCode通过SSH免密远程登录Windows服务器

系列 1.1 VSCode通过SSH远程登录Windows服务器 1.2 VSCode通过SSH免密远程登录Windows服务器 文章目录 系列1 准备工作2 本地电脑配置2.1 生成密钥2.2 VS Code配置密钥 3. 服务端配置3.1 配置SSH服务器sshd_config3.2 复制公钥3.3 配置权限&#xff08;常见问题&#xff09;3.…

大模型训练全流程深度解析

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到网站。https://www.captainbed.cn/north 文章目录 1. 大模型训练概览1.1 训练流程总览1.2 关键技术指标 2. 数据准备2.1 数据收集与清洗2.2 数据…

export、export default 和 module.exports 深度解析

文章目录 1. 模块系统概述1.1 模块系统对比1.2 模块加载流程 2. ES Modules2.1 export 使用2.2 export default 使用2.3 混合使用 3. CommonJS3.1 module.exports 使用3.2 exports 使用 4. 对比分析4.1 语法对比4.2 使用场景 5. 互操作性5.1 ES Modules 中使用 CommonJS5.2 Com…

AI芯片设计

目的&#xff1a;未来的时代&#xff0c;一定会是AI的时代&#xff0c;那么&#xff0c;AI时代的三个重要组成部分&#xff0c;我要参与其中之一&#xff01; 参考视频&#xff1a;AI芯片设计第一讲_哔哩哔哩_bilibili 端处理 云端

动手学深度学习:CNN和LeNet

前言 该篇文章记述从零如何实现CNN&#xff0c;以及LeNet对于之前数据集分类的提升效果。 从零实现卷积核 import torch def conv2d(X,k):h,wk.shapeYtorch.zeros((X.shape[0]-h1,X.shape[1]-w1))for i in range(Y.shape[0]):for j in range(Y.shape[1]):Y[i,j](X[i:ih,j:jw…

【开源代码解读】AI检索系统R1-Searcher通过强化学习RL激励大模型LLM的搜索能力

关于R1-Searcher的报告&#xff1a; 第一章&#xff1a;引言 - AI检索系统的技术演进与R1-Searcher的创新定位 1.1 信息检索技术的范式转移 在数字化时代爆发式增长的数据洪流中&#xff0c;信息检索系统正经历从传统关键词匹配到语义理解驱动的根本性变革。根据IDC的统计…

使用Node的http模块创建web服务,给客户端返回html页面时,css失效的根本原因(有助于理解http)

最近正在尝试使用node写后端&#xff0c;使用node创建http服务的时候&#xff0c;碰到了这样的一个问题&#xff1a; 这是我的源代码&#xff1a; import { createServer } from http import { join, dirname, extname } from path import { fileURLToPath } from url import…

JVM 2015/3/15

定义&#xff1a;Java Virtual Machine -java程序的运行环境&#xff08;java二进制字节码的运行环境&#xff09; 好处&#xff1a; 一次编写&#xff0c;到处运行 自动内存管理&#xff0c;垃圾回收 数组下标越界检测 多态 比较&#xff1a;jvm/jre/jdk 常见的JVM&…

IP风险度自检,互联网的安全“指南针”

IP地址就像我们的网络“身份证”&#xff0c;而IP风险度则是衡量这个“身份证”安全性的重要指标。它关乎着我们的隐私保护、账号安全以及网络体验&#xff0c;今天就让我们一起深入了解一下IP风险度。 什么是IP风险度 IP风险度是指一个IP地址可能暴露用户真实身份或被网络平台…

【鸿蒙】封装日志工具类 ohos.hilog打印日志

封装一个ohos.hilog打印日志 首先要了解hilog四大日志类型&#xff1a; info、debug、warm、error 方法中四个参数的作用 domain: number tag: string format: string ...args: any[ ] 实例&#xff1a; //普通的info日志&#xff0c;使用info方法来打印 //第一个参数 : 0x0…

走路碎步营养补充贴士

走路碎步&#xff0c;这种步伐不稳的现象&#xff0c;在日常生活中并不罕见&#xff0c;特别是对于一些老年人或身体较为虚弱的人来说&#xff0c;更是一种常见的行走状态。然而&#xff0c;这种现象可能不仅仅是肌肉或骨骼的问题&#xff0c;它还可能是身体在向我们发出营养缺…

Python软件和搭建运行环境

目录 一、Python安装全流程&#xff08;Windows/Mac/Linux&#xff09; 1. 下载官方安装包 2. 详细安装步骤&#xff08;以Windows为例&#xff09; 3. 环境变量配置&#xff08;Mac/Linux&#xff09; 二、虚拟环境管理&#xff08;关键&#xff01;&#xff09; 为什么需…

【蓝桥杯】省赛:神奇闹钟

思路 python做这题很简单&#xff0c;灵活用datetime库即可 code import os import sys# 请在此输入您的代码 import datetimestart datetime.datetime(1970,1,1,0,0,0) for _ in range(int(input())):ls input().split()end datetime.datetime.strptime(ls[0]ls[1],&quo…

RabbitMQ (Java)学习笔记

目录 一、概述 ①核心组件 ②工作原理 ③优势 ④应用场景 二、入门 1、docker 安装 MQ 2、Spring AMQP 3、代码实现 pom 依赖 配置RabbitMQ服务端信息 发送消息 接收消息 三、基础 work Queue 案例 消费者消息推送限制&#xff08;解决消息堆积方案之一&#…

HW基本的sql流量分析和wireshark 的基本使用

前言 HW初级的主要任务就是看监控&#xff08;流量&#xff09; 这个时候就需要我们 了解各种漏洞流量数据包的信息 还有就是我们守护的是内网环境 所以很多的攻击都是 sql注入 和 webshell上传 &#xff08;我们不管对面是怎么拿到网站的最高权限的 我们是需要指出它是…

camellia redis proxy v1.3.3对redis主从进行读写分离(非写死,自动识别故障转移)

1 概述 camellia-redis-proxy是一款高性能的redis代理&#xff08;https://github.com/netease-im/camellia&#xff09;&#xff0c;使用netty4开发&#xff0c;主要特性如下&#xff1a; 支持代理到redis-standalone、redis-sentinel、redis-cluster。支持其他proxy作为后端…

贪吃蛇小游戏-简单开发版

一、需求 本项目旨在开发一个经典的贪吃蛇游戏&#xff0c;用户可以通过键盘控制蛇的移动方向&#xff0c;让蛇吃掉随机出现在游戏区域内的食物&#xff0c;每吃掉一个食物&#xff0c;蛇的身体长度就会增加&#xff0c;同时得分也会相应提高。游戏结束的条件为蛇撞到游戏区域的…