Zynq 高层次综合HLS 项目开发经历(4) Vivado IP Integrator

IP Integrator这一块比较难配置,网上的教程比较乱,而且也少有针对AXIS Stream IP核专门配置的教程。所以我尝试着手把手的把配置过程在博客里展现出来。
我使用的开发板是ALinx AX7020,FPGA芯片zynq-xc7z020clg400-2。有些配置与我的板子是高度相关的,如果使用的是其他开发板请参照开发板提供商给的教程或参考资料配置。
继续阅读

Zynq 高层次综合HLS 项目开发经历(3) 代码移植及Vivado HLS的使用

首先简介一下HLS的整套开发流程。详细流程可参考[1][2]。
1、设计C/C++代码以及C/C++的TestBench
2、C Simulation,验证软件代码是否正确
3、C Synthesis,将C/C++化为Verilog/VHDL
4、C/RTL Co-Simulation,将C的TestBench与生成的RTL一起进行Simulation
5、Export RTL,导出IP核
HLS的使用本身不是一个难点,我猜Vivado HLS本身就是一个Eclipse搭起来的IDE,所以只要对Eclipse上手就行。难点在于代码移植。
继续阅读

Zynq 高层次综合HLS 项目开发经历(2) LZ77 x86软件实现

LZ77的x86软件实现repo地址可参见[1]。

要想实现软件,肯定需要把LZ77的算法读懂。慢慢啃就是了。我查询的参考资料可参见[2]。我主要参考了其中的“短语字典的维护”、“压缩和解压缩数据”、“LZ77的接口定义”三部分。而实现与定义则完全由自己制定。
继续阅读

Zynq 高层次综合HLS 项目开发经历(1) 概要

今年寒假我有幸参加了Xilinx的Winter Camp,在Camp上,讲者对HLS整个开发流程做了简单的介绍。为了巩固同学们的学习,主办方特意要求我们在冬令营结束之后自己用HLS做一个项目。本着拿一个证书和提高一下自己知识水平的想法,我也自己给自己定一个选题参加了这个项目。

原本是想要开发一个MD5的硬件实现的,后来在Github上面搜了一下,发现已经有人开发过HLS的MD5算法了。不得已只能放弃。后来经Foobar院猴哥推荐LZ77压缩算法,我发现也比较合适。于是就决定在硬件上实现一个LZ77压缩算法。值得注意的是,LZ77压缩算法只是一个大概的实现模板,具体怎么实现可以根据开发者的想法任意指定规范。所以这个压缩算法中的所有规范都是我自己指定的。因此和其他的压缩软件是不通用的。

现在硬软件设计基本上都已经做得差不多了,这一路走来坑真的不少。所以我决定写一个系列来记录一下自己遇到的坑,也许会对后来的同学有帮助。如果我的文章或者代码中有什么错误,也欢迎在博客里评论、Github上开issue留言给我。

首先,大概说一下高层次综合HLS是个什么东西。通俗易懂的讲,HLS就是可以把C/C++语言转换为Verilog/VHDL/System C的一个处理过程/软件。这样对于软件开发者来说,能够比较轻松简单的开发硬件设计。当然,其中也有很多与全软件开发不完全相同的部分。这些都需要我们另行学习。
然后,简介Zynq的两大部分:PS (Processing System) 指处理系统,它是在FPGA上提前实现好的硬核ARM处理器。PL (Programming Logic) 指FPGA可编程逻辑部分,它就是真正可操纵的硬件部分。PS和PL两部分可以通过预先定义好的总线接口相互连接。

我的开发流程是这样的:先在x86体系结构的软件上实现LZ77压缩算法,经过测试通过之后,转而开始开发硬件部分。首先使用Vivado HLS软件,移植软件C代码到硬件C代码,使用Directive规定接口和优化,生成IP核;打开Vivado,使用其中的IP Integrator,将Zynq处理系统 (Zynq Processing System) 的IP核以及其他各配套IP核与HLS开发的IP核连接在一起。综合实现生成比特流,导出硬件描述文件和比特流(到此为止的代码运行在PL端);切换到Vivado SDK中进行开发,使用硬件描述文件创建Standalone软件项目,编写在PS端运行的程序以及PL与PS的交互。到这里整体的开发流程就算完毕了。

在后面的文章里,我会依次的认真介绍一下各个部分的开发流程。

LZ77硬件实现项目:https://github.com/bjrjk/HLS-LZ77
LZ77软件实现项目:https://github.com/bjrjk/LZ77-Demo
Xilinx HLS设计流程课程:https://github.com/xupgit/High-Level-Synthesis-Flow-on-Zynq-using-Vivado-HLS

