cutlass
CUTLASS
COUTLASS 3.6.0 - Octobre 2024
CUTLASS est une collection d'abstractions de modèles CUDA C++ pour la mise en œuvre de la multiplication matrice-matrice (GEMM) hautes performances et des calculs associés à tous les niveaux et échelles au sein de CUDA. Il intègre des stratégies de décomposition hiérarchique et de mouvement des données similaires à celles utilisées pour implémenter cuBLAS et cuDNN. CUTLASS décompose ces « pièces mobiles » en composants logiciels réutilisables et modulaires abstraits par des classes de modèles C++. Les primitives pour différents niveaux d'une hiérarchie de parallélisation conceptuelle peuvent être spécialisées et ajustées via des tailles de mosaïque personnalisées, des types de données et d'autres politiques algorithmiques. La flexibilité qui en résulte simplifie leur utilisation en tant qu’éléments de base au sein de noyaux et d’applications personnalisés.
Pour prendre en charge une grande variété d'applications, CUTLASS fournit une prise en charge étendue des calculs de précision mixte, fournissant des abstractions spécialisées de déplacement de données et de multiplication-accumulation pour la virgule flottante demi-précision (FP16), BFloat16 (BF16), Tensor Float 32 (TF32), virgule flottante simple précision (FP32), émulation FP32 via une instruction tensor core, types à virgule flottante double précision (FP64), types de données entiers (4b et 8b) et les types de données binaires (1b). CUTLASS démontre des opérations de multiplication matricielle synchrones ciblant les cœurs Tensor programmables à haut débit implémentés par les architectures Volta, Turing, Ampere et Hopper de NVIDIA.
Consultez le Guide de démarrage rapide pour démarrer rapidement.
Consultez la liste des fonctionnalités pour connaître la liste des opérations prises en charge à chaque niveau de la hiérarchie du modèle d'exécution.
CUTLASS 3.0 a introduit une nouvelle bibliothèque principale, CuTe, pour décrire et manipuler les tenseurs de threads et de données. CuTe est une collection d'abstractions de modèles C++ CUDA permettant de définir et d'opérer sur des dispositions hiérarchiquement multidimensionnelles de threads et de données. CuTe fournit des objets Layout et Tensor qui regroupent de manière compacte le type, la forme, l'espace mémoire et la disposition des données, tout en effectuant une indexation compliquée pour l'utilisateur. Cela permet aux programmeurs de se concentrer sur les descriptions logiques de leurs algorithmes pendant que CuTe s'occupe de la comptabilité mécanique à leur place. Avec ces outils, nous pouvons rapidement concevoir, implémenter et modifier toutes les opérations d’algèbre linéaire dense.
Les abstractions principales de CuTe sont des dispositions hiérarchiquement multidimensionnelles qui peuvent être composées de tableaux de données pour représenter les tenseurs. La représentation des mises en page est suffisamment puissante pour représenter presque tout ce dont nous avons besoin pour mettre en œuvre une algèbre linéaire dense efficace. Les mises en page peuvent également être combinées et manipulées via une composition fonctionnelle, sur laquelle nous construisons un large ensemble d'opérations courantes telles que le carrelage et le partitionnement.
CUTLASS 3.0 et au-delà adopte CuTe dans toute la hiérarchie GEMM dans ses modèles. Cela simplifie grandement la conception et améliore la composabilité et la lisibilité du code. Plus de documentation spécifique à CuTe peut être trouvée dans son répertoire de documentation dédié.
En plus des GEMM, CUTLASS implémente une convolution haute performance via l'algorithme implicite GEMM. Le GEMM implicite est la formulation d'une opération de convolution en tant que GEMM tirant ainsi parti du pipeline GEMM modulaire de CUTLASS. Cela permet à CUTLASS de créer des convolutions en réutilisant des composants GEMM hautement optimisés.
CUTLASS 3.6.0 est une mise à jour de CUTLASS ajoutant :
Hopper structuré GEMM clairsemé.
PC16
PC8
INT8
TF32
Une refactorisation de l'API de convolution kernel::ConvUniversal de CUTLASS 3.x pour l'aligner sur gemm::GemmUniversal . Désormais, l'API de convolution 3.x n'est plus considérée comme une API bêta.
Un GEMM à entrée mixte amélioré et une implémentation de table de recherche pour le mode échelle uniquement INT4 x FP8 .
Nœuds EVT pour la sélection Top-K et exemple softmax et GEMM les utilisant.
Programmatic Dependent Launch (PDL) qui exploite une nouvelle fonctionnalité Hopper pour accélérer deux noyaux consécutifs, ainsi que ses documentations correspondantes.
Un nouvel outil de débogage, synclog, pour transférer tous les événements de synchronisation d'un noyau vers un fichier. Veuillez consulter la documentation du synclog pour plus de détails.
Un nouvel épilogue compatible TMA pour les GEMM groupés qui apporte une amélioration significative des performances, ainsi que sa prise en charge EVT.
Un épilogue de tableau de pointeurs compatible SIMT.
Un nouveau planning du noyau Ping-Pong pour Grouped GEMM et quelques autres optimisations.
Une nouvelle stratégie d'instanciation pour les noyaux du profileur CUTLASS ainsi qu'une documentation améliorée pour le niveau d'instanciation dans le profileur CUTLASS.
Un nouveau support matériel pour les comparaisons et les calculs de cutlass::bfloat16_t
Correction de l'utilisation d'isnan sous Windows pour half_t .
Exigences minimales :
Architecture : Volta
Compilateur : doit prendre en charge au moins C++17
Version de la boîte à outils CUDA : 11.4
À partir de CUTLASS 3.0, CUTLASS a supprimé la prise en charge des éléments suivants :
Architectures GPU Maxwell et Pascal
Ubuntu 16.04
CUDA 10.2
Versions du langage C++ inférieures à 17.
Consultez le CHANGELOG pour une liste détaillée des versions et des mises à jour.
Les primitives CUTLASS sont très efficaces. Lorsqu'ils sont utilisés pour construire des noyaux GEMM à l'échelle de l'appareil, ils présentent des performances maximales comparables à celles de cuBLAS pour les calculs GEMM scalaires. La figure ci-dessus montre les améliorations continues des performances de CUTLASS sur un NVIDIA H100 (architecture NVIDIA Hopper) depuis CUTLASS 3.1. CUTLASS 3.5.1 a été compilé avec la boîte à outils CUDA 12.5u1. Les opérations Tensor Core sont implémentées à l'aide des instructions mma et wgmma de CUDA.
Lors de l'utilisation des blocs de construction CUTLASS pour construire des noyaux gemm implicites (Fprop, Dgrad et Wgrad) à l'échelle de l'appareil, les performances de CUTLASS sont également comparables à cuDNN lors de l'exécution de couches Resnet-50 sur un NVIDIA A100, comme le montre la figure ci-dessus. Les opérations Tensor Core sont implémentées à l'aide de l'instruction mma de CUDA.
CUTLASS nécessite un compilateur hôte C++17 et fonctionne mieux lorsqu'il est construit avec le Toolkit CUDA 12.4 . Il est également compatible avec 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 et CUDA 12.3.2.
Nous avons testé les environnements suivants.
| Système opérateur | Compilateur |
|---|---|
| Ubuntu 18.04 | CCG 7.5.0 |
| Ubuntu 20.04 | CCG 10.3.0 |
| Ubuntu 22.04 | CCG 11.2.0 |
| Ubuntu 22.04 | Clang 10.0.0 |
| Ubuntu 22.04 | Clang 14.0.6 |
| Ubuntu 22.04 | Clang 17.0.6 |
| Windows 10.0 | Visual Studio 2019 v16.11.27 |
Remarque : GCC 8.5.0 a connu des régressions concernant les expressions de repli et les opérateurs surchargés. L'utilisation de GCC 7.5.0 ou (de préférence) GCC >= 9 est recommandée.
CUTLASS fonctionne avec succès sur les GPU NVIDIA suivants, et il devrait être efficace sur les GPU NVIDIA basés sur l'architecture Volta, Turing, Ampere, Ada et Hopper.
| GPU | Capacité de calcul CUDA | Boîte à outils CUDA minimale requise par CUTLASS-3 |
|---|---|---|
| GPU NVIDIA V100 Tensor Core | 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 |
| GPU NVIDIA A100 Tensor Core | 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 |
| GPU NVIDIA H100 Tensor Core | 9.0 | 11.8 |
En général, le code PTX généré pour une architecture cible peut être exécuté sur des architectures futures (c'est-à-dire qu'il est compatible avec les versions ultérieures). Cependant, CUDA 12.0 a introduit le concept de « fonctionnalités accélérées par l'architecture » dont le PTX n'a pas de garantie de compatibilité ascendante. Plusieurs instructions Hopper PTX relèvent de cette catégorie de fonctionnalités accélérées par l'architecture et nécessitent donc une architecture cible sm_90a (notez le « a » ajouté). Pour plus de détails à ce sujet et sur d'autres instructions accélérées par l'architecture, veuillez vous référer à la documentation CUDA.
Les informations sur l'architecture cible sont transmises à CUTLASS via l'indicateur cmake CUTLASS_NVCC_ARCHS . Afin de maximiser les performances du Hopper GH100, les utilisateurs doivent construire CUTLASS avec 90a comme architecture cible. Si un utilisateur crée accidentellement un noyau qui utilise les fonctionnalités SM90a (par exemple Hopper Tensor Core Instructions), en utilisant la cible SM90 (notez l'absence de "a"), avec CUDA Toolkit 12 ou 11.8, le noyau devrait échouer avec un runtime erreur.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
Veuillez vous référer à la documentation des fonctionnalités pour plus de détails sur les noyaux qui nécessitent quelles architectures cibles.
CUTLASS est décrit dans les documents suivants et dans la documentation Doxygen qui l'accompagne.
Guide de démarrage rapide - créer et exécuter CUTLASS
Fonctionnalité - résume les fonctionnalités disponibles dans CUTLASS
GEMM efficace dans CUDA - décrit comment les noyaux GEMM peuvent être implémentés efficacement dans CUDA
Conception CUTLASS 3.x - décrit la conception CUTLASS 3.x, ses avantages et comment CuTe nous permet d'écrire des composants beaucoup plus composables
GEMM API 3.x - décrit le modèle GEMM CUTLASS 3.x et les concepts de modèle C++
GEMM API 2.x - décrit le modèle GEMM CUTLASS 2.x et les concepts de modèle C++
Convolution GEMM implicite - décrit la convolution 2D et 3D dans CUTLASS
Code Organisation - décrit l'organisation et le contenu du projet CUTLASS
Terminologie - décrit les termes utilisés dans le code
Directives de programmation - directives pour écrire du CUDA C++ moderne et efficace
Types fondamentaux - décrit les classes C++ de base utilisées dans CUTLASS pour représenter des quantités numériques et des tableaux
Dispositions - décrit les dispositions des matrices et des tenseurs en mémoire
Tile Iterators - décrit les concepts C++ pour itérer sur les tuiles de matrices en mémoire
CUTLASS Profiler - application de profilage pilotée par ligne de commande
Utilitaires CUTLASS - modèles supplémentaires utilisés pour faciliter un développement rapide
Lancement du noyau dépendant - décrit une nouvelle fonctionnalité dans Hopper qui permet de superposer des noyaux dépendants dans le même flux, et comment elle est utilisée dans CUTLASS.
Nous avons également décrit la structure d'un GEMM efficace dans notre intervention à la GPU Technology Conference 2018.
CUTLASS : primitives logicielles pour l'algèbre linéaire dense à tous les niveaux et à toutes les échelles au sein de CUDA
Développement de noyaux CUDA pour pousser les cœurs Tensor à la limite absolue sur NVIDIA A100
Accélération de la convolution avec les cœurs tenseurs dans CUTLASS
Accélération du gradient de données vers l'arrière en augmentant l'utilisation du noyau Tensor dans CUTLASS
CUTLASS : API Python, améliorations et NVIDIA Hopper
CUTLASS est une bibliothèque de modèles d'en-tête uniquement et n'a pas besoin d'être construite pour être utilisée par d'autres projets. Les applications client doivent cibler le répertoire include/ de CUTLASS dans leurs chemins d'inclusion.
Les tests unitaires, les exemples et les utilitaires CUTLASS peuvent être créés avec CMake. La version minimale de CMake est indiquée dans le guide de démarrage rapide. Assurez-vous que la variable d'environnement CUDACXX pointe vers NVCC dans la boîte à outils CUDA installée sur votre système.
$ export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc Créez un répertoire de construction dans le projet CUTLASS, puis exécutez CMake. Par défaut, CUTLASS construira des noyaux pour les versions d'architecture CUDA 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9 et 9.0. Pour réduire le temps de compilation, vous pouvez spécifier les architectures pour lesquelles construire CUTLASS en modifiant le paramètre de configuration CMake CUTLASS_NVCC_ARCHS .
$ mkdir construire && cd construire $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # compile pour l'architecture Ampere de NVIDIA
Depuis le répertoire build/ , compilez et exécutez les tests unitaires CUTLASS en construisant la cible test_unit avec make.
Les tests unitaires sont organisés sous forme de plusieurs binaires reflétant les espaces de noms de niveau supérieur de CUTLASS, et ils peuvent être exécutés en parallèle via l'argument de ligne de commande -j de make.
$ make test_unit -j ... ... ... [---------] Démantèlement de l'environnement de test global [==========] 946 tests issus de 57 scénarios de test ont été exécutés. (10 812 ms au total) [ RÉUSSI ] 946 tests.
Tous les tests doivent réussir sur les plates-formes prises en charge, bien que le nombre exact de tests puisse varier au fil du temps.
CUTLASS est organisé comme une bibliothèque d'en-tête uniquement avec des utilitaires, des outils, des exemples et des tests unitaires. La documentation Doxygen fournit une liste complète des fichiers, classes et concepts de modèles définis dans le projet CUTLASS.
Une explication détaillée de l'organisation du code source peut être trouvée dans la documentation de CUTLASS, mais plusieurs composants principaux sont résumés ci-dessous.
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
Les exemples du SDK CUTLASS appliquent les modèles CUTLASS pour implémenter les calculs de base.
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.
Le répertoire test/unit/ se compose de tests unitaires implémentés avec Google Test qui démontrent l'utilisation de base des composants Core API et des tests complets des calculs CUTLASS GEMM.
Les instructions pour créer et exécuter les tests unitaires sont décrites dans le guide de démarrage rapide.
Le répertoire tools/profiler/ contient un utilitaire de ligne de commande pour lancer chacun des noyaux GEMM. Il peut être construit comme suit :
$ make cutlass_profiler -j16
Par défaut, une seule taille de tuile est instanciée pour chaque type de données, instruction mathématique et disposition. Pour tout instancier, définissez la variable d'environnement suivante lors de l'exécution de CMake à partir d'un répertoire build/ vide. Attention, cela entraîne des dizaines de milliers de noyaux et des temps de construction longs. Cela entraînerait également une grande taille binaire et, sur certaines plates-formes, l'éditeur de liens échouerait lors de la construction de la bibliothèque. Par conséquent, il est fortement recommandé de générer uniquement un sous-ensemble de noyaux, comme démontré dans la sous-section ci-dessous.
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ make cutlass_profiler -j16
Pour compiler strictement un noyau ou un petit ensemble de noyaux, une liste de noms de noyau délimités par des virgules avec des caractères génériques peut être utilisée pour réduire l'ensemble des noyaux. Les exemples suivants montrent la création d'un ou d'un sous-ensemble de noyaux pour l'architecture NVIDIA Ampere et Turing :
Pour compiler un sous-ensemble de noyaux Tensor Core GEMM avec accumulation FP32 et entrée FP16 ciblant l'architecture NVIDIA Ampere et Turing, utilisez la ligne de commande cmake ci-dessous :
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ make cutlass_profiler -j16
Un exemple de ligne de commande pour profiler un sous-ensemble de noyaux Tensor Core GEMM est le suivant :
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
=============================
ID du problème : 1
Fournisseur : COUTLASS
Type d'opération : gemm
Opération : cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Statut : Succès
Vérification : ON
Disposition : Réussi
reference_device : Réussi
cuBLAS : Réussi
Arguments : --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:colonne --B=f16:ligne --C=f32:colonne --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
Octets : 1 18489088 octets
FLOP : 115992428544 flops
Durée d'exécution : 1,55948 ms
Mémoire : 70,7616 Gio/s
Mathématiques : 74378,8 GFLOP/s
=============================
...Pour compiler un noyau SGEMM ciblant l'architecture NVIDIA Ampere et Turing, utilisez la ligne de commande cmake ci-dessous :
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ make cutlass_profiler -j16
Un exemple de ligne de commande pour profiler un seul noyau SGEMM CUDA est le suivant :
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
=============================
ID du problème : 1
Fournisseur : COUTLASS
Type d'opération : gemm
Opération : cutlass_simt_sgemm_128x128_8x2_nn_align1
Statut : Succès
Vérification : ON
Disposition : Réussi
cuBLAS : Réussi
Arguments : --m=3456 --n=4096 --k=4096 --A=f32:colonne --B=f32:colonne --C=f32:colonne --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
Octets : 180355072 octets
FLOP : 115992428544 flops
Durée d'exécution : 6,73655 ms
Mémoire : 24,934 Gio/s
Mathématiques : 17218,4 GFLOP/s
=============================Pour compiler un sous-ensemble de noyaux de convolution Tensor implémentant la propagation directe (fprop) avec accumulation FP32 et entrée FP16 ciblant l'architecture NVIDIA Ampere et Turing, utilisez la ligne de commande cmake ci-dessous :
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $ make cutlass_profiler -j16
Un exemple de ligne de commande pour profiler un sous-ensemble de noyaux de convolution Tensor Core est le suivant :
$ ./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 du problème : 1
Fournisseur : COUTLASS
Type d'opération : conv2d
Opération : cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Statut : Succès
Vérification : ON
Disposition : Réussi
reference_device : Réussi
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=optimisé --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
Octets : 1130659840 octets
FLOP : 118482796544 flops
Durée d'exécution : 0,711496 ms
Mémoire : 1 479,99 Gio/s
Mathématique : 166 526 GFLOP/s
=============================
...Pour compiler et exécuter un noyau de convolution CUDA Core implémentant la propagation directe (fprop) avec accumulation F32 et entrée FP32 ciblant l'architecture NVIDIA Ampere et Turing, utilisez la ligne de commande cmake ci-dessous :
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ make cutlass_profiler -j16
Exemple de ligne de commande pour profiler un noyau de convolution 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 du problème : 1
Fournisseur : COUTLASS
Type d'opération : conv2d
Opération : cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Statut : Succès
Vérification : ON
Disposition : Réussi
reference_device : Réussi
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=optimisé --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1
--eq_gemm_provider=aucun --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
Octets : 2055798784 octets
FLOP : 118482796544 flops
Durée d'exécution : 7,34266 ms
Mémoire : 260,752 Gio/s
Mathématiques : 16 136,2 GFLOP/s
=============================Veuillez suivre les liens pour plus d'exemples CMake sur la compilation sélective des noyaux CUTLASS :
Exemples GEMM CMake
Exemples CMake de convolution GEMM implicite
De plus amples détails sur le profileur CUTLASS sont décrits ici.
CUTLASS est publié par NVIDIA Corporation en tant que logiciel Open Source sous la licence BSD « nouvelle » à 3 clauses.
La liste officielle des développeurs et contributeurs de CUTLASS est disponible ici : CONTRIBUTEURS.
Copyright (c) 2017 - 2024 NVIDIA CORPORATION ET AFFILIÉS. Tous droits réservés. Identifiant de licence SPDX : clause BSD-3
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.
