自动驾驶-芯片-DDR
自动驾驶-芯片-DDR参考文献链接
https://mp.weixin.qq.com/s/QYaxxFL9DgaH8w7-oWTLnQ
https://mp.weixin.qq.com/s/piz05ElprV88xnrHFufnvg
https://mp.weixin.qq.com/s/6H_JsumnYUoyIT8_ovBihQ
https://mp.weixin.qq.com/s/O1jMMqKn-fnQp0T35tMP-w
裁员、卖身、大亏损，西半球刮来自动驾驶第一股寒风
「今年开始，行业大环境突变……」
自动驾驶第一股图森未来原董事长陈默在 7 月份接受媒体采访时的一句话，总结了美国上半年的市场情况，也意外成功预言了下半年的行业趋势。
起初是图森在 3 月份突传出售中国业务，团队随之大调整，陈默也从公司出走。
接着，6 月份，特斯拉关闭位于加州圣马特奥的自动驾驶系统部门，同时裁减约 200 个时薪制岗位。没多久，明星公司 Argo AI 也宣布裁减约 150 名员工。
而最近的消息是，上市不到一年的 Aurora，寻求部分或整体将公司打包出售。要知道，在 Aurora 去年的招股书中，还曾意气风发地把自己和 Waymo、Cruise 并列为自动驾驶三分天下的势力......
毫无疑问，事情正在起变化，而背后的原因是行业正在经历一轮资本退潮：
一边是美联储加息背景下资金成本越来越高，另一边是自动驾驶企业越来越止不住的烧钱窟窿，投资方此时选择站在一处，目睹一场泥沙俱下的行业「盛况」。
自动驾驶企业对此无能为力，在技术、商业化均未取得突破性进展之前，它们只能仰赖外部输血。而当这一过程被突然中断，也只能选择断臂求生，甚至将自己打包出售。
「凛冬已至。」一位业内人士感慨道。值得注意的是，这场看似起源于西半球的寒风，也刮到了国内。在前不久，华为创始人任正非率先感受到寒气，并提出要传给每一个人。
2022，注定是自动驾驶行业空前艰难的一年。
01
自动驾驶堪称碎钞机，
Cruise 每天消耗 500 万美元
3 月份，图森未来传出要将亚太业务（主要是中国业务）以 10 亿美元的价格出售。在外界看来，这是一笔为了符合美国监管而迫不得已达成的交易。
自去年 4 月图森 IPO 以来，美国外国投资委员会 (CFIUS) 就一直以图森未来拥有中国区业务为由对其展开调查，并最终在今年 2 月份与图森达成协议，后者需将自动驾驶卡车业务的一些技术监督权移交给美国相关部门，并定期开会并向美国外资投资委员会报告等。
在这种要求下，为了避免纠纷，图森直接出售亚太业务显得合情合理。然而，陈默表示，此举还有另一层原因：缺钱。「出售亚太业务是为了让图森回笼资金，以支撑美国业务发展。」
据了解，图森上市至今「烧钱」速度惊人。以其最新发布的二季度报为例，实现营收 259.4 万美元，录得亏损 1.11 亿美元（折合约为 7.59 亿人民币），其中研发投入 5.93 亿元人民币，如果按每季度 90 天计算，每天消耗超过 600 万元。
作为上市公司，图森按理来说可以在资本市场寻求募资，然而其股价自高峰一落千丈，市值跌了 80%，陈默坦言：「我们不可能在市场上募集到钱。」
和图森一样陷入财务困境的还有明星公司 Argo AI。
这是一家由福特和大众集团投资的自动驾驶技术企业，由谷歌和 Uber 背景的两位技术大牛在 2016 年成立，在仅有 12 名员工的时候就获得福特 10 亿美元投资，而后又在 2020 年，获得大众 26 亿美元投资。
7 月 7 日，Argo AI 宣布裁减约 150 名员工。尽管该公司发言人旋即表示，此次裁员不会减缓自动驾驶系统开发商业服务的计划，但还是让外界感觉到了自动驾驶行业的一丝寒意。
如今，这股寒意已经传导到下半年。日前有消息称，去年底以 SPAC 方式借壳上市的自动驾驶公司 Aurora，将通过裁员、降低福利等方式削减成本，以度过当前入不敷出的困境。
Aurora 首席执行官克里斯·厄姆森（Chris Urmson）表示，正在考虑出售优质资产如旗下激光雷达业务，或将公司整体打包出售给苹果、微软或其他潜在的第一梯队汽车公司。
有媒体报道，自上市至今，Aurora 已经烧掉了大约 2.3 亿美元的现金，其股价也已经从最高的每股 17.77 美元跌落至最近的 1.43 美元，降幅超 90%。
不过和真正的「烧钱大户」Cruise、Waymo 相比，Aurora 显得「稍逊风骚」。
通用汽车最新财报显示，其自动驾驶子公司 Cruise 在第二季度营收 2500 万美元，却亏损了高达 5.4 亿美元，相当于每天 500 万美元。
如果算上 2022 年整个上半年，累计亏损达 9 亿美元，较去年同期多亏损 3 亿美元。
Waymo 也「不遑多让」，Google 母公司 Alphabet 发布的 2022 财年第二季度财报显示，OtherBets（包括自动驾驶汽车部门 Waymo 和生命科学部门 Verily 等）本季度亏损 16.68 亿美元。
虽然没有直接的数据，但根据此前 Waymo 前员工及其他行业内人士估算，Waymo 每年要在各方面花掉 10 亿美元以上。
在过去，自动驾驶公司「烧钱」并不是什么新鲜事，按照以往业内人的说法，这是赢在未来之前必要的投资，不同玩家或依靠母公司，或引入外部融资等方式进行输血，然而现在，事情开始起变化。
02
技术、商业接连受阻，
资本市场「累觉不爱」
据赛博汽车报道，在自动驾驶干线物流赛道，有企业今年见了 80 多家投资机构后，才在融资方面有了点眉目，而不少企业「已经见了一百多家了，还没有着落」。
按照陈默的说法，美联储多次加息，导致资本市场银根收紧，是重要原因之一。
世界首富马斯克也在推特上直言「美联储大幅加息有通缩的风险」，并接连出售约 800 万股特斯拉股票，累计套现 70 亿美元。
今年以来，美股大多数科技公司都受到加息周期影响，市值接连回调，自动驾驶公司下跌表现尤甚，这让原本看好该领域的投资方也变得迟疑，为了保证自己资金链的安全，他们选择少出手甚至不出手。
以 Mobileye 为例，该公司本来准备在今年上市，一度被市场给出 500 亿美元的估值，有望拿下今年美股最大的 IPO。
2022 年 3 月份，英特尔公司宣布已经向美国证劵交易委员会秘密提交 Form S-1 注册声明草案，拟首次公开发行 Mobileye 新股。按照规划，上市筹措的资金将用于加码 Mobileye 新业务。
然而，仅仅过了半年，Mobileye 估值大砍 40％，英特尔预计，该公司最新估值最高为 300 亿美元，远低于最初期望水平。为此，Mobileye 上市的计划便不了了之。
投资网站 PitchBook 的数据显示，2022 年自动驾驶投资已大幅下滑，第二季度对自动驾驶创业公司的投资下降到 9.58 亿美元，在风险投资中的占比不到 10%。
而自动驾驶公司今年在技术和商业化上不及预期的情况，更加重了资本市场的观望情绪。
以图森未来为例，4 月份，其在美国部署测试的自动驾驶卡车突然左转横切，撞上一个水泥路障。这起交通事故经《华尔街日报》报道后，很快引来当地监管部门对于自动驾驶安全风险的关注。
为此，图森未来停止了整个车队的测试，并进行独立调查，图森未来联合创始人兼 CEO 侯晓迪还表示，在确保车队可以安全运营之前，图森未来不会进行商业化。
Cruise 也在今年经历了一个「多事之秋」。
5 月份，该公司遭内部员工向加州监管机构匿名举报，Cruise 存在故意隐瞒涉及车辆和业务具有潜在破坏性的问题，如安全报告系统仍处于混乱状态，这一消息让外界哗然。
6 月，Cruise 再次卷入风波，先后发生两起交通事故。起先是其 Robotaxi 在左转转弯处与一辆迎面而来的丰田普锐斯发生碰撞，导致两人受轻伤。
为此，Cruise 召回并更新了 80 辆 Robotaxi 的软件。
而后在同一个月内，十几辆 Cruise 自动驾驶车辆无故停在十字路口，导致交通瘫痪了一个多小时，直到人类驾驶员将其开走之后，事情才得到解决。
事后，Cruise 回应称，是「技术问题」导致了上述情况发生。
让公众对自动驾驶技术越发不信任的还包括特斯拉。
7 月 6 日，一辆 2015 款的特斯拉 Model S 从佛罗里达州盖恩斯维尔以南的 75 号州际公路上驶入高速公路休息区停车场，径直撞上了一辆停在那里的牵引车，尽管还未证实在当时自动驾驶功能是否开启，但巧合的是，这次撞上的又是一辆「白车」。
更早之前，根据美国国家公路交通安全管理局（NHTSA）的数据，截至 2022 年 5 月的 10 个月中，有 200 多起车祸都与特斯拉的 Autopilot 软件有关。
特斯拉为此被加州机动车辆管理局（DMV）和美国国家公路交通安全管理局（NHTSA）认定其技术存在虚假宣传。加州参议院也已通过法案，禁止在智能驾驶的广告里包含「自动驾驶」等词汇。
自动驾驶公司在技术上还没达到成熟，已经开始让之前的投资人变得焦急，而直到现在，商业化上也一直没有大的起色，更是让这些人「封住了口袋」。
Aurora 最近宣布，其最快商用的自动驾驶卡车上路的时间，仍旧需要推迟一年，至 2024 年。
该公司创始人兼 CEO 克里斯·厄姆森解释称，背后原因在于 OEM 进展缓慢。由于商业化不及预期，他还表示，「预计未来六个月内，不会出现足够规模的传统融资机会来为公司续命。」
图森未来面临相同的窘境，此前，其与美国第三大卡车制造商 Navistar 达成合作，实现 L4 自动驾驶卡车的量产。
然而随着时间的推移，Navistar 一直在推迟图森未来量产 L4 自动驾驶卡车的进程，从预定的 2024 年，到 2025 年，2026 年……
陈默曾在一次媒体访谈中直接吐槽 Navistar 的自动驾驶量产进度太慢，尤其在欧美市场。「我们原本给投资人的承诺是，图森未来在 2024 年 Q3 实现量产……」
Waymo 和 Cruise 倒是在商业化上传来了一些好消息。
3 月 1 日，美国加州公共事业委员会 (CPUC) 向 Alphabet 和通用汽车发放了自动驾驶客运服务的许可证，允许两家公司旗下的自动驾驶公司在旧金山及周边提供收费客运服务。
此前，Waymo 和 Cruise 在加州只被允许在测试的基础上提供有限载人客运，不准收取费用。
不过这只能算是商业化初期，毕竟 Robotaxi 当前铺开的城市并不多，以 Waymo 为例，目前仅在凤凰城的东谷、凤凰城市中心以及旧金山三地向公众提供出行服务。
此外，Waymo 目前的商业模式还略显鸡肋。根据 CNBC 记者在 2022 年 1 月的体验，Waymo 在凤凰城的 Robotaxi 跑 5 英里需用时 14 分钟，每分钟收费 1 美元。
而跑同样里程的 Uber 不仅快得多，每分钟收费更仅有 0.4 美元。
「旧金山 Robotaxi 的服务价格将定在合理且有竞争力的区间，与此同时，有限用户免费测试/收费商业运作的模式将在其他地方铺开。」Waymo 发言人在一场会议上表达了自己的乐观。
相较而言，更多的人对于自动驾驶却没有这么乐观，他们眼下更关心的是，漫漫「烧钱」之路还要延续多久？
03
烧钱游戏进入倒计时：
没有金山倚靠，或需考虑卖身
截至 2022 年 6 月 30 日，图森未来资产负债表上仅有 11.6 亿美元的现金，而按照其给出的 2022 年调整后的 EBITDA 亏损 3.6-3.8 亿美元来看，在不融资的情况，公司现金仅够支撑 3 年左右。
而现在，图森仍没有开启自主造血的进程，资料显示，其卡车预订总数达到 7485 辆，但至今仍无一辆进入交付量产阶段。
在招股书中，图森曾这样计算，一辆卡车一年可实现 6 万美元营收，当运营数量过 5000 辆时，公司将达到收支平衡，之后才能逐渐实现盈利。
Aurora 则直接给出了时间表，预计 2027 年之前都无法实现盈亏平衡。而 Cruise 首席执行官 Kyle Vogt 在今年 9 月 12 日一次会议上表示，公司目标是到 2025 年时获得 10 亿美元的营收，相当于目前来自通用的年投资水平的一半。
以上这些预示着自动驾驶公司还得至少烧 3-5 年的钱，并且最终还不一定能跑出来，有投资方等不及了。
2022 年 3 月，私募股权巨头软银集团旗下愿景基金向通用汽车出售其持有的 Crusie 股份，价值 21 亿美元。
Vogt 坦言，自动驾驶行业已经从「极端乐观」转向「极度悲观」。这意味着，随着资本退潮，自动驾驶公司已经到了该谋求出路的时刻。
强如 Waymo、Crusie 这类背靠「金山」的企业，似乎还不用特别担心。Cruise 目前还拥有 37 亿美元的资金，并获得了通用汽车金融部门约 50 亿美元的信贷额度。
在一季度的财报发布会上，通用计划 2022 年向 Cruise 支出约 20 亿美元，这比以往任何时候都要高。
而在接下来，通用还承诺将持续为 Cruise 提供资金，确保后者在市场上占据领导地位。
据了解，在接收软银集团转让的股份后，通用 3 月份就向 Cruise 追加了 13.5 亿美元投资。
Waymo 也有母公司 Alphabet 的供养。
在 2020 年之前，谷歌/Alphabet 已经为 Waymo 持续输血 11 年，不过由于投入过于巨大，最终在当年年初首次开始引进外部融资：3 月份筹集了 22.5 亿美元，5 月份筹集了 7.5 亿美元，最终在 73 天内累计完成 30 亿美元的融资。
得益于 Alphabet 强大的品牌效应，Waymo 在 2021 年 6 月，又宣布完成一轮 25 亿美元的投资。
与此同时，知情人士透露，Waymo 还讨论了最终公开上市的计划，以减轻其母公司的资金压力。
不过按当前资本市场的表现，Alphabet 短期内大概率不会推动 Waymo 的 IPO，取而代之的是，将继续用公司的利润去补贴支撑 Waymo 的巨额亏损。
最危险的是没有「靠山」的自动驾驶公司。
正如前文所说，Aurora 已经尝试性迈出第一步：考虑卖身主机厂或科技巨头。
该公司首席执行官克里斯·厄姆森（Chris Urmson）表示，可以将公司私有化，前提是找到一个能够提供 15 亿美元融资的合作伙伴。「鉴于目前的股价，我们应该成为任何希望拥有自动驾驶技术公司的有吸引力的目标。」
戏剧性的是，Aurora 原本也是通过收购自动驾驶公司做大业务的，而不料如今却走向了寻求被收购的道路。
2020 年 12 月，还是初创公司的 Aurora 以 40 亿美元的价格，近乎「蛇吞象」式地收购 Uber 旗下自动驾驶部门 ATG，彼时正值 Uber 面临巨大运营成本压力，其 ATG 部门在 2019 年造成了约 5 亿美元的亏损。
在出售自动驾驶业务后，Uber 净亏损确实得到一定改善。而如今历史重演，Aurora 面临整体出售的暗淡时刻。
尽管如此，出售仍是比破产倒闭更体面的结局，如果能卖给主机厂，甚至可以称得上是个不错的归宿。
随着行业智能化趋势加深，自动驾驶被视作汽车的「灵魂」。
通用、福特、大众等传统主机厂都在这方面加码投入，而对于那些还没开始布局的 OEM，要参与这股浪潮，比较可行的选项便是收购一家自动驾驶公司，而这正给了像 Aurora 这样苦苦挣扎的企业一丝希望。
「网约车第一股」Lyft 旗下自动驾驶部门 Level 5 就是这样完成「和平交接」的。
成立于 2017 年夏天，Level 5 同样辉煌过，在 2018 年已经获得加州公共道路测试许可，并取得了不错的研发、测试数据积累。
然而随着 2020 年疫情突降，母公司 Lyft 遭受巨额业务亏损，无法支撑 Level 5 烧钱继续自动驾驶的梦想。
2021 年 4 月，Lyft 将 Level 5 以 5.5 亿美金的价格出售给丰田旗下的子公司 Woven Planet Holdings，至此每年至少节省 1 亿美元日常开支。
如今，陷入僵局的自动驾驶企业也在谋求卖身主机厂的机会，后者雄厚的资金可以帮助它们渡过 2022 这个难熬的「寒冬」。
日前，就有消息称，吉利控股集团向图森未来美国总部发出收购要约，拟收购图森未来控股亚太地区业务的全部股份，具体收购细节暂未公布。
04
自动驾驶寒风开始刮到东半球：
活下去，才有未来
把视线拉回国内，中国的自动驾驶企业看起来还没遭受到大洋彼岸同行那样的危机，但若仔细看，却也经历了一些困难。
2021 年，境外上市监管环境收紧，小马智行、智加科技等 IPO 计划折戟，融资节奏随之打乱，导致业务发生重大调整，人员也随之流动。
典型如小马智行卡车部门去年遭遇动荡，该公司原自动驾驶技术总负责人潘震皓、国内自动驾驶规划与控制组负责人孙浩文以及负责战略合作和融资的副总裁赵睿璇纷纷出走，各自成立三家卡车自动驾驶公司。
时间来到 2022 年，小马智行卡车业务重整旗鼓：7 月 28 日，小马智行与三一重卡成立合资公司，计划在 2022 年内开始小规模量产交付自动驾驶卡车。
而这背后，是小马智行再次获得输血：3 月 7 日，小马智行宣布完成 D 轮融资的首次交割。
小马智行 CFO 劳伦斯·斯泰恩表示：「公司财务状况十分稳健，为小马智行未来几年的发展提供坚实基础，直至我们开启大规模商业化的进程。」
继终止以 SPAC 方式在纽交所上市的协议后，智加科技同样传出有变。
今年 8 月份，虎嗅一篇题为《图森们的「并购式」跑路与谎言》文章写道，有消息人士爆料，干线物流平台满帮集团或将收购自动驾驶卡车企业智加科技。
不过，这一消息尚未得到证实。
从融资上来看，智加科技在 2021 年创下 4.2 亿美元的融资记录，为干线物流领域最大。
而据最近消息，2022 年 8 月 16 日，智加科技联合挚途科技完成面向荣庆物流的 100 台自动驾驶重卡订单的首批交付。
从小马智行和智加科技来看，国内的玩家似乎平稳度过了资金难关。然而，没过多久，8 月 22 日，华为创始人任正非在内部喊出「寒气论」，让自动驾驶行业再度顿觉凉意。
按照任正非的说法，未来十年是一个非常痛苦的历史时期，全球经济会持续衰退。为了应对挑战，华为将收缩或关闭未来几年内不能产生价值和利润的业务。
巧的是，智能驾驶便是这样的业务。
在 7 月 7 日的一场行业论坛上，华为常务董事、终端 BG CEO、智能汽车解决方案 BU CEO 余承东透露，汽车业务是华为目前唯一亏损的业务。
他表示，华为在该业务上一年投入十几亿美元，直接投入 7000 人，间接投入 1 万人。而其中，又以智能驾驶占大头。「百分之七八十都投入在智能驾驶辅助领域。」
而目前，根据任正非的要求，对于智能汽车解决方案，华为将减少科研预算，加强商业闭环研发。
华为这样年营收 6、7000 亿的公司尚且在收紧对自动驾驶业务的投入，对于那些还没有实现商业化、完成自主造血的企业，未来资金压力可见一斑。
回到当下，美国自动驾驶企业的遭遇已经为国内同行敲响了警钟，如果没有「大树」背靠，按照该领域公司一年一次融资的烧钱节奏，现在还没有融到钱的企业，可以说是十分危险了。
自动驾驶不仅是技术的比拼，更是一场时间的游戏，能抵达终点的不一定是跑得最快的那个，现在活下去才有资格谈未来。
CUDA 矩阵乘法终极优化
单精度矩阵乘法（SGEMM）几乎是每一位学习 CUDA 的同学绕不开的案例，这个经典的计算密集型案例可以很好地展示 GPU 编程中常用的优化技巧。
CUDA 矩阵乘法优化手段详解
Naive 实现的分析：到底差在哪里？
笔者面试过不少具有 CUDA 编程经验的校招同学，当提问使用 CUDA 编写一个 SGEMM Kernel 的时候，通常会获得这么一个答案：
__global__ void matrixMul(const float *A, const float *B, float *C,
int M, int N, int K) {
int tx = blockIdx.x * blockDim.x + threadIdx.x;
int ty = blockIdx.y * blockDim.y + threadIdx.y;
if(ty < M && tx < N) {
float c = 0;
for(int i = 0; i < K; ++i){
c += A * B;
}
C = c;
}
}
这样一个 Naive 的 Kernel 当然不是笔者所期待的，因为这个 Kernel 的性能基本可以断定连 cublas 的 1/10 都不到，显然不符合我们追求高性能的需求。那么这个 Naive 的实现究竟差在哪呢？分析代码我们可以看到，计算一次 FMA（乘累加）之前需要读一次 A 和读一次 B，众所周知，读取 Global Memory 的代价很大，通常都需要几百个 cycle（时钟周期），而计算一次 FMA 通常只需要几个 cycle，大量的时间被花费在了访存上。也许有思维活络的同学立马想到，可以将 A 和 B 矩阵先搬运到 Shared Memory（SM 中低延迟的 on-chip memory，block 内线程共享，附 NVIDIA GPU 内存结构图）中降低访存的开销，这的确是一个很好的思路，但是这只能将访存代价从几百 cycle 降低到几十 cycle，并不改变问题的本质。问题的关键在于主体循环由两条 Load 指令与一条 FMA 指令构成，计算指令只占总体的 1/3，计算访存比过低，最终导致了访存延迟不能被隐藏，从而性能不理想。
让我们打开思路，若一个 thread 并不只计算一个结果，而是计算 4x4 个结果，并且使用 Shared Memory 优化，Hot Loop 会是什么样呢，伪代码如下所示：
float c = {{0}};
float a_reg;
float b_reg;
for(int i = 0; i < K; i += TILE_K){
__syncthreads();
// transfer tile from global mem to shared mem
load_gmem_tile_to_smem(A, i, smemA);
load_gmem_tile_to_smem(B, i, smemB);
__syncthreads();
#pragma unroll
for(int j = 0; j < TILE_K; ++j) {
// load tile from shared mem to register
load_smem_tile_to_reg(smemA, j, a_reg);
load_smem_tile_to_reg(smemB, j, b_reg);
// compute matrix multiply accumulate 4x4
mma4x4(a_reg, b_reg, c)；
}
}
分析可以得出从 smemA 读取到寄存器 a_reg 中，需要进行 4 次访存操作，B 同理，那么主体的计算访存指令比例变成了 16/8，相对于之前的情况，计算指令的占比大大提高了。足够大的计算访存比能提升计算单元的利用率，并能起到隐藏访存延迟的作用。我们可以进一步提升计算访存比，从而使得 kernel 的性能接近理论峰值。
矩阵分块与资源分配
显然我们不能只使用一个 block 计算一个超大矩阵，这样会造成大量 SM（Streaming Multiprocessor）的闲置浪费，这就需要对矩阵进行分块计算，如下图所示：
不同的分块大小在不同 shape 的矩阵乘法应用上性能各有优劣，本文选取 128x128 的分块举例。
从上一小节我们可以看到，提升计算访存比有很大的好处，那么计算访存比可以无限提升吗，答案是否定的。因为要提升计算访存比，单个 thread 就需要计算一个更大的块，这就需要更多的寄存器，但寄存器的个数是有限的。以 Turing 架构的 GPU 为例，单个 SM 的寄存器总量为 65536，因为指令编码的限制，单个 thread 能使用的最大寄存器个数为 255，并且寄存器个数并不是用得越多越好。这里需要引入一个 Occupancy（占用率）的概念，Occupancy 是指每个 SM 中活动线程束（Warp）数量与最大并发线程束数量的比值，高的 Occupancy 不一定意味着高性能，但可以通过切换执行 Warp 来起到一定隐藏延迟的作用。而每个 SM 中的 Active Warp 数量，取决于 block 使用的资源数量，具体为每个线程使用的寄存器个数与 Shared Memory 用量。Occupany可通过 CUDA Toolkit 中提供的 CUDA_Occupancy_Calculator.xls 工具获得。
考虑一个 block 计算 128x128 的分块，若每个线程计算 128 个结果，需要的 block size 为 128，单个线程需要 128 个寄存器储存计算结果，加上所需的 Gmem to Smem，Smem to Reg 等一些所需的寄存器，大概共需要至少 180 多个，计算 Occupany 可知此时的 Active Warp 数只有 8，Occupany 为 25%；若设置 block size 为 256，则每个线程仅需计算 64 个结果，调整寄存器和 Shared Memory 的使用量并观察 Occupany，可知若每个线程只使用 128 个寄存器，block 内的 Shared Memory 使用量限制在 32K，Active Warp 数可以达到 16，是一个更优的选择：
并且此时的配置计算访存比可以达到 64/4（使用向量读取），已经足够隐藏访存延迟。
极致的访存优化
通常情况下，在选取了合适的 block 资源配置，利用 Shared Memory 降低访存延迟，做好循环展开之后，SGEMM Kernel 的性能已经能达到一个不错的水平（80% cublas），但这并不是我们旅程的终点。首先，我们可以使用向量读取指令LDS.128优化 Shared Memory 访问（对应 float4 数据类型），这能大幅减少访存指令的数量，进一步提升计算访存比，由此我们需要将 A 矩阵存入 smemA 之前做一次转置：
同时，我们的 kernel 为 256 个线程计算 128x128 的分块，为了能够合并访问 Shared Memory，我们将 256 个线程划为二维，令：
int tx = threadIdx.x % 16;
int ty = threadIdx.x / 16;
并按照如下方式向量读取 Shared Memory 中的数据：
最终单个线程计算 2x2 个 4x4 的结果，结果布局如图所示：
并且通过 micro benchmark 可以探测出，Turing(Tesla T4) 的 Global Memory 的访存延迟约 300 cycle，Shared Memory 的访存延迟在约 30 cycle，需要充分利用 Prefetch 的思想，隐藏 Global Memory 读入中间寄存器、将来自 Global Memory 的数据块写入 Shared Memory、从 Shared Memory 中读出数据块的访存延迟，以免计算单元因为 stall 而空闲太久，最终的伪代码如下所示：
#define TILE_K 16
__shared__ float4 smemA;
__shared__ float4 smemB;
float4 c = {{make_float4(0.f, 0.f, 0.f, 0.f)}};
float4 ldg_a_reg;
float4 ldg_b_reg;
float4 a_reg;
float4 b_reg;
// transfer first tile from global mem to shared mem
load_gmem_tile_to_reg(A, 0, ldg_a_reg);
load_gmem_tile_to_reg(B, 0, ldg_b_reg);
store_reg_to_smem_tile_transpose(ldg_a_reg, 0, smemA);
store_reg_to_smem_tile(ldg_b_reg, 0, smemB);
__syncthreads();
// load first tile from shared mem to register
load_smem_tile_to_reg(smemA, 0, a_reg);
load_smem_tile_to_reg(smemB, 0, b_reg);
int write_stage_idx = 1; //ping pong switch
do {
i += TILE_K;
// load next tile from global mem
load_gmem_tile_to_reg(A, i, ldg_a_reg);
load_gmem_tile_to_reg(B, i, ldg_b_reg);
int load_stage_idx = write_stage_idx ^ 1;
#pragma unroll
for(int j = 0; j < TILE_K - 1; ++j) {
// load next tile from shared mem to register
load_smem_tile_to_reg(smemA, j + 1, a_reg[(j + 1) % 2]);
load_smem_tile_to_reg(smemB, j + 1, b_reg[(j + 1) % 2]);
// compute matrix multiply accumulate 8x8
mma8x8(a_reg, b_reg, c)；
}
if(i < K) {
// store next tile to shared mem
store_reg_to_smem_tile_transpose(ldg_a_reg, 0, smemA);
store_reg_to_smem_tile(ldg_b_reg, 0, smemB);
// use double buffer, only need one sync
__syncthreads();
// switch
write_stage_idx ^= 1;
}
// load first tile from shared mem to register of next iter
load_smem_tile_to_reg(smemA, 0, a_reg);
load_smem_tile_to_reg(smemB, 0, b_reg);
// compute last tile mma 8x8
mma8x8(a_reg, b_reg, c)；
} while (i < K);
store_c(c, C);
注：此处偷懒假设了 M、N、K 都是 4 的倍数，若非 4 的倍数则 Global Memory 不能使用 float4 进行读取，结果也不能用 float4 进行写回，而且为了合并写回，需要通过 Shared Memory 交换 warp 内的结果，保证每个 warp 执行一条 Store 指令能够写回一片连续的内存空间。
至此我们获得了一个充分优化的 SGEMM Kernel。另外 Ampere GPU 新增了LDGSTS指令，数据块从 Global Memory 到 Shared Memory 的过程不需要经过中间寄存器，可以进一步的优化 SGEMM 的性能。
性能对比
为了避免 cublas 选取到 split K 的 Kernel，我们将 K 固定为 1024，取 M, N = 2048, 4096, 8192 和 16384 作为测试用例，对比了上述 SGEMM Kernel 与 cublas 的性能（测试 GPU 为 Tesla T4，锁定核心频率为 1100)：
可以看到所实现的 SGEMM Kernel 达到了 cublas 平均 97.5% 的性能。
超越 cublas：使用 SASS 调优 Kernel
到这里，可能有同学依然有一个疑问，我们似乎把所有能想到的优化手段都用上了，为什么写出来的 CUDA C Kernel 依然离 cublas 有一定的差距，答案是 cublas 所使用的 kernel 中有一大部分并不是通过 nvcc 编译的 CUDA Kernel，而是使用 NVIDIA GPU 的汇编语言（Shader Assembly，简称 SASS）编写的深度调优版本。
尽管 nvcc 编译器在不断的进步，特别是 CUDA 11 中的 nvcc，所编译的 Kernel 与手工汇编优化版本之间的差距已大幅缩小，但依旧无法完全避免寄存器 Bank conflict 的影响以及充分利用寄存器的 Reuse Cache（这两个概念下面会进行详细的介绍），使得差距依旧存在。即使 PTX 这样的伪汇编语言，也无法精确控制寄存器的分配，和 CUDA C 面临着一样的困境。
所以为了充分挖掘 GPU 的性能极限，需要对 GPU 指令和寄存器进行精确控制，就必须交由 GPU 原生汇编语言 SASS 完成。这方面已经有了很多研究，如出自 Citadel 的深入研究 NV GPU 架构的 Dissecting the NVidia XXX GPU architecture via microbenchmarking 系列论文，这一系列文章对底层架构做了系统的测试、分析和总结，虽然其中某些结论可能并不准确，但总体来讲有很高的参考价值。同时催生了不少开源汇编器如 KeplerAs、maxas（最成熟，影响深远）、turingas 和 CuAssembler 等一系列开源 SASS 汇编器，使得使用 SASS 编写高性能 Kernel 变成了可能。
寄存器 Bank conflict
我们知道 Shared Memory 有 Bank conflict，而寄存器的 Bank conflict 也是类似的概念。NVIDIA GPU 每个 SM 有独立的 Register File，而 Register File 被分为若干个 Bank，以 Maxwell 为例，若一条指令所需的源寄存器有 2 个以上来自于同一 Bank，则会产生 conflict，指令会相当于重发射，浪费一个 cycle。Maxwell/Pascal 的 Register File 的 Bank 数为 4，寄存器的id%4即为该寄存器的所属 bank（如 R0 属于 Bank 0，R5 属于 Bank 1），FFMA R1, R0, R4, R1这样的指令就会产生寄存器 Bank conflict。而 Turing 架构做了改进，Register File 被分为 2 个 Bank，每个 Bank 有 2 个 Port，若非三个源寄存器 id 同奇偶则不会产生冲突，大大缓解了寄存器 Bank conflict。
maxas 中的 Maxwell SGEMM SASS Kernel 为了缓解寄存器 Bank conflict，就对参与 FFMA 计算的寄存器做了精巧的分配（参考 maxas 的 SGEMM 文档），如下图所示：
经过对 C 的巧妙排布，寄存器 Bank conflict 大大减少，但依然无法完全避免（如上图中黑框标识的部分，A/B 所使用的寄存器会产生 Bank conflict），这部分冲突就需要用到寄存器 Reuse 来消除。
Register Reuse
寄存器 Reuse 是 NVIDIA 为了缓解寄存器 Bank conflict 的问题，在 Maxwell 开始引入的一种机制，NVIDIA 在读取指令操作数的 Collector 单元加入了寄存器的 Reuse Cache。Reuse Cache 是只读的，指令获取 Operand 是否通过此 Cache 由该指令的 control code（maxas 的 control code wiki中有详细的介绍）所指定，使用 cuobjdump 反汇编一些 Kernel 可以发现一些寄存器后有 .reuse的 flag，即表示该寄存器从 Reuse Cache 而非 Register File 中取值，从而消除寄存器 Bank conflict：
# Maxwell GPU但是使用 .reuse需要满足一定条件（寄存器将被改写前不能设置 .reuse），胡乱设置 reuse flag 会有可能获取的是历史值，造成计算错误，根据笔者的理解，.reuse 更像是使该寄存器的值在 Reuse Cache 中 hold 住的标识。nvcc 编译 CUDA Kernel 也会使用 Reuse Cache 去规避一些寄存器 Bank conflict，但是因为寄存器分配及指令排布的原因，Reuse 的利用率并不高，反汇编我们刚才写的 SGEMM Kernel，对主循环的所有 FFMA 指令做个统计，可以发现 Reuse Cache 仅达到 20% 左右，而 maxas 的 SASS Kernel 通过设计使得 Reuse 的利用率可以达到 49%。
最终通过 SASS 精细调优的 SGEMM Kernel 的性能可以全面超越 cublas，感兴趣的同学们可以自行编译 maxas 中的 SGEMM Kernel 在 Maxwell 或者 Pascal GPU 上进行测试。最后，虽然使用 SASS 能充分挖掘 GPU 的性能，但面临有三大问题：1. 第三方 NV GPU 汇编器依赖于对 GPU 架构的逆向研究，可能因为没有探究到全部的硬件底层细节而存在未知的 BUG；2. 汇编 Kernel 难于开发，更难于调试；3. NV 每一代 GPU 的 ISA（指令集）都不尽相同，需要不断开发对应的汇编器和汇编 Kernel。正因为这几大问题的存在，使得使用 SASS 编写 Kernel 是个费时费力的工作，除非有追求极致性能的需求，否则不建议轻易尝试。
GEMM 的延伸：优化卷积运算
我们都知道优化卷积运算可以通过 im2col 将卷积映射为矩阵乘法来实现，对于上述 SGEMM Kernel，只需要将 Global Memory 的数据搬运到 Shared Memory 这一过程稍作修改，由对应位置的映射变为 im2col 映射，SGEMM Kernel 就摇身一变成为了计算 Conv 的 Kernel，这即是 cudnn 卷积运算的 Implicit Gemm 算法。而在 im2col 过程中，若直接计算指针的偏移量的话，会引入大量的整数除法和取余运算，这是一笔不小的开销，所以可以将地址的偏移量在 host 端预先计算好，作为 param 传入 kernel 中，则可以在需要时从常量内存中读取，避免整数除法和取余，实现 Implicit Precomp Gemm。
总结
本文详细介绍了如何编写一个高效率的 CUDA SGEMM Kernel，并且介绍了使用 SASS 编程这一极限优化性能的手段，并稍稍延伸展开了通过 Implicit Gemm 优化卷积运算的思路，希望可以给予有志于极致挖掘硬件性能的同学们一定的启发
RISC-V将赢得下一轮架构之争？
在全球市场上，芯片指令集呈现双寡头格局，基于X86和ARM架构的处理器长期占据绝大多数市场份额，X86架构在PC及服务器市场一家独大，移动市场则由ARM架构一统江湖。
在这样一个格局中，中下游厂商大多只能在这二者之间选择，但是ARM授权费用昂贵，传统X86的授权又过于复杂，业界一直期待在CPU架构领域能有更多选择。
随着AIoT时代的到来，RISC-V架构开放、灵活、模块化，特别适合满足AIoT市场场景碎片化、差异化的市场需求，产业界普遍认为它有望成为下一代广泛应用的处理器架构。Semico Research预测，到2025年，RISC-V市场规模将超10亿美元。
近日，行业大牛Jim Keller在一次公开演讲中也直言，未来是属于RISC-V的。
Jim keller认为，RISC-V的开放式生态让其拥有广阔的前景和巨大潜力，相比X86和ARM架构，RISC-V虽然在性能、生态等问题上存在不足，但凭借开放式标准，RISC-V的发展速度远超X86及ARM。
RISC-V采用的开放标准让各个厂商可以在架构基础上进行大幅度的修改和定制，不需要担心后续会因为各种原因而失去授权或影响到产品的生产使用。同时，RISC-V存在架构优势，使其可以同时适应极低功耗及高性能处理器的设计需求，这一点也是X86和ARM无法做到的。
当Jim keller看好RISC-V时，这无疑为其平添了几分可能性。对Jim keller有了解的读者应该知道，Jim keller过去创造过许多成功的处理器方案，几乎每一次跳槽，Jim keller都能刷新对应处理器的行业上限。
“硅仙人”的传奇履历
在其20多年的工作经历中，Jim keller先后在DEC、AMD、博通、苹果、特斯拉、Intel等公司任职，一次又一次创造了被认为是芯片行业中，至关重要的领先部件。
上世纪九十年代，Jim Keller在DEC工作，涉足了Alpha处理器项目的设计，产品包括21164与22164。Jim Keller的本段经历使其深入了解RISC的设计理念，并通过Alpha架构，奠定了架构设计能力。Alpha架构的设计理念深度影响了后续AMD、Intel等芯片厂商的架构设计路线。
1998年，Jim keller从DEC离职，加盟AMD，主导了K7/K8/K12处理器的研发，其中K8处理器让AMD第一次拥有了对抗Intel的资本，X86架构不再是Intel一家独大，同时推动AMD的64位X86架构处理器落地，使得AMD在技术路线上第一次超越Intel。
1999年，Jim Keller加入SiByte，研制MIPS的网络处理器。SiByte后被博通收购，Jim Keller担任首席架构师，直到2004年离职，加入PASemi任工程副总裁，转型研究RISC产品，使MIPS的构建速度可与专用ASIC相媲美。
PASemi被苹果收购后，Jim Keller开始主导苹果的A系列处理器研发计划，为苹果A芯片领先全球，奠定了架构基础。
完成A系列处理器的设计后，Jim keller再次跳槽，重返AMD。自从Jim keller离开AMD后，AMD后续的Bulldozer系列处理器一度被称为最糟糕的X86处理器，AMD依靠K8处理器抢到的市场份额几乎完全丢失，甚至一度濒临破产。
Jim keller回归之后，将Bulldozer系列的后续研发计划全部推倒，重新设计了一个全新的处理器架构，这就是大家都不陌生的Zen架构。Zen架构让AMD一举翻身，在性能排行榜上压制Intel数年，甚至一度在处理器出货量上超越Intel，基于Zen系列，AMD后续还拓展了EYPC等系列处理器，开始全方位与Intel竞争。
第二次拯救AMD于水火之中后，Jim Keller于2016年加入特斯拉担任Autopilot部门的负责人，为其打造自动驾驶AI芯片，并在正式落地量产前的2018年离职。
2018年，Jim Keller来到Intel，担任高级副总裁、系统架构师，受限于保密协议，Jim keller负责的项目没有曝光。但从Intel进展来看，2019年5月，Intel发布全新的ice lake架构。据悉，Intel 12代酷睿所使用的大小核设计就由Jim Keller主导开发，而12代酷睿的诞生被誉为Intel的翻身一战，再次压制了AMD成为x86架构的新王者。
回顾Jim keller的经历，几乎每一款经典处理器都改写了行业上限，从x86到ARM，似乎就没有他不了解的处理器架构。
然而，同样是两年后，Jim keller从Intel离职，这次他选择了创业，担任Tenstorrent的联合创始人及CTO，主攻RISC-V。
也正因如此，当这位曾在 x86 和 ARM 架构芯片领域取得过较高成就的硅谷传奇工程师，选择调转航向，亲自下场押注RISC-V架构，不免让业界对RISC-V多了几分信心和慎重。
卡位RISC-V赛道
中国科学院软件研究所总工程师武延军表示，从2010年发布到现在，RISC-V指令集只走过12年的时间，但它在这期间取得的成就，远远超过同时期的X86和ARM。
RISC-V开源ISA于2016年首次推出，但最初的内核仅适用于微控制器和一些基本的片上系统设计。然而，经过几年的发展，众多芯片开发商已经针对云数据中心、AI 工作负载和高级存储应用等多领域创建了设计，其生态获得了较大的丰富，这足以证明RISC-V的实力。
前不久，RISC-V基金会CEO Calista Redmond表示，RISC-V基金会已经在全球70多个国家吸引了超过3100个会员加入，会员数量在2021年暴增了134%。RISC-V产业在过去几年里取得的初步成果，搭载RISC-V内核的芯片出货量已经突破了100亿。
ARM用了17年才完成这个目标，而RISC-V基金会成立至今仅7年。由此可见，RISC-V这个新兴架构正在吸引着全球众多参与者躬身其中，纷纷卡位RISC-V赛道。
从国外厂商动态来看，苹果积极转向RISC-V阵营，Intel业务也延伸至此，包括Jim Keller坐镇Tenstorrent、Microchip、瑞萨电子、Nordic、ST等国际厂商都在积极布局，开发基于RISC-V指令集的处理器。
同时，还有一众RISC-V IP公司深耕于此，包括SiFive、Codasip、imagination、MIPS等，其中值得注意的是，SiFive IP芯片累计出货量超10亿；搭载Codasip的IP芯片出货量已达20亿颗；作为业内成熟的IP供应商，Imagination进入RISC-V市场再次突显了RISC-V架构的扩展特性，为RISC-V提供了更大的发展动力。
除了欧美国家之外，RISC-V在印度、巴基斯坦等国家也在被广泛采用，行业厂商积极布局。
能够看到，国外厂商在RISC-V上的布局在逐渐呈现，除了芯片设计、IP厂商之外，包括IAR、verify、Imperas、Embecosm等工具链、安全软件算法、模拟器和编译器等产业链企业的支持，都在加速RISC-V架构的发展和成熟。
RISC-V，在国内市场掀起风暴
中国RISC-V市场发展迅猛，凭借国内庞大的半导体市场，中国企业可以迅速对RISC-V芯片进行迭代更新，并且依据用户使用反馈进行优化，这些都是国外初创企业难以做到的。全球100亿颗RISC-V核心出货量中，预计中国占比超过1/3。
中国社区作为工作小组和研发伙伴，一直是RISC-V最强有力的贡献者之一。在RISC-V基金会初期的19家高级会员企业中，有整整12家来自中国，华为、阿里巴巴、紫光展锐、中兴等企业都是早期高级会员，对RISC-V的发展有着举足轻重的作用。
RISC-V对于中国的半导体企业来说是一个难得的机会，因为RISC-V起步晚所以大家都还处在一个起跑线上，而且开放式标准也让国内企业可以更安心的发展RISC-V，不用担心未来某一天突然失去授权。
在日前举办的RISC-V峰会上，倪光南院士在致辞中表示：“目前，主流CPU市场仍被X86、ARM架构垄断，但新兴的开源指令集RISC-V将为我国芯片产业发展提供新机遇，如果抓住机会，就有可能在CPU核心技术上掌握主动权。”
在此机遇和趋势下，国内涌现出一批RISC-V玩家，包括华为海思、全志科技、兆易创新、乐鑫科技、北京君正、中微半导体、汇顶科技、凌思微电子、先楫半导体、华米科技、沁恒微电子、芯晟科技、爱普特微电子、晶视智能、启英泰伦、方寸微电子、中科蓝汛、航顺芯片、飞思灵微电子、博流智能以及中科院计算所等在内的芯片厂商和科研机构都相继推出了基于RISC-V架构的SoC、MCU、DSP芯片等产品。
此外，阿里平头哥、芯来科技、赛昉科技、睿思芯科等可以提供基于RISC-V架构的处理器IP、编译器、工具链等产品。据悉，阿里平头哥玄铁系列的RISC-V芯片已经供应一百五十多个企业和客户，超过500个授权使用，出货量已经达到25亿颗；芯来科技RISC-V CPU IP产品线覆盖从低功耗到高性能的各种应用需求，正式授权客户已超过100家...
一些列动态和成果，呈现出国内RISC-V芯片行业上下游产业链的发展速度和全面布局态势。能够看到，在这些中国技术人和企业的努力下，RISC-V处理器的可用性及覆盖面大大提升，使其从一个象牙塔的基础架构，真正走向市场化发展。
国产RISC-V的优势和挑战
随着半导体市场进入一个新的阶段，中国半导体行业需要抓住一切机遇，寻找突破科技封锁的方法，RISC-V或许将成为一个突破口。
武延军对此也表示认同。他指出，对于中国发展RISC-V来说，很明显的优势首先是一个巨量的市场；然后有大规模的计算机工程师团队和专业人才；同时有足够丰富的应用场景能够充分发挥RISC-V的多样性。
“还有一个独特的优势，那就是我们可能没有国外一些老牌厂商和机构的历史包袱，没有它们已经存在的商业利益或者技能上的惯性，使我们有机会从一个全新的视角、一个全新的架构下面去做一些开创性工作。” 武延军补充道。
不过，相应的劣势就是中国在很多技术领域没有国外巨头在历史上积累的丰富经验。同时，在高端核心人才储备上也没有像国外处理器、基础软件巨头那么多。此外，可能在专利等知识产权方面也存在一定差距。
因此，对于RISC-V在中国的发展，除了需要面对制造的挑战以外，还需要克服处理器IP核和核心基础软件两方面的困难。其中，后者的编译工具链和操作系统是RISC-V生态里面最难突破的两个技术点，跟国际水平存在差距。
平头哥半导体副总裁孟建熠对此指出，RISC-V还需要大家把整个架构不断的做得稳定、可靠；有了“稳定、可靠”之后整个上层的软件适配就会变得更加容易，软件越来越丰富之后要做的事情就是往各个纵深的方向去发展。目前来看，整个发展势头还是非常好的。
生态，RISC-V胜负关键手
RISC-V的潜力，正让很多芯片公司感到兴奋，国内外新兴RISC-V企业和原有科技巨头纷纷布局卡位。
据华泰证券研报，开源指令集架构RISC-V在AIoT时代应用场景及市场空间将快速增长，预计2025年全球RISC-V出货量将超过600亿颗。RISC-V似乎爆发在即，生态建设成为胜负手关键。
从市场现状来看，RISC-V目前只是在MCU领域稍有建树，局限于一些对算力和对软件生态丰富程度要求不那么高的领域。在往更高性能的处理器方面，RISC-V目前似乎与Arm和X86相比，差距甚大。
针对其应用领域，RISC-V基金会CTO Mark Himelstein透露，RISC-V在未来几年将会发力汽车、数据中心和安全等多个应用领域，包括最近备受关注的加速器和智能网卡也是RISC-V未来几年的关注重点。
孟建熠强调，RISC-V现在应该在一个初步证明它商业化可用的这条路上，其发展速度非常快。和Arm一样，RISC-V在早期的时候首先在嵌入式领域先打开一个环节，基本证明可用。接下来，可能就是它要往更加高性能、更加纵深的方向去发展。
比如数据中心，就是RISC-V正在努力的方向。
然而，近日Arm高管在接受采访时试图淡化RISC-V对ARM业务的威胁。Arm基础设施业务高级副总裁兼总经理Chris Bergey表示：“虽然RISC-V可能在某些市场上获得牵引力，但Arm并不担心新贵ISA会侵蚀其在数据中心的新立足点。”
RISC-V要进入数据中心，确实还有很大的距离。孟建熠表示，要往高性能发展，处理器首先本身要在性能上有突破。“性能突破”就需要投入大量的研发精力和资金投入；其次，往高性能走，处理器的“稳定性”也是非常重要的。上层软件栈越厚，软件与硬件的协同优化就越困难，必须要求足够的硬件稳定性，来适配更复杂的应用场景。
武延军也认为，RISC-V往高性能走，对处理器的设计和软件生态都是一个非常大的挑战。
他举例表示，中科院软件所在国际上游开源社区，国内开源欧拉社区、平头哥等都做了大量的RISC-V基础软件适配工作。“但我们可以看到，目前还是有很多核心的基础软件不能很好地跑在RISC-V平台上。”据了解，这可能会与指令集规范还不成熟有关，但更多的是因为这些基础软件包之前都是在X86和ARM上面去跑，从维护者、社区的角度，还没有把RISC-V当成Tier1或者First-Class-Citizen度去对待。这有理念问题，有投入问题，也有商业利益回报的问题。
因此，基础软件的意义是避免生态碎片化，避免大家在同样的指令集架构下还去“造轮子”。因此在一些基础的编译工具链、操作系统，比如GCC/LLVM，Linux等，希望整个行业能够合力去共同打造，而不是说每个RISC-V处理器厂家都要自己做一套，这其实是一个很大的浪费。
目前，国外已经有好多公司在往数据中心的方向发展了，它们也是整个赛道里面非常热门的公司，在大家努力下，整个RISC-V在数据中心的生态正在逐步建设当中。
写在最后
从生态繁荣程度来看，RISC-V目前可能处在ARM大概在2000-2005年的阶段。彼时，有关ARM的各种培训资料、书籍、培训班在大量的出现，市场上对ARM开发者的人才需求也是突然爆增，在嵌入式领域已经跟X86、MIPS已经展开了非常激烈的竞争。
武延军指出，现在RISC-V也差不多到了这个阶段，而且以一个更高的加速度在增长。我们看到不光是教育科研、各种文档材料、人才培训，而且已经有非常好的商用场景、非常高的出货量。当前阶段，可能也是整个生态上下游都可以开始发力的阶段了，如果说之前还在观望的话，我觉得从这个阶段开始大家已经可以不用去观望了，不管是从个人学习能力提升角度，亦或是从公司业务发展的角度，都可以放心地去投入RISC-V领域了。
虽然Arm今天可能不将RISC-V视为威胁，但也正如我们观察到的，随着RISC-V的发展，Arm已经改变了其内核许可的方式。例如，面对免版税、开放和模块化的竞争对手，Arm已采取措施，在某些情况下降低其IP许可的前期成本，并允许被许可人添加自定义指令。
回过头看，无论是Wintel（Windows+Intel），还是AA联盟（Android +Arm），以及那些早已凋零的MIPS、PowerPC等芯片架构，决定竞争终局的或许从来不是技术本身，而是丰富的上层应用、高效的软硬协同以及爆发的市场需求。
总结起来就是两个字：生态。
生态的培育，是一个漫长的过程。对于RISC-V未来的生态发展趋势，孟建熠强调：“从竞争性角度看，Arm架构和x86架构分别在移动终端、PC和服务器市场垄断多年，在这些领域RISC-V新玩家渗透进去还非常需要时日。但是在AIoT、新能源汽车电子、异构计算等新兴领域，RISC-V和其他架构站在同一起跑线，反而具备一些巨头们不具备的新起跑优势。”
展望未来，RISC-V与Arm和X86将会在竞争过程中逐渐找到自己的定位，不断融合、互相借鉴、长期共存，形成“三分天下”的局面，在各自擅长的领域发挥优势。
再谈DDR内存技术原理
从动态的角度来分析时序结构，包括read/write的整个过程到数据返回发生了什么，其中包括以下内容：
·DDR是如何完成Read、Write 等的操作
·DDR的基本命令
·DDR的时序参数
·DDR的性能分析
中国存储器芯片行业概览
1 DRAM基本组成
对于DRAM，其主要由行和列组成，每一个bit中都是由类似右下图的类晶体管的结构组成，对于sdram的数据，可以通过控制column和row就可以访问sdram的随机地址的内容。
·读取某一个bit的状态，就是选中word line，那么图示中的晶体管M1就会导通，通过bit line的sense就可以感知到这个时候电容Cs上的状态，例如，现在如果这个bit的状态为1，那么导通之后就会从bit line上得到1，反之也是同样的道理。
·向某一bit写入1，首先通过row decoder选中word line,将会导致m1导通，那么bit line为1，会导致电容Cs充电，导致其电平为1，如果要写入，那么bit line的电平 为0，将会导致电容Cs放电，致使此时的电平为0。
由上面可以看出一个位只能表示一个bit，那么我们想读取多个位的时候，该怎么办呢？那就出现了ddr中的bank的概念，由多个memory array就组成了一个bank，如下图,一次可以读取2bit/4bit/8bit的数据：
由多个bank就可以组成一个memory device，如下图，一个dram的芯片，由8个banks组成，而每个bank由4个memory array构成，而此时每个bank输出4个bit的I/O bus，那么为什么会出现bank的概念呢？动态内存区别于静态内存要定时刷新，每读取一个状态的时候，都需要重新充电。如果没有采用bank，假设我们现在要读取01-08地址的数据，当读取01的地址后，要等这个bank自刷新后才能读取02地址的值，而采用8个bank之后，没有这类问题，例如我们读取完01地址之后，那么读取02，因为02与01的控制方式不同，所以对于这段时间01可以后台的完成自刷新，依次类推，那么就可以很好的解决动态内存需要刷新的问题。
上图中，对于RANK、DIMM等在深入浅出DDR系列(1)——DDR原理篇中已经详细介绍了。
2 DDR工作原理
了解了DDR的基本组成后，我们来看看DDR如何来完成一次的读写过程。如下图所示，DRAM 的相关操作在内部大概可以分为以下的四个阶段：
·command transport and decode: 在这个阶段，Host 端会通过 Command Bus 和 Address Bus 将具体的 Command 以及相应参数传递给 DRAM。DRAM 接收并解析 Command，接着驱动内部模块进行相应的操作。其中会根据将addr bus上的数据解码成对应的row address和通过bank control解码后得到对应的bank，其次对应的column也会解码得到对应的地址
·in bank data movement: 在这个阶段，第一阶段发送需要读取的 Column 的地址给 DRAM。然后 DRAM 再将 Active Command 所选中的 Row 中，DRAM 就将 Memory Array 中的数据从 DRAM Cells 中读出到 Sense Amplifiers，或者将数据从 Sense Amplifiers 写入到 DRAM Cells。
·in device data movement: 这个阶段中，数据将通过 IO 电路缓存到 Read Latchs 或者通过 IO 电路和 Write Drivers 更新到 Sense Amplifiers。
·system data transport: 在这个阶段，进行读数据操作时，SDRAM 会将数据输出到数据总线上，进行写数据操作时，则是 Host 端的 Controller 将数据输出到总线上。
在上述的四个阶段中，每个阶段都会有一定的耗时，例如数据从 DRAM Cells 搬运到 Read Latchs 的操作需要一定的时间，因此在一个具体的操作需要按照一定时序进行。同时，由于内部的一些部件可能会被多个操作使用，例如读数据和写数据都需要用到部分 IO 电路，因此多个不同的操作通常不能同时进行，也需要遵守一定的时序。此外，某些操作会消耗很大的电流，为了满足 SDRAM 设计上的功耗指标，可能会限制某一些操作的执行频率。
3 DRAM基本命令
对于一款ddr，我们需要知道通过什么样的方式来控制完成我们需要，我们来看看ddr的状态，ddr的工作就是在这几个状态之间切换：
在芯片上电后，完成初始化后，dram处于idle阶段，上图是需要进入各个阶段的时候，应该需要进行那些基本的操作，对于ddr使用比较频繁的几个基本命令访问方式如下
·刷新模式:储体中电容的数据有效是有时间限制的，所以为了保证数据的不丢失，所以要对ddr进行定时的刷新，SDRAM内部有一个行地址生成器（也称刷新计数器）用来自动的依次生成行地址。由于刷新是针对一行中的所有存储体进行。该模式是由Host主动控制DRAM完成刷新，存储体中电容的数据有效保存期上限是64ms（毫秒，1/1000秒），也就是说每一行刷新的循环周期是64ms。
·自我刷新模式：当系统进入低功耗模式，只需要发送一条 SRF指令，主要用于休眠模式低功耗状态下的数据保存，比较常见的应用是STR（Suspend to RAM，休眠挂起于内存）。就进入了该模式，此时不再依靠系统时钟工作，而是根据内部的时钟进行刷新操作。期间除了CKE之外的所有外部信号都是无效的（无需外部提供刷新指令），只有重新使CKE有效才能退出自刷新模式并进入正常操作状态。
·MRS模式(mode register set):模式寄存器中的数据控制着 DDR2 SDRAM的操作模式.它控制着 CAS 延迟, 突发长度, 突发顺序, 测试模式, DLL复位, WR等各种选项，支持着 DDR2 SDRAM 的各种应用. 模式寄存器的默认值没有被定义, 所以上电之后必须按规定的时序规范来设定模式寄存器的值。
·EMRS 扩展模式寄存器:存储着激活或禁止DLL的控制信息, 输出驱动强度, ODT 值的选择 和附加延迟等信息。
·预充电：对一行读写操作后，关闭现有工作行，准备打开新行的操作就是预充电。
·读过程：访问操作开始ACT一个激活命令，主要是激活bank和rol，就等于选通了某一Bank的某一行，接着发送一个read指令，就可以通过数据总线将数据送出去了，然后就进行预充电，恢复到读写的状态，预充电完成后，就恢复到idle状态。
·写过程：与读过程基本类似。
DRAM的基本命令是通过操作各种控制信号/地址信号的组合来完成，下表是DRAM的命令表：
4 DDR的时序参数
4.1 Row Active Command
在进行数据的读写前，Controller 需要先发送 Row Active Command，打开 DRAM Memory Array 中的指定的 Row。Row Active Command 的时序如下图所示：
tRCD：RAS－to－CAS Delay（tRCD），内存行地址传输到列地址的延迟时间。
Row Active Command 通过地址总线指明需要打开某一个 Bank 的某一个 Row。DRAM 在接收到该 Command 后，会打开该 Row 的 Wordline，将其存储的数据读取到 Sense Amplifiers 中，这一时间定义为 tRCD（RCD for Row Address to Column Address Delay）。DRAM 在完成 Row Sense 阶段后，Controller 就可以发送 Read 或 Write Command 进行数据的读写了。这也意味着，Controller 在发送 Row Active Command 后，需要等待 tRCD 时间才能接着发送 Read 或者 Write Command 进行数据的读写。
tRAS: Row Active Time,内存行地址选通延迟
由于 DRAM 的特性，Row 中的数据在被读取到 Sense Amplifiers 后，需要进行 Restore 的操作。Restore 操作可以和数据的读取同时进行，即在这个阶段，Controller 可能发送了 Read Command 进行数据读取。
DRAM 接收到 Row Active Command 到完成 Row Restore 操作所需要的时间定义为 tRAS（RAS for Row Address Strobe）。
Controller 在发出一个 Row Active Command 后，必须要等待 tRAS 时间后，才可以发起另一次的 Precharge 和 Row Access。
4.2 Column Read Command
Controller 发送 Row Active Command 并等待 tRCD 时间后，再发送 Column Write Command 进行数据写入。数据 Burst Length 为 8 时的 Column Write Command 时序如下图所示：
tCWD/tCL/tCWL:内存CAS延迟时间
Column Write Command 通过地址总线 A 指明需要写入数据的 Column 的起始地址。Controller 在发送完 Write Command 后，需要等待 tCWD （CWD for Column Write Delay） 时间后，才可以发送待写入的数据。tCWD 在一些描述中也称为 tCWL（CWL for Column Write Latency）
tWR（WR for Write Recovery）
DRAM 接收完数据后，需要一定的时间将数据写入到 DRAM Cells 中，这个时间定义为 tWR（WR for Write Recovery）。该值说明在一个激活的bank中完成有效的写操作及预充电前，必须等待多少个时钟周期。这段必须的时钟周期用来确保在预充电发生前，写缓冲中的数据可以被写进内存单元中。同样的，过低的tWD虽然提高了系统性能，但可能导致数据还未被正确写入到内存单元中，就发生了预充电操作，会导致数据的丢失及损坏。
4.3 Precharge Command
要访问 DRAM Cell 中的数据，需要先进行 Precharge 操作。相应地，在 Controller 发送 Row Active Command 访问一个具体的 Row 前， Controller 需要发送 Precharge Command 对该 Row 所在的 Bank 进行 Precharge 操作。下面的时序图描述了 Controller 访问一个 Row 后，执行 Precharge，然后再访问另一个 Row 的流程。
DRAM 执行 Precharge Command 所需要的时间定义为 tRP（RP for Row Precharge）。Controller 在发送一个 Row Active Command 后，需要等待 tRC（RC for Row Cycle）时间后，才能发送第二个 Row Active Command 进行另一个 Row 的访问。
从时序图上我们可以看到，tRC = tRAS + tRP，tRC 时间决定了访问 DRAM 不同 Row 的性能。在实际的产品中，通常会通过降低 tRC 耗时或者在一个 Row Cycle 执行尽可能多数据读写等方式来优化性能。
4.4 Row Refresh Command
一般情况下，为了保证 DRAM 数据的有效性，Controller 每隔 tREFI（REFI for Refresh Interval） 时间就需要发送一个 Row Refresh Command 给 DRAM，进行 Row 刷新操作。DRAM 在接收到 Row Refresh Command 后，会根据内部 Refresh Counter 的值，对所有 Bank 的一个或者多个 Row 进行刷新操作。
DRAM 刷新的操作与 Active + Precharge Command 组合类似，差别在于 Refresh Command 是对 DRAM 所有 Bank 同时进行操作的。下图为 DRAM Row Refresh Command 的时序图：
DRAM 完成刷新操作所需的时间定义为 tRFC（RFC for Refresh Cycle）。
tRFC 包含两个部分的时间，一是完成刷新操作所需要的时间，由于 DRAM Refresh 是同时对所有 Bank 进行的，刷新操作会比单个 Row 的 Active + Precharge 操作需要更长的时间；tRFC 的另一部分时间则是为了降低平均功耗而引入的延时，DRAM Refresh 操作所消耗的电流会比单个 Row 的 Active + Precharge 操作要大的多，tRFC 中引入额外的时延可以限制 Refresh 操作的频率。
4.5 Read Cycle
一个完整的 Burst Length 的 Read Cycle 如下图所示：
下面是DDR常见的一些参数及定义如下：
上述的 DRAM Timing 中的一部分参数可以编程设定，例如 tCAS、tAL、Burst Length 等。这些参数通常是在 Host 初始化时，通过 Controller 发起 Load Mode Register Command 写入到 DRAM 的 Mode Register 中。DRAM 完成初始化后，就会按照设定的参数运行。
5性能分析
在学习完DDR的基本操作和时序参数之后，我们就看看性能的影响。当频率和位宽固定后，带宽也就不可更改，但是在内存的工作周期内，不可能总处于数据传输的状态，因为要有命令、寻址等必要的过程。那么这些操作占用的时间越短，内存工作的效率就越高，性能也就越好。
对于我们来说，最好的方法是提高频率，但是提高频率会受多方面的影响，还有什么办法提高内存访问采取速度。
·多通道: 现代内存控制器从北桥移入CPU内部，而内存控制器都可以同时操作多个通道。比如现在的笔记本开始支持双通道、三通道，如果数据分布在不同通道的内存条上，内存控制器就可以不用管上面的这些延迟时序，同时可以读取它们，速度可以增加两倍，甚至三倍。
·交织方式(Interleaving)： 同一块内存分布到不同的通道中去，这样无论Cache命中与否都可以同时存取，多通道的技术才能发挥更大的用处。
·超频内存： 也就是提升DDR的频率来增加速度
6总结
对于DDR的读写以及一些时序参数的原理性知识后，下一步就进入到DDR的驱动调式，主要是对于一款控制器，我们该如何去调试DDR。其实对于DDR的调试，主要的读写的控制，都是由DDR的控制器完成了，我们主要是通过MRS模式/EMRS模式来完成对于DDR参数的配置，而对于MRS模式的使用，已经集成到DDR控制器中完成了，我们只需要根据控制器手册配置相应的寄存器就可以完成对于DDR调试。
参考文献链接
https://mp.weixin.qq.com/s/QYaxxFL9DgaH8w7-oWTLnQ
https://mp.weixin.qq.com/s/piz05ElprV88xnrHFufnvg
https://mp.weixin.qq.com/s/6H_JsumnYUoyIT8_ovBihQ
https://mp.weixin.qq.com/s/O1jMMqKn-fnQp0T35tMP-w
----------------------------- 吴博士，有时间讨论一下内存优化啊[大笑]
页:
[1]