cutlass
CUTLASS
CUTLASS 3.6.0 - 2024 年 10 月
CUTLASS は、CUDA 内のあらゆるレベルとスケールで高性能行列間乗算 (GEMM) および関連する計算を実装するための CUDA C++ テンプレート抽象化のコレクションです。これには、cuBLAS および cuDNN の実装に使用されるものと同様の、階層分解とデータ移動のための戦略が組み込まれています。 CUTLASS は、これらの「可動部分」を、C++ テンプレート クラスによって抽象化された再利用可能なモジュール式ソフトウェア コンポーネントに分解します。 概念的な並列化階層のさまざまなレベルのプリミティブは、カスタム タイル サイズ、データ型、その他のアルゴリズム ポリシーによって特殊化および調整できます。結果として得られる柔軟性により、カスタム カーネルおよびアプリケーション内のビルディング ブロックとしての使用が簡素化されます。
さまざまなアプリケーションをサポートするために、CUTLASS は混合精度計算の広範なサポートを提供し、半精度浮動小数点 (FP16)、BFloat16 (BF16)、Tensor Float 32 (TF32)、単精度浮動小数点 (FP32)、Tensor コア命令による FP32 エミュレーション、倍精度浮動小数点(FP64) 型、整数データ型 (4b および 8b)、およびバイナリ データ型 (1b)。 CUTLASS は、NVIDIA の Volta、Turing、Ampere、および Hopper アーキテクチャによって実装されたプログラマブルな高スループットTensor コアをターゲットとしたワープ同期行列乗算演算を示します。
すぐに開始するには、クイック スタート ガイドを参照してください。
実行モデル階層の各レベルでサポートされる操作のリストについては、機能リストを参照してください。
CUTLASS 3.0 では、スレッドとデータのテンソルを記述および操作するための新しいコア ライブラリ CuTe が導入されました。 CuTe は、スレッドとデータの階層的な多次元レイアウトを定義および操作するための C++ CUDA テンプレート抽象化のコレクションです。 CuTe は、ユーザーのために複雑なインデックス付けを実行しながら、データの型、形状、メモリ空間、レイアウトをコンパクトにパッケージ化するLayoutオブジェクトとTensorオブジェクトを提供します。これにより、プログラマは、CuTe が機械的な簿記を行う一方で、アルゴリズムの論理的な記述に集中できます。これらのツールを使用すると、すべての高密度線形代数演算を迅速に設計、実装、および変更できます。
CuTe の中核となる抽象化は、テンソルを表すデータ配列で構成できる階層的な多次元レイアウトです。レイアウトの表現は、効率的な高密度線形代数を実装するために必要なほぼすべてを表現できるほど強力です。レイアウトは、機能合成によって組み合わせたり操作したりすることもできます。これに基づいて、タイル化やパーティション化などの一般的な操作の大規模なセットが構築されます。
CUTLASS 3.0 以降では、テンプレートの GEMM 階層全体にわたって CuTe が採用されています。これにより、設計が大幅に簡素化され、コードの構成可能性と可読性が向上します。 CuTe に固有のその他のドキュメントは、専用のドキュメント ディレクトリにあります。
GEMM に加えて、CUTLASS は暗黙的な GEMM アルゴリズムを介して高性能の畳み込みを実装します。 Implicit GEMM は、畳み込み演算を GEMM として定式化し、CUTLASS のモジュール式 GEMM パイプラインを利用します。これにより、CUTLASS は高度に最適化された GEMM コンポーネントを再利用して畳み込みを構築できます。
CUTLASS 3.6.0 は、以下を追加する CUTLASS のアップデートです。
ホッパー構造のスパースGEMM。
FP16
FP8
INT8
TF32
CUTLASS 3.x コンボリューションkernel::ConvUniversal API をリファクタリングしてgemm::GemmUniversalと一致させます。現在、3.x コンボリューション API はベータ API とはみなされなくなりました。
改良された混合入力 GEMM とINT4 x FP8スケールオンリー モード用のルックアップ テーブル実装。
Top-K 選択用の EVT ノードと、それらを使用したソフトマックスおよび GEMM の例。
新しいホッパー機能を利用して 2 つの連続したカーネルとそれに対応するドキュメントを高速化するプログラム依存起動 (PDL)。
新しいデバッグ ツール synclog は、カーネル内からすべての同期イベントをファイルにダンプします。詳細については、synclog のドキュメントを参照してください。
グループ化された GEMM 用の新しい TMA 対応エピローグは、パフォーマンスの大幅な向上と EVT サポートをもたらします。
SIMT 対応のポインター配列エピローグ。
Grouped GEMM およびその他の最適化のための新しい Ping-Pong カーネル スケジュール。
CUTLASS プロファイラー カーネルの新しいインスタンス化戦略と、CUTLASS プロファイラーのインスタンス化レベルに関するドキュメントの改善。
cutlass::bfloat16_tの比較と計算のための新しいハードウェア サポート
Windows でのhalf_tに対する isnan の使用を修正しました。
最小要件:
アーキテクチャ: ボルタ
コンパイラ: 少なくとも C++17 をサポートする必要があります
CUDA ツールキットのバージョン: 11.4
CUTLASS 3.0 以降、CUTLASS は以下のサポートを削除しました。
Maxwell および Pascal GPU アーキテクチャ
Ubuntu 16.04
CUDA 10.2
C++ 言語バージョン 17 未満。
リリースとアップデートの詳細なリストについては、CHANGELOG を参照してください。
CUTLASS プリミティブは非常に効率的です。 デバイス全体の GEMM カーネルの構築に使用すると、スカラー GEMM 計算の cuBLAS に匹敵するピーク パフォーマンスを示します。上の図は、CUTLASS 3.1 以降、NVIDIA H100 (NVIDIA Hopper アーキテクチャ) での継続的な CUTLASS パフォーマンスの向上を示しています。 CUTLASS 3.5.1 は CUDA 12.5u1 ツールキットを使用してコンパイルされました。 Tensor コアの操作は、CUDA の mma および wgmma 命令を使用して実装されます。
CUTLASS ビルディング ブロックを使用してデバイス全体の暗黙的な gemm (Fprop、Dgrad、および Wgrad) カーネルを構築する場合、上の図に示すように、NVIDIA A100 で Resnet-50 レイヤーを実行する場合、CUTLASS のパフォーマンスも cuDNN に匹敵します。 Tensor コアの操作は、CUDA の mma 命令を使用して実装されます。
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、CUDA 12.2.2、CUDA 12.3.1、および CUDA 12.3.2 とも互換性があります。
以下の環境で動作確認を行っております。
| オペレーティング·システム | コンパイラ |
|---|---|
| Ubuntu 18.04 | GCC 7.5.0 |
| Ubuntu 20.04 | GCC 10.3.0 |
| Ubuntu 22.04 | GCC 11.2.0 |
| Ubuntu 22.04 | クラン 10.0.0 |
| Ubuntu 22.04 | クラン 14.0.6 |
| Ubuntu 22.04 | クラン 17.0.6 |
| Windows 10.0 | Visual Studio 2019 v16.11.27 |
注: GCC 8.5.0 には、fold 式とオーバーロードされた演算子に関する既知のリグレッションがあります。 GCC 7.5.0 または (推奨) GCC >= 9 を使用することをお勧めします。
CUTLASS は次の NVIDIA GPU で正常に実行され、Volta、Turing、Ampere、Ada、および Hopper アーキテクチャ ベースの NVIDIA GPU で効率的であることが期待されます。
| GPU | CUDA コンピューティング機能 | CUTLASS-3 に必要な最小 CUDA ツールキット |
|---|---|---|
| NVIDIA V100 Tensor コア GPU | 7.0 | 11.4 |
| NVIDIA TitanV | 7.0 | 11.4 |
| NVIDIA GeForce RTX 2080 TI、2080、2070 | 7.5 | 11.4 |
| NVIDIA T4 | 7.5 | 11.4 |
| NVIDIA A100 Tensor コア GPU | 8.0 | 11.4 |
| NVIDIA A10 | 8.6 | 11.4 |
| NVIDIA GeForce RTX 3090 | 8.6 | 11.4 |
| NVIDIA GeForce RTX 4090 | 8.9 | 11.8 |
| NVIDIA L40 | 8.9 | 11.8 |
| NVIDIA H100 Tensor コア GPU | 9.0 | 11.8 |
一般に、1 つのターゲット アーキテクチャに対して生成された PTX コードは、将来のアーキテクチャでも実行できます (つまり、前方互換性があります)。 ただし、CUDA 12.0 では、PTX に前方互換性が保証されていない「アーキテクチャ アクセラレーション機能」の概念が導入されました。いくつかの Hopper PTX 命令は、このカテゴリのアーキテクチャ高速化機能に該当するため、 sm_90aターゲット アーキテクチャが必要です (「a」が追加されていることに注意してください)。この命令およびその他のアーキテクチャ高速化命令の詳細については、CUDA ドキュメントを参照してください。
ターゲット アーキテクチャ情報は、cmake フラグCUTLASS_NVCC_ARCHSを介して CUTLASS に渡されます。 Hopper GH100 のパフォーマンスを最大化するには、ターゲット アーキテクチャとして90aを使用して CUTLASS を構築する必要があります。ユーザーが CUDA Toolkit 12 または 11.8 のいずれかを使用して、SM90 ターゲット (「a」がないことに注意してください) を使用して SM90a 機能 (例: Hopper Tensor Core 命令) を使用するカーネルを誤ってビルドした場合、カーネルはランタイムで失敗することが予想されます。エラー。
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
どのカーネルがどのターゲット アーキテクチャを必要とするかの詳細については、機能ドキュメントを参照してください。
CUTLASS については、次のドキュメントとそれに付随する Doxygen ドキュメントで説明されています。
クイック スタート ガイド - CUTLASS のビルドと実行
機能 - CUTLASS で利用できる機能をまとめています。
CUDA での効率的な GEMM - CUDA で GEMM カーネルを効率的に実装する方法について説明します。
CUTLASS 3.x 設計 - CUTLASS 3.x 設計、その利点、および CuTe を使用してより多くの構成可能なコンポーネントを作成できる方法について説明します。
GEMM API 3.x - CUTLASS 3.x GEMM モデルと C++ テンプレートの概念について説明します
GEMM API 2.x - CUTLASS 2.x GEMM モデルと C++ テンプレートの概念について説明します
暗黙的な GEMM 畳み込み - CUTLASS での 2 次元および 3 次元の畳み込みについて説明します。
コード構成 - CUTLASS プロジェクトの構成と内容を説明します。
用語 - コードで使用される用語について説明します。
プログラミング ガイドライン - 効率的な最新の CUDA C++ を作成するためのガイドライン
基本型 - CUTLASS で数値と配列を表すために使用される基本的な C++ クラスについて説明します。
レイアウト - メモリ内の行列とテンソルのレイアウトを説明します。
タイル イテレーター - メモリ内の行列のタイルを反復処理するための C++ の概念について説明します。
CUTLASS Profiler - コマンドライン駆動のプロファイリング アプリケーション
CUTLASS ユーティリティ - 迅速な開発を促進するために使用される追加のテンプレート
依存カーネルの起動 - 同じストリーム内で依存カーネルのオーバーラップを可能にする Hopper の新機能と、それが CUTLASS でどのように使用されるかについて説明します。
GPU Technology Conference 2018 での講演でも、効率的な GEMM の構造について説明しました。
CUTLASS: CUDA 内のすべてのレベルおよびスケールでの高密度線形代数のためのソフトウェア プリミティブ
NVIDIA A100 で Tensor コアを絶対限界まで押し上げる CUDA カーネルの開発
CUTLASS の Tensor コアによる畳み込みの高速化
CUTLASS で Tensor コアの使用率を高めて後方データ勾配を加速する
CUTLASS: Python API、機能拡張、および NVIDIA ホッパー
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 ビルド && cd ビルド $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # NVIDIA の Ampere アーキテクチャ用にコンパイルします
build/ディレクトリから、make でターゲットのtest_unitビルドすることにより、CUTLASS 単体テストをコンパイルして実行します。
単体テストは、CUTLASS のトップレベル名前空間をミラーリングするいくつかのバイナリとして構成されており、make の-jコマンド ライン引数を介して並列実行できます。
$ make test_unit -j ... ... ... [----------] グローバルテスト環境の解体 [==========] 57 のテスト ケースから 946 のテストが実行されました。 (合計 10812 ミリ秒) [ 合格 ] 946 テスト。
すべてのテストはサポートされているプラットフォームで合格する必要がありますが、正確なテスト数は時間の経過とともに変化する可能性があります。
CUTLASS は、ユーティリティ、ツール、例、単体テストとともにヘッダーのみのライブラリとして配置されています。 Doxygen のドキュメントには、CUTLASS プロジェクトで定義されているファイル、クラス、およびテンプレートの概念の完全なリストが記載されています。
ソース コード構成の詳細な説明は 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 テンプレートを適用して基本的な計算を実装します。
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 で実装された単体テストで構成されており、コア API コンポーネントの基本的な使用法と CUTLASS GEMM 計算の完全なテストを示します。
単体テストの構築と実行の手順については、クイックスタート ガイドに記載されています。
tools/profiler/ディレクトリには、各 GEMM カーネルを起動するためのコマンドライン ユーティリティが含まれています。次のように構築できます。
$ makecutlass_profiler -j16
デフォルトでは、データ型、演算命令、レイアウトごとに 1 つのタイル サイズのみがインスタンス化されます。すべてをインスタンス化するには、空のbuild/ディレクトリから CMake を実行するときに次の環境変数を設定します。これにより、数万のカーネルが必要になり、ビルド時間が長くなることに注意してください。これにより、バイナリ サイズが大きくなり、一部のプラットフォームではリンカーがライブラリのビルドに失敗する可能性があります。したがって、以下のサブセクションで説明するように、カーネルのサブセットのみを生成することを強くお勧めします。
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ makecutlass_profiler -j16
厳密に 1 つのカーネルまたは少数のカーネル セットをコンパイルするには、ワイルドカード文字を含むカーネル名のカンマ区切りのリストを使用して、カーネルのセットを減らすことができます。次の例は、NVIDIA Ampere および Turing アーキテクチャ用のカーネルの 1 つまたはサブセットを構築する方法を示しています。
NVIDIA Ampere および Turing アーキテクチャをターゲットとする FP32 蓄積と FP16 入力を使用して Tensor コア GEMM カーネルのサブセットをコンパイルするには、以下の cmake コマンド ラインを使用します。
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ makecutlass_profiler -j16
Tensor Core GEMM カーネルのサブセットをプロファイリングするためのコマンド ラインの例は次のとおりです。
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
=============================
問題ID: 1
提供者: カトラス
操作種類: gemm
オペレーション:cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
ステータス: 成功
検証：オン
処分: 合格
参照デバイス: 渡されました
cuBLAS: 合格
引数: --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
バイト: 118489088バイト
フロップ数: 115992428544 フロップス
実行時間: 1.55948 ミリ秒
メモリ: 70.7616 GiB/秒
計算: 74378.8 GFLOP/秒
=============================
...NVIDIA Ampere および Turing アーキテクチャをターゲットとする 1 つの SGEMM カーネルをコンパイルするには、以下の cmake コマンド ラインを使用します。
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ makecutlass_profiler -j16
単一の SGEMM CUDA カーネルをプロファイリングするためのコマンド ラインの例は次のとおりです。
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
=============================
問題ID: 1
提供者: カトラス
操作種類: gemm
オペレーション:cutlass_simt_sgemm_128x128_8x2_nn_align1
ステータス: 成功
検証：オン
処分: 合格
cuBLAS: 合格
引数: --m=3456 --n=4096 --k=4096 --A=f32:column --B=f32:column --C=f32:column --alpha=1 --beta=0 --分割k_スライス=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
バイト: 180355072バイト
フロップ数: 115992428544 フロップス
実行時間: 6.73655 ミリ秒
メモリ: 24.934 GiB/秒
計算: 17218.4 GFLOP/秒
=============================NVIDIA Ampere および Turing アーキテクチャをターゲットとした FP32 累積および FP16 入力を備えた順伝播 (fprop) を実装する Tensor コア コンボリューション カーネルのサブセットをコンパイルするには、以下の cmake コマンド ラインを使用します。
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $ makecutlass_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
...
=============================
問題ID: 1
提供者: カトラス
操作の種類: conv2d
オペレーション:cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
ステータス: 成功
検証：オン
処分: 合格
参照デバイス: 渡されました
引数: --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=最適化 --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
バイト: 1130659840バイト
フロップ数: 118482796544 フロップス
実行時間: 0.711496 ミリ秒
メモリ: 1479.99 GiB/秒
数学: 166526 GFLOP/秒
=============================
...NVIDIA Ampere および Turing アーキテクチャをターゲットとする F32 蓄積と FP32 入力を備えた順伝播 (fprop) を実装する 1 つの CUDA コア コンボリューション カーネルをコンパイルして実行するには、以下の cmake コマンド ラインを使用します。
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ makecutlass_profiler -j16
1 つの 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
=============================
問題ID: 1
提供者: カトラス
操作の種類: conv2d
オペレーション:cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
ステータス: 成功
検証：オン
処分: 合格
参照デバイス: 渡されました
引数: --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=最適化 --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
バイト: 2055798784 バイト
フロップ数: 118482796544 フロップス
実行時間: 7.34266 ミリ秒
メモリ: 260.752 GiB/秒
計算: 16136.2 GFLOP/秒
=============================CUTLASS カーネルの選択的コンパイルに関するその他の CMake の例については、リンクをクリックしてください。
GEMM CMake の例
暗黙的な GEMM 畳み込み CMake の例
CUTLASS Profiler の詳細については、ここで説明します。
CUTLASS は、3 条項の「新しい」BSD ライセンスに基づいてオープン ソース ソフトウェアとして NVIDIA Corporation からリリースされています。
CUTLASS 開発者および貢献者の公式リストは、CONTRIBUTORS から入手できます。
著作権 (c) 2017 - 2024 NVIDIA CORPORATION および AFFILIATES。無断転載を禁じます。 SPDX ライセンス識別子: 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.
