• CUTLASS

    文章目录

    1、关于 CUTLASS

    CUDA Templates for Linear Algebra Subroutines

    CUTLASS是CUDAC++模板抽象的集合，用于在CUDA中的所有级别和规模上实现 高性能矩阵矩阵乘法（GEMM）和相关计算。
    它结合了分层分解和数据移动的策略，类似于用于实现cuBLAS和cuDNN的策略。
    CUTLASS将这些“移动部分”分解为可重用的模块化软件组件，由C++模板类抽象出来。
    概念并行层次结构的不同级别的基元可以通过自定义平铺大小、数据类型和其他算法策略进行专门化和调整。
    由此产生的灵活性简化了它们作为自定义内核和应用程序中的构建块的使用。

    为了支持各种各样的应用程序，CUTLASS提供了广泛的支持 混合精度计算，提供专门的数据移动和 半精度浮动的乘累加抽象 点（FP16）、BFloat16（BF16）、张量浮点数32（TF32）、 单精度浮点（FP32）， FP32通过张量核心指令进行仿真， 双精度浮动 点（FP64）类型、整数数据类型（4b和8b）和二进制数据类型（1b）。
    CUTLASS演示经纱同步矩阵乘法运算 针对可编程、高通量的张量内核，由 NVIDIA的Volta、图灵、安培和Hopper架构。

    请参阅快速入门指南以快速入门。

    有关操作列表，请参阅功能列表 执行模型层次结构的每个级别都支持。

    CUTLASS 3.0引入了一个新的核心库CuTe，用于描述和操作线程和数据的张量。
    CuTe是 C++ CUDA 模板抽象的集合，用于定义和操作线程和数据的分层多维布局。
    CuTe提供了Layout和Tensor对象，它们紧凑地封装数据的类型、形状、内存空间和布局，同时为用户执行复杂的索引。
    这让程序员专注于算法的逻辑描述，而CuTe为他们做机械簿记。
    使用这些工具，我们可以快速设计、实现和修改所有密集的线性代数运算。

    CuTe的核心抽象是分层多维布局，可以用数据数组来表示张量。
    布局的表示功能强大，足以表示我们实现高效密集线性代数所需的几乎所有内容。
    布局也可以通过函数组合进行组合和操作，我们在其上构建了大量常见操作，如平铺和分区。

    CUTLASS 3.0及更高版本在其模板中的整个GEMM层次结构中都采用了CuTe。
    这极大地简化了设计 并提高了代码的可组合性和易读性。
    更多特定于CuTe的文档可以在其专用文档目录中找到。

    除了GEMM，CUTLASS还通过隐式GEMM算法实现高性能卷积。
    隐式GEMM是将卷积操作表述为GEMM，从而利用CUTLASS的模块化GEMM管道。
    这允许CUTLASS通过重用高度优化的GEMM组件来构建卷积。

    2、CUTLASS 3.5中的新增功能

    CUTLASS 3.5（2024年4月） 是对CUTLASS的更新，添加了：

    • 通过WGMMA针对Hopper SM90A的隐含GEMM卷积+ MA im2col
    • 支持通过2. x API的Ada（SM89）FP8张量内核。需要CUDA 12.4或更新版本。
    • CuTe 和 CUTLASS 3.x 中的 Ampere gather/scatter convolution example
      • 展示如何使用CUTLASS 3. x和CuTe编写和优化自定义内核，以及将卷积实现为GETT特化的一般策略。
      • 实现粗粒度稀疏聚集/分散内核，在安培类张量内核上实现峰值性能。
    • CUTLASS 2. x 中添加了 32x 和 16x 图块大小，以提高窄高和宽短矩阵的性能。
    • 更新CuTe文档为cute::Tensor<>，MMA原子，和大修CuTe GEMM教程系列。
    • 扩展到CuTe以支持L2预取和TMA存储+减少。
    • 删除一些CUTLASS 2. x API头文件的C++11要求。所有CUTLASS文件现在都需要C++17。
    • 修复以大大减少构建警告。
    • 来自社区的更新和错误修复（谢谢！）

    最低要求：

    • 架构：Volta
    • 编译器：必须支持至少C++17
    • CUDA Toolkit 版本：11.4

    从CUTLASS 3.0开始，CUTLASS删除了对以下内容的支持：

    • Maxwell 和 Pascal GPU架构
    • Ubuntu 16.04
    • CUDA 10.2
    • C++语文版本少于17。

    有关版本和更新的详细列表，请参阅CHANGELOG。

    3、性能

    外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传

    CUTLASS原语非常高效。当用于构建设备范围的GEMM内核时， 它们表现出与标量GEMM的cuBLAS相当的峰值性能 计算。

    上图显示了CUTLASS相对于cuBLAS的性能 对于NVIDIA H100（NVIDIA Hopper架构）上的大矩阵尺寸， 一个NVIDIA L40（NVIDIA Ada架构）， 一个NVIDIA A100（NVIDIA安培架构），
    和英伟达A40（英伟达安培架构）。

    CUTLASS 3.0是使用CUDA 12.0工具包编译的。 张量核心操作是使用CUDA的 mma和 wgmma指令。

    外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传

    使用CUTLASS构建块构建设备范围的隐式gem（Fprop、Dgrad和Wgrad）时 内核，CUTLASS性能也相当于cuDNN当运行Resnet-50层在NVIDIA A100 如上图所示。张量核心操作是使用CUDA的 mma指令。

    4、兼容性

    CUTLASS需要C++17主机编译器和 使用CUDA 12.4工具包构建时表现最佳。

    它还兼容CUDA 11.4、CUDA 11.5、CUDA 11.6、CUDA 11.7、CUDA 11.8、CUDA 12.0、CUDA 12.1、CUDA12.2.2、CUDA12.3.1和CUDA12.3.2。

    4、操作系统

    我们测试了以下环境。

    操作系统编译器
    Ubuntu 18.04GCC7.5.0
    Ubuntu 20.04GCC10.3.0
    Ubuntu 22.04GCC11.2.0
    Ubuntu 22.04Clang10.0.0
    Ubuntu 22.04Clang14.0.6
    Ubuntu 22.04Clang17.0.6
    Windows 10.0Visual Studio 2019 v16.11.27

    注意：GCC 8.5.0 有关于折叠表达式和重载运算符的已知回归。建议使用GCC7.5.0 或（首选）GCC>=9。

    5、硬件

    CUTLASS在以下NVIDIA GPU上成功运行，预计它将在基于Volta、图灵、安培、Ada和Hopper架构的NVIDIA GPU上高效运行。

    GPUCUDA计算能力CUTLASS-3所需的最低CUDA工具包
    NVIDIA V100张量核心GPU7.011.4
    NVIDIA TitanV7.011.4
    NVIDIA GeForce RTX 2080 TI，2080,20707.511.4
    NVIDIA T47.511.4
    NVIDIA A100张量核心GPU8.011.4
    NVIDIA A108.611.4
    NVIDIA GeForce RTX 30908.611.4
    NVIDIA GeForce RTX 40908.911.8
    NVIDIA L408.911.8
    NVIDIA H100张量核心GPU9.011.8

    6、目标架构

    一般来说，为一个目标架构生成的PTX代码可以在未来的架构上运行（即，它是前向兼容的）。
    然而，CUDA 12.0引入了 architecture-accelerated 特征 的概念，其PTX没有前向兼容保证。
    几个Hopper PTX指令属于这一类 architecture-accelerated特征，因此需要一个sm_90a的目标架构（注意所附的“a”）。
    有关此指令和其他architecture-accelerated指令的更多详细信息，请参阅 CUDA文档。

    目标架构信息通过cmake标志 CUTLASS_NVCC_ARCHS 传递给CUTLASS。
    为了最大限度地提高 Hopper GH100 的性能，用户需要以 90a作为目标架构构建CUTLASS。
    如果用户不小心构建了一个使用 SM90a功能（例如Hopper Tensor Core Instructions）的内核，使用SM90目标（注意缺少“a”），CTK 12或11.8，内核可能会因运行时错误而失败。

    cmake .. -DCUTLASS_NVCC_ARCHS="90a"

    有关哪些内核需要哪些目标架构的详细信息，请参阅 功能文档。

    7、文档

    CUTLASS在以下文件和随附的文件中进行了描述 Doxygen 文档。

    8、资源

    我们还描述了高效GEMM的结构 2018 GPU技术大会。

    9、构建 CUTLASS

    CUTLASS是一个仅标题的模板库，不需要构建即可供其他用户使用 客户端应用程序应该以CUTLASS的include/目录为目标 路径。

    CUTLASS单元测试、示例和实用程序可以使用CMake构建。
    CMake的最低版本在快速入门指南中给出。
    确保CUDACXX环境变量指向安装的CUDA工具包中的NVCC 在你的系统上。

    $ export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc

    在CUTLASS项目中创建一个构建目录，然后运行CMake。
    默认情况下，CUTLASS将构建内核 对于CUDA架构版本5.0、6.0、6.1、7.0、7.5、8.0、8.6、8.9和9.0。
    要减少编译时间，您可以指定 通过更改CMake配置 CUTLASS_NVCC_ARCHS 设置来构建CUTLASS的架构 .

    $ mkdir build && cd build

