文章目錄

• 前言
• intro
• programming model
• SIMT core內(nèi)部架構(gòu)
• • instruction fetch loop
  • two-loop approximation
  • three-loop approximation
• cache 架構(gòu)

前言

本文章來自general-purpose graphics processor units architecture

intro

首先GPGPU從傳統(tǒng)上來說用于處理圖形計算，但是隨著時代的發(fā)展，GPGPU越來越用于其他的場景，比如高性能計算，比如人工智能領(lǐng)域，所以本不涉及圖形層次

如果我們將蘋果A8處理器拆出來看，會發(fā)現(xiàn)A8處理器的GPU區(qū)域已經(jīng)大于CPU的區(qū)域，它不僅僅作圖形運算，還做SIMD等類型的workload

進(jìn)幾十年來，一些computing system性能越來越高，為什么？因為器內(nèi)部的晶體管的大小越來越小(工藝的提升)，因為硬件架構(gòu)的提升，因為編譯器技術(shù)的提升，因為算法的提升，最主要是晶體管越來越小，但是從2005年開始盲目的增加晶體管的數(shù)量已經(jīng)不是我們提升計算機架構(gòu)的首選(dannard scaling rule失敗了)，一個非常關(guān)鍵的因素是時鐘頻率提升變慢(因為時鐘頻率越來越高，熱量越來越大，熱量越大散熱變成非常大的一個難題)，所以現(xiàn)在提升計算機架構(gòu)的一個方法是找到一個更合適的硬件架構(gòu)，比如我們將一些應(yīng)用移到GPU上可能獲得500X的能源效率提升，有的通過消除指令集的瓶頸獲得了10X效率的提升

當(dāng)前計算機架構(gòu)的一個最主要的挑戰(zhàn)是，找到一種計算機架構(gòu)，這個架構(gòu)使用了特殊的硬件架構(gòu)獲得了更好的性能，并且適用的應(yīng)用也非常的廣，目前流行的做法是設(shè)計一種專門為某種應(yīng)用而運行的硬件，比如google的tensor processing unit，他專門用于深度神經(jīng)網(wǎng)絡(luò)

經(jīng)常有一個疑問叫做GPU是否可以完全的代替CPU？就目前而言，這非常的難，因為目前GPU并不是一個可以stand-alone computing的設(shè)備，目前的GPU要么和CPU集成在一個芯片上，GPU或者插在主板上通過PCIE和CPU進(jìn)行通訊(PCIE通訊速率也是GPU的瓶頸之一，因為PCIE傳輸對于GPU來說太慢了)，CPU用于初始化計算，并且將數(shù)據(jù)傳入到GPU中(通過PCIE總線)，并且，目前為止CPU中有北橋可以訪問主存，有中斷器可以相應(yīng)input設(shè)備，CPU還提供了一套API去訪問IO，但是GPU雖然說做到這些不難， 但是還沒有廠商實現(xiàn)，并且操作系統(tǒng)也是一個方面，我們的GPU一般并行處理數(shù)據(jù)，但是操作系統(tǒng)要求并行處理數(shù)據(jù)的場景比較少，所以操作系統(tǒng)也是一個問題，最終我們所考慮的形式是GPU和CPU進(jìn)行交互

下面是CPU和GPU如何共存的overlay，左邊是我們GPU直接插入PCIE插槽，右邊是GPU和CPU集成成一個芯片(AMD的APU)

從某個方面來說GPU程序的初始化必須要由CPU來完成，比如我們將數(shù)據(jù)傳到GPU的mem中，需要CPU來完成，我們的kernel(GPU程序)需要launch多少個線程(grim)，CPU調(diào)用GPU的API，在GPU上分配內(nèi)存，這些都要由CPU完成，然后將數(shù)據(jù)傳輸?shù)紾PU中，再由GPU運行

我們都知道GPU由非常多個SIMT Core Cluster組成,SIMT Core組成SIMT Core cluster，SIMT Core在nvidia中叫做streaming multiprocessors(SM),在AMD中叫做compute units(CU)，每個SIMT Core有自己的data / instruction cache(這里后面會展開講，我也會介個GPGPU-Sim的架構(gòu)來介紹)

