Cuda

cuda toolkit 和驱动版本的对应关系: https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index... 查看 cuda toolkit 历史版本：https://developer.nvidia.com/cuda-toolkit-archive nv 驱动下载地址：https://www.nvidia.com/download/index.aspx?lang=en-us

转置是深度模型中利用最宽泛的算子之一。各种深度学习框架崛起的明天，它被“隔离”到了框架的底层，不再显式的沉闷于用户背后。谈及转置，就不得不先聊聊卷积的优化算法以及它们所须要的图像数据格式。AUTHOR：vector  小何 01图像数据格式的前世今生目前业界对于卷积的实现次要有4种形式，大白算法(Direct Convolution)，空洞卷积(Winograd Convolution)，隐式矩阵乘卷积(Implicit GEMM Convolution)和疾速傅里叶变换(FFT Convolution)，前两个须要的数据格式是通道前置(NCHW)，而后两者须要通道后置(NHWC or NCxHWx)。 1.1 Direct Convolution vs Implicit GEMM Convolution大白算法，顾名思义，就是用最直观的形式对原始数据(Feature Map)和卷积核(Filter)做相互关(Cross Correlation)计算。这种算法的实质逻辑就是卷积核在原始数据上一直地滑动计算出后果。从内存加载数据到寄存器的过程中，为了保障访存效率，须要H, W维度间断，因而对于大白算法来说，NCHW数据格式显然占优。隐式矩阵乘卷积，还有另外一个耳熟能详的名字 ------ im2col。既然卷积本质上就是输出通道维度上的累加，那么无妨以此作为优化突破口，把累加维度后置，不便SIMD指令访存计算。对于此种优化算法来说，NHWC数据格式更占优。 1.2 CPU vs GPU对于Intel系列的CPU来说，得益于SIMD指令集SSE系列，AVX系列的加持，隐式矩阵乘在性能上全方位碾压了大白算法。因而许多CPU线性代数计算库会以NHWC作为规范数据格式。 以Nvidia为代表的GPU常常被用于各种高性能计算场景，在volta架构之前，N卡的计算外围是未分家的ALU和FPU，习惯性会将其统称为CUDA Core。既然依然依赖于惯例FPU的计算，那么就脱离不开乘法指令比加法指令执行更慢的魔咒，因而空洞卷积成为了最高效的卷积算法。如果对于精度要求较高或输入输出形态极其非凡的场景来说，高并行度的大白算法是斗争后的抉择。总的来说隐式矩阵乘依然排不上号。 从Volta架构开始，N卡引入了新的硬件构造：张量计算单元(1st generation Tensor Core)，专门用于减速FMA-reduction(Fused Multiply Add with Reduction)类型的算子。同时隔壁Google家也早早地就用上了自研的TPU(Tensor Process Unit)来减速深度模型的训练和推理。这些支流的张量计算单元是通过脉动阵列(Systolic Array)技术实现的，这项古老的技术在寂静了多年之后，直到急需高效FMA-reduction算子的明天再次得以重见天日，随之而来的是隐式矩阵乘算法的疾速崛起。自然而然的，通道后置格局也变得重要了起来。几种数据格式间的高效转换 ------ 转置变得十分重要。 1.3 转置让咱们先来简化一下问题，实际上(N, C, H, W)也能够写作(N, C/x, x, H, W)，而咱们须要的数据排布为(N, C/x, H, W, x)。这个时候咱们把没有发生变化的维度(N和C/x, H和W)合并。咱们就将问题简化成了：咱们不难发现，这个问题的实质实际上就是多个一般的二维矩阵转置。在NCHW转置到NC/xHWx时，针对于不同的数据类型，编译期常量x的值是固定的，罕用的： 因而，对于int8类型的数据来说，咱们待解决的问题就变成了：在这里，咱们次要以针对于int8数据类型的NCHW转置NC/32HW32为切入口，来浅谈CUDA优化。 02初版实现截止到CUDA 11.8 和 Hopper架构之前， CUDA在软件层面分为4个层级网格(Grid)，线程块(ThreadBlock)，线程束(Warp)，线程(Thread)。每个流处理器有4个线程束调度器(Warp Scheduler)，因而每个线程块调配128线程个别是最高效的调配计划。接下来，定好每个层级须要解决哪局部数据即可。在这个问题中，咱们自然而然会想到每个Block解决32行数据，每个线程负责搬运一个数，那么图示如下：（点击查看大图） 接下来，咱们只须要计算出每个线程须要从何取数又放回何处即可。先来定义形容各层级Shape的构造体：（点击查看大图） 在定义计算存取数偏移量的仿函数。抉择封装成仿函数是为了不便偏特化，以便把一些常量计算放到编译期进行，缩小运行时的冗余计算指令：（点击查看大图）（点击查看大图） 依据咱们最后的分块策略，初版的Kernel实现如下：（点击查看大图） 其中 GlobalReadInitializer 和 GlobalWriteInitializer 用来获取每个线程块读与写全局内存的首地址。GlobalReader 和 GlobalWriter 则依据二维排布的行和列来计算每个线程读与写的地址。可能会有同学好奇为何全副应用uint64_t类型，这样做是为了对指针地址做运算时，缩小各种clamp相干和cast相干的指令，能够最大化缩小指令数量。当然对于64bit运算，都是分高下位的，所以胆怯计算周期变久的同学大可放心使用。当输出形态为(1, 32, 576, 960)时，在GeForce 4090上的速度为69.05us，看起来的确对得起它的龟速。性能实现完了，接下来咱们开启性能优化之旅。 ...

