自動(dòng)駕駛-芯片-DDR

果殼醬

自動(dòng)駕駛-芯片-DDR
參考文獻(xiàn)鏈接
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
裁員、賣(mài)身、大虧損，西半球刮來(lái)自動(dòng)駕駛第一股寒風(fēng)
「今年開(kāi)始，行業(yè)大環(huán)境突變……」
自動(dòng)駕駛第一股圖森未來(lái)原董事長(zhǎng)陳默在 7 月份接受媒體采訪時(shí)的一句話，總結(jié)了美國(guó)上半年的市場(chǎng)情況，也意外成功預(yù)言了下半年的行業(yè)趨勢(shì)。

登錄/注冊(cè)后可看大圖

起初是圖森在 3 月份突傳出售中國(guó)業(yè)務(wù)，團(tuán)隊(duì)隨之大調(diào)整，陳默也從公司出走。
接著，6 月份，特斯拉關(guān)閉位于加州圣馬特奧的自動(dòng)駕駛系統(tǒng)部門(mén)，同時(shí)裁減約 200 個(gè)時(shí)薪制崗位。沒(méi)多久，明星公司 Argo AI 也宣布裁減約 150 名員工。
而最近的消息是，上市不到一年的 Aurora，尋求部分或整體將公司打包出售。要知道，在 Aurora 去年的招股書(shū)中，還曾意氣風(fēng)發(fā)地把自己和 Waymo、Cruise 并列為自動(dòng)駕駛?cè)痔煜碌膭?shì)力......
毫無(wú)疑問(wèn)，事情正在起變化，而背后的原因是行業(yè)正在經(jīng)歷一輪資本退潮：
一邊是美聯(lián)儲(chǔ)加息背景下資金成本越來(lái)越高，另一邊是自動(dòng)駕駛企業(yè)越來(lái)越止不住的燒錢(qián)窟窿，投資方此時(shí)選擇站在一處，目睹一場(chǎng)泥沙俱下的行業(yè)「盛況」。
自動(dòng)駕駛企業(yè)對(duì)此無(wú)能為力，在技術(shù)、商業(yè)化均未取得突破性進(jìn)展之前，它們只能仰賴(lài)外部輸血。而當(dāng)這一過(guò)程被突然中斷，也只能選擇斷臂求生，甚至將自己打包出售。
「凜冬已至?！挂晃粯I(yè)內(nèi)人士感慨道。值得注意的是，這場(chǎng)看似起源于西半球的寒風(fēng)，也刮到了國(guó)內(nèi)。在前不久，華為創(chuàng)始人任正非率先感受到寒氣，并提出要傳給每一個(gè)人。
2022，注定是自動(dòng)駕駛行業(yè)空前艱難的一年。
01
自動(dòng)駕駛堪稱(chēng)碎鈔機(jī)，
Cruise 每天消耗 500 萬(wàn)美元
3 月份，圖森未來(lái)傳出要將亞太業(yè)務(wù)（主要是中國(guó)業(yè)務(wù)）以 10 億美元的價(jià)格出售。在外界看來(lái)，這是一筆為了符合美國(guó)監(jiān)管而迫不得已達(dá)成的交易。
自去年 4 月圖森 IPO 以來(lái)，美國(guó)外國(guó)投資委員會(huì) (CFIUS) 就一直以圖森未來(lái)?yè)碛兄袊?guó)區(qū)業(yè)務(wù)為由對(duì)其展開(kāi)調(diào)查，并最終在今年 2 月份與圖森達(dá)成協(xié)議，后者需將自動(dòng)駕駛卡車(chē)業(yè)務(wù)的一些技術(shù)監(jiān)督權(quán)移交給美國(guó)相關(guān)部門(mén)，并定期開(kāi)會(huì)并向美國(guó)外資投資委員會(huì)報(bào)告等。
在這種要求下，為了避免糾紛，圖森直接出售亞太業(yè)務(wù)顯得合情合理。然而，陳默表示，此舉還有另一層原因：缺錢(qián)?！赋鍪蹃喬珮I(yè)務(wù)是為了讓圖森回籠資金，以支撐美國(guó)業(yè)務(wù)發(fā)展?！?br /> 據(jù)了解，圖森上市至今「燒錢(qián)」速度驚人。以其最新發(fā)布的二季度報(bào)為例，實(shí)現(xiàn)營(yíng)收 259.4 萬(wàn)美元，錄得虧損 1.11 億美元（折合約為 7.59 億人民幣），其中研發(fā)投入 5.93 億元人民幣，如果按每季度 90 天計(jì)算，每天消耗超過(guò) 600 萬(wàn)元。
作為上市公司，圖森按理來(lái)說(shuō)可以在資本市場(chǎng)尋求募資，然而其股價(jià)自高峰一落千丈，市值跌了 80%，陳默坦言：「我們不可能在市場(chǎng)上募集到錢(qián)?！?br />

登錄/注冊(cè)后可看大圖

和圖森一樣陷入財(cái)務(wù)困境的還有明星公司 Argo AI。
這是一家由福特和大眾集團(tuán)投資的自動(dòng)駕駛技術(shù)企業(yè)，由谷歌和 Uber 背景的兩位技術(shù)大牛在 2016 年成立，在僅有 12 名員工的時(shí)候就獲得福特 10 億美元投資，而后又在 2020 年，獲得大眾 26 億美元投資。
7 月 7 日，Argo AI 宣布裁減約 150 名員工。盡管該公司發(fā)言人旋即表示，此次裁員不會(huì)減緩自動(dòng)駕駛系統(tǒng)開(kāi)發(fā)商業(yè)服務(wù)的計(jì)劃，但還是讓外界感覺(jué)到了自動(dòng)駕駛行業(yè)的一絲寒意。
如今，這股寒意已經(jīng)傳導(dǎo)到下半年。日前有消息稱(chēng)，去年底以 SPAC 方式借殼上市的自動(dòng)駕駛公司 Aurora，將通過(guò)裁員、降低福利等方式削減成本，以度過(guò)當(dāng)前入不敷出的困境。

登錄/注冊(cè)后可看大圖

Aurora 首席執(zhí)行官克里斯·厄姆森（Chris Urmson）表示，正在考慮出售優(yōu)質(zhì)資產(chǎn)如旗下激光雷達(dá)業(yè)務(wù)，或?qū)⒐菊w打包出售給蘋(píng)果、微軟或其他潛在的第一梯隊(duì)汽車(chē)公司。
有媒體報(bào)道，自上市至今，Aurora 已經(jīng)燒掉了大約 2.3 億美元的現(xiàn)金，其股價(jià)也已經(jīng)從最高的每股 17.77 美元跌落至最近的 1.43 美元，降幅超 90%。
不過(guò)和真正的「燒錢(qián)大戶(hù)」Cruise、Waymo 相比，Aurora 顯得「稍遜風(fēng)騷」。

登錄/注冊(cè)后可看大圖

通用汽車(chē)最新財(cái)報(bào)顯示，其自動(dòng)駕駛子公司 Cruise 在第二季度營(yíng)收 2500 萬(wàn)美元，卻虧損了高達(dá) 5.4 億美元，相當(dāng)于每天 500 萬(wàn)美元。

登錄/注冊(cè)后可看大圖

如果算上 2022 年整個(gè)上半年，累計(jì)虧損達(dá) 9 億美元，較去年同期多虧損 3 億美元。
Waymo 也「不遑多讓」，Google 母公司 Alphabet 發(fā)布的 2022 財(cái)年第二季度財(cái)報(bào)顯示，OtherBets（包括自動(dòng)駕駛汽車(chē)部門(mén) Waymo 和生命科學(xué)部門(mén) Verily 等）本季度虧損 16.68 億美元。
雖然沒(méi)有直接的數(shù)據(jù)，但根據(jù)此前 Waymo 前員工及其他行業(yè)內(nèi)人士估算，Waymo 每年要在各方面花掉 10 億美元以上。
在過(guò)去，自動(dòng)駕駛公司「燒錢(qián)」并不是什么新鮮事，按照以往業(yè)內(nèi)人的說(shuō)法，這是贏在未來(lái)之前必要的投資，不同玩家或依靠母公司，或引入外部融資等方式進(jìn)行輸血，然而現(xiàn)在，事情開(kāi)始起變化。
02
技術(shù)、商業(yè)接連受阻，
資本市場(chǎng)「累覺(jué)不愛(ài)」
據(jù)賽博汽車(chē)報(bào)道，在自動(dòng)駕駛干線物流賽道，有企業(yè)今年見(jiàn)了 80 多家投資機(jī)構(gòu)后，才在融資方面有了點(diǎn)眉目，而不少企業(yè)「已經(jīng)見(jiàn)了一百多家了，還沒(méi)有著落」。
按照陳默的說(shuō)法，美聯(lián)儲(chǔ)多次加息，導(dǎo)致資本市場(chǎng)銀根收緊，是重要原因之一。
世界首富馬斯克也在推特上直言「美聯(lián)儲(chǔ)大幅加息有通縮的風(fēng)險(xiǎn)」，并接連出售約 800 萬(wàn)股特斯拉股票，累計(jì)套現(xiàn) 70 億美元。
今年以來(lái)，美股大多數(shù)科技公司都受到加息周期影響，市值接連回調(diào)，自動(dòng)駕駛公司下跌表現(xiàn)尤甚，這讓原本看好該領(lǐng)域的投資方也變得遲疑，為了保證自己資金鏈的安全，他們選擇少出手甚至不出手。

登錄/注冊(cè)后可看大圖

以 Mobileye 為例，該公司本來(lái)準(zhǔn)備在今年上市，一度被市場(chǎng)給出 500 億美元的估值，有望拿下今年美股最大的 IPO。
2022 年 3 月份，英特爾公司宣布已經(jīng)向美國(guó)證劵交易委員會(huì)秘密提交 Form S-1 注冊(cè)聲明草案，擬首次公開(kāi)發(fā)行 Mobileye 新股。按照規(guī)劃，上市籌措的資金將用于加碼 Mobileye 新業(yè)務(wù)。
然而，僅僅過(guò)了半年，Mobileye 估值大砍 40％，英特爾預(yù)計(jì)，該公司最新估值最高為 300 億美元，遠(yuǎn)低于最初期望水平。為此，Mobileye 上市的計(jì)劃便不了了之。
投資網(wǎng)站 PitchBook 的數(shù)據(jù)顯示，2022 年自動(dòng)駕駛投資已大幅下滑，第二季度對(duì)自動(dòng)駕駛創(chuàng)業(yè)公司的投資下降到 9.58 億美元，在風(fēng)險(xiǎn)投資中的占比不到 10%。
而自動(dòng)駕駛公司今年在技術(shù)和商業(yè)化上不及預(yù)期的情況，更加重了資本市場(chǎng)的觀望情緒。
以圖森未來(lái)為例，4 月份，其在美國(guó)部署測(cè)試的自動(dòng)駕駛卡車(chē)突然左轉(zhuǎn)橫切，撞上一個(gè)水泥路障。這起交通事故經(jīng)《華爾街日?qǐng)?bào)》報(bào)道后，很快引來(lái)當(dāng)?shù)乇O(jiān)管部門(mén)對(duì)于自動(dòng)駕駛安全風(fēng)險(xiǎn)的關(guān)注。
為此，圖森未來(lái)停止了整個(gè)車(chē)隊(duì)的測(cè)試，并進(jìn)行獨(dú)立調(diào)查，圖森未來(lái)聯(lián)合創(chuàng)始人兼 CEO 侯曉迪還表示，在確保車(chē)隊(duì)可以安全運(yùn)營(yíng)之前，圖森未來(lái)不會(huì)進(jìn)行商業(yè)化。
Cruise 也在今年經(jīng)歷了一個(gè)「多事之秋」。
5 月份，該公司遭內(nèi)部員工向加州監(jiān)管機(jī)構(gòu)匿名舉報(bào)，Cruise 存在故意隱瞞涉及車(chē)輛和業(yè)務(wù)具有潛在破壞性的問(wèn)題，如安全報(bào)告系統(tǒng)仍處于混亂狀態(tài)，這一消息讓外界嘩然。
6 月，Cruise 再次卷入風(fēng)波，先后發(fā)生兩起交通事故。起先是其 Robotaxi 在左轉(zhuǎn)轉(zhuǎn)彎處與一輛迎面而來(lái)的豐田普銳斯發(fā)生碰撞，導(dǎo)致兩人受輕傷。

