CMU 15213: Bomb Lab (1)

最近在自学CMU的神课15-213: Introduction to Computer Systems (ICS)。这篇文章主要是记录我完成这门课第二个实验Bomb lab的历程。Bomb lab主要考察我们对汇编语言(assembly Language)的掌握情况。

Bomb lab介绍

我们可以从课程网站下载Bomb lab所需要的材料。这份材料是针对自学的学生的。换句话说,所有自学的学生所要解的炸弹其实都是一样的。但是,CMU的学生似乎每个人解的炸弹都不一样,我想CMU的老师大概开发了什么自动炸弹生成工具,来避免学生互相抄袭吧。

解压缩bomb.tar之后,我们发现里面其实有三个文件:bomb二进制可执行文件,bomb.c源代码文件和一个README(其实啥有效信息都没有)。我粗略了阅读了一下bomb.c这个源代码文件,发现似乎就是要我们输入6行文字。1行文字对应一个炸弹,一旦某行文字输入的不符合要求,对应的炸弹就会爆炸。然而bomb.c这个文件并没有包含如何检查每一行文字是否符合要求的代码,所以我们还必须去研究bomb这个二进制执行文件。

为了完成这个lab,我们主要需要两个工具,第一个是objdump,这个工具可以帮助我们把二进制可执行文件反汇编(disassembly)成程序员可以理解的汇编代码。第二个是gdb,这个工具可以帮我们单步调试二进制可执行文件。我首先使用objdump反汇编bomb,把汇编语言代码保存到assemble.txt这个文件里去。具体命令如下:
$ objdump -d bomb > assemble.txt

Phase 1

如下图所示,bomb.c调用了phase_1这个函数来检查输入,phase_1这个函数仅有的一个输入参数就是input

    /* Hmm...  Six phases must be more secure than one phase! */
    input = read_line();             /* Get input                   */
    phase_1(input);                  /* Run the phase               */
    phase_defused();                 /* Drat!  They figured it out!
                                      * Let me know how they did it. */
    printf("Phase 1 defused. How about the next one?\n");

我们来看一下phase_1这个函数对应的汇编语言代码:

0000000000400ee0 <phase_1>:
  400ee0:	48 83 ec 08          	sub    $0x8,%rsp
  400ee4:	be 00 24 40 00       	mov    $0x402400,%esi
  400ee9:	e8 4a 04 00 00       	callq  401338 <strings_not_equal>
  400eee:	85 c0                	test   %eax,%eax
  400ef0:	74 05                	je     400ef7 <phase_1+0x17>
  400ef2:	e8 43 05 00 00       	callq  40143a <explode_bomb>
  400ef7:	48 83 c4 08          	add    $0x8,%rsp
  400efb:	c3                   	retq

我们都知道,函数的第一个参数,是通过寄存器%rdi传递进来的。所以%rdi其实保存了指向用户输入字符串input的地址。指令400ee4让%rsi保存数值0x402400,然后就调用了(callq)strings_not_equal这个函数。很明显,我们往strings_not_equal这个函数传递了两个参数,一个是保存在寄存器%rdi中的用户输入字符串地址,一个是保存在寄存器%rsi中的数值0x402400。函数strings_not_equal的返回结果保存在寄存器%rax (%eax)中。根据指令400eee,400ef0和400ef2,我们不能看出,如果返回结果为0,我们就可以避免调用explode_bomb这个炸弹爆炸的函数。换而言之,只要用户输入的字符串,和0x402400指向的字符串,完全相等,我们就可以不触发phase 1的炸弹。所以我们需要的就是用gdb打印出0x402400所指向的字符串,具体操作如下:

$ gdb ./bomb
(gdb) break *0x400ee4
Breakpoint 1 at 0x400ee4
(gdb) run
Starting program: ./bomb
Welcome to my fiendish little bomb. You have 6 phases with
which to blow yourself up. Have a nice day!
test

Breakpoint 1, 0x0000000000400ee4 in phase_1 ()
(gdb) print (char*)0x402400
$1 = 0x402400 "Border relations with Canada have never been better."

所以,Phase 1的答案就是字符串”Border relations with Canada have never been better.”。

