共享記憶體啟用塊中線程之間的協作。當乙個塊中的多個執行緒使用全域性記憶體中的相同資料時，共享記憶體只能用於從全域性記憶體訪問一次資料。共享記憶體還可以用來避免未合併的記憶體訪問，方法是從全域性記憶體中載入和儲存合併模式的資料，然後將其重新排列在共享記憶體中。除了儲存體衝突之外，共享記憶體中的變形對非連續或未對齊訪問不會造成任何損失。
對於具有維度mxw的a，具有維度wxn的維度b以及維度mxn的維度c的情況，通過矩陣乘法c = ab的簡單示例來說明共享儲存器的使用。為了簡化核心，m和n是32的倍數，對於計算能力2.0或更高的裝置，w為32。
問題的自然分解是使用wxw執行緒的塊和瓦片大小。因此，就wxw瓦片而言，a是列矩陣，b是行矩陣，c是它們的外積;參見圖9.以m / w塊為單位的n / w網格被啟動，其中每個執行緒塊從a的單個塊和b的單個塊計算c中不同塊的元素
為此，******multiply核心（未優化矩陣乘法）計算矩陣c的乙個瓦片的輸出元素。
未經優化的矩陣乘法：
__global__ void ******multiply(float *a, float* b, float *c,
int n)
c[row*n + col] = sum;
}
c.行和列是c中元素的行和列由特定的執行緒計算。 for迴圈結束我將一行a乘以b的列，然後寫入c.
該核心的有效頻寬在nvidia tesla k20x上僅為6.6gb / s（ecc關閉）。 為了分析效能，有必要考慮warp在for迴圈中如何訪問全域性記憶體。 每個執行緒的變形計算c的乙個圖塊的一行，這取決於a的單個行和b的整個圖塊，如圖10所示。
對於for迴圈的每個迭代i，warp中的執行緒讀取b tile的一行，這是所有計算功能的順序和合併訪問。
但是，對於每次迭代i而言，warp中的所有執行緒都從矩陣a的全域性記憶體中讀取相同的值，因為索引行* tile_dim + i在warp內是不變的。 儘管在計算能力2.0的裝置上這種訪問只需要1次事務 或更高，因為事務中存在浪費的頻寬，因為快取記憶體行中的32個字中只有乙個4位元組的字被使用。 我們可以在迴圈的後續迭代中重用這個快取行，我們最終將使用全部32個字; 然而，當許多warps同時在同乙個多處理器上執行時，就像通常情況一樣，快取記憶體行可能容易從迭代i和i + 1之間的快取記憶體中逐出。
CUDA實踐指南（十三）
裝置記憶體空間 cuda裝置使用多個記憶體空間，這些記憶體空間具有不同的特性，這些特性反映了它們在cuda應用程式中的不同用法 這些記憶體空間包括全域性，本地，共享，紋理和暫存器，如圖2所示 合併的訪問要求取決於裝置的計算能力，可以查閱cuda c程式設計指南中。在這些不同的記憶空間中，全域性記憶是...
CUDA實踐指南（九）
頻寬 頻寬 資料可以傳輸的速率 是效能最重要的門控因素之一。的幾乎所有變化都應該在它們如何影響頻寬的情況下進行。頻寬可能會因資料儲存的記憶體選擇，資料布局方式以及訪問順序以及其他因素而受到顯著影響。要準確測量效能，計算理論和有效頻寬非常有用。當後者遠低於前者時，設計或實現細節可能會降低頻寬，並且應該...
CUDA實踐指南（七）
數值精度和精度 由於計算和儲存浮點值的方式，導致錯誤或意外的結果主要來自浮點精度問題。以下各節介紹了感興趣的主要專案。cuda c程式設計指南的特性和技術規範 以及有關浮點精度和效能的 和附帶網路研討會中提供了浮點運算的其他特性，可從獲得精密performancefloating 點和ieee 75...