CSS (Cascading Style Sheets) 和 SCSS (Sassy CSS) 都是用于样式表的编程语言,用于定义网页的外观和布局。
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.强制对齐。
__syncthreadsself-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.^11 | > curl myip.ipip.net |
1 | # 网页服务直接下载检查内容 |
或者扫描指定端口
1 | # IPV6 也行 |
全部端口,但是会很慢。50分钟
1 | sudo nmap -sT -p- 4.shaojiemike.top |
上方的过滤窗口
1 | tcp.port==80&&(ip.dst==192.168.1.2||ip.dst==192.168.1.3) |
抓包前在capture option中设置,仅捕获符合条件的包,可以避免产生较大的捕获文件和内存占用,但不能完整的复现测试时的网络环境。
1 | host 192.168.1.1 //抓取192.168.1.1 收到和发出的所有数据包 |

传统命令行抓包工具
注意过滤规则间的and
-nn :-i 指定网卡 -D查看网卡-v,-vv 和 -vvv 来显示更多的详细信息port 80 抓取 80 端口上的流量,通常是 HTTP。在前面加src,dst限定词tcpudmp -i eth0 -n arp host 192.168.199 抓取192.168.199.* 网段的arp协议包,arp可以换为tcp,udp等。-A,-X,-xx会逐渐显示包内容更多信息-e : 显示数据链路层信息。1 | 192.168.1.106.56166 > 124.192.132.54.80 |
> 符号代表数据的方向。常见的三次握手 TCP 报文的 Flags:
1 | [S] : SYN(开始连接) |
1 | curl --trace-ascii - www.github.com |
github ip 为 20.205.243.166
ifconfig显示 ibs5的网卡有21TB的带宽上限,肯定是IB卡了。
1 | sudo tcpdump -i ibs5 '((tcp) and (host 20.205.243.166))' |
1 | $ sudo tcpdump -i ibs5 -nn -vvv -e '((port 80) and (tcp) and (host 20.205.243.166))' tcpdump: listening on ibs5, link-type LINUX_SLL (Linux cooked v1), capture size 262144 bytes |
snode0 ip 是 10.1.13.50
mtr = traceroute+ping
1 | $ traceroute www.baid.com |
traceroute命令用于显示数据包到主机间的路径。
1 | # shaojiemike @ snode0 in /etc/NetworkManager [16:49:55] |
好像之前使用过的样子。
1 | # shaojiemike @ snode0 in /etc/NetworkManager [16:56:36] C:127 |
应该是这个 Secure site-to-site connection with Linux IPsec VPN 来设置的
暂无
暂无
FJW说所有网络都是通过NFS一起出去的
https://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有些版本内核会失效
PyG是一个基于PyTorch的用于处理不规则数据(比如图)的库,或者说是一个用于在图等数据上快速实现表征学习的框架。它的运行速度很快,训练模型速度可以达到DGL(Deep Graph Library )v0.2 的40倍(数据来自论文)。除了出色的运行速度外,PyG中也集成了很多论文中提出的方法(GCN,SGC,GAT,SAGE等等)和常用数据集。因此对于复现论文来说也是相当方便。
经典的库才有函数可以支持,自己的模型,自己根据自动微分实现。还要自己写GPU并行。
MessagePassing 是网络交互的核心
torch_geometric.data.Data (下面简称Data) 用于构建图
通过data.face来扩展Data
在 PyG 中,我们使用的不是这种写法,而是在get()函数中根据 index 返回torch_geometric.data.Data类型的数据,在Data里包含了数据和 label。

由于是无向图,因此有 4 条边:(0 -> 1), (1 -> 0), (1 -> 2), (2 -> 1)。每个节点都有自己的特征。上面这个图可以使用 torch_geometric.data.Data来表示如下:
1 | import torch |
注意edge_index中边的存储方式,有两个list,第 1 个list是边的起始点,第 2 个list是边的目标节点。注意与下面的存储方式的区别。
1 | import torch |
这种情况edge_index需要先转置然后使用contiguous()方法。关于contiguous()函数的作用,查看 PyTorch中的contiguous。
1 | import torch |
DataLoader 这个类允许你通过batch的方式feed数据。创建一个DotaLoader实例,可以简单的指定数据集和你期望的batch size。
1 | loader = DataLoader(dataset, batch_size=512, shuffle=True) |
DataLoader的每一次迭代都会产生一个Batch对象。它非常像Data对象。但是带有一个‘batch’属性。它指明了了对应图上的节点连接关系。因为DataLoader聚合来自不同图的的batch的x,y 和edge_index,所以GNN模型需要batch信息去知道那个节点属于哪一图。
1 | for batch in loader: |

