CUDA运行时 Runtime（二）

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（二）相关推荐

CUDA运行时 Runtime（四）
CUDA运行时 Runtime(四) 一. 图图为CUDA中的工作提交提供了一种新的模型.图是一系列操作,如内核启动,由依赖项连接,依赖项与执行分开定义.这允许定义一次图形,然后重复启动.将图的定义 ...
CUDA运行时Runtime（三）
CUDA运行时Runtime(三) 一．异步并发执行 CUDA将以下操作公开为可以彼此并发操作的独立任务: 主机计算: 设备计算: 从主机到设备的内存传输: 从设备到主机的存储器传输: 在给定设备的存 ...
CUDA运行时 Runtime（一）
CUDA运行时 Runtime(一) 一. 概述运行时在cudart库中实现,该库通过静态方式链接到应用程序库cudart.lib和 libcudart.a,或动态通过cudart.dll或者lib ...
Deep Learning部署TVM Golang运行时Runtime
Deep Learning部署TVM Golang运行时Runtime 介绍 TVM是一个开放式深度学习编译器堆栈,用于编译从不同框架到CPU,GPU或专用加速器的各种深度学习模型.TVM支持来自Te ...
“ compiler-rt”运行时runtime库
" compiler-rt"运行时runtime库编译器-rt项目包括: • Builtins-一个简单的库,提供了代码生成和其他运行时runtime组件所需的特定于目标的低级接 ...
iOS运行时Runtime浅析
运行时是iOS中一个很重要的概念,iOS运行过程中都会被转化为runtime的C代码执行.例如[target doSomething];会被转化成objc)msgSend(target,@select ...
深入浅出Substrate：剖析运行时Runtime
基于Substrate开发自己的运行时模块,会遇到一个比较大的挑战,就是理解Substrate运行时(Runtime).本文首先介绍了Runtime的架构,类型,常用宏,并结合一个实际的演示项目,做了 ...
ios runtime重要性_iOS运行时RunTime详解
Objective-C语言是扩展于C语言的一种面向对象的编程语言,然而其方法的调用方式又和大多数面向对象语言大有不同,其采用消息传递.转发的方式进行方法的调用.因此在OC中,对象的真正行为往往发生在运 ...
HotSpot VM运行时01---命令行选项解析
HotSpot VM有3个主要组件:VM运行时(Runtime).JIT编译器(JIT Compiler)以及内存管理器(Memory Manager). HotSpot VM运行时担当许多职责:命令 ...

CUDA运行时 Runtime（二）

CUDA运行时 Runtime（二）相关推荐

最新文章

热门文章