1.3d稀疏卷积——spconv源码剖析(一)
2.3d稀疏卷积——spconv源码剖析(五)
3.3d稀疏卷积——spconv源码剖析(三)
3d稀疏卷积——spconv源码剖析(一)
本文主要阐述卷积的唱歌唱歌基本理论,并以spconv源码为例进行解析。网站网站网页首先,源码源码介绍2D与3D卷积的唱歌唱歌基础知识及其分类。随后,网站网站网页深入探讨3D稀疏卷积的源码源码windows android源码工作原理。
2D卷积涉及卷积核在二维图像空间上的唱歌唱歌滑动操作。它分为单通道卷积与多通道卷积。网站网站网页单通道卷积在输入图像的源码源码单一通道上进行,得到特征图。唱歌唱歌多通道卷积在同一图像中不同通道上进行,网站网站网页每个通道得到一个对应的源码源码新通道,最终通过相加生成特征图。唱歌唱歌springboot源码探究
3D卷积在此基础上扩展到三维空间,网站网站网页涉及单通道与多通道情况。源码源码三维单通道卷积在立方体上进行,而三维多通道卷积则处理拥有多个通道的三维图像。
2D与3D卷积计算涉及输入层、输出层与参数关系的数学公式。考虑偏置参数与计算量,FLOPS(浮点运算量)也在此阶段被计算。
稀疏卷积分为SC(Sparse Convolution)与VSC(Valid Sparse Convolution)两种类型。SC卷积计算激活站点并丢弃非激活站点,而VSC卷积在SC的基础上进行了简化。
卷积神经网络对三维点云数据处理时,腾讯源码接口面临计算量增加的问题,而SC与VSC卷积利用稀疏性实现高效处理。构建输入与输出哈希表,对点云数据进行快速访问。GetOffset()函数用于定位卷积操作的位置,Rulebook用于存储原子操作规则,指导稀疏卷积过程。
稀疏卷积的关键在于构建输入、输出哈希表以及建立两者之间的联系,实现对稀疏数据的有效处理。spconv库中的get_indice_pairs函数通过调用getIndicePairs实现这一过程。
3d稀疏卷积——spconv源码剖析(五)
介绍在构建的vtcp源码分析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,其代码位于spconv/functional.py。
通过Python接口调用底层C++函数可能不够直观,因此使用torch.autograd.Function封装算子底层调用,该类表示PyTorch中的可导函数,具备前向推理和反向传播实现时,即可作为普通PyTorch函数使用。
值得注意的scrm源码 php是,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底层的注册接口(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...