$ cmake .. -DCUTLASS_NVCC_ARCHS=80               # compiles for NVIDIA's Ampere Architecture

    在build/目录中，通过使用make构建目标test_unit来编译并运行CUTLASS单元测试。

    单元测试被组织为几个反映CUTLASS顶级命名空间的二进制文件， 它们可以通过make的-j命令行参数并行执行。

    $ make test_unit -j
...
...
...
[----------] Global test environment tear-down
[==========] 946 tests from 57 test cases ran. (10812 ms total)
[  PASSED  ] 946 tests.

    所有测试都应该通过支持的平台，尽管确切的考试数量可能会随着时间的推移而变化。

    10、项目结构

    CUTLASS与实用程序、工具、示例和单元测试一起被安排为仅标头库。
    Doxygen 文档提供了一个完整的文件列表，类， 和CUTLASS项目中定义的模板概念。

    源代码组织的详细说明可以在 CUTLASS文档，但几个主要组成部分总结如下。

    11、CUTLASS模板库

    include/                     # client applications should target this directory in their build's include paths

  cutlass/                   # CUDA Templates for Linear Algebra Subroutines and Solvers - headers only

    arch/                    # direct exposure of architecture features (including instruction-level GEMMs)

    conv/                    # code specialized for convolution

    epilogue/                # code specialized for the epilogue of gemm/convolution

    gemm/                    # code specialized for general matrix product computations

    layout/                  # layout definitions for matrices, tensors, and other mathematical objects in memory

    platform/                # CUDA-capable Standard Library components

    reduction/               # bandwidth-limited reduction kernels that do not fit the "gemm" model

    thread/                  # simt code that can be performed within a CUDA thread
    
    transform/               # code specialized for layout, type, and domain transformations

    *                        # core vocabulary types, containers, and basic numeric operations

  cute/                      # CuTe Layout, layout algebra, MMA/Copy atoms, tiled MMA/Copy

    algorithm/               # Definitions of core operations such as copy, gemm, and operations on cute::tuples

    arch/                    # Bare bones PTX wrapper structs for copy and math instructions

    atom/                    # Meta-information either link to or built from arch/ operators

      mma_atom.hpp           # cute::Mma_Atom and cute::TiledMma

      copy_atom.hpp          # cute::Copy_Atom and cute::TiledCopy

      *sm*.hpp               # Arch specific meta-information for copy and math operations

    *                        # Core library types such as Shape, Stride, Layout, Tensor, and associated operations

    CUTLASS SDK示例

    CUTLASS SDK示例应用CUTLASS模板来实现基本计算。

    工具

    tools/
  library/                   # CUTLASS Instance Library - contains instantiations of all supported CUTLASS templates
    include/
      cutlass/
        library/

  profiler/                  # CUTLASS Profiler         - command-line utility for executing operations in the
                             #                            CUTLASS Library
  
  util/                      # CUTLASS Utilities        - contains numerous helper classes for
    include/                 #                            manging tensors in device memory, reference
      cutlass/               #                            implementations for GEMM, random initialization
        util/                #                            of tensors, and I/O.

    测试

    test/unit/目录包含了用Google Test实现的单元测试 Core API组件的基本用法和CUTLASS GEMM计算的完整测试。

    构建和运行单元测试的说明在快速入门指南中描述。

    12、性能分析

    tools/profiler/ 目录包含一个命令行实用程序，用于启动每个GEMM内核。 它可以构建如下：

    $ make cutlass_profiler -j16

    13、构建所有GEMM和卷积内核（构建时间长）

    默认情况下，每种数据类型、数学指令和布局只实例化一个磁贴大小。
    要实例化所有内容，请在从空build/目录运行CMake时设置以下环境变量。
    请注意，这会导致数以万计的内核和较长的构建时间。 这也会导致较大的二进制大小，并且在某些平台上链接器在构建库时失败。
    因此，强烈建议只生成内核的子集，如下面的小节所示。

    $ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all
