// 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);
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);
4 DDR的時序參數(shù) 4.1 Row Active Command
在進(jìn)行數(shù)據(jù)的讀寫前，Controller 需要先發(fā)送 Row Active Command，打開 DRAM Memory Array 中的指定的 Row。Row Active Command 的時序如下圖所示：
(, 下載次數(shù): 12)
上傳
點(diǎn)擊文件名下載附件
tRCD：RAS－to－CAS Delay（tRCD），內(nèi)存行地址傳輸?shù)搅械刂返难舆t時間。
Row Active Command 通過地址總線指明需要打開某一個 Bank 的某一個 Row。DRAM 在接收到該 Command 后，會打開該 Row 的 Wordline，將其存儲的數(shù)據(jù)讀取到 Sense Amplifiers 中，這一時間定義為 tRCD（RCD for Row Address to Column Address Delay）。DRAM 在完成 Row Sense 階段后，Controller 就可以發(fā)送 Read 或 Write Command 進(jìn)行數(shù)據(jù)的讀寫了。這也意味著，Controller 在發(fā)送 Row Active Command 后，需要等待 tRCD 時間才能接著發(fā)送 Read 或者 Write Command 進(jìn)行數(shù)據(jù)的讀寫。 tRAS: Row Active Time,內(nèi)存行地址選通延遲
由于 DRAM 的特性，Row 中的數(shù)據(jù)在被讀取到 Sense Amplifiers 后，需要進(jìn)行 Restore 的操作。Restore 操作可以和數(shù)據(jù)的讀取同時進(jìn)行，即在這個階段，Controller 可能發(fā)送了 Read Command 進(jìn)行數(shù)據(jù)讀取。
DRAM 接收到 Row Active Command 到完成 Row Restore 操作所需要的時間定義為 tRAS（RAS for Row Address Strobe）。
Controller 在發(fā)出一個 Row Active Command 后，必須要等待 tRAS 時間后，才可以發(fā)起另一次的 Precharge 和 Row Access。 4.2 Column Read Command
Controller 發(fā)送 Row Active Command 并等待 tRCD 時間后，再發(fā)送 Column Write Command 進(jìn)行數(shù)據(jù)寫入。數(shù)據(jù) Burst Length 為 8 時的 Column Write Command 時序如下圖所示：
(, 下載次數(shù): 12)
上傳
點(diǎn)擊文件名下載附件
tCWD/tCL/tCWL:內(nèi)存CAS延遲時間
Column Write Command 通過地址總線 A[0:9] 指明需要寫入數(shù)據(jù)的 Column 的起始地址。Controller 在發(fā)送完 Write Command 后，需要等待 tCWD （CWD for Column Write Delay） 時間后，才可以發(fā)送待寫入的數(shù)據(jù)。tCWD 在一些描述中也稱為 tCWL（CWL for Column Write Latency） tWR（WR for Write Recovery）
DRAM 接收完數(shù)據(jù)后，需要一定的時間將數(shù)據(jù)寫入到 DRAM Cells 中，這個時間定義為 tWR（WR for Write Recovery）。該值說明在一個激活的bank中完成有效的寫操作及預(yù)充電前，必須等待多少個時鐘周期。這段必須的時鐘周期用來確保在預(yù)充電發(fā)生前，寫緩沖中的數(shù)據(jù)可以被寫進(jìn)內(nèi)存單元中。同樣的，過低的tWD雖然提高了系統(tǒng)性能，但可能導(dǎo)致數(shù)據(jù)還未被正確寫入到內(nèi)存單元中，就發(fā)生了預(yù)充電操作，會導(dǎo)致數(shù)據(jù)的丟失及損壞。 4.3 Precharge Command
要訪問 DRAM Cell 中的數(shù)據(jù)，需要先進(jìn)行 Precharge 操作。相應(yīng)地，在 Controller 發(fā)送 Row Active Command 訪問一個具體的 Row 前， Controller 需要發(fā)送 Precharge Command 對該 Row 所在的 Bank 進(jìn)行 Precharge 操作。下面的時序圖描述了 Controller 訪問一個 Row 后，執(zhí)行 Precharge，然后再訪問另一個 Row 的流程。
(, 下載次數(shù): 11)
上傳
點(diǎn)擊文件名下載附件
DRAM 執(zhí)行 Precharge Command 所需要的時間定義為 tRP（RP for Row Precharge）。Controller 在發(fā)送一個 Row Active Command 后，需要等待 tRC（RC for Row Cycle）時間后，才能發(fā)送第二個 Row Active Command 進(jìn)行另一個 Row 的訪問。
從時序圖上我們可以看到，tRC = tRAS + tRP，tRC 時間決定了訪問 DRAM 不同 Row 的性能。在實(shí)際的產(chǎn)品中，通常會通過降低 tRC 耗時或者在一個 Row Cycle 執(zhí)行盡可能多數(shù)據(jù)讀寫等方式來優(yōu)化性能。 4.4 Row Refresh Command
一般情況下，為了保證 DRAM 數(shù)據(jù)的有效性，Controller 每隔 tREFI（REFI for Refresh Interval） 時間就需要發(fā)送一個 Row Refresh Command 給 DRAM，進(jìn)行 Row 刷新操作。DRAM 在接收到 Row Refresh Command 后，會根據(jù)內(nèi)部 Refresh Counter 的值，對所有 Bank 的一個或者多個 Row 進(jìn)行刷新操作。
DRAM 刷新的操作與 Active + Precharge Command 組合類似，差別在于 Refresh Command 是對 DRAM 所有 Bank 同時進(jìn)行操作的。下圖為 DRAM Row Refresh Command 的時序圖：
(, 下載次數(shù): 14)
上傳
點(diǎn)擊文件名下載附件
DRAM 完成刷新操作所需的時間定義為 tRFC（RFC for Refresh Cycle）。
tRFC 包含兩個部分的時間，一是完成刷新操作所需要的時間，由于 DRAM Refresh 是同時對所有 Bank 進(jìn)行的，刷新操作會比單個 Row 的 Active + Precharge 操作需要更長的時間；tRFC 的另一部分時間則是為了降低平均功耗而引入的延時，DRAM Refresh 操作所消耗的電流會比單個 Row 的 Active + Precharge 操作要大的多，tRFC 中引入額外的時延可以限制 Refresh 操作的頻率。 4.5 Read Cycle
一個完整的 Burst Length 的 Read Cycle 如下圖所示：