第一章
1.2 CUDA支持C与C++两种编程语言,该书中的实例采取的是
Thrust数据并行API,.cu作为CUDA源代码文件,其中编译器为ncvv。
1.3 CUDA提供多种API:
- 数据并行C++ Thrust API
- 可用于C或者C++的Runtime API
- 可用于C或者C++的Driver API
1.4 CUDA的基本概念:
- 用于CUDA的GPU是安装于主机系统(Host)的独立设备,主处理器和所有的GPGPU可以同时处理各自的计算任务。
- 传输方式:cudaMemcpu()显示传输,锁定页内存映射的隐式传输(可以实现零拷贝操作),最底层通过设备驱动程序的软件模块进行交互。
- GPGPU运行在一个和主处理器相隔离的存储空间中。这些存储比传统的主机内存(8~20GB/s)有更大的带宽 (160~200GB/s)
- CUDA kernel是可以在主机代码中调用而在CUDA设备上运行的子程序。kernel没有返回值,通过__global__来定义,可以暂时理解为GPU的main函数。
- Kernel调用是异步的,主机仅仅把要执行的kernel顺序提交给GPGPU,并不等待其完成,然后直接处理后面的任务。
- 由于kernel异步执行,为提高效率,可以创建一个kernel组成的流水线,使得GPGPU尽可能长时间保持忙碌;CUDA提供的同步方式:显示调用cudaThreadSynchronize()和cudaMemcpy()。
- GPU上的基本运行单位是线程。每个线程在运行时都好像独占一个处理器,并同时运行于共享内存环境中。一个kernel利用多个线程完成任务称为线程级并行(TLP),有别于处理器指令间的指令级并行(ILP)
- GPU上最大可共享的内存区域称为全局内存,它是遵循整合访问的硬件,通常对连续的128字节数据块进行内存访问。但其访问速度在GPU中最慢,最快的访问是通过寄存器来进行访问。
- CUDA提供了简单的C语言扩展,使得线程可以通过CUDA共享内存空间(shared memory)或通过院子内存操作(atomic)来通信。
- 环境配置不仅定义执行kernel所需的线程数量,还包括grid网格中维度的分配。
一般变量命名中,设备变量常在前面加"d_";blokcIdx,线程块序号;blockDim,线程块的维度线程数量;threadIdx,定位线程在线程块内的编号。
1.6 GPGPU编程的三条法则
- 将数据放入并始终存储于GPGPU:GPU全局内存带宽比主机带宽快20倍。
- 交给GPGPU足够多的任务:防止启动kernel的时间所占比重过大。
- 注重GPGPU上的数据重用,以避免带宽限制。
大O记号是一种表达问题尺寸增长对算法资源小号的影响的常用方式,一些常见的增长率:
- O(1),无论输入集尺寸如何,算法总消耗固定的资源,具有固定的执行时间。如利用GPU索引数组的元素。
- O(n),资源消耗随着输入尺寸增长成线性增长,如循环访问数据集的算法。
- O(n2),平方比例关系,还可能有n的三次方,n的四次方等
- 级别1:向量-向量操作,数据复杂度O(n),计算复杂度O(n);如两个向量的内积。
- 级别2:矩阵-向量操作,数据复杂度O(n2),计算复杂度O(n2);矩阵与向量乘积。
- 级别3:矩阵-向量操作,数据复杂度O(n2),计算复杂度O(n3);密集矩阵乘法。
1.8 CUDA和Amdahl定律
一个程序中可并行话部分为P(百分比),分配在n个处理器上,则并行化后理想执行时间与程序的关系为
S(n) = 1 / (( 1 - P ) + P / n)
因此在并行的时候应该有两种理念:优化并行部分代码的表达式,尽可能减少串行时间。
1.9 数据并行与任务并行
之前的讨论都是数据并行,而任务并行是指给CPU和GPU分配不同的任务
第二章 CUDA在机器学习与优化中的应用
2.1建模与模拟
数学模型是对现实的抽象,用来对现实进行分析与预测;数值模拟是通过应用程序将数学模型映射到电脑上。通常建模的方法有两种:
- 基于第一性原理分析及其他方法的人工推导模型。
- 基于数据拟合的参数化模型。(如神经网络)
Nelder-Mead方法:一种直接搜索型非线性优化方法,是利用单纯形(N维空间的广义三角形)进行搜索。
Levenberg-Marquardt方法:可信区域算法,用来在参数空间中寻找函数的最佳拟合。
加速算法,优化技术(如共轭梯度法),通过目标函数的导数来加速查找极小值的过程。(GiNaC,C++库,用以进行微分等数学运算)
2.2 机器学习与神经网络
人工神经网络(ANN)是一种基于观察到的数据推测函数的机器学习技术,学习方法有监督学习,非监督学习和随机学习等。
2.3异或逻辑:一个重要的非线性机器学习问题
单层网络无法表达异或逻辑关系,而加入隐藏层可以大大提高运算能力
第四章 CUDA执行模型
4.1 GPU架构综述
GPU
硬件上的大规模并行是通过重复这是多个相同的通用构件块(流多处理器,SM)实现的;线程块是协作线程间合作的容器,只有线程块内部的线程才能
共享数据;每个SM可以被调度执行一个或多个线程块,这中映射:
- 可以透明得扩展到任意数量的SM上
- 对SM的位置没有限制
- 能够将执行中的Kernel与用户参数广播给硬件。*并行广播是扩展性最强、速度最快的通信机制。
把工作分配给SM是千兆线程全局调度器的职责,为每个SM分配线程块的数量取决于一个SM中最多能常驻的线程与线程块的数量。32个线程可以看做一个warp,每个时钟周期中,SM的warp调度器执行哪个warp将从符合下列两项的warp中选择:
- 当前不需要等待数据从设备内存中传来。
- 当前不需要等待前一条指令的完成。
4.1.2 warp 分歧:GPGPU设备被归为单指令多线程(SIMT)设备,warp中个条件子句分支之间是串行执行,因此if语句会降低SM的执行性能。
*warp分歧的避免准则:
- 使用不同的算法重新规划问题
- 将不同计算代价的任务分别列表,每个列表使用一个kernel函数
- 将计算任务排序并分块,块的大小为半个warp的整数倍
- 利用异步kernel执行方式
- 使用主处理器来执行分任务
4.2 Warp调度与TLP(thread level parallel)
*在每个SM上执行多个warp是GPU隐藏ALU与内存延迟、保持执行单元繁忙的唯一方式。
TL
P的思想是给调度器尽可能多的线程以供选择,进而使性能损失的可能性最小。占用率是TLP的一个量度,多处理上并行执行的warp数量除以最多能容纳的warp数量,高占用率意味着SM的调度器有很多warp可供选择,因此隐藏了ALU与数据延迟。
4.3 ILP:高性能低占用率
ILP(instruction level parallel)指令级并行,可以是更少的线程保持SIMD核心忙碌,拥有更多的寄存器,ILP仅通过增多每个线程独立的命令数量就可以提升性能,最佳的性能是在SM上常驻线程数量为576时达到。其中在数据延迟中,通过一系列复杂的设计也可以使得延迟隐藏(P84)。
*一些结论:密集线性代数学方法更适合GPU架构,而不是传统的多核架构;为了充分利用总线带宽,每个SM提交保持30~50个128字节的储存操作。
4.4 Little 法则
从TLP角度来看,运行时储存操作数量N是到达率a(这是期望的指令速率与加载指令密度的成绩)与内存延迟L的乘机(N=aL);从ILP角度来看,独立的访存事件能够被成批处理(N=aL-B),其中B是合并的独立加载指令的数量。
*提高并行性的一些考虑:
- 提升占用率
- 最大化可用的寄存器(让编译器使用__launch__bounds__给每个kernel函数分配额外的寄存器)
- 调整线程块的维度以最好的利用SM Warp调度器
- 修改代码以利用ILP是每个线程处理多个元素
- 不要把工作都塞到一个kernel中:
-
- 尝试创建小线程块,是操作密度均匀
- 不要讲相同类型的操作都聚集到kernel函数的一部分中
- 不要在一个功能单元上出现瓶颈
4.5 检测限制因素的CUDA工具。
第五章 CUDA储存器
5.1 CUDA储存器层次结构
主机和所有设备都拥有各自的独立内存空间,CUDA4.0引入了统一虚拟编址(UVA)来简化GPU编程,开发人员可以直接通过cudaMemcpy()函数实现不同GPU之间的数据传输,但在使用UVA之前要通过cudaHostRegister注册内存区域。最后使用cudaHostUnregister()函数就可以在指定的内存区域终止使用UVA进行数据传输。
5.2 GPU存储器
SM上的存储器是速度最快的,但通常容量只有KB级;全局内存是共享的储存系统,容量可达GB级,但速度最慢。
要实现数据重用,只能通过挖掘数据的局部性,GPGPU支持两种局部性:
- 时间局部性:假设当前访问的数据在不远的将来可能会被再次访问,如满足LRU算法
- 空间局部性:缓存邻近的数据,如渲染操作。
5.3 L2缓存
在2.0以上的设备中,所有数据加载与存储操作都经过L2缓存(包括CPU与GPU之间的数据拷贝)
5.4 L1缓存
2.0的设备具有64KB的L1缓存,说明如下:
- L1缓存是基于空间重用而非时间重用设计
- L1缓存并不会影响全局内存的写操作,这些操作会越过L1缓存。
- L1缓存并不保证一致性
- L1缓存有10~20时钟周期的延迟
- 具有良好的可配置性,可以将其设置为动态缓存和共享缓存
*局域内存是一种用来存放寄存器溢出的局部数据的内存空间,将存放在L1缓存中。
*编译器可以根据代码生成统一加载(LDU)指令访问L1,可是现在SM中高效得向线程广播(broadcast)数据,但需要指针使用const修饰,线程块内所有线程君访问同一地址,如:
__global void kernel (float *g_dst,
const float *g_src)
{
g_dst =
g_src[0] + gsrc[blockIdx.x];
}
5.5 CUDA 内存类型
5.5.1 寄存器
每个SM支持2^15的32位寄存器,每个Kernel可以使用的寄存器数量为63个。
5.5.2 局域内存
对应的是对自动变量的操作。自动变量是在设备端申请的,不包括__device__,__shared__,__constant__等限定符,通常其会放到寄存器中,但通过常亮索引访问数组、体积过大时会放到局域内存中。
5.5.3 共享内存
每个SM可使用的共享内存(有时也被称为smem)为16K或48K,字宽为32位,有以下三种方式进行申请:
- 在kernel中以静态的方式或在文件中以全局变量进行声明。如在kernel函数中调用带有__shared__标识的函数。
- 通过Drived API的函数 cuFuncSetSharedSize在kernel中动态声明。
- 通过执行配置动态声明。
*SM上的共享内存被组织成一些32位的
Bank,不同的线程请求相同Bank会产生Bank冲突,引起性能下降,解决方法是填充数据消耗内存来避免性能的下降。
共享内存还具有广播能力,所一个warp中多个线程访问同一个字,则硬件上秩序一次共享内存读取操作便可以改变多个线程。
5.5.4 常亮内存
1.x设备中,常量内存只有64K,且目的是用于将所有设备广播只读取数据的理想方式;2.x设备中,开放人员访问全局变量时会达到访问常量内存同样的效果。数据必须满足:
- 数据存放与全局内存
- 数据在kernel中是只读的(可以使用const来限制)
- 数据访问与线程ID无关
5.5.5 纹理内存
其通常可以绑定全局内存的数据,并具有一定的缓存功能,主要特征如下:
其通常可以绑定全局内存的数据,并具有一定的缓存功能,主要特征如下:
- 通常用于可视化处理
- 基于2D空间局部性的纹理内存可对缓存进行优化
- 每个SM有8K的缓存空间
- 可以高效的拆解和广播数据
- 具有9位数据计算单元,可以进行越界数据访问处理、插值以及整型到浮点型的转换。
**当仅作为缓存时,cudaMallco()申请的内存较小,不至于浪费;当用作数据处理的时候,需要绑定cudaMallocPitch()申请的线性对齐内存。
从纹理内存中读取数据最简单的形式是使用tex1Dfetch();另外在使用tex1D(),tex2D(),tex3D()处理时,要根据一些可变属性和不可变属性来进行不同纹理坐标的插值处理。Page107
5.6全局内存
内存整合:若干线程的内存请求在硬件上被整合或者合并,成为一次多数据的内存数据传输。若要获得高性能,需要注意:
- 内存地址是对齐的,要保证合适的线程块尺寸(16的倍数),在定义结构体的时候使用限定符__align__(8), __align__(16)
- warp内各个线程所访问的数据在一个连续的区域。
全局内存的数据的读取方式分为缓存读取(L1搜索->L2搜索->128字节的缓存行读取)和非缓存读取(跳过L1搜索,L2搜索->32字节的全局内存搜索)。
全局内存的申请:
静态申请:利用cudaMalloc(),申请到的空间是256字节对齐的;动态申请:在CUDA线程中使用malloc,申请到的空间是16字节对齐的。
第六章 高效使用CUDA存储 器
6.1 归约
归于运算时求一个向量的最小值、最大值或者加和等操作的总称。
6.2 使用非规则数据结构
6.3稀疏矩阵和CUSP支持库
CUSP库是一个基于Thurst API实现的用于GPU上完成稀疏矩阵和图论计算的项目。
6.4 图论算法
6.5 AoS(Array of Structures)、SoA(Structure of Arrays)以及其他数据结构
struct S{float x; float y;};
struct S myDatas[N];
相比于上面,SoA避免了内存整合问题
struct S{float x[N]; float y[N];};
struct S myDatas;
6.6 分片和分块
分片和分块是用于创建多为分解网络的一种抽象,这些抽象可帮编程者将数据访问组织为通用性是,并定义用于线程块内各线程之间通信的共享内存。
**Thrust API将CUDA的复杂性封装入一个简单的API接口。
第七章 提高并行度的技巧
设计CUDA:充分开发GPU内部大规模并行,通过并发执行的流来实现多GPU使用,异步数据传输以及单设备内不同的kernel函数的并行。
7.1 CUDA上下文环境对并行度的扩展
CUDA在第一次调用一个改变驱动状态的函数时默认创建一个上下文环境,例如cudaMalloc()就是改变上下文的一个调用。而同一时刻只能有一个上下文处于活动状态。一般来说,上下文环境通常在GPU 0上创建,但可以通过cudaSetCevice()来选择另一个GPU。
7.2 流与上下文环境
CUDA 应用程序通过将操作排队到流中来管理任务并行,kernel函数会在当前设备相关的流中列入一个kernel调用操作。
kernel<<<nBlocks, nThreadsPerBlocks, 0, stream[i]>>>(parameters)
队列充当一个先入先出缓冲区,在设备井并发执行或者在单个设备上并发执行多个kernel函数时,需要使用多个流。
7.2.1 多GPU使用
在多个设备上创建上下文环境:
cudaGetDeviceCount(&nGPU);
int * d_A[nGPU];
for(int i = 0; i < nGPU; ++i)
{
cudaSetDevice(i);
cudaMalloc(&d_A[i], n*sizeof(int));
}
除非在API调用或者kernel函数中显示地定义一个流,否则GPU操作都会在cudaSetDevice()之后隐式地在与设备相关联的流中进行排列。
7.2.2 显示同步
事件是在流内创建占位标志的一种方法。其中0号流内的事件在所有流中所有任务都完成后才能完成。
一些显示同步流与事件的运行时的方法课参考P140
7.2.3 隐式同步
以下主机操作要格外注意,它们有可能停止所有并发操作并给程序性能带来负面影响
- 锁存页主机内存的分配
- 设备内存分配
- 设备内存设置
- 设备到设备的内存拷贝
- L1缓存与共享内存之间的配置转换
7.2.4 统一虚拟地址空间(64位)
7.2.5 一个简单示例
在每个GPU上分配n个证书空间,并行的向每个GPU中填充向量,从而组成一个连续整数的大向量,GPU内存向主机内存异步传输数据。代码如下:
#include <stdio.h>
__global__ void fillKernel(int * a, int n, int offset)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid < n)
{
for(int i = 0; i < 100; ++i)
{
a[tid] = offset + tid;
}
}
}
int main(int argc, char * agrv[])
{
int nGPU;
int n = 1000000;
int size = n * sizeof(int);
cudaGetDeviceCount(&nGPG);
int *d_A[nGPU];
for(int i = 0; i < nGPU; i++)
{
cudaSetDevice(i);
cudaMalloc(&d_A[i], size);
}
int * h_A;
cudaHostAlloc(&h_A, nGPU * sizeof(int), cudaHostAllocPortable);
for(int i = 0; i < nGPU; i++)
{
int nThreadsPerBlock = 512;
int nBlocks = n / nThreadsPerBlock + ((n % nThreadsPerBlock) ? 1 : 0);
cudaSetDevice(i);
fillKernel<<<nBlocks, nThreadsPerBlock>>>(d_A[i], n, i * n);
cudaMemcpyAsync(&h_A[i * n], d_A[i], size, cudaMemcpyDeviceToHost);
}
cudaDeviceSynchronize();
//检查正确性
for(int i = 0; i < nGPU * n; i++)
{
if(h_A[i] != i)
{
printf("Error h_A[%d] = %d\n", i, h_A[i]);
exit(1);
}
}
printf("Success.\n");
cudaFreeHost(h_A);
for(int i = 0; i < nGPU; i++)
{
cudaSetDevice(i);
cudaFree(d_A[i]);
}
return 0 ;
}
__global__ void fillKernel(int * a, int n, int offset)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid < n)
{
for(int i = 0; i < 100; ++i)
{
a[tid] = offset + tid;
}
}
}
int main(int argc, char * agrv[])
{
int nGPU;
int n = 1000000;
int size = n * sizeof(int);
cudaGetDeviceCount(&nGPG);
int *d_A[nGPU];
for(int i = 0; i < nGPU; i++)
{
cudaSetDevice(i);
cudaMalloc(&d_A[i], size);
}
int * h_A;
cudaHostAlloc(&h_A, nGPU * sizeof(int), cudaHostAllocPortable);
for(int i = 0; i < nGPU; i++)
{
int nThreadsPerBlock = 512;
int nBlocks = n / nThreadsPerBlock + ((n % nThreadsPerBlock) ? 1 : 0);
cudaSetDevice(i);
fillKernel<<<nBlocks, nThreadsPerBlock>>>(d_A[i], n, i * n);
cudaMemcpyAsync(&h_A[i * n], d_A[i], size, cudaMemcpyDeviceToHost);
}
cudaDeviceSynchronize();
//检查正确性
for(int i = 0; i < nGPU * n; i++)
{
if(h_A[i] != i)
{
printf("Error h_A[%d] = %d\n", i, h_A[i]);
exit(1);
}
}
printf("Success.\n");
cudaFreeHost(h_A);
for(int i = 0; i < nGPU; i++)
{
cudaSetDevice(i);
cudaFree(d_A[i]);
}
return 0 ;
}
7.3 使用多个流乱序执行
*在同一GPU内并发执行kernel函数的建议:
- 所有独立操作应该在依赖操作之前执行
- 任何同步应该尽量延后
考虑在一个CUDA流内完全归约至一个值,可以选择:写一个独立的kernel并在functionReduce()之后执行;或者使用原子操作在kernel内部进行同步操作。
7.4 将数据捆绑计算
CUDA编程者可以选择以下方式来使数据用于多个GPU上执行的kernel:
- 将主机内存映射到所有GPU设备的内存空间。
- 手动分配空间并传递数据。
- 对于将在分布式MPI(Message Passing Interface)环境中运行的程序,在各个设备上分配空间,并通过MIP发送与接收调用将数据传给GPU。
其中手动分割数据是最灵活,最具扩展性,性能最高的数据捆绑计算方法;内存映射将提供便利性。
将内存映射到系统中的设备方便使用:
- 无需分割数据
- 无需在设备内存中分配空间或者手动拷贝数据
- 无需使用流将数据传输与kernel执行进行重叠
但在使用内存映射的时候要注意:主机内存页的
对齐(使用cudaAllocHost())以及在映射内存更改后,程序必须使用流或事件同步内存访问。
*映射内存的重要特性:它允许多个设备正确地更新连续的、无重叠的内存设置。
第八章 CUDA在所有GPU与CPU程序中的应用
8.3 支持库
8.3.1 CUBLAS:
基本线性代数子程序(BLAS),CUDA通过自身的CUBLAS在支持这个库。
其模型是:在GPU内存空间中创建矩阵与向量实例,用数据填充这些实例,调用一些列CUBLAS函数,将结果返回给主机。
8.3.2 CUFFT:
提供一个基于GPU的快速Fourier变换的实现,支持double类型(需要在nvcc命令行中指定-D REAL = double)
8.3.3 MAGMA
8.3.4 phiGEMM支持库
能在具有多个GPU与多核处理器的异构系统中进行矩阵间乘法运算,其扩展了Fatica映射,用来支持单精度、双精度和复数矩阵。
8.3.5 CURAND
生成高质量的伪随机数或者拟随机数
第九章 CUDA与图形渲染混合编程
第十章 在云计算和集群环境中使用CUDA
第十一章 CUDA在现实问题中的应用