1.FasterTransformer Decoding 源码分析(三)-LayerNorm介绍
2.银河系CUDA编程指南(2.5)——NVCC与PTX
3.3d稀疏卷积——spconv源码剖析（三）
4.CUDA学习：基础知识小结
5.[fastllm]cuda-kernels源码解析
6.[技术随笔]🛠🛠从源码安装Pytorch3D详细记录及学习资料

FasterTransformer Decoding 源码分析(三)-LayerNorm介绍

       本文深入探讨FasterTransformer中LayerNormalization（层归一化）的源码源码实现与优化。作为深度学习中的解析关键技术，层归一化可确保网络中各层具有相似的源码分布，从而加速训练过程并改善模型性能。解析背景介绍部分详细解释了层归一化的源码工作原理，强调其在神经网络中的解析网络祭祀 源码高效并行特性与广泛应用。文章从代码起点开始剖析，源码具体路径位于解码过程的解析核心部分。调用入口展示了传入参数，源码包括数据描述和关键参数gamma、解析beta、源码eps，解析简洁直观，源码符合公式定义。解析深入源码的源码解析揭示了优化点，特别是针对特定数据类型和维度，使用了定制化内核。此设计针对高效处理半精度数据样本，减少判断指令，实现加速运算，且对偶数维度数据进行调整以最大化Warp特性利用。接下来，内核实现的详细描述，强调了通过共享内存与block、warp级归约实现公式计算的高效性。这部分以清晰的代码结构和可视化说明，解释了块级别与Warp级归约在单个块处理多个数据点时的协同作用，以及如何通过巧妙编程优化数据处理效率。文章总结了FasterTransformer中LayerNormalization的整体优化策略，强调了在CUDA开发中基础技巧的应用，并指出与其他优化方案的比较。此外，文章还推荐了OneFlow的性能优化实践，为读者提供了一个深入探索与对比学习的资源。

银河系CUDA编程指南(2.5)——NVCC与PTX

       在构建了一个以cuDNN和cuBLAS为基础的简单深度学习框架后，我已将其开源，并鼓励大家参与交流学习。未来计划逐步完善框架，将尝试使用纯CUDA C实现，并与cuDNN进行性能比较。关于cuDNN的使用，我也会后续专门撰写文章进行详细介绍。

       NVCC，CUDA的编译器，其核心是NVVM优化器，基于LLVM编译器结构。它本质上是一个集合，调用gcc、cicc、ptxas等工具编译CUDA源代码，区分主机代码（用ANSI C编写）和设备代码（CUDA扩展语言编写）。

       NVCC的编译过程分为离线编译和即时编译，通过预处理将源代码分为两部分，分别由不同编译器处理，最终合并为单个object文件。TBK线源码例如，通过dryrun选项可以查看编译步骤，包括头文件配置、CUDA设备代码编译等。

       PTX是CUDA的编程模型和指令集，是一种虚拟架构汇编，允许跨GPU优化。NVCC通过虚拟架构编译生成PTX，然后在实际GPU上执行为SASS。编译时，需设置虚拟和实际GPU架构以保证兼容性。

       Separate Compilation允许在编译阶段将device code分开处理，形成relocatable代码，然后在链接阶段定位到最终的host object。这与Whole Program Compilation不同，后者直接编译为executable device code。

       以cudnn-learning-framework的Makefile为例，需配置CUDA相关路径，添加cuDNN库，并调整编译生成部分，确保链接所有需要的.o文件。NVCC命令在编译时会执行链接任务。

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的volley内部源码卷积核，其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...

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__决定函数的获取sidphp源码执行环境。CUDA程序通常分为数据准备、执行kernel、数据交换和错误处理等步骤，其中数据拷贝是一个关键环节，使用cudaMemcpy进行，它支持多种数据传输方向。

       计时是性能评估的重要手段，CUDA通过事件来实现。Debug功能则涉及API错误检测和kernel function的异步执行错误检查。CUDA-MEMCHECK工具集是调试工具的重要组成部分，可以通过特定命令调用。