登錄/注冊(cè)后可看大圖

為此，Cruise 召回并更新了 80 輛 Robotaxi 的軟件。
而后在同一個(gè)月內(nèi)，十幾輛 Cruise 自動(dòng)駕駛車(chē)輛無(wú)故停在十字路口，導(dǎo)致交通癱瘓了一個(gè)多小時(shí)，直到人類(lèi)駕駛員將其開(kāi)走之后，事情才得到解決。
事后，Cruise 回應(yīng)稱(chēng)，是「技術(shù)問(wèn)題」導(dǎo)致了上述情況發(fā)生。
讓公眾對(duì)自動(dòng)駕駛技術(shù)越發(fā)不信任的還包括特斯拉。

登錄/注冊(cè)后可看大圖

7 月 6 日，一輛 2015 款的特斯拉 Model S 從佛羅里達(dá)州蓋恩斯維爾以南的 75 號(hào)州際公路上駛?cè)敫咚俟沸菹^(qū)停車(chē)場(chǎng)，徑直撞上了一輛停在那里的牽引車(chē)，盡管還未證實(shí)在當(dāng)時(shí)自動(dòng)駕駛功能是否開(kāi)啟，但巧合的是，這次撞上的又是一輛「白車(chē)」。
更早之前，根據(jù)美國(guó)國(guó)家公路交通安全管理局（NHTSA）的數(shù)據(jù)，截至 2022 年 5 月的 10 個(gè)月中，有 200 多起車(chē)禍都與特斯拉的 Autopilot 軟件有關(guān)。
特斯拉為此被加州機(jī)動(dòng)車(chē)輛管理局（DMV）和美國(guó)國(guó)家公路交通安全管理局（NHTSA）認(rèn)定其技術(shù)存在虛假宣傳。加州參議院也已通過(guò)法案，禁止在智能駕駛的廣告里包含「自動(dòng)駕駛」等詞匯。
自動(dòng)駕駛公司在技術(shù)上還沒(méi)達(dá)到成熟，已經(jīng)開(kāi)始讓之前的投資人變得焦急，而直到現(xiàn)在，商業(yè)化上也一直沒(méi)有大的起色，更是讓這些人「封住了口袋」。
Aurora 最近宣布，其最快商用的自動(dòng)駕駛卡車(chē)上路的時(shí)間，仍舊需要推遲一年，至 2024 年。
該公司創(chuàng)始人兼 CEO 克里斯·厄姆森解釋稱(chēng)，背后原因在于 OEM 進(jìn)展緩慢。由于商業(yè)化不及預(yù)期，他還表示，「預(yù)計(jì)未來(lái)六個(gè)月內(nèi)，不會(huì)出現(xiàn)足夠規(guī)模的傳統(tǒng)融資機(jī)會(huì)來(lái)為公司續(xù)命?！?br /> 圖森未來(lái)面臨相同的窘境，此前，其與美國(guó)第三大卡車(chē)制造商 Navistar 達(dá)成合作，實(shí)現(xiàn) L4 自動(dòng)駕駛卡車(chē)的量產(chǎn)。
然而隨著時(shí)間的推移，Navistar 一直在推遲圖森未來(lái)量產(chǎn) L4 自動(dòng)駕駛卡車(chē)的進(jìn)程，從預(yù)定的 2024 年，到 2025 年，2026 年……
陳默曾在一次媒體訪談中直接吐槽 Navistar 的自動(dòng)駕駛量產(chǎn)進(jìn)度太慢，尤其在歐美市場(chǎng)?！肝覀?cè)窘o投資人的承諾是，圖森未來(lái)在 2024 年 Q3 實(shí)現(xiàn)量產(chǎn)……」
Waymo 和 Cruise 倒是在商業(yè)化上傳來(lái)了一些好消息。
3 月 1 日，美國(guó)加州公共事業(yè)委員會(huì) (CPUC) 向 Alphabet 和通用汽車(chē)發(fā)放了自動(dòng)駕駛客運(yùn)服務(wù)的許可證，允許兩家公司旗下的自動(dòng)駕駛公司在舊金山及周邊提供收費(fèi)客運(yùn)服務(wù)。
此前，Waymo 和 Cruise 在加州只被允許在測(cè)試的基礎(chǔ)上提供有限載人客運(yùn)，不準(zhǔn)收取費(fèi)用。
不過(guò)這只能算是商業(yè)化初期，畢竟 Robotaxi 當(dāng)前鋪開(kāi)的城市并不多，以 Waymo 為例，目前僅在鳳凰城的東谷、鳳凰城市中心以及舊金山三地向公眾提供出行服務(wù)。
此外，Waymo 目前的商業(yè)模式還略顯雞肋。根據(jù) CNBC 記者在 2022 年 1 月的體驗(yàn)，Waymo 在鳳凰城的 Robotaxi 跑 5 英里需用時(shí) 14 分鐘，每分鐘收費(fèi) 1 美元。
而跑同樣里程的 Uber 不僅快得多，每分鐘收費(fèi)更僅有 0.4 美元。
「舊金山 Robotaxi 的服務(wù)價(jià)格將定在合理且有競(jìng)爭(zhēng)力的區(qū)間，與此同時(shí)，有限用戶(hù)免費(fèi)測(cè)試/收費(fèi)商業(yè)運(yùn)作的模式將在其他地方鋪開(kāi)?！筗aymo 發(fā)言人在一場(chǎng)會(huì)議上表達(dá)了自己的樂(lè)觀。
相較而言，更多的人對(duì)于自動(dòng)駕駛卻沒(méi)有這么樂(lè)觀，他們眼下更關(guān)心的是，漫漫「燒錢(qián)」之路還要延續(xù)多久？
03
燒錢(qián)游戲進(jìn)入倒計(jì)時(shí)：
沒(méi)有金山倚靠，或需考慮賣(mài)身
截至 2022 年 6 月 30 日，圖森未來(lái)資產(chǎn)負(fù)債表上僅有 11.6 億美元的現(xiàn)金，而按照其給出的 2022 年調(diào)整后的 EBITDA 虧損 3.6-3.8 億美元來(lái)看，在不融資的情況，公司現(xiàn)金僅夠支撐 3 年左右。
而現(xiàn)在，圖森仍沒(méi)有開(kāi)啟自主造血的進(jìn)程，資料顯示，其卡車(chē)預(yù)訂總數(shù)達(dá)到 7485 輛，但至今仍無(wú)一輛進(jìn)入交付量產(chǎn)階段。
在招股書(shū)中，圖森曾這樣計(jì)算，一輛卡車(chē)一年可實(shí)現(xiàn) 6 萬(wàn)美元營(yíng)收，當(dāng)運(yùn)營(yíng)數(shù)量過(guò) 5000 輛時(shí)，公司將達(dá)到收支平衡，之后才能逐漸實(shí)現(xiàn)盈利。
Aurora 則直接給出了時(shí)間表，預(yù)計(jì) 2027 年之前都無(wú)法實(shí)現(xiàn)盈虧平衡。而 Cruise 首席執(zhí)行官 Kyle Vogt 在今年 9 月 12 日一次會(huì)議上表示，公司目標(biāo)是到 2025 年時(shí)獲得 10 億美元的營(yíng)收，相當(dāng)于目前來(lái)自通用的年投資水平的一半。
以上這些預(yù)示著自動(dòng)駕駛公司還得至少燒 3-5 年的錢(qián)，并且最終還不一定能跑出來(lái)，有投資方等不及了。
2022 年 3 月，私募股權(quán)巨頭軟銀集團(tuán)旗下愿景基金向通用汽車(chē)出售其持有的 Crusie 股份，價(jià)值 21 億美元。
Vogt 坦言，自動(dòng)駕駛行業(yè)已經(jīng)從「極端樂(lè)觀」轉(zhuǎn)向「極度悲觀」。這意味著，隨著資本退潮，自動(dòng)駕駛公司已經(jīng)到了該謀求出路的時(shí)刻。
強(qiáng)如 Waymo、Crusie 這類(lèi)背靠「金山」的企業(yè)，似乎還不用特別擔(dān)心。Cruise 目前還擁有 37 億美元的資金，并獲得了通用汽車(chē)金融部門(mén)約 50 億美元的信貸額度。
在一季度的財(cái)報(bào)發(fā)布會(huì)上，通用計(jì)劃 2022 年向 Cruise 支出約 20 億美元，這比以往任何時(shí)候都要高。
而在接下來(lái)，通用還承諾將持續(xù)為 Cruise 提供資金，確保后者在市場(chǎng)上占據(jù)領(lǐng)導(dǎo)地位。
據(jù)了解，在接收軟銀集團(tuán)轉(zhuǎn)讓的股份后，通用 3 月份就向 Cruise 追加了 13.5 億美元投資。
Waymo 也有母公司 Alphabet 的供養(yǎng)。
在 2020 年之前，谷歌/Alphabet 已經(jīng)為 Waymo 持續(xù)輸血 11 年，不過(guò)由于投入過(guò)于巨大，最終在當(dāng)年年初首次開(kāi)始引進(jìn)外部融資：3 月份籌集了 22.5 億美元，5 月份籌集了 7.5 億美元，最終在 73 天內(nèi)累計(jì)完成 30 億美元的融資。
得益于 Alphabet 強(qiáng)大的品牌效應(yīng)，Waymo 在 2021 年 6 月，又宣布完成一輪 25 億美元的投資。
與此同時(shí)，知情人士透露，Waymo 還討論了最終公開(kāi)上市的計(jì)劃，以減輕其母公司的資金壓力。
不過(guò)按當(dāng)前資本市場(chǎng)的表現(xiàn)，Alphabet 短期內(nèi)大概率不會(huì)推動(dòng) Waymo 的 IPO，取而代之的是，將繼續(xù)用公司的利潤(rùn)去補(bǔ)貼支撐 Waymo 的巨額虧損。
最危險(xiǎn)的是沒(méi)有「靠山」的自動(dòng)駕駛公司。
正如前文所說(shuō)，Aurora 已經(jīng)嘗試性邁出第一步：考慮賣(mài)身主機(jī)廠或科技巨頭。
該公司首席執(zhí)行官克里斯·厄姆森（Chris Urmson）表示，可以將公司私有化，前提是找到一個(gè)能夠提供 15 億美元融資的合作伙伴。「鑒于目前的股價(jià)，我們應(yīng)該成為任何希望擁有自動(dòng)駕駛技術(shù)公司的有吸引力的目標(biāo)?！?br /> 戲劇性的是，Aurora 原本也是通過(guò)收購(gòu)自動(dòng)駕駛公司做大業(yè)務(wù)的，而不料如今卻走向了尋求被收購(gòu)的道路。

登錄/注冊(cè)后可看大圖

2020 年 12 月，還是初創(chuàng)公司的 Aurora 以 40 億美元的價(jià)格，近乎「蛇吞象」式地收購(gòu) Uber 旗下自動(dòng)駕駛部門(mén) ATG，彼時(shí)正值 Uber 面臨巨大運(yùn)營(yíng)成本壓力，其 ATG 部門(mén)在 2019 年造成了約 5 億美元的虧損。
在出售自動(dòng)駕駛業(yè)務(wù)后，Uber 凈虧損確實(shí)得到一定改善。而如今歷史重演，Aurora 面臨整體出售的暗淡時(shí)刻。
盡管如此，出售仍是比破產(chǎn)倒閉更體面的結(jié)局，如果能賣(mài)給主機(jī)廠，甚至可以稱(chēng)得上是個(gè)不錯(cuò)的歸宿。
隨著行業(yè)智能化趨勢(shì)加深，自動(dòng)駕駛被視作汽車(chē)的「靈魂」。
通用、福特、大眾等傳統(tǒng)主機(jī)廠都在這方面加碼投入，而對(duì)于那些還沒(méi)開(kāi)始布局的 OEM，要參與這股浪潮，比較可行的選項(xiàng)便是收購(gòu)一家自動(dòng)駕駛公司，而這正給了像 Aurora 這樣苦苦掙扎的企業(yè)一絲希望。

登錄/注冊(cè)后可看大圖

