每周乱搞日记 (5)

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: