简介

• Pass就是“遍历一遍IR，可以同时对它做一些操作”的意思。翻译成中文应该叫“传递”。
• 在实现上，LLVM的核心库中会给你一些 Pass类 去继承。你需要实现它的一些方法。
  • ModulePass , CallGraphSCCPass, FunctionPass , or LoopPass, or RegionPass
• 最后使用LLVM的编译器会把它翻译得到的IR传入Pass里，给你遍历和修改。

作用

1. 插桩： 在Pass遍历LLVM IR的同时，自然就可以往里面插入新的代码。
2. 机器无关的代码优化：编译原理一课说：IR在被翻译成机器码前会做一些机器无关的优化。 但是不同的优化方法之间需要解耦，所以自然要各自遍历一遍IR，实现成了一个个LLVM Pass。 最终，基于LLVM的编译器会在前端生成LLVM IR后调用一些LLVM Pass做机器无关优化， 然后再调用LLVM后端生成目标平台代码。
3. 静态分析： 像VSCode的C/C++插件就会用LLVM Pass来分析代码，提示可能的错误 (无用的变量、无法到达的代码等等)。

理解 llvm Pass

理解Pass API

　　Pass类是实现优化的主要资源。然而，我们从不直接使用它，而是通过清楚的子类使用它。当实现一个Pass时，你应该选择适合你的Pass的最佳粒度，适合此粒度的最佳子类，例如基于函数、模块、循环、强联通区域，等等。常见的这些子类如下：

• ModulePass：这是最通用的Pass；它一次分析整个模块，函数的次序不确定。它不限定使用者的行为，允许删除函数和其它修改。为了使用它，你需要写一个类继承ModulePass，并重载runOnModule()方法。
• FunctionPass：这个子类允许一次处理一个函数，处理函数的次序不确定。这是应用最多的Pass类型。它禁止修改外部函数、删除函数、删除全局变量。为了使用它，需要写一个它的子类，重载runOnFunction()方法。
• BasicBlockPass：这个类的粒度是基本块。FunctionPass类禁止的修改在这里也是禁止的。它还禁止修改或者删除外部基本块。使用者需要写一个类继承BasicBlockPass，并重载它的runOnBasicBlock()方法。

　　被重载的入口函数runOnModule()、runOnFunction()、runOnBasicBlock()返回布尔值false，如果被分析的单元（模块、函数和基本块）保持不变，否则返回布尔值true。

Pass的执行顺序/依赖

• ChatGPT说默认顺序是：FunctionPass -> Module Pass -> LoopPass ?
• 当然我们是可以修改插入Pass的执行顺序的。

1
2
3
4
5
6
7
8
9
10
11
12
13
14

char PIMProf::AnnotationInjection::ID = 0;
// 注册 llvm pass
static RegisterPass<PIMProf::AnnotationInjection> RegisterMyPass(
    "AnnotationInjection", "Inject annotators to uniquely identify each basic block.");

static void loadPass(const PassManagerBuilder &,
                           legacy::PassManagerBase &PM) {
    PM.add(new PIMProf::AnnotationInjection());
}

// Ox 的代码 llvm pass 在EP_OptimizerLast 位置load
static RegisterStandardPasses clangtoolLoader_Ox(PassManagerBuilder::EP_OptimizerLast, loadPass);
// O0 的代码 llvm pass EP_EnabledOnOptLevel0 位置load
static RegisterStandardPasses clangtoolLoader_O0(PassManagerBuilder::EP_EnabledOnOptLevel0, loadPass);

流程

1. 编写LLVM pass代码
2. 配置编译环境(cmake or make)
3. 运行(opt or clang)

1 代码框架

最简单框架hello.cpp如下，注意Important一定需要：

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

#include "llvm/Pass.h"
#include "llvm/IR/Function.h"
#include "llvm/Support/raw_ostream.h"
#include "llvm/IR/LegacyPassManager.h"
#include "llvm/Transforms/IPO/PassManagerBuilder.h"

using namespace llvm;

namespace {
 // Important
  struct Hello : public FunctionPass {
    static char ID;
    Hello() : FunctionPass(ID) {}
 // Important
    bool runOnFunction(Function &F) override {
      errs() << "Hello: ";
      errs().write_escaped(F.getName()) << '\n';
      return false;
    }
  };
}

char Hello::ID = 0;

// Important:Register for opt
static RegisterPass<Hello> X("hello", "Hello World Pass");

// Important:Register for clang
static RegisterStandardPasses Y(PassManagerBuilder::EP_EarlyAsPossible,
  [](const PassManagerBuilder &Builder, legacy::PassManagerBase &PM) {
    PM.add(new Hello());
  });

2 编译动态库

使用cmake

参考官方文档。

An example of a project layout is provided below.

1
2
3
4
5
6
7
8

<project dir>/
    |
    CMakeLists.txt
    <pass name>/
        |
        CMakeLists.txt
        Pass.cpp
        ...

Contents of <project dir>/CMakeLists.txt:

1
2
3
4
5
6
7

find_package(LLVM REQUIRED CONFIG)

separate_arguments(LLVM_DEFINITIONS_LIST NATIVE_COMMAND ${LLVM_DEFINITIONS})
add_definitions(${LLVM_DEFINITIONS_LIST})
include_directories(${LLVM_INCLUDE_DIRS})

add_subdirectory(<pass name>)

Contents of <project dir>/<pass name>/CMakeLists.txt:

1

add_library(LLVMPassname MODULE Pass.cpp)

运行cmake编译。产生LLVMPassname.so文件

1
2

mkdir build && cd build
cmake .. && make

使用命令行

请阅读知乎的文章

3 使用

opt加载Pass

1
2

clang -c -emit-llvm main.c -o main.bc # 随意写一个C代码并编译到bc格式
opt -load path/to/LLVMHello.so -hello main.bc -o /dev/null

把源代码编译成IR代码，然后用opt运行Pass实在麻烦且无趣。

clang加载Pass

1
2
3

clang -Xclang -load -Xclang path/to/LLVMHello.so main.c -o main
# or
clang++ -Xclang -load -Xclang ./build/hello/libLLVMPassname.so test.cpp -o main

实践

插入代码

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

void InjectSimMagic2(Module &M, Instruction *insertPt, uint64_t arg0, uint64_t arg1, uint64_t arg2)
{
    LLVMContext &ctx = M.getContext();
    std::vector<Type *> argtype {
        Type::getInt64Ty(ctx), Type::getInt64Ty(ctx), Type::getInt64Ty(ctx)
    };
    FunctionType *ty = FunctionType::get(
        Type::getVoidTy(ctx), argtype, false
    );
    // template of Sniper's SimMagic0
    InlineAsm *ia = InlineAsm::get(
        ty,
        "\tmov $0, %rax \n"
        "\tmov $1, %rbx \n"
        "\tmov $2, %rcx \n"
        "\txchg %bx, %bx\n",
        "imr,imr,imr,~{rax},~{rbx},~{rcx},~{dirflag},~{fpsr},~{flags}",
        true
    );
    Value *val0 = ConstantInt::get(IntegerType::get(ctx, 64), arg0);
    Value *val1 = ConstantInt::get(IntegerType::get(ctx, 64), arg1);
    Value *val2 = ConstantInt::get(IntegerType::get(ctx, 64), arg2);
    std::vector<Value *> arglist {val0, val1, val2};
    CallInst::Create(
            ia, arglist, "", insertPt);
}

