皮皮网

1.CUDA学习：基础知识小结
2.银河系CUDA编程指南(2.5)——NVCC与PTX
3.GPU编程3：CUDA环境安装和IDE配置
4.Windows用Visual Studio 2022编译支持CUDA12的工程gromacs 2024.2教程
5.入坑指南| OpenCV4.8 + CUDA + 扩展模块支持编译
6.CUDA编程OneFlow Softmax 算子源码解读之WarpSoftmax

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通过事件来实现。reviewboard 源码Debug功能则涉及API错误检测和kernel function的异步执行错误检查。CUDA-MEMCHECK工具集是调试工具的重要组成部分，可以通过特定命令调用。

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

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

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

       NVCC的编译过程分为离线编译和即时编译，通过预处理将源代码分为两部分，分别由不同编译器处理，最终合并为单个object文件。例如，通过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库，并调整编译生成部分，vbole源码确保链接所有需要的.o文件。NVCC命令在编译时会执行链接任务。

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`，clahe源码简化编译过程。

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

Windows用Visual Studio 编译支持CUDA的gromacs .2教程

       为了在 Windows 上使用 Visual Studio 编译 Gromacs .2 版本支持 CUDA，您需要遵循以下步骤。

       首先，安装 Visual Studio ，无论是企业版、专业版还是社区版均可，确保在安装时选择使用 C++ 的桌面开发组件。

       其次，下载并安装 CUDA ，从官方 CUDA Toolkit Archive 获取。

       接着，下载并安装 FFTW3.3.，从 fftw.org 下载相应的库。

       打开命令提示符，解压 FFTW3.3. 的源码，并在目录中建立 build 文件夹。

       进入 build 文件夹，然后在命令提示符中执行编译安装命令。

       修改 CUDA 头文件中的 host_config.h，定位到大约第 行，将版本号从 改为 ，确保编译过程顺利。

       下载 Gromacs .2 的编译源码，从提供的链接获取。

       下载完成后，解压缩源码，进入 build 目录，执行 cmake 命令进行配置。

       在 cmake 配置时，选择合适的 GMX_CUDA_TARGET_SM 参数，根据您的显卡选择 sm_, sm_, sm_, sm_, sm_, sm_, sm_, sm_, sm_, sm_, sm_ 中的一个，我以 sm_ 为例，即 -DGMX_CUDA_TARGET_SM=。

       编译时可能会遇到错误，如 nvcc fatal 错误或 CMake 错误。解决这类问题需要耐心，确保按照配置正确地执行编译过程。

       如果需要比较修改的代码，可以使用 Beyond_Compare 工具进行代码对比，下载地址为提供的链接。

入坑指南| OpenCV4.8 + CUDA + 扩展模块支持编译

       欢迎关注最新版OpenCV4.8的CUDA加速编译指南。在深度学习模型部署速度提升方面，rps 源码CUDA加速技术发挥着关键作用。为了顺利进行编译，首先需确保软件版本满足要求：使用CMake3.或更高版本，以及VS专业版或以上版本。配合ti显卡和CUDA.3版本，将OpenCV4.8源码包与扩展模块源码包解压至D盘。

       操作路径如下：D:\opencv-4.8.0与D:\opencv_contrib-4.8.0。启动CMake进行配置生成，点击Finish完成配置后，选择Generate，生成项目文件。设置扩展模块路径，并在搜索CUDA关键字后，勾选相关选项，再次点击Configure，生成配置文件。

       打开OpenCV.sln文件，切换到Release模式，生成安装文件。预计编译时间较长，耐心等待即可。最终，编译完成后的结果展示了CUDA加速下的性能提升。

       技巧提示：在编译过程中，注意避免下载第三方文件，可手动放置到source/.cache文件夹内的相应位置。无需编译xFeature2D、FaceModel、Test等模块，避免不必要的依赖下载。

       配置完成后，可运行人脸检测案例和YOLOv8姿态评估模型，验证CUDA加速效果。使用以下代码启用GPU推理执行。

       总之，OpenCV4.8通过CUDA支持实现GPU加速，显著提升图像处理程序性能。请参考本人新书《OpenCV应用开发：入门、进阶与工程化实践》获取更多CUDA函数使用知识。

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环境下进行优化的策略。

一文读懂cuda代码编译流程

       cuda代码编译流程详解

       在仅需在服务器上本地编译GPU程序且不考虑跨平台兼容性和程序大小时，使用默认的nvcc命令即可。但若要考虑程序的可移植性和编译后的文件大小，就需要深入理解nvcc编译指令。本文将逐步解析.cu源代码如何转化为可执行文件，揭示GPU与CPU之间的交互。

       以名为simple_add.cu的简单示例程序为例，我们可以通过命令nvcc simple_add.cu -o simple_add编译生成可执行程序。另外，为了保存编译过程，可以使用mkdir simple_add_tmp && nvcc simple_add.cu -o simple_add -keep -keep-dir=./simple_add_tmp，这会将中间文件保存在simple_add_tmp目录中。

       打开目录，可以看到一系列文件，它们与官方编译流程图相对应，包括编译过程的详细记录和生成的临时文件。例如，通过nvcc simple_add.cu -o simple_add -keep -keep-dir=./simple_add_tmp -dryrun命令，可以获取生成每个文件的详细指令。

       在编译过程中，CUDA代码主要处理CUDA kernel的定义和调用，比如__global__ void add(int *a, int *b, int *c, int n)和add<<>>>();。非kernel部分是标准的C++代码。首先，nvcc会将CUDA代码分解到simple_add.cudafe1.cpp中，然后处理kernel的调用，如通过__cudaPushCallConfiguration存储参数，调用`add`函数，该函数最终由`__device_stub__Z3addPiS_S_i`等辅助函数执行。

       CPU编译时，会处理kernel的overhead，也就是kernel的启动开销。在simple_add.cudafe1.stub.c文件中，可以看到`__cudaLaunch`函数被调用，这代表了CUDA运行时如何查找并执行kernel。

       对于GPU编译，.cudafe1.gpu文件包含了CUDA kernel的源代码，经过cicc编译成ptx，再通过ptxas生成cubin，最终整合成fatbin，形成GPU可执行的二进制文件。这些二进制内容存储在可执行文件的.nv_fatbin部分，可以通过工具如`cuobjdump -h`查看。

       GPU程序版本管理很重要，不同的GPU架构（如compute_、V等）需要不同的编译选项。理解这些版本对应关系，可以帮助我们选择正确的编译参数，平衡程序大小和性能需求。

       在实际编译时，如选择compute_，可能有三种生成方式。为了支持多种GPU，需要考虑多个版本的ptx和cubin代码，并确保nvcc支持的编译选项与当前环境兼容。

       总结来说，理解CUDA代码编译过程涉及从源代码到二进制文件的转换，以及如何根据不同GPU版本进行优化。在打包和发布时，需根据实际需求平衡兼容性和性能，这需要对编译选项有深入的理解。

[技术随笔]🛠🛠从源码安装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的功能和应用，进一步拓展其在计算机图形学、三维重建和深度学习等领域的应用。

OpenCV OpenCV 源码编译并实现 CUDA 加速 (Windows)

       本文介绍了如何在Windows系统上使用OpenCV源码自行编译代码文件，实现CUDA加速，以满足对处理时间要求较高的场景。OpenCV是一个跨平台的计算机视觉和机器学习软件库，支持Linux、Windows、Android和Mac OS等操作系统。

       在实际使用中，OpenCV处理数据可能无法满足某些高速场景的需求，这时可以结合CUDA加速。为了实现CUDA加速，需要自行编译支持CUDA的依赖包。在本次文章中，我们将演示如何在Windows环境下使用CMake-gui + VS进行OpenCV源码的编译。

       首先，确保环境准备充分。本次编译平台是Windows 系统，使用CMake-gui + VS进行编译。需要下载两个源码，分别是opencv和opencv_contrib，并保证版本一致，本文使用的版本为4.8.0。将两个文件解压到同一文件夹下。

       然后，利用CMake创建并配置项目。在CMake软件中设置项目源码路径，并创建build文件夹，进行配置。在第一次配置后，输出编译平台选择，本文选择Visual Studio ，编译平台为x。配置中添加opencv_contrib模块引用，选择WITH_CUDA和OPENCV_DNN_CUDA，以及其他相关选项。配置完成后，检查异常并解决。

       在解决异常后，使用Visual Studio打开生成的OpenCV.sln解决方案文件，并运行ALL_BUILD项目。编译完成后，将获得包含依赖项的install文件夹和python_loader文件夹，用于支持Python API和C++ API的使用。

       项目编译完成后，通过cv2.cuda.getCudaEnabledDeviceCount()接口方法检查CUDA设备是否存在。输出结果为1，表明CUDA设备已正确安装，项目编译成功。

       总结，通过本文的步骤，实现了OpenCV源码编译并结合CUDA加速，提高了处理时间要求较高的场景的性能。后续将结合所编译的库进行项目开发与性能对比。