[fastllm]cuda-kernels源码解析

       在fastllm中，CUDA-kernels的使用是关键优化点之一，主要涉及以下几个高频率使用的kernel：gemv_int4、gemv_int8、gemm_int8、RMSNorm、softmax、RotatePosition2D、swiglu等。其中，gemm是计算密集型的，而其余大部分都是内存受限型。利用量化bit进行计算，比原始的torch转为浮点数更快，同时，没有进行融合操作，为后续优化留下了空间。

       gemv_int4 kernel：主要用于实现float*int4的GEMV乘积，其中偏置值设定为最小值。在计算中，矩阵被划分为不同的tile，不同tile之间并行操作。在遍历m/2的过程中，找到对应int4值的位置，通过保存的mins找到最小值minv。同一组的两个int4值共享同一个minv，计算结果的最终和被保存在sdata[0]上，用于更新对应m列位置的output值。结果向量为n*1。

       gemv_int8 kernel：在功能上与gemv_int4类似，但偏置值由保存的minv变为了zeros。

       gemm_int8 kernel：此kernel负责计算n*m矩阵与m*k矩阵的乘积。计算过程涉及多个tile并行，block内部保存的是部分和。考虑到线程数量限制，通常会有优化空间。最终结果通过为单位进行更新。

       layerNorm实现：此kernel实现layernorm计算，通过计算均值和方差来调整数据分布。计算中，sdata存储所有和，sdata2存储平方和。每个block内计算部分和后，关闭串口源码规约得到全局的均值和方差，从而更新output。

       RMS kernels解析：RMSNorm kernel实现RMS归一化，通过计算输入的平方和和均值，进而更新output。

       softmax kernels解析：计算输入的softmax值，涉及最大值查找、指数计算和规约求和等步骤，以防止浮点数下溢。

       RotatePosition2D Kernels解析：用于旋转位置编码，线程展开成三层循环。LlamaRotatePosition2D、NearlyRotatePosition和RotatePosition2D在旋转方式上有所区别，体现在不同的位置上进行计算。

       AttentionMask Kernels解析：对输入按照mask掩码置值，普通mask直接置为maskv，而Alibimask则是置为相对位置的值之和。具体含义可能涉及空间上的概念，但文中未详细说明。

       swiglu kernels解析：作为激活函数，这些kernel在原地操作中执行常见函数，线程足够使用，直接按照公式计算即可。

       综上所述，fastllm中CUDA-kernels的使用旨在通过优化计算过程和内存操作，提升模型的计算效率，实现更高效的推理和训练。

[技术随笔]🛠🛠从源码安装Pytorch3D详细记录及学习资料

       在启动安装Pytorch3D之前，首要任务是选择合适的pytorch基础镜像。我选择了包含CUDA组件和驱动的pytorch 1.9的devel版本，以确保满足Pytorch3D对于pytorch和cuda版本的要求。我使用的是python 3.7、pytorch 1.9和cuda.2，前提是你已经在宿主机上配置好了显卡驱动和nvidia-docker，以便在容器内映射宿主机的显卡信息。

       在安装前，确保nvcc编译器、CUDA工具箱和驱动正常运行，并且安装了git、vim、sudo和curl等基础工具。

       下一步是配置CUB工具。按照Pytorch3D的安装文档，为了支持CUDA，需要先配置CUB，并设置CUB_HOME环境变量。由于选择的镜像包含CUDA，编译过程中会自动包含cuda。为保险起见，可以指定FORCE_CUDA环境变量为1。

       从源码编译Pytorch3D时，避免了使用conda可能遇到的依赖冲突问题。在确认前两步没有问题后，编译过程通常顺利。安装完成后，检查日志和pytorch3d的版本信息。

       为了验证Pytorch3D的正常运行，从ARkit中导出BS系数，尝试使用它渲染一个简单的白模，并利用GPU。观察到显卡被充分利用，表明设置正确，可以进行后续操作。

       在完成安装并验证Pytorch3D的功能后，可以参考收集的资料来探索其更高级的用法。以下是几个示例：

       从Pytorch3D文档中获取的教程和代码示例。

       开源社区的讨论和问题解答，特别是与Pytorch3D相关的话题。

       个人经验分享和案例研究，可以在GitHub、Stack Overflow等平台找到。

       通过这些资源，您可以深入学习Pytorch3D的功能和应用，进一步拓展其在计算机图形学、三维重建和深度学习等领域的应用。

CUDA编程OneFlow Softmax 算子源码解读之WarpSoftmax

       深度学习框架中的Softmax操作在模型中扮演关键角色，尤其在多分类任务中，其用于将logits映射成概率分布，或在Transformer结构中衡量query与key的相似度。Softmax的CUDA实现直接关系到模型训练效率。本文以OneFlow框架中的一种优化Softmax实现为例，即Warp级别的Softmax，特别适用于矩阵宽度不超过的场景。

       Softmax操作的计算公式如下：

       [公式]

       为解决数值溢出问题，通常先减去向量的最大值。优化后的公式为：

       [公式]

       Softmax计算涉及五个关键步骤：reduceMax、broadcastSub、exp、reduceSum、broadcastDiv。本篇文章将深入探讨OneFlow源码中的实现技巧。

       OneFlow采用分段函数优化SoftmaxKernel，针对不同数量的列选择不同实现策略，以适应各种场景。为实现优化，OneFlow提供三种Softmax实现方式，以期在所有情况下达到较高的有效带宽。

       对于WarpSoftmax分支，源码中函数调用关系清晰，实现细节分为四部分：数据Pack、调用链、DispatchSoftmaxWarpImpl、DispatchSoftmaxWarpImplCols、DispatchSoftmaxWarpImplPadding、LaunchSoftmaxWarpImpl。各部分分别专注于提升访问带宽、确定函数参数、实现核心计算逻辑。

       在WarpSoftmax的核函数SoftmaxWarpImpl中，重点实现以下步骤：核函数启动参数确定、线程网格形状定义、数据加载到寄存器、计算最大值、计算指数和、规约操作、通信优化等。实现过程中，OneFlow通过优化数据访问模式、利用寄存器存储中间结果、并行规约操作，以及束内通信，提升了计算效率。

       总结WarpSoftmax源码中的关键点，本文详细解读了其优化策略与实现细节，旨在提高模型训练速度。通过深入分析OneFlow框架中的Softmax实现，读者可以更全面地理解深度学习框架在CUDA环境下进行优化的策略。

