Microarchitecture: Zero (one) idioms & Mov Elimination

微架构的关系

寄存器重命名是乱序执行Tomasulo算法的一部分

寄存器重命名可以实现:

  1. 部分mov消除
  2. NOPs
  3. zero (one) idioms
    对于这些指令,无序发射到scheduler。可以直接在reorder buffer写入结果。

Zero (one) idioms

Zero (one) idioms 是不管原寄存器src值是什么,结果/目的寄存器dst一直/一定是0 (1)的一类指令。比如:XOR一个寄存器和自己。

  1. 由于是在寄存器重命名阶段(Rename)时实现的
    1. 所以不需要发射到port执行单元执行,占用硬件资源。也没有延迟
    2. 但是需要划分前面部分的decode的带宽,和ROB(reorder buffer)的资源
      1
      2
      sub eax, eax
      xor eax, eax

例子

使用uarch-bench

1
2
3
xor eax, eax
dec edi
jnz .loop

由于第一条指令是Zero idioms;后两条指令可以macro-fusion。

所以各部分平均执行次数为

指令个数 UOPS_ISSUED UOPS_EXECUTED UOPS_RETIRED
3 2 1 2

特殊的情况

有些架构可能不支持srcImm0-dstReg的指令的Zero idioms

1
mov eax, 0 

mov Elimination

  1. 由于是在寄存器重命名阶段(Rename)时实现的
    1. 所以不需要发射到port执行单元执行,占用硬件资源。也没有延迟
    2. 但是需要划分前面部分的decode的带宽,和ROB(reorder buffer)的资源

例子

1
2
3
4
5
add eax,4
mov ebx,eax ; //寄存器重命名,ebx指向eax即可
sub ebx,ecx
dec edi
jnz .loop

由于第二条指令是mov Elimination;后两条指令可以macro-fusion。

所以各部分平均执行次数为

指令个数 UOPS_ISSUED UOPS_EXECUTED UOPS_RETIRED
5 4 3 4

被覆盖的结果是否能消除

1
2
3
4
mov eax, 1 ; will be eliminated?
mov eax, 2
dec edi
jnz .loop

第一个mov被覆盖了。这是属于编译器的工作。CPU做不到这点(即使做得到,为了实现这点设计的硬件开销也会很大,不值得)

无效操作是否能消除

一般和0的立即数作用有关

1
2
xor eax, eax 
sub ebx, eax ; will be eliminated? (eax is always 0)

第二条指令在IvyBridge也不会消除。这同样是编译器的工作

但是llvm-mca通过ZeroRegister的实现,可以消除。

类似的还有

1
2
3
mov eax, 0
mov ebx, 0
cmp eax, ebx ; eax and ebx are always equal

一般也不会消除。这同样是编译器的工作

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://randomascii.wordpress.com/2012/12/29/the-surprising-subtleties-of-zeroing-a-register/

https://easyperf.net/blog/2018/04/22/What-optimizations-you-can-expect-from-CPU

https://zh.m.wikipedia.org/zh-hans/%E5%AF%84%E5%AD%98%E5%99%A8%E9%87%8D%E5%91%BD%E5%90%8D

Git Push 2 Homepage

