cuda性能分析工具nsys使用

发表于 2024-05-19 更新于 2024-12-27 分类于 gpu
本文字数： 1.5k 阅读时长 ≈ 6 分钟

cuda性能分析工具nsys使用

安装 nsys 命令行工具

一般这个工具是随着 cuda toolkit 一起安装的。安装后可以尝试查看：

@└────> # nsys status -e
Timestamp counter supported: Yes

CPU Profiling Environment Check
Root privilege: disabled
Linux Kernel Paranoid Level = 4
Linux Distribution = Ubuntu
Linux Kernel Version = 5.15.0-105-generic: OK
Linux perf_event_open syscall available: Fail
Sampling trigger event available: Fail
Intel(c) Last Branch Record support: Not Available
CPU Profiling Environment (process-tree): Fail
CPU Profiling Environment (system-wide): Fail

See the product documentation at https://docs.nvidia.com/nsight-systems for more information,
including information on how to set the Linux Kernel Paranoid Level.

在可以使用 nsys 工具后，可以使用 nsys 来查看一些 kernel 的性能。
举个例子，我们有以下 cuda 代码：

#include <bits/stdc++.h>
#include <cuda.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <time.h>
#include <sys/time.h>

#define THREAD_PER_BLOCK 256

// baseline
__global__ void reduce0(float* d_in, float* d_out) {
    __shared__ float sdata[THREAD_PER_BLOCK];

    // each thread loads one element from global to shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    sdata[tid] = d_in[i];
    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s = 1; s < blockDim.x; s *= 2) {
        if (tid % (2 * s) == 0) {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) {
        d_out[blockIdx.x] = sdata[0];
    }
}

// bank conflict
__global__ void reduce1(float* d_in, float* d_out) {
    __shared__ float sdata[THREAD_PER_BLOCK];

    // each thread loads one element from global to shared mem
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    sdata[tid] = d_in[i];
    __syncthreads();

    // do reduction in shared mem
    for (unsigned int s = 1; s < blockDim.x; s *= 2) {
        int index = 2 * s * tid;
        if (index < blockDim.x) {
            sdata[index] += sdata[index + s];
        }
        __syncthreads();
    }

    // write result for this block to global mem
    if (tid == 0) {
        d_out[blockIdx.x] = sdata[0];
    }
}

bool check(float* out, float* res, int n) {
    for (int i = 0; i < n; i++) {
        if (out[i] != res[i]) {
            return false;
        }
    }
    return true;
}

int main() {
    const int N = 32 * 1024 * 1024;
    float* a = (float*)malloc(N * sizeof(float));
    float* d_a;
    cudaMalloc((void**)&d_a, N * sizeof(float));

    int block_num = N / THREAD_PER_BLOCK;
    float* out = (float*)malloc((N / THREAD_PER_BLOCK) * sizeof(float));
    float* d_out;
    cudaMalloc((void**)&d_out, (N / THREAD_PER_BLOCK) * sizeof(float));
    float* res = (float*)malloc((N / THREAD_PER_BLOCK) * sizeof(float));

    for (int i = 0; i < N; i++) {
        a[i] = 1;
    }

    for (int i = 0; i < block_num; i++) {
        float cur = 0;
        for (int j = 0; j < THREAD_PER_BLOCK; j++) {
            cur += a[i * THREAD_PER_BLOCK + j];
        }
        res[i] = cur;
    }

    cudaMemcpy(d_a, a, N * sizeof(float), cudaMemcpyHostToDevice);

    dim3 Grid(N / THREAD_PER_BLOCK, 1);
    dim3 Block(THREAD_PER_BLOCK, 1);

    reduce0<<<Grid, Block>>>(d_a, d_out);
    cudaMemcpy(out, d_out, block_num * sizeof(float), cudaMemcpyDeviceToHost);
    if (check(out, res, block_num)) { 
        printf("the ans is right\n");
    } else {
        printf("the ans is wrong\n");
        for (int i = 0; i < block_num; i++) {
            printf("%lf ", out[i]);
        }
        printf("\n");
    }

    reduce1<<<Grid, Block>>>(d_a, d_out);
    cudaMemcpy(out, d_out, block_num * sizeof(float), cudaMemcpyDeviceToHost);
    if (check(out, res, block_num)) { 
        printf("the ans is right\n");
    } else {
        printf("the ans is wrong\n");
        for (int i = 0; i < block_num; i++) {
            printf("%lf ", out[i]);
        }
        printf("\n");
    }

    cudaFree(d_a);
    cudaFree(d_out);
}

以上是一个 reduce_sum 的例子，有两个性能不同的核函数 reduce0 和 reduce1。可以使用命令行工具来看这两个工具的性能。

# 一个过渡方式，沿用之前 nvprof 的用法。
@└────> # nsys nvprof ./a.out

# 后续支持的方式
@└────> # nsys profile --stats=true ./a.out
...
[4/8] Executing 'osrtsum' stats report

 Time (%)  Total Time (ns)  Num Calls    Avg (ns)       Med (ns)     Min (ns)     Max (ns)      StdDev (ns)            Name         
 --------  ---------------  ---------  -------------  -------------  ---------  -------------  -------------  ----------------------
     49.8    6,560,704,940         41  160,017,193.7  100,127,248.0      2,946  3,656,490,088  561,147,066.6  poll                  
     40.5    5,331,726,596      1,591    3,351,179.5       34,054.0      1,200    405,442,288   30,016,785.5  ioctl                 
      4.7      614,593,088         60   10,243,218.1   10,295,123.0  1,532,148     14,425,405    2,880,533.4  waitpid               
      2.6      341,072,143         60    5,684,535.7    5,470,557.0    576,832     10,285,481    1,599,344.7  fork                  
      2.4      312,589,312        113    2,766,277.1       14,265.0      4,724    267,144,344   25,109,089.4  open64                
      0.0        5,684,320        144       39,474.4       11,704.5      1,108      3,826,795      317,952.4  fopen                 
      0.0        2,034,643         38       53,543.2       10,234.5      3,945      1,236,239      198,417.6  mmap64                
      0.0          607,662         10       60,766.2       56,248.0     42,036        112,182       19,745.1  sem_timedwait         
      0.0          405,925        123        3,300.2        2,289.0      1,006         77,642        7,014.0  fclose                
      0.0          391,343          4       97,835.8       80,675.0     58,047        171,946       53,486.5  pthread_create        
      0.0          154,726         19        8,143.5        4,985.0      1,004         47,713       10,904.9  mmap                  
      0.0           82,936          1       82,936.0       82,936.0     82,936         82,936            0.0  pthread_cond_wait     
      0.0           78,328          8        9,791.0        4,877.5      2,074         40,677       12,827.6  munmap                
      0.0           71,358          7       10,194.0        9,879.0      3,968         14,551        3,520.2  open                  
      0.0           51,719          3       17,239.7       13,793.0      3,372         34,554       15,874.2  fread                 
      0.0           42,644         29        1,470.5        1,295.0      1,000          5,372          801.5  fcntl                 
      0.0           41,205          1       41,205.0       41,205.0     41,205         41,205            0.0  fgets                 
      0.0           36,755         15        2,450.3        2,072.0      1,083          7,000        1,497.0  read                  
      0.0           33,701         12        2,808.4        2,360.0      1,384          6,014        1,224.1  write                 
      0.0           30,013          3       10,004.3       11,976.0      5,640         12,397        3,785.5  pipe2                 
      0.0           26,410          2       13,205.0       13,205.0     10,057         16,353        4,451.9  socket                
      0.0           12,546          2        6,273.0        6,273.0      5,777          6,769          701.4  fwrite                
      0.0           10,654          2        5,327.0        5,327.0      4,160          6,494        1,650.4  pthread_cond_broadcast
      0.0           10,412          1       10,412.0       10,412.0     10,412         10,412            0.0  pthread_mutex_trylock 
      0.0            9,158          1        9,158.0        9,158.0      9,158          9,158            0.0  connect               
      0.0            5,650          1        5,650.0        5,650.0      5,650          5,650            0.0  bind                  
      0.0            3,085          1        3,085.0        3,085.0      3,085          3,085            0.0  listen                

