leenldk 发布的文章

6.30

开始看适之学长给的论文
introduction 中提到bert模型的pre_train过程代价极高
1024 V100 1day
bert large 很难在12GB ~ 16GB 的显卡上reproduce 结果

bert有multiple layers 的双向 transformers
每个 transformer 有一个 multi-head self-attention层,position-wise feed-forward层

防止假期摸鱼,记一下每天都干了什么
这个假期主要需要搞TVM那边的工作

6.11

$SHELL环境变量是当前的shell路径

6.12

tmux调整当前窗口全屏 : ctrl+b z
tmux新建窗口: ctrl+b c ,下一个窗口:ctrl+b n
ctrl+b [ : 页面滚动
任务提交dgx-dev : qsub -I -q dgx-dev

6.13

cmake添加前缀路径:CMAKE_PREFIX_PATH = /opt/fftw:/opt/cuda cmake ..

6.21

ssh -L [localport]:[host]:[hostport] [host] 将本地端口localport ssh forward到host的hostport上
ssh -D [localport] [host] 将host转发到localport

/sys : 在内存中
/var : variable,可能grow in size的内容

6.22

CPU flops计算:
对于nico的Intel® Xeon® Gold 5218 Processor:
每个核有1个AVX-512 FMA unit (单个时钟周期内两次512bit 的乘法和加法)
睿频3.9GHz,32位GFLOPS = 64 * 512 / 32 * 2 * 3.9 = 7987 => 8TFLOPS
单V100 single precision 14TFLOPS
带宽900GB/sec

nico1上v100的capability 7.0
参考一些 《CUDA C权威编程指南》里面的内容

atomicCAS(int *address, int compare, int val)
(*address == compare ? val : *address)
可以通过它对 capability < 6.0 设备实现 atomicAdd

一个有趣的东西 warp vote function:__all_sync/__any_sync...(unsigned mask, int predicate)
对wrap中 mask中每个线程传入一个predicate与0比较,返回结果到每个wrap

nvidia-smi -L : 查看所有GPU的设备ID
nvidia-smi -q -i 0 : 查看GPU0详细信息

一个block只能在一个SM上被调度
major revision number对应架构: Volta 7, Pascal 6, Maxwell 5, Kepler 3, Fermi 2, Tesla 1
Turing 7.5 (基于Volta的incremental)

-arch=compute_30 : 编译选项中使用3.0 capability

6.23

Fermi架构中shared memory 和 L1 cache 共享
Kepler : 每个SM中包含192个单精单元,64双精单元,32特殊单元,32加载存储单元
每个SM中有4个 wrap 调度器,8个指令调度器。可同时调度64个wrap,2048个线程
双精1TFLOPS
动态并行,GPU可以启动嵌套kernel

6.24

V100 release时间是2017.5
我觉得目前可以做的是研究V100的tensor core
PTX = parallel thread execution
active threads : 在所在wrap当前执行路径上的threads

ASpT的代码中大量用到了__shfl原语
在一个wrap的threads之间交换一个variable

int __shfl(int var, int srcLane, int width=warpSize);

返回srcLane中var的值
width必须是power of 2
width<warpSize时每width个thread被视为一个子wrap

matrix multiply and accumulate (D = A*B + C)

6.25

register bank conflict : volta架构中寄存器被分为两个64bit bank
一个volta指令每个cycle只能访问每个bank的64 bit

x86_64中,依次使用rdi, rsi, rdx传递函数参数
A100 GPU : 54 billion 晶体管 826 mm^2
ampere架构

6.26

intel cascade lake : 基于 sky lake 优化

5.11

在搞大物实验和数设项目

5.12

let binding :

let %x : Tensor[(10, 10), float32] = Constant(1, (10, 10), float32);
%x + %x

let表达式绑定的变量只作用在表达式body内

5.13

在看TVM优化卷积的应用
TVM使用tensor表示计算的量
tensor的表示形式在AI的一些如CNN等数据比较规整的领域较为合理,目前在考虑能否在不规整数据类型的应用中找到研究点
比如当前考虑的ASpT应用就是一种不规整数据类型

一个结论是在搞一个项目的时候首先应该尽量从已知的文档中获取信息,或者说make the best of前人的工作

tvm.te.schedule :

s = te.create_schedule(B.op)
WW = s.cache_read(W, "shared", [B])
## cache_read(tensor, scope, readers) 
## 将tensor cache到scope中,readers会读取cache的内容

5.14

vtune intel性能分析器

5.17

新学到的cuda相关东西:
shared memory以bank方式存储
为达到最大性能,一个wrap内32个线程:
访问不同bank 或 访问同一地址

同一个wrap中线程的不同执行路径需要串行执行。

为了push自己开了这个坑
主要记录一周里面每天在 实验室 和 超算 那边干的事情
以及大作业等方面的见闻
避免划水

5.5

尝试去装超算的Gromacs
看起来是一个分子模拟的应用
需要用到 intel 的 MKL 库
尝试去官网里面下,由于需要登录信息,因此wget的时候需要指定cookie wget --load-cooikes=cookies.txt [site]
安装时运行install.sh,自选了不安装Fortan部分
upd : 这个在编译的时候报了迷之错误,可能是MKL库没有链接好的问题,之后去问下学长吧

5.6

今天应该以TVM为主
首先在看TVM的developer guide,对TVM的语法结构更深地认识一下。
compiler stack中所有language object都是Object的子类
每个object有一个string type_key表示object的type
ObjectRef 相当于Object的sharedptr
每个Object的子类需要包含VisitAttr函数,通过重载VisitAttr访问成员
跑docker需要先加到docker组里面

和适之学长讨论了用TVM实现aspt的问题
目前的问题在于即使实现了thread间通过atomic_add的reduce功能,由于每个行块中列块数量不同,因此无法直接将列块数量视为一个额外的维度处理进行reduce
目前的方法是对每个行块中的所有列块用一个kernel进行reduce,但这种实现方法要求这些kernel之间并行执行
为此需要使用不同的stream执行这些kernel
那么现在的问题就是研究tvm使用stream的机制

TVM计算图中的node是placeholder或computational node
每个node包含:op(operation type),name等

TVM基于dmlc/HalideIR
HalideIR中loop_type(a,b)表示区间[a,a+b)
TVM中loop_type(a,b)表示区间[a,b)
HalideIR中的四种loop形式:
* A for loop. Execute the 'body' statement for all values of the * variable loop_var from 'min' to 'min + extent'. There are four * types of For nodes. A 'Serial' for loop is a conventional * one. In a 'Parallel' for loop, each iteration of the loop * happens in parallel or in some unspecified order. In a * 'Vectorized' for loop, each iteration maps to one SIMD lane, * and the whole loop is executed in one shot. For this case, * 'extent' must be some small integer constant (probably 4, 8, or * 16). An 'Unrolled' for loop compiles to a completely unrolled * version of the loop. Each iteration becomes its own * statement. Again in this case, 'extent' should be a small * integer constant.

5.7

开始调研搜索引擎大作业的技术栈
看bert-as-service的时候发现pretrained BERT Chinese是character-based。
目前的初步想法:从训练集中提取关键词,所有处理都基于这些关键词

5.8

白天在搞搜索引擎大作业
目前实现了从xml的value里面提取关键词,构造出一个jieba的关键词表。
发现正则处理简单xml的效果还不错
关键词主要有value里面出现过的,以及《》中的法律条文
然后用jieba对QW(全文)部分进行分词
下一步是通过word2vec构建词向量
晚上的时候主要focus在gormacs的编译上
折腾了好久,甚至在尝试把编译栈转移到icc上面,但后来发现他还依赖cuda,也就是说转移的话需要用icc编译的cuda
编译了一个icc的openmpi,icc编译的时候很慢,但编译后文件体积不大
最后和家傲学长的编译命令对比才发现,我看文档的时候不够仔细,文档的说明里面已经写了load mkl或者 加一个mkl的选项,但我load之后又加了选项,导致gg
目前为止的结论是cmake多生成的都会堆到当前文件夹下
所以把build删了就可以重新编译

5.9

在使用word2vec训练词向量
默认的epoch是5个

gromacs文件:
.pdb : protein databank file 描述分子结构
.top : topology 定义分子 包含nonbonded parameters (atom types and charges)和bonded parameters (bonds, angles, and dihedrals),topol.top : system topology
.gro : GROMACS-formatted structure file ,force field中的所有原子

5.10

早上起来继续看gromacs
.mdp : molecular dynamics parameter
.tpr : 系统中所有原子的所有参数,可通过.mdp生成

尝试跑了一下lignocellulose-rf.tpr的测例来测试性能
发现gromacs有一个GMX_THREAD_MPI的选项,不支持多机,但据说可以比MPI快一点
实测并没有太大用处
mpi版本命令:mpirun -n 4 gmx_mpi mdrun -s lignocellulose-rf.tpr -v -deffnm temp -ntomp 4 -nsteps 1000
threadMPI版本命令:gmx mdrun -s lignocellulose-rf.tpr -v -deffnm temp -ntomp 4 -ntmpi 4 -nsteps 1000

发现当前的编译选项用的是FFTW的数学库
目前感觉如果换一个intel-mkl的话可能可以获得性能提升
但现在mkl找library的时候还会遇到之前的问题
家傲学长说cmake有一个findmkl的子模块,之后去了解一下

counter <= (0      => '1',
            4      => '1',
            others => '0') // signal赋值
rising_edge(clk) //检测上升沿函数