ibug的网站部署思路

  1. 基于ibug.github.io
  2. 图片markdown两个仓库
  3. 对于acsa的网站
    1. 设置了action产生public/*.html
    2. 通过webhook来实现,服务器接收仓库的event信息。
      1. acsa的nginx接收location转发snode5
      2. snode5的nginx转发到127.0.0.2:9000上
      3. webhook.service接收到信息,然后git clone。并返回信息

hugo网站的action文件

根据公开的仓库,hugo的html文件会产生在gh-pages分支下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
name: build

on:
push:
branches: [master]

jobs:
build:
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@v2
with:
#submodules: true # Fetch Hugo themes (true OR recursive)
fetch-depth: 0 # Fetch all history for .GitInfo and .Lastmod

- name: Setup Hugo
uses: peaceiris/actions-hugo@v2
with:
hugo-version: 'latest'
#extended: true

- name: Build
run: hugo --minify

- name: Deploy
uses: peaceiris/actions-gh-pages@v3
if: github.ref == 'refs/heads/master'
with:
github_token: ${{ secrets.GITHUB_TOKEN }}
publish_dir: ./public

webhook的实现

接收端转发到内网的机器上(通过修改vim /etc/nginx/sites-enabled/default

1
2
3
4
5
6
server{
location /_webhook/ {
proxy_pass http://snode5.swangeese.fun;
proxy_set_header Host $http_host;
}
}

记得reload systemctl reload nginx

Nginx中location的作用是根据Url来决定怎么处理用户请求(转发请求给其他服务器处理或者查找本地文件进行处理)。location支持正则表达式,配置十分灵活。我们可以在一个虚拟主机(nginx中的一个server节点)下配置多个location以满足如动静分离,防盗链等需求。

在snode5上nginx也需要转发

1
2
3
location /_webhook/ {
proxy_pass http://127.0.0.2:9000;
}

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

Golang Syntax

为什么要学习go语言

  1. 同步方式轻松实现高并发,充分利用多核
  2. 基于消息传递的通信方式
  3. 适合服务器和网络编程
  4. 有垃圾回收机制
  5. 静态语言,有编译过程,和独立的静态可执行文件,只依赖glibc
    1. 不像python要安装各种库,java也要JRE
  6. 兼顾python的易开发性和c的性能
  7. 内存占用极小,支持10W+的并行

一些缺点

  1. 实际运行时,由于GC的影响,延迟会比较严重
  2. 代码会有很多重复的地方

有趣的工具

  1. gofmt
  2. gofix
  3. govet

数据类型

  • int8类型 表示 -128~127
  • Channel 类型
  • 切片类型 (可变长数组

变量声明

第一种,指定变量类型,如果没有初始化,则变量默认为零值

1
2
//var v_name v_type
var b, c int = 1, 2
1
2
3
4
5
6
7
//特殊
var a *int
var a []int
var a map[string] int
var a chan int
var a func(string) int
var a error // error 是接口

第二种,根据值自行判定变量类型。

1
2
//var v_name = value
var d = true

第三种,使用声明符号:=

但是如果变量已经使用 var 声明过了,再使用 := 声明变量,就产生编译错误,格式:

1
v_name := value

循环语句

1
2
3
for key, value := range oldMap {
newMap[key] = value
}

并发和通道通讯

go函数

Go 语言支持并发,我们只需要通过 go 关键字来开启 goroutine 即可。

goroutine 是轻量级线程,goroutine 的调度是由 Golang 运行时进行管理的。

goroutine 语法格式:go 函数名( 参数列表 )

Go 允许使用 go 语句开启一个新的运行期线程, 即 goroutine,以一个不同的、新创建的 goroutine 来执行一个函数。 同一个程序中的所有 goroutine 共享同一个地址空间。

通道(channel)

通道可用于两个 goroutine 之间通过传递一个指定类型的值来同步运行和通讯。操作符 <- 用于指定通道的方向,发送或接收。如果未指定方向,则为双向通道。

1
2
3
ch <- v    // 把 v 发送到通道 ch
v := <-ch // 从 ch 接收数据
// 并把值赋给 v

声明一个通道很简单,我们使用chan关键字即可,通道在使用前必须先创建:

1
ch := make(chan int)

example

1

1
2
3
4
5
6
7
8
9
10
11
12
13
func countGoodRectangles(rectangles [][]int) int {
cnt, maxLen := 0, 0
for _, rectangle := range rectangles {
k := int(math.Min(float64(rectangle[0]), float64(rectangle[1])))
if k == maxLen {
cnt++
}
if k > maxLen {
maxLen, cnt = k, 1
}
}
return cnt
}

webhook

https://github.com/swangeese/acsa-web/tree/webhook

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://www.runoob.com/go/go-concurrent.html

Go Install and Command

Install

1
2
3
4
5
6
wget https://go.dev/dl/go1.18.3.linux-amd64.tar.gz
rm -rf /usr/local/go && tar -C /usr/local -xzf go1.18.3.linux-amd64.tar.gz
(maybe need sudo)
sudo rm -rf /usr/local/go && sudo tar -C /usr/local -xzf go1.18.3.linux-amd64.tar.gz
export PATH=$PATH:/usr/local/go/bin
go version

Command usage

1
2
3
4
5
6
$ cd $HOME/go/src/hello
$ go run main.go #直接运行
Hello, World!!
$ go build # 产生可执行文件
$ ./hello
Hello, World!!

包管理

Packages

Go packages are folders that contain one more go files.

Modules

A modules (starting with vgo and go 1.11) is a versioned collection of packages.

1
2
go get github.co­m/a­nda­nhm­/go­-pr­ett­ytimee
go mod init github.co­m/a­nda­nhm­/go­-pr­ett­ytime

go list -m -u all 来检查可以升级的package,

使用go get -u need-upgrade-package 升级后会将新的依赖版本更新到go.mod

也可以使用 go get -u 升级所有依赖

作者:若与
链接:https://www.jianshu.com/p/760c97ff644c
来源:简书
著作权归作者所有。商业转载请联系作者获得授权,非商业转载请注明出处。

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://devhints.io/go

Go mod

简介

go modules 是 golang 1.11 新加的特性。现在1.12 已经发布了,是时候用起来了。Modules官方定义为:

模块是相关Go包的集合。modules是源代码交换和版本控制的单元。 go命令直接支持使用modules,包括记录和解析对其他模块的依赖性。modules替换旧的基于GOPATH的方法来指定在给定构建中使用哪些源文件。

使用

初始化项目

1
2
3
mkdir Gone
cd Gone
go mod init Gone

对应go.mod文件

1
2
module Gone
go 1.14

go.mod文件一旦创建后,它的内容将会被go toolchain全面掌控。

go toolchain会在各类命令执行时,比如go get、go build、go mod等修改和维护go.mod文件。

go.mod 提供了module, require、replace和exclude 四个命令

module 语句指定包的名字(路径)
require 语句指定的依赖项模块
replace 语句可以替换依赖项模块
exclude 语句可以忽略依赖项模块

自动添加依赖

对于main.go里的import

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
package main

import (
"crypto/hmac"
"crypto/sha1"
"encoding/hex"
"encoding/json"
"fmt"
"io/ioutil"
"log"
"net/http"
"os"
"os/exec"
"strings"
)

……

执行 go run main.go 运行代码会发现 go mod 会自动查找依赖自动下载,并修改go.mod(安装 package 的原則是先拉最新的 release tag,若无tag则拉最新的commit)

自己发布module包

结合github很简单实现

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://www.jianshu.com/p/760c97ff644c

Game Streaming & Video Streaming

视频传输编码

在初步接触了视频传输编码之后,我开始好奇Streaming采用的哪种视频编码呢?

详见 Streaming Protocol一文

Moonlight for IPAD

Nvidia Geforece界面

  1. 常规中开启分享
  2. SHEILD 开启
    1. 添加,按照C:\Windows\System32\mstsc.exe这个地址,将mstsc.exe添加进去,mstsc.exe就是你的桌面,等会串流,可以用手机直接操控你的电脑桌面。

串流画面问题

moonlight找不到电脑

https://www.bilibili.com/read/cv10239020

1
2
netsh advfirewall firewall add rule name="GameStream UDP" dir=in protocol=udp localport=5353,47995,47998-48010 action=allow
netsh advfirewall firewall add rule name="GameStream TCP" dir=in protocol=tcp localport=47984,47989,47995,48010 action=allow


还是不行,猜测是

但是这个是wifi6 11ax

尝试安装Internet-Hosting-Tool,运行有提示失败,建议重装也不行。

sjf的解决办法

  1. 卸载当前版本,然后安装3.19
  2. 打开服务
    1. 找到属性 单击打开。找到登录 复选框 把里面的登录身份选择成 本地系统账户 运行服务与桌面交换。勾选然后应用
    2. moonlight可以搜索到电脑
    3. https://pan.baidu.com/s/1x83Uk4kkYQritiNAqg_vLg [/url]提取码:1111 获得NvContainerNetworkService服务注册表文件
  3. 官网下载更新GF到最新

实际解决

  1. 官网下载更新GF到最新
  2. 通过上面的注册表添加NvContainerNetworkService服务,启动
    1. 任务计划程序里设置, 设置开机启动moonlightNVNetStart任务
    2. 程序"C:\Program Files\NVIDIA Corporation\NvContainer\nvcontainer.exe"
    3. 参数-s NvContainerNetworkService -f "C:\ProgramData\NVIDIA\NvContainerNetworkService.log" -l 3 -d "C:\Program Files\NVIDIA Corporation\NvContainer\plugins\NetworkService" -r -p 30000 -st "C:\Program Files\NVIDIA Corporation\NvContainer\NvContainerTelemetryApi.dll"
  3. 成功
  4. 修改分辨率为ipad分辨率,全屏应用
  5. 修改英伟达控制面板的分辨率为IPAD 2388*1688 macbook 2560*1600

问题:Nvidia控制面板没有显示一项

  • 如果显卡驱动装好,且显卡都开了,但就是没有显示选项。
  • 打开服务,找到NVDisplay.ContainerLocalSystem,点登录项,将“允许服务与桌面交互(W)”前的勾打上,重启NVDisplay.ContainerLocalSystem服务,
  • 返回桌面,右键-显示设置,将分辩率任意改一个可用的-应用,
  • 桌面右键-N..控制面板,就有了显示选项,可以改2K分辩率啦。

IPAD moonlight 串流控制

type ESC and mouse scroll

对于实体键盘可以修改映射, 但是滚轮就不好用了。

与其这样不如换个思路,添加手柄,看其能不能支持滚轮和ESC。初步尝试,滚轮可以只是灵敏度有点低。AntiMicroX完美解决了这个问题,配置文件路径 I:\BT\GAME\x18Game\moonlightAntiMicroX.gamecontroller.amgp

体验十分丝滑,任意程序也可以添加。支持PS4手柄(长按PS和share键配对)

share时一定要登录steam

任意应用全屏

应用串流至少将某应用窗口转发,所以只需要停止流式传输,然后调整分辨率就行了。

晚上关闭屏幕,不休眠

  1. 方法一:管理员运行代码 @powercfg -h off

云原神测试

高画质60帧 1.6MB/s

最低画质30帧 500KB/s

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

Crawler

如何获取请求链接

这个api是怎么来的呢?
lesson_info_url = "https://www.eeo.cn/saasajax/webcast.ajax.php?action=getLessonLiveInfo"

感谢大佬回答

输入

返回数据

PHP源文件

PHP是后端语言,前端是无法查看的,前端看到的是最终运算之后的结果,PHP源代码是无法查看的。

使用

header改一下就能用了,注意不要开代理

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
from requests import Session

session = Session()
lesson_info_url = "https://www.eeo.cn/saasajax/webcast.ajax.php?action=getLessonLiveInfo"

headers = {
'Accept': 'text/html,application/xhtml+xml,application/xml;q=0.9,*/*;q=0.8',
'Accept-Language': 'zh-CN,zh;q=0.9,zh-TW;q=0.8,en;q=0.7',
'User-Agent': 'Mozilla/5.0 (Windows NT 10.0; Win64; x64) AppleWebKit/537.36 (KHTML, like Gecko) Chrome/80.0.3987.149 Safari/537.36'
}