每個SIMT Core中真正SIMD的是一個叫做warp的東西，nVidia一般是32個線程組成一個warp，換句話將32個線程同一時間執(zhí)行一個指令

每個SIMT Core可以跑非常多的線程，我們上面提到SIMT Core中有自己的data cache自己的instruction cache，每個SIMT Core Cluster也有一個cache(last level cache)供本集群內(nèi)所有的SIMT Core共享，每個SIMT Core通過on-chip的通道(cross bar也是上圖的interconnection network)訪問這個last-level cache，以上都是on-chip的cache，最后就是我們的顯存AKA device memory，這個device memory其實是off-chip的

曾經(jīng)在2009年Gus et al做了一個模型去分析GPU和CPU的性能，當(dāng)大量的cache(on-cache)被小量的線程共享(模擬CPU，且cache是定量)，然后慢慢的增加線程量，最后發(fā)現(xiàn)在cache不夠hold整個work set的時候性能就會急劇下降(因為on-chip cache不夠了會根據(jù)驅(qū)逐策略將一些內(nèi)容“驅(qū)逐到”off-chip中)然后慢慢的上升(因為線程數(shù)量慢慢的增多慢慢的隱藏了off-chip的延遲)，當(dāng)線程增加到一定程度就變成GPU形式了，然后隨著線程繼續(xù)增加性能會再遇到一個瓶頸，如下圖，MC(Multi Core表示模擬CPU)，MT(Multithread表示模擬GPU)

programming model

CUDA的文章太多了這里就不詳細(xì)的講解CUDA怎么寫，怎么使用shared_mem等知識，我們直接開始匯編
我們知道高級語言代碼經(jīng)過編譯器翻譯成匯編代碼，且匯編代碼本質(zhì)上是處理器ISA規(guī)定了底層硬件接口的格式，匯編代碼調(diào)用這些接口操作機器，但是我們GPU是另外的一個處理器，所以GPU也有自己的匯編語言，我們寫cuda的時候kernel部分被nvcc翻譯成PTX(或者叫做parallel thread execution ISA)，也就是nvidia GPU的高級虛擬指令(high level virtual instructor)，且這個匯編指令還是抽象的，因為nvidia產(chǎn)品線不同，每個GPU可能有著不同的"更底層的匯編"AKA SASS(streaming ASSembler)，而PTX這個匯編語言是一個抽象的高級匯編語言可以適用于所有的GPU產(chǎn)品，PTX通過nvidia的匯編器翻譯成"更底層的匯編語言"SASS

PTX和SASS打個比方，我們x86和MISP和RISC，ARM可能指令都不一樣，因為要設(shè)計處理器，首先就需要有指令集，規(guī)定處理器相應(yīng)操作，通過指令集去控制處理器實現(xiàn)相應(yīng)功能。所以SASS可以看成MISP，RISC，ARM等不同指令架構(gòu)的處理器，當(dāng)然他們因為功能不同也有不同的匯編，那么此時我們想讓我們的匯編可以跑在任何一個架構(gòu)處理器上，需要抽象出一個高層次的匯編，這里PTX就是高層次的匯編

PTX匯編如下(向量相加)

