leenldk 发布的文章

compiler

函数 getGroup:规划出一个 group,返回一个 GateGroup

schedule

GateGroup:一组门
1. relatedQubits
2. state
3. cuttPlans

compile

backend 选项:group, mix, blas

cublas

首先创建 cublas handle

#include <cublas_v2.h>
#define checkCudaErrors(status) do {                                   \
    std::stringstream _error;                                          \
    if (status != 0) {                                                 \
      _error << "Cuda failure: " << status;                            \
      FatalError(_error.str());                                        \
    }                                                                  \
} while(0)


cublasHandle_t cublasH;
checkCudaErrors(cublasCreate(&cublasH));
// 之后的 library function call 显式传入 handle
cublasDestroy(cublasH);

curand

#include <curand.h>
curandGenerator_t curand;
curandCreateGenerator(&curand, CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(curand, 123ULL);
curandGenerateUniform(curand, p, size);

NCCL

nvidia 集合通信库。多 GPU 多节点通信原语。
支持 all-reduce, all-gather 等

vtune

intel profiler

source /home/leenldk/intel/oneapi/vtune/2021.2.0/env/vars.sh  #加载

gprof

gcc 开源 profile 工具

编译时添加 -pg 选项进行插装
运行后生成 gmon.out
通过 gprof 输出 profiling 文件

gcc example.c -o temp -g -pg
./temp
gprof temp > profiling.out

nvprof

update : nvprof 已经不再支持最新 GPU,请使用 nsys 和 ncu

cuda toolkit 中自带工具
使用:

nvprof ./gemm # 输出 prof 结果
# 在使用了 unified memory 时可能需要 添加 --unified-memory-profiling off
nvprof --unified-memory-profiling off ./gemm

-o prof.nvvp : 输出为 nvvp 文件
--metrics [all/gld_throughput] : profile 所有参数/Global Load Throughput (可能需要 sudo)

可视化:使用 x11 forwarding nvvp prof.out
cuda 11 版本可能有 java 问题,此时需要
sudo apt install openjdk-8-jdk
nvvp -vm /usr/lib/jvm/java-8-openjdk-amd64/jre/bin/java prof.out

windows :
.\nvvp.exe -vm 'D:\Program Files\Java\jdk1.8.0_311\jre\bin\java.exe'

nsys (nsight system)

粗粒度 timeline profile

ncu (nsight compute)

细粒度单个 kernel 级别 profile
ncu --list-sets 获取支持的 metric section set

--set full
-o file

1.14

为了防止自己在康paper时睡着来记一点笔记
目前在重新(?)看TACO的文章

TACO 支持在 CPU 上生成稀疏和稠密 tensor 表达式的代码
文章做的几点事情:
1. tensor 存储方法
2. iteration graph
3. merge lattices
4. 代码生成算法

由用户指定 merge 几个简单的 kernel 或者当成一个整体计算
稀疏kernel计算复杂原因:
1. sparse 数据结构维度的不同排布方式
2. sparse 下标合并

Gustafson's law:
任务总工作量为 W ,其中 p 比例的工作量可以通过并行加速,1-p 比例无法通过并行加速
当并行加速的加速比为 s 时,理论工作量为:
$$ W(s) = (1-p)W + spW $$

upd: 这个东西似乎应该这么理解:
Amdahl's law:
设串行需要时间 t, 其中 f 比例可以并行,则 N 线程加速比为:
$$S = \frac{t}{(1-f)t + \frac{f}{N} t} = \frac{1}{(1-f) + \frac{f}{N}}$$

Gustafson's law:
设 N 线程并行需要时间 ts, 其中 f 比例以并行执行,则相比单线程加速比为:
$$S = \frac{(1-f)t_s + N f t_s}{t_s} = (1-f) + N f$$

二者区别应该在于 Gustafson 从并行角度、(强调并行可以提高可解决问题规模,对应弱扩展性),而 Amdahl 从串行角度(强调并行瓶颈在于串行部分,对应强扩展性)

1.15

GPU中所有 active thread 被分配了单独的寄存器,当切换线程时不需要交换寄存器

cuda 中所有 kernel launch 都为异步
在开始和结束 CPU timer 之前都需要调用 cudaDeviceSynchronize

1.16

GPU 理论带宽计算:
V100 使用 HBM2 (double data rate) RAM, 时钟 877MHz, 4096位内存接口
理论带宽为:$$(0.877 \times 10^9 \times (4096 / 8) \times 2)\div 10^9 = 898GB/s$$

cudaMallocManaged( void** devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal ) :申请 unified memory,可以从 device 和 host 上访问

Requested Global Load/Store Throughput :
kernel 需要的 gobal memory throughput,对应于等效带宽
Global Load/Store Throughput: 最小内存传输块较大,实际传输量可能超过 kernel 的需求量,记为 Global Load/Store Throughput

GPU architecture

在 V100 中 device memory 到 GPU 理论峰值带宽为 898GB/s
host memory 到 device memory 峰值带宽为 16GB/s

page locked (pinned) memory 可以获得较高host to device 带宽,可达约 12GB/s
使用 cudaHostAlloc() 进行分配
使用 cudaHostRegister() pin已经分配的内存
pinned memory 不能过量使用,分配是 heavyweight 操作

cudaMemcpy() 是阻塞操作
cudaMemcpyAsync() 非阻塞,需要 pinned host memory,需要指定 stream ID,可以与 host 的 cpu function overlap,但不能与同 stream 的 kernel overlap
两个不同的非 default stream 可以 overlap

zero copy : 需要 mapped pinned (non-pageable) memory

1.18

CUDA memory space: