Shared Memory

時間  2020-12-30

1、引言

在global memory部分，數據對齊和連續是提升性能的很重要的因素，當使用L1 cache的時候，對齊問題不再是問題，但是非連續的獲取內存依然會降低性能。依賴於算法本質，某些情況下，非連續訪問是不可避免的。使用shared memory是另一種提高性能的方式。

2、Introduction CUDA Shared Memory

GPU上的memory有兩種：

1、On-board memory

2、On-chip memory

global memory就是一塊很大的on-board memory，並且有很高的latency；而shared memory正好相反，是一塊很小、低延遲的on-chip memory，比global memory擁有高得多的帶寬。我們可以把他當做可編程的cache，其主要作用有：

1、An intra-block thread communication channel   一個block中線程間交流通道

2、A program-managed cache for global memory data  可編程的cache

3、Scratch pad memory for transforming data to improve global memory access patterns 

shared memory（SMEM）是GPU的重要組成之一。物理上，每個SM包含一個當前正在執行的block中所有thread共享的低延遲的內存池。SMEM使得同一個block中的thread能夠相互合作，重用on-chip數據，並且能夠顯著減少kernel需要的global memory帶寬。由於CUDA可以直接顯式的操作SMEM的內容，所以又被稱爲可編程緩存。
由於shared memory和L1要比L2和global memory更接近SM，shared memory的延遲比global memory低20到30倍，帶寬大約高10倍。
CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀
當一個block開始執行時，GPU會分配其一定數量的shared memory，這個shared memory的地址空間會由block中的所有thread 共享。shared memory是劃分給SM中駐留的所有block的，也是GPU的稀缺資源。所以，使用越多的shared memory，能夠並行的active thread blocks就越少。
關於Program-Managed Cache：在C語言編程裏，循環（loop transformation）一般都使用cache來優化。在循環遍歷的時候使用重新排列的迭代順序可以很好利用cache局部性。在算法層面上，我們需要手動調節循環來達到令人滿意的空間局部性，同時還要考慮cache size。cache對於程序員來說是透明的，編譯器會處理所有的數據移動，我們沒有能力控制cache的行爲。shared memory則是一個可編程可操作的cache，程序員可以完全控制其行爲。

Shared Memory Allocation

我們可以動態或者靜態的分配shared memory，其聲明即可以在kernel內部也可以作爲全局變量。CUDA支持聲明1D、2D和3D的shared memory數組。其標識符爲： __shared__。
靜態聲明2D浮點型數組： 
 
  
__shared__ float tile[size_y][size_x];
 
   如果在kernel中聲明的話，其作用域就是kernel內，否則是對所有kernel有效。如果shared memory的大小在編譯期未知的話，可以使用 
  extern關鍵字修飾，例如下面聲明一個未知大小的1D數組： 
 
 
 
 
  
 
  
extern __shared__ int tile[];
由於其大小在編譯期未知，我們需要在每個kernel調用時，動態的分配其shared memory，也就是最開始提及的第三個參數： 
 
 
 
 
  
 
  
kernel<<<grid, block, isize * sizeof(int)>>>(...)
注意： 
  只有1D數組才能這樣動態使用。 
 
 
 
 
  
Shared Memory Banks and Access Mode
 
 
 
 
 
   當優化內存性能時，有兩個重要的因素來量化：latency和bandwidth。shared memory能夠用來隱藏由於latency和bandwidth對性能的影響。下面將解釋shared memory的組織方式，以便研究其對性能的影響。 
 
 
 
 
  (1) Memory Banks 
 
 
 
 
   爲了獲得高帶寬，shared memory被分成32（ 
  計算能力1.x的device劃分爲16個banks）個相等大小的內存塊，每塊大小32-bit(4 bytes)，他們可以被同時訪問。不同的計算能力的device，shared memory以不同的模式映射到不同的塊（稍後詳解）。如果warp訪問shared memory，對於每個bank只訪問不多於一個內存地址，那麼只需要一次內存傳輸就可以了，否則需要多次傳輸，因此會降低內存帶寬的使用。 
 
 
 
 
  (2) Banks Conflict 
 
 
 
 
   當一個warp中多個地址請求落在同一個bank中就會發生bank conflict，從而導致請求多次執行。硬件會把這類請求分散到儘可能多的沒有conflict的那些傳輸操作裏面，降低有效帶寬的因素是被分散到的傳輸操作個數。 warp有三種典型的獲取shared memory的模式： 
 
 
 
 
  
