CUDA运行时 Runtime（一）
CUDA運行時 Runtime（一）
一. 概述
運行時在cudart庫中實現(xiàn)，該庫通過靜態(tài)方式鏈接到應用程序庫cudart.lib和 libcudart.a，或動態(tài)通過cudart.dll或者libcudart.so. 需要cudart.dll和/或libcudart。索對于動態(tài)鏈接，通常將它們作為應用程序安裝包的一部分包括在內(nèi)。只有在鏈接到CUDA運行時的同一實例的組件之間傳遞CUDA運行時符號的地址才是安全的。
它的所有入口點都以cuda為前綴。
正如在異構編程中所提到的，CUDA編程模型假設一個系統(tǒng)由一個主機和一個設備組成，每個主機和設備都有各自獨立的內(nèi)存。設備內(nèi)存概述了用于管理設備內(nèi)存的運行時函數(shù)。
共享內(nèi)存說明了如何使用線程層次結(jié)構中引入的共享內(nèi)存來最大限度地提高性能。
頁面鎖定主機內(nèi)存引入了頁面鎖定主機內(nèi)存，它需要將內(nèi)核執(zhí)行與主機和設備內(nèi)存之間的數(shù)據(jù)傳輸重疊起來。
異步并發(fā)執(zhí)行描述了用于在系統(tǒng)的各個級別上啟用異步并發(fā)執(zhí)行的概念和API。
多設備系統(tǒng)顯示了編程模型如何擴展到多個設備連接到同一主機的系統(tǒng)。
錯誤檢查描述如何正確檢查運行時生成的錯誤。
調(diào)用堆棧提到用于管理CUDA C++調(diào)用堆棧的運行時函數(shù)。
紋理和表面存儲器提供了紋理和表面存儲器空間，它們提供了訪問設備存儲器的另一種方式；它們還公開了GPU紋理硬件的一個子集。
圖形互操作性引入了運行時提供的與兩個主要圖形api（OpenGL和Direct3D）互操作的各種功能。
二．初始化
運行時沒有顯式的初始化函數(shù)；它在第一次調(diào)用運行時函數(shù)時初始化（更具體地說，除了參考手冊的錯誤處理和版本管理部分中的函數(shù)以外的任何函數(shù)）。在計時運行時函數(shù)調(diào)用和將錯誤代碼從第一次調(diào)用解釋到運行時時，需要記住這一點。
在初始化期間，運行時為系統(tǒng)中的每個設備創(chuàng)建一個CUDA上下文（有關CUDA上下文的更多詳細信息，請參閱上下文）。此上下文是此設備的主上下文，它在應用程序的所有主機線程之間共享。作為此上下文創(chuàng)建的一部分，如果需要，設備代碼將及時編譯（請參閱及時編譯）并加載到設備內(nèi)存中。這一切都是透明的。如果需要，例如對于驅(qū)動程序API互操作性，可以從驅(qū)動程序API訪問設備的主上下文，如運行時API和驅(qū)動程序API互操作性中所述。
當主機線程調(diào)用cudaDeviceReset（）時，這會破壞主機線程當前操作的設備（即設備選擇中定義的當前設備）的主上下文。任何將此設備作為當前設備的主機線程進行的下一次運行時函數(shù)調(diào)用將為此設備創(chuàng)建新的主上下文。
注意：CUDA接口使用全局狀態(tài)，全局狀態(tài)在主機程序啟動時初始化，在主機程序終止時銷毀。CUDA運行時和驅(qū)動程序無法檢測此狀態(tài)是否無效，因此在程序啟動或在主程序之后終止期間使用這些接口（隱式或顯式）將導致未定義的行為。
三．設備存儲器
正如在異構編程中所提到的，CUDA編程模型假設一個系統(tǒng)由一個主機和一個設備組成，每個主機和設備都有各自獨立的內(nèi)存。內(nèi)核在設備內(nèi)存中運行，因此運行時提供分配、取消分配和復制設備內(nèi)存以及在主機內(nèi)存和設備內(nèi)存之間傳輸數(shù)據(jù)的功能。
設備存儲器可以作為線性存儲器或CUDA陣列分配。
CUDA陣列是為紋理提取優(yōu)化的不透明內(nèi)存布局。它們在紋理和表面記憶中被描述。
線性存儲器被分配在一個統(tǒng)一的地址空間中，這意味著單獨分配的實體可以通過指針相互引用，例如，在二叉樹或鏈表中。地址空間的大小取決于主機系統(tǒng)（CPU）和所用GPU的計算能力：
注意：在具有計算能力5.3（Maxwell）和更早版本的設備上，CUDA驅(qū)動程序創(chuàng)建一個未提交的40位虛擬地址存儲，以確保內(nèi)存分配（指針）落入支持的范圍。此存儲顯示為存儲虛擬內(nèi)存，但在程序?qū)嶋H分配內(nèi)存之前不會占用任何物理內(nèi)存。
線性內(nèi)存通常使用cudaMalloc（）分配，使用cudaFree（）釋放，主機內(nèi)存和設備內(nèi)存之間的數(shù)據(jù)傳輸通常使用cudaMemcpy（）完成。在內(nèi)核的矢量加法代碼示例中，需要將矢量從主機內(nèi)存復制到設備內(nèi)存：
// 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
…
}
線性內(nèi)存也可以通過cudamalocpatch（）和cudamaloc3d（）分配。建議將這些函數(shù)用于2D或3D數(shù)組的分配，因為它確保適當?shù)靥畛浞峙湟詽M足設備內(nèi)存訪問中描述的對齊要求，從而確保在訪問行地址或在2D數(shù)組和設備內(nèi)存的其他區(qū)域之間執(zhí)行復制時（使用cudammcpy2d（）和cudammcpy3d（）函數(shù)）。返回的力度（或步幅）必須用于訪問數(shù)組元素。下面的代碼示例分配一個寬x高的浮點值二維數(shù)組，并演示如何在設備代碼中的數(shù)組元素上循環(huán)：
// 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深度的三維數(shù)組，并演示如何在設備代碼中的數(shù)組元素上循環(huán)：
// 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（）分配的線性內(nèi)存、使用cudamalocpitch（）或cudamaloc3d（）分配的線性內(nèi)存、CUDA數(shù)組和為全局或恒定內(nèi)存空間中聲明的變量分配的內(nèi)存之間復制內(nèi)存的所有各種函數(shù)。
下面的代碼示例演示了通過運行時API訪問全局變量的各種方法：下面的代碼示例分配了一個寬x高x深的浮點值三維數(shù)組，并演示了如何在設備代碼中循環(huán)數(shù)組元素：
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（）用于檢索指向為全局內(nèi)存空間中聲明的變量分配的內(nèi)存的地址。分配的內(nèi)存大小是通過cudaGetSymbolSize（）獲得的。
四．共享內(nèi)存
如變量內(nèi)存空間說明符中所述，共享內(nèi)存是使用共享內(nèi)存空間說明符分配的。
共享內(nèi)存預計將比線程層次結(jié)構中提到的和共享內(nèi)存中詳細描述的全局內(nèi)存快得多。它可以用作scratchpad內(nèi)存（或軟件管理的緩存），以最小化來自CUDA塊的全局內(nèi)存訪問，如下面的矩陣乘法示例所示。
下面的代碼示例是不利用共享內(nèi)存的矩陣乘法的直接實現(xiàn)。每個線程讀取A的一行和B的一列，并計算C的相應元素，如圖9所示。因此，A從全局內(nèi)存中讀取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. 無共享內(nèi)存的矩陣乘法
總結(jié)
以上是生活随笔為你收集整理的CUDA运行时 Runtime（一）的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： CUDA C++编程接口：编译
- 下一篇： CUDA运行时 Runtime（二）