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

MPI profile

在单节点 profile 中，nsys 可以在 mpirun 之前：
nsys profile [nsys args] mpirun [mpirun args] ...

在多节点 profile 中，nsys 必须在 mpirun 之后：
mpirun [mpirun args] nsys profile [nsys args] ...

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

proxy

proxy server :
https://github.com/tinyproxy/tinyproxy
默认端口 8888，后台运行

ssh tunneling

local port forwarding :
本地命令：
ssh -L 8181:192.168.0.135:3389 pi@192.168.0.135
此时连接本地 localhost:8181 相当于连接 192.168.0.135:3389 且只有 localhost:22 的连接通过防火墙

dynamic port forwarding :
本地命令：
ssh -D 8181 pi@192.168.0.135
设置本地代理为 localhost:8181，把本地所有流量传到远端

remote port forwarding :
本地命令：
ssh -R 8181:localhost:3389 pi@192.168.0.135
将本地 3389 端口传输至远端8181端口，远端通过连接 localhost:8181 可以连接本地 3389 端口

cmd 使用 ss 代理：

set HTTP_PROXY=socks5://127.0.0.1:10808
set HTTPS_PROXY=socks5://127.0.0.1:10808

网关：连接不同类型的网络

route -n #查看路由表

修改网卡设置：/etc/network/interfaces

auto lo
iface lo inet loopback
auto eth0
iface eth0 inet dhcp

arp : ip 和 mac 映射表

arp -a #查看 arp 表缓存
arp -d #删除 arp 表缓存

arp 欺骗：

sudo arpspoof -i [网卡] -t [目标ip] [网关ip]

nmap :

nmap -sP 192.168.40.0/24 #扫描网段中的 IP

VMware :
桥接模式：与主机处在同一网段，虚拟机拥有独立 IP，所有虚拟机可以和主机相互访问

10.10

在研究 pruning 时 bert 是一个主要的 benchmark
使用了 pytorch-pretrained-bert 的代码
在 i1 和 nico 上进行了测试，选用的 benchmark 为 SQuAD
- 其中一个奇怪的事情是在 i1 双卡上的用时略小于 nico 8卡上的用时

两个 epoch

i1 双卡： epoch1 48.37, epoch2 41.36
准确率： EM 81.164, f1 88.464
nico 8卡：epoch1 54.20, epoch2 54.02
准确率： EM 81.362, f1 88.475

在nico上测试了一个 epoch：
准确率：EM 80.142, f1 87.566

对 bert base uncased 进行了调研：

超参数：

{
  "attention_probs_dropout_prob": 0.1,
  "hidden_act": "gelu",
  "hidden_dropout_prob": 0.1,
  "hidden_size": 768,
  "initializer_range": 0.02,
  "intermediate_size": 3072,
  "max_position_embeddings": 512,
  "num_attention_heads": 12,
  "num_hidden_layers": 12,
  "type_vocab_size": 2,
  "vocab_size": 30522
}

总参数：109M (109483778)
其中 bert 部分参数数量: 109M (109482240)
bert部分：
embeddings : 23M（23837184）
encoder : 85M (85054464)
encoder 层为 12 个 Bertlayer 层的叠加：
每个 layer 参数数量为 7087872
每个layer中：
attention : 2363904
intermediate : 2362368
output : 2361600
三层有大致等量的参数
其中 attention 又分为 self 和 output 两部分
self 包含 q,k,v 主要 attention 部分，参数量: 1771776
output : 592128
所有layer中这三层的参数占模型参数总量 77.7%
embedding中参数占模型参数总量 21.8%
目前还不知道 embedding 中参数是否可prune

参数大致来源如下：

attention.self : hidden * hidden * 3
attention.output : hidden * hidden
intermediate : hidden * intermediate
output : intermediate * hidden

11.19

（这都隔了一个多月了呀喂）
在看一篇在 TASO 上做 sparse 的文章
"A sparse iteration space transformation framework for sparse tensor algebra"
是自动生成 sparse 的 CPU 和 GPU 操作的文章

另一种可能的 sparse op ： MTTKRP
$$ A = B_{(1)} (D \dot C) $$
$$ A_{ij} = B_{ikl} \dot D_{lj} \dot C_{kj} $$
其中 A,D,C 为二维矩阵，B为三维 tensor
在存储 sparse tensor 时可以引入新格式 CSF

cold cache : 冷缓存？
在小矩阵情况下 CPU 性能优于 GPU
TACO : 计算 tensor expression 的 C++ library

可以生成 atomic 操作 （相比 TVM 优点）