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 memory
Matrix 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 column
int blockRow = blockIdx.y;
int blockCol = blockIdx.x;
// Each thread block computes one sub-matrix Csub of C
Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
// Each thread computes one element of Csub
// by accumulating results into Cvalue
float Cvalue = 0;
// Thread row and column within Csub
int 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 results
for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {
// Get sub-matrix Asub of A
Matrix Asub = GetSubMatrix(A, blockRow, m);
// Get sub-matrix Bsub of B
Matrix 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-matrix
As[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 together
for (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 element
SetElement(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字节的写入。