· Parallel access：多個地址分散在多個bank。
 
  
· Serial access：多個地址落在同一個bank。
 
  
· Broadcast access：一個地址讀操作落在一個bank。
 
 
 
 
 
  

    Parallel access是最通常的模式，這個模式表示，一些（也可能是全部）地址請求能夠被一次傳輸解決。理想情況是，獲取無conflict的shared memory的時，每個地址都在落在不同的bank中。 
  
 
  

    Serial access是最壞的模式，如果warp中的32個thread都訪問了同一個bank中的不同位置，那就是32次單獨的請求，而不是同時訪問了。 
  
 
  

    Broadcast access也是隻執行一次傳輸，然後傳輸結果會廣播給所有發出請求的thread。這樣的話就會導致帶寬利用率低。 
  
 
 
 
 

   下圖是最優情況的訪問圖示，每個線程訪問一個32-bit的數據，不存在bank conflict： 
 
 
 
 
  
 
   CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
  
 下圖是不規律的隨機訪問模式，因爲每個thread訪問不同的bank，因此也沒有衝突： 
 
 
 
 
  
 
   CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
  
 下圖是bank衝突的情況，幾個threads訪問同一個bank，會產生下列兩種行爲： 
 
 
 
 
  
· Conflict-free broadcast access if threads access the same address within a bank
 
  
· Bank conflict access if threads access different addresses within a bank
 
 
 
 
 
  
 
   CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
  
  
  (3) Access Mode 
 
 
 
 
   根據device不同的計算能力版本，bank的大小配置也不同，具體爲： 
 
 
 
 
  
· 4 bytes for devices of compute capability 2.x
 
  
· 8 bytes for devices of compute capability 3.x
 
 
 
 

   以Fermi的GPU爲例，它有32個banks，每個bank 32-bit，即4 bytes，每個bank的帶寬是32bits每兩個cycle。連續的32位數據映射到連續的bank中，也就是說，bank的索引和shared memory地址的映射關係如下： 
 
 
 
 
  
 
  
 bank index = (byte address ÷ 4 bytes/bank) % 32 banks
下圖是Fermi的地址映射關係，注意到，bank中每個地址相差32，相鄰的word分到不同的bank中以便使warp能夠獲得更多的並行獲取內存操作（獲取連續內存時，連續地址分配到了不同bank中）。 
 
 
 
 
  CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
 
 
 

   當同一個warp的兩個thread要獲取同一個地址（注意是同一個地址，同一個bank會造成衝突）的時候並不發生bank conflict。對於讀操作，會用一次transaction獲得結果後廣播給所有請求，當寫操作時，只有一個thread會真正去寫，但是哪個thread執行是無法確定的。 
 
 
 
 
   對於Kepler設備來說，shared memory有兩種地址模式的32個banks： 
 
 
 
 
  
· 64-bit mode
 
  
· 32-bit mode
 
 
 
 

   在64-bit模式中，連續的64-bits字會映射到連續的bank。每個bank帶寬是64bite/1個clock。其映射關係公式： 
 
 
 
 
  
 
  
 bank index = (byte address ÷ 8 bytes/bank) % 32 banks
