CUDA程式設計指南閱讀筆記（六）

接上文。

共享記憶體是cuda裝置中非常重要的乙個儲存區域，有效地使用共享記憶體可以充分利用cuda裝置的潛能，極大提公升程式效能。那麼，共享記憶體有哪些特點呢？

1、共享記憶體（shared memory）是整合在gpu處理器晶元上的（on-chip），因此相比於存在於視訊記憶體顆粒中的全域性記憶體（global memory）和本地記憶體（local memory），它具有更高的傳輸頻寬，一般情況下，共享記憶體的頻寬大約是全域性記憶體頻寬的7-10倍。

2、共享記憶體的容量很小。根據nvidia官方文件的說法，在計算能力1.x的裝置中，每乙個流多處理器（streaming multiprocessor）上的共享記憶體容量為16kb。對於計算能力2.x、3.0及3.5的裝置該引數為48kb。因此共享記憶體是稀有資源。

3、共享記憶體在物理上被劃分為很多塊，每一塊被稱為乙個儲存體（bank）。在同一時刻，cuda裝置可以同時訪問多個儲存體。因此，如果一次針對共享記憶體的訪存操作需要讀取n個位址，而這n個位址恰好分布在n個不同的儲存體（bank）中，那麼只需要乙個訪問週期就可以完成n個位址的訪存任務了。對於計算能力1.x的裝置，共享記憶體被平均劃分為16個儲存體。而對於計算能力2.x、3.0及3.5的裝置此引數為32。在共享記憶體中，相鄰兩塊32bit的資料分別屬於相鄰的兩個儲存體。儲存體每兩個時鐘週期可以傳輸32位資料。

4、共享記憶體既可以靜態分配，也可以動態分配。

從共享記憶體的這些特點中我們可以看出，它實際上相當於乙個程式設計師可以操控的快取（cache），下面，我們使用矩陣乘法的例子來說明如何有效使用共享記憶體。

首先，我們使用最直觀的方法來完成矩陣乘法c = a x b：讀取a的每一行和b的每一列，順次完成計算任務。矩陣乘法的示意圖如下所示：

下面是矩陣乘法的cuda c主要實現**：

[cpp]view plain

copy

// matrices are stored in row-major order:

// m(row, col) = *(m.elements + row * m.width + col)

typedef

struct

matrix;

// thread block size

#define block_size 16

// forward declaration of the matrix multiplication kernel

__global__ void

matmulkernel(

const

matrix,

const

matrix, matrix);

// matrix multiplication - host code

// matrix dimensions are assumed to be multiples of block_size

void

matmul(

const

matrix a,

const

matrix b, matrix c)

// matrix multiplication kernel called by matmul()

__global__ void

matmulkernel(matrix a, matrix b, matrix c)

可以看出，為了計算矩陣c的任何乙個元素，程式都需要從全域性記憶體（global memory）中獲得矩陣a的一行和矩陣b的一列。因此，完成這一計算矩陣a被讀取了b.width次，矩陣b被讀取了a.height次。

現在我們來使用共享記憶體（shared memory）實現矩陣乘法。假設矩陣c可以被劃分為若干個較小的子方陣c

sub，我們使用乙個執行緒塊（thread block）來負責某一子方陣的計算，執行緒塊中的每乙個執行緒（thread）正好負責子方陣c

sub中乙個元素的計算。這樣劃分後，任何乙個結果子方陣c

sub'（尺寸為block_size * block_size）都是與該方陣具有相同行索引的尺寸為a.width * block_size的a的子矩陣a

sub和與該方陣具有相同列索引的尺寸為block_size * b.height的b的子矩陣b

sub相乘所得到。

為了匹配裝置的計算資源，兩個子矩陣asub和bsub被劃分為盡可能多的分離的維度為block_size的子方陣，csub的值便是這些子矩陣相乘後相加所得到的結果。子矩陣乘法的執行順序都是首先將它們從全域性記憶體（global memory）拷貝到共享記憶體（shared memory）（執行緒塊中的每乙個執行緒正好負責方陣乙個元素的拷貝），然後由執行緒自己完成相應元素的計算任務，利用暫存器儲存區域性結果，最後將暫存器的內容與新得到的計算結果依此累加起來得到最終運算結果並將其傳輸到全域性記憶體（global memory）中。

通過使用這種分治的計算策略，共享記憶體得到了很好的利用，採用這種方案計算完成時全域性記憶體中矩陣a被訪問的次數為b.width / block_size，矩陣b被訪問的次數為a.height / block_size，很明顯，這為我們節省了非常多的全域性記憶體頻寬。優化後的矩陣計算示意圖如下所示：

為了提公升計算效率，我們為型別matrix增加了乙個成員變數stride。__device__函式用來獲得和設定子矩陣的元素。下面是優化後的**：

[cpp]view plain

copy

// matrices are stored in row-major order;

// m(row, col) = *(m.elements + row * m.stride + col)

typedef

struct

matrix;

// get a matrix element

__device__ float

getelement(

const

matrix a,

introw,

intcol)

// set a matrix element

__device__ void

setelement(matrix a,

introw,

intcol,

float

value)

// get the block_sizexblock_size sub-matrix asub of a that is

// located col sub-matrices to the right and row sub-matrices down

// from the upper-left corner of a

__device__ matrix getsubmatrix(matrix a, int

row,

intcol)

// thread block size

#define block_size 16

// forward declaration of the matrix multiplication kernel

__global__ void

matmulkernel(

const

matrix,

const

matrix, matrix);

// matrix multiplication - host code

// matrix dimensions are assumed to be multiples of block_size

void

matmul(

const

matrix a,

const

matrix b, matrix c)

// matrix multiplication kernel called by matmul()

__global__ void

matmulkernel(matrix a, matrix b, matrix c)

// write csub to device memory

// each thread writes one element

setelement(csub, row, col, cvalue); }

同乙個block中的執行緒在95行的for迴圈中獲取到的asub，bsub，csub是一樣的，每個執行緒就負責csub內元素的計算

CUDA程式設計指南閱讀筆記（六）

CUDA程式設計指南閱讀筆記（六）

CUDA程式設計指南閱讀筆記（四）

《HTTP權威指南》閱讀筆記（六）

CUDA程式設計指南閱讀筆記（六）

CUDA程式設計指南閱讀筆記（六）

CUDA程式設計指南閱讀筆記（四）

《HTTP權威指南》閱讀筆記（六）

相關推薦