CUDA上的自旋锁(SpinLock)实现问题

最近在研究CUDA。其中有一个需求是要对一段代码加锁。CUDA里只提供了现成的对单个变量进行操作的原子操作库[1],而没有提供对一段代码加锁的机制,于是我就想怎么改造一下这个原子库。

一开始自己写了一段代码,差不多是这个样子的[2]。

1
2
3
4
5
6
__device__ inline void lock(int* p_mutex ){
    while(atomicCAS(p_mutex, 0, 1) != 0); // compare and swap
}
__device__ inline void unlock(int* p_mutex ){
    atomicExch(p_mutex, 0);// exchange
}

调用了之后发现从加锁的循环里压根就跑不出来。自己想了半天死活也搞不懂。然后从StackOverFlow上查到资料。发现这个和CUDA的Warp有关[3]。每个Warp由32个线程组成,它们必须同时执行相同的指令。如果用上述代码实现自旋锁。假定Warp中有一个Thread获得了锁,需要向下执行才能够释放锁。但是其他Thread却没得到锁,因此整个Warp就不能向下执行,始终卡在while循环里形成死锁。这种死锁叫SIMD死锁(SIMD DeadLock)。

之后翻了一大堆文献和网站,里面成堆的人都是实现同上,没有一个能用的。最后翻到了这篇文章[4],解决了问题。整理出的自旋锁的模板如下:

继续阅读

第二届BJDCTF DreamerJack题目出题人题解

下面是DreamerJack同学的出题人题解。

Programming: Strenuous_Huffman

本题的题解需要占个坑。主要是因为吧。。。题目描述里说的是真的。。这个是本人的数据结构作业,老师还没审查,直接把整套源码发出来我作业就不用交了,所以先占个坑。后续整套源码会发布到Github上。不过可以先说下本题的难点在哪里。其实这个题只是我数据结构作业题的一部分,只涉及解压缩。难度相对于压缩部分较小。正统解法是BitMap+Trie Tree解压。如果暴力一点也可以不用Trie。难度就在于BitMap的编写上。因为这个压缩的时候都是按位压缩的,最低的操作单位是比特。而咱们C/C++操作的最低单位都是按字节。所以要求代码编写者对位运算有比较深刻的认识才能够写得出来。

主要的解题思路如下。编写一个BitMap类,要支持从文件加载内容到内部数组和将内部数组存储到文件中,还要支持按位对数组的访问。然后按照压缩编码和原始编码的对应关系,逐个的将压缩后的比特信息翻译为原始信息并保存。

整套源码发布遥遥无期(看老师啥时候查完作业),想要源码的同学可持续关注本文章。
PS:要提醒的一点是:我自己开发的压缩算法经过测试,确实压缩效果较差。本源码仅供教学用途使用。不要异想天开真的拿这玩意去压缩东西。。。

2020.05.09:此代码已开源:https://github.com/bjrjk/JackZIPPacker

Reverse: 8086 ASM

为什么想要出这道题目呢?我校最近正在学习8086汇编,因此出这道题目。
源码如下:

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
31
32
33
34
DATAS SEGMENT
    str1 DB 5dH,55H,5bH,64H,75H,7eH,7cH,74H
    DB 40H,7bH,7aH,40H,77H,6aH,2eH,7dH,2eH
    DB 7eH,71H,40H,67H,6aH,7aH,7bH,7aH,40H
    DB 77H,7aH,71H,57H,7eH,2fH,62H,59
DATAS ENDS

STACKS SEGMENT
STACKS ENDS

CODES SEGMENT
    ASSUME CS:CODES,DS:DATAS,SS:STACKS
loc:
    jmp loc
    mov CX,34
    lea BX,str1
   
lop:
    mov DI,CX
    dec DI
    xor BYTE PTR [BX][DI],31
    loop lop
    lea DX,str1
    mov AH,09H
    int 21H
    ret
START:
    MOV AX,DATAS
    MOV DS,AX
    call loc
    MOV AH,4CH
    INT 21H
CODES ENDS
    END START

本题其实很简单,str1存放的是一个MSDOS格式的以$结尾的经过异或加密的字符串。汇编代码就是把这玩意解密了一下然后syscall输出。jmp loc的死循环就是用来干扰人的。如果想要不劳而获放到DOS里直接跑会死循环。放到IDA里看的话,IDA会把lop下面那段代码当作是数据。需要把他们强转一下成代码才行。本题也可以用调试器手动改EIP执行。顺便宣传一下自己做的8086调试器套件:DOSBox-8086Assembly
PS:我汇编学的还行是真的^v^,死循环是我故意写的^v^

