cutlass
CUTLASS
CORTE 3.6.0 - Octubre 2024
CUTLASS es una colección de abstracciones de plantillas CUDA C++ para implementar la multiplicación matriz-matriz (GEMM) de alto rendimiento y cálculos relacionados en todos los niveles y escalas dentro de CUDA. Incorpora estrategias de descomposición jerárquica y movimiento de datos similares a las utilizadas para implementar cuBLAS y cuDNN. CUTLASS descompone estas "partes móviles" en componentes de software modulares reutilizables abstraídos por clases de plantilla de C++. Las primitivas para diferentes niveles de una jerarquía de paralelización conceptual se pueden especializar y ajustar mediante tamaños de mosaico personalizados, tipos de datos y otras políticas algorítmicas. La flexibilidad resultante simplifica su uso como bloques de construcción dentro de aplicaciones y núcleos personalizados.
Para admitir una amplia variedad de aplicaciones, CUTLASS proporciona un amplio soporte para cálculos de precisión mixta, proporcionando movimiento de datos especializado y abstracciones de acumulación múltiple para punto flotante de media precisión (FP16), BFloat16 (BF16), Tensor Float 32 (TF32), punto flotante de precisión simple (FP32), emulación FP32 mediante instrucción de núcleo tensor, tipos de punto flotante de doble precisión (FP64), tipos de datos enteros (4b y 8b), y tipos de datos binarios (1b). CUTLASS demuestra operaciones de multiplicación de matrices síncronas por deformación dirigidas a los Tensor Cores programables y de alto rendimiento implementados por las arquitecturas Volta, Turing, Ampere y Hopper de NVIDIA.
Consulte la Guía de inicio rápido para comenzar rápidamente.
Consulte la lista de funciones para conocer la lista de operaciones admitidas en cada nivel de la jerarquía del modelo de ejecución.
CUTLASS 3.0 introdujo una nueva biblioteca central, CuTe, para describir y manipular tensores de hilos y datos. CuTe es una colección de abstracciones de plantillas CUDA de C++ para definir y operar en diseños jerárquicamente multidimensionales de subprocesos y datos. CuTe proporciona objetos Layout y Tensor que empaquetan de forma compacta el tipo, la forma, el espacio de memoria y el diseño de los datos, mientras realiza la complicada indexación para el usuario. Esto permite a los programadores centrarse en las descripciones lógicas de sus algoritmos mientras CuTe realiza la contabilidad mecánica por ellos. Con estas herramientas, podemos diseñar, implementar y modificar rápidamente todas las operaciones de álgebra lineal densa.
Las abstracciones centrales de CuTe son diseños jerárquicamente multidimensionales que se pueden componer con matrices de datos para representar tensores. La representación de diseños es lo suficientemente potente como para representar casi todo lo que necesitamos para implementar un álgebra lineal densa eficiente. Los diseños también se pueden combinar y manipular mediante la composición funcional, sobre la cual construimos un gran conjunto de operaciones comunes, como mosaico y partición.
CUTLASS 3.0 y posteriores adoptan CuTe en toda la jerarquía GEMM en sus plantillas. Esto simplifica enormemente el diseño y mejora la componibilidad y legibilidad del código. Puede encontrar más documentación específica de CuTe en su directorio de documentación dedicado.
Además de GEMM, CUTLASS implementa convolución de alto rendimiento a través del algoritmo GEMM implícito. GEMM implícito es la formulación de una operación de convolución como GEMM, aprovechando así la canalización modular GEMM de CUTLASS. Esto permite a CUTLASS crear convoluciones reutilizando componentes GEMM altamente optimizados.
CUTLASS 3.6.0 es una actualización de CUTLASS que agrega:
Tolva estructurada GEMM dispersa.
FP16
8PM
INT8
TF32
Una refactorización de la API del kernel::ConvUniversal para alinearla con gemm::GemmUniversal . Ahora la API de convolución 3.x ya no se considera una API beta.
Un GEMM de entrada mixta mejorado y una implementación de tabla de búsqueda para el modo de solo escala INT4 x FP8 .
Nodos EVT para la selección Top-K y ejemplo de softmax y GEMM usándolos.
Lanzamiento dependiente programático (PDL) que aprovecha una nueva característica Hopper para acelerar dos kernels consecutivos y sus documentaciones correspondientes.
Una nueva herramienta de depuración, synclog, para descargar todos los eventos de sincronización desde un kernel a un archivo. Consulte la documentación de sincronización para obtener más detalles.
Un nuevo epílogo habilitado para TMA para GEMM agrupado que aporta una mejora significativa del rendimiento, así como su soporte EVT.
Un epílogo de matriz de punteros habilitado para SIMT.
Un nuevo calendario de kernel de Ping-Pong para GEMM agrupados y algunas otras optimizaciones.
Una nueva estrategia de creación de instancias para los núcleos del perfilador CUTLASS junto con documentación mejorada para el nivel de creación de instancias en el perfilador CUTLASS.
Un nuevo soporte de hardware para comparaciones y cálculos de cutlass::bfloat16_t
Se corrigió el uso de isnan en Windows para half_t .
Requisitos mínimos:
Arquitectura: Volta
Compilador: debe admitir al menos C++17
Versión del kit de herramientas CUDA: 11.4
A partir de CUTLASS 3.0, CUTLASS eliminó la compatibilidad con lo siguiente:
Arquitecturas de GPU Maxwell y Pascal
ubuntu 16.04
CUDA 10.2
Versiones del lenguaje C++ inferiores a 17.
Consulte el REGISTRO DE CAMBIOS para obtener una lista detallada de lanzamientos y actualizaciones.
Las primitivas CUTLASS son muy eficientes. Cuando se utilizan para construir núcleos GEMM para todo el dispositivo, exhiben un rendimiento máximo comparable al cuBLAS para cálculos GEMM escalares. La figura anterior muestra las continuas mejoras de rendimiento de CUTLASS en una NVIDIA H100 (arquitectura NVIDIA Hopper) desde CUTLASS 3.1. CUTLASS 3.5.1 se compiló con el kit de herramientas CUDA 12.5u1. Las operaciones de Tensor Core se implementan utilizando las instrucciones mma y wgmma de CUDA.
Cuando se utilizan bloques de construcción CUTLASS para construir núcleos gemm implícitos (Fprop, Dgrad y Wgrad) en todo el dispositivo, el rendimiento de CUTLASS también es comparable al de cuDNN cuando se ejecutan capas Resnet-50 en una NVIDIA A100, como se muestra en la figura anterior. Las operaciones de Tensor Core se implementan utilizando la instrucción mma de CUDA.
CUTLASS requiere un compilador host C++17 y funciona mejor cuando se compila con CUDA 12.4 Toolkit . También es compatible con 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 y CUDA 12.3.2.
Hemos probado los siguientes entornos.
| Sistema operativo | Compilador |
|---|---|
| ubuntu 18.04 | CCG 7.5.0 |
| ubuntu 20.04 | CCG 10.3.0 |
| ubuntu 22.04 | CCG 11.2.0 |
| ubuntu 22.04 | Sonido metálico 10.0.0 |
| ubuntu 22.04 | Sonido metálico 14.0.6 |
| ubuntu 22.04 | Sonido metálico 17.0.6 |
| Ventanas 10.0 | Estudio Visual 2019 v16.11.27 |
Nota: GCC 8.5.0 tiene regresiones conocidas con respecto a expresiones de plegado y operadores sobrecargados. Se recomienda utilizar GCC 7.5.0 o (preferido) GCC >= 9.
CUTLASS se ejecuta correctamente en las siguientes GPU NVIDIA y se espera que sea eficiente en las GPU NVIDIA basadas en arquitectura Volta, Turing, Ampere, Ada y Hopper.
| GPU | Capacidad de computación CUDA | Kit de herramientas CUDA mínimo requerido por 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 con núcleo tensor | 9.0 | 11.8 |
En general, el código PTX generado para una arquitectura de destino se puede ejecutar en arquitecturas futuras (es decir, es compatible con versiones posteriores). Sin embargo, CUDA 12.0 introdujo el concepto de "características aceleradas por arquitectura" cuyo PTX no tiene garantías de compatibilidad futura. Varias instrucciones Hopper PTX se incluyen en esta categoría de características de arquitectura acelerada y, por lo tanto, requieren una arquitectura de destino sm_90a (tenga en cuenta la "a" adjunta). Para obtener más detalles sobre esta y otras instrucciones aceleradas por arquitectura, consulte la documentación de CUDA.
La información de la arquitectura de destino se pasa a CUTLASS a través del indicador cmake CUTLASS_NVCC_ARCHS . Para maximizar el rendimiento en Hopper GH100, los usuarios deben crear CUTLASS con 90a como arquitectura de destino. Si un usuario construye accidentalmente un kernel que usa características SM90a (por ejemplo, instrucciones Hopper Tensor Core), usando el destino SM90 (tenga en cuenta la falta de "a"), ya sea con CUDA Toolkit 12 o 11.8, se espera que el kernel falle con un tiempo de ejecución. error.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
Consulte la documentación de funcionalidad para obtener detalles sobre qué núcleos requieren qué arquitecturas de destino.
CUTLASS se describe en los siguientes documentos y en la documentación adjunta de Doxygen.
Guía de inicio rápido: compilar y ejecutar CUTLASS
Funcionalidad: resume la funcionalidad disponible en CUTLASS
GEMM eficiente en CUDA: describe cómo los núcleos GEMM se pueden implementar de manera eficiente en CUDA
Diseño CUTLASS 3.x: describe el diseño CUTLASS 3.x, sus beneficios y cómo CuTe nos permite escribir componentes mucho más componibles.
GEMM API 3.x: describe el modelo GEMM de CUTLASS 3.x y los conceptos de plantilla de C++
GEMM API 2.x: describe el modelo CUTLASS 2.x GEMM y los conceptos de plantilla C++
Convolución GEMM implícita: describe la convolución 2D y 3D en CUTLASS
Organización del código: describe la organización y el contenido del proyecto CUTLASS.
Terminología: describe los términos utilizados en el código.
Pautas de programación: pautas para escribir CUDA C++ moderno y eficiente
Tipos fundamentales: describe las clases básicas de C++ utilizadas en CUTLASS para representar cantidades numéricas y matrices.
Diseños: describe diseños de matrices y tensores en la memoria.
Iteradores de mosaicos: describe conceptos de C++ para iterar sobre mosaicos de matrices en la memoria.
CUTLASS Profiler: aplicación de creación de perfiles basada en línea de comandos
Utilidades CUTLASS: plantillas adicionales utilizadas para facilitar el desarrollo rápido
Lanzamiento del kernel dependiente: describe una nueva característica en Hopper que permite superponer kernels dependientes en la misma secuencia y cómo se usa en CUTLASS.
También describimos la estructura de un GEMM eficiente en nuestra charla en la GPU Technology Conference 2018.
CUTLASS: Primitivas de software para álgebra lineal densa en todos los niveles y escalas dentro de CUDA
Desarrollo de kernels CUDA para llevar los núcleos tensoriales al límite absoluto en NVIDIA A100
Acelerando la convolución con núcleos tensoriales en CUTLASS
Acelerar el gradiente de datos hacia atrás aumentando la utilización del núcleo tensorial en CUTLASS
CUTLASS: API de Python, mejoras y NVIDIA Hopper
CUTLASS es una biblioteca de plantillas de solo encabezado y no es necesario crearla para que la utilicen otros proyectos. Las aplicaciones cliente deben apuntar al directorio include/ de CUTLASS en sus rutas de inclusión.
Las pruebas unitarias, ejemplos y utilidades de CUTLASS se pueden crear con CMake. La versión mínima de CMake se proporciona en la guía de inicio rápido. Asegúrese de que la variable de entorno CUDACXX apunte a NVCC en el kit de herramientas CUDA instalado en su sistema.
$ exportar CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc Cree un directorio de compilación dentro del proyecto CUTLASS y luego ejecute CMake. De forma predeterminada, CUTLASS creará kernels para las versiones de arquitectura CUDA 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9 y 9.0. Para reducir el tiempo de compilación, puede especificar las arquitecturas para las que construir CUTLASS cambiando la configuración de CMake CUTLASS_NVCC_ARCHS .
$ mkdir compilación && compilación de cd $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # compila para la arquitectura Ampere de NVIDIA
Desde el directorio build/ , compila y ejecuta las pruebas unitarias CUTLASS construyendo la test_unit de destino con make.
Las pruebas unitarias están organizadas como varios archivos binarios que reflejan los espacios de nombres de nivel superior de CUTLASS, y pueden ejecutarse en paralelo mediante el argumento de línea de comando -j de make.
$ hacer unidad_prueba -j ... ... ... [----------] Desmontaje del entorno de prueba global [==========] Se ejecutaron 946 pruebas de 57 casos de prueba. (10812 ms en total) [APROBADO] 946 pruebas.
Todas las pruebas deben aprobarse en plataformas compatibles, aunque el número exacto de pruebas puede variar con el tiempo.
CUTLASS está organizado como una biblioteca de solo encabezado junto con utilidades, herramientas, ejemplos y pruebas unitarias. La documentación de Doxygen proporciona una lista completa de archivos, clases y conceptos de plantilla definidos en el proyecto CUTLASS.
Puede encontrar una explicación detallada de la organización del código fuente en la documentación de CUTLASS, pero a continuación se resumen varios componentes principales.
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
Los ejemplos de CUTLASS SDK aplican plantillas CUTLASS para implementar cálculos básicos.
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.
El directorio test/unit/ consta de pruebas unitarias implementadas con Google Test que demuestran el uso básico de los componentes Core API y pruebas completas de los cálculos de CUTLASS GEMM.
Las instrucciones para crear y ejecutar las pruebas unitarias se describen en la guía de inicio rápido.
El directorio tools/profiler/ contiene una utilidad de línea de comandos para iniciar cada uno de los núcleos GEMM. Se puede construir de la siguiente manera:
$ hacer cutlass_profiler -j16
De forma predeterminada, solo se crea una instancia de un tamaño de mosaico para cada tipo de datos, instrucción matemática y diseño. Para crear una instancia de todo, configure la siguiente variable de entorno cuando ejecute CMake desde un directorio build/ vacío. Tenga cuidado, esto da como resultado decenas de miles de núcleos y tiempos de compilación prolongados. Esto también daría como resultado un tamaño binario grande y, en algunas plataformas, el enlazador fallaría al construir la biblioteca. Por lo tanto, se recomienda generar solo un subconjunto de núcleos como se demuestra en la subsección siguiente.
$ cmake.. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=todos ... $ hacer cutlass_profiler -j16
Para compilar estrictamente un núcleo o un pequeño conjunto de núcleos, se puede utilizar una lista delimitada por comas de nombres de núcleos con caracteres comodín para reducir el conjunto de núcleos. Los siguientes ejemplos muestran la construcción de exactamente uno o un subconjunto de núcleos para la arquitectura NVIDIA Ampere y Turing:
Para compilar un subconjunto de kernels Tensor Core GEMM con acumulación FP32 y entrada FP16 dirigidas a la arquitectura NVIDIA Ampere y Turing, use la siguiente línea de comando cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ hacer cutlass_profiler -j16
La línea de comando de ejemplo para crear perfiles de un subconjunto de núcleos GEMM de Tensor Core es la siguiente:
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
==============================
ID del problema: 1
Proveedor: CORTE
Tipo de operación: gemm
Operación: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Estado: Éxito
Verificación: ENCENDIDO
Disposición: Aprobada
dispositivo_referencia: Aprobado
cuBLAS: Aprobado
Argumentos: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:columna --B=f16:fila --C=f32:columna --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 --etapas=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
FLOP: 115992428544 fracasos
Tiempo de ejecución: 1,55948 ms
Memoria: 70,7616 GiB/s
Matemáticas: 74378,8 GFLOP/s
==============================
...Para compilar un kernel SGEMM dirigido a la arquitectura NVIDIA Ampere y Turing, utilice la siguiente línea de comando cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ hacer cutlass_profiler -j16
La línea de comando de ejemplo para crear perfiles de un solo kernel SGEMM CUDA es la siguiente:
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
==============================
ID del problema: 1
Proveedor: CORTE
Tipo de operación: gemm
Operación: cutlass_simt_sgemm_128x128_8x2_nn_align1
Estado: Éxito
Verificación: ENCENDIDO
Disposición: Aprobada
cuBLAS: Aprobado
Argumentos: --m=3456 --n=4096 --k=4096 --A=f32:columna --B=f32:columna --C=f32:columna --alpha=1 --beta=0 -- split_k_rebanadas=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
FLOP: 115992428544 fracasos
Tiempo de ejecución: 6,73655 ms
Memoria: 24.934 GiB/s
Matemáticas: 17218,4 GFLOP/s
==============================Para compilar un subconjunto de núcleos de convolución centrales de Tensor que implementan propagación hacia adelante (fprop) con acumulación de FP32 y entrada de FP16 dirigida a la arquitectura NVIDIA Ampere y Turing, use la siguiente línea de comando cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $ hacer cutlass_profiler -j16
La línea de comando de ejemplo para crear perfiles de un subconjunto de núcleos de convolución de Tensor Core es la siguiente:
$ ./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 del problema: 1
Proveedor: CORTE
Tipo de operación: conv2d
Operación: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Estado: Éxito
Verificación: ENCENDIDO
Disposición: Aprobada
dispositivo_referencia: Aprobado
Argumentos: --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 --Activación=f16:nhwc --Filter=f16:nhwc --Salida=f32:nhwc
--conv_mode=cross --iterator_algorithm=optimizado --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
FLOP: 118482796544 fracasos
Tiempo de ejecución: 0,711496 ms
Memoria: 1479,99 GiB/s
Matemáticas: 166526 GFLOP/s
==============================
...Para compilar y ejecutar un kernel de convolución CUDA Core que implemente la propagación hacia adelante (fprop) con acumulación F32 y entrada FP32 dirigida a la arquitectura NVIDIA Ampere y Turing, use la siguiente línea de comando cmake:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ hacer cutlass_profiler -j16
Línea de comando de ejemplo para crear perfiles de un núcleo de convolución 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 del problema: 1
Proveedor: CORTE
Tipo de operación: conv2d
Operación: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Estado: Éxito
Verificación: ENCENDIDO
Disposición: Aprobada
dispositivo_referencia: Aprobado
Argumentos: --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 --Activación=f32:nhwc --Filter=f32:nhwc --Salida=f32:nhwc
--conv_mode=cross --iterator_algorithm=optimizado --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
FLOP: 118482796544 fracasos
Tiempo de ejecución: 7,34266 ms
Memoria: 260.752 GiB/s
Matemáticas: 16136,2 GFLOP/s
==============================Siga los enlaces para obtener más ejemplos de CMake sobre cómo compilar selectivamente kernels CUTLASS:
Ejemplos de GEMM CMake
Ejemplos de CMake de convolución GEMM implícita
Aquí se describen más detalles sobre CUTLASS Profiler.
CUTLASS es lanzado por NVIDIA Corporation como software de código abierto bajo la licencia BSD "Nueva" de 3 cláusulas.
La lista oficial de desarrolladores y contribuyentes de CUTLASS está disponible aquí: COLABORADORES.
Copyright (c) 2017 - 2024 NVIDIA CORPORATION & AFILIADOS. Reservados todos los derechos. Identificador de licencia SPDX: Cláusula 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.