GPU编程3：CUDA环境安装和IDE配置

       本文指导如何在个人机器上安装CUDA环境，结合集成开发环境Clion进行配置，以方便后续CUDA编程学习。

       安装CUDA环境如下：

       1. 针对显卡型号，从官方下载相应驱动。

       示例显卡型号：小米pro寸，GF MX 。

       参考链接：nvidia.cn/Download/index.aspx

       2. 阻止或卸载nouveau驱动。

       3. 通过控制台进入文本界面，安装NVIDIA驱动。

       步骤示例：调整引导运行级别，以便开机进入文本界面。

       网址参考：jingyan.baidu.com/article/0abcb0fbdf.html

       4. 确认驱动安装。

       5. 尽量与CUDA版本匹配安装NVIDIA驱动。

       6. 进行CUDA测试。

       CUDA代码编译与运行：

       编译CUDA源码时，包含两个部分：CUDA设备函数与主机函数，它们分开独立编译。CUDA 5.0+支持文件间设备代码独立编译，而整体编译是默认模式。

       编译三个文件（a.cu, b.cu, c.cpp），其中a.cu调用了b.cu中定义的设备代码，可以使用独立编译方式实现。

       详细编译步骤：使用nvcc编译设备函数，普通C/C++编译器编译主机代码。

       举例：`nvcc a.cu`编译设备文件。

       实际工程中，为了优化编译效率，常采用`makefile`或`CMake`工具配置源码编译。

       `nvcc`支持多种快捷开关，如`-arch=sm_`编译特定架构。

       基于Clion的CUDA配置流程：

       1. 遇到Clion创建CUDA可执行文件失败问题。

       检查是否已安装NVCC。

       验证机器安装GPU卡。

       检查安装路径：执行`which nvcc`，若未找到，则进行安装。

       确认安装位置：输入`nvcc`显示默认路径，通常为`/usr/bin/nvcc`。

       2. 利用Clion新建CUDA项目，并设置CMake。

       配置CMake代替`makefile`，简化编译过程。

       输出及结果：提供示例链接供参考。

Python语言学习（三）：Tensorflow_gpu搭建及convlstm核心源码解读

       在探索深度学习领域，使用Python语言进行编程无疑是一条高效且灵活的途径。尤其在科研工作或项目实施中，Python以其丰富的库资源和简单易用的特性，成为了许多专业人士的首选。本文旨在分享在Windows系统下使用Anaconda搭建TensorFlow_gpu环境及解读ConvLSTM核心源码的过程。在提供具体步骤的同时，也期待读者的反馈，以持续改进内容。

       为了在Windows系统下搭建适合研究或项目的TensorFlow_gpu环境，首先需要确认TensorFlow_gpu版本及其对应的cuDNN和CUDA版本。访问相关网站，以获取适合自身硬件配置的版本信息。以TensorFlow_gpu2.为例，进行环境搭建。

       在Anaconda环境下，通过命令行操作来创建并激活特定环境，如`tensorflow-gpu`环境，选择Python3.版本。接着，安装cuDNN8.1和CUDA.2。推荐使用特定命令确保安装过程顺利，亲测有效。随后，使用清华镜像源安装TensorFlow_gpu=2..0。激活虚拟环境后，使用Python环境验证安装成功，通常通过特定命令检查GPU版本是否正确。

       为了在Jupyter Notebook中利用该环境，需要安装ipykernel，并将环境写入notebook的kernel中。激活虚拟环境并打开Jupyter Notebook，通过命令确保内核安装成功。

       对于ConvLSTM核心源码的解读，重点在于理解模型的构建与参数设置。模型核心代码通常包括输入数据维度、模型结构、超参数配置等。以官方样例为例，构建模型时需关注样本整理、标签设置、卷积核数量等关键参数。例如，输入数据维度为（None，，，1），输出数据维度为（None，None，，，）。通过返回序列设置，可以控制模型输出的形态，是返回单个时间步的输出还是整个输出序列。

       在模型改造中，将彩色图像预测作为目标，需要调整模型的最后层参数，如将`return_sequence`参数更改为`False`，同时将`Conv3D`层修改为`Conv2D`层以适应预测彩色图像的需求。此外，选择合适的损失函数（如MAE）、优化器（如Adam）以及设置Metrics（如MAE）以便在训练过程中监控模型性能。

       通过上述步骤，不仅能够搭建出适合特定研究或项目需求的TensorFlow_gpu环境，还能够深入理解并灵活应用ConvLSTM模型。希望本文内容能够为读者提供有价值的指导，并期待在后续过程中持续改进和完善。