1 .visible .entry _Z5saxpyifPfS_( 2 .param .u32 _Z5saxpyifPfS__param_0, 3 .param .f32 _Z5saxpyifPfS__param_1, 4 .param .u64 _Z5saxpyifPfS__param_2, 5 .param .u64 _Z5saxpyifPfS__param_3 6 ) 7 { 8 .reg .pred %p<2>; 9 .reg .f32 %f<5>; 10 .reg .b32 %r<6>; 11 .reg .b64 %rd<8>; 12 13 14 ld.param.u32 %r2, [_Z5saxpyifPfS__param_0]; 15 ld.param.f32 %f1, [_Z5saxpyifPfS__param_1]; 16 ld.param.u64 %rd1, [_Z5saxpyifPfS__param_2]; 17 ld.param.u64 %rd2, [_Z5saxpyifPfS__param_3]; 18 mov.u32 %r3, %ctaid.x; 19 mov.u32 %r4, %ntid.x; 20 mov.u32 %r5, %tid.x; 21 mad.lo.s32 %r1, %r4, %r3, %r5; 22 setp.ge.s32 %p1, %r1, %r2; 23 @%p1 bra BB0_2; 24 25 cvta.to.global.u64 %rd3, %rd2; 26 cvta.to.global.u64 %rd4, %rd1; 27 mul.wide.s32 %rd5, %r1, 4; 28 add.s64 %rd6, %rd4, %rd5; 29 ld.global.f32 %f2, [%rd6]; 30 add.s64 %rd7, %rd3, %rd5; 31 ld.global.f32 %f3, [%rd7]; 32 fma.rn.f32 %f4, %f2, %f1, %f3; 33 st.global.f32 [%rd7], %f4; 34 35 BB0_2: 36 ret; 37 }

因為PTX是高級抽象匯編語言，所以其寄存器是無限量的，而SASS則不是

SASS匯編代碼如下(nvidia fermi架構(gòu))

Address Dissassembly Encoded Instruction 2 ======== =============================================== ======================== 3 /*0000*/ MOV R1, c[0x1][0x100]; /* 0x2800440400005de4 */ 4 /*0008*/ S2R R0, SR_CTAID.X; /* 0x2c00000094001c04 */ 5 /*0010*/ S2R R2, SR_TID.X; /* 0x2c00000084009c04 */ 6 /*0018*/ IMAD R0, R0, c[0x0][0x8], R2; /* 0x2004400020001ca3 */ 7 /*0020*/ ISETP.GE.AND P0, PT, R0, c[0x0][0x20], PT; /* 0x1b0e40008001dc23 */ 8 /*0028*/ @P0 BRA.U 0x78; /* 0x40000001200081e7 */ 9 /*0030*/ @!P0 MOV32I R5, 0x4; /* 0x18000000100161e2 */ 10 /*0038*/ @!P0 IMAD R2.CC, R0, R5, c[0x0][0x28]; /* 0x200b8000a000a0a3 */ 11 /*0040*/ @!P0 IMAD.HI.X R3, R0, R5, c[0x0][0x2c]; /* 0x208a8000b000e0e3 */ 12 /*0048*/ @!P0 IMAD R4.CC, R0, R5, c[0x0][0x30]; /* 0x200b8000c00120a3 */ 13 /*0050*/ @!P0 LD.E R2, [R2]; /* 0x840000000020a085 */ 14 /*0058*/ @!P0 IMAD.HI.X R5, R0, R5, c[0x0][0x34]; /* 0x208a8000d00160e3 */ 15 /*0060*/ @!P0 LD.E R0, [R4]; /* 0x8400000000402085 */ 16 /*0068*/ @!P0 FFMA R0, R2, c[0x0][0x24], R0; /* 0x3000400090202000 */ 17 /*0070*/ @!P0 ST.E [R4], R0; /* 0x9400000000402085 */ 18 /*0078*/ EXIT; /* 0x8000000000001de7 */

因為我們說了nvidia不同的產(chǎn)品的底層SASS不同所以相同代碼(向量相加)，fermi架構(gòu)如上，pascal架構(gòu)的SASS如下