...
$ make cutlass_profiler -j16

    14、构建GEMM和卷积内核的子集（减少构建时间）

    要严格编译一个内核或一小部分内核，可以使用带有通配符的逗号分隔内核名称列表来减少内核集。以下示例显示了为NVIDIA Ampere和图灵架构构建一个或一个内核子集：

    构建子集Tensor Core GEMM内核

    要编译具有FP32累积和针对NVIDIA Ampere和图灵架构的FP16输入的Tensor Core GEMM内核子集，请使用以下cmake命令行：

    $ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8
...
$ make cutlass_profiler -j16

    用于分析Tensor Core GEMM内核子集的示例命令行如下：

    ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096

...
=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: gemm
       Operation: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8

          Status: Success
    Verification: ON
     Disposition: Passed

reference_device: Passed
          cuBLAS: Passed

       Arguments: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:column --B=f16:row --C=f32:column --alpha=1  \
                  --beta=0 --split_k_slices=1 --batch_count=1 --op_class=tensorop --accum=f32 --cta_m=256 --cta_n=128  \
                  --cta_k=32 --stages=2 --warps_m=4 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=8 --min_cc=75  \
                  --max_cc=1024

           Bytes: 118489088  bytes
           FLOPs: 115992428544  flops

         Runtime: 1.55948  ms
          Memory: 70.7616 GiB/s

            Math: 74378.8 GFLOP/s