「網(wǎng)約車(chē)第一股」Lyft 旗下自動(dòng)駕駛部門(mén) Level 5 就是這樣完成「和平交接」的。
成立于 2017 年夏天，Level 5 同樣輝煌過(guò)，在 2018 年已經(jīng)獲得加州公共道路測(cè)試許可，并取得了不錯(cuò)的研發(fā)、測(cè)試數(shù)據(jù)積累。
然而隨著 2020 年疫情突降，母公司 Lyft 遭受巨額業(yè)務(wù)虧損，無(wú)法支撐 Level 5 燒錢(qián)繼續(xù)自動(dòng)駕駛的夢(mèng)想。
2021 年 4 月，Lyft 將 Level 5 以 5.5 億美金的價(jià)格出售給豐田旗下的子公司 Woven Planet Holdings，至此每年至少節(jié)省 1 億美元日常開(kāi)支。
如今，陷入僵局的自動(dòng)駕駛企業(yè)也在謀求賣(mài)身主機(jī)廠的機(jī)會(huì)，后者雄厚的資金可以幫助它們渡過(guò) 2022 這個(gè)難熬的「寒冬」。
日前，就有消息稱(chēng)，吉利控股集團(tuán)向圖森未來(lái)美國(guó)總部發(fā)出收購(gòu)要約，擬收購(gòu)圖森未來(lái)控股亞太地區(qū)業(yè)務(wù)的全部股份，具體收購(gòu)細(xì)節(jié)暫未公布。
04
自動(dòng)駕駛寒風(fēng)開(kāi)始刮到東半球：
活下去，才有未來(lái)
把視線拉回國(guó)內(nèi)，中國(guó)的自動(dòng)駕駛企業(yè)看起來(lái)還沒(méi)遭受到大洋彼岸同行那樣的危機(jī)，但若仔細(xì)看，卻也經(jīng)歷了一些困難。
2021 年，境外上市監(jiān)管環(huán)境收緊，小馬智行、智加科技等 IPO 計(jì)劃折戟，融資節(jié)奏隨之打亂，導(dǎo)致業(yè)務(wù)發(fā)生重大調(diào)整，人員也隨之流動(dòng)。
典型如小馬智行卡車(chē)部門(mén)去年遭遇動(dòng)蕩，該公司原自動(dòng)駕駛技術(shù)總負(fù)責(zé)人潘震皓、國(guó)內(nèi)自動(dòng)駕駛規(guī)劃與控制組負(fù)責(zé)人孫浩文以及負(fù)責(zé)戰(zhàn)略合作和融資的副總裁趙睿璇紛紛出走，各自成立三家卡車(chē)自動(dòng)駕駛公司。

登錄/注冊(cè)后可看大圖

時(shí)間來(lái)到 2022 年，小馬智行卡車(chē)業(yè)務(wù)重整旗鼓：7 月 28 日，小馬智行與三一重卡成立合資公司，計(jì)劃在 2022 年內(nèi)開(kāi)始小規(guī)模量產(chǎn)交付自動(dòng)駕駛卡車(chē)。
而這背后，是小馬智行再次獲得輸血：3 月 7 日，小馬智行宣布完成 D 輪融資的首次交割。
小馬智行 CFO 勞倫斯·斯泰恩表示：「公司財(cái)務(wù)狀況十分穩(wěn)健，為小馬智行未來(lái)幾年的發(fā)展提供堅(jiān)實(shí)基礎(chǔ)，直至我們開(kāi)啟大規(guī)模商業(yè)化的進(jìn)程?！?br /> 繼終止以 SPAC 方式在紐交所上市的協(xié)議后，智加科技同樣傳出有變。
今年 8 月份，虎嗅一篇題為《圖森們的「并購(gòu)式」跑路與謊言》文章寫(xiě)道，有消息人士爆料，干線物流平臺(tái)滿(mǎn)幫集團(tuán)或?qū)⑹召?gòu)自動(dòng)駕駛卡車(chē)企業(yè)智加科技。
不過(guò)，這一消息尚未得到證實(shí)。

登錄/注冊(cè)后可看大圖

從融資上來(lái)看，智加科技在 2021 年創(chuàng)下 4.2 億美元的融資記錄，為干線物流領(lǐng)域最大。
而據(jù)最近消息，2022 年 8 月 16 日，智加科技聯(lián)合摯途科技完成面向榮慶物流的 100 臺(tái)自動(dòng)駕駛重卡訂單的首批交付。
從小馬智行和智加科技來(lái)看，國(guó)內(nèi)的玩家似乎平穩(wěn)度過(guò)了資金難關(guān)。然而，沒(méi)過(guò)多久，8 月 22 日，華為創(chuàng)始人任正非在內(nèi)部喊出「寒氣論」，讓自動(dòng)駕駛行業(yè)再度頓覺(jué)涼意。
按照任正非的說(shuō)法，未來(lái)十年是一個(gè)非常痛苦的歷史時(shí)期，全球經(jīng)濟(jì)會(huì)持續(xù)衰退。為了應(yīng)對(duì)挑戰(zhàn)，華為將收縮或關(guān)閉未來(lái)幾年內(nèi)不能產(chǎn)生價(jià)值和利潤(rùn)的業(yè)務(wù)。
巧的是，智能駕駛便是這樣的業(yè)務(wù)。
在 7 月 7 日的一場(chǎng)行業(yè)論壇上，華為常務(wù)董事、終端 BG CEO、智能汽車(chē)解決方案 BU CEO 余承東透露，汽車(chē)業(yè)務(wù)是華為目前唯一虧損的業(yè)務(wù)。

登錄/注冊(cè)后可看大圖

他表示，華為在該業(yè)務(wù)上一年投入十幾億美元，直接投入 7000 人，間接投入 1 萬(wàn)人。而其中，又以智能駕駛占大頭?！?strong>百分之七八十都投入在智能駕駛輔助領(lǐng)域?！?br /> 而目前，根據(jù)任正非的要求，對(duì)于智能汽車(chē)解決方案，華為將減少科研預(yù)算，加強(qiáng)商業(yè)閉環(huán)研發(fā)。
華為這樣年?duì)I收 6、7000 億的公司尚且在收緊對(duì)自動(dòng)駕駛業(yè)務(wù)的投入，對(duì)于那些還沒(méi)有實(shí)現(xiàn)商業(yè)化、完成自主造血的企業(yè)，未來(lái)資金壓力可見(jiàn)一斑。
回到當(dāng)下，美國(guó)自動(dòng)駕駛企業(yè)的遭遇已經(jīng)為國(guó)內(nèi)同行敲響了警鐘，如果沒(méi)有「大樹(shù)」背靠，按照該領(lǐng)域公司一年一次融資的燒錢(qián)節(jié)奏，現(xiàn)在還沒(méi)有融到錢(qián)的企業(yè)，可以說(shuō)是十分危險(xiǎn)了。
自動(dòng)駕駛不僅是技術(shù)的比拼，更是一場(chǎng)時(shí)間的游戲，能抵達(dá)終點(diǎn)的不一定是跑得最快的那個(gè)，現(xiàn)在活下去才有資格談未來(lái)。
CUDA 矩陣乘法終極優(yōu)化

單精度矩陣乘法（SGEMM）幾乎是每一位學(xué)習(xí) CUDA 的同學(xué)繞不開(kāi)的案例，這個(gè)經(jīng)典的計(jì)算密集型案例可以很好地展示 GPU 編程中常用的優(yōu)化技巧。
CUDA 矩陣乘法優(yōu)化手段詳解
Naive 實(shí)現(xiàn)的分析：到底差在哪里？

筆者面試過(guò)不少具有 CUDA 編程經(jīng)驗(yàn)的校招同學(xué)，當(dāng)提問(wèn)使用 CUDA 編寫(xiě)一個(gè) SGEMM Kernel 的時(shí)候，通常會(huì)獲得這么一個(gè)答案：
__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[ty * K + i] * B[i * N + tx];
        }
        C[ty * N + tx] = c;
    }
}
這樣一個(gè) Naive 的 Kernel 當(dāng)然不是筆者所期待的，因?yàn)檫@個(gè) Kernel 的性能基本可以斷定連 cublas 的 1/10 都不到，顯然不符合我們追求高性能的需求。那么這個(gè) Naive 的實(shí)現(xiàn)究竟差在哪呢？分析代碼我們可以看到，計(jì)算一次 FMA（乘累加）之前需要讀一次 A 和讀一次 B，眾所周知，讀取 Global Memory 的代價(jià)很大，通常都需要幾百個(gè) cycle（時(shí)鐘周期），而計(jì)算一次 FMA 通常只需要幾個(gè) cycle，大量的時(shí)間被花費(fèi)在了訪存上。也許有思維活絡(luò)的同學(xué)立馬想到，可以將 A 和 B 矩陣先搬運(yùn)到 Shared Memory（SM 中低延遲的 on-chip memory，block 內(nèi)線程共享，附 NVIDIA GPU 內(nèi)存結(jié)構(gòu)圖）中降低訪存的開(kāi)銷(xiāo)，這的確是一個(gè)很好的思路，但是這只能將訪存代價(jià)從幾百 cycle 降低到幾十 cycle，并不改變問(wèn)題的本質(zhì)。問(wèn)題的關(guān)鍵在于主體循環(huán)由兩條 Load 指令與一條 FMA 指令構(gòu)成，計(jì)算指令只占總體的 1/3，計(jì)算訪存比過(guò)低，最終導(dǎo)致了訪存延遲不能被隱藏，從而性能不理想。

登錄/注冊(cè)后可看大圖

讓我們打開(kāi)思路，若一個(gè) thread 并不只計(jì)算一個(gè)結(jié)果，而是計(jì)算 4x4 個(gè)結(jié)果，并且使用 Shared Memory 優(yōu)化，Hot Loop 會(huì)是什么樣呢，偽代碼如下所示：
        float c[4][4] = {{0}};
    float a_reg[4];
    float b_reg[4];
    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 中，需要進(jìn)行 4 次訪存操作，B 同理，那么主體的計(jì)算訪存指令比例變成了 16/8，相對(duì)于之前的情況，計(jì)算指令的占比大大提高了。足夠大的計(jì)算訪存比能提升計(jì)算單元的利用率，并能起到隱藏訪存延遲的作用。我們可以進(jìn)一步提升計(jì)算訪存比，從而使得 kernel 的性能接近理論峰值。
矩陣分塊與資源分配

顯然我們不能只使用一個(gè) block 計(jì)算一個(gè)超大矩陣，這樣會(huì)造成大量 SM（Streaming Multiprocessor）的閑置浪費(fèi)，這就需要對(duì)矩陣進(jìn)行分塊計(jì)算，如下圖所示：

登錄/注冊(cè)后可看大圖

不同的分塊大小在不同 shape 的矩陣乘法應(yīng)用上性能各有優(yōu)劣，本文選取 128x128 的分塊舉例。
從上一小節(jié)我們可以看到，提升計(jì)算訪存比有很大的好處，那么計(jì)算訪存比可以無(wú)限提升嗎，答案是否定的。因?yàn)橐嵘?jì)算訪存比，單個(gè) thread 就需要計(jì)算一個(gè)更大的塊，這就需要更多的寄存器，但寄存器的個(gè)數(shù)是有限的。以 Turing 架構(gòu)的 GPU 為例，單個(gè) SM 的寄存器總量為 65536，因?yàn)橹噶罹幋a的限制，單個(gè) thread 能使用的最大寄存器個(gè)數(shù)為 255，并且寄存器個(gè)數(shù)并不是用得越多越好。這里需要引入一個(gè) Occupancy（占用率）的概念，Occupancy 是指每個(gè) SM 中活動(dòng)線程束（Warp）數(shù)量與最大并發(fā)線程束數(shù)量的比值，高的 Occupancy 不一定意味著高性能，但可以通過(guò)切換執(zhí)行 Warp 來(lái)起到一定隱藏延遲的作用。而每個(gè) SM 中的 Active Warp 數(shù)量，取決于 block 使用的資源數(shù)量，具體為每個(gè)線程使用的寄存器個(gè)數(shù)與 Shared Memory 用量。Occupany可通過(guò) CUDA Toolkit 中提供的 CUDA_Occupancy_Calculator.xls 工具獲得。
考慮一個(gè) block 計(jì)算 128x128 的分塊，若每個(gè)線程計(jì)算 128 個(gè)結(jié)果，需要的 block size 為 128，單個(gè)線程需要 128 個(gè)寄存器儲(chǔ)存計(jì)算結(jié)果，加上所需的 Gmem to Smem，Smem to Reg 等一些所需的寄存器，大概共需要至少 180 多個(gè)，計(jì)算 Occupany 可知此時(shí)的 Active Warp 數(shù)只有 8，Occupany 為 25%；若設(shè)置 block size 為 256，則每個(gè)線程僅需計(jì)算 64 個(gè)結(jié)果，調(diào)整寄存器和 Shared Memory 的使用量并觀察 Occupany，可知若每個(gè)線程只使用 128 個(gè)寄存器，block 內(nèi)的 Shared Memory 使用量限制在 32K，Active Warp 數(shù)可以達(dá)到 16，是一個(gè)更優(yōu)的選擇：