1 Address Dissassembly Encoded Instruction 2 ======== =============================================== ======================== 3 /* 0x001c7c00e22007f6 */ 4 /*0008*/ MOV R1, c[0x0][0x20]; /* 0x4c98078000870001 */ 5 /*0010*/ S2R R0, SR_CTAID.X; /* 0xf0c8000002570000 */ 6 /*0018*/ S2R R2, SR_TID.X; /* 0xf0c8000002170002 */ 7 /* 0x001fd840fec20ff1 */ 8 /*0028*/ XMAD.MRG R3, R0.reuse, c[0x0] [0x8].H1, RZ; /* 0x4f107f8000270003 */ 9 /*0030*/ XMAD R2, R0.reuse, c[0x0] [0x8], R2; /* 0x4e00010000270002 */ 10 /*0038*/ XMAD.PSL.CBCC R0, R0.H1, R3.H1, R2; /* 0x5b30011800370000 */ 11 /* 0x081fc400ffa007ed */ 12 /*0048*/ ISETP.GE.AND P0, PT, R0, c[0x0][0x140], PT; /* 0x4b6d038005070007 */ 13 /*0050*/ @P0 EXIT; /* 0xe30000000000000f */ 14 /*0058*/ SHL R2, R0.reuse, 0x2; /* 0x3848000000270002 */ 15 /* 0x081fc440fec007f5 */ 16 /*0068*/ SHR R0, R0, 0x1e; /* 0x3829000001e70000 */ 17 /*0070*/ IADD R4.CC, R2.reuse, c[0x0][0x148]; /* 0x4c10800005270204 */ 18 /*0078*/ IADD.X R5, R0.reuse, c[0x0][0x14c]; /* 0x4c10080005370005 */ 19 /* 0x0001c800fe0007f6 */ 20 /*0088*/ IADD R2.CC, R2, c[0x0][0x150]; /* 0x4c10800005470202 */ 21 /*0090*/ IADD.X R3, R0, c[0x0][0x154]; /* 0x4c10080005570003 */ 22 /*0098*/ LDG.E R0, [R4]; } /* 0xeed4200000070400 */ 23 /* 0x0007c408fc400172 */ 24 /*00a8*/ LDG.E R6, [R2]; /* 0xeed4200000070206 */ 25 /*00b0*/ FFMA R0, R0, c[0x0][0x144], R6; /* 0x4980030005170000 */ 26 /*00b8*/ STG.E [R2], R0; /* 0xeedc200000070200 */ 27 /* 0x001f8000ffe007ff */ 28 /*00c8*/ EXIT; /* 0xe30000000007000f */ 29 /*00d0*/ BRA 0xd0; /* 0xe2400fffff87000f */ 30 /*00d8*/ NOP; /* 0x50b0000000070f00 */ 31 /* 0x001f8000fc0007e0 */ 32 /*00e8*/ NOP; /* 0x50b0000000070f00 */ 33 /*00f0*/ NOP; /* 0x50b0000000070f00 */ 34 /*00f8*/ NOP; /* 0x50b0000000070f00 */

SIMT core內(nèi)部架構(gòu)

上述的pipeline可以分成2部分(綠色黃色)，分別是Front-end，Back-end
我們要注意的是這里有3個“l(fā)oop”

• instruction fetch loop

  包括fetch指令，緩存指令到I-Cache中，decode I-Cache中的指令且緩存到I-Buffer中

• instruction issue loop

  包含I-Buffer，score Board，SIMT-Stack，Issue

• register access scheduling loop

  包含ALU，memory 等等

本章就圍繞這三部分講

instruction fetch loop

首先我們要明白GPU為了提升效率，其調(diào)度的最小單元是Warp，每個cycle GPU調(diào)度一個warp去執(zhí)行,按照上面的圖，我們知道指令被取值單元取后，緩存在I-Cache中，然后decode，decode后緩存到I-Buffer中。I-buffer中會round robin的為每一個decode后的指令分配warp，假如指令通過scoreboard的檢測(檢測WAW，RAW之類的問題)那么就會在I-Buffer的對應(yīng)字段設(shè)置r(ready)
SIMT stack就是為了處理分支條件判斷的情況而設(shè)計的，假如我們有以下的cuda代碼