data = {
'lessonKey': lessonKey
}

resp = session.post(url=lesson_info_url, headers=headers, data=data)
text = resp.json()
CourseName = text['data']['courseName']

urllib.request下载视频

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
from urllib import request

base_url = 'https://f.us.sinaimg.cn/001KhC86lx07laEy0PtC01040200y8vC0k010.mp4?label=mp4_hd&template=640x360.28&Expires=1528689591&ssig=qhWun5Mago&KID=unistore,video'
#下载进度函数
def report(a,b,c):
'''
a:已经下载的数据块
b:数据块的大小
c:远程文件的大小
'''
per = 100.0 * a * b / c
if per > 100:
per = 100
if per % 1 == 1:
print ('%.2f%%' % per)
#使用下载函数下载视频并调用进度函数输出下载进度
request.urlretrieve(url=base_url,filename='weibo/1.mp4',reporthook=report,data=None)

例子一

小白尝试 学校的资源网址(http://wlkt.ustc.edu.cn/)

爬取List读取

正则匹配video/detail出视频网址后缀

网页视频位置

正则匹配mp4.php得到视频位置http://wlkt.ustc.edu.cn/mp4.php?file=HXMEV11IQNB2ZXPM6BVWY77AJ2HZTM4U

但是不打开网站没有php返回,网页只能得到。

可通过下面API返回需要的, 可以见github代码

1
2
3
4
5
6
7
8
9
opener = urllib.request.FancyURLopener({})
f = opener.open(taskUrl)
content = f.read()

#1.得到beautifulsoup对象
soup = BeautifulSoup(content,'html.parser')

#通过指定的 属性获取对象
ic(soup.find(id=glv._get(taskType)["data1id"]).attrs['value'])#单个对象


data输入

返回数据

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

科大BB clashIn 想爬录像。但是网上的两个都用不了了,想自学,改一下

https://github.com/aoxy/ClassIn-Video-Download

https://github.com/JiangGua/classin-downloader

参考文献

https://blog.csdn.net/qq_37275405/article/details/80780925

EpicUnrealEngine

我跪了,看来垃圾电脑玩不来,官方光明山脉demo要64GB内存和200GB储存。而且打开渲染超级慢

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

  1. 想实现美少女跳舞
    1. 其实好像Unity 3D更简单

参考文献

Cuda Optimize : Vectorized Memory Access

baseline

1
2
3
4
5
6
7
8
9
10
11
12
13
__global__ void device_copy_scalar_kernel(int* d_in, int* d_out, int N) { 
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N; i += blockDim.x * gridDim.x) {
d_out[i] = d_in[i];
}
}

