cutlass
CUTLASS
CUTLASS 3.6.0 – Oktober 2024
CUTLASS ist eine Sammlung von CUDA-C++-Vorlagenabstraktionen zur Implementierung der Hochleistungs-Matrix-Matrix-Multiplikation (GEMM) und verwandter Berechnungen auf allen Ebenen und Skalen innerhalb von CUDA. Es umfasst Strategien zur hierarchischen Zerlegung und Datenbewegung, die denen ähneln, die zur Implementierung von cuBLAS und cuDNN verwendet werden. CUTLASS zerlegt diese „beweglichen Teile“ in wiederverwendbare, modulare Softwarekomponenten, die durch C++-Vorlagenklassen abstrahiert werden. Primitive für verschiedene Ebenen einer konzeptionellen Parallelisierungshierarchie können über benutzerdefinierte Kachelgrößen, Datentypen und andere algorithmische Richtlinien spezialisiert und optimiert werden. Die daraus resultierende Flexibilität vereinfacht ihre Verwendung als Bausteine in benutzerdefinierten Kerneln und Anwendungen.
Um eine Vielzahl von Anwendungen zu unterstützen, bietet CUTLASS umfassende Unterstützung für Berechnungen mit gemischter Genauigkeit und bietet spezielle Datenbewegungs- und Multiplikations-Akkumulations-Abstraktionen für Gleitkommazahlen mit halber Genauigkeit (FP16), BFloat16 (BF16), Tensor Float 32 (TF32). Gleitkomma-Typen mit einfacher Genauigkeit (FP32), FP32-Emulation über Tensor-Kernbefehl, Gleitkomma-Typen mit doppelter Genauigkeit (FP64), ganzzahlige Datentypen (4b und 8b) und binäre Datentypen (1b). CUTLASS demonstriert Warp-synchrone Matrixmultiplikationsoperationen, die auf die programmierbaren Tensorkerne mit hohem Durchsatz abzielen, die von NVIDIAs Volta-, Turing-, Ampere- und Hopper-Architekturen implementiert werden.
Sehen Sie sich die Kurzanleitung an, um schnell loszulegen.
Die Liste der Vorgänge, die auf jeder Ebene der Ausführungsmodellhierarchie unterstützt werden, finden Sie in der Funktionsliste.
Mit CUTLASS 3.0 wurde eine neue Kernbibliothek, CuTe, eingeführt, um Tensoren von Threads und Daten zu beschreiben und zu manipulieren. CuTe ist eine Sammlung von C++-CUDA-Vorlagenabstraktionen zum Definieren und Bearbeiten hierarchisch mehrdimensionaler Layouts von Threads und Daten. CuTe stellt Layout und Tensor Objekte bereit, die Typ, Form, Speicherplatz und Layout von Daten kompakt zusammenfassen und gleichzeitig die komplizierte Indizierung für den Benutzer durchführen. Dadurch können sich Programmierer auf die logischen Beschreibungen ihrer Algorithmen konzentrieren, während CuTe die mechanische Buchhaltung für sie übernimmt. Mit diesen Werkzeugen können wir alle Operationen der dichten linearen Algebra schnell entwerfen, implementieren und ändern.
Die Kernabstraktionen von CuTe sind hierarchisch mehrdimensionale Layouts, die mit Datenarrays zur Darstellung von Tensoren zusammengestellt werden können. Die Darstellung von Layouts ist leistungsstark genug, um nahezu alles darzustellen, was wir zur Implementierung einer effizienten dichten linearen Algebra benötigen. Layouts können auch über die funktionale Komposition kombiniert und manipuliert werden, auf der wir eine Vielzahl gängiger Operationen wie Kacheln und Partitionieren aufbauen.
CUTLASS 3.0 und höher übernehmen CuTe in der gesamten GEMM-Hierarchie in seinen Vorlagen. Dies vereinfacht das Design erheblich und verbessert die Zusammensetzbarkeit und Lesbarkeit des Codes. Weitere Dokumentation speziell für CuTe finden Sie im entsprechenden Dokumentationsverzeichnis.
Zusätzlich zu GEMMs implementiert CUTLASS Hochleistungsfaltung über den impliziten GEMM-Algorithmus. Implizites GEMM ist die Formulierung einer Faltungsoperation als GEMM und nutzt dabei die Vorteile der modularen GEMM-Pipeline von CUTLASS. Dies ermöglicht es CUTLASS, Faltungen durch die Wiederverwendung hochoptimierter GEMM-Komponenten zu erstellen.
CUTLASS 3.6.0 ist ein Update für CUTLASS und fügt Folgendes hinzu:
Hopper strukturiert spärlich GEMM.
FP16
FP8
INT8
TF32
Eine Umgestaltung der CUTLASS 3.x- kernel::ConvUniversal -API, um sie mit gemm::GemmUniversal in Einklang zu bringen. Jetzt gilt die 3.x-Faltungs-API nicht mehr als Beta-API.
Ein verbessertes Mixed-Input-GEMM und eine Lookup-Table-Implementierung für den Nur-Skalierungsmodus INT4 x FP8 .
EVT-Knoten für die Top-K-Auswahl und Softmax- und GEMM-Beispiel mit diesen.
Programmatic Dependent Launch (PDL), das eine neue Hopper-Funktion nutzt, um zwei aufeinanderfolgende Kernel und die entsprechenden Dokumentationen zu beschleunigen.
Ein neues Debugging-Tool, synclog, zum Ausgeben aller Synchronisierungsereignisse innerhalb eines Kernels in eine Datei. Weitere Informationen finden Sie in der Synclog-Dokumentation.
Ein neuer TMA-fähiger Epilog für gruppiertes GEMM, der eine deutliche Leistungsverbesserung sowie EVT-Unterstützung bietet.
Ein SIMT-fähiger Zeiger-Array-Epilog.
Ein neuer Ping-Pong-Kernel-Zeitplan für Grouped GEMM und einige andere Optimierungen.
Eine neue Instanziierungsstrategie für CUTLASS-Profiler-Kernel zusammen mit einer verbesserten Dokumentation für die Instanziierungsebene im CUTLASS-Profiler.
Eine neue Hardwareunterstützung für Vergleiche und Berechnungen von cutlass::bfloat16_t
Die Verwendung von isnan unter Windows für half_t wurde korrigiert.
Mindestanforderungen:
Architektur: Volta
Compiler: Muss mindestens C++17 unterstützen
CUDA Toolkit-Version: 11.4
Ab CUTLASS 3.0 hat CUTLASS die Unterstützung für Folgendes entfernt:
Maxwell- und Pascal-GPU-Architekturen
Ubuntu 16.04
CUDA 10.2
C++-Sprachversionen kleiner als 17.
Eine detaillierte Liste der Veröffentlichungen und Updates finden Sie im CHANGELOG.
CUTLASS-Grundelemente sind sehr effizient. Wenn sie zum Erstellen geräteweiter GEMM-Kernels verwendet werden, weisen sie eine mit cuBLAS vergleichbare Spitzenleistung für skalare GEMM-Berechnungen auf. Die obige Abbildung zeigt die kontinuierlichen CUTLASS-Leistungsverbesserungen auf einem NVIDIA H100 (NVIDIA Hopper-Architektur) seit CUTLASS 3.1. CUTLASS 3.5.1 wurde mit dem CUDA 12.5u1 Toolkit kompiliert. Tensor-Core-Operationen werden mithilfe der mma- und wgmma-Anweisungen von CUDA implementiert.
Bei Verwendung von CUTLASS-Bausteinen zum Erstellen geräteweiter impliziter Gemm-Kernel (Fprop, Dgrad und Wgrad) ist die Leistung von CUTLASS auch mit cuDNN vergleichbar, wenn Resnet-50-Schichten auf einem NVIDIA A100 ausgeführt werden, wie in der obigen Abbildung dargestellt. Tensor-Core-Operationen werden mithilfe der mma-Anweisung von CUDA implementiert.
CUTLASS erfordert einen C++17-Host-Compiler und bietet die beste Leistung, wenn es mit dem CUDA 12.4 Toolkit erstellt wird. Es ist auch kompatibel mit 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 und CUDA 12.3.2.
Wir haben die folgenden Umgebungen getestet.
| Betriebssystem | Compiler |
|---|---|
| Ubuntu 18.04 | GCC 7.5.0 |
| Ubuntu 20.04 | GCC 10.3.0 |
| Ubuntu 22.04 | GCC 11.2.0 |
| Ubuntu 22.04 | Klirren 10.0.0 |
| Ubuntu 22.04 | Klirren 14.0.6 |
| Ubuntu 22.04 | Klirren 17.0.6 |
| Windows 10.0 | Visual Studio 2019 v16.11.27 |
Hinweis: GCC 8.5.0 weist bekannte Regressionen in Bezug auf Faltausdrücke und überladene Operatoren auf. Es wird empfohlen, GCC 7.5.0 oder (bevorzugt) GCC >= 9 zu verwenden.
CUTLASS läuft erfolgreich auf den folgenden NVIDIA-GPUs und wird voraussichtlich auf NVIDIA-GPUs, die auf der Volta-, Turing-, Ampere-, Ada- und Hopper-Architektur basieren, effizient sein.
| GPU | CUDA-Rechenfähigkeit | Mindestens erforderliches CUDA-Toolkit für CUTLASS-3 |
|---|---|---|
| NVIDIA V100 Tensor Core 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 Core 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 Core GPU | 9.0 | 11.8 |
Im Allgemeinen kann PTX-Code, der für eine Zielarchitektur generiert wurde, auf zukünftigen Architekturen ausgeführt werden (dh er ist vorwärtskompatibel). Mit CUDA 12.0 wurde jedoch das Konzept der „architekturbeschleunigten Funktionen“ eingeführt, deren PTX keine Vorwärtskompatibilitätsgarantien bietet. Mehrere Hopper-PTX-Befehle fallen in diese Kategorie architekturbeschleunigter Funktionen und erfordern daher eine sm_90a Zielarchitektur (beachten Sie das angehängte „a“). Weitere Einzelheiten zu dieser und anderen architekturbeschleunigten Anweisungen finden Sie in der CUDA-Dokumentation.
Die Informationen zur Zielarchitektur werden über das cmake-Flag CUTLASS_NVCC_ARCHS an CUTLASS weitergegeben. Um die Leistung des Hopper GH100 zu maximieren, müssen Benutzer CUTLASS mit 90a als Zielarchitektur erstellen. Wenn ein Benutzer versehentlich einen Kernel erstellt, der SM90a-Funktionen verwendet (z. B. Hopper Tensor Core-Anweisungen), und dabei das SM90-Ziel verwendet (beachten Sie das Fehlen von „a“), entweder mit CUDA Toolkit 12 oder 11.8, wird erwartet, dass der Kernel mit einer Laufzeit fehlschlägt Fehler.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
Einzelheiten dazu, welche Kernel welche Zielarchitekturen erfordern, finden Sie in der Funktionsdokumentation.
CUTLASS wird in den folgenden Dokumenten und der begleitenden Doxygen-Dokumentation beschrieben.
Kurzanleitung – CUTLASS erstellen und ausführen
Funktionalität – fasst die in CUTLASS verfügbaren Funktionen zusammen
Effizientes GEMM in CUDA – beschreibt, wie GEMM-Kernel effizient in CUDA implementiert werden können
CUTLASS 3.x Design – beschreibt das CUTLASS 3.x Design, seine Vorteile und wie CuTe es uns ermöglicht, viel besser zusammensetzbare Komponenten zu schreiben
GEMM API 3.x – beschreibt das CUTLASS 3.x GEMM-Modell und C++-Vorlagenkonzepte
GEMM API 2.x – beschreibt das CUTLASS 2.x GEMM-Modell und C++-Vorlagenkonzepte
Implizite GEMM-Faltung – beschreibt die 2D- und 3D-Faltung in CUTLASS
Code-Organisation – beschreibt die Organisation und Inhalte des CUTLASS-Projekts
Terminologie – beschreibt die im Code verwendeten Begriffe
Programmierrichtlinien – Richtlinien zum Schreiben effizienter moderner CUDA C++
Grundlegende Typen – beschreibt grundlegende C++-Klassen, die in CUTLASS zur Darstellung numerischer Mengen und Arrays verwendet werden
Layouts – beschreibt Layouts von Matrizen und Tensoren im Speicher
Tile Iterators – beschreibt C++-Konzepte für die Iteration über Kacheln von Matrizen im Speicher
CUTLASS Profiler – befehlszeilengesteuerte Profilerstellungsanwendung
CUTLASS-Dienstprogramme – zusätzliche Vorlagen zur Erleichterung einer schnellen Entwicklung
Abhängiger Kernel-Start – beschreibt eine neue Funktion in Hopper, die überlappende abhängige Kernel im selben Stream ermöglicht, und wie sie in CUTLASS verwendet wird.
Den Aufbau eines effizienten GEMM haben wir auch in unserem Vortrag auf der GPU Technology Conference 2018 beschrieben.
CUTLASS: Software-Primitive für die dichte lineare Algebra auf allen Ebenen und Skalen innerhalb von CUDA
Entwicklung von CUDA-Kerneln, um Tensorkerne auf NVIDIA A100 an ihre absolute Grenze zu bringen
Beschleunigte Faltung mit Tensorkernen in CUTLASS
Beschleunigung des Rückwärtsdatengradienten durch Erhöhung der Tensorkernauslastung in CUTLASS
CUTLASS: Python-API, Verbesserungen und NVIDIA Hopper
CUTLASS ist eine reine Header-Vorlagenbibliothek und muss nicht erstellt werden, um von anderen Projekten verwendet zu werden. Clientanwendungen sollten in ihren Include-Pfaden auf include/ -Verzeichnis von CUTLASS abzielen.
CUTLASS-Komponententests, Beispiele und Dienstprogramme können mit CMake erstellt werden. Die Mindestversion von CMake ist in der Schnellstartanleitung angegeben. Stellen Sie sicher, dass die Umgebungsvariable CUDACXX im auf Ihrem System installierten CUDA Toolkit auf NVCC verweist.
$ export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc Erstellen Sie ein Build-Verzeichnis im CUTLASS-Projekt und führen Sie dann CMake aus. Standardmäßig erstellt CUTLASS Kernel für die CUDA-Architekturversionen 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9 und 9.0. Um die Kompilierungszeit zu verkürzen, können Sie die Architekturen angeben, für die CUTLASS erstellt werden soll, indem Sie die CMake-Konfigurationseinstellung CUTLASS_NVCC_ARCHS ändern.
$ mkdir build && cd build $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # kompiliert für NVIDIAs Ampere-Architektur
Kompilieren Sie im Verzeichnis build/ die CUTLASS-Komponententests und führen Sie sie aus, indem Sie mit make die Zieleinheit test_unit erstellen.
Die Komponententests sind als mehrere Binärdateien organisiert, die die Top-Level-Namespaces von CUTLASS widerspiegeln, und können über das Befehlszeilenargument -j von make parallel ausgeführt werden.
$ make test_unit -j ... ... ... [----------] Abbau der globalen Testumgebung [==========] 946 Tests aus 57 Testfällen wurden ausgeführt. (insgesamt 10812 ms) [BESTANDEN] 946 Tests.
Alle Tests sollten auf unterstützten Plattformen bestanden werden, die genaue Anzahl der Tests kann jedoch im Laufe der Zeit variieren.
CUTLASS ist als reine Header-Bibliothek zusammen mit Dienstprogrammen, Tools, Beispielen und Komponententests angeordnet. Die Doxygen-Dokumentation bietet eine vollständige Liste der im CUTLASS-Projekt definierten Dateien, Klassen und Vorlagenkonzepte.
Eine ausführliche Erläuterung der Quellcode-Organisation finden Sie in der CUTLASS-Dokumentation, einige Hauptkomponenten werden jedoch nachstehend zusammengefasst.
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-Beispiele wenden CUTLASS-Vorlagen an, um grundlegende Berechnungen zu implementieren.
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.
Das Verzeichnis test/unit/ besteht aus mit Google Test implementierten Unit-Tests, die die grundlegende Verwendung von Core-API-Komponenten demonstrieren, sowie vollständige Tests der CUTLASS GEMM-Berechnungen.
Anweisungen zum Erstellen und Ausführen der Unit-Tests finden Sie in der Schnellstartanleitung.
Das Verzeichnis tools/profiler/ enthält ein Befehlszeilendienstprogramm zum Starten der einzelnen GEMM-Kernel. Es kann wie folgt aufgebaut werden:
$ make cutlass_profiler -j16
Standardmäßig wird für jeden Datentyp, jede Mathematikanweisung und jedes Layout nur eine Kachelgröße instanziiert. Um alles zu instanziieren, legen Sie die folgende Umgebungsvariable fest, wenn Sie CMake aus einem leeren build/ -Verzeichnis ausführen. Beachten Sie, dass dies zu Zehntausenden von Kerneln und langen Build-Zeiten führt. Dies würde auch zu einer großen Binärgröße führen und auf einigen Plattformen würde der Linker beim Erstellen der Bibliothek fehlschlagen. Daher wird dringend empfohlen, nur eine Teilmenge der Kernel zu generieren, wie im folgenden Unterabschnitt gezeigt.
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ make cutlass_profiler -j16
Um nur einen Kernel oder einen kleinen Kernelsatz zu kompilieren, kann eine durch Kommas getrennte Liste von Kernelnamen mit Platzhalterzeichen verwendet werden, um den Kernelsatz zu reduzieren. Die folgenden Beispiele zeigen die Erstellung genau eines Kernels oder einer Teilmenge von Kerneln für die NVIDIA Ampere- und Turing-Architektur:
Um eine Teilmenge der Tensor Core GEMM-Kernel mit FP32-Akkumulation und FP16-Eingabe für die NVIDIA Ampere- und Turing-Architektur zu kompilieren, verwenden Sie die folgende cmake-Befehlszeile:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ make cutlass_profiler -j16
Ein Beispiel für eine Befehlszeile zum Profilieren einer Teilmenge der Tensor Core GEMM-Kernel lautet wie folgt:
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
=============================
Problem-ID: 1
Anbieter: CUTLASS
Betriebsart: gemm
Operation: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Status: Erfolg
Überprüfung: EIN
Disposition: Bestanden
reference_device: Bestanden
cuBLAS: Bestanden
Argumente: --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
Laufzeit: 1,55948 ms
Speicher: 70,7616 GiB/s
Mathematik: 74378,8 GFLOP/s
=============================
...Um einen SGEMM-Kernel für die NVIDIA Ampere- und Turing-Architektur zu kompilieren, verwenden Sie die folgende cmake-Befehlszeile:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ make cutlass_profiler -j16
Ein Beispiel für eine Befehlszeile zum Profilieren eines einzelnen SGEMM-CUDA-Kernels lautet wie folgt:
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
=============================
Problem-ID: 1
Anbieter: CUTLASS
Betriebsart: gemm
Operation: cutlass_simt_sgemm_128x128_8x2_nn_align1
Status: Erfolg
Überprüfung: EIN
Disposition: Bestanden
cuBLAS: Bestanden
Argumente: --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
Laufzeit: 6,73655 ms
Speicher: 24,934 GiB/s
Mathematik: 17218,4 GFLOP/s
=============================Um eine Teilmenge der Tensor-Kern-Faltungskerne zu kompilieren, die Vorwärtsausbreitung (fprop) mit FP32-Akkumulation und FP16-Eingabe für die NVIDIA Ampere- und Turing-Architektur implementieren, verwenden Sie die folgende cmake-Befehlszeile:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $ make cutlass_profiler -j16
Ein Beispiel für eine Befehlszeile zum Profilieren einer Teilmenge von Tensor Core-Faltungskernen lautet wie folgt:
$ ./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
Anbieter: CUTLASS
Operationsart: conv2d
Operation: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Status: Erfolg
Überprüfung: EIN
Disposition: Bestanden
reference_device: Bestanden
Argumente: --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
Laufzeit: 0,711496 ms
Speicher: 1479,99 GiB/s
Mathematik: 166526 GFLOP/s
=============================
...Um einen CUDA Core-Faltungskern zu kompilieren und auszuführen, der Vorwärtspropagation (fprop) mit F32-Akkumulation und FP32-Eingabe für die NVIDIA Ampere- und Turing-Architektur implementiert, verwenden Sie die folgende cmake-Befehlszeile:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ make cutlass_profiler -j16
Beispiel einer Befehlszeile zum Profilieren eines CUDA Core-Faltungskerns:
$ ./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
Anbieter: CUTLASS
Operationsart: conv2d
Operation: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Status: Erfolg
Überprüfung: EIN
Disposition: Bestanden
reference_device: Bestanden
Argumente: --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
Laufzeit: 7,34266 ms
Speicher: 260,752 GiB/s
Mathematik: 16136,2 GFLOP/s
=============================Bitte folgen Sie den Links für weitere CMake-Beispiele zum selektiven Kompilieren von CUTLASS-Kerneln:
GEMM CMake-Beispiele
CMake-Beispiele für implizite GEMM-Faltung
Weitere Details zum CUTLASS Profiler finden Sie hier.
CUTLASS wird von der NVIDIA Corporation als Open-Source-Software unter der 3-Klausel „New“ BSD-Lizenz veröffentlicht.
Die offizielle Liste der CUTLASS-Entwickler und Mitwirkenden finden Sie hier: MITARBEITER.
Copyright (c) 2017–2024 NVIDIA CORPORATION & AFFILIATES. Alle Rechte vorbehalten. SPDX-Lizenzkennung: BSD-3-Klausel
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.