do{t1 = tid*N; //At2 = t1 + i;t3 = datal[t2];t4 = 0;if( t3 != t4 ){t5 = data2[t2]; //Bif( t5 != t4 ){x += 1; //C}else{y += 2; //D}}else{z += 3; //F}i++; //G }while( i < N )

上述代碼的PTX匯編(做了改動，不是反編譯后直接給你寫上來)

A: mul.lo.u32 t1, tid, N; //tid * N存入到t1里面格式是u32,在local存儲區(qū)域(每個線程獨享)add.u32 t2, t1, i;ld.global.u32 t3, [t2];mov.u32 t4, 0;setp.eq.u32 p1, t3, t4; @p1 bra F; B: ld.global.u32 t5, [t2];setp.eq.u32 p2, t5, t4; @p2 bra D; C: add.u32 x, x, 1;bra E; D: add.u32 y, y, 2; E: bra G; F: add.u32 z, z, 3; G: add.u32 i, i, 1;setp.le.u32 p3, i, N; @p3 bra A;

參考下面的鏈接
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html
https://downloads.ti.com/docs/esd/SLAU131L/Content/SLAU131L_HTML/assembler_directives.html
為什么nvidia指令會有.lo.u32這樣的東西出現(xiàn)？首先指令后面的lo代表在local存儲空間開辟.u32大小的空間存入數(shù)據(jù)，數(shù)據(jù)格式是u32大小，中間的.lo .global代表我們存放到那個cache中，最后的.u32 .b32代表我們最后存入的格式，cache類別如下

namedescription

.reg	Registers, fast.
.sreg	Special registers. Read-only; pre-defined; platform-specific.
.const	Shared, read-only memory.
.global	Global memory, shared by all threads.
.local	Local memory, private to each thread.
.param	Kernel parameters, defined per-grid; or Function or local parameters, defined per-thread.
.shared	Addressable memory, defined per CTA, accessible to all threads in the cluster throughout the lifetime of the CTA that defines it.
.tex	Global texture memory (deprecated).

首先根據(jù)上述的代碼我們直接列出他們的分支圖

當(dāng)我們執(zhí)行完上述代碼第六行的時候(第一個分支),SIMT stack如下

首先假設(shè)warp只有4個線程，TOS是棧頂，出棧是先從棧頂出，1110代表warp的4個線程block了最后一個，換句話說warp執(zhí)行到這個分支的時候只會有前三個warp執(zhí)行

繼續(xù)執(zhí)行到第八行的時候又要遇到分支(此時第八行已經(jīng)執(zhí)行完畢)，此時會彈出上圖SIMT stack的TOS，然后又從TOS入棧，此時的SIMT stack如下Ⅰ，Ⅱ，Ⅲ都是新入棧的entire，第8行執(zhí)行的時候要么走C，要么走D，C和D走完走E，所以TOS是C代表先走C

當(dāng)C走完后代表著不用執(zhí)行D了，那么SIMT stack連續(xù)出棧2次(POP C,D)，最后SIMT Stack如下

小心SIMT死鎖！
關(guān)于atomic的article看這里https://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/memorystatisticsatomics.htm
我們接下來用的atomic函數(shù)是atomicCAS(A, B, num)這個是典型的自旋鎖，他的意思是對比A和B假如A和B相等就講num寫入A中，這個是原子的，且返回A的原值，下面代碼會引起deadlock

mutex = 0; while(!atomicCAS(mutex, 0, 1)); // critical section automicExch(mutex, 0);

為什么會發(fā)生死鎖？，首先我們的warp執(zhí)行到while的時候發(fā)現(xiàn)有一個atomic語句，則warp內(nèi)所有的線程序列化的執(zhí)行(這里不是lock-step，效率會變慢)，當(dāng)一個線程將mutex改為1的時候，此線程會執(zhí)行while內(nèi)部的語句(當(dāng)然要等warp內(nèi)其他的線程完成atomic操作)，當(dāng)其他的線程準(zhǔn)備完成atomic操作的時候發(fā)現(xiàn)mutex不是等于0，而是等于1，此時其他的線程會跳過while里面的句子，執(zhí)行下面的atomicExch

warp調(diào)度
假設(shè)我們在一個理想的的內(nèi)存環(huán)境下，內(nèi)存在某個固定的延遲下響應(yīng)內(nèi)存請求，那么我們可以設(shè)計充足的warp，用細(xì)顆粒讀的多線程來隱藏這種延遲(當(dāng)發(fā)生stall就開啟一個新的warp來調(diào)度)

two-loop approximation

假設(shè)我們有2個指令A(yù)和B，B必須要先于A執(zhí)行(因為A執(zhí)行需要B的結(jié)果)，此時B先issue給warp了，但是還沒有執(zhí)行完畢，A就執(zhí)行了，此時在我們GPU架構(gòu)中是萬萬不可得，因此在one-loop approximation中就這個問題解決，首先指令issue邏輯是我們只需要access這個warp的tid，還有address of next instruction，我們并不知道next instruction和這次執(zhí)行的instruction是否相關(guān)，不知道著此執(zhí)行的instruction是否執(zhí)行完畢，所以在第一次fetch指令(從內(nèi)存中)的時候dependency information就變得至關(guān)重要(假如拖到I-cache中就不好解決dependency的問題，因為此時已經(jīng)將decoded的指令綁定了warp，此時只需要檢查一些bit位是否合法就直接issue，scoreboard只檢查WAW，RAW等問題)，有了這些信息我們就知道這些hazards是否存在

在GPU中解決方法是我們搞了一個I-Buffer，I-Buffer在I-cache之前，I-Buffer有一個單獨的調(diào)度器，這個調(diào)度器來決定，哪一些指令(I-Buffer中的指令)應(yīng)該發(fā)布到余下的pipeline中，所以在I-Buffer的調(diào)度器中就解決了指令依賴的問題

一般來說I-Buffer是FLI(first-level instruction),I-buffer聯(lián)合MSHRs來幫助我們用來降低i-cache的latency，對于I-Buffer的結(jié)構(gòu)最通用的是一個warp對應(yīng)多個instruction

那么我們?nèi)绾翁綔y一個warp中的instruction有依賴呢？在傳統(tǒng)CPU架構(gòu)上有2種方法

