【NVIDIA CUDA】2023 CUDA夏令营编程模型(二)

博主未授权任何人或组织机构转载博主任何原创文章,感谢各位对原创的支持!
博主链接

本人就职于国际知名终端厂商,负责modem芯片研发。
在5G早期负责终端数据业务层、核心网相关的开发工作,目前牵头6G算力网络技术标准研究。


博客内容主要围绕:
       5G/6G协议讲解
       算力网络讲解(云计算,边缘计算,端计算)
       高级C语言讲解
       Rust语言讲解

文章目录

  • CUDA编程模型——共享内存
    • 一、多种CUDA存储单元介绍
      • 1.1 共享内容介绍
      • 1.2 配方式
      • 1.3 bank竞争
      • 1.4 如何避免冲突
    • 1.5 共享内存优化



CUDA编程模型——共享内存

一、多种CUDA存储单元介绍

在这里插入图片描述

内存访问速度(由快到慢):

  • Register file
  • Shared Memory
  • Constant Memory
  • Texture Memory
  • Local Memory and Global Memory:位于Device memory中,空间最大,latency最大,是GPU最基础的内存;

1.1 共享内容介绍

实际驻留在GPU芯片上的内存只有两种类型:寄存器和共享内存。所以,Shared Memory是目前最快的可以让多个线程通信的地方。那么,就有可能会出现同时有很多线程访问Shared Memory上的数据。为了克服这个同时访问的瓶颈,Shared Memory被分成32个逻辑块,称为bank。

  • Shared Memory可以被设置成16KB,32KB ,48KB…剩下的给L1缓存;
  • 带宽可以使32bit 或者 64 bit;
  • 可以被多线程同时访问,因此存储器被划分为 banks;
  • 连续的 32-bit 访存被分配到连续的 banks;
  • 每个 bank 每个周期可以响应一个地址;
  • 如果有多个bank的话可以同时响应更多地址申请;

1.2 配方式

静态分配:

  • __shared__ int s[64];
    动态分配:
  • dynamicKernel<<<1, n, n*sizeof(int)>>>(d_d, n);
    extern __shared__ int s[];

1.3 bank竞争

  1. 同常量内存一样,当一个 warp 中的所有线程访问同一地址的共享内存时,会触发一个广播(broadcast)机制到
    warp 中所有线程,这是最高效的;
  2. 如果同一个 half-warp/warp 中的线程访问同一个 bank中的不同地址时将发生 bank conflict;
  3. 每个 bank 除了能广播(broadca st)还可以多播(mutilcast)(计算能力 >= 2.0),也就是说,如果一个 warp 中的多个线程访问同一个 bank 的同一个地址时(其他线程也没有访问同一个bank 的不同地址)不会发生 bank
    conflict;
  4. 即使同一个 warp 中的线程随机的访问不同的 bank,只要没有访问同一个 bank 的不同地址就不会发生 bank conflict;

在这里插入图片描述

如果没有bank冲突的话,Shared memory 跟 registers 一样快:

  • 快速情况:
    • warp 内所有线程访问 不同 banks, 没有冲突
    • warp 内所有线程读取同一地址,没有冲突(广播)
  • 慢速情况:
    • Bank Conflict: warp 内多个线程访问同一个bank
    • 访存必须串行化

1.4 如何避免冲突

先看一个有bank冲突的例子:

在这里插入图片描述
在这里插入图片描述
一个warp中的线程会访问,同一列中的数据,产生了bank冲突。

解决方法:

  • memory padding方法
    在这里插入图片描述

    使用了上面的内存padding方法之后,访问顺序编程了右图所示的“斜线”的顺序,代码如下:

    在这里插入图片描述

1.5 共享内存优化



在这里插入图片描述

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

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

相关文章

Vue3 select循环多个,选项option不能重复被选

Vue3 select循环多个&#xff0c;选项option不能重复被选 环境&#xff1a;vue3tsviteelement plus 实现目标&#xff1a;Vue3 select循环多个&#xff0c;当其中一个option值被选后&#xff0c;其他select里面不能再重复选择该option值。第二种&#xff0c;当其中一个option值…

SNN论文总结

Is SNN a great work ? Is SNN a convolutional work ? ANN的量化在SNN中是怎么体现的&#xff0c;和threshold有关系吗&#xff0c;threshold可训练和这个有关吗&#xff08;应该无关&#xff09; 解决过发放不发放的问题。 Intuation SNN编码方式 Image to spike patter…

后端面试话术集锦第 十 篇:springMVC面试话术

这是后端面试集锦第十篇博文——springMVC面试话术❗❗❗ 1. 介绍一下springMVC springmvc是一个视图层框架,通过MVC模型让我们很方便的接收和处理请求和响应。 我给你说说他里边的几个核心组件吧: 它的核心控制器是DispatcherServlet,他的作用是接收用户请求,然后给用户…

SASS常用内置函数

1&#xff0c;math 引入&#xff1a;use "sass:math"; 常用函数&#xff1a; ceil($number) - 向上取整floor($number) - 向下取整round($number) - 四舍五入max($number...) - 比较若干数值并取最大值min($number...) - 比较若干数值并取最小值random() - 取0~1之…

C++项目:网络版本在线五子棋对战

目录 1.项目介绍 2.开发环境 3.核心技术 4. 环境搭建 5.websocketpp 5.1原理解析 5.2报文格式 5.3websocketpp常用接口介绍 5.4websocket服务器 6.JsonCpp使用 6.1Json数据格式 6.2JsonCpp介绍 7.MySQL API 7.1MySQL API介绍 7.2MySQL API使用 7.3实现增删改查…

