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

在 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:

fzf

zsh : 在.zshrc 中加入 plugins=(fzf)

fish

zsh

vim-plug