这段代码使用内联汇编嵌入到 LLVM IR 中，指令如下：

1
2
3
4

mov $0, %rax
mov $1, %rbx
mov $2, %rcx
xchg %bx, %bx

其中：

• mov $0, %rax 将立即数 arg0 装载到通用寄存器 %rax 中。
• mov $1, %rbx 将立即数 arg1 装载到通用寄存器 %rbx 中。
• mov $2, %rcx 将立即数 arg2 装载到通用寄存器 %rcx 中。
• xchg %bx, %bx 是一条无操作指令，用于保证该汇编代码的原子性。

打印每个BBL内的汇编指令

由于直接打印的是llvm IR的表示，想要打印特定架构比如x86的汇编代码，其实需要进行llvm后端的转换。（取巧，可执行文件反汇编，然后根据插入的汇编桩划分）

1

需要进一步的研究学习

暂无

遇到的问题

暂无

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

复现PIMProf论文时，用到了使用 llvm pass来插入特殊汇编

参考文献

https://www.llvm.org/docs/WritingAnLLVMPass.html

https://zhuanlan.zhihu.com/p/122522485

GNU汇编语法

伪指令

• 指示（Directives）: 以点号开始，用来指示对编译器，连接器，调试器有用的结构信息。指示本身不是汇编指令。

• 语法：.type name , description
  • description取值如下：
    • %function 表示该符号用来表示一个函数名
    • %object 表示该符号用来表示一个数据对象

至于其它的指示你可以忽略。

实践：阅读汇编文件

从最简单的C文件入手

1
2
3

int main(){
 return 0;
}

运行gcc -S -O3 main.c -o main.s，得到main.s文件

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

 .file "simple.cpp"
 .text
 .section .text.startup,"ax",@progbits
 .p2align 4
 .globl main
 .type main, @function
main:
.LFB0:
 .cfi_startproc
 endbr64
 xorl %eax, %eax
 ret
 .cfi_endproc
.LFE0:
 .size main, .-main
 .ident "GCC: (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0"
 .section .note.GNU-stack,"",@progbits
 .section .note.gnu.property,"a"
 .align 8
 .long  1f - 0f
 .long  4f - 1f
 .long  5
0:
 .string  "GNU"
1:
 .align 8
 .long  0xc0000002
 .long  3f - 2f
2:
 .long  0x3
3:
 .align 8
4:

• 下面回答来自ChatGPT-3.5，暂时没有校验其可靠性(看上去貌似说得通)。

section

• .section .rodata.str1.1,"aMS",@progbits,1
  • rodata.str1.1是一个标号（label）, 意思是只读数据段的字符串常量
  • aMS是一个属性值：
    • 可分配的（allocatable），即程序运行时需要动态分配空间才能分配该代码段，
    • 不可执行 （M），
    • 数据的类型为串（S）
    • 其余属性值：对齐方式的通常为 b（byte对齐），w（word对齐），或者其他更大的对齐单位，例如 d（double word对齐）。
  • @progbits: 表示该段的类型是程序数据段（PROGBITS），这种类型的段包含程序的代码和数据。
  • 1: 表示该段的对齐方式是2^1 = 2个字节（按字节对齐）。如果不写这个数字，默认对齐到当前机器的字长。
• .section .text.startup,"ax",@progbits 其中ax表示该段是可分配的（allocatable）和可执行的（executable）。
• “.section .note.GNU-stack“指令用于告诉链接器是否允许在堆栈上执行代码。
• “.section .note.gnu.property“指令用于指定一些属性，这里是一个GNU特性标记。

汇编的入口

• 汇编的执行流程：入口函数在哪里
• 入口函数在该文件中的名称为“main”，定义于“.text.startup” section，其首地址为“.globl main”。

1
2
3
4

.section .text.startup,"ax",@progbits
.p2align 4
.globl main
.type main, @function

构造函数

• 为了确保这些初始化操作可以在程序启动时正确执行，编译器将把这些构造函数和析构函数的调用代码打包成若干个函数，统一放在名字为“_GLOBAL__sub_I_xxx”的section中。
  • 因为在C++程序编译后的二进制文件中，全局变量、静态变量和全局对象等信息都需要进行初始化操作，包括构造函数（初始化对象）和析构函数（清理对象）。
  • 在这段汇编代码中，也就是那个”_GLOBAL__sub_I_main”函数，它是C++全局变量和静态变量的构造函数，它调用了预初始化函数 “ios_base::Init()“，并注册了一个在程序退出时调用的析构函数 “__cxa_atexit“。
• 在”.init_array“ section中，定义了一个”_GLOBAL__sub_I_main”的地址，这是在程序启动时需要调用的所有C++全局和静态对象的初始化函数列表，编译器链接这个列表并在程序启动时依次调用这些初始化函数。
• 总之，这两个section的存在是为了保证C++全局变量和静态变量的正确初始化。

其中四条指令都定义了一些符号或变量，并分配了一些内存空间，这些在程序里的意义如下：

1. “.quad _GLOBAL__sub_I_main“:

在程序启动时，将调用所有全局静态对象的构造函数。这些构造函数被放在一个名为”_GLOBAL__sub_I_xxx”的section中，而每个section都是由一个指向该section所有对象的地址列表所引用。这里的”.quad _GLOBAL__sub_I_main”是为了将”_GLOBAL__sub_I_main”函数的地址添加到该列表中。

2. “.local _ZStL8__ioinit“:

这条指令定义了一个本地符号”_ZStL8__ioinit”，它表示C++标准输入输出的初始化过程。由于该符号是一个本地符号，所以只能在编辑该文件的当前单元中使用该符号。

3. “.comm _ZStL8__ioinit,1,1“:

这条指令定义了一个名为”_ZStL8__ioinit”的未初始化的弱符号，并为该符号分配了1个大小的字节空间。这个弱符号定义了一个C++标准输入输出部分的全局状态对象。在全用动态库时，不同的动态库可能有自己的IO状态，所以为了确保C++输入输出的状态正确，需要为其指定一个单独的段来存储这些状态数据。在这里，”.comm _ZStL8__ioinit,1,1”将会为”_ZStL8__ioinit”符号分配一个字节大小的空间。

4. “.hidden __dso_handle“:

这条指令定义了一个隐藏的符号 “__dso_handle”。这个符号是一个链接器生成的隐式变量，其定义了一个指向被当前动态库使用的全局数据对象的一个指针。该符号在被链接进来的库中是隐藏的，不会被其他库或者main函数本身调用，但是在main返回后，可以用来检查库是否已经被卸载。

末尾的元数据

这段代码是一些特殊的指令和数据，主要是用于向可执行文件添加一些元数据（metadata）。这些元数据可能包含各种信息，如调试信息、特定平台的指令集支持等等。

具体来说：

• “.long”指令用于定义一个长整型数值，这里用来计算地址之间的差值。
  • 例如，第一行”.long 1f - 0f“建立了一个长整型数值，表示”1:”标签相对于当前指令地址（即0f）的偏移量。偏移量可以用来计算标签对应的指令地址，从而可用于跳转或计算指针偏移量。
  • “4f - 1f“，即”4:”标签相对于”1:”标签的偏移量；