登錄/注冊(cè)后可看大圖

并且此時(shí)的配置計(jì)算訪存比可以達(dá)到 64/4（使用向量讀取），已經(jīng)足夠隱藏訪存延遲。
極致的訪存優(yōu)化

通常情況下，在選取了合適的 block 資源配置，利用 Shared Memory 降低訪存延遲，做好循環(huán)展開(kāi)之后，SGEMM Kernel 的性能已經(jīng)能達(dá)到一個(gè)不錯(cuò)的水平（80% cublas），但這并不是我們旅程的終點(diǎn)。首先，我們可以使用向量讀取指令LDS.128優(yōu)化 Shared Memory 訪問(wèn)（對(duì)應(yīng) float4 數(shù)據(jù)類(lèi)型），這能大幅減少訪存指令的數(shù)量，進(jìn)一步提升計(jì)算訪存比，由此我們需要將 A 矩陣存入 smemA 之前做一次轉(zhuǎn)置：

登錄/注冊(cè)后可看大圖

同時(shí)，我們的 kernel 為 256 個(gè)線程計(jì)算 128x128 的分塊，為了能夠合并訪問(wèn) Shared Memory，我們將 256 個(gè)線程劃為二維，令：
    int tx = threadIdx.x % 16;   
    int ty = threadIdx.x / 16;
并按照如下方式向量讀取 Shared Memory 中的數(shù)據(jù)：

登錄/注冊(cè)后可看大圖

最終單個(gè)線程計(jì)算 2x2 個(gè) 4x4 的結(jié)果，結(jié)果布局如圖所示：

登錄/注冊(cè)后可看大圖

并且通過(guò) micro benchmark 可以探測(cè)出，Turing(Tesla T4) 的 Global Memory 的訪存延遲約 300 cycle，Shared Memory 的訪存延遲在約 30 cycle，需要充分利用 Prefetch 的思想，隱藏 Global Memory 讀入中間寄存器、將來(lái)自 Global Memory 的數(shù)據(jù)塊寫(xiě)入 Shared Memory、從 Shared Memory 中讀出數(shù)據(jù)塊的訪存延遲，以免計(jì)算單元因?yàn)?stall 而空閑太久，最終的偽代碼如下所示：
    #define TILE_K 16
    __shared__ float4 smemA[2][TILE_K * 128 / 4];
    __shared__ float4 smemB[2][TILE_K * 128 / 4];
    float4 c[8][2] = {{make_float4(0.f, 0.f, 0.f, 0.f)}};
    float4 ldg_a_reg[2];
    float4 ldg_b_reg[2];
    float4 a_reg[2][2];
    float4 b_reg[2][2];

    // 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[0]);
    store_reg_to_smem_tile(ldg_b_reg, 0, smemB[0]);
    __syncthreads();

    // load first tile from shared mem to register
    load_smem_tile_to_reg(smemA[0], 0, a_reg[0]);
    load_smem_tile_to_reg(smemB[0], 0, b_reg[0]);

    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[load_stage_idx], j + 1, a_reg[(j + 1) % 2]);
            load_smem_tile_to_reg(smemB[load_stage_idx], j + 1, b_reg[(j + 1) % 2]);
            // compute matrix multiply accumulate 8x8
            mma8x8(a_reg[j % 2], b_reg[j % 2], c)；
        }

        if(i < K) {
            // store next tile to shared mem
            store_reg_to_smem_tile_transpose(ldg_a_reg, 0, smemA[write_stage_idx]);
            store_reg_to_smem_tile(ldg_b_reg, 0, smemB[write_stage_idx]);
            // 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[load_stage_idx ^ 1], 0, a_reg[0]);
        load_smem_tile_to_reg(smemB[load_stage_idx ^ 1], 0, b_reg[0]);
        // compute last tile mma 8x8
        mma8x8(a_reg[1], b_reg[1], c)；
    } while (i < K);

    store_c(c, C);
    注：此處偷懶假設(shè)了 M、N、K 都是 4 的倍數(shù)，若非 4 的倍數(shù)則 Global Memory 不能使用 float4 進(jìn)行讀取，結(jié)果也不能用 float4 進(jìn)行寫(xiě)回，而且為了合并寫(xiě)回，需要通過(guò) Shared Memory 交換 warp 內(nèi)的結(jié)果，保證每個(gè) warp 執(zhí)行一條 Store 指令能夠?qū)懟匾黄B續(xù)的內(nèi)存空間。
至此我們獲得了一個(gè)充分優(yōu)化的 SGEMM Kernel。另外 Ampere GPU 新增了LDGSTS指令，數(shù)據(jù)塊從 Global Memory 到 Shared Memory 的過(guò)程不需要經(jīng)過(guò)中間寄存器，可以進(jìn)一步的優(yōu)化 SGEMM 的性能。
性能對(duì)比

為了避免 cublas 選取到 split K 的 Kernel，我們將 K 固定為 1024，取 M, N = 2048, 4096, 8192 和 16384 作為測(cè)試用例，對(duì)比了上述 SGEMM Kernel 與 cublas 的性能（測(cè)試 GPU 為 Tesla T4，鎖定核心頻率為 1100)：

登錄/注冊(cè)后可看大圖

可以看到所實(shí)現(xiàn)的 SGEMM Kernel 達(dá)到了 cublas 平均 97.5% 的性能。
超越 cublas：使用 SASS 調(diào)優(yōu) Kernel
到這里，可能有同學(xué)依然有一個(gè)疑問(wèn)，我們似乎把所有能想到的優(yōu)化手段都用上了，為什么寫(xiě)出來(lái)的 CUDA C Kernel 依然離 cublas 有一定的差距，答案是 cublas 所使用的 kernel 中有一大部分并不是通過(guò) nvcc 編譯的 CUDA Kernel，而是使用 NVIDIA GPU 的匯編語(yǔ)言（Shader Assembly，簡(jiǎn)稱(chēng) SASS）編寫(xiě)的深度調(diào)優(yōu)版本。
盡管 nvcc 編譯器在不斷的進(jìn)步，特別是 CUDA 11 中的 nvcc，所編譯的 Kernel 與手工匯編優(yōu)化版本之間的差距已大幅縮小，但依舊無(wú)法完全避免寄存器 Bank conflict 的影響以及充分利用寄存器的 Reuse Cache（這兩個(gè)概念下面會(huì)進(jìn)行詳細(xì)的介紹），使得差距依舊存在。即使 PTX 這樣的偽匯編語(yǔ)言，也無(wú)法精確控制寄存器的分配，和 CUDA C 面臨著一樣的困境。
所以為了充分挖掘 GPU 的性能極限，需要對(duì) GPU 指令和寄存器進(jìn)行精確控制，就必須交由 GPU 原生匯編語(yǔ)言 SASS 完成。這方面已經(jīng)有了很多研究，如出自 Citadel 的深入研究 NV GPU 架構(gòu)的 Dissecting the NVidia XXX GPU architecture via microbenchmarking 系列論文，這一系列文章對(duì)底層架構(gòu)做了系統(tǒng)的測(cè)試、分析和總結(jié)，雖然其中某些結(jié)論可能并不準(zhǔn)確，但總體來(lái)講有很高的參考價(jià)值。同時(shí)催生了不少開(kāi)源匯編器如 KeplerAs、maxas（最成熟，影響深遠(yuǎn)）、turingas 和 CuAssembler 等一系列開(kāi)源 SASS 匯編器，使得使用 SASS 編寫(xiě)高性能 Kernel 變成了可能。
寄存器 Bank conflict

我們知道 Shared Memory 有 Bank conflict，而寄存器的 Bank conflict 也是類(lèi)似的概念。NVIDIA GPU 每個(gè) SM 有獨(dú)立的 Register File，而 Register File 被分為若干個(gè) Bank，以 Maxwell 為例，若一條指令所需的源寄存器有 2 個(gè)以上來(lái)自于同一 Bank，則會(huì)產(chǎn)生 conflict，指令會(huì)相當(dāng)于重發(fā)射，浪費(fèi)一個(gè) cycle。Maxwell/Pascal 的 Register File 的 Bank 數(shù)為 4，寄存器的id%4即為該寄存器的所屬 bank（如 R0 屬于 Bank 0，R5 屬于 Bank 1），FFMA R1, R0, R4, R1這樣的指令就會(huì)產(chǎn)生寄存器 Bank conflict。而 Turing 架構(gòu)做了改進(jìn)，Register File 被分為 2 個(gè) Bank，每個(gè) Bank 有 2 個(gè) Port，若非三個(gè)源寄存器 id 同奇偶則不會(huì)產(chǎn)生沖突，大大緩解了寄存器 Bank conflict。
maxas 中的 Maxwell SGEMM SASS Kernel 為了緩解寄存器 Bank conflict，就對(duì)參與 FFMA 計(jì)算的寄存器做了精巧的分配（參考 maxas 的 SGEMM 文檔），如下圖所示：

登錄/注冊(cè)后可看大圖

經(jīng)過(guò)對(duì) C 的巧妙排布，寄存器 Bank conflict 大大減少，但依然無(wú)法完全避免（如上圖中黑框標(biāo)識(shí)的部分，A/B 所使用的寄存器會(huì)產(chǎn)生 Bank conflict），這部分沖突就需要用到寄存器 Reuse 來(lái)消除。
Register Reuse

寄存器 Reuse 是 NVIDIA 為了緩解寄存器 Bank conflict 的問(wèn)題，在 Maxwell 開(kāi)始引入的一種機(jī)制，NVIDIA 在讀取指令操作數(shù)的 Collector 單元加入了寄存器的 Reuse Cache。Reuse Cache 是只讀的，指令獲取 Operand 是否通過(guò)此 Cache 由該指令的 control code（maxas 的 control code wiki中有詳細(xì)的介紹）所指定，使用 cuobjdump 反匯編一些 Kernel 可以發(fā)現(xiàn)一些寄存器后有 .reuse的 flag，即表示該寄存器從 Reuse Cache 而非 Register File 中取值，從而消除寄存器 Bank conflict：
# Maxwell GPU但是使用 .reuse需要滿(mǎn)足一定條件（寄存器將被改寫(xiě)前不能設(shè)置 .reuse），胡亂設(shè)置 reuse flag 會(huì)有可能獲取的是歷史值，造成計(jì)算錯(cuò)誤，根據(jù)筆者的理解，.reuse 更像是使該寄存器的值在 Reuse Cache 中 hold 住的標(biāo)識(shí)。nvcc 編譯 CUDA Kernel 也會(huì)使用 Reuse Cache 去規(guī)避一些寄存器 Bank conflict，但是因?yàn)榧拇嫫鞣峙浼爸噶钆挪嫉脑?，Reuse 的利用率并不高，反匯編我們剛才寫(xiě)的 SGEMM Kernel，對(duì)主循環(huán)的所有 FFMA 指令做個(gè)統(tǒng)計(jì)，可以發(fā)現(xiàn) Reuse Cache 僅達(dá)到 20% 左右，而 maxas 的 SASS Kernel 通過(guò)設(shè)計(jì)使得 Reuse 的利用率可以達(dá)到 49%。
最終通過(guò) SASS 精細(xì)調(diào)優(yōu)的 SGEMM Kernel 的性能可以全面超越 cublas，感興趣的同學(xué)們可以自行編譯 maxas 中的 SGEMM Kernel 在 Maxwell 或者 Pascal GPU 上進(jìn)行測(cè)試。最后，雖然使用 SASS 能充分挖掘 GPU 的性能，但面臨有三大問(wèn)題：1. 第三方 NV GPU 匯編器依賴(lài)于對(duì) GPU 架構(gòu)的逆向研究，可能因?yàn)闆](méi)有探究到全部的硬件底層細(xì)節(jié)而存在未知的 BUG；2. 匯編 Kernel 難于開(kāi)發(fā)，更難于調(diào)試；3. NV 每一代 GPU 的 ISA（指令集）都不盡相同，需要不斷開(kāi)發(fā)對(duì)應(yīng)的匯編器和匯編 Kernel。正因?yàn)檫@幾大問(wèn)題的存在，使得使用 SASS 編寫(xiě) Kernel 是個(gè)費(fèi)時(shí)費(fèi)力的工作，除非有追求極致性能的需求，否則不建議輕易嘗試。
GEMM 的延伸：優(yōu)化卷積運(yùn)算