• scoreboard
• reservation stations

這里直接講scoreboard，這個東西支持in-order excution和out of order execution(比如CDC 6600)，另一方面，scoreboard對于單線程in-order CPU實現(xiàn)非常簡單，CPU上的每個寄存器在scoreboard種都用單bit表示，當(dāng)指令想寫/讀這個寄存器的時候，再set這個bit，當(dāng)另一個instruction想寫這個寄存器，必須等這個寄存器對應(yīng)的bit clear后才能set(如果沒有clear就stall)，這樣有效的預(yù)防了RAW，WAW等危害，但是在in-order multi warp就有了不小的挑戰(zhàn)，挑戰(zhàn)如下

• 假設(shè)一個寄存器分配一個bit，如果我們GPU每個warp有128個寄存器，一共64個warp(遠(yuǎn)遠(yuǎn)不止)，那么我們一共有8192個寄存器，scoreboard為每一個寄存器分配一個bit那么就是8192個bit
• 還有一個是，假如A依賴B，只有B運行結(jié)束后，將其結(jié)果寫入寄存器后，A才能訪問B結(jié)果寫入的寄存器，取其值使用，不然A會一直查找其操作數(shù)(operands)，假設(shè)在多線程的環(huán)境下，那么多個線程可能早就被stall住，等待更早的指令完成(那么多線程怎么確定指令有沒有完成呢？不斷地探測scoreboard)，此時問題來了，假設(shè)GPU支持64個warp，一個warp最多支持4個操作數(shù)(通常來說操作數(shù)存儲在寄存器，內(nèi)存中)，那么scoreboard要有256個讀取端口(讓線程訪問)

three-loop approximation

首先我們要隱藏延遲，我們應(yīng)該支持單個core多warp，并且支持warp之間的切換，且我們要有一個大的register file，這個register file被分割成多個物理寄存器分配給每個warp，我們知道寄存器存儲的有操作數(shù)，當(dāng)我們在一個周期內(nèi)一個指令想訪問這個存器的操作數(shù)就需要一個端口(或者通道)，但是每個周期不是只有一個指令訪問一個寄存器，所以一個端口是遠(yuǎn)遠(yuǎn)不夠的，但是端口做多了會使寄存器變得面積非常大，所以主流的做法是用multiple banks 模擬一個非常大數(shù)量的端口，比如operand collector如下圖顯示了一個原本架構(gòu)的，下面的架構(gòu)現(xiàn)實了遞增的register file bandwidth