void device_copy_scalar(int* d_in, int* d_out, int N)
{
int threads = 128;
int blocks = min((N + threads-1) / threads, MAX_BLOCKS);
device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

简单的分块拷贝。

通过cuobjdump -sass executable.得到对应的标量copy对应的SASS代码

1
2
3
4
5
6
/*0058*/ IMAD R6.CC, R0, R9, c[0x0][0x140]                
/*0060*/ IMAD.HI.X R7, R0, R9, c[0x0][0x144]
/*0068*/ IMAD R4.CC, R0, R9, c[0x0][0x148]
/*0070*/ LD.E R2, [R6]
/*0078*/ IMAD.HI.X R5, R0, R9, c[0x0][0x14c]
/*0090*/ ST.E [R4], R2

(SASS不熟悉,请看SASS一文)

其中4条IMAD指令计算出读取和存储的指令地址R6:R7R4:R5。第4和6条指令执行32位的访存命令。

Vector way1: CUDA C/C++ standard headers

通过使用int2, int4, or float2

比如将int的指针d_in类型转换然后赋值。

1
2
3
reinterpret_cast<int2*>(d_in)
// simple in C99
(int2*(d_in))

但是需要注意对齐问题,比如

1
reinterpret_cast<int2*>(d_in+1)

这样是非法的。

Vector way2: structures

通过使用对齐的结构体来实现同样的目的。

1
2
3
4
struct Foo {int a, int b, double c}; // 16 bytes in size
Foo *x, *y;

x[i]=y[i];

实际修改LD.E.64

执行for循环次数减半,注意边界处理。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
__global__ void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
}

