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...