這裏，如果兩個thread訪問同一個64-bit中的任意一個sub-word（1byte）也不會導致bank conflict，因爲一次64-bit（bank帶寬64bit/cycle）的讀操作就可以滿足請求了。也就是說， 
  同等情況下，64-bit模式一般比32-bit模式更少碰到bank conflict。 
 
 
 
 
   下圖表示了32-bit模式下從字節地址到bank索引的映射關係圖。上面表示用字節地址和4-byte word索引標籤的共享內存，下面表示從4-byte word索引到bank索引的映射關係。儘管word 0和word 32都在bank0中，同時讀這兩個word也不會導致bank conflict（64-bit/cycle）： 
 
 
 
 
  CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
 
 
 

   下圖是64-bit模式下，conflict-free的情況，每個thread獲取不同的bank： 
 
 
 
 
  
 
   CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
  
下圖是另一種conflict-free情況，兩個thread或獲取同一個bank中的word： 
 
 
 
 
  
 
   
 
    CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
   
下圖紅色箭頭是三路bank conflict發生的情況： 
  
  
  CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
 
 
   
 (3) Memory Padding 
 
 
   memory padding是一種避免bank conflict的方法，如下圖所示，所有的thread分別訪問了bank0的5個不同的word，這時就會導致bank conflict，我們採取的方法就是在每N（bank數目）個word後面加一個word，這樣就如下面右圖那樣，原本bank0的每個word轉移到了不同的bank中，從而避免了bank conflict。 
 
 
 
 
  
 
   CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
  
 
   增加的這些word不會用來存儲數據，其唯一的作用就是移動原始bank中的word，避免衝突。使用memory padding會導致block可獲得shared memory中有用的數量減少。內存填充後，需要重新計算數組索引以確保可以訪問到正確的數據。注意Fermi和Kepler的bank寬度不同，所以針對於Fermi的填充模型用於Kepler可能導致bank conflict。 
 
 
 

     
  (3) Access Mode Configuration 
 
 
 
 
   對Kepler來說，默認情況是4-byte模式，可以用下面的API來查看： 
 
 
 
 
  
 
  
 cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig *pConfig);
返回結果放在pConfig中，其結果可以是下面兩種： 
 
 
 
 
  
cudaSharedMemBankSizeFourByte
 
  
cudaSharedMemBankSizeEightByte
 
 
 
 

   可以使用下面的API來設置bank的大小： 
 
 
 
 
  
 
  
 cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);
 
 
 
 

   bank的配置參數如下三種： 
 
 
 
 
  
cudaSharedMemBankSizeDefault
 
  
cudaSharedMemBankSizeFourByte
 
  
cudaSharedMemBankSizeEightByte
 
 
 
 

   在啓動不同的kernel之間修改bank配置會有一個隱式的device同步。 
  修改shared memory的bank大小不會增加shared memory的利用率或者影響kernel的Occupancy，但是對性能是一個主要的影響因素。一個大的bank會產生較高的帶寬，但是鑑於不同的access pattern，可能導致更多的bank conflict。 
 
 
 
 
  
Configuring the Amount of Shared Memory
 
 
 
 
 
   每個SM擁有64 KB的片上內存，shared memory和L1 cache共享這塊內存。CUDA提供了兩種方式配置它們各自的大小，參考 
  CUDA學習----Memory Model，裏面提到的是 
  Per-kernel configuration的情況。 
 
 
 
 
  
Per-device configuration
 
  
Per-kernel configuration
 
 
 
 
 
  
 
  
cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);
 
 
 
 

   對於 
  Per-device configuration的情況用以上API，兩種情況類似，只是作用範圍不同。配置方式孰優孰劣取決於kernel用的shared memory的多少。 
 
 
 
 
  
Synchronization
 
 
 
 
 
   因爲shared memory可以被同一個block中的不同的thread同時訪問，當同一個地址的值被多個thread修改就導致了inter-thread conflict，所以我們需要同步操作。CUDA提供了兩類block內部的同步操作，即： 
 
 
 
 
  