=============================
...

    构建一个CUDA Core GEMM内核

    要编译一个针对NVIDIA Ampere和图灵架构的SGEMM内核，请使用以下cmake命令行：

    $ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1
...
$ make cutlass_profiler -j16

    用于分析单个SGEMM CUDA内核的示例命令行如下：

    $ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096

=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: gemm
       Operation: cutlass_simt_sgemm_128x128_8x2_nn_align1

          Status: Success
    Verification: ON
     Disposition: Passed

          cuBLAS: Passed

       Arguments: --m=3456 --n=4096 --k=4096 --A=f32:column --B=f32:column --C=f32:column --alpha=1 --beta=0 --split_k_slices=1  \
                  --batch_count=1 --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4  \
                  --warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024

           Bytes: 180355072  bytes
           FLOPs: 115992428544  flops

         Runtime: 6.73655  ms
          Memory: 24.934 GiB/s

            Math: 17218.4 GFLOP/s

=============================

    构建张量核心卷积内核的子集

    要编译实现前向传播（fprop）的Tensor核心卷积内核子集，使用FP32累加和针对NVIDIA Ampere和图灵架构的FP16输入，请使用以下cmake命令行：

    $ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16
...
$ make cutlass_profiler -j16

    用于分析Tensor Core卷积内核子集的示例命令行如下：

    $ ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*fprop_optimized_f16 --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3