我們都知道優(yōu)化卷積運(yùn)算可以通過(guò) im2col 將卷積映射為矩陣乘法來(lái)實(shí)現(xiàn)，對(duì)于上述 SGEMM Kernel，只需要將 Global Memory 的數(shù)據(jù)搬運(yùn)到 Shared Memory 這一過(guò)程稍作修改，由對(duì)應(yīng)位置的映射變?yōu)?im2col 映射，SGEMM Kernel 就搖身一變成為了計(jì)算 Conv 的 Kernel，這即是 cudnn 卷積運(yùn)算的 Implicit Gemm 算法。而在 im2col 過(guò)程中，若直接計(jì)算指針的偏移量的話，會(huì)引入大量的整數(shù)除法和取余運(yùn)算，這是一筆不小的開(kāi)銷(xiāo)，所以可以將地址的偏移量在 host 端預(yù)先計(jì)算好，作為 param 傳入 kernel 中，則可以在需要時(shí)從常量?jī)?nèi)存中讀取，避免整數(shù)除法和取余，實(shí)現(xiàn) Implicit Precomp Gemm。
總結(jié)

本文詳細(xì)介紹了如何編寫(xiě)一個(gè)高效率的 CUDA SGEMM Kernel，并且介紹了使用 SASS 編程這一極限優(yōu)化性能的手段，并稍稍延伸展開(kāi)了通過(guò) Implicit Gemm 優(yōu)化卷積運(yùn)算的思路，希望可以給予有志于極致挖掘硬件性能的同學(xué)們一定的啟發(fā)
RISC-V將贏得下一輪架構(gòu)之爭(zhēng)？

在全球市場(chǎng)上，芯片指令集呈現(xiàn)雙寡頭格局，基于X86和ARM架構(gòu)的處理器長(zhǎng)期占據(jù)絕大多數(shù)市場(chǎng)份額，X86架構(gòu)在PC及服務(wù)器市場(chǎng)一家獨(dú)大，移動(dòng)市場(chǎng)則由ARM架構(gòu)一統(tǒng)江湖。
在這樣一個(gè)格局中，中下游廠商大多只能在這二者之間選擇，但是ARM授權(quán)費(fèi)用昂貴，傳統(tǒng)X86的授權(quán)又過(guò)于復(fù)雜，業(yè)界一直期待在CPU架構(gòu)領(lǐng)域能有更多選擇。
隨著AIoT時(shí)代的到來(lái)，RISC-V架構(gòu)開(kāi)放、靈活、模塊化，特別適合滿(mǎn)足AIoT市場(chǎng)場(chǎng)景碎片化、差異化的市場(chǎng)需求，產(chǎn)業(yè)界普遍認(rèn)為它有望成為下一代廣泛應(yīng)用的處理器架構(gòu)。Semico Research預(yù)測(cè)，到2025年，RISC-V市場(chǎng)規(guī)模將超10億美元。

登錄/注冊(cè)后可看大圖

近日，行業(yè)大牛Jim Keller在一次公開(kāi)演講中也直言，未來(lái)是屬于RISC-V的。

登錄/注冊(cè)后可看大圖

Jim keller認(rèn)為，RISC-V的開(kāi)放式生態(tài)讓其擁有廣闊的前景和巨大潛力，相比X86和ARM架構(gòu)，RISC-V雖然在性能、生態(tài)等問(wèn)題上存在不足，但憑借開(kāi)放式標(biāo)準(zhǔn)，RISC-V的發(fā)展速度遠(yuǎn)超X86及ARM。
RISC-V采用的開(kāi)放標(biāo)準(zhǔn)讓各個(gè)廠商可以在架構(gòu)基礎(chǔ)上進(jìn)行大幅度的修改和定制，不需要擔(dān)心后續(xù)會(huì)因?yàn)楦鞣N原因而失去授權(quán)或影響到產(chǎn)品的生產(chǎn)使用。同時(shí)，RISC-V存在架構(gòu)優(yōu)勢(shì)，使其可以同時(shí)適應(yīng)極低功耗及高性能處理器的設(shè)計(jì)需求，這一點(diǎn)也是X86和ARM無(wú)法做到的。
當(dāng)Jim keller看好RISC-V時(shí)，這無(wú)疑為其平添了幾分可能性。對(duì)Jim keller有了解的讀者應(yīng)該知道，Jim keller過(guò)去創(chuàng)造過(guò)許多成功的處理器方案，幾乎每一次跳槽，Jim keller都能刷新對(duì)應(yīng)處理器的行業(yè)上限。
“硅仙人”的傳奇履歷
在其20多年的工作經(jīng)歷中，Jim keller先后在DEC、AMD、博通、蘋(píng)果、特斯拉、Intel等公司任職，一次又一次創(chuàng)造了被認(rèn)為是芯片行業(yè)中，至關(guān)重要的領(lǐng)先部件。
上世紀(jì)九十年代，Jim Keller在DEC工作，涉足了Alpha處理器項(xiàng)目的設(shè)計(jì)，產(chǎn)品包括21164與22164。Jim Keller的本段經(jīng)歷使其深入了解RISC的設(shè)計(jì)理念，并通過(guò)Alpha架構(gòu)，奠定了架構(gòu)設(shè)計(jì)能力。Alpha架構(gòu)的設(shè)計(jì)理念深度影響了后續(xù)AMD、Intel等芯片廠商的架構(gòu)設(shè)計(jì)路線。
1998年，Jim keller從DEC離職，加盟AMD，主導(dǎo)了K7/K8/K12處理器的研發(fā)，其中K8處理器讓AMD第一次擁有了對(duì)抗Intel的資本，X86架構(gòu)不再是Intel一家獨(dú)大，同時(shí)推動(dòng)AMD的64位X86架構(gòu)處理器落地，使得AMD在技術(shù)路線上第一次超越Intel。
1999年，Jim Keller加入SiByte，研制MIPS的網(wǎng)絡(luò)處理器。SiByte后被博通收購(gòu)，Jim Keller擔(dān)任首席架構(gòu)師，直到2004年離職，加入PASemi任工程副總裁，轉(zhuǎn)型研究RISC產(chǎn)品，使MIPS的構(gòu)建速度可與專(zhuān)用ASIC相媲美。
PASemi被蘋(píng)果收購(gòu)后，Jim Keller開(kāi)始主導(dǎo)蘋(píng)果的A系列處理器研發(fā)計(jì)劃，為蘋(píng)果A芯片領(lǐng)先全球，奠定了架構(gòu)基礎(chǔ)。
完成A系列處理器的設(shè)計(jì)后，Jim keller再次跳槽，重返A(chǔ)MD。自從Jim keller離開(kāi)AMD后，AMD后續(xù)的Bulldozer系列處理器一度被稱(chēng)為最糟糕的X86處理器，AMD依靠K8處理器搶到的市場(chǎng)份額幾乎完全丟失，甚至一度瀕臨破產(chǎn)。
Jim keller回歸之后，將Bulldozer系列的后續(xù)研發(fā)計(jì)劃全部推倒，重新設(shè)計(jì)了一個(gè)全新的處理器架構(gòu)，這就是大家都不陌生的Zen架構(gòu)。Zen架構(gòu)讓AMD一舉翻身，在性能排行榜上壓制Intel數(shù)年，甚至一度在處理器出貨量上超越Intel，基于Zen系列，AMD后續(xù)還拓展了EYPC等系列處理器，開(kāi)始全方位與Intel競(jìng)爭(zhēng)。
第二次拯救AMD于水火之中后，Jim Keller于2016年加入特斯拉擔(dān)任Autopilot部門(mén)的負(fù)責(zé)人，為其打造自動(dòng)駕駛AI芯片，并在正式落地量產(chǎn)前的2018年離職。
2018年，Jim Keller來(lái)到Intel，擔(dān)任高級(jí)副總裁、系統(tǒng)架構(gòu)師，受限于保密協(xié)議，Jim keller負(fù)責(zé)的項(xiàng)目沒(méi)有曝光。但從Intel進(jìn)展來(lái)看，2019年5月，Intel發(fā)布全新的ice lake架構(gòu)。據(jù)悉，Intel 12代酷睿所使用的大小核設(shè)計(jì)就由Jim Keller主導(dǎo)開(kāi)發(fā)，而12代酷睿的誕生被譽(yù)為Intel的翻身一戰(zhàn)，再次壓制了AMD成為x86架構(gòu)的新王者。

登錄/注冊(cè)后可看大圖

回顧Jim keller的經(jīng)歷，幾乎每一款經(jīng)典處理器都改寫(xiě)了行業(yè)上限，從x86到ARM，似乎就沒(méi)有他不了解的處理器架構(gòu)。
然而，同樣是兩年后，Jim keller從Intel離職，這次他選擇了創(chuàng)業(yè)，擔(dān)任Tenstorrent的聯(lián)合創(chuàng)始人及CTO，主攻RISC-V。
也正因如此，當(dāng)這位曾在 x86 和 ARM 架構(gòu)芯片領(lǐng)域取得過(guò)較高成就的硅谷傳奇工程師，選擇調(diào)轉(zhuǎn)航向，親自下場(chǎng)押注RISC-V架構(gòu)，不免讓業(yè)界對(duì)RISC-V多了幾分信心和慎重。
卡位RISC-V賽道
中國(guó)科學(xué)院軟件研究所總工程師武延軍表示，從2010年發(fā)布到現(xiàn)在，RISC-V指令集只走過(guò)12年的時(shí)間，但它在這期間取得的成就，遠(yuǎn)遠(yuǎn)超過(guò)同時(shí)期的X86和ARM。
RISC-V開(kāi)源ISA于2016年首次推出，但最初的內(nèi)核僅適用于微控制器和一些基本的片上系統(tǒng)設(shè)計(jì)。然而，經(jīng)過(guò)幾年的發(fā)展，眾多芯片開(kāi)發(fā)商已經(jīng)針對(duì)云數(shù)據(jù)中心、AI 工作負(fù)載和高級(jí)存儲(chǔ)應(yīng)用等多領(lǐng)域創(chuàng)建了設(shè)計(jì)，其生態(tài)獲得了較大的豐富，這足以證明RISC-V的實(shí)力。
前不久，RISC-V基金會(huì)CEO Calista Redmond表示，RISC-V基金會(huì)已經(jīng)在全球70多個(gè)國(guó)家吸引了超過(guò)3100個(gè)會(huì)員加入，會(huì)員數(shù)量在2021年暴增了134%。RISC-V產(chǎn)業(yè)在過(guò)去幾年里取得的初步成果，搭載RISC-V內(nèi)核的芯片出貨量已經(jīng)突破了100億。

登錄/注冊(cè)后可看大圖