装置minianaconda创立虚拟环境 conda create -n yolov5-torch python=3.8激活与退出环境 conda activate yolov5-torchconda deactivate装置pytorch抉择什么版本的cuda？torch官网：https://pytorch.org/# 查看显卡驱动对应的cuda:nvidia-smi# 查问后果为：11.7# 查看runtime对应的cudanvcc --version# 查问后果：没有装置！！！装置nvcc，理论就是装置cudacuda下载地址：https://developer.nvidia.com/...cuda版本抉择：不得高于显卡驱动版本对应的cuda版本起因：https://stackoverflow.com/que...抉择 runfile版本，不会替换掉之前的高版本驱动sudo bash cuda_11.7.0_515.43.04_linux.run# 在这里始终失败，提醒没有装置驱动#卸载之前的驱动sudo service lightdm status #敞开X服务sudo service lightdm stopsudo service lightdm statussudo apt-get --purge remove nvidia*sudo apt-get autoremove nvidia*sudo bash cuda_11.7.0_515.43.04_linux.run# 依然失败sudo reboot# 反复之前的步骤胜利~/.bashrcexport PATH=/usr/local/cuda/bin:$PATHexport LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH#验证source ~/.bashrcnvcc -Vnvidia-smi# 装置多版本cuda，比方装置cuda11.6sudo bash cuda_11.6.0_510.39.01_linux.run # 装置过程中去掉驱动项即可，其余照常装置# 装置实现后 /usr/local 呈现 cuda-11.6/#切换不同版本cudacd /usr/localsudo rm -rf cudasudo ln -s /usr/local/cuda-11.6 /usr/local/cuda# 验证nvcc -V装置torch conda install pytorch torchvision torchaudio pytorch-cuda=11.7 -c pytorch -c nvidia下载yolov5源码 ...

撰文｜郑泽康 InsightFace模型里大量应用了PReLU激活函数，而PReLU的工作模式有两种： PReLU(1)，此时权重alpha的形态为(1, )，等价于一个Elementwise操作。PReLU(channels)，此时权重alpha的形态为(channels, )，和输出特色(N, C, H, W)中C的大小是对应的。此时PReLU等价于一个Binary Broadcast操作。InsightFace模型里的PReLU工作模式是第二种，之前曾经介绍过CUDA Elementwise操作优化，而在Broadcast情景下也存在肯定的优化机会。 1 奢侈实现一个奢侈实现的思维就是在循环外部，依据以后元素的索引，推算出该元素对应须要应用的alpha权重的索引。而后判断以后元素x是否大于0，若大于0则返回x，小于0则返回alpha*x。对应代码如下： template<typename T>__global__ void PReluForwardGpu(const int32_t elem_cnt, const int32_t alpha_size, const int32_t inner_size, const T* x, const T* alpha, T* y) { CUDA_1D_KERNEL_LOOP(i, elem_cnt) { const T x_i = x[i]; const T alpha_i = alpha[(i / inner_size) % alpha_size]; y[i] = x_i > 0 ? x_i : x_i * alpha_i; }}其中： inner_size示意的是通道维前面维度乘积，以NCHW格局为例，inner_size=H*Walpha_size示意通道维大小在CUDA中，整数除法的计算代价是比拟低廉的（https://docs.nvidia.com/cuda/...）对于计算指令耗时这一章中有提到： Integer division and modulo operation are costly as they compile to up to 20 instructions. ...