...
=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: conv2d
       Operation: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc

          Status: Success
    Verification: ON
     Disposition: Passed

reference_device: Passed

       Arguments: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q=224 --pad_h=1 --pad_w=1  \
                  --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f16:nhwc --Filter=f16:nhwc --Output=f32:nhwc  \
                  --conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1  \
                  --eq_gemm_provider=none --op_class=tensorop --accum=f32 --cta_m=128 --cta_n=128 --cta_k=32 --stages=5  \
                  --warps_m=2 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=16 --min_cc=80 --max_cc=1024

           Bytes: 1130659840  bytes
           FLOPs: 118482796544  flops

         Runtime: 0.711496  ms
          Memory: 1479.99 GiB/s

            Math: 166526 GFLOP/s

=============================
...

    构建一个卷积CUDA内核

    要编译和运行一个实现前向传播（fprop）的CUDA Core卷积内核，该内核具有F32累加和针对NVIDIA Ampere和图灵架构的FP32输入，请使用以下cmake命令行：

    $ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
...
$ make cutlass_profiler -j16

    用于分析一个CUDA Core卷积内核的示例命令行：

    $ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3

=============================
  Problem ID: 1

        Provider: CUTLASS
   OperationKind: conv2d
       Operation: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc

          Status: Success
    Verification: ON
     Disposition: Passed

reference_device: Passed

       Arguments: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q=224 --pad_h=1 --pad_w=1  \
                  --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc  \
                  --conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1  \
                  --eq_gemm_provider=none --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4  \
                  --warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024

           Bytes: 2055798784  bytes
           FLOPs: 118482796544  flops

         Runtime: 7.34266  ms
          Memory: 260.752 GiB/s

            Math: 16136.2 GFLOP/s

=============================

    15、有关编译CUTLASS内核 和 CUTLASS Profiler的更多详细信息

    关于

    CUTLASS由NVIDIA Corporation作为开源软件发布 3条款“新”BSD许可证。

    贡献者

    CUTLASS开发者和贡献者的官方列表可在此处获得：贡献者。

    版权所有

    版权所有（c）2017-2024 NVIDIA CORPORATION&AFFILIATES。保留所有权利。SPDX-License-Identifier：BSD-3-Clause

      Redistribution and use in source and binary forms, with or without
  modification, are permitted provided that the following conditions are met:

  1. Redistributions of source code must retain the above copyright notice, this
  list of conditions and the following disclaimer.

  2. Redistributions in binary form must reproduce the above copyright notice,
  this list of conditions and the following disclaimer in the documentation
  and/or other materials provided with the distribution.

  3. Neither the name of the copyright holder nor the names of its
  contributors may be used to endorse or promote products derived from
  this software without specific prior written permission.

  THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
  AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
  IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
  DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
  FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
  DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
  SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
  CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
  OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
  OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

    2024-07-14（日）

  • 相关阅读:
    depot_tools原理和实现
    爆肝！阿里最新版的Spring Security源码手册，强行霸占GitHub榜首
    excel文档打不开怎么修复?
    java多线程深度修炼
    python必会的10个知识点
    java代码proguard代码混淆GUI使用，附带混淆map映射
    Celery笔记九之task运行结果查看
    Spring Boot简介
    C++学习寄录（八.继承）
    架构体系-黑马学习2：-业务幂等性技术架构体系(1)
  • 原文地址：https://blog.csdn.net/lovechris00/article/details/140413267