CUDA運行時 Runtime（二）
一. 概述
下面的代碼示例是利用共享內存的矩陣乘法的實現。在這個實現中，每個線程塊負責計算C的一個方子矩陣C sub，塊內的每個線程負責計算Csub的一個元素。如圖10所示，Csub等于兩個矩形矩陣的乘積：與Csub具有相同行索引的維度A（A.width，block_size）的子矩陣和與Csub具有相同列索引的維度B（block_size，A.width）的子矩陣。為了適應設備的資源，這兩個矩形矩陣根據需要被劃分為任意多個尺寸塊的正方形矩陣，并且計算Csub作為這些正方形矩陣的乘積之和。這些產品中的每一個都是通過首先將兩個對應的方陣從全局內存加載到共享內存，其中一個線程加載每個矩陣的一個元素，然后讓每個線程計算產品的一個元素來執行的。每個線程將這些產品的結果累積到一個寄存器中，并在完成后將結果寫入全局內存。

通過這種方式阻塞計算，我們利用快速共享內存并節省大量全局內存帶寬，因為a僅從全局內存讀取（B.width/block_size）次，B讀取（a.height/block_size）次。

前一個代碼示例中的矩陣類型增加了一個跨距字段，這樣子矩陣就可以用相同的類型有效地表示。__device_函數用于獲取和設置元素，并從矩陣構建任何子矩陣。

// Matrices are stored in row-major order:

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

typedef struct {

int width;int height;int stride; float* elements;

} Matrix;

// Get a matrix element

device float GetElement(const Matrix A, int row, int col)

{
return A.elements[row * A.stride + col];
}

// Set a matrix element

device void SetElement(Matrix A, int row, int col, float value)
{
A.elements[row * A.stride + col] = 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, int col)

{

Matrix Asub;Asub.width    = BLOCK_SIZE;Asub.height   = BLOCK_SIZE;Asub.stride   = A.stride;Asub.elements =

&A.elements[A.stride * BLOCK_SIZE * row

• BLOCK_SIZE * col];

  return Asub;
  }

// 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)

{

// Load A and B to device memoryMatrix d_A;d_A.width = d_A.stride = A.width; d_A.height = A.height;size_t size = A.width * A.height

• sizeof(float);

  cudaMalloc(&d_A.elements, size);

  cudaMemcpy(d_A.elements, A.elements, size,
  cudaMemcpyHostToDevice);

  Matrix d_B;

  d_B.width = d_B.stride = B.width; d_B.height = B.height;

  size = B.width * B.height * sizeof(float);

  cudaMalloc(&d_B.elements, size);

  cudaMemcpy(d_B.elements, B.elements, size,

  cudaMemcpyHostToDevice);
  // Allocate C in device memory

  Matrix d_C;

  d_C.width = d_C.stride = C.width; d_C.height = C.height;

  size = C.width * C.height * sizeof(float);

  cudaMalloc(&d_C.elements, size);
  // Invoke kernel

  dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);

  dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);

  MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
  // Read C from device memory

  cudaMemcpy(C.elements,
  d_C.elements, size, cudaMemcpyDeviceToHost);

  // Free device memory

  cudaFree(d_A.elements);

  cudaFree(d_B.elements);

  cudaFree(d_C.elements);

}
// Matrix multiplication kernel called by MatMul()

global void MatMulKernel(Matrix A, Matrix B, Matrix C)