其中,x 表示表格节点的 embedding,e 表示边的特征,ϕ 表示 message 函数,□ 表示聚合 aggregation 函数,γ 表示 update 函数。上标表示层的 index,比如说,当 k = 1 时,x 则表示所有输入网络的图结构的数据。
为了实现这个,我们需要定义:
[2, num_messages], where messages from nodes in edge_index[0] are sent to nodes in edge_index[1]会根据 flow=“source_to_target”和if flow=“target_to_source”或者x_i,x_j,来区分处理的边。
x_j表示提升张量,它包含每个边的源节点特征,即每个节点的邻居。通过在变量名后添加_i或_j,可以自动提升节点特征。事实上,任何张量都可以通过这种方式转换,只要它们包含源节点或目标节点特征。
_j表示每条边的起点,_i表示每条边的终点。x_j表示的就是每条边起点的x值(也就是Feature)。如果你手动加了别的内容,那么它的_j, _i也会自动进行处理,这个自己稍微单步执行一下就知道了
在实现message的时候,节点特征会自动map到各自的source and target nodes。
aggregation scheme 只需要设置参数就好,“add”, “mean”, “min”, “max” and “mul” operations
aggregation 输出作为第一个参数,后面的参数是 propagate()的
$$
\mathbf{x}i^{(k)} = \sum{j \in \mathcal{N}(i) \cup { i }} \frac{1}{\sqrt{\deg(i)} \cdot \sqrt{\deg(j)}} \cdot \left( \mathbf{\Theta}^{\top} \cdot \mathbf{x}_j^{(k-1)} \right)
$$
该式子先将周围的节点与权重矩阵\theta相乘, 然后通过节点的度degree正则化,最后相加
步骤可以拆分如下
步骤1 和 2 需要在message passing 前被计算好。 3 - 5 可以torch_geometric.nn.MessagePassing 类。
添加self-loop的目的是让featrue在聚合的过程中加入当前节点自己的feature,没有self-loop聚合的就只有邻居节点的信息。
1 | import torch |
所有的逻辑代码都在forward()里面,当我们调用propagate()函数之后,它将会在内部调用message()和update()。
1 | conv = GCNConv(16, 32) |

聚合函数(aggregation)我们用最大池化(max pooling),这样上述公示中的 AGGREGATE 可以写为:
上述公式中,对于每个邻居节点,都和一个 weighted matrix 相乘,并且加上一个 bias,传给一个激活函数。相关代码如下(对应第二个图):
1 | class SAGEConv(MessagePassing): |
对于 update 方法,我们需要聚合更新每个节点的 embedding,然后加上权重矩阵和偏置(对应第一个图第二行):
1 | class SAGEConv(MessagePassing): |
综上所述,SageConv 层的定于方法如下:
1 | import torch |
GNN的batch实现和传统的有区别。
将网络复制batch次,batchSize的数据产生batchSize个Loss。通过Sum或者Max处理Loss,整体同时更新所有的网络参数。至于网络中循环输入和输出的H^(t-1)和H^t。(感觉直接平均就行了。
有几个可能的问题
通过 rescaling or padding(填充) 将相同大小的网络复制,来实现新添加维度。而新添加维度的大小就是batch_size。
但是由于图神经网络的特殊性:边和节点的表示。传统的方法要么不可行,要么会有数据的重复表示产生的大量内存消耗。
为此引入了ADVANCED MINI-BATCHING来实现对大量数据的并行。
https://pytorch-geometric.readthedocs.io/en/latest/notes/batching.html

可以实现将多个图batch成一个大图。 通过重写collate()来实现,并继承了pytorch的所有参数,比如num_workers.
在合并的时候,除开edge_index [2, num_edges]通过增加第二维度。其余(节点)都是增加第一维度的个数。
1 | # 原本是[2*4] |
1 | from torch_geometric.data import Data |
1 | from torch_geometric.datasets import TUDataset |
暂无
暂无