TC框架也是一個能夠降低高性能機器學習代碼編寫門檻的工具,它能夠直接將高級語言生成GPU代碼哈肖。(填補研究人員于數(shù)學運算領(lǐng)域的溝通鴻溝)
適用情況
1吻育、你的 Pytorch 層又大又慢,你打算為此寫一段 C++ 或者 CUDA 的優(yōu)化代碼牡彻,但是你又不擅長扫沼。
2、你寫好了 CUDA 代碼庄吼,但是你還要花費大量時間去優(yōu)化缎除。你希望可以在最短時間內(nèi)搞定這些。
3总寻、為了加速器罐,你想在網(wǎng)絡(luò)中融合多個層,例如 Conv-ReLU-BatchNorm 或者 Linear-ReLU-Linear-ReLU渐行。但是這很難理解轰坊。
4、你的研究涉及 CuDNN 和 MKL 未能優(yōu)化的不尋常的張量祟印。例如肴沫,你要使用 13 x 24 的卷積核對 143 x 55 的輸入圖像進行卷積。你試著用 CUDNN 跑蕴忆,并且發(fā)現(xiàn)它慢的超乎想象颤芬。
5、通過調(diào)整 Tensors 以適應(yīng)特定的內(nèi)存布局套鹅,你的代碼會變慢站蝠。你希望編寫在你的輸入布局上高效運行的自定義代碼很容易。
背景:大量的具有卷積和循環(huán)網(wǎng)絡(luò)的DL模型卓鹿,諸如TensorFlow菱魔、Pytorch這些框架也分別就可用性和表達性之間做了不同權(quán)衡。它們在計算運算符的DAG上操作吟孙,包裝成高性能庫澜倦,如CUDNN,自動進行內(nèi)存分配杰妓、同步和分發(fā)肥隆。但如果計算不適合現(xiàn)有的庫調(diào)用,則需要自定義運算符稚失,通常需要高昂的工程成本。
論文TC:Framework-Agnostic High-Performance Machine Learning Abstractions提出貢獻:(TC適合例如GPU這種有分區(qū)內(nèi)存的硬件設(shè)備)
1恰聘、提出了TC(Tensor Comprehensions):一個用于ML tensor計算的高級語言句各。(語法是泛化愛因斯坦Einstein表示法)吸占。支持形狀和大小推理;
2凿宾、端到端編譯流程:能夠?qū)C“降低”到高效的GPU代碼矾屯;
3、Polyhedral編譯算法:與通用的并行編譯器不同初厚,我們主要通過kernel融合優(yōu)化以減少啟動和同步的開銷件蚕,還支持多級并行和提升到內(nèi)存層次的更深層次;
4产禾、Autotuning框架:利用JIT編譯和代碼緩存的自動調(diào)整框架排作。它包括非標準大小的specialization,消除了控制和地址生成邏輯亚情,并擁有從ML框架到代碼生成器的所有優(yōu)化旋鈕 (knobs)妄痪。
5、Polyhedral JIT(實時)編譯器:用于將深度學習DAG的數(shù)學描述轉(zhuǎn)換為CUDA kernel楞件,并提供了諸如算子融合(operator fusion)和的specialization衫生;
6、集成到兩個通用的ML框架PyTorch和Caffe2土浸。并且在原則上罪针,這個系統(tǒng)是足夠通用集成到其他ML框架中的。
TC框架自動結(jié)合了仿射affine變換黄伊,包括層次平鋪 (hierarchical tiling)泪酱、映射 (mapping)、移動 (shifting)毅舆、融合 (fusion)西篓、分布 (distribution)、交換 (interchange)憋活,涉及參數(shù)化或完全實例化的問題岂津,這些問題對Halide、Latte或XLA都是不可達的悦即。
Tensor Comprehension 將 Halide 編譯器作為調(diào)用庫吮成。FAIR 研究員構(gòu)建了 Halide 的中間表征(IR)和分析工具,并與多面編譯進行技術(shù)配對辜梳,因此粱甫,用戶可以在無需理解運行原理的條件下使用相似的高級語法編寫層。此外作瞄,F(xiàn)AIR 研究員也找到了簡化語言的方法茶宵,不需要為縮減運算制定循環(huán)邊界。
發(fā)展
Phase1:現(xiàn)有的active library或者BTO (built-to-order)都是孤立地調(diào)整庫內(nèi)核宗挥,因此會錯過上下文依賴乌庶,并且創(chuàng)建一個新的庫种蝶,包含各個獨立kernel的所有組合是不可行的。
Phase2:因此瞒大,Halide這種domain-specific的語言被生成螃征,在imaging方面取得了成功,因為它能在不混淆底層算法的情況下融合大型的流水線pipelines透敌。然而盯滚,在GPU上使用使用Halide時,所有調(diào)度轉(zhuǎn)換都必須手動指定酗电,并且通過正確的組合實現(xiàn)高性能超出了大多數(shù)用戶的能力魄藕。(Halide 是一種最近在高性能圖像處理領(lǐng)域頗受歡迎的語言,它采用類似的高級函數(shù)語法來描述一個圖像處理的 pipeline顾瞻,隨后在單獨代碼塊中調(diào)度到硬件上泼疑,并且詳細到如何平鋪、矢量化荷荤、并行化和融合退渗。對于具有專業(yè)知識的人而言,這是一種非常高效的語言蕴纳;但對于機器學習從業(yè)者來說会油,這一難度并不小。Halide 的自動調(diào)度在研究上非彻琶活躍翻翩,但對于 GPU 上運行的機器學習代碼,目前還沒有很好的解決方案稻薇。)
Phase3:XLA和Latte:它們結(jié)合了計算圖和算子嫂冻,允許跨運算符進行優(yōu)化以及利用數(shù)據(jù)大小進行優(yōu)化。然而塞椎,GPU上的性能還不行桨仿。這些框架的轉(zhuǎn)換語言無法表示復(fù)雜調(diào)度 (complex schedule)和映射轉(zhuǎn)換 (mapping transformations),而這些轉(zhuǎn)換對于使用分區(qū)內(nèi)存架構(gòu)的GPU來說往往至關(guān)重要案狠。
綜上服傍,計算圖引擎的有效編程語言必須同時解決以下兩個難題:
1、確保abstraction不僅能提高程序員的工作效率骂铁,而且能使編譯器及其支持執(zhí)行環(huán)境消除與目標平臺無關(guān)的顧慮吹零,通過靠近機器的中間表示來改進代碼,并自動探索優(yōu)化空間拉庵。換言之灿椅,該系統(tǒng)必須能夠提供"abstraction without regret",同時傳遞編譯時可用的豐富的語義信息;
2阱扬、選擇適當?shù)?b>中間表示和優(yōu)化算法泣懊,處理深度并行性和內(nèi)存層次,以及硬件功能麻惶,如向量指令和專用內(nèi)存。
TC語法
1信夫、在表達式中使用索引變量隱式定義窃蹋,從索引內(nèi)容推斷索引變量的范圍;
2静稻、表達式右邊出現(xiàn)但左邊沒有的索引假設(shè)為reduction維度警没;?
3、迭代空間中各點的評估順序不影響輸出振湾。
實例
1杀迹、矩陣向量積
在TC中,定義 +=! 來替代 += 押搪,這樣就可以避免初始化操作树酪,即
函數(shù)mv大州,引入了兩個索引變量i和k续语。變量未在任何地方定義,隱式地成為索引變量厦画。通過索引使用的方式推斷它們的范圍疮茄;并且,因為k只出現(xiàn)在表達式右側(cè)部分根暑,所以存入C的存儲將對k進行reduce通過reduction運算符“+”力试。Reductions可以發(fā)生在多個變量之間,但是它們都共享同一類型的關(guān)聯(lián)和交換運算符(例如排嫌, +=)畸裳,以確保計算順序不影響計算值。
TC支持in-place就地更新躏率,但要求在LHS左手邊上分配任何元素之前讀取完整的RHS右邊表達式躯畴。如果LHS張量也在RHS中出現(xiàn):編譯器負責檢查元素依賴項的就地更新的因果關(guān)系。只允許在點態(tài)定義和張量收縮上進行就地更新(即a(i,j) = a(j,i)是不可以的)薇芝。
2蓬抄、SGEMM矩陣乘法、Relu夯到、conv2D嚷缭、max_pool等。注意where 語句用來確定變量的變化范圍,從而程序不會出越界錯阅爽。(如果沒有where語句路幸,那么循環(huán)將默認從0開始。)下標表達式可以是迭代器的任何仿射函數(shù)付翁,或下標表達式的下標及其組合简肴。后者捕獲與數(shù)據(jù)相關(guān)的訪問,如收集操作:
TC還可以很簡潔的實現(xiàn) strided convolution(跳格卷積百侧,即步長不為1)砰识,sh代表豎直步長,sw代表水平步長佣渴,H辫狼、W分別代表圖像數(shù)據(jù)(I)的長、寬辛润。
范圍推導(dǎo)
——在表達式中使用索引變量隱式定義膨处,從索引內(nèi)容推斷索引變量的范圍。以矩陣向量積為例砂竖,output的C的尺寸是通過i真椿,k可取的最大值推理而來的,即M和K晦溪,所以C[M*K]瀑粥。
全局數(shù)據(jù)布局轉(zhuǎn)換
TC使全局數(shù)據(jù)布局轉(zhuǎn)換變得非常容易。]ML研究者對于這類原語的主要用法是算法平鋪和層次分解三圆。前者與數(shù)據(jù)平鋪具有強連接狞换,現(xiàn)在在頻域中實現(xiàn)高性能卷積以控制內(nèi)存占用【83】或適合快速本地內(nèi)存【2】中無處不在。TC支持廣義張量轉(zhuǎn)置(即舟肉,應(yīng)用n-D置換矩陣修噪,其中n>2),并且可以通過簡單地重塑張量和調(diào)整TC索引表達式來實現(xiàn)數(shù)據(jù)平鋪路媚。范圍推斷和檢查保證了這種重塑將始終在整個TC一致黄琼。結(jié)構(gòu)數(shù)組到結(jié)構(gòu)數(shù)組的轉(zhuǎn)換以及短向量上的類似操作是可用的副產(chǎn)品:它們是維交換和數(shù)據(jù)平鋪的特殊情況。此時整慎,數(shù)據(jù)布局和TC轉(zhuǎn)換由領(lǐng)域?qū)<以谠创a級別進行脏款, TC推理過程保證了表達式具有適當?shù)姆秶N覀儗⒃趯硪隩C中的自動數(shù)據(jù)布局轉(zhuǎn)換時重新評估這一假設(shè)裤园。
高層工作流程說明
首先將TC表達式集成到ML框架中撤师。我們選擇了“in process”方式的實現(xiàn),以簡化與基于它們的計算圖引擎和ML應(yīng)用程序的交互拧揽,這是全自動調(diào)度(fully-automated scheduling)和映射流(mappling flow)的獨特功能剃盾。我們提供了一個輕API腺占,可將特定的張量對象模型轉(zhuǎn)換為我們自己的模型。算子定義被重寫以生成TC痒谴,而不是框架的后端實現(xiàn)衰伯,并為用戶提供編寫自己的TC的能力。在這種情況下积蔚,單個TC可能對應(yīng)于ML框架中一個運算符DAG(有向無環(huán)圖意鲸?)。目前尽爆,此匹配是手動完成的临扮。(以后將會由我們的編譯流程自動DAG分區(qū),匹配和重寫教翩。)然后,將TC進行JIT編譯贪壳,如圖所示饱亿。
從具有特定tensor大小和strides跨度的TC開始蚂且,我們將其降低為參數(shù)Halide表達式配猫。直接將Halide-IR降低為Poly表示形式。
使用 Halide 和多面編譯技術(shù)杏死,Tensor Comprehensions 能通過委托內(nèi)存管理和同步功能以自動合成 CUDA 內(nèi)核泵肄。這一編譯能夠針對特定尺寸對一般操作符進行融合,對快速本地內(nèi)存淑翼、快速縮減和 Just-in-Time 專業(yè)化都能優(yōu)化腐巢。因不嘗試涉足內(nèi)存管理,因此這一流程能夠輕松集成到機器學習的任意框架玄括,以及任何允許調(diào)用 C++ 函數(shù)的語言中冯丙。與經(jīng)典編譯器技術(shù)和庫方法相反的是,多面編譯能讓張量理解為每個新網(wǎng)絡(luò)按需調(diào)度單個張量元素的計算遭京。在 CUDA 層面胃惜,它結(jié)合了仿射循環(huán)轉(zhuǎn)換、融合/分裂和自動并行處理哪雕,同時確保數(shù)據(jù)在存儲器層次結(jié)構(gòu)中正確移動船殉。
Polyhedral JIT編譯器
編譯器彌合了高層張量操作的邏輯layout(尺寸排序)與Poly代碼生成器期望的數(shù)據(jù)格式之間的不匹配。 “下降”步驟可確保大小和跨度的組合對應(yīng)于非混疊數(shù)組和子數(shù)組語法热监; 它還確保沒有邊界訪問捺弦,分析訪問關(guān)系和推斷的張量范圍;它發(fā)出張量聲明和重新排序表達式以匹配目標語言的數(shù)據(jù)模型,即行主數(shù)組。
請注意列吼,張量規(guī)范可能具有輸入和輸出參數(shù)別名的特征幽崩,以便就地計算和隱式轉(zhuǎn)換高維張量。 我們認為寞钥,這樣的規(guī)范應(yīng)該導(dǎo)致每個別名場景的多版本化和專業(yè)化慌申。 而且,TC基于范圍推斷所采用的語義不同于XLA理郑,PyTorch和MXNet以某種形式或其他形式采用的NumPy風格的“廣播”語義蹄溉。6TC不需要這種隱式語法糖。 例如您炉,與所謂的外積矩陣乘積[p柒爵,q,r] matmul [1赚爵,s棉胀,r,t]→[p冀膝,s唁奢,q,t]對應(yīng)的TC很簡單:
如果需要窝剖,可以進一步轉(zhuǎn)換布局并導(dǎo)出QPTS版本(以輸出尺寸的順序命名)麻掸,而不是PSQT。
其他降低步驟包括卷積表達式的正向替換(存儲/計算折衷)赐纱,零填充脊奋,鏡像和裁剪。
多面體框架提供了基于循環(huán)的程序的最先進的內(nèi)部編譯器表示形式千所,這些循環(huán)程序“大多是規(guī)則的” [27]狂魔。它以最基本的形式描述了由循環(huán)和分支包圍的算術(shù)表達式,這些分支和分支的控制流可以靜態(tài)確定淫痰。因此最楷,據(jù)說多面體框架可在程序的靜態(tài)控制部分(SCoP)上運行。更具體地說待错,循環(huán)邊界和分支條件是外部循環(huán)迭代器的仿射表達式和象征性對待并稱為參數(shù)的靜態(tài)未知常數(shù)值籽孙。使用數(shù)組元素上的算術(shù)表達式描述計算,并且對下標的限制與對循環(huán)界限的約束相同火俄。存在擴展以通過過度逼近[10]或用戶定義的注釋[3]處理不規(guī)則性犯建。盡管該框架看似簡單,但它捕獲了大量的計算密集型代碼瓜客,它在域和數(shù)組大小上是參數(shù)化的适瓦,并且比諸如Halide的特定于域的表示更具表達力竿开。
Poly框架特別適合于深層神經(jīng)網(wǎng)絡(luò),它與具有長依賴性鏈和不均勻或全部模式的大而深嵌套的循環(huán)相關(guān)聯(lián)(由完全連接的層以及張量收縮和換位引起)玻熙。這些功能將優(yōu)化問題推到了與Halide用于圖像處理的啟發(fā)式空間不同的地方否彩,并且比單獨的線性代數(shù)擴大了空間。
調(diào)度樹
Affine仿射圖可以組成調(diào)度樹嗦随,以促進從高級語言(TC)到下游優(yōu)化器的屬性通信列荔。調(diào)度樹介紹特定的節(jié)點類型。一個band節(jié)點通過在迭代域上定義的一個或多個分段仿射函數(shù)來定義partial執(zhí)行順序枚尼。即可置換調(diào)度band贴浙,一維調(diào)度功能的元組,可以在保留程序語義的同時自由互換署恍。band中的仿射功能通常稱為band member或schedule dimension崎溃。過濾器節(jié)點(filter nodes)的集合劃分了迭代空間,將其子樹綁定到迭代域的子集盯质。根據(jù)執(zhí)行順序是否必須為正確性而受約束笨奠,可以將它們安排在set或sequence節(jié)點中。上下文context節(jié)點提供有關(guān)變量和參數(shù)的其他信息唤殴,例如張量范圍或GPU塊的大小到腥;它們還可能在子樹中引入局部作用域和常量參數(shù)朵逝,這在將induction變量映射到CUDA中的塊和線程標識符時很有用。最后乡范,擴展extension節(jié)點引入了不屬于原始迭代域的輔助計算配名,這對于例如引入將數(shù)據(jù)復(fù)制到shared memory。
TC的規(guī)范調(diào)度樹由外部Sequence節(jié)點定義晋辆,然后是每個TC語句的Filter節(jié)點渠脉。 在每個filter的分支內(nèi),Band節(jié)點定義一個身份調(diào)度瓶佳,該身份調(diào)度的一維調(diào)度功能與語句的循環(huán)迭代器一樣多芋膘。 根據(jù)TC語義,隱式循環(huán)形成一個可置換band霸饲。 圖3.a顯示了sgemm TC的規(guī)范調(diào)度樹为朋,參數(shù)聲明(N,M厚脉,K)→{...}习寸。
可以識別2-D嵌套初始化語句,然后識別3-D嵌套更新語句傻工。調(diào)度可以是輸入大小的參數(shù)霞溪,也可以具有有關(guān)張量大小的其他上下文信息孵滞。在Band節(jié)點未定義內(nèi)射injective調(diào)度的情況下,語句實例將按照其域坐標的字典順序進行調(diào)度鸯匹。Poly模型中的程序轉(zhuǎn)換涉及定義不同的調(diào)度坊饶,該調(diào)度對應(yīng)于遍歷迭代域的不同(部分或全部)順序。例如忽你,觀察到C張量在兩個嵌套之間被重用幼东,可以構(gòu)造圖3.b中的調(diào)度以利用訪問局部性并提高性能。該樹的特征是帶有兩個循環(huán)的i和j循環(huán)的外帶節(jié)點科雳,這對應(yīng)于循環(huán)融合根蟹。序列節(jié)點確保在啟用正確初始化的T的各個實例之前執(zhí)行S的實例。第二band僅適用于T糟秘,并且對應(yīng)于最里面的(減少)回路k简逮。此外,樹還引入了一個Context節(jié)點來陳述有關(guān)參數(shù)值的假設(shè)尿赚。
越界訪問
Poly模型允許對張量訪問進行相關(guān)編碼散庶。 將這些迭代域與表示為集合的迭代域進行組合,可以計算所有訪問的張量元素的集合(即語句的足跡)凌净,并檢查其是否適合(指定的或推斷的)張量大小悲龟。 特別地,訪問關(guān)系使得能夠檢測越界訪問冰寻。 從張量形狀推斷出的屬于足跡F = D.A而不屬于張量元素T集合的元素對應(yīng)于出界訪問须教。 因此,(D.A)\ T =?是不存在越界訪問的條件斩芭。
Polyhedral Transformation and Mapping
以并行架構(gòu)為目標時轻腺,程序轉(zhuǎn)換涉及調(diào)度以及映射策略的改變,必須遵守依賴關(guān)系划乖。除了保證轉(zhuǎn)換的有效性之外贬养,依賴關(guān)系還可以用于公開并行性(獨立實例可以并行執(zhí)行)或改善數(shù)據(jù)訪問局部性(依賴實例在近距離執(zhí)行)。存在幾種有效的調(diào)度算法琴庵,著眼于并行性误算,局部性和矢量化的組合。我們的轉(zhuǎn)換機制基于四個元素:
1迷殿、 isl提供核心的Poly調(diào)度:它可以自動優(yōu)化(外部)循環(huán)并行性和局部性尉桩;我們調(diào)整了仿射調(diào)度啟發(fā)法,將完整的TC程序折疊到單個GPU內(nèi)核中贪庙;
2蜘犁、調(diào)度被進一步平鋪以促進在GPU的深度并行性和內(nèi)存層次上的映射和時間重用;
3止邮、借鑒PPCG實現(xiàn)的映射到GPU算法:并帶有附加擴展以支持ML內(nèi)核中出現(xiàn)的更復(fù)雜和不完美的嵌套結(jié)構(gòu)这橙。
4奏窑、內(nèi)存提升通過顯式數(shù)據(jù)往返于共享shared和私有private內(nèi)存空間的傳輸。
Scheduling
isl調(diào)度程序的核心部分迭代解決整數(shù)線性規(guī)劃問題屈扎,以計算形成調(diào)度帶的分段仿射函數(shù)埃唯。它還確保這些band是可置換的,并且可以進一步平鋪鹰晨。在內(nèi)部墨叛,它構(gòu)建一個數(shù)據(jù)依賴圖(DDG),其中節(jié)點對應(yīng)于語句模蜡,邊表示它們之間的依賴漠趁。每個邊緣都帶有一組“類型化”的依賴關(guān)系注釋拢操。 isl調(diào)度程序設(shè)計用于更好的可伸縮性迁杨,因為在最壞的情況下整數(shù)線性編程具有指數(shù)復(fù)雜性缠导。特別是妖枚,它介紹了仿射聚類技術(shù),該技術(shù)基于分別為DDG的各個強連接組件計算調(diào)度范圍驯击,然后迭代地對這些組件進行聚類并相對于彼此進行調(diào)度携兵。聚類不僅減少了調(diào)度程序必須解決的線性問題的大小蕾哟,而且還為isl的循環(huán)融合啟發(fā)式算法奠定了基礎(chǔ)则披。
Poly擴展了isl調(diào)度程序共缕,以便為調(diào)用方提供對調(diào)度過程的更細粒度的控制。對于仿射變換士复,用戶可以提供其他任意約束以插入線性程序骄呼。對于聚類,在調(diào)度程序證明其有效之后判没,用戶可以為成對依賴圖組件組合提供決策函數(shù)。這些配置點用作調(diào)度策略的基礎(chǔ)隅茎。利用這些策略澄峰,仿射變換可以被限制為:(1)非負調(diào)度系數(shù)and/or(2)非偏斜變換(即,循環(huán)交換和移位)辟犀。聚類決策允許控制常規(guī)的最小和最大融合目標俏竞,此外,最大融合保留了至少三個嵌套的并行循環(huán)(映射到CUDA塊和線程)堂竟』昊伲可以通過自動調(diào)整過程配置和選擇調(diào)度策略。在所有情況下出嘹,我們都強制生成一個GPU內(nèi)核席楚。
Imperfectly Nested Loop Tiling不完全嵌套循環(huán)平鋪
進行調(diào)度后,循環(huán)平鋪將作為一個單獨的步驟實施税稼,并作為調(diào)度樹轉(zhuǎn)換執(zhí)行烦秩。本質(zhì)上垮斯,它將可轉(zhuǎn)換的調(diào)度band轉(zhuǎn)換為兩個band的鏈chain,其中外部頻段包含tile loops只祠,而內(nèi)部頻段包含point loops兜蠕,且行程數(shù)固定。這可以看作是常規(guī)的strip-mine和sink transformation抛寝。例如熊杨,圖3.c顯示了融合和平鋪的sgemm的進度樹。
除了常規(guī)的循環(huán)嵌套平鋪盗舰,調(diào)度樹表示還允許對不完全嵌套的循環(huán)進行平鋪晶府。該技術(shù)基于以下觀察:如果一個循環(huán)不攜帶依賴項(即并行),則它可能會沉入任何其他循環(huán)下面岭皂。在有效的調(diào)度中郊霎,所有依賴關(guān)系都由某個循環(huán)來承載(或滿足),沿著這些循環(huán)具有正距離爷绘。僅當依賴關(guān)系在某個循環(huán)中被另一個循環(huán)承載之前為負值時书劝,才會違反依賴關(guān)系。并行循環(huán)從定義上講不會攜帶依賴關(guān)系土至,因此不會影響依賴關(guān)系的滿足或違反购对。因此,不完全嵌套的切片是通過首先隔離所有band陶因,然后將并行點循環(huán)沉入樹中來實現(xiàn)的骡苞。在此過程中,將point band復(fù)制到序列(或集合)節(jié)點下的每個子樹中楷扬,并更新其調(diào)度解幽,以使其僅包含相關(guān)的域點。
sgemm的調(diào)度樹故意有兩個不完全嵌套的帶烘苹。相關(guān)性分析表明躲株,循環(huán)i和j是并行的。因此镣衡,我們可以將它們平鋪霜定,并將點環(huán)沉入約簡k環(huán)的帶以下,從而生成圖3.d中的調(diào)度樹廊鸥。檢查排列性后望浩,可以將帶點環(huán)的最內(nèi)層嵌套帶合并為一個band。
Mapping to Blocks and Threads
調(diào)度樹也可以用于表示到加速器的映射惰说,特別是具有多個塊和線程的GPU磨德。通過將某些調(diào)度程序band members和相應(yīng)的循環(huán)與線程或塊索引相關(guān)聯(lián)來執(zhí)行此操作。然后吆视,如果可能剖张,Poly代碼生成器將省略循環(huán)切诀,并相應(yīng)地重寫索引表達式。我們的映射算法是從最初為PPCG設(shè)計的算法衍生而來的搔弄,其中網(wǎng)格grid和塊的大小獨立于圖塊大小而指定幅虑。由于塊和線程的語義,只能映射屬于可置換調(diào)度范圍的并行循環(huán)顾犹。如果將點循環(huán)映射到線程倒庵,則切片大小和塊大小之間的比率將控制每個線程執(zhí)行的迭代次數(shù)。請注意炫刷,如果塊的點循環(huán)最終被映射到線程擎宝,則小于塊大小的切片大小會導(dǎo)致某些線程不執(zhí)行任何計算。
我們要求調(diào)度樹至少具有一個最外層的band浑玛,其外部尺寸平行绍申。與PPCG相反,PPCG將可能的多個外部帶中的每個都映射到塊和線程(平鋪后)顾彰,我們的映射算法將單個帶映射到塊极阅,以便生成ML框架所期望的單個內(nèi)核,而可以映射多個帶線程涨享。 (單個)最外面的band的平行尺寸映射到GPU塊筋搏。在每個調(diào)度樹分支中,通常由點循環(huán)組成的最里面的可替換帶被映射到具有以下限制的GPU線程厕隧”计辏跨分支的映射維數(shù)必須相等。這可能需要在某些分支中將某些線程尺寸映射為零吁讨。
映射本身是通過插入特殊名稱髓迎,通過上下文節(jié)點與代碼生成器通信以及將它們與filter節(jié)點中的band維度相關(guān)聯(lián)來執(zhí)行的。對于矩陣乘法示例建丧,我們的映射策略在圖3.e中生成了調(diào)度樹排龄。我們在調(diào)度樹中引入了一個上下文節(jié)點,以指示參數(shù)的有效大小以及網(wǎng)格和塊大胁杈椤(分別表示為bx,by和tx艰亮,ty闭翩,它們代表由blockIdx.xy和threadIdx可能采用的值。 xy)迄埃。當已知有效張量大小時疗韵,將立即執(zhí)行此插入操作。還要注意引用bx侄非,by蕉汪,tx和ty參數(shù)的Filter節(jié)點流译。這些節(jié)點表示到GPU的映射。
Memory Promotion
我們有興趣將部分張量提升為shared或private GPU內(nèi)存者疤。 雖然提升決策是通過試探法做出的福澡,并且在稍后的階段生成了相應(yīng)的命令性代碼,但是調(diào)度樹提供了一個方便的接口驹马,用于附加與內(nèi)存相關(guān)的信息革砸。 內(nèi)存提升基于數(shù)組平鋪array tile(一種用于軟件控制的本地內(nèi)存的數(shù)據(jù)切片形式)的概念。 它是數(shù)組中大小恒定的潛在跨步塊糯累,覆蓋了給定(調(diào)度)圖塊中訪問的所有元素算利。 我們重新訪問并擴展了PPCG對內(nèi)存提升的支持。
間接訪問數(shù)組的提升
我們的內(nèi)存提升方法還可以處理間接訪問的數(shù)組泳姐。不失一般性效拭,考慮訪問O [l + Idx [i] [j]] [k]。我們將O稱為外部數(shù)組胖秒,將Idx稱為索引數(shù)組缎患。如果是嵌套的間接尋址,則從最內(nèi)層到最外層將迭代處理外/索引對扒怖。雖然外部數(shù)組的第一個索引表達式所取的值是靜態(tài)未知的较锡,但我們?nèi)钥梢詫⑺鼈儽镜鼐彺鏋閟hared_O [l] [i] [j] [k] = O [l + Idx [i] [j]] [k]。由于某些值可以重復(fù)盗痒,因此僅當外部數(shù)組和索引數(shù)組均被讀取時才可以間接升級蚂蕴,因此對它們的寫入可能會導(dǎo)致無法簡單合并的不同值。通常俯邓,我們要求索引數(shù)組具有一個數(shù)組切片骡楼,即僅訪問其固定大小的塊。在為外部數(shù)組計算數(shù)組切片時稽鞭,我們忽略下標的間接部分(仿射部分照常處理)鸟整。然后,我們在提升的外部數(shù)組中引入與索引數(shù)組中一樣多的其他索引表達式朦蕴。這些新維度上的數(shù)組范圍完全對應(yīng)于索引數(shù)組的數(shù)組切片大小篮条。因此,升級后的數(shù)組的元素包含將使用給定索引數(shù)組訪問的全局數(shù)組元素的副本吩抓。間接下標僅在從全局存儲器復(fù)制時使用涉茧,而所有其他訪問均在代碼生成中重寫。如果存在多個共享子表達式并在相應(yīng)維度上具有相同圖塊大小的間接索引表達式疹娶,則對于所有相同的子表達式伴栓,在提升的數(shù)組中引入單個索引表達式就足夠了。
提升啟發(fā)法
如果存在固定大小的數(shù)組磁貼,單個元素被多次訪問并且至少其中一個訪問不具有內(nèi)存合并功能钳垮,則直接訪問的數(shù)組將提升為共享內(nèi)存惑淳。 從訪問關(guān)系和應(yīng)用于該域的調(diào)度可以看到后者:最后一個訪問維應(yīng)與映射到x線程的調(diào)度維對齊。 對于間接數(shù)組饺窿,由于存在這些情況所帶來的額外的長內(nèi)存依賴關(guān)系歧焦,因此刪除了合并要求。 共享內(nèi)存的總量是固定的短荐,如果所需的共享內(nèi)存數(shù)量會超出可用數(shù)量倚舀,我們將采用簡單的貪婪啟發(fā)式方法并拒絕升級。
Mapping to Registers
當前沒有方法實現(xiàn)比PPCG先前支持的更復(fù)雜的寄存器提升策略忍宋『勖玻基于以下觀察,這是一個臨時的糠排,務(wù)實的選擇:
1舵稠、無經(jīng)驗證據(jù)表明,根據(jù)高級規(guī)范自動生成低級代碼會比匯編產(chǎn)生顯著的性能提升入宦;另外哺徊,現(xiàn)代的組裝的generators現(xiàn)在已經(jīng)公開可用并且可以重新定位;?
2乾闰、將寄存器優(yōu)化泛化和重新定向到具有多個向量長度落追,對齊約束和模式的不同體系結(jié)構(gòu)并不是一件容易的事;
3涯肩、存在許多不同的策略轿钠,包括在寄存器級別進行流水線處理、寄存器輪換病苗、多緩沖疗垛、具有集體樣式語義的新ISA,其中有些甚至不可用沒有明顯的變化;
4硫朦、通常贷腕,此類策略是首先以內(nèi)部和概念驗證的方式實施的,并且可以使用PeachPy之類的工具輕松地進行模板化和重用咬展。?
實施的這種臨時狀態(tài)會對性能結(jié)果產(chǎn)生輕微的影響泽裳。我們還計劃依靠外部庫調(diào)用,例如破婆,以加快減少速度涮总。