作者：陈振寰 | 旷视科技 MegEngine 架构师背景近年来，主动混合精度（Auto Mixed-Precision，AMP）技术在各大深度学习训练框架中作为一种应用简略、代价低廉、效果显著的训练减速伎俩，被越来越宽泛地利用到算法钻研中。然而大部分对于混合精度训练的文章个别停留在框架接口介绍、如何防止 FP16 类型带来的精度损失以及如何避免出现 NaN 等根底原理和应用技巧方面，对于将深度学习框架视为黑盒工具的研究员来说的确足够了，然而如果想要再往下多走一步，理解一点更底层的减速细节，那么 GPU 显卡架构、CUDA 编程里的一个个专业名词就很容易让不足背景常识的人摸不着头脑。 本文会以混合精度训练背地波及的 Tensor Core 为终点，联合代码实例，帮忙读者对框架层面应用 Tensor Core 进行训练减速的细节乃至 CUDA 编程有一些根本的意识。 Tensor Core 原理首先还是简略介绍一下 混合精度 和 Tensor Core 是什么。混合精度是指在底层硬件算子层面，应用半精度（FP16）作为输出和输入，应用全精度（FP32）进行两头后果计算从而不损失过多精度的技术，而不是网络层面既有 FP16 又有 FP32。这个底层硬件层面其实指的就是 Tensor Core，所以 GPU 上有 Tensor Core 是应用混合精度训练减速的必要条件。 图 1 Tensor Core 直译为张量外围，其物理含意是 NVIDIA GPU 上一块非凡的区域（如图 2 中大块深绿色局部所示），与其位置相似的有一般的 CUDA Core（浅绿色和小块深绿色局部）以及最新的 RT Core（Ray Tracing，光追外围，浅黄色局部）。CUDA Core 个别蕴含多个数据类型，每个数据类型蕴含多个小外围，比方图中的 INT32 Core 和 FP32 Core 就各有 4×16 个，在计算专用卡上还可能会蕴含 FP64 Core（比方 V100 和 A100 显卡），而 Tensor Core 在架构图和接口上则没有具体的辨别，能够视作 GPU 上一块较为独立的计算单元（尽管理论外部有肯定的辨别）。 ...

作者：马骏 | 旷视 MegEngine 架构师前言单精度矩阵乘法（SGEMM）简直是每一位学习 CUDA 的同学绕不开的案例，这个经典的计算密集型案例能够很好地展现 GPU 编程中罕用的优化技巧，而是否写出高效率的 SGEMM Kernel，也是反映一位 CUDA 程序员对 GPU 体系结构的了解水平的优良考题。本文将具体介绍 CUDA SGEMM 的优化伎俩，适宜认真浏览过 《CUDA C++ Programming Guide》，具备肯定 CUDA 编程根底的同学浏览，心愿能给谋求极致性能的同学们一些启发。 CUDA 矩阵乘法优化伎俩详解Naive 实现的剖析：到底差在哪里？笔者面试过不少具备 CUDA 编程教训的校招同学，当发问应用 CUDA 编写一个 SGEMM Kernel 的时候，通常会取得这么一个答案： __global__ void matrixMul(const float *A, const float *B, float *C, int M, int N, int K) { int tx = blockIdx.x * blockDim.x + threadIdx.x; int ty = blockIdx.y * blockDim.y + threadIdx.y; if(ty < M && tx < N) { float c = 0; for(int i = 0; i < K; ++i){ c += A[ty * K + i] * B[i * N + tx]; } C[ty * N + tx] = c; }}这样一个 Naive 的 Kernel 当然不是笔者所期待的，因为这个 Kernel 的性能根本能够判定连 cublas 的 1/10 都不到，显然不合乎咱们谋求高性能的需要。那么这个 Naive 的实现到底差在哪呢？ ...