· Barriers
 
  
· Memory fences
 
 
 
 

   對於barriers，所有thread會等待其他threads到達barrier point；對於Memory fence，所有threads會被阻塞直到所有修改memory的操作對其他threads可見。下面解釋下CUDA需要同步的主要原因：weakly-ordered。 
 
 
 
 
  (1) Weakly-Ordered Memory Model 
 
 
 
 
  
 
    現代內存架構有非常寬鬆的內存模式，也就是意味着，memory的獲取不必按照程序中的順序來執行。CUDA採用了一種叫做weakly-ordered內存模型來獲取更激進的編譯器優化。 
  
 
  
 
    GPU thread寫數據到不同的memory的順序（比如shared memory，global memory，page-locked host memory或者另一個device上的memory）同樣沒必要跟程序裏面順序相同。一個thread的讀操作的順序對其他thread可見時也可能與實際執行寫操作的thread順序不一致。 
  
 
  
 
    爲了顯式的強制程序以一個確切的順序運行，就需要用到fence和barrier。它們也是唯一的方式能保證kernel與其它線程分享資源時對memory有正確行爲。 
  
 
 
 
 
 
  (2) Explicit Barrier 
 
 
 
 
   可以在kernel中設置一個barrier point通過調用下列函數： 
 
 
 
 
  
 
  
void __syncthreads();
 
  

    __syncthreads就是作爲一個barrier point起作用， 
   block中的threads必須等待所有thread都到達這個point後才能繼續下一步。這也保證了所有在這個point之前獲取global memory和shared memory的操作對同一個block中所有thread可見。 
    __syncthreads被用來協作同一個block中的thread通信。當一些thread獲取memory相同的地址時，就會導致潛在的問題（read-after-write，write-after-read和write-after-write）從而引起未定義行爲狀態，此時就可以使用__syncthreads來避免這種情況。 
  
 
  
 
    使用__syncthreads要相當小心，只有在所有thread都會到達這個point時纔可以調用這個同步，顯而易見， 
   如果同一個block中的某些thread永遠都不能到達該point，那麼程序將一直等下去，下面代碼就是一種錯誤的使用方式： 
  
 
 
 
 
 
  
 
  
if (threadID % 2 == 0)
 __syncthreads();
 
  
else
 __syncthreads();
如果在block之間不同步的話，thread blocks可能以任意順序，並行或者串行，在任意的SM上被執行。如果一個CUDA kernel需要全局同步，可以通過在同步點分割kernel和啓動多個kernel來達到這種期望的行爲。 
 
 
 
 
  (3) Memory Fence 
 
 
 
 
   這種方式保證了任何在fence之前的memory寫操作對fence之後thread都可見，也就是，fence之前寫完了，fence之後其它thread就都知道這塊memory寫後的值了。fence的設置範圍比較廣，分爲：block，grid和system。 可以通過下面的API來設置fence： 
 
 
 
 
 
void __threadfence_block();
 
 
void __threadfence();
 
 
void __threadfence_system();
 
 其中，第一個函數是對應的block範圍的，也就是保證同一個block中thread在fence之前寫完的值對block中其它的thread可見，不同於barrier，該函數不需要所有的thread都執行；第二個函數是對應grid範圍的；第三個對用system的，其範圍針對整個系統，包括device和host。 
 
 
 
 
  (4) Volatile Qualifier 
 
 
 
 
   聲明一個使用global memory或者shared memory的變量，用volatile修飾符來修飾該變量的話，會組織編譯器做一個該變量的cache優化。使用該修飾符後，編譯器就會認爲該變量可能在某一時刻被別的thread改變，如果使用cache優化的話，得到的值就缺乏時效，因此使用volatile強制每次都到global 或者shared memory中去讀取其絕對有效值。 
 
 
 
3、Checking the Data Layout of Shared Memory
 
 
 
   我們在設計使用shared Memory的時候應該關注下面的信息： 
 
 
 
 
  
· Mapping data elements across Memory banks
 
  
· Mapping from thread index to shared Memory offset
 
 
 
 

   搞明白這兩點，就可以掌握shared memory的使用了，從而構建出高性能的代碼。 
 
 
 
 
  