{

// Block row and columnint blockRow = blockIdx.y;int blockCol = blockIdx.x;// Each thread block computes one sub-matrix Csub of CMatrix Csub = GetSubMatrix(C, blockRow, blockCol);// Each thread computes one element of Csub// by accumulating results into Cvaluefloat Cvalue = 0;// Thread row and column within Csubint row = threadIdx.y;int col = threadIdx.x;// Loop over all the sub-matrices of A and B that are// required to compute Csub// Multiply each pair of sub-matrices together// and accumulate the resultsfor (int m = 0; m < (A.width / BLOCK_SIZE); ++m) 
{// Get sub-matrix Asub of AMatrix Asub = GetSubMatrix(A, blockRow, m);// Get sub-matrix Bsub of BMatrix Bsub = GetSubMatrix(B, m, blockCol);// Shared memory used to store Asub and Bsub respectively__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];// Load Asub and Bsub from device memory to shared memory// Each thread loads one element of each sub-matrixAs[row][col] = GetElement(Asub, row, col);Bs[row][col] = GetElement(Bsub, row, col);// Synchronize to make sure the sub-matrices are loaded// before starting the computation__syncthreads();// Multiply Asub and Bsub togetherfor (int e = 0; e < BLOCK_SIZE; ++e)Cvalue += As[row][e] * Bs[e][col];// Synchronize to make sure that the preceding// computation is done before loading two new// sub-matrices of A and B in the next iteration__syncthreads();}// Write Csub to device memory// Each thread writes one elementSetElement(Csub, row, col, Cvalue);

}

圖10. 共享內存矩陣乘法

二. 頁面鎖定主機內存

運行時提供允許使用頁鎖定（也稱為固定）主機內存（而不是malloc（）分配的常規可分頁主機內存）的函數：

cudaHostAlloc（）和cudaFreeHost（）分配并釋放頁面鎖定的主機內存；

cudaHostRegister（）頁鎖定malloc（）分配的內存范圍（有關限制，請參閱參考手冊）。

使用頁鎖定主機內存有幾個好處：

對于異步并發執行中提到的某些設備，頁面鎖定的主機內存和設備內存之間的復制可以與內核執行同時執行。

在某些設備上，頁鎖定的主機內存可以映射到設備的地址空間，從而無需將其復制到設備內存或從設備內存復制，如映射內存中所術。

在具有前端總線的系統上，如果主機內存被分配為頁鎖定，則主機內存和設備內存之間的帶寬更高，如果另外它被分配為寫入合并，則帶寬更高，如寫入合并內存中所述。

但是，頁鎖定的主機內存是一種稀缺資源，因此，在頁鎖定內存中的分配將在可分頁內存中的分配之前很長一段時間開始失敗。此外，通過減少操作系統可用于分頁的物理內存量，消耗過多的頁鎖定內存會降低總體系統性能。

注意：頁面鎖定的主機內存不緩存在非I/O一致的Tegra設備上。此外，非I/O相干Tegra設備不支持cudaHostRegister（）。

簡單的零拷貝CUDA示例附帶了一個關于頁面鎖定內存api的詳細文檔。

三． 便攜式存儲器

頁面鎖定內存塊可以與系統中的任何設備一起使用（有關多設備系統的詳細信息，請參閱多設備系統），但默認情況下，使用上面描述的頁鎖定內存的好處僅與分配塊時的當前設備（以及所有設備共享相同的統一地址空間（如果有的話，如統一虛擬地址空間中所述）結合使用。要使這些優勢對所有設備都可用，需要通過將標志cudaHostAllocPortable傳遞給cudaHostAlloc（）來分配塊，或者通過將標志cudaHostRegisterPortable傳遞給cudaHostRegister（）來鎖定頁。

四． 寫入組合存儲器

默認情況下，頁鎖定的主機內存被分配為可緩存的。通過將標志cudaHostAllocWriteCombined傳遞給cudaHostAlloc（），可以選擇將其分配為寫合并。寫組合內存釋放主機的一級和二級緩存資源，使更多緩存可用于應用程序的其余部分。此外，在跨PCI Express總線傳輸期間，不會窺探寫合并內存，這可以將傳輸性能提高高達40%。

從主機讀取寫組合內存的速度非常慢，因此通常應將寫組合內存用于主機只寫入的內存。

五 . 映射內存

也可以通過傳遞標記cudahostallocmaped to cudaHostAlloc（）或傳遞標記cudahostragistermapped
to cudahostratrigister（）將頁鎖定主機內存塊映射到設備的地址空間。因此，這樣的塊通常有兩個地址：一個在主機內存中，由cudaHostAlloc（）或malloc（）返回；另一個在設備內存中，可以使用cudaHostGetDevicePointer（）檢索，然后用于從內核中訪問塊。唯一的例外是使用cudaHostAlloc（）分配的指針，以及在統一虛擬地址空間中為主機和設備使用統一地址空間時。

直接從內核中訪問主機內存不會提供與設備內存相同的帶寬，但確實有一些優點：

l 不需要在設備內存中分配一個塊并在這個塊和主機內存中的塊之間復制數據；數據傳輸是根據內核的需要隱式執行的；

l 不需要使用流（參見并發數據傳輸）來將數據傳輸與內核執行重疊；內核發起的數據傳輸會自動與內核執行重疊。

但是，由于映射的頁鎖定內存在主機和設備之間共享，應用程序必須使用流或事件同步內存訪問（請參閱異步并發執行），以避免任何潛在的讀后寫、讀后寫或寫后寫危險。
要能夠檢索指向任何映射的頁鎖定內存的設備指針，在執行任何其他CUDA調用之前，必須通過使用cudaDeviceMapHost標志調用cudaSetDeviceFlags（）來啟用頁鎖定內存映射。否則，cudaHostGetDevicePointer（）將返回錯誤。

如果設備不支持映射頁鎖定的主機內存，cudaHostGetDevicePointer（）也會返回錯誤。應用程序可以通過檢查canMapHostMemory設備屬性（請參閱設備枚舉）來查詢此功能，對于支持映射頁鎖定主機內存的設備，該屬性等于1。

請注意，從主機或其他設備的角度來看，在映射頁鎖定內存上運行的原子函數（請參閱原子函數）不是原子函數。

還要注意，CUDA運行時要求從主機和其他設備的角度，將從設備啟動的1字節、2字節、4字節和8字節自然對齊的加載和存儲保存為單個訪問。在某些平臺上，內存原子可能會被硬件分解為單獨的加載和存儲操作。這些組件加載和存儲操作對保持自然對齊的訪問具有相同的要求。例如，CUDA運行時不支持PCI Express總線拓撲，其中PCI Express網橋在設備和主機之間將8字節自然對齊的寫入拆分為兩個4字節的寫入。

總結

以上是生活随笔為你收集整理的CUDA运行时 Runtime（二）的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。