ARM用了17年才完成這個(gè)目標(biāo)，而RISC-V基金會(huì)成立至今僅7年。由此可見(jiàn)，RISC-V這個(gè)新興架構(gòu)正在吸引著全球眾多參與者躬身其中，紛紛卡位RISC-V賽道。
從國(guó)外廠商動(dòng)態(tài)來(lái)看，蘋(píng)果積極轉(zhuǎn)向RISC-V陣營(yíng)，Intel業(yè)務(wù)也延伸至此，包括Jim Keller坐鎮(zhèn)Tenstorrent、Microchip、瑞薩電子、Nordic、ST等國(guó)際廠商都在積極布局，開(kāi)發(fā)基于RISC-V指令集的處理器。
同時(shí)，還有一眾RISC-V IP公司深耕于此，包括SiFive、Codasip、imagination、MIPS等，其中值得注意的是，SiFive IP芯片累計(jì)出貨量超10億；搭載Codasip的IP芯片出貨量已達(dá)20億顆；作為業(yè)內(nèi)成熟的IP供應(yīng)商，Imagination進(jìn)入RISC-V市場(chǎng)再次突顯了RISC-V架構(gòu)的擴(kuò)展特性，為RISC-V提供了更大的發(fā)展動(dòng)力。
除了歐美國(guó)家之外，RISC-V在印度、巴基斯坦等國(guó)家也在被廣泛采用，行業(yè)廠商積極布局。
能夠看到，國(guó)外廠商在RISC-V上的布局在逐漸呈現(xiàn)，除了芯片設(shè)計(jì)、IP廠商之外，包括IAR、verify、Imperas、Embecosm等工具鏈、安全軟件算法、模擬器和編譯器等產(chǎn)業(yè)鏈企業(yè)的支持，都在加速RISC-V架構(gòu)的發(fā)展和成熟。
RISC-V，在國(guó)內(nèi)市場(chǎng)掀起風(fēng)暴
中國(guó)RISC-V市場(chǎng)發(fā)展迅猛，憑借國(guó)內(nèi)龐大的半導(dǎo)體市場(chǎng)，中國(guó)企業(yè)可以迅速對(duì)RISC-V芯片進(jìn)行迭代更新，并且依據(jù)用戶(hù)使用反饋進(jìn)行優(yōu)化，這些都是國(guó)外初創(chuàng)企業(yè)難以做到的。全球100億顆RISC-V核心出貨量中，預(yù)計(jì)中國(guó)占比超過(guò)1/3。
中國(guó)社區(qū)作為工作小組和研發(fā)伙伴，一直是RISC-V最強(qiáng)有力的貢獻(xiàn)者之一。在RISC-V基金會(huì)初期的19家高級(jí)會(huì)員企業(yè)中，有整整12家來(lái)自中國(guó)，華為、阿里巴巴、紫光展銳、中興等企業(yè)都是早期高級(jí)會(huì)員，對(duì)RISC-V的發(fā)展有著舉足輕重的作用。
RISC-V對(duì)于中國(guó)的半導(dǎo)體企業(yè)來(lái)說(shuō)是一個(gè)難得的機(jī)會(huì)，因?yàn)镽ISC-V起步晚所以大家都還處在一個(gè)起跑線上，而且開(kāi)放式標(biāo)準(zhǔn)也讓國(guó)內(nèi)企業(yè)可以更安心的發(fā)展RISC-V，不用擔(dān)心未來(lái)某一天突然失去授權(quán)。
在日前舉辦的RISC-V峰會(huì)上，倪光南院士在致辭中表示：“目前，主流CPU市場(chǎng)仍被X86、ARM架構(gòu)壟斷，但新興的開(kāi)源指令集RISC-V將為我國(guó)芯片產(chǎn)業(yè)發(fā)展提供新機(jī)遇，如果抓住機(jī)會(huì)，就有可能在CPU核心技術(shù)上掌握主動(dòng)權(quán)?！?br /> 在此機(jī)遇和趨勢(shì)下，國(guó)內(nèi)涌現(xiàn)出一批RISC-V玩家，包括華為海思、全志科技、兆易創(chuàng)新、樂(lè)鑫科技、北京君正、中微半導(dǎo)體、匯頂科技、凌思微電子、先楫半導(dǎo)體、華米科技、沁恒微電子、芯晟科技、愛(ài)普特微電子、晶視智能、啟英泰倫、方寸微電子、中科藍(lán)汛、航順芯片、飛思靈微電子、博流智能以及中科院計(jì)算所等在內(nèi)的芯片廠商和科研機(jī)構(gòu)都相繼推出了基于RISC-V架構(gòu)的SoC、MCU、DSP芯片等產(chǎn)品。
此外，阿里平頭哥、芯來(lái)科技、賽昉科技、睿思芯科等可以提供基于RISC-V架構(gòu)的處理器IP、編譯器、工具鏈等產(chǎn)品。據(jù)悉，阿里平頭哥玄鐵系列的RISC-V芯片已經(jīng)供應(yīng)一百五十多個(gè)企業(yè)和客戶(hù)，超過(guò)500個(gè)授權(quán)使用，出貨量已經(jīng)達(dá)到25億顆；芯來(lái)科技RISC-V CPU IP產(chǎn)品線覆蓋從低功耗到高性能的各種應(yīng)用需求，正式授權(quán)客戶(hù)已超過(guò)100家...
一些列動(dòng)態(tài)和成果，呈現(xiàn)出國(guó)內(nèi)RISC-V芯片行業(yè)上下游產(chǎn)業(yè)鏈的發(fā)展速度和全面布局態(tài)勢(shì)。能夠看到，在這些中國(guó)技術(shù)人和企業(yè)的努力下，RISC-V處理器的可用性及覆蓋面大大提升，使其從一個(gè)象牙塔的基礎(chǔ)架構(gòu)，真正走向市場(chǎng)化發(fā)展。
國(guó)產(chǎn)RISC-V的優(yōu)勢(shì)和挑戰(zhàn)
隨著半導(dǎo)體市場(chǎng)進(jìn)入一個(gè)新的階段，中國(guó)半導(dǎo)體行業(yè)需要抓住一切機(jī)遇，尋找突破科技封鎖的方法，RISC-V或許將成為一個(gè)突破口。
武延軍對(duì)此也表示認(rèn)同。他指出，對(duì)于中國(guó)發(fā)展RISC-V來(lái)說(shuō)，很明顯的優(yōu)勢(shì)首先是一個(gè)巨量的市場(chǎng)；然后有大規(guī)模的計(jì)算機(jī)工程師團(tuán)隊(duì)和專(zhuān)業(yè)人才；同時(shí)有足夠豐富的應(yīng)用場(chǎng)景能夠充分發(fā)揮RISC-V的多樣性。
“還有一個(gè)獨(dú)特的優(yōu)勢(shì)，那就是我們可能沒(méi)有國(guó)外一些老牌廠商和機(jī)構(gòu)的歷史包袱，沒(méi)有它們已經(jīng)存在的商業(yè)利益或者技能上的慣性，使我們有機(jī)會(huì)從一個(gè)全新的視角、一個(gè)全新的架構(gòu)下面去做一些開(kāi)創(chuàng)性工作?！?武延軍補(bǔ)充道。
不過(guò)，相應(yīng)的劣勢(shì)就是中國(guó)在很多技術(shù)領(lǐng)域沒(méi)有國(guó)外巨頭在歷史上積累的豐富經(jīng)驗(yàn)。同時(shí)，在高端核心人才儲(chǔ)備上也沒(méi)有像國(guó)外處理器、基礎(chǔ)軟件巨頭那么多。此外，可能在專(zhuān)利等知識(shí)產(chǎn)權(quán)方面也存在一定差距。
因此，對(duì)于RISC-V在中國(guó)的發(fā)展，除了需要面對(duì)制造的挑戰(zhàn)以外，還需要克服處理器IP核和核心基礎(chǔ)軟件兩方面的困難。其中，后者的編譯工具鏈和操作系統(tǒng)是RISC-V生態(tài)里面最難突破的兩個(gè)技術(shù)點(diǎn)，跟國(guó)際水平存在差距。
平頭哥半導(dǎo)體副總裁孟建熠對(duì)此指出，RISC-V還需要大家把整個(gè)架構(gòu)不斷的做得穩(wěn)定、可靠；有了“穩(wěn)定、可靠”之后整個(gè)上層的軟件適配就會(huì)變得更加容易，軟件越來(lái)越豐富之后要做的事情就是往各個(gè)縱深的方向去發(fā)展。目前來(lái)看，整個(gè)發(fā)展勢(shì)頭還是非常好的。
生態(tài)，RISC-V勝負(fù)關(guān)鍵手
RISC-V的潛力，正讓很多芯片公司感到興奮，國(guó)內(nèi)外新興RISC-V企業(yè)和原有科技巨頭紛紛布局卡位。
據(jù)華泰證券研報(bào)，開(kāi)源指令集架構(gòu)RISC-V在AIoT時(shí)代應(yīng)用場(chǎng)景及市場(chǎng)空間將快速增長(zhǎng)，預(yù)計(jì)2025年全球RISC-V出貨量將超過(guò)600億顆。RISC-V似乎爆發(fā)在即，生態(tài)建設(shè)成為勝負(fù)手關(guān)鍵。
從市場(chǎng)現(xiàn)狀來(lái)看，RISC-V目前只是在MCU領(lǐng)域稍有建樹(shù)，局限于一些對(duì)算力和對(duì)軟件生態(tài)豐富程度要求不那么高的領(lǐng)域。在往更高性能的處理器方面，RISC-V目前似乎與Arm和X86相比，差距甚大。
針對(duì)其應(yīng)用領(lǐng)域，RISC-V基金會(huì)CTO Mark Himelstein透露，RISC-V在未來(lái)幾年將會(huì)發(fā)力汽車(chē)、數(shù)據(jù)中心和安全等多個(gè)應(yīng)用領(lǐng)域，包括最近備受關(guān)注的加速器和智能網(wǎng)卡也是RISC-V未來(lái)幾年的關(guān)注重點(diǎn)。

登錄/注冊(cè)后可看大圖