TCP之超时重传、流量控制和拥塞控制

一、超时重传 TCP超时重传是TCP协议中的一种机制&#xff0c;用于在发生丢包或数据包未及时确认的情况下&#xff0c;重新发送未确认的数据段。 当发送方发送一个数据段后&#xff0c;会启动一个定时器&#xff08;称为超时计时器&#xff09;&#xff0c;等待接收方的确认。…

kafka如何避免消费组重平衡

目录 前言&#xff1a; 协调者 重平衡的影响 避免重平衡 重平衡发生的场景 参考资料 前言&#xff1a; Rebalance 就是让一个 Consumer Group 下所有的 Consumer 实例就如何消费订阅主题的所有分区达成共识的过程。在 Rebalance 过程中&#xff0c;所有 Consumer 实例…

大数据精准营销获客能为企业带来哪些东西?

广告圈里一句名言:我知道我的广告浪费了一半&#xff0c;但我不知道浪费了哪一半。当前&#xff0c;越来越多的企业在大数据思维指导下进行广告投放&#xff0c;广告能通过对人群的定向&#xff0c;投放给准确的目标顾客&#xff0c;特别是互联网广告现在能够做到根据不同的人向…

XSSchallenge1-20

test1 第一题直接在test插入XSS代码即可 test2 第二关对内容进行”“包裹 这里可以采用”>来绕过 test3 代码审计发现这里用了htmlspecialchars函数&#xff0c;这个函数对<>和’ “等进行了转义&#xff0c;这里可以用事件来绕过 test4 这里用了str_replace&a…

GIT使用教程(超详细)

目录 前言 1 git安装 2 增加git账户 3 git全局参数配置 4 创建本地仓库 5 关联远程分支 6 删除远程分支关联 7. 删除分支 8 git stash 9 git reset 10 git checkout 11 合并 12 git log 13 git提交模板 &#x1f388;个人主页&#x1f388;&#xff1a;li…

14家展商集中亮相!8月30-31日,智能汽车软件与座舱车联大会

过去几年&#xff0c;在特斯拉及新势力的带动下&#xff0c;车企的盈利模式正在寻求从“一次售卖”转变为“硬件预埋&#xff0b;软件付费解锁”&#xff0c;背后是驱动汽车软件架构的迭代&#xff0c;即从面向信号的软件架构&#xff0c;过渡至面向服务的SOA架构。 同时&#…

Flutter小功能实现-咖啡店

1 导航栏实现 效果图&#xff1a; 1.Package google_nav_bar: ^5.0.6 使用文档&#xff1a; google_nav_bar | Flutter Package 2.Code //MyBottomNavBar class MyBottomNavBar extends StatelessWidget {void Function(int)? onTabChange;MyBottomNavBar({super.key, …

JVM下篇知识

第01章&#xff1a;概述篇 第02章&#xff1a;JVM监控及诊断工具-命令行篇 第03章&#xff1a;JVM监控及诊断工具-GUI篇 第04章&#xff1a;JVM运行时参数 第05章&#xff1a;分析GC日志

Apple Configurator iphone ipad 设备管控 描述文件使用方法

一、准备 App Store 下载安装 Apple Configurator 二、Apple Configurator 注册组织&#xff0c; -----------这个组织可以是个人&#xff0c;或者其它组织导出-------再导入进来&#xff1a; 三、描述文件配置&#xff1a;“” 根据管控需求进行配置 “” 四、使用 Ap…

Docker consul容器服务自动发现和更新

目录 一、什么是服务注册与发现 二、Docker-consul集群 1.Docker-consul 2.registrator 3.Consul-template 三、Docker-consul实现过程 四、Docker-consul集群配置 1.下载consul服务 2.web服务器启动多例nginx容器&#xff0c;使用registrator自动发现 3.使用…

vant2 van-calendar组件增加清除按钮和确定按钮

利用自定义插槽增加一个清除按钮 <van-calendar ref"fTime1" select"selectTimePicker" confirm"changeTimePicker" :default-date"null" :show-confirm"false" v-model"timePickerShow" type"range&quo…

Docker之私有仓库 RegistryHarbor

目录 一、Docker私有仓库&#xff08;Registry&#xff09; 1.1 Registry的介绍 二、搭建本地私有仓库 2.1首先下载 registry 镜像 2.2在 daemon.json 文件中添加私有镜像仓库地址 2.3运行 registry 容器 2.4Docker容器的重启策略 2.5为镜像打标签 2.6上传到私有仓库 2…

Qt+C++桌面计算器源码

程序示例精选 QtC桌面计算器源码 如需安装运行环境或远程调试&#xff0c;见文章底部个人QQ名片&#xff0c;由专业技术人员远程协助&#xff01; 前言 这篇博客针对<<QtC桌面计算器源码>>编写代码&#xff0c;代码整洁&#xff0c;规则&#xff0c;易读。 学习与…

VUE 开发知识个人记录

以下内容是本人在开发过程中实际验证过的一些固定写法&#xff0c;记录下来&#xff0c;方便自己后续开发过程中直接拷贝粘贴。 1、div 设置宽度自适应文本内容&#xff1a;设置 div 节点的 class 为下面 auto-adjust-text-width .auto-adjust-text-width {display: inline-b…

微签京瓷合作,亮相2023办公行业博览会

武汉&#xff0c;2023年8月8日至8月10日&#xff0c;2023中国现代办公行业年会暨中国智能办公行业博览会在武汉光谷科技会展中心盛大开幕。在这场行业盛会上&#xff0c;微签与京瓷合作打造的OA数字化管理系统重磅亮相&#xff0c;向广大消费者展示了微签在办公设备领域的转型升…