cutlass
CUTLASS
CUTLASS 3.6.0 - Oktober 2024
CUTLASS adalah kumpulan abstraksi template CUDA C++ untuk mengimplementasikan perkalian matriks-matriks kinerja tinggi (GEMM) dan komputasi terkait di semua level dan skala dalam CUDA. Ini menggabungkan strategi untuk dekomposisi hierarki dan pergerakan data yang serupa dengan yang digunakan untuk mengimplementasikan cuBLAS dan cuDNN. CUTLASS menguraikan "bagian bergerak" ini menjadi komponen perangkat lunak modular yang dapat digunakan kembali yang diabstraksi oleh kelas templat C++. Primitif untuk berbagai tingkat hierarki paralelisasi konseptual dapat dispesialisasikan dan disesuaikan melalui ukuran petak khusus, tipe data, dan kebijakan algoritmik lainnya. Fleksibilitas yang dihasilkan menyederhanakan penggunaannya sebagai blok penyusun dalam kernel dan aplikasi khusus.
Untuk mendukung beragam aplikasi, CUTLASS memberikan dukungan ekstensif untuk komputasi presisi campuran, menyediakan pergerakan data khusus dan abstraksi akumulasi ganda untuk floating point setengah presisi (FP16), BFloat16 (BF16), Tensor Float 32 (TF32), floating point presisi tunggal (FP32), emulasi FP32 melalui instruksi inti tensor, tipe floating point presisi ganda (FP64), tipe data integer (4b dan 8b), dan tipe data biner (1b). CUTLASS mendemonstrasikan operasi penggandaan matriks warp-synchronous yang menargetkan Tensor Cores dengan throughput tinggi yang dapat diprogram dan diterapkan oleh arsitektur Volta, Turing, Ampere, dan Hopper NVIDIA.
Lihat Panduan Memulai Cepat untuk memulai dengan cepat.
Lihat daftar fungsionalitas untuk daftar operasi yang didukung di setiap tingkat hierarki model eksekusi.
CUTLASS 3.0 memperkenalkan pustaka inti baru, CuTe, untuk mendeskripsikan dan memanipulasi tensor thread dan data. CuTe adalah kumpulan abstraksi template C++ CUDA untuk mendefinisikan dan mengoperasikan tata letak thread dan data yang multidimensi secara hierarki. CuTe menyediakan objek Layout dan Tensor yang mengemas tipe, bentuk, ruang memori, dan tata letak data secara ringkas, sekaligus melakukan pengindeksan yang rumit bagi pengguna. Hal ini memungkinkan pemrogram fokus pada deskripsi logis dari algoritme mereka sementara CuTe melakukan pembukuan mekanis untuk algoritme tersebut. Dengan alat ini, kita dapat dengan cepat merancang, mengimplementasikan, dan memodifikasi semua operasi aljabar linier padat.
Abstraksi inti CuTe adalah tata letak multidimensi hierarki yang dapat disusun dengan larik data untuk mewakili tensor. Representasi tata letak cukup kuat untuk mewakili hampir semua yang kita perlukan untuk mengimplementasikan aljabar linier padat yang efisien. Tata letak juga dapat digabungkan dan dimanipulasi melalui komposisi fungsional, yang menjadi dasar pembuatan serangkaian besar operasi umum seperti ubin dan partisi.
CUTLASS 3.0 dan seterusnya mengadopsi CuTe di seluruh hierarki GEMM dalam templatnya. Ini sangat menyederhanakan desain dan meningkatkan komposisi dan keterbacaan kode. Dokumentasi lebih spesifik untuk CuTe dapat ditemukan di direktori dokumentasi khusus.
Selain GEMM, CUTLASS mengimplementasikan konvolusi kinerja tinggi melalui algoritma GEMM implisit. GEMM implisit adalah formulasi operasi konvolusi sebagai GEMM sehingga memanfaatkan pipeline GEMM modular CUTLASS. Hal ini memungkinkan CUTLASS membangun konvolusi dengan menggunakan kembali komponen GEMM yang sangat optimal.
CUTLASS 3.6.0 adalah pembaruan untuk CUTLASS yang menambahkan:
GEMM jarang berstruktur hopper.
FP16
FP8
INT8
TF32
Pemfaktoran ulang pada kernel::ConvUniversal agar sejalan dengan gemm::GemmUniversal . Sekarang API konvolusi 3.x tidak lagi dianggap sebagai API beta.
GEMM masukan campuran yang ditingkatkan dan implementasi tabel pencarian untuk mode skala saja INT4 x FP8 .
Node EVT untuk pemilihan Top-K dan contoh softmax dan GEMM menggunakannya.
Peluncuran Ketergantungan Terprogram (PDL) yang memanfaatkan fitur Hopper baru untuk mempercepat dua kernel berturut-turut, dan dokumentasi terkait.
Alat debugging baru, synclog, untuk membuang semua peristiwa sinkronisasi dari dalam kernel ke file. Silakan lihat dokumentasi sinkronisasi untuk detailnya.
Epilog berkemampuan TMA baru untuk GEMM berkelompok yang menghadirkan peningkatan kinerja signifikan, serta dukungan EVT-nya.
Epilog array penunjuk berkemampuan SIMT.
Jadwal kernel Ping-Pong baru untuk Grouped GEMM dan beberapa optimasi lainnya.
Strategi instantiasi baru untuk kernel profiler CUTLASS bersama dengan dokumentasi yang ditingkatkan untuk tingkat instantiasi di profiler CUTLASS.
Dukungan perangkat keras baru untuk perbandingan dan perhitungan cutlass::bfloat16_t
Memperbaiki penggunaan isnan di Windows untuk half_t .
Persyaratan minimal:
Arsitektur: Volta
Kompiler: Harus mendukung setidaknya C++17
Versi Perangkat CUDA: 11.4
Mulai dari CUTLASS 3.0, CUTLASS menghapus dukungan untuk hal berikut:
Arsitektur GPU Maxwell dan Pascal
Ubuntu 16.04
CUDA 10.2
Versi bahasa C++ kurang dari 17.
Lihat CHANGELOG untuk daftar rinci rilis dan pembaruan.
Primitif CUTLASS sangat efisien. Saat digunakan untuk membuat kernel GEMM di seluruh perangkat, kernel tersebut menunjukkan kinerja puncak yang sebanding dengan cuBLAS untuk komputasi skalar GEMM. Gambar di atas menunjukkan peningkatan kinerja CUTLASS yang berkelanjutan pada NVIDIA H100 (arsitektur NVIDIA Hopper) sejak CUTLASS 3.1. CUTLASS 3.5.1 dikompilasi dengan Toolkit CUDA 12.5u1. Operasi Tensor Core diimplementasikan menggunakan instruksi mma dan wgmma CUDA.
Saat menggunakan blok penyusun CUTLASS untuk membuat kernel permata implisit (Fprop, Dgrad, dan Wgrad) di seluruh perangkat, performa CUTLASS juga sebanding dengan cuDNN saat menjalankan lapisan Resnet-50 pada NVIDIA A100 seperti yang ditunjukkan pada gambar di atas. Operasi Tensor Core diimplementasikan menggunakan instruksi mma CUDA.
CUTLASS memerlukan kompiler host C++17 dan berkinerja paling baik bila dibuat dengan CUDA 12.4 Toolkit . Ini juga kompatibel dengan 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 dan CUDA 12.3.2.
Kami telah menguji lingkungan berikut.
| Sistem Operasi | Penyusun |
|---|---|
| Ubuntu 18.04 | GCC 7.5.0 |
| Ubuntu 20.04 | GCC 10.3.0 |
| Ubuntu 22.04 | GCC 11.2.0 |
| Ubuntu 22.04 | Dentang 10.0.0 |
| Ubuntu 22.04 | Dentang 14.0.6 |
| Ubuntu 22.04 | Dentang 17.0.6 |
| jendela 10.0 | Visual Studio 2019 v16.11.27 |
Catatan: GCC 8.5.0 telah mengetahui regresi terkait ekspresi lipatan dan operator yang kelebihan beban. Disarankan menggunakan GCC 7.5.0 atau (lebih disukai) GCC >= 9.
CUTLASS berjalan dengan sukses pada GPU NVIDIA berikut, dan diharapkan efisien pada GPU NVIDIA berbasis arsitektur Volta, Turing, Ampere, Ada, dan Hopper.
| GPU | Kemampuan Komputasi CUDA | Toolkit CUDA Minimum yang Diperlukan oleh CUTLASS-3 |
|---|---|---|
| GPU Inti Tensor NVIDIA V100 | 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 Inti Tensor NVIDIA A100 | 8.0 | 11.4 |
| NVIDIA A10 | 8.6 | 11.4 |
| NVIDIAGeForce RTX 3090 | 8.6 | 11.4 |
| NVIDIAGeForce RTX 4090 | 8.9 | 11.8 |
| NVIDIA L40 | 8.9 | 11.8 |
| GPU Inti Tensor NVIDIA H100 | 9.0 | 11.8 |
Secara umum, kode PTX yang dihasilkan untuk satu arsitektur target dapat dijalankan pada arsitektur masa depan (yaitu kompatibel ke depan). Namun, CUDA 12.0 memperkenalkan konsep "fitur akselerasi arsitektur" yang PTX-nya tidak memiliki jaminan kompatibilitas ke depan. Beberapa instruksi Hopper PTX termasuk dalam kategori fitur akselerasi arsitektur ini, dan karenanya memerlukan arsitektur target sm_90a (perhatikan tambahan "a"). Untuk detail lebih lanjut tentang hal ini dan instruksi akselerasi arsitektur lainnya, silakan merujuk ke Dokumentasi CUDA.
Informasi arsitektur target diteruskan ke CUTLASS melalui flag cmake CUTLASS_NVCC_ARCHS . Untuk memaksimalkan kinerja pada Hopper GH100, pengguna diharuskan membangun CUTLASS dengan 90a sebagai arsitektur target. Jika pengguna secara tidak sengaja membuat kernel yang menggunakan fitur SM90a (misalnya Petunjuk Inti Hopper Tensor), menggunakan target SM90 (perhatikan kekurangan "a"), dengan CUDA Toolkit 12 atau 11.8, kernel diperkirakan akan gagal saat runtime kesalahan.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
Silakan merujuk ke dokumentasi fungsionalitas untuk mengetahui detail tentang kernel mana yang memerlukan arsitektur target tertentu.
CUTLASS dijelaskan dalam dokumen berikut dan dokumentasi Doxygen yang menyertainya.
Panduan Memulai Cepat - buat dan jalankan CUTLASS
Fungsionalitas - merangkum fungsionalitas yang tersedia di CUTLASS
GEMM yang efisien di CUDA - menjelaskan bagaimana kernel GEMM dapat diimplementasikan secara efisien di CUDA
Desain CUTLASS 3.x - menjelaskan desain CUTLASS 3.x, manfaatnya, dan bagaimana CuTe memungkinkan kita menulis lebih banyak komponen yang dapat disusun
GEMM API 3.x - menjelaskan model CUTLASS 3.x GEMM dan konsep template C++
GEMM API 2.x - menjelaskan model CUTLASS 2.x GEMM dan konsep template C++
Konvolusi GEMM Implisit - menjelaskan konvolusi 2-D dan 3-D di CUTLASS
Organisasi Kode - menjelaskan organisasi dan konten proyek CUTLASS
Terminologi - menjelaskan istilah yang digunakan dalam kode
Pedoman Pemrograman - pedoman untuk menulis CUDA C++ modern yang efisien
Tipe fundamental - menjelaskan kelas C++ dasar yang digunakan di CUTLASS untuk merepresentasikan kuantitas dan array numerik
Tata Letak - menjelaskan tata letak matriks dan tensor dalam memori
Tile Iterators - menjelaskan konsep C++ untuk melakukan iterasi pada ubin matriks di memori
CUTLASS Profiler - aplikasi pembuatan profil berbasis baris perintah
CUTLASS Utilities - templat tambahan yang digunakan untuk memfasilitasi pengembangan cepat
Peluncuran kernel dependen - menjelaskan fitur baru di Hopper yang memungkinkan kernel dependen tumpang tindih dalam aliran yang sama, dan cara penggunaannya di CUTLASS.
Kami juga telah menjelaskan struktur GEMM yang efisien dalam pembicaraan kami di Konferensi Teknologi GPU 2018.
CUTLASS: Perangkat Lunak Primitif untuk Aljabar Linier Padat di Semua Tingkat dan Skala dalam CUDA
Mengembangkan Kernel CUDA untuk Mendorong Inti Tensor ke Batas Absolut pada NVIDIA A100
Mempercepat Konvolusi dengan Tensor Cores di CUTLASS
Mempercepat Gradien Data Mundur dengan Meningkatkan Pemanfaatan Inti Tensor di CUTLASS
CUTLASS: Python API, Penyempurnaan, dan NVIDIA Hopper
CUTLASS adalah pustaka templat khusus header dan tidak perlu dibuat untuk digunakan oleh proyek lain. Aplikasi klien harus menargetkan direktori include/ CUTLASS di jalur penyertaannya.
Pengujian unit CUTLASS, contoh, dan utilitas dapat dibuat dengan CMake. Versi minimum CMake diberikan dalam panduan Memulai Cepat. Pastikan variabel lingkungan CUDACXX menunjuk ke NVCC di CUDA Toolkit yang diinstal pada sistem Anda.
$ ekspor CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc Buat direktori build dalam proyek CUTLASS, lalu jalankan CMake. Secara default CUTLASS akan membangun kernel untuk arsitektur CUDA versi 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9, dan 9.0. Untuk mengurangi waktu kompilasi, Anda dapat menentukan arsitektur yang akan dibangun CUTLASS dengan mengubah pengaturan konfigurasi CMake CUTLASS_NVCC_ARCHS .
$ mkdir membangun && cd membangun $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # kompilasi untuk Arsitektur Ampere NVIDIA
Dari direktori build/ , kompilasi dan jalankan pengujian unit CUTLASS dengan membuat target test_unit dengan make.
Pengujian unit diatur sebagai beberapa binari yang mencerminkan namespace tingkat atas CUTLASS, dan dapat dieksekusi secara paralel melalui argumen baris perintah -j make.
$ buat unit_tes -j ... ... ... [----------] Penghancuran lingkungan pengujian global [==========] 946 pengujian dari 57 kasus uji dijalankan. (total 10812 ms) [LULUS] 946 tes.
Semua tes harus lulus pada platform yang didukung, meskipun jumlah pasti tes dapat bervariasi dari waktu ke waktu.
CUTLASS disusun sebagai perpustakaan khusus header bersama dengan Utilitas, Alat, Contoh, dan pengujian unit. Dokumentasi Doxygen menyediakan daftar lengkap file, kelas, dan konsep templat yang ditentukan dalam proyek CUTLASS.
Penjelasan rinci tentang organisasi kode sumber dapat ditemukan dalam dokumentasi CUTLASS, namun beberapa komponen utama dirangkum di bawah ini.
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
Contoh CUTLASS SDK menerapkan template CUTLASS untuk mengimplementasikan komputasi dasar.
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.
Direktori test/unit/ terdiri dari pengujian unit yang diimplementasikan dengan Google Test yang mendemonstrasikan penggunaan dasar komponen Core API dan menyelesaikan pengujian komputasi CUTLASS GEMM.
Petunjuk untuk membuat dan menjalankan pengujian Unit dijelaskan dalam panduan Memulai Cepat.
Direktori tools/profiler/ berisi utilitas baris perintah untuk meluncurkan setiap kernel GEMM. Itu dapat dibangun sebagai berikut:
$ membuat cutlass_profiler -j16
Secara default, hanya satu ukuran ubin yang dipakai untuk setiap tipe data, instruksi matematika, dan tata letak. Untuk membuat instance semuanya, atur variabel lingkungan berikut saat menjalankan CMake dari direktori build/ yang kosong. Hati-hati, ini menghasilkan puluhan ribu kernel dan waktu pembuatan yang lama. Hal ini juga akan mengakibatkan ukuran biner yang besar dan pada beberapa platform linker gagal dalam membangun perpustakaan. Oleh karena itu, sangat disarankan untuk menghasilkan hanya sebagian dari kernel seperti yang ditunjukkan pada sub-bagian di bawah.
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=semua ... $ membuat cutlass_profiler -j16
Untuk mengkompilasi hanya satu kernel atau sekumpulan kecil kernel, daftar nama kernel yang dipisahkan koma dengan karakter wildcard dapat digunakan untuk mengurangi kumpulan kernel. Contoh berikut menunjukkan pembuatan satu atau sebagian kernel untuk arsitektur NVIDIA Ampere dan Turing:
Untuk mengkompilasi subset kernel Tensor Core GEMM dengan akumulasi FP32 dan input FP16 yang menargetkan arsitektur NVIDIA Ampere dan Turing, gunakan baris perintah cmake di bawah ini:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ membuat cutlass_profiler -j16
Contoh baris perintah untuk membuat profil subset kernel Tensor Core GEMM adalah sebagai berikut:
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
==============
ID Masalah: 1
Penyedia: CUTLASS
Jenis Operasi: permata
Operasi: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Status: Sukses
Verifikasi: AKTIF
Disposisi: Lulus
reference_device: Lulus
cuBLAS: Lulus
Argumen: --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 --tahapan=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
Byte: 118489088 byte
FLOP: 115992428544 gagal
Waktu proses: 1,55948 mdtk
Memori: 70,7616 GiB/dtk
Matematika: 74378,8 GFLOP/s
==============
...Untuk mengkompilasi satu kernel SGEMM yang menargetkan arsitektur NVIDIA Ampere dan Turing, gunakan baris perintah cmake di bawah ini:
$ cmake.. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ membuat cutlass_profiler -j16
Contoh baris perintah untuk membuat profil kernel SGEMM CUDA tunggal adalah sebagai berikut:
$ ./tools/profiler/cutlass_profiler --kernel=sgemm --m=3456 --n=4096 --k=4096
==============
ID Masalah: 1
Penyedia: CUTLASS
Jenis Operasi: permata
Operasi: cutlass_simt_sgemm_128x128_8x2_nn_align1
Status: Sukses
Verifikasi: AKTIF
Disposisi: Lulus
cuBLAS: Lulus
Argumen: --m=3456 --n=4096 --k=4096 --A=f32:kolom --B=f32:kolom --C=f32:kolom --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
Byte: 180355072 byte
FLOP: 115992428544 gagal
Waktu proses: 6,73655 mdtk
Memori: 24,934 GiB/dtk
Matematika: 17218,4 GFLOP/s
==============Untuk mengkompilasi subset kernel konvolusi inti Tensor yang mengimplementasikan propagasi maju (fprop) dengan akumulasi FP32 dan input FP16 yang menargetkan arsitektur NVIDIA Ampere dan Turing, gunakan baris perintah cmake di bawah ini:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $ membuat cutlass_profiler -j16
Contoh baris perintah untuk membuat profil subset kernel konvolusi Tensor Core adalah sebagai berikut:
$ ./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 Masalah: 1
Penyedia: CUTLASS
Jenis Operasi: konv2d
Operasi: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Status: Sukses
Verifikasi: AKTIF
Disposisi: Lulus
reference_device: Lulus
Argumen: --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 --Aktivasi=f16:nhwc --Filter=f16:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=dioptimalkan --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1
--eq_gemm_provider=tidak ada --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
Byte: 1130659840 byte
FLOP: 118482796544 gagal
Waktu proses: 0,711496 mdtk
Memori: 1479,99 GiB/dtk
Matematika: 166526 GFLOP/s
==============
...Untuk mengkompilasi dan menjalankan satu kernel konvolusi CUDA Core yang mengimplementasikan propagasi maju (fprop) dengan akumulasi F32 dan input FP32 yang menargetkan arsitektur NVIDIA Ampere dan Turing, gunakan baris perintah cmake di bawah ini:
$ cmake.. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ membuat cutlass_profiler -j16
Contoh baris perintah untuk membuat profil satu kernel konvolusi 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 Masalah: 1
Penyedia: CUTLASS
Jenis Operasi: konv2d
Operasi: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Status: Sukses
Verifikasi: AKTIF
Disposisi: Lulus
reference_device: Lulus
Argumen: --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 --Aktivasi=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=dioptimalkan --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1
--eq_gemm_provider=tidak ada --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
Byte: 2055798784 byte
FLOP: 118482796544 gagal
Waktu proses: 7,34266 mdtk
Memori: 260,752 GiB/dtk
Matematika: 16136,2 GFLOP/s
==============Silakan ikuti tautan untuk melihat lebih banyak contoh CMake tentang kompilasi selektif kernel CUTLASS:
GEMM CBuat Contoh
Contoh CMake konvolusi GEMM implisit
Detail lebih lanjut tentang CUTLASS Profiler dijelaskan di sini.
CUTLASS dirilis oleh NVIDIA Corporation sebagai perangkat lunak Open Source di bawah lisensi BSD "Baru" 3-klausul.
Daftar resmi pengembang dan kontributor CUTLASS tersedia di sini: KONTRIBUTOR.
Hak Cipta (c) 2017 - 2024 NVIDIA CORPORATION & AFILIASI. Semua hak dilindungi undang-undang. Pengenal-Lisensi SPDX: Klausul 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.