Misc: A Beautiful Picture

很简单的一个png隐写。说实话,图画的不怎么好看。。一看这图片长1000,高900,有没有感觉有点不对劲呢?用010HexEditor,直接把png图片高度改成1000,flag就出来了。。

更新:题目复现地址: https://buuoj.cn/
20200510更新:本人CTFd已下线。
题目下载链接:https://renjikai.com/wp-content/uploads/2020/03/BJD2_DreamerJack.zip

洛谷 P3369 【模板】普通平衡树

手写Splay成功。庆祝能写出来的第一个BBST!
代码如下:

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
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
#include<iostream>
#include<cassert>
using namespace std;
class Splay {
private:
    struct Node {
        int v; //值
        Node* f, * l, * r; //父,左孩子,右孩子
        int cnts; //本子树下包含自身的所有**元素**数目
        int rep; //该值重复了多少次
        Node(int v, Node* f) :v(v), f(f), l(NULL), r(NULL), cnts(1), rep(1) {}
    };
    Node* _root;
#define root _root->r
    void update(Node* p) { //更新cnts值
        p->cnts = p->rep;
        if (p->l)p->cnts += p->l->cnts;
        if (p->r)p->cnts += p->r->cnts;
    }
    bool identify(Node* s, Node* f) { //确定本节点是左孩子还是右孩子,左孩子false,右孩子true
        return !f || s == f->r;
    }
    void connect(Node* s, Node* f, bool r) { //儿子地址,父亲地址,连接到左儿子还是右儿子
        if (f)(r ? f->r : f->l) = s; //父亲更新新儿子
        if (s)s->f = f; //儿子更新新父亲
    }
    void rotate(Node* s) { //儿子节点的旋转
        Node* f = s->f, * gf = s->f->f;
        if (!identify(s, f)) { //左孩子
            connect(s->r, f, false);
            connect(f, s, true);
            connect(s, gf, identify(f, gf));
        }
        else { //右孩子
            connect(s->l, f, true);
            connect(f, s, false);
            connect(s, gf, identify(f, gf));
        }
        update(f);
        update(s);
    }
    void splay(Node* s, Node* e) { //伸展操作,将节点s旋转到节点e所在的位置
        e = e->f;
        while (s->f != e) {
            Node* f = s->f;
            if (f->f == e)rotate(s); //s是e的直系儿子,只需做单旋
            else if (identify(f, f->f) == identify(s, f)) { //Zig-Zig或Zag-Zag,需要先旋父亲节点,再旋儿子节点
                rotate(f);
                rotate(s);
            }
            else { //Zig-Zag或Zag-Zig
                rotate(s);
                rotate(s);
            }
        }
    }

public:
    Splay() {
        _root = new Node(0, NULL);
    }
    ~Splay() {
        delete _root;
    }
    Node* find(int v) {
        Node* cur = root;
        if (!cur)return NULL;
        while (1) {
            if (cur->v == v) break;
            Node* next = v < cur->v ? cur->l : cur->r;
            if (!next)break;
            cur = next;
        }
        splay(cur, root);
        root = cur;
        if (cur->v == v)return cur;
        else return NULL;
    }
    void del(int v) {
        Node* cur = find(v);
        if (!cur)return;
        if (cur->rep > 1) { //节点个数出现多于1次
            cur->rep--;
            cur->cnts--;
            return;
        }
        if (!cur->l && !cur->r) { //删除最后一个仅剩的节点
            root = NULL;
        }
        else if (!cur->l) { //无左子树时直接把右子树拼到根
            root = cur->r;
            root->f = _root;
        }
        else { //有左子树时,把左子树的最大值旋到根的左子,将根的右子放到根的左子的右子,删根后补左子树
            Node* l = cur->l;
            while (l->r)l = l->r;
            splay(l, cur->l);
            Node* r = cur->r;
            connect(r, l, true);
            root = l;
            root->f = _root;
            update(l);
        }
        delete cur;
    }
    void insert(int v) {
        Node* cur = find(v);
        if (!root) { //特判空树
            root = new Node(v, _root);
            return;
        }
        if (cur && cur->v == v) { //元素存在,直接把次数+1
            cur->rep++;
            cur->cnts++;
            return;
        }
        Node* newNode = new Node(v, _root);
        if (root->v < v) { //将v接入右侧
            connect(root, newNode, false);
            connect(root->r, newNode, true);
            root->r = NULL;
        }
        else { //将v接入左侧
            connect(root, newNode, true);
            connect(root->l, newNode, false);
            root->l = NULL;
        }
        update(root);
        update(newNode);
        root = newNode;
        newNode->f = _root;
    }
    int rank(int v) {
        Node* cur = find(v);
        if (!cur)return -1;
        int lCnts = cur->l ? cur->l->cnts : 0;
        return lCnts + 1;
    }
    int atrank(int rank) {
        Node* cur = root;
        while (cur) {
            int lCnts = cur->l ? cur->l->cnts : 0;
            if (lCnts < rank && rank <= lCnts + cur->rep) {
                splay(cur, root);
                return cur->v;
            }
            if (rank <= lCnts)cur = cur->l;
            else {
                rank -= lCnts + cur->rep;
                cur = cur->r;
            }
        }
        return -1;
    }
    int upper(int v) {
        Node* cur = find(v);
        int lCnts = root->l ? root->l->cnts : 0;
        if (root->v <= v)return atrank(lCnts + root->rep + 1);
        return root->v;
    }
    int lower(int v) {
        Node* cur = find(v);
        int lCnts = root->l ? root->l->cnts : 0;
        if (root->v >= v)return atrank(lCnts);
        return root->v;
    }
#undef root
};
int main() {
    Splay splay;
    int n;
    cin >> n;
    while (n--) {
        int op1, op2;
        cin >> op1 >> op2;
        if (op1 == 1) {
            splay.insert(op2);
        }
        else if (op1 == 2) {
            splay.del(op2);
        }
        else if (op1 == 3) {
            cout << splay.rank(op2) << endl;
        }
        else if (op1 == 4) {
            cout << splay.atrank(op2) << endl;
        }
        else if (op1 == 5) {
            cout << splay.lower(op2) << endl;
        }
        else if (op1 == 6) {
            cout << splay.upper(op2) << endl;
        }
    }
}