• “.long 0xc0000002“表示这是一个特殊的属性标记，标识这个文件可以在Linux平台上执行。它是用来告诉操作系统这个程序是用特定指令集编译的。
• “.long 0x3“表示另一个属性标记，表示这个文件可以加载到任意地址。

总之，这些元数据可能对程序运行起到关键作用，但在大多数情况下可能都没有明显的作用，因此看起来没有用。

比较汇编的debugging symbols

执行gcc -S -g testBigExe.cpp -o testDebug.s，对比之前的汇编文件，由72行变成9760行。

.loc

1
2
3
4

.LBE32:
 .file 3 "/usr/include/c++/9/bits/char_traits.h"
 .loc 3 342 2 is_stmt 1 view .LVU4
 .loc 1 5 11 is_stmt 0 view .LVU5

• 第一行：.loc 3 342 2 表示当前指令对应的源代码文件ID为3，在第342行，第2列（其中第1列是行号，第2列是第几个字符），同时is_stmt为1表示这条指令是语句的起始位置。
• 第二行：.loc 1 5 11 表示当前指令对应的源代码文件ID为1，在第5行，第11列，同时is_stmt为0表示这条指令不是语句的起始位置。
• view .LVU4 表示当前指令所处的作用域（scope）是.LVU4。作用域是指该指令所在的函数、代码块等一段范围内的所有变量和对象的可见性。在这个例子中，.LVU4 是一个局部变量作用域，因为它是位于一个C++标准库头文件中的一个函数的起始位置。

debug section

新增的这些 section 存储了 DWARF 调试信息。DWARF（Debugging With Attributed Record Formats）是一种调试信息的标准格式，包括代码中的变量、类型、函数、源文件的映射关系，以及代码的编译相关信息等等。

具体来说，这些 section 存储的内容如下：

• .debug_info：包含程序的调试信息，包括编译单元、类型信息、函数和变量信息等。
• .debug_abbrev：包含了 .debug_info 中使用到的所有缩写名称及其对应的含义，用于压缩格式和提高效率。
• .debug_loc：存储每个程序变量或表达式的地址范围及其地址寄存器、表达式规则等信息。在调试时用来确定变量或表达式的值和范围。
• .debug_aranges：存储简化版本的地址范围描述，允许调试器加速地定位代码和数据的位置。
• .debug_ranges：存储每个编译单元（CU）的地址范围，每个范围都是一个有限开区间。
• .debug_line：存储源代码行号信息，包括每行的文件、行号、是否为语句起始位置等信息。
• .debug_str：包含了所有字符串，如文件名、函数名等，由于每个调试信息的数据都是字符串，因此这是所有调试信息的基础。

需要注意的是，这些 section 中的信息是根据编译器的配置和选项生成的，因此不同编译器可能会生成略有不同的调试信息。

需要进一步的研究学习

• 在编译的过程中，哪个阶段 label会变成真实执行地址

遇到的问题

暂无

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

参考文献

https://www.cnblogs.com/zuofaqi/articles/12853734.html

https://www.cnblogs.com/zuofaqi/articles/12853734.html

常见的web服务器

常见的web服务器有Apache、nginx、IIS

1. Apache
   1. Apache音译为阿帕奇, 是全世界最受欢迎的web服务器，因其快速、可靠并且可通过简单的API扩充，能将Python\Perl等解释器部署在其上面等优势，受到广泛的关注与使用。
   2. 但是现在用的人少了，而且性能没nginx好
2. nginx
   1. Apache的致命缺陷就是在同时处理大量的（一万个以上）请求时，显得有些吃力，所以“战斗民族”的人设计的一款轻量级的web服务器——nginx, 在高并发下nginx 能保持比Apache低资源低消耗高性能 ，
3. IIS
   1. iis是Internet Information Services的缩写，意为互联网信息服务，是由微软公司提供的基于运行Microsoft Windows的互联网基本服务,

合法使用

虽然 Ryujinx 模拟器项目本身是开源免费且合法的，但它默认情况下并不能直接运行市面上发行的各种商业游戏，因为它并不包含 Switch 系统固件，也没有游戏 ROM。

而按照国外的法规，如果你用户购买了主机和游戏，将其内容 DUMP (提取) 出来自己使用是合法的。

所以，无论是 Ryujinx 还是 Yuzu 等模拟器，想要开玩都需要先完成

• 安装系统固件和
• prod.keys 密钥等步骤。
  • 如果游戏要求最新的key和firmware你需要去更新

基本流程

1. 下载 Ryujinx 模拟器主程序
2. 下载 prod.keys 密钥文件以及 Switch 的系统固件 (Firmware)
3. 自行网上搜索下载你喜欢的任天堂游戏文件，支持 .NSP 或者 .XCI 格式

密钥文件以及 Switch 的系统固件

• 通过decode 知道 网站中编码的下载地址
  • 为key

Ryujinx (龙神)模拟器

率先支持了ARM和苹果 M 系列芯片

简易步骤教程：

1. 下载模拟器对应github
2. 将 Ryujinx 模拟器主程序解压到「不包含中文的路径」下。
3. 首次启动 Ryujinx 模拟器后，会提示找不到 key 文件的错误
   1. 点击菜单 文件 (File) → 打开 Ryujinx 文件夹 (Open Ryujinx Folder)，会弹出模拟器数据所在的目录。
   2. 将 prod.keys 文件放进到 Ryujinx 目录中的 system 文件夹里，重启模拟器
4. 放置好 keys 文件之后就开始安装固件 (Firmware) 了，
   1. 点击菜单 工具 (Tools)→安装固件 (Install Firmware)→从 XCI 或 ZIP 安装固件 (Install a firmware from XCI or ZIP)，选择你下载到的 Firmware 压缩包(不需解压)，
   2. 模拟器就会开始安装，直到显示安装完成即可。
5. 此时已经可以运行 Switch 游戏了，
   1. 点击菜单 选项 (Options) → 设置 (Settings)，在 用户界面 (General) → 游戏目录 (Game Directories) 下点击 添加 (Add) 来「添加一个游戏 ROM 文件存放目录」。
   2. 此后，模拟器会自动加载出存放在这些文件夹里的游戏列表。

安装nsp DLC文件

不是

而是

性能优化

YUZU(柚子)模拟器

yuzu，奶刃2好像都是用这个

密钥缺失

将你原来的User文件夹和ROM文件夹拖到新版模拟器文件夹的根目录即可。

将 prod.keys 文件放在Yuzu\user\keys

更新密钥和固件

• 下面两个目录的内容都可以删除，然后解压
• keys解压到Yuzu\user\keys
• 固件解压到Yuzu\user\nand\system\Contents\registered

YUZU模拟器使用教程

https://www.playdanji.com/yuzu

YUZU金手指

https://www.playdanji.com/yuzujinshouzhi

YUZU软件升级方法

Early Access版本是需要花钱订阅才能下载的。

github的release直接下载zip后解压替换即可。

YUZU存档位置

游戏右键，打开存档位置。

退出全屏

F11

YUZU 安装 nsp和xci文件

key和中文的问题，建议用之前好的文件

https://zhuanlan.zhihu.com/p/406048136

安装升级补丁nsp文件和DLC