Phase 2

我们首先来看一下phast_2这个函数对应的汇编代码:

0000000000400efc <phase_2>:
  400efc:	55                   	push   %rbp
  400efd:	53                   	push   %rbx
  400efe:	48 83 ec 28          	sub    $0x28,%rsp
  400f02:	48 89 e6             	mov    %rsp,%rsi
  400f05:	e8 52 05 00 00       	callq  40145c <read_six_numbers>
  400f0a:	83 3c 24 01          	cmpl   $0x1,(%rsp)
  400f0e:	74 20                	je     400f30 <phase_2+0x34>
  400f10:	e8 25 05 00 00       	callq  40143a <explode_bomb>
  400f15:	eb 19                	jmp    400f30 <phase_2+0x34>
  400f17:	8b 43 fc             	mov    -0x4(%rbx),%eax
  400f1a:	01 c0                	add    %eax,%eax
  400f1c:	39 03                	cmp    %eax,(%rbx)
  400f1e:	74 05                	je     400f25 <phase_2+0x29>
  400f20:	e8 15 05 00 00       	callq  40143a <explode_bomb>
  400f25:	48 83 c3 04          	add    $0x4,%rbx
  400f29:	48 39 eb             	cmp    %rbp,%rbx
  400f2c:	75 e9                	jne    400f17 <phase_2+0x1b>
  400f2e:	eb 0c                	jmp    400f3c <phase_2+0x40>
  400f30:	48 8d 5c 24 04       	lea    0x4(%rsp),%rbx
  400f35:	48 8d 6c 24 18       	lea    0x18(%rsp),%rbp
  400f3a:	eb db                	jmp    400f17 <phase_2+0x1b>
  400f3c:	48 83 c4 28          	add    $0x28,%rsp
  400f40:	5b                   	pop    %rbx
  400f41:	5d                   	pop    %rbp
  400f42:	c3                   	retq  

我们发现,指令400f05调用一个叫read_six_numbers的函数,看来phase 2的输入应该是6个数字,我们需要深入read_six_numbers这个函数看看6个数字储存在哪里。在研究read_six_numbers的汇编代码前,我们首先看看这个函数的输入参数(arguments)。众所周知,函数的第一个参数应该被保存在寄存器%rdi里面。从指令400efc400f02,我们并没有看到涉及到%rdi的操作。因此,我们可以判断出,read_six_numbers的第一个输入参数其实也是phase_2的第一个输入参数,即输入的字符串。指令400f02显示,函数read_six_numbers的第二个输入参数其实是栈指针%rsp,结合read_six_numbers要读6个数字的任务,我猜测第二个参数其实是一个数组地址,指令400efe是让数组预留出足够的空间(40个字节)来存储6个数字。

在理解了read_six_numbers的输入参数后,我们来看看它的汇编代码吧。

000000000040145c <read_six_numbers>:
  40145c:	48 83 ec 18          	sub    $0x18,%rsp
  401460:	48 89 f2             	mov    %rsi,%rdx
  401463:	48 8d 4e 04          	lea    0x4(%rsi),%rcx
  401467:	48 8d 46 14          	lea    0x14(%rsi),%rax
  40146b:	48 89 44 24 08       	mov    %rax,0x8(%rsp)
  401470:	48 8d 46 10          	lea    0x10(%rsi),%rax
  401474:	48 89 04 24          	mov    %rax,(%rsp)
  401478:	4c 8d 4e 0c          	lea    0xc(%rsi),%r9
  40147c:	4c 8d 46 08          	lea    0x8(%rsi),%r8
  401480:	be c3 25 40 00       	mov    $0x4025c3,%esi
  401485:	b8 00 00 00 00       	mov    $0x0,%eax
  40148a:	e8 61 f7 ff ff       	callq  400bf0 <__isoc99_sscanf@plt>
  40148f:	83 f8 05             	cmp    $0x5,%eax
  401492:	7f 05                	jg     401499 <read_six_numbers+0x3d>
  401494:	e8 a1 ff ff ff       	callq  40143a <explode_bomb>
  401499:	48 83 c4 18          	add    $0x18,%rsp
  40149d:	c3                   	retq  

