1.3d稀疏卷积——spconv源码剖析(五)
2.3d稀疏卷积——spconv源码剖析(三)
3.SourceMODSourceMOD插件编译
4.CUDA学习:基础知识小结
5.3d稀疏卷积——spconv源码剖析(一)
6.软件sp和cc是源码什么意思
3d稀疏卷积——spconv源码剖析(五)
介绍在构建的Rulebook指导下执行特定的稀疏卷积计算,关注于类SparseConvolution,编译其代码位于spconv/conv.py。源码
Fsp.indice_subm_conv和Fsp.indice_conv经过spconv/functional.py中的编译SubMConvFunction和SparseConvFunction对象转换,最终会调用spconv/ops.py模块中的源码indice_conv等函数。
专注于子流线卷积接口:indice_subm_conv,编译cas server 源码下载其代码位于spconv/functional.py。源码
通过Python接口调用底层C++函数可能不够直观,编译因此使用torch.autograd.Function封装算子底层调用,源码该类表示PyTorch中的编译可导函数,具备前向推理和反向传播实现时,源码即可作为普通PyTorch函数使用。编译
值得注意的源码是,Function类在模型部署中具有优势,编译若定义了symbolic静态方法,源码此Function在执行torch.onnx.export()时,可依据symbolic定义规则转换为ONNX算子。
apply方法是torch.autograd.Function的一部分,此方法负责在前向推理或反向传播时的调度工作。通过将indice_subm_conv = SubMConvFunction.apply简化为indice_subm_conv接口,简化了算子使用,屏蔽了SubMConvFunction的具体实现。
SubMConvFunction的前向传播方法forward调用spconv/ops.py的indice_conv函数。在src/spconv/all.cc文件中,题库移动版源码通过PyTorch提供的OP Register对底层C++API进行注册。
通过torch.ops.load_library加载.so文件,使用torch.ops.spconv.indice_conv调用src/spconv/spconv_ops.cc文件中的indiceConv函数。
深入探索src/spconv/spconv_ops.cc文件中的indiceConv函数。
代写部分代码内容...
3d稀疏卷积——spconv源码剖析(三)
构建Rulebook
下面看ops.get_indice_pairs,位于:spconv/ops.py
构建Rulebook由ops.get_indice_pairs接口完成
get_indice_pairs函数具体实现:
主要就是完成了一些参数的校验和预处理。首先,对于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底层的impalajdbc41 源码注册接口(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表示)到输出序号之间的bbi ebbi附图源码一张哈希表
见: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...
SourceMODSourceMOD插件编译
SourceMOD插件的基础文件结构包含源码文件,其后缀名为.sp,以及头文件,后缀为.inc。当你完成源码编写后,需要将其编译成可执行的插件,这时的文件后缀会变为.smx。
编译SourceMOD插件的过程与AMXX插件类似,提供了四种编译方法供开发者选择。首先,你可以直接将源码文件放入与编译器compile.exe相同的目录,然后双击compile.exe执行编译,编译后的目标文件会出现在compiled目录中。
第二种方法是更为便捷的拖放操作。只需将源码文件拖动到compile.exe上,编译器会自动在源码同级目录下生成编译后的易语言源码注册文件。
对于那些使用PawnStudio这款SourceMOD插件的专用编辑器的开发者,可以利用其提供的便利。在PawnStudio中设置好编译器路径和目标文件输出路径,就能在编辑器内直接进行插件的编写和编译,无需离开编辑环境。
最后,如果你想要在线编译,官方也提供了相应的服务,只需将你的源码上传到官方平台,即可在线完成编译过程,省去了在本地环境配置的步骤。
以上就是SourceMOD插件的编译方法,无论你选择哪种方式,关键步骤都是将源码转换为可在游戏中使用的插件文件。
CUDA学习:基础知识小结
在CUDA学习中,理解编译流程是基础。首先,让我们深入探讨nvcc的编译过程。在将源代码SRC编译为PTX时,编译选项中的XY代表虚拟架构的计算能力,它限制了代码可以利用的CUDA特性。接着,ZW在PTX到cubin的转换中表示真实架构的计算能力,确保执行文件能适配的GPU性能,必须注意ZW应大于等于XY。示例编译选项如下:
除了常规编译,JIT编译方式会在执行时动态编译PTX,这里也有一个JIT编译选项示例。简化编译选项如-arch=sm_XY,等同于指定虚拟架构。
CUDA编程中,SM、SP、grid、block、warp和thread等概念是关键。从软件角度看,它们之间有明确的关系。例如,grid和block的维度可以是三维的,而thread的索引通过维度转换来获取。这里有一张图展示了这些概念的关系。
kernel function是CUDA程序的核心,它的定义和使用有一些限制。要正确调用,需要指定grid_size和block_size,它们对应于block的数量和thread的数量。这里详细解释了kernel function的定义和调用方式。
CUDA函数的修饰词__host__、__device__、__global__决定函数的执行环境。CUDA程序通常分为数据准备、执行kernel、数据交换和错误处理等步骤,其中数据拷贝是一个关键环节,使用cudaMemcpy进行,它支持多种数据传输方向。
计时是性能评估的重要手段,CUDA通过事件来实现。Debug功能则涉及API错误检测和kernel function的异步执行错误检查。CUDA-MEMCHECK工具集是调试工具的重要组成部分,可以通过特定命令调用。
3d稀疏卷积——spconv源码剖析(一)
本文主要阐述卷积的基本理论,并以spconv源码为例进行解析。首先,介绍2D与3D卷积的基础知识及其分类。随后,深入探讨3D稀疏卷积的工作原理。
2D卷积涉及卷积核在二维图像空间上的滑动操作。它分为单通道卷积与多通道卷积。单通道卷积在输入图像的单一通道上进行,得到特征图。多通道卷积在同一图像中不同通道上进行,每个通道得到一个对应的新通道,最终通过相加生成特征图。
3D卷积在此基础上扩展到三维空间,涉及单通道与多通道情况。三维单通道卷积在立方体上进行,而三维多通道卷积则处理拥有多个通道的三维图像。
2D与3D卷积计算涉及输入层、输出层与参数关系的数学公式。考虑偏置参数与计算量,FLOPS(浮点运算量)也在此阶段被计算。
稀疏卷积分为SC(Sparse Convolution)与VSC(Valid Sparse Convolution)两种类型。SC卷积计算激活站点并丢弃非激活站点,而VSC卷积在SC的基础上进行了简化。
卷积神经网络对三维点云数据处理时,面临计算量增加的问题,而SC与VSC卷积利用稀疏性实现高效处理。构建输入与输出哈希表,对点云数据进行快速访问。GetOffset()函数用于定位卷积操作的位置,Rulebook用于存储原子操作规则,指导稀疏卷积过程。
稀疏卷积的关键在于构建输入、输出哈希表以及建立两者之间的联系,实现对稀疏数据的有效处理。spconv库中的get_indice_pairs函数通过调用getIndicePairs实现这一过程。
软件sp和cc是什么意思
软件SP是一种软件产品的简写,是英文单词“service pack”的缩写。它是一组微软公司推出的升级补丁,用于提供操作系统的安全性和稳定性。通常情况下,软件SP会把多种更新、修补、补丁和安全修复程序打包成一个升级文件,用户可以通过安装该文件来更新其计算机上的操作系统。
软件CC是一些软件程序经过编译后生成的可执行文件的后缀名。在编译时,源代码会被转化为计算机可识别的二进制代码,并将其转换为可执行文件。这些可执行文件最终会使用软件CC的后缀名来命名。常见的软件CC有C++、Java、Python等,这些编译器通过转换源代码为机器码,使得程序具有了运行的能力。
软件SP和CC在计算机软件领域中具有非常广泛的应用。软件SP通过打包多种更新和修复程序,可以快速成批更新操作系统,提高了计算机的安全性和稳定性,减少了计算机遭受黑客攻击和崩溃的几率。而软件CC则是编译器生成的可执行文件后缀名,是许多编程语言的标志。程序员们在编写代码时,需要使用到相关软件CC将源代码编译成可执行文件,便于程序的调试和运行。