导入NAND文件即可

异度之刃2

打包本体1

30帧720P

https://switch520.com/23050.html

推荐OpenGL模式

v模式，暗场景会过曝。

画质补丁

贴吧老哥的放入目录

1

E:\gamePojie\NaiRen2\A3285 v2.0.2_yuzuEA2077\user\sdmc\atmosphere\contents

但是没什么用。

ini配置原理是，如下图对应配置目录

放入如下修改的ini文件来修改画质

如下图成功

贴吧10楼：刚试了下，我把属性的mod选项关掉，然后把0100F3400332C000的画质mod删掉，效果一样有，所以效果应该只能在0100F3400332D001\画质mod\romfs\monolib\shader下的lib_nx.ini里改，其他的都没用

1
2
3
4
5
6
7
8
9
10
11
12

red_sclX=2.0
red_sclY=2.0
red_hdsclX=2.0
red_hdsclY=2.0


red_Auto=on
red_AtMaxX=2.0
red_AtMaxY=2.0
red_AtMinX=2.0
red_AtMinY=2.0
red_AtRate=100.0

2.0就是1440p，1.0就是原版720p。你们可以试试改其他的

帧数补丁

60帧补丁实际效果远没有60帧而且一堆副作用，不用浪费时间了

bug花屏

按照B站设置，主要改了GLSL

塞尔达-旷野之息

游侠论坛的cemu模拟的效果就不错

塞尔达-王国之泪

• yuzu 1414以上
• keys firmware 16.0.2以上

龙神模拟器，会经常闪退，暂时不知道解决办法(Cache PPTC rebuild?)。yuzu没有闪退的问题

需要进一步的研究学习

暂无

遇到的问题

暂无

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

之前买了正版的switch，游戏也入了两三千。旷/荒野之息，奥德赛，奶刃2都通关了。可惜被妈妈没收了~

想研究一下，PC模拟，记录踩坑过程

参考文献

BT简述

BitTorrent (简称 BT) 协议是和点对点（point-to-point）的协议程序不同，它是用户群对用户群（peer-to-peer, 或简写为 P2P) 传输协议, 它被设计用来高效地分发文件 (尤其是对于大文件、多人同时下载时效率非常高)。该协议基于HTTP协议，属于TCP/IP应用层。

将文件划分成多块(默认256Kb一块)，每块可以从网络中不同的用户的BT客户端处并行下载。

1
2
3

BT 下载的文件都是别人上传给你的。
BT 下载速度均来自其他下载同一资源的用户上传速度。
上传的用户越多，你的下载速度越快，相反没用户上传你就没有下载速度。

比特彗星，包括其他 BT 软件（迅雷除外，迅雷不是会员会限速，高速通道下载提高的速度一部分就是接触限速后获得的）都不会限制下载速度。

BT分享规则

与迅雷不同，BT旨在“人人为我，我为人人”。用户和用户之间对等交换自己手中已有的资源。如果任何一方试图白嫖另外一方的资源，而自己不愿意上传自己的资源，那么那方就会被人视作吸血者而被踢出这个交换，下场是没有人会愿意和你交换数据，你的下载速度也就归零。

如果把上传速度限制为了10KB/s，10KB/s是BitComet上传最低限速，很大时候就这10KB会被包含DHT查询、向Tracker服务器注册，连接用户所产生的上传全部占满。在下载种子的时候，其他用户连上你是只能拿到1～2KB/s甚至一点都没有的。

现在的BT下载客户端都可以做到智能反吸血，所以基本想和交换数据的用户都把你当作Leecher（吸血鬼）Ban（封禁）处理了，故没有下载速度不足为奇。

一般来说，只要预留50KB/s的上传给其他网页浏览、聊天就可以了，在下载时应该尽量把上传留给那些和你交换资源的用户，这样才不会被他们视作你在吸血进而屏蔽你。

如果上传不足，就应该主动限制自己的下载速度，否则单位时间下载量远超过上传量反而会遭来更多的屏蔽，对下载速度提升更加不利。

BT基本流程

.torrent 种子文件本质上是文本文件，包含Tracker信息和文件信息两部分。Tracker信息主要是BT下载中需要用到的Tracker服务器的地址和针对Tracker服务器的设置。

1. 下载时，BT客户端首先解析种子文件得到Tracker地址，然后连接Tracker服务器。
2. Tracker服务器回应下载者的请求，提供下载者其他下载者（包括发布者）的IP。
3. 下载者再连接其他下载者，根据种子文件，两者分别告知对方自己已经有的块，然后交换对方所没有的数据。
   1. 此时不需要其他服务器参与，分散了单个线路上的数据流量，因此减轻了服务器负担。
4. 下载者每得到一个块，需要算出下载块的Hash验证码与种子文件中的对比，如果一样则说明块正确，不一样则需要重新下载这个块。这种规定是为了解决下载内容准确性的问题。

常规部署

• 安装qBittorrent-nox, tmux下运行，在8080端口挂载WebUI
• 安装qbittorrent, 用户运行qbittorrent, x11弹出应用窗口
  • 问题：
    • 怎么维持窗口？sudo XAUTHORITY=/home/qcjiang/.Xauthority qbittorrent
    • 关于写文件权限，如何写网络硬盘
    • node5 上传很快, 网络原因？种子原因？(不是，是因为网络硬盘，所以下载多少要占用多少上传)

docker部署

以qBit的docker为例，参考linuxsever的docker-compose如下：(qBit相对于Transmission有多线程IO的优势) 也可以使用其余docker镜像

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

---
version: "2.1"
services:
  qbittorrent:
    image: lscr.io/linuxserver/qbittorrent:latest
    container_name: qbittorrent
    environment:
      - PUID=0 # 这里是root 如果想以其他用戶A修改文件，可以改成其他用户的UID。通過id A 查看
      - PGID=0
      - TZ=Etc/UTC
      - WEBUI_PORT=8080
    volumes:
      - /addDisk/DiskNo4/qBit:/config
      - /addDisk/DiskNo4/bt:/downloads

    network_mode: host
    restart: unless-stopped

默认账号 admin 默认密码 adminadmin
然后通过webUI http://222.195.72.218:8080/管理。

如果不想网络通过wireguard，而是本地可以如下设置

同一台机器实现两个账号做种

• 思路：两个docker，同一个网络出口
• 問題：
  • qBit要有权限写文件
  • 两个docker兼容性：
    • WebUI有bug，不一定会显示。重启刷新即可
    • 其次有一个docker会没有网络。
      • 这是由于端口随机错误导致的，会导致下面连接状态显示火焰。换端口刷新即可解决。
      • 
• 关于封号
  • 由于国内环境下载一般都是大内网。同一个网络出去应该没有问题。
  • 大多数做种是通过ipv6，会被检测出同一个机器多个账户做种，导致封号
  • 还有关于盒子刷上传，一方面通过速度，另一方面由于盒子的ip是固定的，所以会被检测出重复导致封号。

测速

1. 先用 https://www.speedtest.net 测速
   1. 
2. 考虑热门种子测试 http://releases.ubuntu.com/19.10/ubuntu-19.10-desktop-amd64.iso.torrent
   1. 
   2. 没通过代理能找到的用户变少，速度也变慢了。
3. 如果跑不到网络上限，可能和软件设置有关