孟建熠強(qiáng)調(diào)，RISC-V現(xiàn)在應(yīng)該在一個(gè)初步證明它商業(yè)化可用的這條路上，其發(fā)展速度非?？?。和Arm一樣，RISC-V在早期的時(shí)候首先在嵌入式領(lǐng)域先打開(kāi)一個(gè)環(huán)節(jié)，基本證明可用。接下來(lái)，可能就是它要往更加高性能、更加縱深的方向去發(fā)展。
比如數(shù)據(jù)中心，就是RISC-V正在努力的方向。
然而，近日Arm高管在接受采訪時(shí)試圖淡化RISC-V對(duì)ARM業(yè)務(wù)的威脅。Arm基礎(chǔ)設(shè)施業(yè)務(wù)高級(jí)副總裁兼總經(jīng)理Chris Bergey表示：“雖然RISC-V可能在某些市場(chǎng)上獲得牽引力，但Arm并不擔(dān)心新貴ISA會(huì)侵蝕其在數(shù)據(jù)中心的新立足點(diǎn)?！?br /> RISC-V要進(jìn)入數(shù)據(jù)中心，確實(shí)還有很大的距離。孟建熠表示，要往高性能發(fā)展，處理器首先本身要在性能上有突破?！靶阅芡黄啤本托枰度氪罅康难邪l(fā)精力和資金投入；其次，往高性能走，處理器的“穩(wěn)定性”也是非常重要的。上層軟件棧越厚，軟件與硬件的協(xié)同優(yōu)化就越困難，必須要求足夠的硬件穩(wěn)定性，來(lái)適配更復(fù)雜的應(yīng)用場(chǎng)景。
武延軍也認(rèn)為，RISC-V往高性能走，對(duì)處理器的設(shè)計(jì)和軟件生態(tài)都是一個(gè)非常大的挑戰(zhàn)。
他舉例表示，中科院軟件所在國(guó)際上游開(kāi)源社區(qū)，國(guó)內(nèi)開(kāi)源歐拉社區(qū)、平頭哥等都做了大量的RISC-V基礎(chǔ)軟件適配工作?！暗覀兛梢钥吹剑壳斑€是有很多核心的基礎(chǔ)軟件不能很好地跑在RISC-V平臺(tái)上?！睋?jù)了解，這可能會(huì)與指令集規(guī)范還不成熟有關(guān)，但更多的是因?yàn)檫@些基礎(chǔ)軟件包之前都是在X86和ARM上面去跑，從維護(hù)者、社區(qū)的角度，還沒(méi)有把RISC-V當(dāng)成Tier1或者First-Class-Citizen度去對(duì)待。這有理念問(wèn)題，有投入問(wèn)題，也有商業(yè)利益回報(bào)的問(wèn)題。
因此，基礎(chǔ)軟件的意義是避免生態(tài)碎片化，避免大家在同樣的指令集架構(gòu)下還去“造輪子”。因此在一些基礎(chǔ)的編譯工具鏈、操作系統(tǒng)，比如GCC/LLVM，Linux等，希望整個(gè)行業(yè)能夠合力去共同打造，而不是說(shuō)每個(gè)RISC-V處理器廠家都要自己做一套，這其實(shí)是一個(gè)很大的浪費(fèi)。
目前，國(guó)外已經(jīng)有好多公司在往數(shù)據(jù)中心的方向發(fā)展了，它們也是整個(gè)賽道里面非常熱門(mén)的公司，在大家努力下，整個(gè)RISC-V在數(shù)據(jù)中心的生態(tài)正在逐步建設(shè)當(dāng)中。
寫(xiě)在最后
從生態(tài)繁榮程度來(lái)看，RISC-V目前可能處在ARM大概在2000-2005年的階段。彼時(shí)，有關(guān)ARM的各種培訓(xùn)資料、書(shū)籍、培訓(xùn)班在大量的出現(xiàn)，市場(chǎng)上對(duì)ARM開(kāi)發(fā)者的人才需求也是突然爆增，在嵌入式領(lǐng)域已經(jīng)跟X86、MIPS已經(jīng)展開(kāi)了非常激烈的競(jìng)爭(zhēng)。
武延軍指出，現(xiàn)在RISC-V也差不多到了這個(gè)階段，而且以一個(gè)更高的加速度在增長(zhǎng)。我們看到不光是教育科研、各種文檔材料、人才培訓(xùn)，而且已經(jīng)有非常好的商用場(chǎng)景、非常高的出貨量。當(dāng)前階段，可能也是整個(gè)生態(tài)上下游都可以開(kāi)始發(fā)力的階段了，如果說(shuō)之前還在觀望的話，我覺(jué)得從這個(gè)階段開(kāi)始大家已經(jīng)可以不用去觀望了，不管是從個(gè)人學(xué)習(xí)能力提升角度，亦或是從公司業(yè)務(wù)發(fā)展的角度，都可以放心地去投入RISC-V領(lǐng)域了。
雖然Arm今天可能不將RISC-V視為威脅，但也正如我們觀察到的，隨著RISC-V的發(fā)展，Arm已經(jīng)改變了其內(nèi)核許可的方式。例如，面對(duì)免版稅、開(kāi)放和模塊化的競(jìng)爭(zhēng)對(duì)手，Arm已采取措施，在某些情況下降低其IP許可的前期成本，并允許被許可人添加自定義指令。
回過(guò)頭看，無(wú)論是Wintel（Windows+Intel），還是AA聯(lián)盟（Android +Arm），以及那些早已凋零的MIPS、PowerPC等芯片架構(gòu)，決定競(jìng)爭(zhēng)終局的或許從來(lái)不是技術(shù)本身，而是豐富的上層應(yīng)用、高效的軟硬協(xié)同以及爆發(fā)的市場(chǎng)需求。
總結(jié)起來(lái)就是兩個(gè)字：生態(tài)。
生態(tài)的培育，是一個(gè)漫長(zhǎng)的過(guò)程。對(duì)于RISC-V未來(lái)的生態(tài)發(fā)展趨勢(shì)，孟建熠強(qiáng)調(diào)：“從競(jìng)爭(zhēng)性角度看，Arm架構(gòu)和x86架構(gòu)分別在移動(dòng)終端、PC和服務(wù)器市場(chǎng)壟斷多年，在這些領(lǐng)域RISC-V新玩家滲透進(jìn)去還非常需要時(shí)日。但是在AIoT、新能源汽車(chē)電子、異構(gòu)計(jì)算等新興領(lǐng)域，RISC-V和其他架構(gòu)站在同一起跑線，反而具備一些巨頭們不具備的新起跑優(yōu)勢(shì)?！?br /> 展望未來(lái)，RISC-V與Arm和X86將會(huì)在競(jìng)爭(zhēng)過(guò)程中逐漸找到自己的定位，不斷融合、互相借鑒、長(zhǎng)期共存，形成“三分天下”的局面，在各自擅長(zhǎng)的領(lǐng)域發(fā)揮優(yōu)勢(shì)。
再談DDR內(nèi)存技術(shù)原理

從動(dòng)態(tài)的角度來(lái)分析時(shí)序結(jié)構(gòu)，包括read/write的整個(gè)過(guò)程到數(shù)據(jù)返回發(fā)生了什么，其中包括以下內(nèi)容：
·DDR是如何完成Read、Write 等的操作
·DDR的基本命令
·DDR的時(shí)序參數(shù)
·DDR的性能分析
中國(guó)存儲(chǔ)器芯片行業(yè)概覽
1 DRAM基本組成

登錄/注冊(cè)后可看大圖

對(duì)于DRAM，其主要由行和列組成，每一個(gè)bit中都是由類(lèi)似右下圖的類(lèi)晶體管的結(jié)構(gòu)組成，對(duì)于sdram的數(shù)據(jù)，可以通過(guò)控制column和row就可以訪問(wèn)sdram的隨機(jī)地址的內(nèi)容。
·讀取某一個(gè)bit的狀態(tài)，就是選中word line，那么圖示中的晶體管M1就會(huì)導(dǎo)通，通過(guò)bit line的sense就可以感知到這個(gè)時(shí)候電容Cs上的狀態(tài)，例如，現(xiàn)在如果這個(gè)bit的狀態(tài)為1，那么導(dǎo)通之后就會(huì)從bit line上得到1，反之也是同樣的道理。
·向某一bit寫(xiě)入1，首先通過(guò)row decoder選中word line,將會(huì)導(dǎo)致m1導(dǎo)通，那么bit line為1，會(huì)導(dǎo)致電容Cs充電，導(dǎo)致其電平為1，如果要寫(xiě)入，那么bit line的電平 為0，將會(huì)導(dǎo)致電容Cs放電，致使此時(shí)的電平為0。
由上面可以看出一個(gè)位只能表示一個(gè)bit，那么我們想讀取多個(gè)位的時(shí)候，該怎么辦呢？那就出現(xiàn)了ddr中的bank的概念，由多個(gè)memory array就組成了一個(gè)bank，如下圖,一次可以讀取2bit/4bit/8bit的數(shù)據(jù)：

登錄/注冊(cè)后可看大圖

由多個(gè)bank就可以組成一個(gè)memory device，如下圖，一個(gè)dram的芯片，由8個(gè)banks組成，而每個(gè)bank由4個(gè)memory array構(gòu)成，而此時(shí)每個(gè)bank輸出4個(gè)bit的I/O bus，那么為什么會(huì)出現(xiàn)bank的概念呢？動(dòng)態(tài)內(nèi)存區(qū)別于靜態(tài)內(nèi)存要定時(shí)刷新，每讀取一個(gè)狀態(tài)的時(shí)候，都需要重新充電。如果沒(méi)有采用bank，假設(shè)我們現(xiàn)在要讀取01-08地址的數(shù)據(jù)，當(dāng)讀取01的地址后，要等這個(gè)bank自刷新后才能讀取02地址的值，而采用8個(gè)bank之后，沒(méi)有這類(lèi)問(wèn)題，例如我們讀取完01地址之后，那么讀取02，因?yàn)?2與01的控制方式不同，所以對(duì)于這段時(shí)間01可以后臺(tái)的完成自刷新，依次類(lèi)推，那么就可以很好的解決動(dòng)態(tài)內(nèi)存需要刷新的問(wèn)題。

登錄/注冊(cè)后可看大圖

上圖中，對(duì)于RANK、DIMM等在深入淺出DDR系列(1)——DDR原理篇中已經(jīng)詳細(xì)介紹了。
2 DDR工作原理
了解了DDR的基本組成后，我們來(lái)看看DDR如何來(lái)完成一次的讀寫(xiě)過(guò)程。如下圖所示，DRAM 的相關(guān)操作在內(nèi)部大概可以分為以下的四個(gè)階段：

登錄/注冊(cè)后可看大圖

·command transport and decode: 在這個(gè)階段，Host 端會(huì)通過(guò) Command Bus 和 Address Bus 將具體的 Command 以及相應(yīng)參數(shù)傳遞給 DRAM。DRAM 接收并解析 Command，接著驅(qū)動(dòng)內(nèi)部模塊進(jìn)行相應(yīng)的操作。其中會(huì)根據(jù)將addr bus上的數(shù)據(jù)解碼成對(duì)應(yīng)的row address和通過(guò)bank control解碼后得到對(duì)應(yīng)的bank，其次對(duì)應(yīng)的column也會(huì)解碼得到對(duì)應(yīng)的地址
·in bank data movement: 在這個(gè)階段，第一階段發(fā)送需要讀取的 Column 的地址給 DRAM。然后 DRAM 再將 Active Command 所選中的 Row 中，DRAM 就將 Memory Array 中的數(shù)據(jù)從 DRAM Cells 中讀出到 Sense Amplifiers，或者將數(shù)據(jù)從 Sense Amplifiers 寫(xiě)入到 DRAM Cells。
·in device data movement: 這個(gè)階段中，數(shù)據(jù)將通過(guò) IO 電路緩存到 Read Latchs 或者通過(guò) IO 電路和 Write Drivers 更新到 Sense Amplifiers。
·system data transport: 在這個(gè)階段，進(jìn)行讀數(shù)據(jù)操作時(shí)，SDRAM 會(huì)將數(shù)據(jù)輸出到數(shù)據(jù)總線上，進(jìn)行寫(xiě)數(shù)據(jù)操作時(shí)，則是 Host 端的 Controller 將數(shù)據(jù)輸出到總線上。
在上述的四個(gè)階段中，每個(gè)階段都會(huì)有一定的耗時(shí)，例如數(shù)據(jù)從 DRAM Cells 搬運(yùn)到 Read Latchs 的操作需要一定的時(shí)間，因此在一個(gè)具體的操作需要按照一定時(shí)序進(jìn)行。同時(shí)，由于內(nèi)部的一些部件可能會(huì)被多個(gè)操作使用，例如讀數(shù)據(jù)和寫(xiě)數(shù)據(jù)都需要用到部分 IO 電路，因此多個(gè)不同的操作通常不能同時(shí)進(jìn)行，也需要遵守一定的時(shí)序。此外，某些操作會(huì)消耗很大的電流，為了滿(mǎn)足 SDRAM 設(shè)計(jì)上的功耗指標(biāo)，可能會(huì)限制某一些操作的執(zhí)行頻率。
3 DRAM基本命令
對(duì)于一款ddr，我們需要知道通過(guò)什么樣的方式來(lái)控制完成我們需要，我們來(lái)看看ddr的狀態(tài)，ddr的工作就是在這幾個(gè)狀態(tài)之間切換：

登錄/注冊(cè)后可看大圖