参考资料:
https://www.luogu.com.cn/blog/user19027/solution-p3369
非常感谢rentenglong同学

CTF RSA 高精度整数开根 解一元二次方程

做BJDCTF第二届出的RSA题目,需要解高精度整数的一元二次方程,参考着网上写了个Python脚本。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
def HPSqrt(a,b): #高精度开整数根,a为开几次根,b为要开根的数
    l=0
    r=1
    while(r**a<=b):
        l=r
        r=r*2  
    while(l+1<r):  
        mid=(l+r)//2
        if (mid**a<=b): l=mid
        else: r=mid
    if (l**a<=b): return l
    else: return r

def HPSolve12(a,b,c): #高精度解整数一元二次方程,解RSA用
    key=HPSqrt(2,b*b-4*a*c)
    return ((-b+key)//2//a,(-b-key)//2//a)

a=1
b=-2
c=1
print(HPSolve12(a,b,c)) #将方程整理成ax^2+bx+c=0的形式代入即可

参考资料:高精度Python开根 https://www.luogu.com.cn/blog/wjy666/solution-p2293

Dockerfile编写使用cron后出现无法执行定时任务的解决方法

最近给自己的项目编写一个Dockerfile,这个项目需要用到Cron定时任务。在Dockerfile里apt安装cron后不论如何使用RUN命令启动cron都无效(包括使用service和直接执行cron命令)。后来将所有任务集成在了一个sh脚本里解决了问题。举例,我要运行的任务是apache。那么我可以写一个脚本startup.sh,把Dockerfile的ENTRYPOINT设为这个脚本。

1
2
3
#!/bin/sh
cron
apache2-foreground

原因是什么呢?

方法1和方法2不能成功,是因为docker只是一个进程隔离的沙箱环境,并不是真正的虚拟机。而service xxx start 和systemctl start xxx 分别是upstart和systemd这两个/sbin/init进程的替代者的服务管理命令。而upstart和systemd都要求系统必须是物理机或虚拟机,并不支持作为container的init进程。方法3存在问题是因为,在正常的系统中,init进程永远占用PID=1的位置,回收僵尸进程、处理未处理的信号等都是由init进程帮我们完成的,一个子进程如果失去了父进程,也会由init进程接管。但是在container中,init进程并不存在,PID=1的进程是我们在Dockerfile中定义的Entrypoint或最后一个CMD指定的命令。

参考资料:https://www.asuri.org/2018/08/25/run-multi-service-in-one-container/

而直接执行cron命令为何失败呢?我推测,Dockerfile设计RUN命令的本义是对文件系统进行一定的操作。而非监控Container在启动时需要开启什么进程。所以一定要把所有需要运行的程序都写到一个shell脚本里并且把其设为入口点才行。