上传速度

m站刷上传的时候，发现基本都是对方基本都是通过ipv6下载

uTP是一种基于UDP的协议，它可以根据网络拥塞情况自动调节传输速度，从而减少对其他网络应用的影响。

BT连接是一种基于TCP的协议，它可以保证数据的完整性和可靠性，但是可能会占用较多的网络带宽和资源。

在qBittorrent中，标志U K E P分别表示以下含义：

1
2
3
4

U：表示你正在上传数据给对方，或者对方正在从你那里下载数据。
K：表示对方支持uTP协议，即基于UDP的传输协议。
E：表示对方使用了加密连接，即通过加密算法保护数据的安全性。
P：表示对方使用了代理服务器或VPN服务，即通过第三方网络隐藏自己的真实IP地址。

PT设置

• 虽然说PT下载，客户端要关闭DHT,PeX, LSD。
• 但是其实种子是默认关闭的，无序额外设置，北洋和M站一样。
• 

基本概念

Tracker

收集下载者信息的服务器，并将此信息提供给其他下载者，使下载者们相互连接起来，传输数据。

种子

指一个下载任务中所有文件都被某下载者完整的下载，此时下载者成为一个种子。发布者本身发布的文件就是原始种子。

做种

发布者提供下载任务的全部内容的行为；下载者下载完成后继续提供给他人下载的行为。

分享率

上傳資料量 / 下傳資料量的比率,是一種BT的良心度,沒實際作用.(一般为了良心，至少大于1)

长期种子

BitComet的概念，相对于种子任务的上传能够控制。

长效种子就是你不开启任务做种，只要你启动了比特彗星，软件挂后台，当有其他用户也是用比特彗星下载你列表里的存在的文件时候就会被认为是长效种子 。

DHT

.DHT全称叫分布式哈希表(Distributed Hash Table)，是一种分布式存储方法。在不需要服务器的情况下，每个客户端负责一个小范围的路由，并负责存储一小部分数据，从而实现整个DHT网络的寻址和存储。新版BitComet允许同行连接DHT网络和Tracker，也就是说在完全不连上Tracker服务器的情况下，也可以很好的下载，因为它可以在DHT网络中寻找下载同一文件的其他用户。

类似Tracker的根据种子特征码返回种子信息的网络。

在BitComet中，无须作任何设置即可自动连接并使用DHT网络，完全不需要用户干预。

用户交换Pex

Peer Exchange (PEX)， 每个peer客户端的用户列表，可以互相交换通用。可以将其理解为“节点信息交换”。前面说到了 DHT 网络是没有中心服务器的，那么我们的客户端总不能满世界去喊：“我在下载这个文件，快来连我吧.”（很大声）。所以就通过各个 BT 客户端自带的节点去同步路由表实现 DHT 网络连接。

本地用户发现

LSD（LPD）就是本地网络资源，内网下载，没什么几把用的东西，可能学校等私有网络好使

ISP

網絡業務提供商(Internet Service Provider，簡稱ISP)，互聯網服務提供商，即向廣大用户綜合提供互聯網接入業務、信息業務、和增值業務的電信運營商。

反吸血机制

基本原理

1. 根据流量： 默认设置为120秒，持续对某个peer产生上传，并且从该peer用户获取的下载流量没有超过1KB文件（1024字节）大小，即拉黑该peer，预警颜色为黄色，合法为绿色，红色为封禁。
2. 可组合检测指定客户端进行反吸血，比如说指定屏蔽qbittorrent、utorrent等吸血客户端选择（可多选客户端，可对下载任务和做种任务生效）
3. 可组合检测客户端连接端口号进行反吸血，比如说指定屏蔽15000迅雷X版本客户端等吸血端口（可多选端口号，可对下载任务和做种任务生效）
4. 可组合检测客户端连接peer_id标志符进行反吸血屏蔽，例如屏蔽XL0018客户端等吸血标志符（可多选标志符，可对下载任务和做种任务生效）

高级设置

1
2
3
4
5

bittorrent.anti_leech_min_byte
设定反吸血保护流量：要求对方在指定时间（秒）内需要上传的最少流量（byte）， 取值范围：1-10000。

bittorrent.anti_leech_min_stable_sec
设定反吸血保护时间：指定与对方连接多长时间（秒）后开始检查流量（byte），取值范围：1-10000。

常见问题

需要软件开着吗？

需要

原文件改名或者移动，还会上传吗？

文件下载后不能移动，不能删除，不能重命名（但可以在软件内改）。 一但BT 软件找不到文件，或删除了任务，就无法做种上传了。

晚上避免上传

可以在Bitcomet高级设置里设置时段限速

对硬盘损害大吗？

分享上传也需要频繁读取硬盘。

以Bitcomet为例，该软件就是通过磁盘缓存技术减小频繁随机读写对硬盘的损伤。

磁盘缓存就是利用物理内存作为缓冲，将下载下的数据先存放于内存中，然后定期的一次性写入硬盘，以减少对硬盘的写入操作，很大的程度上降低了磁盘碎片。

因为通常我们设置内存（磁盘缓存）为每任务XX兆，意味着，这个缓冲区可以存放数兆甚至几十兆的“块”，基本上可以杜绝碎片了。

现在BT软件都是自动设置缓存的，它是根据你物理内存的大小分配的。

注意设置

1. 设置了“反吸血”，应对迅雷
2. 校园网设置9，不限制P2P
   1. 国内各省份不同运营商限速策略（QOS）

需要进一步的研究学习

路由器下载？

参考文献

简介

• Cloudflare 作为全球最大的网络服务提供商，全球大约有30%的网络服务由它提供。
• WARP是Cloudflare提供的免费的，没有带宽上限或限制的VPN服务，可以让你在任何地方访问互联网，而不会受到地域限制。
  • WARP软件是基于wireguard魔改的。
  • WARP有安卓和Windows的客户端，但是使用人数过多，体验并不好
  • Linux下通过WARP代理能实现20MB/s的下载速度。

WARP on Linux

安装配置

• 缘由：WARP下的PT做种快得多。(不是，是因为网络硬盘，所以下载多少要占用多少上传)
• 参考教程。
• 脚本和所需文件在Z:\shaojiemike\Documents\文献\计算机网络目录下。这里先使用fjw的脚本。

1. 通过注册脚本register.py，获得私钥和分配的ip
2. 配置wg。其中Endpoint端口可从官方文档中找到，默认的2408很可能被封。WARP with firewall · Cloudflare Zero Trust docs

   1 2 3 4 5 6 7 8 9 10

   [Interface] PrivateKey = xxxx #register.py的私钥 Address = xxx/32,xxx/128 #register.py的ipv4和ipv6 Table = off # off禁止wg修改路由表 [Peer] PublicKey = bmXOC+F1FxEMF9dyiK2H5/1SUtzH0JuVo51h2wPfgyo= AllowedIPs = 0.0.0.0/0,::/0 Endpoint = [2606:4700:d0::a29f:c001]:500 #2408, 1701, 500, 4500 PersistentKeepalive = 25