[5/8] Executing 'cudaapisum' stats report

 Time (%)  Total Time (ns)  Num Calls    Avg (ns)       Med (ns)      Min (ns)      Max (ns)      StdDev (ns)            Name         
 --------  ---------------  ---------  -------------  -------------  -----------  -------------  -------------  ----------------------
     77.1    1,961,601,754          2  980,800,877.0  980,800,877.0  674,626,715  1,286,975,039  432,995,652.3  cudaMalloc            
     22.2      564,147,432          2  282,073,716.0  282,073,716.0  275,230,367    288,917,065    9,677,957.0  cudaFree              
      0.7       17,105,648          3    5,701,882.7      657,298.0      434,348     16,014,002    8,931,253.0  cudaMemcpy            
      0.0          415,480          2      207,740.0      207,740.0       42,497        372,983      233,688.9  cudaLaunchKernel      
      0.0            1,368          1        1,368.0        1,368.0        1,368          1,368            0.0  cuModuleGetLoadingMode

[6/8] Executing 'gpukernsum' stats report

 Time (%)  Total Time (ns)  Instances  Avg (ns)   Med (ns)   Min (ns)  Max (ns)  StdDev (ns)      GridXYZ          BlockXYZ               Name           
 --------  ---------------  ---------  ---------  ---------  --------  --------  -----------  ----------------  --------------  -------------------------
     62.5          556,575          1  556,575.0  556,575.0   556,575   556,575          0.0  131072    1    1   256    1    1  reduce0(float *, float *)
     37.5          334,111          1  334,111.0  334,111.0   334,111   334,111          0.0  131072    1    1   256    1    1  reduce1(float *, float *)

[7/8] Executing 'gpumemtimesum' stats report

 Time (%)  Total Time (ns)  Count    Avg (ns)      Med (ns)     Min (ns)    Max (ns)   StdDev (ns)      Operation     
 --------  ---------------  -----  ------------  ------------  ----------  ----------  -----------  ------------------
     99.7       15,820,690      1  15,820,690.0  15,820,690.0  15,820,690  15,820,690          0.0  [CUDA memcpy HtoD]
      0.3           46,176      2      23,088.0      23,088.0      23,040      23,136         67.9  [CUDA memcpy DtoH]

[8/8] Executing 'gpumemsizesum' stats report

 Total (MB)  Count  Avg (MB)  Med (MB)  Min (MB)  Max (MB)  StdDev (MB)      Operation     
 ----------  -----  --------  --------  --------  --------  -----------  ------------------
    134.218      1   134.218   134.218   134.218   134.218        0.000  [CUDA memcpy HtoD]
      1.049      2     0.524     0.524     0.524     0.524        0.000  [CUDA memcpy DtoH]

可以看见，步骤 4 中调用的是 osrtsum（OS Runtime Summary），关注操作系统层面的性能数据。而如果在命令行中加入

@└────> # nsys profile --stats=true --trace=cuda,nvtx,cudnn,cublas ./a.out
...
[4/7] Executing 'cudaapisum' stats report

 Time (%)  Total Time (ns)  Num Calls    Avg (ns)       Med (ns)     Min (ns)   Max (ns)     StdDev (ns)            Name         
 --------  ---------------  ---------  -------------  -------------  --------  -----------  -------------  ----------------------
     94.6      282,480,873          2  141,240,436.5  141,240,436.5   269,662  282,211,211  199,362,781.2  cudaMalloc            
      5.2       15,504,039          3    5,168,013.0      826,866.0   541,742   14,135,431    7,767,320.2  cudaMemcpy            
      0.1          392,989          2      196,494.5      196,494.5    96,115      296,874      141,958.1  cudaFree              
      0.1          318,607          2      159,303.5      159,303.5    28,009      290,598      185,678.5  cudaLaunchKernel      
      0.0            1,184          1        1,184.0        1,184.0     1,184        1,184            0.0  cuModuleGetLoadingMode

则使用的是 cudaapisum（CUDA API Summary），关注的是 CUDA API 层面的性能数据。
除此之外，可以看到 gpukernsum 中核函数执行时间。reduce1 函数性能优于 reduce0，他们的执行次数，执行时间最大值最小值和平均值。

若算子是用 pybind 绑定，用 python 调用的，可以使用 torch 的函数来只监控该算子。

import torch
# warmup
for i in range(10):
    _ = torch.matmul(x, y) # x 和 y 是矩阵，这里就不展开了

# start profiling
torch.cuda.cudart().cudaProfilerStart()

### benchmarking
for i in range(100): # 测试 100 次
    torch.cuda.nvtx.range_push("your_ops_name")
    _ = your_ops(x, y)
    torch.cuda.nvtx.range_pop()
torch.cuda.synchronize()

# stop profiling
torch.cuda.cudart().cudaProfilerStop()

在进行完命令行分析后，会生成一个报告文件，结尾是 .nsys-rep。这个文件可以下载下来，丢进 nsight-system 可视化软件，在软件中可以看到更加详细的数据以及程序执行的时间线。
从上图可以看出 cuda hardware 的函数执行时间情况和 cpu 侧的执行时间情况。下面是核函数的发射时间，上面 device 侧是核函数实际执行时间，和 gpukernsum 统计的时间一致。
上图是在软件中可以看见的较为详细的统计数据，和命令行结果一致。

Llama结构解析

发表于 2023-12-06 更新于 2023-12-25 分类于 AI
本文字数： 179 阅读时长 ≈ 1 分钟

模型打印

已Llama-7B hugging face版本为例：

import torch
from transformers import AutoTokenizer, AutoModelForCausalLM
device = torch.device("cuda:{}".format(gpu))
tokenizer = AutoTokenizer.from_pretrained(model_name, trust_remote_code = True)
model = AutoModelForCausalLM.from_pretrained(model_name, trust_remote_code = True).half().to(device)
print(model)