上圖的register file由4個單port邏輯bank寄存器組成，注意這里寫的是邏輯bank，他可能實際由非常多的物理bank組成，

• pipeline register是用來緩存operand，用于發(fā)送給SIMD執(zhí)行單元的
• arbitrator用來訪問每個單獨的bank，并且route訪問到的結(jié)果給對應(yīng)的pipeline register
• bank(single-ported registerd file bank)就是我們傳統(tǒng)意義上的寄存器

下圖顯示了bank的邏輯layout

如上圖register 4被warp1調(diào)用且存儲在bank0中
這樣設(shè)計有什么不好，我們假設(shè)有以下的2個指令

i1 mad r2, r5, r4, r6 i2 add r5, r5, r1

第一個指令是將r5(bank1中)，r4(bank0)，r6(bank2)都相加寫入(write back)到r2(bank2)中
第二個指令是將r5(bank1中)，r1(bank1)相加寫入(write back)到r5中

假設(shè)decode的cycle如下，且warp操作和寫回之間至少間隔一個周期(就是在沒有stall的情況下從讀寄存器到寫回到寄存器最少3個cycle)

那么實際的執(zhí)行cycle如下(一定要記住bank是單port，一個cycle只能有一個指令訪問)

我們發(fā)現(xiàn)第一個cycle還挺正常，第一條指令完美的在bank0，bank1，bank2中執(zhí)行訪問，第二個周期就開始不正常了，第二個周期執(zhí)行第二條指令(i2),要訪問寄存器r5和r1，但是r5和r1都在bank1中，同一個cycle一個bank只能被一條指令訪問，因為bank是single port的，所以指令只能訪問bank1中的一個寄存器，那就r1，第三個周期訪問bank1的第二個寄存器r5，并且在cycle3中第一條指令完成writeback進(jìn)bank2的r2寄存器中(因為在周期3bank2的port是空的，所以可以訪問)

綜上所屬，瓶頸就在存放寄存器的bank中，因為這4個bank，每個bank只允許在一個cycle中一個指令的操作(單端口)，而bank又由多個register組成，假如操作同一個bank的不同register，則需要多個cycle

---

此時一種新的operand 架構(gòu)由2008年提出來的如下圖

上圖最大的改變是用于緩存即將發(fā)往SIMD 執(zhí)行單元的operand的staging register改成了collector units，為了減少bank的沖突，還有一個ideal就是allocate equivalent register from different warps in different banks翻譯成中文就是同一個register要分配給不同bank中的warp

為什么寄存器1 register1既可以存在bank1和bank2中？因為我們這里討論的是邏輯寄存器，而非物理寄存器，邏輯寄存器根本上來說是邏輯的，他可能對應(yīng)不同的物理寄存器，上述新ideal就是在一個bank中不同的邏輯寄存器不能分配給同一個warp(假設(shè)寄存器1分配給了warp1，那么后續(xù)在這個bank中寄存器1就不能再分配給warp1只能分配給不同的warp)，這樣大大的減少了不同warp訪問同一個bank引發(fā)的沖突

如下圖

然后我們指令如下

i1: add r1, r2, r5 i2: mad r4, r3, r7, r1

我們指令的cycle如下

對應(yīng)的bank訪問如下

cache 架構(gòu)

首先我們給出GPU的cache 組成架構(gòu)，這個架構(gòu)實現(xiàn)了shared memory和L1 data cache

首先instruction pipeline的Load/Store unit(ldst)發(fā)送過來memory request，memory request是由多個memory 地址組成,
從上圖中我們可以看到LDST單元發(fā)送過來的請求先給arbiter，arbiter先探測，LDST請求的地址會不會發(fā)生bank conflicts(因為有可能一個warp中不同的thread訪問不同的線程)

這里詳細(xì)說一下，我們在cuda編程的時候會提到coalesced和uncoalesced，這里coalesced就是一個warp內(nèi)的所有線程都訪問到一個L1data cache block中，不管命沒命中，都是coalesced的，因為，命中萬事大吉直接返回，沒有命中就會讓其中一個線程去L2cache或者global mem中取(MSHR!),假設(shè)一個warp中的所有thread訪問的是不同的cache block，此時會形成多個memory access，這個就是uncoalesced

