本篇为学习笔记,学习内容为2019年参加英伟达GTC会议的课程
需要提下学习CUDA的目的,就是为了加速自己的应用,相比于CPU-only的应用程序,可以用GPU实现较大加速,当然程序首先是计算密集型而非IO密集型
GPU加速系统,又被称异构系统(Heterogeneous),由CPU和GPU组成
如果熟悉C编程,可以很快上手CUDA编程,两者在代码形式上有很多类似地方,一个比较重要概念是GPU的launch kernel
C代码用gcc编译,cuda代码用nvcc编译,nvcc内部会调用gcc
启动核函数的配置 <<>> thread是最小执行单位,由threads组成block,多个block组成grid;kernel只能运行在一个grid
一般最简单的加速示例就是一个CPU的循环,执行简单的算术运算;主要是暗示我们什么类型的程序适合GPU加速
关于threads:
cuda6之后的版本可以分配出CPU/GPU都能访问的内存,API接口为:cudaMallocManaged
关于异常处理:
迭代设计过程:
APOD:Assess Parallelize Optimize Deploy
评估->并行->优化->部署
使用Nsight命令行工具nsys来评估性能,确定优化机会
nsys基本用法: nsys profile —stats=true exe
它会生成qdrep报告,包含诸多信息:API统计,kernel执行统计,内存拷贝的大小和时间等
程序优化方法之一:更改kernel launch的配置参数
Streaming Multiprocessor 流处理单元SM
为了获取SM的数量,调用API:
int deviceId;
cudaGetDevice(&deviceId);
cudaDeviceProp props;
cudaGetDeviceProperties(&props, deviceId);
SMs = props.multiProcessorCount;
基础知识:CUDA’s Unified Memory
Asynchronous Memory Prefetching 异步内存预取:减小页错误和按需内存迁移的间接开销的技术,提高性能
cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId); // Prefetch to GPU device.
cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId); // Prefetch to host. `cudaCpuDeviceId`
Nsight Systems 可视化的性能分析工具,其可以直接打开nsys生成的qdrep文件
Concurrent CUDA Streams 并发流;流是一系列顺序执行的命令,kernel的执行,和许多内存迁移都是发生在流内,不指定的情况下使用default stream
关于控制流的几个规则:
API:
cudaStream_t stream; // CUDA streams are of type `cudaStream_t`.
cudaStreamCreate(&stream); // Note that a pointer must be passed to `cudaCreateStream`.
someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>(); // `stream` is passed as 4th EC argument.
cudaStreamDestroy(stream); // Note that a value, not a pointer, is passed to `cudaDestroyStream`.
第三个参数是每个block允许使用的shared memory的bytes,默认为0
profile driven and iterative 配置文件驱动和迭代
当确定数据只在device使用,最好只分配device的内存,减小数据迁移的开销,API:
Using Streams to Overlap data transfers and code execution
只要CPU内存是锁页内存,就可以使用cudaMemcpyAsync()来进行异步拷贝,另一个条件就是使用非默认流
默认情况下GPU函数执行时对CPU函数是异步的,而异步拷贝,不仅对CPU,对GPU的kernel也是异步的,可以达到边计算边拷贝数据的目的,从而掩盖数据传输时间,尽量挖掘GPU计算能力