LlamaForCausalLM(
  (model): LlamaModel(
    (embed_tokens): Embedding(32000, 4096)
    (layers): ModuleList(
      (0-31): 32 x LlamaDecoderLayer(
        (self_attn): LlamaFlashAttention2(
          (q_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (k_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (v_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (o_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (rotary_emb): LlamaRotaryEmbedding()
        )
        (mlp): LlamaMLP(
          (gate_proj): Linear(in_features=4096, out_features=11008, bias=False)
          (up_proj): Linear(in_features=4096, out_features=11008, bias=False)
          (down_proj): Linear(in_features=11008, out_features=4096, bias=False)
          (act_fn): SiLUActivation()
        )
        (input_layernorm): LlamaRMSNorm()
        (post_attention_layernorm): LlamaRMSNorm()
      )
    )
    (norm): LlamaRMSNorm()
  )
  (lm_head): Linear(in_features=4096, out_features=32000, bias=False)
)

从结构可以看出来，模型参数量为 $32,000 * 4,096 + 32 * (4,096 * 4,096 * 4 + 4,096 * 11,008 * 3) + 4,096 * 32,000 = 6,738,149,376$ 。所以约为 $7B$ 。

模型图解

以输入为 10 个 token 为例：

cpp attribute 之 visibility

发表于 2023-11-10 更新于 2023-12-25 分类于 cpp
本文字数： 470 阅读时长 ≈ 2 分钟

attribute ((visibility(“”)))

是 gcc 的编译器指令，用于设置在 shared object 中所修饰的符号对外的可见性。该修饰对 .a 文件不生效，只对 .so 库生效。

attribute ((visibility(“default”)))

该修饰用于修饰符号的可见性为默认对外可见。意思是通过该符号修饰的函数可以在 so 文件外访问到。

func.cpp:

1
2
3

__attribute__ ((visibility("default"))) void func1(int a) {
        cout << a << endl;
}

main.cpp:

#include <iostream>
using namespace std;
extern void func1(int a);
int main() {
    func1(10);
}

1 2	@└────> g++ func.cpp --shared -fPIC -o libfunc.so @└────> g++ main.cpp -L./ -lfunc

之后发现是可以编译成功的。因为该符号是可见的。

1 2	@└────> nm libfunc.so \| grep func 0000000000001179 T _Z5func1i

大写的 T 表示定义在 text 段，并且可被外部引用。如果你是通过编译 .o 文件再链接为 .so 文件的，还可以使用 readelf -s 查看 .o 文件的可见性.

attribute ((visibility(“hidden”)))

该修饰用于修饰符号的可见性为默认对外不可见。意思是通过该符号修饰的函数不可以在 so 文件外访问到,只能在 so 文件内部访问到。

func.cpp:

1
2
3

__attribute__ ((visibility("hidden"))) void func2(int a) {
        cout << a << endl;
}

main.cpp:

#include <iostream>
using namespace std;
extern void func2(int a);
int main() {
    func2(10);
}

@└────> g++ func.cpp --shared -fPIC -o libfunc.so
@└────> g++ main.cpp -L./ -lfunc
/usr/bin/ld: /tmp/cc7GABC5.o: in function `main':
fstream.cpp:(.text+0xe): undefined reference to `func2(int)'
collect2: error: ld returned 1 exit status

之后发现是可以编译失败,因为符号不可见.

1 2	@└────> nm libfunc.so \| grep func 00000000000011ef T _Z5func2i

其他

1 2	__attribute__ ((visibility("internal"))) __attribute__ ((visibility("protected")))

上述两种一样是用于修饰符号, internal 对外不可见，而 protected 对外可见。
此外，在编译 so 文件时可以通过指定 -fvisibility=xxx 来指定默认的没有给出修饰的符号属性。
如：

1	@└────> gcc -fPIC -shared -o libtest.so -fvisibility=hidden test.c

这样在 test.c 中没用经过修饰的符号对外都不可见，而修饰为 default 的依旧对外可见。

plt表链接过程

发表于 2023-06-13 更新于 2023-12-25 分类于 compiler
本文字数： 1.7k 阅读时长 ≈ 6 分钟

概念

在 ELF 文件中，查看可以获得它的节的名字。其中有几个带有 plt 和 got 的节。

在此处，给出各节的定义如下：

.got：Global Offset Table，全局偏移表。这是链接器为外部符号填充的实际偏移表。
.plt：Procedure Linkage Table，程序链接表。他有两个作用，要么在 .got.plt 中拿到链接地址跳转，要么触发链接器去寻找地址。
.got.plt：是 .got 的一部分（但是是两个不同的节），是 got 专门为 plt 准备的节，包含了 plt 表需要的地址。（新版 gcc 可能将他叫为 .plt.got）
.rela.plt：程序链接表的重定位表，记录所有全局函数的动态链接信息，用于在程序加载时修正 plt 表中的跳转指针，使它们指向正确的地址。

实验

接下来将使用 gdb 一步一步跟着汇编走完动态链接的过程。

准备工作

实验代码如下：

#include <stdio.h>
int main() {
    puts("hello");
    printf("hello");
    return 0;
}

查看节的地址与大小：

@└────> # objdump -h plt | grep -E "plt|got"
plt:     file format elf64-x86-64
  9 .rela.plt     00000030  0000000000400468  0000000000400468  00000468  2**3
 11 .plt          00000030  00000000004004c0  00000000004004c0  000004c0  2**4
 20 .got          00000020  0000000000600fe0  0000000000600fe0  00000fe0  2**3
 21 .got.plt      00000028  0000000000601000  0000000000601000  00001000  2**3

查看需要动态链接的符号：

@└────> # readelf -r plt

Relocation section '.rela.dyn' at offset 0x408 contains 4 entries:
  Offset          Info           Type           Sym. Value    Sym. Name + Addend
000000600fe0  000100000006 R_X86_64_GLOB_DAT 0000000000000000 _ITM_deregisterTMClone + 0
000000600fe8  000400000006 R_X86_64_GLOB_DAT 0000000000000000 __libc_start_main@GLIBC_2.2.5 + 0
000000600ff0  000500000006 R_X86_64_GLOB_DAT 0000000000000000 __gmon_start__ + 0
000000600ff8  000600000006 R_X86_64_GLOB_DAT 0000000000000000 _ITM_registerTMCloneTa + 0

Relocation section '.rela.plt' at offset 0x468 contains 2 entries:
  Offset          Info           Type           Sym. Value    Sym. Name + Addend
000000601018  000200000007 R_X86_64_JUMP_SLO 0000000000000000 puts@GLIBC_2.2.5 + 0
000000601020  000300000007 R_X86_64_JUMP_SLO 0000000000000000 printf@GLIBC_2.2.5 + 0

反汇编查看 plt 相关函数：

@└────> # objdump -d plt
Disassembly of section .plt:

00000000004004c0 <.plt>:
  4004c0:       ff 35 42 0b 20 00       pushq  0x200b42(%rip)        # 601008 <_GLOBAL_OFFSET_TABLE_+0x8>
  4004c6:       ff 25 44 0b 20 00       jmpq   *0x200b44(%rip)        # 601010 <_GLOBAL_OFFSET_TABLE_+0x10>
  4004cc:       0f 1f 40 00             nopl   0x0(%rax)

00000000004004d0 <puts@plt>:
  4004d0:       ff 25 42 0b 20 00       jmpq   *0x200b42(%rip)        # 601018 <puts@GLIBC_2.2.5>
  4004d6:       68 00 00 00 00          pushq  $0x0
  4004db:       e9 e0 ff ff ff          jmpq   4004c0 <.plt>

00000000004004e0 <printf@plt>:
  4004e0:       ff 25 3a 0b 20 00       jmpq   *0x200b3a(%rip)        # 601020 <printf@GLIBC_2.2.5>
  4004e6:       68 01 00 00 00          pushq  $0x1
  4004eb:       e9 d0 ff ff ff          jmpq   4004c0 <.plt>

开始

首先断点到 puts 函数，查看调用处：

@(gdb) disassemble main
Dump of assembler code for function main:
   0x00000000004005d6 <+0>:     push   %rbp
   0x00000000004005d7 <+1>:     mov    %rsp,%rbp
=> 0x00000000004005da <+4>:     mov    $0x400698,%edi
   0x00000000004005df <+9>:     callq  0x4004d0 <puts@plt>
   0x00000000004005e4 <+14>:    mov    $0x400698,%edi
   0x00000000004005e9 <+19>:    mov    $0x0,%eax
   0x00000000004005ee <+24>:    callq  0x4004e0 <printf@plt>
   0x00000000004005f3 <+29>:    mov    $0x0,%eax
   0x00000000004005f8 <+34>:    pop    %rbp
   0x00000000004005f9 <+35>:    retq
End of assembler dump.

可以看到，调用处实际上是使用 call 指令走到 puts 的代码段。下面的 printf 也是如出一辙。

查看 puts@plt 的汇编指令

@(gdb) disassemble
Dump of assembler code for function puts@plt:
=> 0x00000000004004d0 <+0>:     jmpq   *0x200b42(%rip)        # 0x601018 <puts@got.plt>
   0x00000000004004d6 <+6>:     pushq  $0x0
   0x00000000004004db <+11>:    jmpq   0x4004c0
End of assembler dump.

可以看到，在汇编中，他首先要跳转到 0x601018 地址的位置。这个地址内容是个全局变量，实际上根据节的地址位置和大小可以判断，是处于 .got.plt 的位置内（ 0x601000 ~ 0x601028）。所以可以认为，在 .got.plt 中，存在了 puts 函数的地址。

查看 .got.plt

@(gdb) x/16x 0x601018
0x601018 <puts@got.plt>:        0x004004d6      0x00000000      0x004004e6      0x00000000
0x601028:       0x00000000      0x00000000      0x00000000      0x00000000
0x601038:       0x00000000      0x00000000      0x00000000      0x00000000
0x601048:       0x00000000      0x00000000      0x00000000      0x00000000

查看表中内容，发现跳转的地址是 0x4004d6，这不就是我们跳转之前的下一个地址吗！（puts@plt 的第二条指令） 同理，printf 函数也是如此（0x4004e6）。这是因为，之前没有调用过 puts 函数，第一次查找的时候，.got.plt 表中找不到函数的地址，那就先返回继续执行去调用链接器获取地址。

准备调用链接器

00000000004004d0 <puts@plt>:
  4004d0:       ff 25 42 0b 20 00       jmpq   *0x200b42(%rip)        # 601018 <puts@GLIBC_2.2.5>
  4004d6:       68 00 00 00 00          pushq  $0x0
  4004db:       e9 e0 ff ff ff          jmpq   4004c0 <.plt>

首先 pushq $0x0，这个是在 got.plt 中的编号，如 puts 是 0，printf 是 1。这个参数是给后续链接器使用的。然后跳到了 .plt 的位置执行（0x4004c0）。可以看到，printf@plt 函数最后也是跳到这个位置执行。

调用链接器

00000000004004c0 <.plt>:
  4004c0:       ff 35 42 0b 20 00       pushq  0x200b42(%rip)        # 601008 <_GLOBAL_OFFSET_TABLE_+0x8>
  4004c6:       ff 25 44 0b 20 00       jmpq   *0x200b44(%rip)        # 601010 <_GLOBAL_OFFSET_TABLE_+0x10>
  4004cc:       0f 1f 40 00             nopl   0x0(%rax)

首先 push 了 0x601008 到栈中，这是 .got.plt 表中的一个地址。之后跳转到 0x601010 所存储的地址去执行相应的代码。不难看出，0x601010 也是存储在 .got.plt 表中的。查看一下存储的内容：

@(gdb) x/10x 0x601010
0x601010:       0xf7de64a0      0x00007fff      0x004004d6      0x00000000
0x601020 <printf@got.plt>:      0x004004e6      0x00000000      0x00000000      0x00000000
0x601030:       0x00000000      0x00000000

可以看到，是让我们跳转到 0x00007ffff7de64a0 去执行相应的代码。那么这块代码是什么呢？

@(gdb) info sharedlibrary
From                To                  Syms Read   Shared Object Library
0x00007ffff7dd0fa0  0x00007ffff7df2cd4  Yes (*)     /lib64/ld-linux-x86-64.so.2
0x00007ffff7a2cb90  0x00007ffff7b798ad  Yes (*)     /lib64/libc.so.6
(*): Shared library is missing debugging information.

可以看到，该地址是 ld-linux-x86-64.so 加载的位置。说明执行的是链接器的代码。

1: x/5i $pc
=> 0x7ffff7de64a0 <_dl_runtime_resolve_xsavec>: endbr64
   0x7ffff7de64a4 <_dl_runtime_resolve_xsavec+4>:       push   %rbx
   0x7ffff7de64a5 <_dl_runtime_resolve_xsavec+5>:       mov    %rsp,%rbx
   0x7ffff7de64a8 <_dl_runtime_resolve_xsavec+8>:       and    $0xffffffffffffffc0,%rsp
   0x7ffff7de64ac <_dl_runtime_resolve_xsavec+12>:
    sub    0x21616d(%rip),%rsp        # 0x7ffff7ffc620 <_rtld_local_ro+384>
@(gdb) bt
#0  0x00007ffff7de64a0 in _dl_runtime_resolve_xsavec () from /lib64/ld-linux-x86-64.so.2
#1  0x00000000004005e4 in main () at plt.c:3

可以看到这里代码执行的是 ld 中的 _dl_runtime_resolve_xsavec 函数是第一次函数调用时用于查找函数符号的，并且在结尾处会直接去调用找到的函数符号（本文中为 puts 函数）。

写回 .got.plt 表
在 puts 上打个断点，这样继续的话就是执行完 _dl_runtime_resolve_xsavec 还未执行 puts 的状态了。

@(gdb) bt
#0  0x00007ffff7a7d8c0 in puts () from /lib64/libc.so.6
#1  0x00000000004005e4 in main () at plt.c:3
@(gdb) x/10x 0x601018
0x601018 <puts@got.plt>:        0xf7a7d8c0      0x00007fff      0x004004e6      0x00000000
0x601028:       0x00000000      0x00000000      0x00000000      0x00000000
0x601038:       0x00000000      0x00000000

可以看到，此时，got.plt 表中的地址已经被写为 puts 函数实际的地址了（0x00007ffff7a7d8c0 在 0x00007ffff7a2cb90 ~ 0x00007ffff7b798ad 范围内，属于 /lib64/libc.so.6），这样下次调用 puts 就不用再次调用链接器了。

题外话

其实看一下 .got.plt 表的内容，会发现明明 puts 是第一个需要被链接的函数，为什么第一个却不是它呢？

@(gdb) x/10x 0x601000
0x601000:       0x0000000000600e10      0x00007ffff7ffe1d0
0x601010:       0x00007ffff7de64a0      0x00007ffff7a7d8c0
0x601020 <printf@got.plt>:      0x00000000004004e6      0x0000000000000000
0x601030:       0x0000000000000000      0x0000000000000000
0x601040:       0x0000000000000000      0x0000000000000000

puts 地址实际上是 got[3]:0x00007ffff7a7d8c0，前面还有 3 项。其中：

got[0]:0x0000000000600e10 自身模块 dynamic 段地址

1 2	@(gdb) info symbol 0x0000000000600e10 _DYNAMIC in section .dynamic of /root/xxx/plt

got[1]:0x00007ffff7ffe1d0 本模块的 link_map 的地址。编译期间会初始化为 0。link_map 是一个双向链表的入口，链接进程所有加载的动态库。当链接器查找符号时，通过遍历该链表找到对应的符号。
got[2]:0x00007ffff7de64a0 _dl_runtime_resolve_xsavec 的地址。

1 2	@(gdb) info symbol 0x00007ffff7de64a0 _dl_runtime_resolve_xsavec in section .text of /lib64/ld-linux-x86-64.so.2

_dl_runtime_resolve 格式：

//调用形式为：
_dl_runtime_resolve((link_map*)(got[1]), 0);
// 第二个参数 0，为 <puts@plt>:中的 pushq $0x0；
// 同理如果是 printf，就是<printf@plt>:中 pushq $0x1；

总结

虚拟地址空间内流程图：

第二次调用：

内嵌汇编

发表于 2023-06-09 更新于 2023-12-25 分类于 linux
本文字数： 469 阅读时长 ≈ 2 分钟

为什么要使用内嵌汇编？

内嵌汇编通常用于在程序中实现一些高效、精确的操作。例如，在嵌入式平台上运行的程序，如果需要代码占用内存更小、程序运行的效率更高或需要准确地操作寄存器时，嵌入汇编会是不错的选择。

基本语法

asm("assembly code"        /* 汇编代码 */
    :output_operand        /* 输出参数列表 */          
    :input_operand         /* 输入参数列表 */          
    :clobbered_operand     /* 被改变的操作对象列表 */      
);

// 举例
static int value_assignment(int input) {
    int ret = 0;
    asm volatile(
        "movl %1, %0\n"   // 超过一条指令就要用 \n 来分割，排版整齐还要加 \t
        :"=r"(ret)
        :"r"(input)
    );
    return ret;
}

被改变的操作对象列表

在被改变的参数列表 clobbered_operand 中有一个比较有用的标识符：memory。指定 memory，相当于对编译器形成了一个内存读写的屏障，保证在内联汇编执行前，编译器将某些寄存器里的值刷新进内存，同时在内联汇编执行后，编译器重新加载相关变量的值。
所以我们可以见到这样的代码：

1	asm volatile ("" ::: "memory");

作为内存屏障，保证编译器的优化不会跨过这道屏障。加上 volatile 告诉编译器不要优化汇编。

修饰符

修饰符一般跟在参数列表前面。

修饰符	含义
=	只写，常用于修饰所有输出操作数
无	只读
+	可读可写
r	可以是任意通用寄存器存储其值
m	一个有效的内存地址
i	是立即数
%	被修饰的操作数可以和下一个互换
&	只能做输出，一般和 “=” 一起使用，如 “=&r(val)”
x	只能做输入

占位符

%0 表示输入和输出列表合并的第 1 个操作数，%1 表示第 2 个，以此类推。

内存模型与内存序

发表于 2023-06-07 更新于 2023-12-25 分类于 compiler
本文字数： 1.9k 阅读时长 ≈ 7 分钟

硬件结构

现代机器都是多个处理器，每个处理器有自己的 cache。这个结构如下所示:

可以看到，每个 CPU 都有自己的缓存，之后再写到内存中。并且由于编译器的优化，你写的代码可能和你执行的代码顺序有所不同。他们优化的规则是：保证对于一个单核情况下，执行结果不会发生变化。但是多线程就不一定了。

那么在多线程情况下，如何协调这些 CPU 缓存的数据一致性就成了一个问题。

常见优化

再谈保证数据的一致性之前，先谈谈编译器能做的优化。

重排 Reordering

编译器和 CPU 都会发生重排，为了提升代码的效率。采用乱序执行、流水线、分支预测以及多级缓存等方法来提升程序性能。编译器会基于这些规则来提升自己代码的速度，所以就会对指令进行优化。例子如下：

int a = 0;
int b = 0;

void fun() {
    a = b + 1; // L5
    b = 1;     // L6
}

int main() {
    fun();
    return 0;
}

@└────> # gcc 1.c -O0 -g
@└────> # objdump -d a.out
0000000000400536 <fun>:
  400536:       55                      push   %rbp
  400537:       48 89 e5                mov    %rsp,%rbp
  40053a:       8b 05 e4 0a 20 00       mov    0x200ae4(%rip),%eax        # 601024 <b>
  400540:       83 c0 01                add    $0x1,%eax
  400543:       89 05 d7 0a 20 00       mov    %eax,0x200ad7(%rip)        # 601020 <__TMC_END__>
  400549:       c7 05 d1 0a 20 00 01    movl   $0x1,0x200ad1(%rip)        # 601024 <b>
  400550:       00 00 00 
  400553:       90                      nop
  400554:       5d                      pop    %rbp
  400555:       c3                      retq

@└────> # gcc 1.c -O2 -g
@└────> # objdump -d a.out
0000000000400560 <fun>:
  400560:       8b 05 ba 0a 20 00       mov    0x200aba(%rip),%eax        # 601020 <__TMC_END__>
  400566:       c7 05 b0 0a 20 00 01    movl   $0x1,0x200ab0(%rip)        # 601020 <__TMC_END__>
  40056d:       00 00 00 
  400570:       83 c0 01                add    $0x1,%eax
  400573:       89 05 ab 0a 20 00       mov    %eax,0x200aab(%rip)        # 601024 <a>
  400579:       c3                      retq   
  40057a:       66 0f 1f 44 00 00       nopw   0x0(%rax,%rax,1)

对于 O0 等级的优化，执行顺序是 L5->L6。
但是对于 O2 等级的优化，执行顺序是 L6->L5，但是结果是不影响的。

为什么要这么做呢？因为 CPU 读取数据从 cache 中读取。如果不优化的话，先读 b，再读 a 的时候可能把 b 的缓存换出去了，那么再写 b 的时候还需要把 b 换进来。但是如果优化了，就是读 b，写 b，再写 a，就不存在缓存的换入换出了。

插入 Invention

假设有如下代码：

1
2
3

for (int i = 0; i < n; ++i) {
    x[i] = y[i] + z[i];
}

可能优化成如下：

for (int i = 0; i < n; ++i) {
    __builtin_prefetch(&y[i + 16]);
    __builtin_prefetch(&z[i + 16]);
    x[i] = y[i] + z[i];
}

预读取这些数据来减少缓存未命中次数。

删除 Removal

删除很好理解了，删除没用的变量赋值。

int x = 1;
int y = 2;
int z = x + y;
x = 3;
y = 4;

优化后：

int x;
int y;
x = 3;
y = 4;

关系术语

sequence-before

sequence-before 是对一个线程内，求值顺序关系的描述：

A sequence-before B，先对 A 求值，再对 B 求值。
A not sequence-before B，并且 B not sequence-before A，那么 A 和 B 谁先求值是未知的。

synchronizes-with

描述的是不同线程内的执行关系。在两个线程分别执行时，即使线程 A 先执行，线程 B 后执行，A 中写了某个共享变量，由于指令重排或者写到了 cache 或寄存器没来得及写入内存导致 B 读到了错误的值。

A synchronizes-with B，在线程 A 中的写操作结果对线程 B 可见。

happens-before

是 sequence-before 的扩展，包括了不同线程的关系。

A happens-before B，那么不但 A 先于 B 执行，并且 A 的结果对 B 可见。
- 同线程：和 sequence-before 一样。
- 不同线程：和 synchronizes-with 一样。

内存序

C++11 中引入了 6 种内存序：

typedef enum memory_order {
    memory_order_relaxed,
    memory_order_consume,
    memory_order_acquire,
    memory_order_release,
    memory_order_acq_rel,
    memory_order_seq_cst
} memory_order;

内存序类型	用于读/写	含义
memory_order_relaxed	读/写	仅要求原子性和内存一致性
memory_order_consume	读	读操作所在线程该操作后面的和该变量有依赖关系的读写操作不会被优化到先于该操作执行
memory_order_acquire	读	读操作所在线程该操作后面的读写操作不会被优化到先于该操作执行
memory_order_release	写	写操作所在线程该操作前面的读写操作不会被优化到后于该操作执行
memory_order_acq_rel	读/写	是 memory_order_acquire 和 memory_order_release 组成的双向屏障，上下皆不能跨过该指令
memory_order_seq_cst	读/写	双向屏障，并且该线程所有原子指令并且也指定为 memory_order_seq_cst 的都已全局内存修改顺序为参照

值得一提的是，若一个原子变量在一个线程中施加了 memory_order_release，但是在其他线程中没有使用 memory_order_acquire 或 memory_order_consume 读取，那么他就不会具备 memory_order_release 所赋予的屏障功能。（即只有被观测才会起作用，读操作也是如此）

如上图所示，就像加锁一样会构成临界区。但是外面的变量可以移入临界区，却不能移出去，所以称 memory_order_acquire 和 memory_order_release 如同单向屏障一般。

内存模型

一言以蔽之，引入内存模型的原因，有以下几个原因：

编译器优化：在某些情况下，即使是简单的语句，也不能保证是原子操作。
CPU out-of-order：CPU 为了提升计算性能，可能会调整指令的执行顺序。
CPU Cache 不一致：在 CPU Cache 的影响下，在某个 CPU 下执行了指令，不会立即被其它 CPU 所看到。

从上面的内存序中，按照访问控制的角度可以分为三种模型：

Sequential Consistency 模型
Acquire-Release 模型
Relax 模型

其中，Sequential Consistency 模型约束最强，Acquire-Release 次之，Relax 模型最弱。

Sequential Consistency 模型

对应 memory_order_seq_cst 内存序。Sequential Consistency 模型有以下特点：

每个线程的执行顺序与代码顺序严格一致
线程的执行顺序可能会交替进行，但是从单个线程的角度来看，仍然是顺序执行

例如：

x = y = 0;

thread1:
x = 1;
r1 = y;

thread2:
y = 1;
r2 = x;

那么可能的执行顺序为：

可能性	第一步	第二步	第三步	第四步
1	x = 1	r1 = y	y = 1	r2 = x
2	y = 1	r2 = x	x = 1	r1 = y
3	x = 1	y = 1	r1 = y	r2 = x
4	x = 1	r2 = x	y = 1	r1 = y
5	y = 1	x = 1	r1 = y	r2 = x
6	y = 1	x = 1	r2 = x	r1 = y

std::atomic 默认值都是使用 memory_order_seq_cst，保证不出错。但是相对的，限制了 CPU 并行处理的能力，会降低效率。这个模型的所有线程都参考全局的内存修改顺序。因此，我们可认为所有变量的读写都直接从内存进行，从而完全不用考虑 Cache，Store Buffer 这些因素。

Acquire-Release 模型

对应 memory_order_consume、memory_order_acquire、memory_order_release、memory_order_acq_rel 内存序。对于一个原子变量 A，对 A 的写操作（Release）和读操作（Acquire）之间进行同步，并建立排序约束关系，即对于写操作（release）X，在写操作 X 之前的所有读写指令都不能放到写操作 X 之后；对于读操作（acquire）Y，在读操作 Y 之后的所有读写指令都不能放到读操作 Y 之前。

Relax 模型

对应的是 memory_order_relaxed 内存序。其对于内存序的限制最小，也就是说这种方式只能保证当前的数据访问是原子操作（不会被其他线程的操作打断），但是对内存访问顺序没有任何约束，也就是说对不同的数据的读写可能会被重新排序。

x86_64常用汇编

发表于 2023-05-30 更新于 2023-12-25 分类于 linux
本文字数： 1.8k 阅读时长 ≈ 6 分钟

本文用以记录常用汇编指令以供快速查找回忆，仅限于 X86_64 的 AT&T 格式。

语法格式

1. 引用寄存器前加 %。如

1	mov %rsp, %rbp

2. 指令长度后缀

对于访问内存的数据，指令后加上 b w l q，操作 1 2 4 8 字节。如

movb   $0x1,0x201c3f(%rip)
nopw   %cs:0x0(%rax,%rax,1)
movl   $0x5,-0xc(%rbp)
movq   $0x400b30,-0x18(%rbp)

3. 立即数前加 $。16 进制数用 0x 开头。如

1 2	movl $1, %eax mov $0x0,%eax

4. 注释可以用 ! 开头，也可以用 ;

5. 操作数顺序

从源操作数到目的操作数，如下将 %rsp 寄存器中的数传给 %rbp 寄存器。

1	mov %rsp,%rbp

6. 数据声明

命令	数据类型
.ascii	文本字符串
.asciz	以空字符串结尾的文本字符串
.byte	字节值
.double	双精度浮点数
.float	单精度浮点数
.single	单精度浮点数同上
.int	32位整数
.long	32位整数同上
.octa	16字节整数
.quad	8字节整数
.short	16位整数
.comm	声明未初始化的数据的通用内存区域
.lcomm	声明未初始化的数据的本地通用内存区域

7. 文件组成

命令	作用
.org	定义当前汇编位置
.globl	让段全局可见
.text	存放代码指令正文段
.bss	存放未初始化的全局和静态变量，运行时该区域初始化为 0
.rodata	read only data
.data	可读可写的数据段

8. 寻址方式

直接寻址：把某个地址上的值放到寄存器中

1	mov $0x8000,%eax

间址寻址：把寄存器上的值所代表的地址所指向的值放到寄存器中

1 2	movl $0x8000,%ebx movl (%ebx),%eax ; 间址寻址, 把地址 0x8000（在寄存器 %ebx 中）上的值放到 %eax 中

基址寻址：以寄存器里的数值作为基址，加上一个常数得到最终地址，把地址上的值放到寄存器中

1 2	movl $0x8000,%eax movl 4(%eax),%ebx ; 基址寻址, 把地址 0x8004（0x8000+4）上的值放到 %eax 中

变址寻址：以两个寄存器里的数值之和加上一个常数得到最终地址，把地址上的值放到寄存器中

movl   $0x8000,%eax
movl   $0x4,%ebx
movl   (%eax,%ebx),%ecx    ; 变址寻址, 把地址 0x8004（0x8000+4）上的值放到 %ecx 中
movl   4(%eax,%ebx),%ecx   ; 变址寻址, 把地址 0x8008（0x8000+4+4）上的值放到 %ecx 中

比例变址寻址：以一个寄存器里的数值加上另一个寄存器里的数字，乘以一个比例因子（1,2,4,8）再加上一个常数得到最终地址，把地址上的值放到寄存器中

movl   $0x2000,%eax   
movl   $0x2,%ebx   
movl   (,%eax,4),%ecx      ; 比例变址寻址, 把地址 0x8000（0 + 0x2000*4）上的值放到 %ecx 中
movl   6(,%eax,4), %ecx    ; 比例变址寻址, 把地址 0x8006（0 + 0x2000*4 + 6）上的值放到 %ecx 中
movl   (%ebx,%eax,4),%ecx  ; 比例变址寻址, 把地址 0x8002（0x2 + 0x2000*4）上的值放到 %ecx 中
movl   6(%ebx,%eax,4),%ecx ; 比例变址寻址, 把地址 0x8008（0x2 + 0x2000*4 + 6）上的值放到 %ecx 中

常见指令

1. mov 用于将源操作数移动到目的操作数

1	mov %rsp,%rbp ; %rbp = %rsp

2. add 用于将源操作数加给目的操作数

1	addl %eax,%ebx ; %ebx = %ebx + %eax

3. sub 用于将两个数相减

1	subl %eax,%ebx ; %ebx = %ebx - %eax

4. inc 用于加一

1	incl %eax ; %eax = %eax + 1

5. dec 用于减一

1	decl %eax ; %eax = %eax - 1

6. push 用于将数据压入栈

1	pushl %eax ; 入栈，%esp = %esp - 0x4, %esp = %eax

7. pop 用于将数据出栈

1	popl %eax ; 出栈，%eax = %esp, %esp = %esp + 0x4

8. jmp 跳转

jmp    label          ; 无条件跳转为 label, %rip = label
je     label          ; 相等 ZF = 1, %rip = label
jne    label          ; 不相等 ZF = 0, %rip = label
jg     label          ; 大于 %rip = label
jge    label          ; 大于等于 %rip = label
jl     label          ; 小于 %rip = label
jle    label          ; 小于等于 %rip = label
ja     label          ; 无符号比较 大于 %rip = label
jae    label          ; 无符号比较 大于等于 %rip = label
jb     label          ; 无符号比较 小于 %rip = label
jbe    label          ; 无符号比较 小于等于 %rip = label

9. mul 乘法

1 2	imull %eax,%ebx ; %ebx = %eax * %ebx 用于有符号数 mull %eax,%ebx ; %ebx = %eax * %ebx 用于无符号数

10. div 除法

1 2	idivl %ebx ; %edx = %eax % %ebx, %eax = %eax / %ebx 用于有符号数 divl %ebx ; %edx = %eax % %ebx, %eax = %eax / %ebx 用于无符号数

11. and 按位与

1	andl %eax,%ebx ; %ebx = %ebx & %eax

12. or 按位或

1	orl %eax,%ebx ; %ebx = %ebx \| %eax

13. xor 按位异位

1	xorl %eax,%ebx ; %ebx = %eax ^ %ebx

14. shl 和 sal 位左移

1 2	shll $1,%eax ; %eax = %eax << 1 逻辑左移，填充 0 sall $1,%eax ; %eax = %eax << 1 算数左移，填充 0

15. shr 和 sar 位右移

1 2	shrl $1,%eax ; %eax = %eax >> 1 逻辑右移，填充 0 sarl $1,%eax ; %eax = %eax >> 1 算数右移，填充符号位

16. lea 装载有效地址

1	leal 8(%ebx),%eax ; %eax = 8 + %ebx 可理解为 %eax = &(*(%ebx)) + 8

17. call 函数调用

1	call func_name ; 将下一条指令的 %rip push 到栈中，之后 %rip = func_name

18. ret 函数返回

1	ret ; 将函数返回地址的下一条要执行指令的值赋值给 %rip，push %rip

19. test 与运算并设置标志寄存器

testl  %eax,%ebx      ; %eax & %ebx，不会改变这两个寄存器值，改变标志寄存器零标志位（ZF）、符号标志位（SF）、奇偶标志位（PF）和进位标志位（CF），但不会影响溢出标志位（OF）

20. cmp 比较操作数大小

1	cmpl %eax,%ebx ; 根据 %ebx - %eax 的值来改变零标志位（ZF）、符号标志位（SF）、奇偶标志位（PF）、进位标志位（CF）和溢出标志位（OF）

21. rep 重复执行指令直到某一条件

1 2	repz movsb ; 重复执行 movsb 直到 ZF = 0 repne scasb ; 重复执行 scasb 直到 ZF = 1

22. lock 锁定总线

1	lock addl $1,(%eax) ; 锁定总线，并使 (%eax) = (%eax) + 1，因为总线是锁定的，不会被其他处理器打断

23. xadd 交换两个操作数值，使他们相加

1	xaddl %eax,%ebx ; tmp = %eax，%eax = %ebx，%ebx = tmp + %ebx 交换两个数，并将和写到 %ebx

24. nop 空操作

1	nop ; 什么都不做，充当占位符或者插入延迟

25. hlt 使处理器暂停直到收到中断信号

hlt                   ; 使处理器进入暂停状态，直到发生外部中断。它通常用于操作系统内核中，以降低功耗和发热量。只有特权级别为 0 （内核态） 才能使用，否则会导致异常

26. xchg 交换两个操作数的值

1	xchgl %eax,%ebx ; tmp = %eax，%eax = %ebx，%ebx = tmp

27. cld 清除方向寄存器（DF）

1	cld ; 清除方向寄存器，使 %rdi 递增

28. movsb 移动字符串

1	movsb ; 以 %rsi 为源地址，%rdi 为目的地址，将字符以一个字节拷贝。每次执行 movsb，%rsi 和 %rdi 以方向标志寄存器（DF）自动递增或递减

29. scasb 查找字符

1	scasb ; 将被查找字符放到 %al 中，与 %rdi 地址的字符串依次比较，根据比较结果设置标志寄存器

30. cli 禁用所有中断

1 2	cli ; 禁用所用中断 hlt ; 使处理器保持暂停状态，直到中断被重新启用

进程代码段地址空间随机加载

发表于 2023-05-26 更新于 2023-12-25 分类于 linux
本文字数： 778 阅读时长 ≈ 3 分钟

问题

今天调试代码的时候看到地址的时候突然感到奇怪：我记得我之前看到的代码地址空间好多都是 0x400xxx 开头的，怎么这次的地址空间是 0x5562b845axxx 呢？是什么导致了这个差异？

我换了地址空间为 0x400xxx 开头的机器，准备了相同的代码，在两台不同的机器上编译：

#include <stdio.h>
int main() {
    printf("%p\n", main);
}

这个简单的程序可以打出 main 函数的地址。经测试，在不同的机上打出的结果有很大差异。

@└────> # ./a.out 
0x5562b845a649

@└────> # ./b.out 
0x400596

答案

经查阅资料，这个问题是 Linux 的 ASLR （Address Space Layout Randomization）导致的。这项技术会在装载时，装载到随机地址，防止黑客利用固定地址注入恶意代码。对于 b.out，没有使用该技术。所以 b.out 的代码段虚拟地址一直是 0x400000 开头。

@└────> # readelf -h b.out 
ELF Header:
  Magic:   7f 45 4c 46 02 01 01 00 00 00 00 00 00 00 00 00 
  Class:                             ELF64
  Data:                              2's complement, little endian
  Version:                           1 (current)
  OS/ABI:                            UNIX - System V
  ABI Version:                       0
  Type:                              EXEC (Executable file)        // 这里是 EXEC
  Machine:                           Advanced Micro Devices X86-64
  Version:                           0x1
  Entry point address:               0x4004b0                      // 这里是 _start 的绝对地址
  Start of program headers:          64 (bytes into file)
  Start of section headers:          15608 (bytes into file)
  Flags:                             0x0
  Size of this header:               64 (bytes)
  Size of program headers:           56 (bytes)
  Number of program headers:         9
  Size of section headers:           64 (bytes)
  Number of section headers:         30
  Section header string table index: 29

可以看到，对于 b.out，他的文件类型是 Executable file，_start 的地址是 0x400xxx 开头。这种就是没有使用 ASLR 技术的。而对于 a.out，结果如下：

@└────> # readelf -h a.out 
ELF Header:
  Magic:   7f 45 4c 46 02 01 01 00 00 00 00 00 00 00 00 00 
  Class:                             ELF64
  Data:                              2's complement, little endian
  Version:                           1 (current)
  OS/ABI:                            UNIX - System V
  ABI Version:                       0
  Type:                              DYN (Shared object file)      // 这里是 DYN
  Machine:                           Advanced Micro Devices X86-64
  Version:                           0x1
  Entry point address:               0x560                         // 这里是 _start 的相对地址
  Start of program headers:          64 (bytes into file)
  Start of section headers:          12744 (bytes into file)
  Flags:                             0x0
  Size of this header:               64 (bytes)
  Size of program headers:           56 (bytes)
  Number of program headers:         9
  Size of section headers:           64 (bytes)
  Number of section headers:         31
  Section header string table index: 30

对于 a.out，文件类型为 Shared object file，而且 _start 的地址是个相对地址。就是这个导致的这个差异。每次装载 a.out 时，代码会被加载到随机的位置。可以看到，每次运行，得到的地址都不同。

@└────> # ./a.out 
0x559536d9d649
@└────> # ./a.out 
0x559a7a6df649
@└────> # ./a.out 
0x55ca5dbd4649

发生根因

之所以发生这个原因，是因为操作系统版本导致的。低版本操作系统默认不使用 ASLR。想要在不同的操作系统上复现这两个方式也很简单：

1	@└────> # gcc 1.c -fPIC -pie

这种方式编译出来的就是使用了 ASLR 技术的。其中 -pie 的意思是 position-independent executable，位置无关的可执行文件。编译时还需要加上 -fPIC （Position-Independent Code）生成位置无关代码。而

1	@└────> # gcc 1.c -no-pie

方式编出来的就是固定地址。有些工具必须使用 -no-pie 才可以使用。这样固定的情况也比较好调试，因为虚拟地址固定。

linux中>&和&>区别

发表于 2023-05-22 更新于 2023-12-25 分类于 linux
本文字数： 422 阅读时长 ≈ 2 分钟

Linux 中常用的文件描述符

0 文件描述符，表示标准输入。
1 文件描述符，表示标准输出。
2 文件描述符，表示标准错误。

标准情况下，这些文件描述符和以下设备关联：

0 文件描述符关联键盘，并返回给前端。
1 正确返回值，返回给前端。
2 错误返回值，返回给前端。

> 符号

在 shell 中，我们经常使用 > 符号，把输出重定位到一个文件。例如：

1	cat /proc/xxx/maps > memory.txt

以上输出是把某个进程的内存布局重定向到一个文件。其中，> 是 1> 的简写，实际意思是把标准输出重定向到后面的文件。这样屏幕上就不会有打印了，打印会重定向到文件中。

>& 符号

本质上，>& 符号不是一个符号。我们经常见到 2>&1 符号，实际意义是，将标准错误重新定位到标准输出。那为什么要加个 & 呢？因为不加 & 的话操作系统不会认为你是想把标准错误重定位给标准输出，而是想重定向到一个叫 “1” 的文件。所以 &1 表示 1 输出通道。举例，strace 命令可以查看系统调用，这个结果是输出到标准错误的。

1	strace ls > log 2>&1

将标准输出重定向到 log 文件，并将标准错误重定向到标准输出。这样标准错误也会被重定向到 log 文件。

&> 符号

&> 意思是把标准错误和标准输出都重定向到某个文件。

1	strace ls &> log

写起来比较简单，且省力。

递归转迭代

发表于 2023-05-17 更新于 2023-12-25 分类于 algorithm ， os
本文字数： 1.2k 阅读时长 ≈ 4 分钟

函数调用时发生了什么

在 C 语言中，函数调用在底层汇编究竟发生了什么呢？示例如下：

int add(int a, int b) {
    return a + b;
}
int main() {
    int a = 9;
    int b = 10;
    int c = 11;
    int d = add(a, b);
    return 0;
}

对上面这个文件编译的结果进行反汇编，这两个函数的反汇编结果如下：

@└────> # objdump -d a.out

0000000000400536 <add>:
  400536:       55                      push   %rbp              // rbp 是被调用者保存，把上一个函数的栈基址保存
  400537:       48 89 e5                mov    %rsp,%rbp         // 将栈顶设为新的栈基址，至此本函数的栈已经初始化好了
  40053a:       89 7d fc                mov    %edi,-0x4(%rbp)   // 将局部变量 a 放在栈上
  40053d:       89 75 f8                mov    %esi,-0x8(%rbp)   // 将局部变量 b 放在栈上
  400540:       8b 55 fc                mov    -0x4(%rbp),%edx
  400543:       8b 45 f8                mov    -0x8(%rbp),%eax
  400546:       01 d0                   add    %edx,%eax         // 相加两个数，并将返回值放在 eax
  400548:       5d                      pop    %rbp              // 恢复 上一个函数的栈基址
  400549:       c3                      retq                     // 恢复 rip 寄存器，返回调用该函数的下一条指令地址

000000000040054a <main>:
  40054a:       55                      push   %rbp
  40054b:       48 89 e5                mov    %rsp,%rbp         // 同上，栈初始化
  40054e:       48 83 ec 10             sub    $0x10,%rsp        // 栈的空间先摆好
  400552:       c7 45 fc 09 00 00 00    movl   $0x9,-0x4(%rbp)   // a 初始化
  400559:       c7 45 f8 0a 00 00 00    movl   $0xa,-0x8(%rbp)   // b 初始化
  400560:       c7 45 f4 0b 00 00 00    movl   $0xb,-0xc(%rbp)   // c 初始化
  400567:       8b 55 f8                mov    -0x8(%rbp),%edx
  40056a:       8b 45 fc                mov    -0x4(%rbp),%eax
  40056d:       89 d6                   mov    %edx,%esi         // 第二个入参，因为从右往左放入
  40056f:       89 c7                   mov    %eax,%edi         // 第一个入参
  400571:       e8 c0 ff ff ff          callq  400536 <add>
  400576:       89 45 f0                mov    %eax,-0x10(%rbp)  // 返回值放入 d
  400579:       b8 00 00 00 00          mov    $0x0,%eax         // return 0
  40057e:       c9                      leaveq 
  40057f:       c3                      retq

可以看到，在函数调用时，发生了不同函数栈的切换。其中涉及到一些相关汇编指令。

汇编指令中的 push pop call 和 ret

push 和 pop

如上图所示，%rbp 寄存器和 %rsp 寄存器表示这一个栈的 “基址” 和 “栈顶”。因为栈是从高向低生长的，所以基址在栈顶上面。push 操作相当于，先将栈顶向下移动（因为存数据是向 %rsp 指向的位置写数据，而 %rsp 指向的是目前已有的栈顶数据），再将数据写入。pop 操作则相反。通过这两个寄存器，可以维护一个栈的存在。调用函数设计到栈的切换，所以就是通过改变这两个寄存器的值来达到切换的。

call 和 ret

如上图所示，调用函数涉及到执行指令的切换。将要执行的下一条指令又由 %rip 寄存器指出。

call 指令可以改变 %rip 寄存器的值：将 call 汇编指令的下一条指令地址 push 到栈中，之后改变 %rip 为 call 的指令地址。

ret 指令也可以改变 %rip 寄存器的值：将栈顶元素 pop 给 %rip，将其改为之前 call 时存的下一条要执行的指令地址。

二叉树遍历示例

那么操作系统这一系列操作对于我们将递归算法转换为迭代算法有什么启发吗？

vector<int> ret;
void helper(TreeNode* root) {
    if (root == nullptr) {
        return;
    }
    helper(root->left);
    ret.push_back(root->val);
    helper(root->right);
}
vector<int> inorderTraversal(TreeNode* root) {
    ret.clear();
    helper(root);
    return ret;
}

上面是二叉树中序遍历的递归写法，非常简洁。其中可以看到在递归函数 helper 中，又调用了两次 helper。我们知道每次调用 helper，都有它自己的调用栈。那么我们向迭代转换时，就要注意不能让他们的调用栈互相冲突。那么我们可以自己定义一个状态变量，用来存储当前执行的函数栈。除此之外，还要知道当前栈执行到本次函数的哪个位置（%rip)。两次 helper 的执行，可以将整个调用结构分为三部分：

执行 helper(root->left) 之前
执行 helper(root->left) 之后，执行 helper(root->right) 之前
执行 helper(root->right) 之后

我们定义一个 state 结构体，用于充当存储函数局部变量和执行位置。改写后函数如下：

typedef struct state {
    TreeNode* node;
    int stage;
};
vector<int> inorderTraversal(TreeNode* root) {
    stack<state> s;
    vector<int> ret;
    s.push({root, 0});
    while (!s.empty()) {
        state top = s.top();
        s.pop();
        if (top.stage == 0) {                  // 状态 1
            if (top.node == nullptr) {
                continue;
            }
            s.push({top.node, 1});             // 将本次调用栈状态修改为 1 再推回去（其实就不应该出栈）
            s.push({top.node->left, 0});       // 新开的函数调用栈 helper(root->left)
        } else if (top.stage == 1) {           // 状态 2
            s.push({top.node, 2});             // 将本次调用栈状态修改为 2 再推回去（其实就不应该出栈）
            ret.push_back(top.node->val);
            s.push({top.node->right, 0});      // 新开的函数调用栈 helper(root->right)
        } else if (top.stage == 2) {           // 状态 3
            // do nothing
        }
    }
    return ret;
}

cuda性能分析工具nsys使用

安装 nsys 命令行工具

安装 Nsight Systems 可视化工具

模型打印

模型图解

attribute ((visibility(“”)))

attribute ((visibility(“default”)))

attribute ((visibility(“hidden”)))

其他

概念

实验

准备工作

实验代码如下：

查看节的地址与大小：

查看需要动态链接的符号：

反汇编查看 plt 相关函数：

开始

题外话

总结

为什么要使用内嵌汇编？

基本语法

被改变的操作对象列表

修饰符

占位符

硬件结构

常见优化

重排 Reordering

插入 Invention

删除 Removal

关系术语

sequence-before

synchronizes-with

happens-before

内存序

内存模型

Sequential Consistency 模型

Acquire-Release 模型

Relax 模型

语法格式

1. 引用寄存器前加 %。如

2. 指令长度后缀

3. 立即数前加 $。16 进制数用 0x 开头。如

4. 注释可以用 ! 开头，也可以用 ;

5. 操作数顺序

6. 数据声明

7. 文件组成

8. 寻址方式

常见指令

1. mov 用于将源操作数移动到目的操作数

2. add 用于将源操作数加给目的操作数

3. sub 用于将两个数相减

4. inc 用于加一

5. dec 用于减一

6. push 用于将数据压入栈

7. pop 用于将数据出栈

8. jmp 跳转

9. mul 乘法

10. div 除法

11. and 按位与

12. or 按位或

13. xor 按位异位

14. shl 和 sal 位左移

15. shr 和 sar 位右移

16. lea 装载有效地址

17. call 函数调用

18. ret 函数返回

19. test 与运算并设置标志寄存器

20. cmp 比较操作数大小

21. rep 重复执行指令直到某一条件

22. lock 锁定总线

23. xadd 交换两个操作数值，使他们相加

24. nop 空操作

25. hlt 使处理器暂停直到收到中断信号

26. xchg 交换两个操作数的值

27. cld 清除方向寄存器（DF）

28. movsb 移动字符串

29. scasb 查找字符

30. cli 禁用所有中断

问题

答案