git branch -d $BRANCH_NAME # 删除分支
git reset head (--{file})  # 将已经stage的文件丢弃到working directory
git stash save "message" # stash当前已经stage的文件，以message保存
git clone -b [branch] [repo] #clone指定branch
git branch -a # 列出包括远程分支在内所有分支
git branch --set-upstream-to=my_origin/master master #将master track到my_origin/master
git clean -fd -n #查看clean会删除的文件
git clean -fdx #删除untracked file
git checkout -- [file] #将文还原为上次提交的版本（若在暂存区中则还原暂存区中版本，否则还原版本库版本）
git remote -v 查看所有远程分支指向
git remote 

clone 后 单独初始化 submodule :

git submodule init
git submodule update

MPI : Message Passing Interface
多线程编程中进程间消息传递

进程（process） 线程(thread)
threads 之间共享 process 的 address space 和resource

thread safe：线程安全，多个线程访问时不会出现race condition

两种common通信协议
Eager协议：发送进程主动发送，不考虑接收进程是否有能力接收，适合发送小消息
Rendezvous协议：在接收端协调缓存接收信息。适合发送较大消息

rank：每个进程的编号，通过指定rank进行进程间通信
displacement,lb,rb单位均为byte

#include "mpi.h"
MPI_Init(&argc, &argv); //初始化MPI
MPI_Comm_rank(MPI_COMM_WORLD, &rank); // 获取rank
MPI_Get_processor_name(name, &namelen); // 节点hostname
MPI_Comm_size(MPI_COMM_WORLD, &size); // process数量
MPI_Finalize();

MPI_Reduce：将多个进程的结果合并到主进程

MPI_Reduce(
    void* send_data,
    void* recv_data,
    int count,
    MPI_Datatype datatype,
    MPI_Op op,
    int root,
    MPI_Comm communicator)

send_data ：需要reduce的data，type为datatype
recv_data：rank root接收的信息，大小sizeof(datatype) * count
op : reduce操作符 e.g. : MPI_MAX, MPI_MIN ....

MPI_Allreduce 将多个进程结果reduce后传给每个进程

MPI_Allreduce(
    void* send_data,
    void* recv_data,
    int count,
    MPI_Datatype datatype,
    MPI_Op op,
    MPI_Comm communicator)

同reduce，不需要root

MPI_Iprobe：非阻塞对信息检验

int MPI_Iprobe(
    int source, 
    int tag, 
    MPI_Comm comm, 
    int *flag, 
    MPI_Status * status)

在接收前检测接收的信息，不需要真正接收
使用MPI_Recv接收信息

MPI_Probe：阻塞的对信息检验

int MPI_Probe(
    int source, 
    int tag, 
    MPI_Comm comm, 
    MPI_Status *status)

同MPI_Iprobe，阻塞检验

MPI_Recv：阻塞接收信息

int MPI_Recv(
    void *buf, 
    int count, 
    MPI_Datatype datatype,
    int source, 
    int tag, 
    MPI_Comm comm, 
    MPI_Status *status)

阻塞接收信息，在写完buf后返回，可以先于对应send返回

MPI_Send(&outMsg, msgLen. MPI_CHAR, receiver, tag, MPI_COMM_WORLD);
MPI_Recv(&inMsg, msgLen, MPI_CHAR, sender, tag, MPI_COMM_WORLD, &stat);

MPI_Bcast 一对多广播

MPI_Bcast(
void* buffer, 
int count, 
MPI_Datatype datatype, 
int root, 
MPI_Comm comm)

count : 信息长度

MPI_Scatter 一对多发送不同消息
MPI_Gather 多对一收集不同消息
MPI_Scatterv 一对多发送不同长度消息
MPI_Gatherv 一对多接收不同长度消息
MPI_Scan 相当于结果为前缀和的all_reduce

MPI_Init(&argc, &argv);
MPI_Finalize();

open mpi