值得注意的是shared memory也會發(fā)生bank conflict，首先shared memory會被化成多個bank，這里注意的是每個bank只有一個讀port，一個寫port，一個cycle只允許一個thread讀，寫，假設(shè)多個thread同時寫一個bank，或者讀一個bank，那么就會發(fā)生bank conflict，發(fā)生了conflict會怎么辦呢？講其他的thread延后，放在后面的cycle執(zhí)行

在寫cuda的時候要避免bank conflict和uncoalesced

arbiter探測是否LDST單元發(fā)過來的一系列地址會發(fā)生shared memory bank conflict，假如發(fā)生了一個或者多個bank conflict，arbiter會講請求split成2部分，首先的一部分是warp中thread所訪問內(nèi)存不會引起conflict的地址，這一部分的地址訪問請求會發(fā)送給接下來的步驟，而第二部分也就是剩下的thread所需要訪問的地址可能會引起bank conflict，這一部分的地址訪問請求會被打回instruction pipeline中，并且要求重新執(zhí)行，剩下已經(jīng)被arbiter接收的shared mem 請求會在tag unit中尋找對應(yīng)的shared mem是否被標(biāo)記成bypass
tag unit還確定那個thread request map到了那個bank中，所以為了控制這些address crossbar(這些crossbar 分發(fā)訪問地址到每個data array的bank中)，且每一個data array中的ban都是32bit寬，且每個bank有自己的decoder，就是為了不相關(guān)的access不同的row in each bank，我們的數(shù)據(jù)從data crossbar中返回存儲到register file中
我們的L1cache block假如有128bytes(fermi和kepler就是的，cache line就是cache block!!!)，且被劃分成4個32byte的sector，因為32bytes大小的sector對應(yīng)著DRAM芯片(DDR5)一次最小讀取的數(shù)量,且一個bank有32bit大小，一個L1 cache block有32個bank(一個sector8個bank)
LDST 單元負(fù)責(zé)warp的coalescing，LDST單元將warp的memory access 分成多個Coalescing access，然后交給arbiter，假設(shè)一些資源不可用那么arbiter就會reject，比如cache line busy(一個cycle只允許一個線程讀或者寫，因為cache data只有2個端口，一個負(fù)責(zé)讀，一個負(fù)責(zé)寫)，且上圖的pending request table滿了(pending request table就是緩存請求)，還有一個情況是cache data可用，但是cache miss，此時，arbiter會返回一些writeback到instruction pipeline的相應(yīng)寄存器中，告訴instruction pipeline在未來再訪問相同的數(shù)據(jù)需要花費固定的cycle才能cache hit，arbiter如何判定request會miss或者h(yuǎn)it？通過tag unit…
當(dāng)cache hit后上圖的data array中會返回對應(yīng)的row給instruction pipeline的register file
假設(shè)cache miss，首先由arbiter探測到，當(dāng)cache miss發(fā)生后，arbiter會同時干2個事情，第一是informs LDST unit(通知LDST單元發(fā)生cache miss了)，第二件事情是將請求的信息緩存到PRT(上圖的pending request table)，此時你會發(fā)現(xiàn)PRT有點像CPU中的MSHR(其實一些GPGPU架構(gòu)中也有MSHR，比如GPGPU-Sim中就有…)，在nvidia中關(guān)于這個PRT至少有2個版本，第一個有點像傳統(tǒng)的MSHR

傳統(tǒng)的MSHR每一個entries包含cache miss的那個cache block(cache line) address

下面開始將L1data cache，在學(xué)習(xí)L1data cache之前建議先看一下這篇微信公眾號的文章，了解什么是VIVT，VIPT，PIPT等概念
CPU大多數(shù)的實現(xiàn)方式是VIPT，也就是Virtual Index Physical tag(取虛擬地址的index，然后轉(zhuǎn)換成物理地址取tag)，

總結(jié)

以上是生活随笔為你收集整理的GPGPU 架构的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。