1.8086模拟器8086tiny源码分析(5)执行mov指令(二)
2.MLIR Operation 与 Op
3.3d稀疏卷积——spconv源码剖析(三)
4.Linux内核中的源码open方法
8086模拟器8086tiny源码分析(5)执行mov指令(二)
本文继续解析tiny模拟器中的MO指令,集中于MOV reg,源码 r/m指令的实现。首先,源码通过xlat_opcode_id赋值为9,源码额外参数extra设置为8,源码为后续解析打下基础。源码旋转音乐源码核心部分在于理解OP(=)的源码操作,其完成的源码是寄存器与内存或另一个寄存器间的数据移动。
进一步分析,源码MEM_OP和R_M_OP分别对应内存操作与寄存器与内存间的源码拷贝,前者是源码基本内存操作,后者完成具体数据移动任务。源码而op_to_addr和op_from_addr则是源码关键变量,前者代表目的源码位置,后者代表源位置。源码具体赋值依赖于是否需要解码mod、rm、reg三个指令字段。公交 小程序源码
当i_mod_size为真时,解码这三个字段,并结合d和w字段,确定操作数。这由DECODE_RM_REG宏完成。在这里,op_to_addr是目的位置(寄存器或内存),op_from_addr是源位置。指令数据移动方向的关键在于i_d变量。如果该变量为真,则表示源操作数与目的操作数需进行交换。
至此,对MOV reg, r/m指令的解析告一段落。通过明确指令字段、操作变量的赋值与交换条件,tiny模拟器成功实现这一重要指令的执行,为深入理解架构与模拟器设计提供了基础。qt源码动态 编译
MLIR Operation 与 Op
MLIR的核心组件之一是Operation,它是Dialect中的基本语义单元,就像方言中的词汇一样,构成了整个代码表示的基石。深入理解MLIR的运作,关键在于掌握Operation类和Op类的运作机制。
Operation类是MLIR中的核心数据结构,它在 mlir/include/mlir/IR/Operation.h 中被定义,为操作提供了丰富的接口,包括静态create函数,用于创建不同类型的Operation实例。例如,当我们创建一个FuncOp时,会先通过Operation::create调用,再通过cast转换为具体类。此外,Operation还负责操作数、div阻止查看源码结果、属性和blocks的管理,以及操作的遍历等。
Op类则是所有具体Operation的基类,它在mlir/include/mlir/IR/OpBase.td中定义,这个文件实际上是TableGen语言编写的。在定义新的Dialect时,开发者需要在Ops.td文件中引入OpBase.td,基于这个基类扩展出符合特定Dialect需求的Op类。
整个过程涉及到TableGen工具的编译和链接,将这些定义转化为编译时可用的代码。通过理解并操作这些基础类,开发者可以构建和优化复杂的编译逻辑。更多关于Op类的具体细节和功能,还需进一步研究mlir的源码和官方文档。
3d稀疏卷积——spconv源码剖析(三)
构建Rulebook
下面看ops.get_indice_pairs,位于:spconv/ops.py
构建Rulebook由ops.get_indice_pairs接口完成
get_indice_pairs函数具体实现:
主要就是bwapp哪里查看源码完成了一些参数的校验和预处理。首先,对于3d普通稀疏卷积,根据输入shape大小,kernel size,stride等参数计算出输出输出shape,子流行稀疏卷积就不必计算了,输出shape和输入shape一样大小
准备好参数之后就进入最核心的get_indice_pairs函数。因为spconv通过torch.ops.load_library加载.so文件注册,所以这里通torch.ops.spconv.get_indice_pairs这种方式来调用该函数。
算子注册:在src/spconv/all.cc文件中通过Pytorch提供的OP Register(算子注册的方式)对底层c++ api进行了注册,可以python接口形式调用c++算子
同C++ extension方式一样,OP Register也是Pytorch提供的一种底层扩展算子注册的方式。注册的算子可以通过 torch.xxx或者 tensor.xxx的方式进行调用,该方式同样与pytorch源码解耦,增加和修改算子不需要重新编译pytorch源码。用该方式注册一个新的算子,流程非常简单:先编写C++相关的算子实现,然后通过pytorch底层的注册接口(torch::RegisterOperators),将该算子注册即可。
构建Rulebook实际通过python接口get_indice_pairs调用src/spconv/spconv_ops.cc文件种的getIndicePairs函数
代码位于:src/spconv/spconv_ops.cc
分析getIndicePairs直接将重心锁定在GPU逻辑部分,并且子流行3d稀疏卷积和正常3d稀疏卷积分开讨论,优先子流行3d稀疏卷积。
代码中最重要的3个变量分别为:indicePairs,indiceNum和gridOut,其建立过程如下:
indicePairs代表了稀疏卷积输入输出的映射规则,即Input Hash Table 和 Output Hash Table。这里分配理论最大的内存,它的shape为{ 2,kernelVolume,numAct},2表示输入和输出两个方向,kernelVolume为卷积核的volume size。例如一个3x3x3的卷积核,其volume size就是(3*3*3)。numAct表示输入有效(active)特征的数量。indiceNum用于保存卷积核每一个位置上的总的计算的次数,indiceNum对应中的count
代码中关于gpu建立rulebook调用create_submconv_indice_pair_cuda函数来完成,下面具体分析下create_submconv_indice_pair_cuda函数
子流线稀疏卷积
子流线稀疏卷积是调用create_submconv_indice_pair_cuda函数来构建rulebook
在create_submconv_indice_pair_cuda大可不必深究以下动态分发机制的运行原理。
直接将重心锁定在核函数:
prepareSubMGridKernel核函数中grid_size和block_size实则都是用的整形变量。其中block_size为tv::cuda::CUDA_NUM_THREADS,在include/tensorview/cuda_utils.h文件中定义,大小为。而grid_size大小通过tv::cuda::getBlocks(numActIn)计算得到,其中numActIn表示有效(active)输入数据的数量。
prepareSubMGridKernel作用:建立输出张量坐标(通过index表示)到输出序号之间的一张哈希表
见:include/spconv/indice.cu.h
这里计算index换了一种模板加递归的写法,看起来比较复杂而已。令:new_indicesIn = indicesIn.data(),可以推导得出index为:
ArrayIndexRowMajor位于include/tensorview/tensorview.h,其递归调用写法如下:
接着看核函数getSubMIndicePairsKernel3:
位于:include/spconv/indice.cu.h
看:
上述写法类似我们函数中常见的循环的写法,具体可以查看include/tensorview/kernel_utils.h
NumILP按默认值等于1的话,其stride也是gridDim.x*blockDim.x。索引最大值要小于该线程块的线程上限索引blockDim.x * gridDim.x,功能与下面代码类似:
参考: blog.csdn.net/ChuiGeDaQ...
Linux内核中的open方法
在Linux系统中,使用open方法打开文件是一个基本操作,本文将从源码角度解析Linux内核中的open方法是如何实现文件打开功能的。
在Linux内核中,open方法的实现主要涉及几个关键步骤。首先,调用do_sys_open方法作为主要入口。该方法内进行了一系列初始化操作,主要包括:
1. 调用build_open_flags方法,初始化struct open_flags实例op。
2. 调用getname方法,分配并初始化struct filename实例tmp。
3. 调用get_unused_fd_flags方法获取一个未被使用的文件描述符fd。
4. 调用do_filp_open方法,继续执行open操作,并将返回值赋值给类型为struct file的实例指针f。
5. 如果do_filp_open成功,则调用fd_install方法,建立从fd到struct file的对应关系。
6. 最后,返回fd给用户。
在do_filp_open方法中,进一步执行open操作。主要步骤包括:
1. 调用set_nameidata方法,初始化struct nameidata类型实例nd。
2. 调用path_openat方法继续执行open操作。
3. 在path_openat方法内,首先调用alloc_empty_file方法分配一个空的struct file实例。
4. 然后调用path_init、link_path_walk、do_last等方法执行后续的open操作。
5. 其中,path_init方法初始化struct nameidata实例中的path、root、inode等字段。
6. link_path_walk方法处理路径组件,最终使struct nameidata实例指向目标路径。
7. 最后,do_last方法完成剩余的open操作,如查找路径中的最后一个组件,执行open逻辑。
整个open过程涉及多层函数调用,从系统调用到内核源码的详细实现,体现了Linux内核设计的复杂性和层次性。通过逐步解析这些步骤,可以更直观地理解Linux内核如何实现文件打开功能。