mpirun -n 16 sst emberLoad.py #指定mpi以16个rank执行

map(func, a) 对a中元素应用func,python3返回一个迭代器
对迭代器x通过 next(x)获取元素
list(x) 将剩余元素转为列表
dir(x) 获取所有属性与方法
id(x)
type(x) x类型

python -i [a.py] 在interactive shell中运行a.py
add = lambda x, y : x + y lambda函数
b = filter(lambda x : x > 1, a) 对list a进行filter

for i, num in enumerate(a):
    print(i, num)

i为编号，num为a中元素

ord(x) 将字符转为asacii码
zip(l1, l2, l3 {...}) 依次提取l1, l2, l3 ...中元素，打包为元组

func.__code__.co_argcount 函数func参数数量
func.__code__.co_varnames 函数参数列表

- 阅读剩余部分 -

reminder

SST和所有外部组件用相同编译器编译

术语

SM : streaming Multiprocessor
基础模块，有4block ( 16 fp32, 8 fp64, 16 int32, 128K L1 cache, 64k register）

SIMT : single instruction multiple threads
一组32 cores执行相同指令不同数据

gridDim.x
blockIdx.x
blockDim.x
threadIdx.x

nvcc --ptxas-options=-v : 显示寄存器和 shared mem 信息

概念

• Page-Locked Memory : Unified Virtual Addressing (UVA)，CPU 和 GPU 共享 VA，实现 zero copy，所有对 CPU 内存访问走 PCIe 在 CPU 使用 page-locked host memory 使用 cudaHostAlloc 和 cudaFreeHost 分配新内存 使用 cudaHostRegister 将 malloc 的内存变为 page-locked
• Unified Memory : GPU 自动管理内存，page fault 时从 CPU 拷贝到 GPU

usage

golbal : host调用device执行函数
__syncthreads() 同步一个block内的thread
shared int a[] : 每个block一个共享内存中的a
T __ldg(const T address) 从address读取一个数据T (从read-only data cache中读取)

thrust:

cuda 中类似 C++ STL

thrust::device_ptr<int> d_arr = thrust::device_pointer_cast<int>(d_bin_counter);
// host指针转device_ptr
thrust::exclusive_scan(d_arr, d_arr + SEGBIN_NUM, d_arr); // 每个位置变为自己之前不包括自己的前缀和

atomicAdd(int* address, int val) 将 val 原子加到 address

cudaDeviceSynchronize() block直到所有设备同步

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

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

cuda 中 default stream (stream 0) 可以保证所有之前任何 stream 的调用均在前执行，且所有之后任何 stream 的调用均在后执行 (serializing)

事件计时：

cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,
NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );

只能在 stream 0 中用此方法计时，其他 stream 无法保证顺序

多GPU

int deviceCount;
cudaGetDeviceCount(&deviceCount); // device 数量
cudaGetDeviceProperties(&deviceProp, 1); // 获取 device 1 的 properties
cudaSetDevice(0); // 设置当前 device 为 0，之后的内存分配和 kernel launch 等操作均在 device 0 上进行

特性

在 CUDA 中用户可以创建多个 stream，但同时不会有超过 32 个 kernel 执行
在一些 device 上， L1 cache 和 shared memory 共享硬件资源

optimization

#pragma unroll [factor] : 强制循环展开，factor为展开参数，factor=1为不展开
restrict : 指针重名优化（Pointer Aliasing），显示指定没有重名指针

A100

tensor core 理论 315TFLOPS
A100 dram 1500GB/s
A100 L2 40MB 所有SM共享 V100 L2 6MB

compute capability :
A100 : 8.0
V100 : 7.0
P100 : 6.0

额，请在编译时指定正确的 compute capability

hip

hipLaunchKernelGGL(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream*/, a, b, c, n /*kernel函数参数*/);
// Replace MyKernel<<<dim3(gridDim), dim3(gridDim), 0, 0>>> (a,b,c,n);