太长不看版python3.8 (必须应用ubuntu 2004)看前一篇文章。python3.6 (ubuntu 2004/1804均可)下载文件 https://tmp.link/room/6113c0214ddfb应用pip装置 pip install --user ./*.whl其余版本python不反对其余操作系统兴许能够霸王硬上弓心愿晓得文件起源或者离线应用看注释注释前一篇文章写了python3.8/CUDA 11的环境下装置，尽管非常不便，然而这种办法局限于python3.8，不反对其余任意版本的python。如果强行pip install nvidia-tensorflow，会提醒无奈找到该库。 对于python3.6，笔者找到了一些文章，间接应用whl文件装置tensorflow并分享了这些whl文件。认真看文件列表，能够发现这些轮子就是pip install的时候下载的文件，这些文件很有可能就是晚期官网保护的然而不再公布的whl文件。 通过检索，笔者找到了这些文件的官网发行记录，为英伟达公布的tensorflow wheel包。网页上记录了每个发行版本所依赖的各个库的版本，以及反对的操作系统、python解释器版本。 翻阅记录可知，老黄官网保护的这一个版本只反对3.6/3.8两个版本的python解释器，以及2004/1804两个版本的ubuntu。在20.11版本及以前，只反对python3.6，ubuntu反对1804/2004两个版本；从20.12版本开始，只反对python3.8以及ubuntu 2004。然而这个页面并没有给出下载链接，所有的超链接都很诚恳的链接到了版本记录。 通过进一步检索，笔者找到了其中一个whl的下载地址： https://developer.download.nvidia.cn/compute/redist/nvidia-tensorflow/关上链接能够看到全版本的whl，依据发行日志下载本人所需的版本即可。对于列表中所列举的其余依赖，只须要将链接中的nvidia-tensorflow替换为对应的名称即可，例如 https://developer.download.nvidia.cn/compute/redist/nvidia-cuda-nvcc/然而列表中有一个小坑，下载nvidia-cublas-cupti时，须要替换的内容为nvidia-cuda-cupti，否则会显示404. 因为一一替换下载工作量较大，就参考前人把文件下载之后分享给大家，点击此处下载即可。 下载完所有文件后，进入文件夹，应用pip装置即可。过程中依然须要下载一些罕用依赖，如果网络不佳，须要换源。 # 强烈建议虚拟环境操作，或应用--user参数pip install ./*.whl最初关上python验证即可 import tensorflow as tftf.enable_eager_execution()a = tf.random.uniform([1000, 1000])b = tf.random.uniform([1000, 1000])tf.matmul(a, b)

作者：章晓 | 旷视 MegEngine 架构师 一、前言2020 年 5 月 Nvidia 公布了新一代的 GPU 架构安培（Ampere）。其中和深度学习关系最亲密的莫过于性能强劲的第三代的 TensorCore ，新一代的 TensorCore 反对了更为丰盛的 DL（Deep Learning）数据类型，包含了新的 TesorFloat-32（TF32），Bfloat16（BF16）计算单元以及 INT8, INT4 和 INT1 的计算单元，这些计算单元为 DL 推理提供了全面的反对。为了施展这些计算单元的能力，以往会由资深的 HPC 工程师手写 GPU 汇编实现的卷积、矩阵乘算子来开掘硬件的能力。然而凭借人力手工优化算子的形式曾经没有方法应答如此多的数据类型，因而对于 DL 利用的优化慢慢地越来越依赖一些自动化的工具，例如面向深度学习畛域的编译器。在这样的趋势下， Nvidia 开发了线性代数模板库 CUTLASS ，形象了一系列高性能的根本组件，能够用于生成各种数据类型，各种计算单元的卷积、矩阵乘算子。 MegEngine 在 CUTLASS 的根底上进行了二次开发，能够高效地开发新的高性能的算子，疾速地迁徙到新的 GPU 架构。在上一篇 文章 中，咱们曾经简略介绍了 MegEngine 的底层卷积算子实现的应用办法，而本文将会深刻介绍 MegEngine CUDA 平台的底层卷积算子的实现原理，并将会对 Nvidia CUTLASS 的 Implicit GEMM 卷积 文档 进行解读和补充。 因而，读者在浏览本文之前必须要理解的 CUDA 常识有： 拜访全局存储（Global Memory）时，同一 Warp 中的相邻线程拜访间断的地址，访存申请会被合并，合并的访存可能最大化 Global Memory 的吞吐。拜访 Global Memory 时，尽可能应用最宽的数据类型（float4）进行拜访，这样能够最大化访存指令的利用率。CUDA 的共享存储（Shared Memory）依照每 4Bytes 划分为一个 bank，共分为 32 个 bank。当同一 Warp 中的线程拜访同一 bank 的不同地址时会发生冲突（bank conflict）。无 bank conflict 的访存模式能力最大化 Shared Memory 的吞吐。GPU 有显存（Global Memory）、L2、L1（Shared Memory）、寄存器 4 个档次的存储，间接拜访显存的提早很高，在优化 GEMM、Convolution 这样的计算密集型算子时，须要 ...