肯定有人问,这不是计算机博客吗?哦!诶!我就不,我想怎么写就怎么写😋
Piano: Transcribe Piano Sheet Music from Video using AI model
肯定有人问,这不是计算机博客吗?哦!诶!我就不,我想怎么写就怎么写😋
有对应的PPT,代码。
最终将1000ms程序优化到1~2ms
乔良师兄有根据知乎介绍如何利用寄存器文件缓存
导致Embarrassingly Parallel Problems
计算某点的梯度,需要前后的function值。
问题:
对于边界上的cells,需要访问相邻区域的元素。
解决办法:
将他们也加入进当前block的SMEM
s_idx = threadIdx.x + RAD;
1 | int main() { |
Kernel Launching
1 | #define TPB 64 |
Kernel Definition
1 | __global__ void ddKernel(float *d_out, const float *d_in, int size, float h) { |
暂无
暂无
研一下USTC并行计算自己的选题
CSS (Cascading Style Sheets) 和 SCSS (Sassy CSS) 都是用于样式表的编程语言,用于定义网页的外观和布局。
如果我们数据只会在 GPU 产生和使用,我们不需要来回进行拷贝。
https://migocpp.wordpress.com/2018/06/08/cuda-memory-access-global-zero-copy-unified/
简而言之,在 host 使用命令:cudaHostRegisterMapped
之后用 cudaHostGetDevicePointer 进行映射
最后解除绑定 cudaHostUnregister
即,
1 | // First, pin the memory (or cudaHostAlloc instead) |
只要两个thread在 同一个warp中,允许thread直接读其他thread的寄存器值,这种比通过shared Memory进行thread间的通讯效果更好,latency更低,同时也不消耗额外的内存资源来执行数据交换。ref
对齐(Starting address for a region must be a multiple of region size)集体访问,有数量级的差异Coalesced
利用好每个block里的thread,全部每个线程各自读取自己对齐(Starting address for a region must be a multiple of region size 不一定是自己用的)数据到shared memory开辟的总空间。由于需要的数据全部合力读取进来了,计算时正常使用需要的读入的数据。
特别是对于结构体使用SoA(structure of arrays)而不是AoS(array of structures),
如果结构体实在不能对齐, 可以使用 __align(X)
, where X = 4, 8, or 16.强制对齐。
__syncthreads
self-tuning
出来占用率是指每个多处理器(Streaming Multiprocessor,SM)的实际的活动warps数量与最大理论的warps数量的比率。
高的占用率不一定能提升性能,因为这一般意味着每个线程分配的寄存器和shared memory变少。但低的占用率会导致内存延迟无法隐藏。
实际需要计算每个线程大概需要的shared memory和register数量
https://www.cnblogs.com/1024incn/p/4541313.html
https://www.cnblogs.com/1024incn/p/4545265.html
通过SMEM实现coalescing access
原本代码
1 | _global__ void transpose_naive(float *odata, float *idata, int width, int height) |
思想:将大矩阵划分成方块,并且存储在SMEM里。不仅SMEM速度更快,而且每行元素个数变少,跨行访问的间距变小,局部性增强。而且对于大矩阵加速效果会更明显。
1 | __global__ void transpose(float *odata, float *idata, int width, int height) |
when Block/tile dimensions are multiples of 16 ???
https://developer.nvidia.com/blog/efficient-matrix-transpose-cuda-cc/
对于一个32 × 32个元素的共享内存块,一列数据中的所有元素都映射到相同的SMEM bank ,导致bank conflict 的最坏情况:读取一列数据会导致32路的存储库冲突。
幸运的是,只需要将tile的元素宽度改为33,而不是32就行。
具体问题:将长数组的所有元素,归约求和为一个结果。[^1][^2]
为了避免全局同步的巨大开销,采取分级归约
由于归约的计算密度低
1 flop per element loaded (bandwidth-optimal)
所以优化目标是将访存带宽用满。
1 | 384-bit memory interface, 900 MHz DDR |
1 | __global__ void reduce0(int *g_idata, int *g_odata) { |
工作的线程越来越少。一开始是全部,最后一次只有thread0.
Just replace divergent branch With strided index and non-divergent branch,但是会带来bank conflict。
原理和Warp发射有关,假如在这里每个Warp并行的线程是2。一个Warp运行耗时为T.
Step0: 4+4+2+1=11T
Step1: 4+2+1+1=8T
1 | for (unsigned int s=1; s < blockDim.x; s *= 2) { |
1 | for (unsigned int s=blockDim.x/2; s>0; s>>=1) { |
原本寻址
现在寻址有一边连续了
方法: 在load SMEM的时候提前做一次规约加法,通过减少一半的block数,将原本两个block里的值load+add存储在sum里。
1 | // perform first level of reduction, |
当s< 32的时候,就只有一个Warp工作了。
使用warp的SIMD还省去了__syncthreads()
的麻烦
1 | for (unsigned int s=blockDim.x/2; s>32; s>>=1) |
为了保持整洁,最后一个if还做了无效的计算。eg, Warp里的最后一个线程只有第一句命令有用。
由于for循环里是二分的,而且小于32的单独处理了,导致for循环里实际运行代码最多就3句。
利用代码模板和编译器的自动优化实现:
1 | template <unsigned int blockSize> |
红色代码会在编译时自动优化。
加速级联??
Cost= processors × time complexity
我们知道N个元素直接二叉树归约是O(log N)
时间 Cost=N*O(log N).
但是假如只有P个线程先做N/P的串行加法, 然后是log(P)的归约。
总cost=P(N/P+log(P))
当P=N/log(N), cost=O(N)
each thread should sum O(log n) elements来设置
比如,1024 or 2048 elements per block vs. 256 线程。每个sum n=4个元素。 具体参数要perf
1 | unsigned int tid = threadIdx.x; |
1 | template <unsigned int blockSize> |
有if语句是没问题的,只要运行的时候全部执行if或者else就行。不要有些执行if,有些执行else,这才会等待。
说不定也不是全部执行if或者else就行,只需要连续32个Thread Index,是相同的执行就行。(猜想,需要测试。
通过增加block里的线程数,并且同时读取来隐藏延迟。 不仅可以隐藏Global Memory的延迟,还可以隐藏写后读的延迟
线程太多会导致分配到每一个的寄存器和SMEM变少
通过编译时加-cubin
选项,.cubin
文件前几行会显示
1 | architecture {sm_10} |
[^1]: SC07 Optimizing Parallel Reduction in CUDA - Mark Harris
[^2]: 2009 清华 邓仰东 cuda lecture pdf 注意也是参考的SC07 Nvidia。
1 | $ which nvprof |
命令行直接运行
1 | nvprof ./myApp |
1 | nvprof --print-gpu-trace ./myApp |
1 | sudo /usr/local/cuda/bin/nvprof --log-file a.log --metrics achieved_occupancy /staff/shaojiemike/github/cutests/22-commonstencil/common |
1 | nvprof --export-profile timeline.prof <app> <app args> |
1 | sudo /usr/local/cuda/bin/ncu -k stencil_kernel -s 0 -c 1 /staff/shaojiemike/github/cutests/22-commonstencil/best |
ncu-ui是可视化界面,但是没弄懂
1 | # shaojiemike @ snode0 in ~/github/cuda-samples-11.0 [16:02:08] $ ./bin/x86_64/linux/release/bandwidthTest [CUDA Bandwidth Test] - Starting... Running on... Device 0: Tesla P40 Quick Mode Host to Device Bandwidth, 1 Device(s) PINNED Memory Transfers Transfer Size (Bytes) Bandwidth(GB/s) 32000000 11.8 Device to Host Bandwidth, 1 Device(s) PINNED Memory Transfers Transfer Size (Bytes) Bandwidth(GB/s) 32000000 13.0 Device to Device Bandwidth, 1 Device(s) PINNED Memory Transfers Transfer Size (Bytes) Bandwidth(GB/s) 32000000 244.3 Result = PASS NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled. # shaojiemike @ snode0 in ~/github/cuda-samples-11.0 [16:03:24] $ ./bin/x86_64/linux/release/p2pBandwidthLatencyTest |
nvprof通过指定与dram,L1或者L2 的metrics来实现。具体解释可以参考官网
在 Maxwell 和之后的架构中 L1 和 SMEM 合并
Metric Name | 解释 |
---|---|
achieved_occupancy | 活跃cycle是 Warps 活跃的比例 |
dram_read_throughput | |
dram_utilization | 在0到10的范围内,相对于峰值利用率,设备内存的利用率水平 |
shared_load_throughput | |
shared_utilization | |
l2_utilization |
暂无
暂无
Liquid
as its templating language. Hugo uses Go
templating. Most people seem to agree that it is a little bit easier to learn Jekyll’s syntax than Hugo’s.^1https://cloud.tencent.com/developer/article/1448642
1 | shaojiemike@snode6:~$ groups shaojiemike |
1 | cat /etc/group |
/etc/passwd
如果发现自己不在/etc/passwd里,很可能使用了ldap 集中身份认证。可以在多台机器上实现分布式账号登录,用同一个账号。
1 | getent passwd |
1 | ctrl + alt + F3 |
宕机一般是爆内存,进程分配肯定会注意不超过物理核个数。
在zshrc里写入 25*1024*1024 = 25GB的内存上限
1 | ulimit -v 26214400 |
当前shell程序超内存,会输出Memory Error
结束。
1 | with open("/home/shaojiemike/test/DynamoRIO/OpenBLASRawAssembly/openblas_utest.log", 'r') as f: |
有文章说Linux有些版本内核会失效
分布式、多平台集成认证系统
https://ibug.io/blog/2022/03/linux-openldap-server/
https://harrychen.xyz/2021/01/17/openldap-linux-auth/
https://www.cnblogs.com/dufeixiang/p/11624210.html
复杂还有bug,我还是改profile吧
https://ibug.io/blog/2022/03/linux-openldap-server/#user-chsh
挂在同一个地方,肯定是一样的
1 | # shaojiemike @ snode2 in ~ [20:18:20] |
tmpfs是磁盘里的虚拟内存的意思。
具体设置要登录到中央机器上去
1 | # shaojiemike @ hades1 in ~ [20:41:06] |
1 | # shaojiemike @ snode0 in ~ [20:36:26] |
暂无