Square Shared Memory
 
 
 
 
 
   我們可以以一種直接的方式用shared memory緩存全局內存中的方陣。下圖展示了一個每一維度有32個元素並以row-major存儲在shared memory的數組，圖的最上方是該矩陣實際的一維存儲圖示，下方是通過映射4-byte數據和banks關係的邏輯二維shared memory： 
 
 
 
 
  
 
   CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
  
 我們可以使用下面的語句靜態聲明一個2D的shared memory變量： 
 
 
 
 
  
 
  
__shared__ int tile[N][N];
因爲是方陣，可以從2D線程塊中以相鄰的thread獲取相鄰的元素的方式訪問數據： 
 
 
 
 
  
 tile[threadIdx.y][threadIdx.x]
 
  
 tile[threadIdx.x][threadIdx.y]
 
 
 
 
 
   上面兩種方式哪個性能更好呢？這就需要注意thread和bank的映射關係了，我們最希望看到的是同一個warp中的thread獲取的是不同的bank。同一個warp中的thread可以使用連續的threadIdx.x來確定。不同bank中的元素同樣是連續存儲的，以word大小作爲偏移。因此，最好是讓連續的thread（由連續的threadIdx.x確定）獲取shared memory中連續的地址， 
  由此得知，tile[threadIdx.y][threadIdx.x]應該展現出更好的性能以及更少的bank conflict。 
 
 
 
 
  (1) Access Row-Major versus Column-Major 
 
 
 
 
   假設我們的grid有2D的block(32,32)，定義如下： 
 
 
 
 
  
 
  
#define BDIMX 32
#define BDIMY 32
dim3 block(BDIMX,BDIMY);
dim3 grid(1,1);
我們對這個kernel有如下兩個操作： 
 
 
 
 
  
· 將thread索引以row-major寫到2D的shared memory數組中；
 
  
· 從shared memory中讀取這些值並寫入到global memory中。
 
 
 
 

   kernel代碼如下： 
 
 
 
 
  
 
  
__global__ void setRowReadRow(int *out) {
 // declare static 2D shared memory
    __shared__ int tile[BDIMY][BDIMX];
 // 因爲block只有一個
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
 // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;
 // 這裏同步是爲了使下面shared memory的獲取以row-major執行
 // 避免若有的線程未完成，而其他線程已經在讀shared memory的情況
 __syncthreads();
 // shared memory load operation
    out[idx] = tile[threadIdx.y][threadIdx.x] ;
} 
 
  

    此段有三個內存操作：(1)向shared Memory存數據；(2)從shared Memor取數據；(3)向global Memory存數據。 
  
 
 
 
 
 
   因爲在同一個warp中的thread使用連續的threadIdx.x來檢索title，該kernel是沒有bank conflict的。如果交換上述代碼threadIdx.y和threadIdx.x的位置，就變成了column-major的順序。每個shared memory的讀寫都會導致Fermi上32-way的bank conflict或者在Kepler上16-way的bank conflict。 
 
 
 
 
  
 
  
__global__ void setColReadCol(int *out) {
 // static shared memor
    __shared__ int tile[BDIMX][BDIMY];
 // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
 // shared memory store operation
    tile[threadIdx.x][threadIdx.y] = idx;
 // wait for all threads to complete
 __syncthreads();
 // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
} 
編譯運行結果如下(在K40上以4-byte模式運行)： 
 
 
 
 
  
 
  
$ nvcc checkSmemSquare.cu –o smemSquare
$ nvprof ./smemSquare
 
  
./smemSquare at device 0 of Tesla K40c with Bank Mode:4-byte
<<< grid (1,1) block (32,32)>>
Time(%) Time Calls Avg Min Max Name
13.25% 2.6880us 1 2.6880us 2.6880us 2.6880us setColReadCol(int*)
11.36% 2.3040us 1 2.3040us 2.3040us 2.3040us setRowReadRow(int*)
 
  從結果可以看出，row-major的kernel表現更出色。 
 
 
 

   然後 
  使用nvprof命令的下面的兩個參數來衡量相應的bank-conflict： 
 
 
 
 
  
 
  
 shared_load_transactions_per_request shared_store_transactions_per_request