3. warp使用的不是标准的wg协议，root下运行，需要通过一个nft脚本main.sh修改包的3个字节。
   1. 安装nft apt-get install nftables
   2. 需要/etc/default/warp-helper文件填写对应的
   3. ROUTING_ID对应register.py的ROUTING_ID。注意三个数之间没空格
   4. UPSTREAM对应wg-conf里Endpoint。比如：

   5.   ROUTING_ID=11,45,14
        UPSTREAM=[2606:4700:d0::a29f:c001]:500
        
      1
      2
      3
      4
      5
      6
      7
      8
      9
      10
      11
      12
      13
      14
      15
      16
      17
      18
      19
      20
      21
      2. 最后开启路由表,Root权限运行` ip route add default dev warp proto static scope link table default`
      
      ## WARP on OpenWRT 
      
      * 目的：为了防止大量流量通过WARP，导致被官方封禁，所以只在OpenWRT上配置WARP分流github的流量。
      * 实现思路：
        * 运行python脚本，通过github的API获得所有的github域名ip，
        * 使用iptables的warp_out表，将目的地址为github域名ip路由到WARP的虚拟网卡上。
      
      ### WARP Wireguard Establishment
      
      ```bash
      python register.py #自动生成warp-op.conf,warp.conf和warp-helper
      mv warp-helper /etc/default
      # cat main.sh
      # cat warp-op.conf
      vim /etc/config/network #填写warp-op.conf内容，默认只转发172.16.0.0/24来测试连接
      ifup warp #启动warp, 代替wg-quick up warp.conf
      bash main.sh #启动防火墙实现报文头关键三字节修改
      nft list ruleset #查看防火墙，是否配置成功
      wg #查看warp状态，测试是否连接成果
      
      

这时还没创建warp_out路由表，所以还不能通过WARP出数据。

1
2
3
4
5
6
7
8

#/etc/config/network
config interface 'warp'
        option proto 'wireguard'
        option private_key 'wKpvFCOk4sf8d/RD001TF7sNQ39bWnllpqaFf8QnHG4='
        option listen_port '51825'
        list addresses '172.16.0.2/32'
        list addresses '2606:4700:110:8466:d4ea:ffb8:10da:470f/128'
        #option disabled '1' 

然后WebUI点击apply 或者命令行运行ifconfig warp down && ifup

Network planning and design

添加了WARP的网络出口后，路由器不在只是通过WAN出数据。防火墙需要更新：

• 原路返回规则。
  • 针对有公网ip的接口，需要原路返回。
    • 配置来自wan和WARP的信报，使用wan和WARP的路由表，优先级3
    • 来自wan的比如来自外部的ssh，为了防止失联。
    • 来自WARP的比如wget --bind-address=WARP_ip来模拟
  • 内网地址没有必要配置，因为通过内网地址访问host，则dst必然也是内网地址。因此会匹配main中的内网地址规则。
• wan和WARP的路由表内各自走wan和WARP的网卡
• 为了使得原本wg正常运行，10: from all lookup main suppress_prefixlength 1
  • 假如warp_out是defualt规则，该项也是为了防止失联。
• 创建warp_out的空路由表1000: from all lookup warp_out,优先级1000

1
2
3
4
5
6
7
8

root@tsjOp:~/warp# ip rule
0:      from all lookup local
3:      from 114.214.233.141/22 iif eth1 lookup wan
3:      from 172.16.0.2 iif warp lookup warp
10:     from all lookup main suppress_prefixlength 1
1000:   from all lookup warp_out
32766:  from all lookup main
32767:  from all lookup default

填充warp_out路由表

1
2
3

cd ip_route
mv ../github_ipv4.txt .
python fill_ip_table.py --table warp_out --iface warp --p2p -f github_ipv4.txt

对所有github域名的ip执行类似ip ro add 192.30.252.0/22 dev warp proto static table warp_out操作。

测试

1
2
3

mtr www.github.com
ssh -vT git@github.com
git clone https://github.com/llvm/llvm-project

添加到启动项

修改/etc/rc.local

1
2
3
4
5
6
7
8
9

# Put your custom commands here that should be executed once
# the system init finished. By default this file does nothing.


sleep 30 && cd /root/warp/ip_route && python fill_ip_table.py --table warp_out --iface warp --p2p -f github_ipv4.txt

/root/warp/main.sh #重新添加防火墙

exit 0

WARP on Windows

基于1.1.1.1 的安装windows版本直接白嫖

需要进一步的研究学习

暂无

遇到的问题

暂无

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

参考文献

https://gist.github.com/iBug/3107fd4d5af6a4ea7bcea4a8090dcc7e

glados

区别

Service

历史上，Linux 的启动一直采用init进程。
下面的命令用来启动服务。

1
2
3

$ sudo /etc/init.d/apache2 start
# 或者
$ service apache2 start

这种方法有两个缺点。

一是启动时间长。init进程是串行启动，只有前一个进程启动完，才会启动下一个进程。

二是启动脚本复杂。init进程只是执行启动脚本，不管其他事情。脚本需要自己处理各种情况，这往往使得脚本变得很长。

Systemd

Systemd 就是为了解决这些问题而诞生的。它的设计目标是，为系统的启动和管理提供一套完整的解决方案。

根据 Linux 惯例，字母d是守护进程（daemon）的缩写。 Systemd 这个名字的含义，就是它要守护整个系统。

使用了 Systemd，就不需要再用init了。Systemd 取代了initd，成为系统的第一个进程（PID 等于 1），其他进程都是它的子进程。

Systemd 的优点是功能强大，使用方便，缺点是体系庞大，非常复杂。事实上，现在还有很多人反对使用 Systemd，理由就是它过于复杂，与操作系统的其他部分强耦合，违反”keep simple, keep stupid”的Unix 哲学。

systemctl

systemctl是 Systemd 的主命令，用于管理系统。

systemctl - Control the systemd system and service manager

1
2
3
4
5
6
7
8
9

systemctl is-enabled servicename.service #查询服务是否开机启动
systemctl enable *.service #开机运行服务
systemctl disable *.service #取消开机运行
systemctl start *.service #启动服务
systemctl stop *.service #停止服务
systemctl restart *.service #重启服务
systemctl reload *.service #重新加载服务配置文件
systemctl status *.service #查询服务运行状态
systemctl --failed #显示启动失败的服务

1
2
3
4
5
6
7
8
9
10
11
12

# 查看相关unit
shaojiemike@tsjNas:~$ systemctl|grep wg
  pkg-wg-quick@wg0.service
      loaded active exited  WireGuard via wg-quick(8) for wg0

# 查看已经启动的服务
systemctl list-unit-files --state=enabled|grep wg
pkg-wg-quick@.service                     enabled //由于我删除了wg0，所以.service前没wg0

# 删除服务
sh-4.4#  systemctl disable pkg-wg-quick@.service
Removed symlink /etc/systemd/system/syno-low-priority-packages.target.wants/pkg-wg-quick@wg0.service.

开机启动原理

Systemd 默认从目录/etc/systemd/system/读取配置文件。但是，里面存放的大部分文件都是符号链接，指向目录/usr/lib/systemd/system/，真正的配置文件存放在那个目录。

systemctl enable命令用于在上面两个目录之间，建立符号链接关系。

1
2
3

$ sudo systemctl enable clamd@scan.service
# 等同于
$ sudo ln -s '/usr/lib/systemd/system/clamd@scan.service' '/etc/systemd/system/multi-user.target.wants/clamd@scan.service'

如果配置文件里面设置了开机启动，systemctl enable命令相当于激活开机启动。

与之对应的，systemctl disable命令用于在两个目录之间，撤销符号链接关系，相当于撤销开机启动。

1

> $ sudo systemctl disable clamd@scan.service

配置文件的后缀名，就是该 Unit 的种类，比如sshd.socket。如果省略，Systemd 默认后缀名为.service，所以sshd会被理解成sshd.service。

实例

1
2
3
4
5
6
7
8
9

[Unit]
Description=Webhook receiver for GitHub

[Service]
Type=simple
ExecStart=/usr/local/bin/webhook

[Install]
WantedBy=multi-user.target

1
2

systemctl start nexus.service  #启动服务
systemctl enable nexus.service #设置开机启动

1
2
3
4

# root @ snode0 in /etc/systemd/system [17:57:48]
$ journalctl -u webhook.service
-- Logs begin at Mon 2022-06-06 15:54:50 CST, end at Tue 2022-06-28 17:57:50 CST. --
Jun 28 17:30:53 snode0 systemd[1]: Started Webhook receiver for GitHub.

问题

1
2
3
4
5
6
7
8
9
10

$ systemctl reload webhook.service
==== AUTHENTICATING FOR org.freedesktop.systemd1.manage-units ===
Authentication is required to reload 'webhook.service'.
Multiple identities can be used for authentication:
 1.  Jun Shi (shijun)
 2.  Shaojie Tan (shaojiemike)
Choose identity to authenticate as (1-2): 2
Password:
==== AUTHENTICATION COMPLETE ===
Failed to reload webhook.service: Job type reload is not applicable for unit webhook.service.

Simple 类型不能reload

参考文献

https://blog.csdn.net/qq_40741855/article/details/104984071

https://www.ruanyifeng.com/blog/2016/03/systemd-tutorial-commands.html

Nsight system compute & Graph 的关系

Nsight Systems

All developers should start with Nsight Systems to identify the largest optimization opportunities. Nsight Systems provides developers a system-wide visualization of an applications performance. Developers can optimize bottlenecks to scale efficiently across any number or size of CPUs and GPUs; from large servers to our smallest SoC. For further optimizations to compute kernels developers should use Nsight Compute or to further optimize a graphics workloads, use Nsight Graphics.

Nsight Compute

Nsight Compute is an interactive kernel profiler for CUDA applications. It provides detailed performance metrics and API debugging via a user interface and command line tool. Nsight Compute also provides customizable and data-driven user interface and metric collection that can be extended with analysis scripts for post-processing results.

Nsight Graphics

Nsight Graphics is a standalone application for the debugging, profiling, and analysis of graphics applications on Microsoft Windows and Linux. It allows you to optimize the performance of your Direct3D 11, Direct3D 12, DirectX Raytracing 1.1, OpenGL, Vulkan, and KHR Vulkan Ray Tracing Extension based applications.

Install Nsight local

1. check the perf config To collect thread scheduling data and IP (instruction pointer) samples
   1. cat /proc/sys/kernel/perf_event_paranoid
   2. 如果大于2，临时改变 sudo sh -c 'echo 2 >/proc/sys/kernel/perf_event_paranoid'重启会重置
   3. 永久修改 sudo sh -c 'echo kernel.perf_event_paranoid=2 > /etc/sysctl.d/local.conf'
2. 下载Nsight
   1. 但是单独下载要会员
   2. 下载cuda toolkit,有集成

Nsight System

目标与功能

运行 nsight-sys，可以从整体上看GPU,CPU资源的使用情况，和分辨出热点函数和kernel，但是对于为什么是热点给不出具体分析。

基本使用

勾选了CUDA-trace, GPU Metrics选项

1
2
3
4
5

GPU Metrics [0]: The user running Nsight Systems does not have permission to access NVIDIA GPU Performance Counters on the target device. For more details, please visit https://developer.nvidia.com/ERR_NVGPUCTRPERM
  - API function: NVPW_GPU_PeriodicSampler_GetCounterAvailability(&params)
  - Error code: 17
  - Source function: static std::vector<unsigned char> QuadDDaemon::EventSource::GpuMetricsBackend::Impl::CounterConfig::GetCounterAvailabilityImage(uint32_t)
  - Source location: /dvs/p4/build/sw/devtools/Agora/Rel/DTC_F/QuadD/Target/quadd_d/quadd_d/jni/EventSource/GpuMetricsBackend.cpp:609

Profile 速度

大致2到3倍时间：默认采样率，单独运行52s, Nsight-sys模拟需要135s。

HPC APP : PCIE, GPU DRAM Bandwidth, Warp

GPU Metrics选项能看出 PCIE, GPU DRAM Bandwidth, Warp的使用情况。

Compute Warps in Flight

将鼠标放在上面会有具体的数值或者名称的解释，(正在使用的Warps)

Unallocated Warps in Active SMs

• Definition: This metric represents the number of warps that are not actively executing but are assigned to an active Streaming Multiprocessor (SM).
• Interpretation: In CUDA, SMs are the fundamental processing units on the GPU. Each SM can execute multiple warps concurrently. “Unallocated Warps in Active SMs” indicates the number of warps that are ready to be scheduled on an SM but are currently waiting due to resource contention or other factors. A high number may suggest that there is room for additional work but available resources are not fully utilized.

NVTX

由于没有根据kernel function区分，很难读。为此提供了NVTX来给代码打标签

1

#include <nvToolsExt.h>

1
2
3
4

nvtxRangePush("checkResult"); //nvtxRangePushA,nvtxRangePushW,nvtxRangePushEx 好像都差不多
checkResult<<<dim3(row_num / TPBX, col_num / TPBY, 1), dim3(TPBX, TPBY, 1)>>>(row_num, col_num, result);
cudaDeviceSynchronize(); 
nvtxRangePop();

AI APP: Stable Diffusion XL

具体分析见 Deploy Stable Diffusion to A100

Nsight Compute

• Nsight Systems 就是nvprof的继任者，NVIDIA最新的用于监测 kernel timeline的工具。
• NVIDIA 计算能力7.5及以上的GPU设备(从A100开始)不再支持nvprof工具进行性能剖析，提示使用Nsight Compute作为替代品.

目标与功能

默认kernel模式，会根据 function的调度关系，将程序划分为kernel

1. Summary： 给出in-order执行的每个kernel的参数，时间，资源占用(寄存器，计算访存单元)信息。
   1. Detail: 对于被选择的kernel给出， NV的优化建议
   2. Source：对于被选择的kernel给出， 给出源代码

基本使用

1
2
3

# recommand running under sudo
ncu # 命令行 Nsight Compute CLI(ncu)
ncu-ui # GUI

Profile速度

目测模拟时间慢百倍。

使用Nsight Compute CLI (nv-nsight-cu-cli / ncu) 输出数据

nv-nsight-cu-cli -> ncu

下面是一个使用样例：

1

/usr/local/NVIDIA-Nsight-Compute/nv-nsight-cu-cli -o mnist -f --csv --profile-from-start off /usr/bin/python3 mnist.py

其中-o是为了输出.nsight-cuprof-report文件用于后续的可视化查看，-f为强制覆盖原有文件，–csv可是在console输出除 timeline 以外数据的时候以逗号分隔数据，方便拷贝至csv文件， –profile-from-start的使用方法和Nsight System以及nvprof一样。其余flag选项可见文档。

上面的例子会生成mnist.nsight-cuprof-report文件。

注意

最前面的可执行文件需要绝对路径，如上面的python3需要使用 /usr/bin/python3。
生成过程中可能会产生很大的临时文件（几十G）。如果本次磁盘空间不够，可以设置如下环境变量来调整存储临时文件的地址。没有找到能直接使用 Nsight Compute 修改临时文件地址的方式。

1

export /TMPDIR=/path/for/tmp

ncu与nvprof命令行抓取参数的映射表

https://www.freesion.com/article/34871449930/

ncu-ui教程

为了显示原代码makefile添加 -g -G选项
对应CmakeList.txt

1
2
3

target_compile_options(better PUBLIC $<$<COMPILE_LANGUAGE:CUDA>:--extended-lambda
    -G -src-in-ptx
    >)

https://blog.csdn.net/yan31415/article/details/109491749

ncu-ui表格&图

我不明白我的SMEM怎么不是从DRAM来的， 而且峰值怎么这么低？

这个错误也是令人迷惑
The memory access pattern for loads from L1TEX to L2 is not optimal. The granularity of an L1TEX request to L2 is a 128 byte cache line. That is 4 consecutive 32-byte sectors per L2 request. However, this kernel only accesses an average of 3.7 sectors out of the possible 4 sectors per cache line. Check the Source Counters section for uncoalesced loads and try to minimize how many cache lines need to be accessed per memory request.

不知道为什么有1%和2% 的bank conflict

可以看到 SMEM， Register，Block Size是怎么影响GPU Warp的分配调度的。

上图没有拖累，吃满了64个warp。

关于if语句

if语句只要warp里执行相同就行。

可以提示出不连续访问的地方。(这里是这样设计的，已经避免了绝大部分的不连续访问)

显示stall最多的指令是什么以及在等待什么。还有执行最多的指令

假如 file mismatched 手动选择文件就行

stall的信息，感觉就这些有点用。(其中sb是scoreboard的意思)

ncu-ui 分析汇编

PTX&SASS汇编说明

有两种汇编

请看PTX SASS一文

基本说明

可以通过指令执行数或者采样率来得知，执行最多的指令。

鼠标悬停可以知道具体命令的含义

Ex1: for循环头

Ex2: for-loop kernel

1

sdata[Regular_local_index]=arr_data[Regular_global_index];

该从DRAM里读取到SMEM的指令对应的PTX和SASS代码

1

cvt.f32.u16 d, a;   // convert 16-bit unsigned to 32-bit float

问题：无效self-mov？

为了隐藏延迟？

直接原因是PTX翻译成SASS。一条mov变多条了

CUDA Visual Profiler

老一代debugger工具，逐渐被Nsight淘汰

1
2

nvprof # 命令行,nsys 之前的名称叫做 nvprof
nvvp

在more里有建议

nvprof捕获信息存储

1
2
3

nvprof --analysis-metrics -o  nbody-analysis.nvprof ./nbody --benchmark -numdevices=2 -i=1
# 下面输出 .qdrep 文件
nsys profile --stats=true --force-overwrite=true  -o baseline-report ./single-thread-vector-add

CUDA Visual Profiler 问题

Nsight Compute 问题

OpenGL 没有安装

1
2
3
4
5
6

Warning: Failed to get OpenGL version. OpenGL version 2.0 or higher is required.
OpenGL version is too low (0). Falling back to Mesa software rendering.
qt.qpa.plugin: Could not load the Qt platform plugin "xcb" in "" even though it was found.
This application failed to start because no Qt platform plugin could be initialized. Reinstalling the application may fix this problem.

Available platform plugins are: offscreen, wayland-egl, wayland, wayland-xcomposite-egl, wayland-xcomposite-glx, xcb.

解决办法

1
2

sudo apt-get install libxcb-xinerama0
sudo apt install libxcb-icccm4 libxcb-image0 libxcb-keysyms1 libxcb-render-util0

Qt插件缺失

1
2
3
4
5
6
7
8
9

qt.qpa.plugin: Could not load the Qt platform plugin "xcb" in "" even though it was found.
This application failed to start because no Qt platform plugin could be initialized. Reinstalling the application may fix this problem.

Available platform plugins are: xcb.

Application could not be initialized!
    This is likely due to missing Qt platform dependencies.
    For a list of dependencies, please refer to https://doc.qt.io/qt-5/linux-requirements.html
    To view missing libraries, set QT_DEBUG_PLUGINS=1 and re-run the application.

按照说明 export QT_DEBUG_PLUGINS=1再次运行, 显示具体问题

1

Cannot load library /staff/shaojiemike/Install/cuda_11.7.0_515.43.04_linux/nsight-compute-2022.2.0/host/linux-desktop-glibc_2_11_3-x64/Plugins/platforms/libqxcb.so: (libxcb-xinput.so.0: cannot open shared object file: No such file or directory)

解决 sudo apt-get install libxcb-xinput0

kernel没权限profile

ERR_NVGPUCTRPERM - The user does not have permission to profile on the target device

要用sudo，或者最新的NV

could not connect to display localhost:10.0 under sudo

1
2
3
4
5
6

$ sudo ncu-ui
MobaXterm X11 proxy: Authorisation not recognised
qt.qpa.xcb: could not connect to display localhost:10.0

MobaXterm X11 proxy: Unsupported authorisation protocol
Error: Can't open display: localhost:10.0

解决办法(原因是sudo相当于切换到root用户，丢失了xauth信息)

1
2
3
4
5
6
7
8
9
10

$ xauth list
snode0/unix:12  MIT-MAGIC-COOKIE-1  84941f1f8be97d19436356685f75b884
snode0/unix:13  MIT-MAGIC-COOKIE-1  5172ee2c7364b055cd37538b460f7741
snode0/unix:11  MIT-MAGIC-COOKIE-1  589f3b5ab852f24ca3710c53e6439260
hades1/unix:10  MIT-MAGIC-COOKIE-1  9346adec202bd65250f3d21239025750
snode0/unix:10  MIT-MAGIC-COOKIE-1  52285c563f1688741fa1b434ed2b7b2c

sudo -s # 切换
xauth add snode0/unix:10  MIT-MAGIC-COOKIE-1  52285c563f1688741fa1b434ed2b7b2c # 补全xauth
# 正常执行 xauth有用的总是最后一个

GPU Metrics [0]: Sampling buffer overflow.

1. 只勾选CUDA Metrics 和 GPU Metrics
2. 降低采样频率

Error 0: UnsupportedGpu

原因是 软件对GPU的支持是逐步的需要安装最新的。

不支持的Nsight的可以尝试老的debugger工具 CUDA Visual Profiler

Error: Profiling is not supported on this device

Pascal support was deprecated, then dropped from Nsight Compute after Nsight Compute 2019.5.1.

The profiling tools that support Pascal in the CUDA Toolkit 11.1 and later are nvprof and visual profiler.

需要进一步的研究学习

暂无

遇到的问题

NVTX问题

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

参考文献

https://developer.nvidia.com/tools-overview

https://www.365seal.com/y/zyn1yxJQn3.html

[^1]: Usage of NVTX