在指令40148a这一步,我们发现函数read_six_numbers其实是调用了函数sscanf来读取6个数字的。我猜测具体的调用形式应该是sscanf(line, "%d %d %d %d %d %d", &array[0], &array[1], &array[2], &array[3], &array[4], &array[5])sscanf一共需要8个输入参数。明白了这个,我们就比较容易理解指令401460401480发生了啥,这些指令本质是设置sscanf的输入参数,8个参数的具体设置如下:

参数index 存储位置 存储内容 指令
1 寄存器 %rdi input string
2 寄存器 %rsi input format 0x4025c3 401480
3 寄存器 %rdx 数组地址 401460
4 寄存器 %rcx 数组地址 + 4 401463
5 寄存器 %r8 数组地址 + 8 40147c
6 寄存器 %r9 数组地址 + 12 401478
7 内存 (%rsp) 数组地址 + 16 401470,401474
8 内存 8(%rsp) 数组地址 + 20 401467,40146b

从指令40148f开始,我们开始处理函数sscanf的返回结果,如果返回结果(成功读取的参数数目)大于5,则函数read_six_numbers正常返回,否则则触发炸弹。

我们现在回到phase_2这个函数,我们注意到在调用read_six_numbers前,phase_2做了一个操作:mov %rsp, %rsi。所以,我们读到的六个数字,其实是存在内存地址:%rsp, %rsp + 4, %rsp + 8, %rsp + 12, %rsp + 16, %rsp + 20

根据指令400f0a,400f0e400f10,我们发现,如果第一个数字不是1的话,那么phase_2就会触发炸弹,因此,第一个数字一定要是1。

现在我们跳转到指令400f30, 接下来,程序让%rbx = %rsp + 4, %rbp = %rsp + 24。接下来我们跳回到指令400f17,我们让%eax = (%rbx - 4) = (%rsp), 这样%eax其实就等于第一个数字了。然后在400f1a指令里,我们把%eax的数值翻倍,这时候%eax的数值等于第一个数字的两倍。在指令400f1c里面,我们把%eax(%rbx)做比较,其实是拿第一个数字的两倍和第二个数字做比较,如果不相等,炸弹爆炸(指令400f20),这就说明第二个数字必须等于2。

如果第一个数字的两倍等于第二个数字,我们会让%rbx = %rbx + 4 = %rsp + 8,然后如果%rbx不等于%rsp + 24,我们就跳回到400f17进行下一轮比较。我们现在可以发现规律了,这个比较总是把读进来的第N-1个数字乘2,然后和第N个数字比较大小。于是我们可以推测出,这六个数字,是一个等比数列。

所以最终的答案是1 2 4 8 16 32。

一个volatile引发的CUDA程序的血案

最近一直在学习一些优化CUDA程序的技术。在阅读了参考资料1, 我了解了如何一步一步的优化一个reduce函数。其中,让我感觉最精彩的一步优化其实是unroll the last warp。因为GPU的Streaming Multiprocessor一般采用SIMT (Single-Instruction, Multiple-Thread)架构,所以一个warp里的线程理论上说在每一个指令上都是自然保持同步的,无须额外的_syncthreads()。这就是所谓的warp synchronicity

具体来说,我们的reduce的kernel函数代码如下所示,它已经经历了avoid thread divergence, avoid shared memory bank conflicts还有first add during load三步优化(关于这三步优化具体细节还请读者阅读参考资料1)。

__global__ void reduce_kernel(int *d_out, int *d_in)
{
    extern __shared__ int s_data[];

    // thread ID inside the block
    unsigned int tid = threadIdx.x;
    // global ID across all blocks
    unsigned int gid = blockIdx.x * blockDim.x * 2 + threadIdx.x;

    // perform first level of reduction,
    // reading from global memory, writing to shared memory
    s_data[tid] = d_in[gid] + d_in[gid + blockDim.x];
    // Ensure all elements have been copied into shared memory
    __syncthreads();    

    // s = blockDim.x / 2, ....., 8, 4, 2, 1
    for (unsigned int s = (blockDim.x >> 1); s >= 1; s >>= 1) {
        if (tid < s) {
            s_data[tid] += s_data[tid + s];
        }
        // Ensure all threads in the block finish add in this round
        __syncthreads();
    }

    if (tid == 0) {
        d_out[blockIdx.x] = s_data[0];
    }    
}