運行結果如下(K40，8-byte模式下)，row-major只有一次transaction，而column-major需要16次，如果4-byte模式下，可能需要32次： 
 
 
 
 
  
 
  
Kernel:setColReadCol (int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 16.000000
Kernel:setRowReadRow(int*)
1 shared_load_transactions_per_request 1.000000
1 shared_store_transactions_per_request 1.000000
Writing Row-Major and Reading Column-Major
 
  (2) Writing Row-Major and Reading Column-Major 
 
 
   
 下面代碼實現以row-major寫shared memory，以column-major讀shared memory： 
 
 
  
 
  
__global__ void setRowReadCol(int *out) {
 // static shared memory
    __shared__ int tile[BDIMY][BDIMX];
 // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
 // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;
 // wait for all threads to complete
 __syncthreads();
 // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
} 
下圖展示了用簡單的5路bank shared memory實現兩種內存操作： 
 
 
 
 
  
 
   CUDA性能優化----Shared Memory - 樂不思蜀 - 樂不思蜀 
  
 用nvprof命令查看相關bank conflict情況： 
 
 
 
 
  
 
  
Kernel:setRowReadCol (int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 1.000000
從結果可以看出：寫操作是沒有conflict的，讀操作則引起了一個16次的transaction。 
 
 
 
 
  
 
 
 
 
Dynamic Shared Memory
 
 
正如前文所說，我們可以全局範圍的動態聲明shared Memory，也可以在kernel內部動態聲明一個局部範圍的shared Memory。注意，動態聲明必須是未確定大小一維數組，因此，我們就需要重新計算索引。因爲我們將要以row-major寫，以colu-major讀，所以就需要保持下面兩個索引值：
 
 
· row_idx：1D row-major 內存的偏移
 
 
· col_idx：1D column-major內存偏移
 
 
kernel代碼：
 
 
 
  
 
   複製代碼 
  
 
  
__global__ void setRowReadColDyn(int *out) {
    // dynamic shared memory
    extern __shared__ int tile[];
    // mapping from thread index to global memory index
    unsigned int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned int col_idx = threadIdx.x * blockDim.y + threadIdx.y;
    // shared memory store operation
    tile[row_idx] = row_idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[row_idx] = tile[col_idx];
}            
 
  
 
   複製代碼 
  
 
 
 
 
kernel調用時配置的shared Memory：
 
 
setRowReadColDyn<<<grid, block, BDIMX * BDIMY * sizeof(int)>>>(d_C);
 
 
查看transaction：
 
 
 
  
Kernel: setRowReadColDyn(int*)
1 shared_load_transactions_per_request 16.000000
1 shared_store_transactions_per_request 1.000000
 
 
 
 
該結果和之前的例子相同，不過這裏使用的是動態聲明。
 
 
Padding Statically Declared Shared Memory
 
 
直接看kernel代碼：
 
 
 
  
 
   複製代碼 
  
 
  
__global__ void setRowReadColPad(int *out) {
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX+IPAD];
    // mapping from thread index to global memory offset
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;
    // wait for all threads to complete
    __syncthreads();
    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}                            
 
  
 
   複製代碼 
  
 
 
 
 
改代碼是setRowReadCol的翻版，查看結果：
 
 
 
  
Kernel: setRowReadColPad(int*)
1 shared_load_transactions_per_request 1.000000
1 shared_store_transactions_per_request 1.000000
 
 
 
 
正如期望的那樣，load的bank_conflict已經消失。在Fermi上，只需要加上一列就可以解決bank-conflict，但是在Kepler上卻不一定，這取決於2D shared Memory的大小，因此對於8-byte模式，可能需要多次試驗才能得到正確結果。
 
 
 
 
 
  
 
 
 
 
原文鏈接：http://blog.163.com/[email protected]/blog/static/71988399201712735436357/
 
 
                http://www.javashuo.com/article/p-syidsyra-ge.html
 


                




                

            

        

    










    

        联系我们
        沪ICP备13005482号-10