CUDA運行時 Runtime（一）
一. 概述
運行時在cudart庫中實現，該庫通過靜態方式鏈接到應用程序庫cudart.lib和libcudart.a，或動態通過cudart.dll或者libcudart.so. 需要cudart.dll和/或libcudart。索對於動態鏈接，通常將它們作為應用程序安裝包的一部分包括在內。只有在鏈接到CUDA運行時的同一實例的組件之間傳遞CUDA運行時符號的地址才是安全的。
它的所有入口點都以cuda為前綴。
正如在異構編程中所提到的，CUDA編程模型假設一個系統由一個主機和一個設備組成，每個主機和設備都有各自獨立的內存。設備內存概述了用於管理設備內存的運行時函數。
共享內存說明了如何使用線程層次結構中引入的共享內存來最大限度地提高性能。
頁面鎖定主機內存引入了頁面鎖定主機內存，它需要將內核執行與主機和設備內存之間的數據傳輸重疊起來。
異步並發執行描述了用於在系統的各個級別上啟用異步並發執行的概念和API。
多設備系統顯示了編程模型如何擴展到多個設備連接到同一主機的系統。
錯誤檢查描述如何正確檢查運行時生成的錯誤。
調用堆棧提到用於管理CUDA C++調用堆棧的運行時函數。
紋理和表面存儲器提供了紋理和表面存儲器空間，它們提供了訪問設備存儲器的另一種方式；它們還公開了GPU紋理硬件的一個子集。
圖形互操作性引入了運行時提供的與兩個主要圖形api（OpenGL和Direct3D）互操作的各種功能。
二．初始化
運行時沒有顯式的初始化函數；它在第一次調用運行時函數時初始化（更具體地說，除了參考手冊的錯誤處理和版本管理部分中的函數以外的任何函數）。在計時運行時函數調用和將錯誤代碼從第一次調用解釋到運行時時，需要記住這一點。
在初始化期間，運行時為系統中的每個設備創建一個CUDA上下文（有關CUDA上下文的更多詳細信息，請參閱上下文）。此上下文是此設備的主上下文，它在應用程序的所有主機線程之間共享。作為此上下文創建的一部分，如果需要，設備代碼將及時編譯（請參閱及時編譯）並加載到設備內存中。這一切都是透明的。如果需要，例如對於驅動程序API互操作性，可以從驅動程序API訪問設備的主上下文，如運行時API和驅動程序API互操作性中所述。
當主機線程調用cudaDeviceReset（）時，這會破壞主機線程當前操作的設備（即設備選擇中定義的當前設備）的主上下文。任何將此設備作為當前設備的主機線程進行的下一次運行時函數調用將為此設備創建新的主上下文。
注意：CUDA接口使用全局狀態，全局狀態在主機程序啟動時初始化，在主機程序終止時銷毀。CUDA運行時和驅動程序無法檢測此狀態是否無效，因此在程序啟動或在主程序之后終止期間使用這些接口（隱式或顯式）將導致未定義的行為。
三．設備存儲器
正如在異構編程中所提到的，CUDA編程模型假設一個系統由一個主機和一個設備組成，每個主機和設備都有各自獨立的內存。內核在設備內存中運行，因此運行時提供分配、取消分配和復制設備內存以及在主機內存和設備內存之間傳輸數據的功能。
設備存儲器可以作為線性存儲器或CUDA陣列分配。
CUDA陣列是為紋理提取優化的不透明內存布局。它們在紋理和表面記憶中被描述。 線性存儲器被分配在一個統一的地址空間中，這意味着單獨分配的實體可以通過指針相互引用，例如，在二叉樹或鏈表中。地址空間的大小取決於主機系統（CPU）和所用GPU的計算能力：
注意：在具有計算能力5.3（Maxwell）和更早版本的設備上，CUDA驅動程序創建一個未提交的40位虛擬地址存儲，以確保內存分配（指針）落入支持的范圍。此存儲顯示為存儲虛擬內存，但在程序實際分配內存之前不會占用任何物理內存。
線性內存通常使用cudaMalloc（）分配，使用cudaFree（）釋放，主機內存和設備內存之間的數據傳輸通常使用cudaMemcpy（）完成。在內核的矢量加法代碼示例中，需要將矢量從主機內存復制到設備內存：
// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N) C[i] = A[i] + B[i];
}
// Host code
int main()
{
int N = ...;
size_t size = N * sizeof(float);
// Allocate input vectors h_A and h_B in host memory
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
// Initialize input vectors ... // Allocate vectors in device memory
float* d_A;
cudaMalloc(&d_A, size);
float* d_B;
cudaMalloc(&d_B, size);
float* d_C;
cudaMalloc(&d_C, size); // Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice); // Invoke kernel
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
// Copy result from device memory to host memory // h_C contains the result in host memory
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost); // Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C); // Free host memory
...
}
線性內存也可以通過cudamalocpatch（）和cudamaloc3d（）分配。建議將這些函數用於2D或3D數組的分配，因為它確保適當地填充分配以滿足設備內存訪問中描述的對齊要求，從而確保在訪問行地址或在2D數組和設備內存的其他區域之間執行復制時（使用cudammcpy2d（）和cudammcpy3d（）函數）。返回的力度（或步幅）必須用於訪問數組元素。下面的代碼示例分配一個寬x高的浮點值二維數組，並演示如何在設備代碼中的數組元素上循環：
// Host code
int width = 64, height = 64;
float* devPtr; size_t pitch;
cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height);
MyKernel<<<100, 512>>>(devPtr, pitch, width, height); // Device code
__global__ void MyKernel(float* devPtr, size_t pitch, int width, int height)
{
for (int r = 0; r < height; ++r)
{
float* row = (float*)((char*)devPtr + r * pitch);
for (int c = 0; c < width; ++c)
{
float element = row[c];
}
}
}
下面的代碼示例為浮點值分配一個寬度x高度x深度的三維數組，並演示如何在設備代碼中的數組元素上循環：
// Host code
int width = 64, height = 64, depth = 64;
cudaExtent extent = make_cudaExtent(width * sizeof(float), height, depth);
cudaPitchedPtr devPitchedPtr; cudaMalloc3D(&devPitchedPtr, extent);
MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);
// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr, int width, int height, int depth)
{
char* devPtr = devPitchedPtr.ptr;
size_t pitch = devPitchedPtr.pitch;
size_t slicePitch = pitch * height;
for (int z = 0; z < depth; ++z)
{
char* slice = devPtr + z * slicePitch;
for (int y = 0;y < height; ++y)
{
float* row =(float*)(slice+ y * pitch);
for (int x = 0; x < width; ++x)
{
float element = row[x];
}
}
}
}
參考手冊列出了用於在使用cudaMalloc（）分配的線性內存、使用cudamalocpitch（）或cudamaloc3d（）分配的線性內存、CUDA數組和為全局或恆定內存空間中聲明的變量分配的內存之間復制內存的所有各種函數。
下面的代碼示例演示了通過運行時API訪問全局變量的各種方法：下面的代碼示例分配了一個寬x高x深的浮點值三維數組，並演示了如何在設備代碼中循環數組元素：
__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));
__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));
__device__ float* devPointer; float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));
cudaGetSymbolAddress（）用於檢索指向為全局內存空間中聲明的變量分配的內存的地址。分配的內存大小是通過cudaGetSymbolSize（）獲得的。
四．共享內存
如變量內存空間說明符中所述，共享內存是使用共享內存空間說明符分配的。
共享內存預計將比線程層次結構中提到的和共享內存中詳細描述的全局內存快得多。它可以用作scratchpad內存（或軟件管理的緩存），以最小化來自CUDA塊的全局內存訪問，如下面的矩陣乘法示例所示。
下面的代碼示例是不利用共享內存的矩陣乘法的直接實現。每個線程讀取A的一行和B的一列，並計算C的相應元素，如圖9所示。因此，A從全局內存中讀取B.width times，B從A.height times中讀取。
// Matrices are stored in row-major order: // M(row, col) = *(M.elements + row *M.width + col)
typedef struct
{
int width;
int height;
float*elements;
} 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)
{
// Load A and B to device memory
Matrix d_A;
d_A.width = 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 = 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 = 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, Cd.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)
{
// Each thread computes one element of C // by accumulating results into Cvalue
float Cvalue= 0;
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
for (int e = 0; e < A.width; ++e)
Cvalue += A.elements[row * A.width + e] * B.elements[e * B.width + col];
C.elements[row * C.width + col] = Cvalue;
}
圖9. 無共享內存的矩陣乘法