对于上述代码,我们发现,当for循环里面的s等于32的时候,整个block里面活跃的线程其实只剩一个warp了(warp的大小一般为32)。在这个情况下,根据warp synchronicity,我们for循环里面的_syncthreads()其实就没有多少意义了。另外,我们也不需要继续减少s的数值以及if (tid < s)这样的判断了。这些代码其实只会我们增加GPU的工作量。根据参考资料1的第21和22页,我们可以用如下代码来展开(unroll)最后的6次迭代。

    // s = blockDim.x / 2, ....., 128, 64
    for (unsigned int s = (blockDim.x >> 1); s > 32; s >>= 1) {
        if (tid < s) {
            s_data[tid] += s_data[tid + s];
        }
        // Ensure all threads in the block finish add in this round
        __syncthreads();
    }

    if (tid < 32) {
        s_data[tid] += s_data[tid + 32];
        s_data[tid] += s_data[tid + 16];
        s_data[tid] += s_data[tid + 8];
        s_data[tid] += s_data[tid + 4];
        s_data[tid] += s_data[tid + 2];
        s_data[tid] += s_data[tid + 1];
    }

然而,当我在GPU上运行上述代码的时候,我发现reduce kernel函数总是得到错误的结果。奇怪的是,当我在展开的代码里面插入_syncthreads()之后(如下所示),我就可以得到正确的结果了。这些奇怪的现象不禁让我质疑自己对warp synchronicity的理解是否正确了。

    if (tid < 32) {
        s_data[tid] += s_data[tid + 32];
        __syncthreads();
        s_data[tid] += s_data[tid + 16];
        __syncthreads();
        s_data[tid] += s_data[tid + 8];
        __syncthreads();
        s_data[tid] += s_data[tid + 4];
        __syncthreads();
        s_data[tid] += s_data[tid + 2];
        __syncthreads();
        s_data[tid] += s_data[tid + 1];
        __syncthreads();
    }

在经过一番疯狂的谷歌后,我发现其实很多人都跟我遇到了相似的问题。在一个stackoverflow问题(参考资料3)下,有人说这是因为没有加volatile关键字的原因,并提供了一个更新的优化reduce的参考资料。根据参考资料2的第22页,正确的代码展开方式应该是如下这样的。经过我的测试,这段代码确实可以得到正确的reduce结果。

__device__ void warpReduce(volatile int* s_data, int tid) 
{
    s_data[tid] += s_data[tid + 32];
    s_data[tid] += s_data[tid + 16];
    s_data[tid] += s_data[tid + 8];
    s_data[tid] += s_data[tid + 4];
    s_data[tid] += s_data[tid + 2];
    s_data[tid] += s_data[tid + 1];
}