在芯片上電后，完成初始化后，dram處于idle階段，上圖是需要進(jìn)入各個(gè)階段的時(shí)候，應(yīng)該需要進(jìn)行那些基本的操作，對(duì)于ddr使用比較頻繁的幾個(gè)基本命令訪問(wèn)方式如下
·刷新模式:儲(chǔ)體中電容的數(shù)據(jù)有效是有時(shí)間限制的，所以為了保證數(shù)據(jù)的不丟失，所以要對(duì)ddr進(jìn)行定時(shí)的刷新，SDRAM內(nèi)部有一個(gè)行地址生成器（也稱(chēng)刷新計(jì)數(shù)器）用來(lái)自動(dòng)的依次生成行地址。由于刷新是針對(duì)一行中的所有存儲(chǔ)體進(jìn)行。該模式是由Host主動(dòng)控制DRAM完成刷新，存儲(chǔ)體中電容的數(shù)據(jù)有效保存期上限是64ms（毫秒，1/1000秒），也就是說(shuō)每一行刷新的循環(huán)周期是64ms。
·自我刷新模式：當(dāng)系統(tǒng)進(jìn)入低功耗模式，只需要發(fā)送一條 SRF指令，主要用于休眠模式低功耗狀態(tài)下的數(shù)據(jù)保存，比較常見(jiàn)的應(yīng)用是STR（Suspend to RAM，休眠掛起于內(nèi)存）。就進(jìn)入了該模式，此時(shí)不再依靠系統(tǒng)時(shí)鐘工作，而是根據(jù)內(nèi)部的時(shí)鐘進(jìn)行刷新操作。期間除了CKE之外的所有外部信號(hào)都是無(wú)效的（無(wú)需外部提供刷新指令），只有重新使CKE有效才能退出自刷新模式并進(jìn)入正常操作狀態(tài)。
·MRS模式(mode register set):模式寄存器中的數(shù)據(jù)控制著 DDR2 SDRAM的操作模式.它控制著 CAS 延遲, 突發(fā)長(zhǎng)度, 突發(fā)順序, 測(cè)試模式, DLL復(fù)位, WR等各種選項(xiàng)，支持著 DDR2 SDRAM 的各種應(yīng)用. 模式寄存器的默認(rèn)值沒(méi)有被定義, 所以上電之后必須按規(guī)定的時(shí)序規(guī)范來(lái)設(shè)定模式寄存器的值。
·EMRS 擴(kuò)展模式寄存器:存儲(chǔ)著激活或禁止DLL的控制信息, 輸出驅(qū)動(dòng)強(qiáng)度, ODT 值的選擇 和附加延遲等信息。
·預(yù)充電：對(duì)一行讀寫(xiě)操作后，關(guān)閉現(xiàn)有工作行，準(zhǔn)備打開(kāi)新行的操作就是預(yù)充電。
·讀過(guò)程：訪問(wèn)操作開(kāi)始ACT一個(gè)激活命令，主要是激活bank和rol，就等于選通了某一Bank的某一行，接著發(fā)送一個(gè)read指令，就可以通過(guò)數(shù)據(jù)總線將數(shù)據(jù)送出去了，然后就進(jìn)行預(yù)充電，恢復(fù)到讀寫(xiě)的狀態(tài)，預(yù)充電完成后，就恢復(fù)到idle狀態(tài)。
·寫(xiě)過(guò)程：與讀過(guò)程基本類(lèi)似。
DRAM的基本命令是通過(guò)操作各種控制信號(hào)/地址信號(hào)的組合來(lái)完成，下表是DRAM的命令表：

登錄/注冊(cè)后可看大圖

4 DDR的時(shí)序參數(shù)
4.1 Row Active Command
在進(jìn)行數(shù)據(jù)的讀寫(xiě)前，Controller 需要先發(fā)送 Row Active Command，打開(kāi) DRAM Memory Array 中的指定的 Row。Row Active Command 的時(shí)序如下圖所示：

登錄/注冊(cè)后可看大圖

tRCD：RAS－to－CAS Delay（tRCD），內(nèi)存行地址傳輸?shù)搅械刂返难舆t時(shí)間。
Row Active Command 通過(guò)地址總線指明需要打開(kāi)某一個(gè) Bank 的某一個(gè) Row。DRAM 在接收到該 Command 后，會(huì)打開(kāi)該 Row 的 Wordline，將其存儲(chǔ)的數(shù)據(jù)讀取到 Sense Amplifiers 中，這一時(shí)間定義為 tRCD（RCD for Row Address to Column Address Delay）。DRAM 在完成 Row Sense 階段后，Controller 就可以發(fā)送 Read 或 Write Command 進(jìn)行數(shù)據(jù)的讀寫(xiě)了。這也意味著，Controller 在發(fā)送 Row Active Command 后，需要等待 tRCD 時(shí)間才能接著發(fā)送 Read 或者 Write Command 進(jìn)行數(shù)據(jù)的讀寫(xiě)。
tRAS: Row Active Time,內(nèi)存行地址選通延遲
由于 DRAM 的特性，Row 中的數(shù)據(jù)在被讀取到 Sense Amplifiers 后，需要進(jìn)行 Restore 的操作。Restore 操作可以和數(shù)據(jù)的讀取同時(shí)進(jìn)行，即在這個(gè)階段，Controller 可能發(fā)送了 Read Command 進(jìn)行數(shù)據(jù)讀取。
DRAM 接收到 Row Active Command 到完成 Row Restore 操作所需要的時(shí)間定義為 tRAS（RAS for Row Address Strobe）。
Controller 在發(fā)出一個(gè) Row Active Command 后，必須要等待 tRAS 時(shí)間后，才可以發(fā)起另一次的 Precharge 和 Row Access。
4.2 Column Read Command
Controller 發(fā)送 Row Active Command 并等待 tRCD 時(shí)間后，再發(fā)送 Column Write Command 進(jìn)行數(shù)據(jù)寫(xiě)入。數(shù)據(jù) Burst Length 為 8 時(shí)的 Column Write Command 時(shí)序如下圖所示：

登錄/注冊(cè)后可看大圖

tCWD/tCL/tCWL:內(nèi)存CAS延遲時(shí)間
Column Write Command 通過(guò)地址總線 A[0:9] 指明需要寫(xiě)入數(shù)據(jù)的 Column 的起始地址。Controller 在發(fā)送完 Write Command 后，需要等待 tCWD （CWD for Column Write Delay） 時(shí)間后，才可以發(fā)送待寫(xiě)入的數(shù)據(jù)。tCWD 在一些描述中也稱(chēng)為 tCWL（CWL for Column Write Latency）
tWR（WR for Write Recovery）
DRAM 接收完數(shù)據(jù)后，需要一定的時(shí)間將數(shù)據(jù)寫(xiě)入到 DRAM Cells 中，這個(gè)時(shí)間定義為 tWR（WR for Write Recovery）。該值說(shuō)明在一個(gè)激活的bank中完成有效的寫(xiě)操作及預(yù)充電前，必須等待多少個(gè)時(shí)鐘周期。這段必須的時(shí)鐘周期用來(lái)確保在預(yù)充電發(fā)生前，寫(xiě)緩沖中的數(shù)據(jù)可以被寫(xiě)進(jìn)內(nèi)存單元中。同樣的，過(guò)低的tWD雖然提高了系統(tǒng)性能，但可能導(dǎo)致數(shù)據(jù)還未被正確寫(xiě)入到內(nèi)存單元中，就發(fā)生了預(yù)充電操作，會(huì)導(dǎo)致數(shù)據(jù)的丟失及損壞。
4.3 Precharge Command
要訪問(wèn) DRAM Cell 中的數(shù)據(jù)，需要先進(jìn)行 Precharge 操作。相應(yīng)地，在 Controller 發(fā)送 Row Active Command 訪問(wèn)一個(gè)具體的 Row 前， Controller 需要發(fā)送 Precharge Command 對(duì)該 Row 所在的 Bank 進(jìn)行 Precharge 操作。下面的時(shí)序圖描述了 Controller 訪問(wèn)一個(gè) Row 后，執(zhí)行 Precharge，然后再訪問(wèn)另一個(gè) Row 的流程。

登錄/注冊(cè)后可看大圖

DRAM 執(zhí)行 Precharge Command 所需要的時(shí)間定義為 tRP（RP for Row Precharge）。Controller 在發(fā)送一個(gè) Row Active Command 后，需要等待 tRC（RC for Row Cycle）時(shí)間后，才能發(fā)送第二個(gè) Row Active Command 進(jìn)行另一個(gè) Row 的訪問(wèn)。
從時(shí)序圖上我們可以看到，tRC = tRAS + tRP，tRC 時(shí)間決定了訪問(wèn) DRAM 不同 Row 的性能。在實(shí)際的產(chǎn)品中，通常會(huì)通過(guò)降低 tRC 耗時(shí)或者在一個(gè) Row Cycle 執(zhí)行盡可能多數(shù)據(jù)讀寫(xiě)等方式來(lái)優(yōu)化性能。
4.4 Row Refresh Command
一般情況下，為了保證 DRAM 數(shù)據(jù)的有效性，Controller 每隔 tREFI（REFI for Refresh Interval） 時(shí)間就需要發(fā)送一個(gè) Row Refresh Command 給 DRAM，進(jìn)行 Row 刷新操作。DRAM 在接收到 Row Refresh Command 后，會(huì)根據(jù)內(nèi)部 Refresh Counter 的值，對(duì)所有 Bank 的一個(gè)或者多個(gè) Row 進(jìn)行刷新操作。
DRAM 刷新的操作與 Active + Precharge Command 組合類(lèi)似，差別在于 Refresh Command 是對(duì) DRAM 所有 Bank 同時(shí)進(jìn)行操作的。下圖為 DRAM Row Refresh Command 的時(shí)序圖：

登錄/注冊(cè)后可看大圖

DRAM 完成刷新操作所需的時(shí)間定義為 tRFC（RFC for Refresh Cycle）。
tRFC 包含兩個(gè)部分的時(shí)間，一是完成刷新操作所需要的時(shí)間，由于 DRAM Refresh 是同時(shí)對(duì)所有 Bank 進(jìn)行的，刷新操作會(huì)比單個(gè) Row 的 Active + Precharge 操作需要更長(zhǎng)的時(shí)間；tRFC 的另一部分時(shí)間則是為了降低平均功耗而引入的延時(shí)，DRAM Refresh 操作所消耗的電流會(huì)比單個(gè) Row 的 Active + Precharge 操作要大的多，tRFC 中引入額外的時(shí)延可以限制 Refresh 操作的頻率。
4.5 Read Cycle
一個(gè)完整的 Burst Length 的 Read Cycle 如下圖所示：

登錄/注冊(cè)后可看大圖

下面是DDR常見(jiàn)的一些參數(shù)及定義如下：

登錄/注冊(cè)后可看大圖

上述的 DRAM Timing 中的一部分參數(shù)可以編程設(shè)定，例如 tCAS、tAL、Burst Length 等。這些參數(shù)通常是在 Host 初始化時(shí)，通過(guò) Controller 發(fā)起 Load Mode Register Command 寫(xiě)入到 DRAM 的 Mode Register 中。DRAM 完成初始化后，就會(huì)按照設(shè)定的參數(shù)運(yùn)行。
5性能分析
在學(xué)習(xí)完DDR的基本操作和時(shí)序參數(shù)之后，我們就看看性能的影響。當(dāng)頻率和位寬固定后，帶寬也就不可更改，但是在內(nèi)存的工作周期內(nèi)，不可能總處于數(shù)據(jù)傳輸?shù)臓顟B(tài)，因?yàn)橐忻?、尋址等必要的過(guò)程。那么這些操作占用的時(shí)間越短，內(nèi)存工作的效率就越高，性能也就越好。
對(duì)于我們來(lái)說(shuō)，最好的方法是提高頻率，但是提高頻率會(huì)受多方面的影響，還有什么辦法提高內(nèi)存訪問(wèn)采取速度。
·多通道: 現(xiàn)代內(nèi)存控制器從北橋移入CPU內(nèi)部，而內(nèi)存控制器都可以同時(shí)操作多個(gè)通道。比如現(xiàn)在的筆記本開(kāi)始支持雙通道、三通道，如果數(shù)據(jù)分布在不同通道的內(nèi)存條上，內(nèi)存控制器就可以不用管上面的這些延遲時(shí)序，同時(shí)可以讀取它們，速度可以增加兩倍，甚至三倍。
·交織方式(Interleaving)： 同一塊內(nèi)存分布到不同的通道中去，這樣無(wú)論Cache命中與否都可以同時(shí)存取，多通道的技術(shù)才能發(fā)揮更大的用處。
·超頻內(nèi)存： 也就是提升DDR的頻率來(lái)增加速度
6總結(jié)
對(duì)于DDR的讀寫(xiě)以及一些時(shí)序參數(shù)的原理性知識(shí)后，下一步就進(jìn)入到DDR的驅(qū)動(dòng)調(diào)式，主要是對(duì)于一款控制器，我們?cè)撊绾稳フ{(diào)試DDR。其實(shí)對(duì)于DDR的調(diào)試，主要的讀寫(xiě)的控制，都是由DDR的控制器完成了，我們主要是通過(guò)MRS模式/EMRS模式來(lái)完成對(duì)于DDR參數(shù)的配置，而對(duì)于MRS模式的使用，已經(jīng)集成到DDR控制器中完成了，我們只需要根據(jù)控制器手冊(cè)配置相應(yīng)的寄存器就可以完成對(duì)于DDR調(diào)試。
參考文獻(xiàn)鏈接
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

-----------------------------

鄧瑋

吳博士，有時(shí)間討論一下內(nèi)存優(yōu)化啊[大笑]