极智开发 | 解读英伟达软件生态 一切的基础CUDA

       欢迎关注我的公众号 [极智视界]，获取我的更多经验分享

       大家好，我是极智视界，本文来介绍一下 解读英伟达软件生态 一切的基础CUDA。

       邀您加入我的知识星球「极智视界」，星球内有超多好玩的项目实战源码下载，链接： t.zsxq.com/0aiNxERDq

       CUDA，全称为 Compute Unified Device Architecture，是英伟达于 年推出的一个平行计算平台和应用编程接口 API 模型。CUDA 之于英伟达的重要性主要体现在下面几个方面：

       所以，解读英伟达软件生态，必须要从 CUDA 说起。虽然 CUDA 再往下还有如 PTX 的指令集加速层级，但是PTX 的普及程度其实并不高，甚至可能很多朋友都没听说过 PTX，其实也算正常，因为基本上的 N 卡开发者，根本没必要接触到 PTX，把 CUDA 学好就足够够的了。

       把 CUDA 作为标杆，似乎是很多 AI 芯片厂商 "共同的做法"，比较有代表性的是升腾的 Ascend C、寒武纪的 Bang C，但是其实这几个之间有相似但又有不相似的地方。相似的地方在于不管是升腾还是寒武纪都想提供一套类似 CUDA 的可以充分调用自己 NPU 硬件加速的对外接口，提高客制化的灵活性。不相似的地方一方面在于我们是在学人家，很多接口其实是为了贴近 CUDA 的接口而进行的高级封装，毕竟大部分开发者其实已经形成了 CUDA 的开发习惯，这个时候让大家切换起来更加顺手的做法就是 "模仿"，而要做这种程度的 "模仿" 势必要协调好硬件架构和软件接口的映射；不相似的另外一方面体现在软件生态的层次清晰度，这个拿升腾来专门说，升腾 Ascend C 的发布时间在 年 5 月 6 日，而反观英伟达 CUDA 的发布时间是 年。什么意思呢，很明显可以看到英伟达的软件生态是以 CUDA 为基础然后层层往上叠的，而升腾是先有了 CANN，先有了 MindSpore 这些 "高层建筑"，然后往下才有了 Ascend C，这种软件生态的层次结构就没有那么清晰，当然这种说法也只是基于时间上的，这并不影响它在空间上还是具备不错的软件生态层次结构。

       再回到 CUDA 本身，需要清楚的一点是，CUDA 其实一开始主要是面向优化计算密集型计算 (Compute-Bound)，因为不管是最开始的通用科学计算还是后来的以 CNN 为主流的深度学习计算都是计算密集型，但是后来 Transformer 又逐渐流行，所以 CUDA 也是不断在 "与时俱进" 在做平衡、做兼顾，比如在 A 开始，CUDA 新增了从 L1 Cache 到 HBM Global Memory 数据直接异步拷贝的指令，其实也是在丰富自身对于访存密集型计算 (Memory-Bound) 的优化。

       我之前写过挺多关于 CUDA 的分享，罗列一些，

       CUDA 是一种硬件强相关的编程模型，要掌握好 CUDA，需要先看懂 GPU 硬件架构，从而映射到 CUDA 内存模型、线程模型上，这点跟 C 语言、跟 C++ 这类 "高级" 编程语言就很不一样，所以很多朋友会觉得 CUDA C 比较难写，特别是要写出高性能的 CUDA C，比较难。确实，这是事实，特别是对于写出高性能的 CUDA C，会涉及资源的高效调度，比如 Shared Memory、L1 Cache 等的调度；会涉及适应硬件架构超参的配置，比如 Thread、Block、Grid 等的配置。

       总之，对于 CUDA 的深入学习，是一门 "稳挣不亏" 的 "买卖"，原因不再过多赘述，主要体现在它的重要性上。

       好了，以上分享了 解读英伟达软件生态 一切的基础CUDA，希望我的分享能对你的学习有一点帮助。

       公众号传送