// in only one thread, process final element (if there is one)
if (idx==N/2 && N%2==1)
d_out[N-1] = d_in[N-1];
}

void device_copy_vector2(int* d_in, int* d_out, int n) {
threads = 128;
blocks = min((N/2 + threads-1) / threads, MAX_BLOCKS);

device_copy_vector2_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

对应汇编可以看出

1
2
3
4
5
6
/*0088*/                IMAD R10.CC, R3, R5, c[0x0][0x140]              
/*0090*/ IMAD.HI.X R11, R3, R5, c[0x0][0x144]
/*0098*/ IMAD R8.CC, R3, R5, c[0x0][0x148]
/*00a0*/ LD.E.64 R6, [R10]
/*00a8*/ IMAD.HI.X R9, R3, R5, c[0x0][0x14c]
/*00c8*/ ST.E.64 [R8], R6

变成了LD.E.64

实际修改LD.E.128

执行for循环次数减半,注意边界处理。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
__global__ void device_copy_vector4_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for(int i = idx; i < N/4; i += blockDim.x * gridDim.x) {
reinterpret_cast<int4*>(d_out)[i] = reinterpret_cast<int4*>(d_in)[i];
}

// in only one thread, process final elements (if there are any)
int remainder = N%4;
if (idx==N/4 && remainder!=0) {
while(remainder) {
int idx = N - remainder--;
d_out[idx] = d_in[idx];
}
}
}

void device_copy_vector4(int* d_in, int* d_out, int N) {
int threads = 128;
int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);

device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

对应汇编可以看出

1
2
3
4
5
6
/*0090*/                IMAD R10.CC, R3, R13, c[0x0][0x140]              
/*0098*/ IMAD.HI.X R11, R3, R13, c[0x0][0x144]
/*00a0*/ IMAD R8.CC, R3, R13, c[0x0][0x148]
/*00a8*/ LD.E.128 R4, [R10]
/*00b0*/ IMAD.HI.X R9, R3, R13, c[0x0][0x14c]
/*00d0*/ ST.E.128 [R8], R4

变成了LD.E.128

summary

(个人感觉,提升也不大吗?也没有两倍和四倍的效果)

绝大部分情况,向量比标量好, increase bandwidth, reduce instruction count, and reduce latency. 。

但是会增加额外的寄存器(SASS里也没有看到??)和降低并行性(什么意思???)

参考文献

https://developer.nvidia.com/blog/cuda-pro-tip-increase-performance-with-vectorized-memory-access/#entry-content-comments