1.3d稀疏卷积——spconv源码剖析（一）
2.3d稀疏卷积——spconv源码剖析（五）
3.3d稀疏卷积——spconv源码剖析（三）

3d稀疏卷积——spconv源码剖析（一）

       本文主要阐述卷积的唱歌唱歌基本理论，并以spconv源码为例进行解析。网站网站网页首先，源码源码介绍2D与3D卷积的唱歌唱歌基础知识及其分类。随后，网站网站网页深入探讨3D稀疏卷积的源码源码Ios系统iptables源码工作原理。

       2D卷积涉及卷积核在二维图像空间上的唱歌唱歌滑动操作。它分为单通道卷积与多通道卷积。网站网站网页单通道卷积在输入图像的源码源码单一通道上进行，得到特征图。唱歌唱歌多通道卷积在同一图像中不同通道上进行，网站网站网页每个通道得到一个对应的源码源码新通道，最终通过相加生成特征图。唱歌唱歌ecilpes 编译源码

       3D卷积在此基础上扩展到三维空间，网站网站网页涉及单通道与多通道情况。源码源码三维单通道卷积在立方体上进行，而三维多通道卷积则处理拥有多个通道的三维图像。

       2D与3D卷积计算涉及输入层、输出层与参数关系的数学公式。考虑偏置参数与计算量，FLOPS（浮点运算量）也在此阶段被计算。

       稀疏卷积分为SC（Sparse Convolution）与VSC（Valid Sparse Convolution）两种类型。SC卷积计算激活站点并丢弃非激活站点，而VSC卷积在SC的基础上进行了简化。

       卷积神经网络对三维点云数据处理时，c prinf 源码面临计算量增加的问题，而SC与VSC卷积利用稀疏性实现高效处理。构建输入与输出哈希表，对点云数据进行快速访问。GetOffset()函数用于定位卷积操作的位置，Rulebook用于存储原子操作规则，指导稀疏卷积过程。

       稀疏卷积的关键在于构建输入、输出哈希表以及建立两者之间的联系，实现对稀疏数据的有效处理。spconv库中的get_indice_pairs函数通过调用getIndicePairs实现这一过程。

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，其代码位于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底层的注册接口（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...