__global__ void reduce_kernel(int *d_out, int *d_in)
{
    // later ....
    if (tid < 32) {
        warpReduce(s_data, tid);
    }

于是,问题就来了,究竟volatile发挥了什么作用,让我们的程序运行正常了呢?我们都知道,现代的编译器都会对代码做很多优化,一种最常用的优化方法是将某个内存变量缓存到寄存器,需要读写变量的时候直接访问寄存器而不是内存。volatile的中文意思是“易变的,不稳定的”,对于用volatile修饰的变量,编译器对访问该变量的代码不再优化,总是从它所在的内存读取数据

对于上面这个例子来说,如果我们不使用volatile,对于一个线程来说(假设线程ID就是tid),它的s_data[tid]可能会被缓存在寄存器里面,且在某个时刻寄存器和shared memory里面s_data[tid]的数值还是不同的。当另外一个线程读取s_data[tid]做加法的时候,也许直接就从shared memory里面读取了旧的数值,从而导致了错误的结果。

根据参考资料4,volatile关键字确保了编译器不会把s_data[tid]的数值缓存到寄存器里面,而是每次去从GPU的shared memory里面去读取s_data[tid]

写完这么多,其实我心里还有一个疑问,为什么我之前那种加入_syncthreads()的方法可以解决这个问题?要解答这个问题,估计只能对比有无_syncthreads()后程序的编译结果了。

References
1. http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf
2. https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf
3. https://stackoverflow.com/questions/45227732/cuda-threads-in-a-warp-appear-to-be-not-in-synchronization
4. https://stackoverflow.com/questions/21205471/cuda-in-warp-reduction-and-volatile-keyword?noredirect=1&lq=1

Local memory in CUDA

QQ图片20171128201847

I used to think local memory is one of the fastest memories in CUDA, like L1 cache in CPU. However, recently, after I read the following references, I realized I misunderstood this. As shown in the above figure, “Local memory” in CUDA is actually global memory (and should really be called “thread-local global memory”) with interleaved addressing . Like global memory, local memory is also off-chip. As a result, local memory accesses have same high latency and low bandwidth as global memory accesses.

According to CUDA C Programming Guide 5.3.2, local memory accesses only occur for some automatic variables as mentioned in Variable Type Qualifiers. Automatic variables that the compiler is likely to place in local memory are:

  • Arrays for which it cannot determine that they are indexed with constant quantities,
  • Large structures or arrays that would consume too much register space,
  • Any variable if the kernel uses more registers than available (this is also known as register spilling).

References
http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#maximize-memory-throughput
https://stackoverflow.com/questions/10297067/in-a-cuda-kernel-how-do-i-store-an-array-in-local-thread-memory
https://graphics.cg.uni-saarland.de/fileadmin/cguds/courses/ss14/pp_cuda/slides/02_-_CUDA_Memory_Hierarchy.pdf

关于同步/异步 VS 阻塞/非阻塞的一点体会

长期以来,同步(synchronous),异步(asynchronous),阻塞(blocking)和非阻塞(non-blocking)这几个概念一直困扰着我。我以前一直简单地以为同步等于阻塞,异步等于非阻塞。直到最近读完UNIX Network Programming (Volume 1)的第6章之后,并查阅了大量网上的资料之后,我才对这个问题有了一个比较清楚的认识。

在讨论这四个概念的区别之前,我们首先要确定一下我们讨论的上下文(context),那就是Linux的network IO。对于一个网络IO来说(以read作为例子),其执行过程通常可以分为两个阶段。第一阶段,等待数据从网络中到达,并被拷贝到内核中某个缓冲区(Waiting for the data to be ready)。第二阶段,把数据从内核态的缓冲区拷贝到用户态的应用进程缓冲区来(Copying the data from the kernel to the process)。

1. 阻塞/非阻塞 IO
根据一个高票的知乎回答,阻塞和非阻塞关注的是程序在等待调用结果(消息,返回值)时的状态。阻塞调用是指调用结果返回之前,当前线程会被挂起。调用线程只有在得到结果之后才会返回。非阻塞调用指在不能立刻得到结果之前,该调用不会阻塞当前线程。

在Linux下,一个socket的文件描述符(file descriptor)默认就是阻塞模式的。在这种模式下,即便这个socket压根没有收到任何数据,我们的read调用也会一直阻塞在那里,无法返回,直到有数据到达为止。

如果我们把这个socket的文件描述符用fcntl设置为非阻塞的。在这种模式下,如果这个socket没有收到任何数据,我们的read调用会立刻返回一个错误。这个时候,我们的程序就知道目前没法从这个socket里读到数据了,索性去干点别的事情,过段时间再调用read。当一个应用进程对一个非阻塞的文件描述符循环调用read时,我们称之为轮询(polling)。

现在再让我们回想网络IO的两个阶段,阻塞和非阻塞主要区别其实是在第一阶段等待数据的时候但是在第二阶段,阻塞和非阻塞其实是没有区别的。程序必须等待内核把收到的数据复制到进程缓冲区来。换句话说,非阻塞也不是真的一点都不”阻塞”,只是在不能立刻得到结果的时候不会傻乎乎地等在那里而已。

2. 同步/异步 IO
对于这两个东西,POSIX其实是有官方定义的。
A synchronous I/O operation causes the requesting process to be blocked until that I/O operation completes;
An asynchronous I/O operation does not cause the requesting process to be blocked;
根据这个定义,不管是blocking IO还是non-blocking IO,其实都是synchronous IO。因为它们一定都会阻塞在第二阶段拷贝数据那里。

3. IO复用
这时候有些人会问了,IO复用(multiplexing)算是什么类型的IO呀。不同于这篇IBM的文章的观点,我个人认为,IO复用是阻塞同步IO

跟传统的阻塞IO不同,IO复用可以阻塞在多个socket文件描述符上。当其中任何一个socket有数据可读的时候(或者超时),IO复用的函数(select, poll,epoll)才会返回。然后进程可以逐一处理可读的socket文件描述符。

4. 真正的异步IO
对此我只知道Linux AIO。可惜本人才疏学浅,并没有Linux AIO的实际开发经验,在此就不详细介绍了,有兴趣的读者可以自己去尝试。

参考资料:
https://www.zhihu.com/question/19732473/answer/20851256
http://lifeofzjs.com/blog/2014/03/29/sycron-vs-block/
http://blog.csdn.net/historyasamirror/article/details/5778378

How to set processor affinity on Linux using taskset

Today’s computers typically adopt multiple CPU cores. A process/thread can be executed on any of those CPU cores (determined by OS scheduling). Hence, performance optimization in such multi-core architecture is crucial.

Processor affinity, or CPU pinning is an important technique for above purpose. It enables the binding and unbinding of a process or a thread to a CPU or a range of CPUs, so that the process or thread will execute only on the designated CPU or CPUs rather than any CPU.

Processor affinity takes advantage of the fact that remnants of a process that was run on a given processor may remain in that processor’s state (for example, data in the cache memory) after another process was run on that processor. Therefore, it can effectively reduce cache miss problems. Also, when two processes communicate via shared memory intensively, scheduling both processes on the cores in the same NUMA domain would speed up their performance.

Now, let’s see how to set processor affinity on Linux. To this end, there are several approaches. To set the CPU affinity of a process, you can use taskset program and sched_setaffinity system call. To set the CPU affinity of a thread, you can use pthread_setaffinity_np and pthread_attr_setaffinity_np. In this article, I want to introduce the usage of taskset.

taskset is used to set or retrieve the CPU affinity of a running process given its PID or to launch a new COMMAND with a given CPU affinity.

Read the CPU Affinity of a Running Process

To retrieve the CPU affinity of a process, you can use the following command.
taskset -p [PID]

For example, to check the CPU affinity of a process with PID 1141.
$ taskset -p 1141
pid 1141's current affinity mask: ffffffff

The return value ffffffff is essentially a hexadecimal bitmask, corresponding to 1111 1111 1111 1111 1111 1111 1111 1111. Each bit in the bitmask corresponds to a CPU core. The bit value 1 means that the process can be executed on the corresponding CPU core. Therefore, in above example, pid 1141 can be executed on CPU core 0-31.

You may think that bitmask is a little hard to understand. Don’t worry. taskset can also show CPU affinity as a list of processors instead of a bitmask using “-c” option. An example is given as follows.
$ taskset -cp 1141
pid 1141's current affinity list: 0-31

Pin a Running Process to Particular CPU Core(s)

You can also use taskset to pin a running process to particular CPU core(s). The command formats are given as follows.
taskset -p [CORE-LIST] [PID]
taskset -cp [CORE-LIST] [PID]

For example, to assign process 1141 to CPU core 0 and 4:
$ taskset -p 0x11 1141

Launch a Program on Specific CPU Cores

tasklet also allows us to launch a program running on specific CPU cores. The command is given as follows.
taskset [COREMASK] [EXECUTABLE]

For example, to launch a ping program (destination 8.8.8.8) on a CPU core 0, use the following command.
$ taskset 0x1 ping 8.8.8.8

References
https://en.wikipedia.org/wiki/Processor_affinity
http://xmodulo.com/run-program-process-specific-cpu-cores-linux.html
https://linux.die.net/man/1/taskset