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. 无共享内存的矩阵乘法