cutlass
CUTLASS
커틀라스 3.6.0 - 2024년 10월
CUTLASS는 CUDA 내의 모든 수준과 규모에서 고성능 GEMM(행렬-행렬 곱셈) 및 관련 계산을 구현하기 위한 CUDA C++ 템플릿 추상화 모음입니다. 여기에는 cuBLAS 및 cuDNN을 구현하는 데 사용되는 것과 유사한 계층적 분해 및 데이터 이동 전략이 포함되어 있습니다. CUTLASS는 이러한 "움직이는 부분"을 C++ 템플릿 클래스로 추상화된 재사용 가능한 모듈식 소프트웨어 구성 요소로 분해합니다. 개념적 병렬화 계층의 다양한 수준에 대한 기본 요소는 사용자 정의 타일링 크기, 데이터 유형 및 기타 알고리즘 정책을 통해 전문화되고 조정될 수 있습니다. 결과적인 유연성으로 인해 사용자 정의 커널 및 애플리케이션 내에서 빌딩 블록으로 사용이 단순화됩니다.
다양한 애플리케이션을 지원하기 위해 CUTLASS는 혼합 정밀도 계산에 대한 광범위한 지원을 제공하여 반 정밀도 부동 소수점(FP16), BFloat16(BF16), Tensor Float 32(TF32)에 대한 특수 데이터 이동 및 곱셈 누산 추상화를 제공합니다. 단정밀도 부동 소수점(FP32), 텐서 코어 명령어를 통한 FP32 에뮬레이션, 배정밀도 부동 소수점(FP64) 유형, 정수 데이터 유형 (4b 및 8b) 및 이진 데이터 유형(1b). CUTLASS는 NVIDIA의 Volta, Turing, Ampere 및 Hopper 아키텍처로 구현된 프로그래밍 가능하고 처리량이 높은 Tensor 코어를 대상으로 하는 워프 동기 행렬 곱셈 작업을 보여줍니다.
빠르게 시작하려면 빠른 시작 가이드를 참조하세요.
실행 모델 계층의 각 수준에서 지원되는 작업 목록은 기능 목록을 참조하세요.
CUTLASS 3.0에는 스레드와 데이터의 텐서를 설명하고 조작하기 위한 새로운 핵심 라이브러리인 CuTe가 도입되었습니다. CuTe는 스레드와 데이터의 계층적 다차원 레이아웃을 정의하고 작동하기 위한 C++ CUDA 템플릿 추상화 모음입니다. CuTe는 사용자를 위해 복잡한 인덱싱을 수행하면서 데이터의 유형, 모양, 메모리 공간 및 레이아웃을 컴팩트하게 패키지화하는 Layout 및 Tensor 객체를 제공합니다. 이를 통해 프로그래머는 알고리즘의 논리적 설명에 집중할 수 있고 CuTe는 기계적인 장부를 수행합니다. 이러한 도구를 사용하면 모든 조밀한 선형 대수 연산을 신속하게 설계, 구현 및 수정할 수 있습니다.
CuTe의 핵심 추상화는 텐서를 나타내는 데이터 배열로 구성될 수 있는 계층적 다차원 레이아웃입니다. 레이아웃 표현은 효율적인 조밀 선형 대수를 구현하는 데 필요한 거의 모든 것을 표현할 수 있을 만큼 강력합니다. 레이아웃은 타일링 및 파티셔닝과 같은 대규모 공통 작업 세트를 구축하는 기능적 구성을 통해 결합 및 조작될 수도 있습니다.
CUTLASS 3.0 이상은 템플릿의 GEMM 계층 구조 전반에 걸쳐 CuTe를 채택합니다. 이는 디자인을 크게 단순화하고 코드 구성성과 가독성을 향상시킵니다. CuTe와 관련된 추가 문서는 전용 문서 디렉터리에서 찾을 수 있습니다.
GEMM 외에도 CUTLASS는 암시적 GEMM 알고리즘을 통해 고성능 컨볼루션을 구현합니다. 암시적 GEMM은 CUTLASS의 모듈식 GEMM 파이프라인을 활용하여 컨볼루션 작업을 GEMM으로 공식화한 것입니다. 이를 통해 CUTLASS는 고도로 최적화된 GEMM 구성 요소를 재사용하여 컨볼루션을 구축할 수 있습니다.
CUTLASS 3.6.0은 CUTLASS에 다음을 추가하는 업데이트입니다.
호퍼 구조의 스파스 GEMM.
FP16
FP8
INT8
TF32
CUTLASS 3.x 컨볼루션 kernel::ConvUniversal API를 리팩터링하여 gemm::GemmUniversal 과 일치시킵니다. 이제 3.x 컨볼루션 API는 더 이상 베타 API로 간주되지 않습니다.
INT4 x FP8 스케일 전용 모드에 대한 향상된 혼합 입력 GEMM 및 조회 테이블 구현입니다.
Top-K 선택을 위한 EVT 노드와 이를 사용하는 소프트맥스 및 GEMM 예제입니다.
새로운 Hopper 기능을 활용하여 2개의 연속 커널과 해당 문서의 속도를 높이는 PDL(Programmatic dependency Launch)입니다.
커널 내에서 파일로 모든 동기화 이벤트를 덤프하기 위한 새로운 디버깅 도구인 synclog입니다. 자세한 내용은 synclog 설명서를 참조하세요.
EVT 지원은 물론 상당한 성능 향상을 제공하는 그룹화된 GEMM을 위한 새로운 TMA 지원 에필로그입니다.
SIMT 지원 포인터 배열 에필로그.
그룹화된 GEMM을 위한 새로운 Ping-Pong 커널 일정 및 기타 최적화.
CUTLASS 프로파일러의 인스턴스화 수준에 대한 향상된 문서와 함께 CUTLASS 프로파일러 커널을 위한 새로운 인스턴스화 전략입니다.
cutlass::bfloat16_t 의 비교 및 계산을 위한 새로운 하드웨어 지원
Windows에서 half_t 에 대한 isnan 사용을 수정했습니다.
최소 요구사항:
아키텍처: 볼타
컴파일러: 최소한 C++17을 지원해야 합니다.
CUDA 툴킷 버전: 11.4
CUTLASS 3.0부터 CUTLASS는 다음에 대한 지원을 제거했습니다.
Maxwell 및 Pascal GPU 아키텍처
우분투 16.04
쿠다 10.2
C++ 언어 버전 17 미만.
릴리스 및 업데이트의 자세한 목록은 변경 로그를 참조하세요.
CUTLASS 기본 요소는 매우 효율적입니다. 장치 전체의 GEMM 커널을 구성하는 데 사용되면 스칼라 GEMM 계산을 위한 cuBLAS에 필적하는 최고 성능을 나타냅니다. 위 그림은 CUTLASS 3.1 이후 NVIDIA H100(NVIDIA Hopper 아키텍처)의 지속적인 CUTLASS 성능 개선을 보여줍니다. CUTLASS 3.5.1은 CUDA 12.5u1 툴킷으로 컴파일되었습니다. Tensor Core 작업은 CUDA의 mma 및 wgmma 명령어를 사용하여 구현됩니다.
CUTLASS 빌딩 블록을 사용하여 장치 전체의 암시적 gemm(Fprop, Dgrad 및 Wgrad) 커널을 구성할 때 CUTLASS 성능은 위 그림에 표시된 대로 NVIDIA A100에서 Resnet-50 레이어를 실행할 때 cuDNN과 비슷합니다. Tensor Core 작업은 CUDA의 mma 명령어를 사용하여 구현됩니다.
CUTLASS에는 C++17 호스트 컴파일러가 필요하며 CUDA 12.4 툴킷 으로 빌드할 때 최고의 성능을 발휘합니다. 또한 CUDA 11.4, CUDA 11.5, CUDA 11.6, CUDA 11.7, CUDA 11.8, CUDA 12.0, CUDA 12.1, CUDA 12.2.2, CUDA 12.3.1 및 CUDA 12.3.2와도 호환됩니다.
우리는 다음 환경을 테스트했습니다.
| 운영 체제 | 컴파일러 |
|---|---|
| 우분투 18.04 | GCC 7.5.0 |
| 우분투 20.04 | GCC 10.3.0 |
| 우분투 22.04 | GCC 11.2.0 |
| 우분투 22.04 | 클랭 10.0.0 |
| 우분투 22.04 | 클랭 14.0.6 |
| 우분투 22.04 | 클랭 17.0.6 |
| 윈도우 10.0 | 비주얼 스튜디오 2019 v16.11.27 |
참고: GCC 8.5.0에는 접기 표현식 및 오버로드된 연산자와 관련된 알려진 회귀가 있습니다. GCC 7.5.0 또는 (선호) GCC >= 9를 사용하는 것이 좋습니다.
CUTLASS는 다음 NVIDIA GPU에서 성공적으로 실행되며 Volta, Turing, Ampere, Ada 및 Hopper 아키텍처 기반 NVIDIA GPU에서 효율적일 것으로 예상됩니다.
| GPU | CUDA 컴퓨팅 기능 | CUTLASS-3에 필요한 최소 CUDA 툴킷 |
|---|---|---|
| NVIDIA V100 텐서 코어 GPU | 7.0 | 11.4 |
| 엔비디아 타이탄V | 7.0 | 11.4 |
| 엔비디아 지포스 RTX 2080TI, 2080, 2070 | 7.5 | 11.4 |
| 엔비디아 T4 | 7.5 | 11.4 |
| NVIDIA A100 텐서 코어 GPU | 8.0 | 11.4 |
| 엔비디아 A10 | 8.6 | 11.4 |
| 엔비디아 지포스 RTX 3090 | 8.6 | 11.4 |
| 엔비디아 지포스 RTX 4090 | 8.9 | 11.8 |
| 엔비디아 L40 | 8.9 | 11.8 |
| NVIDIA H100 텐서 코어 GPU | 9.0 | 11.8 |
일반적으로 하나의 대상 아키텍처에 대해 생성된 PTX 코드는 향후 아키텍처에서 실행될 수 있습니다(즉, 향후 호환 가능). 그러나 CUDA 12.0에는 PTX에 향후 호환성이 보장되지 않는 "아키텍처 가속 기능"이라는 개념이 도입되었습니다. 여러 Hopper PTX 명령어는 이 아키텍처 가속 기능 범주에 속하므로 sm_90a 대상 아키텍처가 필요합니다(추가된 "a" 참고). 이에 대한 자세한 내용과 기타 아키텍처 가속 지침은 CUDA 문서를 참조하세요.
대상 아키텍처 정보는 cmake 플래그 CUTLASS_NVCC_ARCHS 를 통해 CUTLASS에 전달됩니다. Hopper GH100의 성능을 최대화하려면 사용자는 90a 대상 아키텍처로 사용하여 CUTLASS를 구축해야 합니다. 사용자가 SM90a 기능(예: Hopper Tensor Core Instructions)을 사용하는 커널을 실수로 CUDA Toolkit 12 또는 11.8과 함께 SM90 대상("a" 부족 참고)을 사용하여 빌드하는 경우 커널은 런타임과 함께 실패할 것으로 예상됩니다. 오류.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
어떤 커널에 어떤 대상 아키텍처가 필요한지에 대한 자세한 내용은 기능 문서를 참조하세요.
CUTLASS는 다음 문서와 함께 제공되는 Doxygen 문서에 설명되어 있습니다.
빠른 시작 가이드 - CUTLASS 빌드 및 실행
기능 - CUTLASS에서 사용할 수 있는 기능을 요약합니다.
CUDA의 효율적인 GEMM - CUDA에서 GEMM 커널을 효율적으로 구현하는 방법을 설명합니다.
CUTLASS 3.x 디자인 - CUTLASS 3.x 디자인, 그 이점 및 CuTe를 통해 훨씬 더 많은 구성 가능한 구성 요소를 작성할 수 있는 방법을 설명합니다.
GEMM API 3.x - CUTLASS 3.x GEMM 모델 및 C++ 템플릿 개념을 설명합니다.
GEMM API 2.x - CUTLASS 2.x GEMM 모델 및 C++ 템플릿 개념을 설명합니다.
암시적 GEMM 컨볼루션 - CUTLASS의 2D 및 3D 컨볼루션을 설명합니다.
코드 구성 - CUTLASS 프로젝트의 구성과 내용을 설명합니다.
용어 - 코드에 사용된 용어를 설명합니다.
프로그래밍 지침 - 효율적인 최신 CUDA C++ 작성을 위한 지침
기본 유형 - 숫자 수량 및 배열을 표현하기 위해 CUTLASS에서 사용되는 기본 C++ 클래스를 설명합니다.
레이아웃 - 메모리의 행렬 및 텐서 레이아웃을 설명합니다.
타일 반복자(Tile Iterator) - 메모리의 행렬 타일을 반복하는 C++ 개념을 설명합니다.
CUTLASS 프로파일러 - 명령줄 기반 프로파일링 애플리케이션
CUTLASS 유틸리티 - 신속한 개발을 촉진하는 데 사용되는 추가 템플릿
종속 커널 실행 - 동일한 스트림에서 종속 커널을 겹칠 수 있는 Hopper의 새로운 기능과 CUTLASS에서 이 기능이 사용되는 방법을 설명합니다.
우리는 또한 GPU Technology Conference 2018에서의 강연에서 효율적인 GEMM의 구조를 설명했습니다.
CUTLASS: CUDA 내의 모든 수준과 규모에서 조밀한 선형 대수학을 위한 소프트웨어 프리미티브
NVIDIA A100의 절대 한계까지 Tensor 코어를 확장하기 위한 CUDA 커널 개발
CUTLASS의 Tensor 코어로 컨볼루션 가속화
CUTLASS에서 Tensor Core 활용도를 높여 역방향 데이터 기울기 가속화
CUTLASS: Python API, 향상된 기능 및 NVIDIA Hopper
CUTLASS는 헤더 전용 템플릿 라이브러리이므로 다른 프로젝트에서 사용하기 위해 빌드할 필요가 없습니다. 클라이언트 애플리케이션은 포함 경로에서 CUTLASS의 include/ 디렉터리를 대상으로 해야 합니다.
CUTLASS 단위 테스트, 예제 및 유틸리티는 CMake를 사용하여 빌드할 수 있습니다. CMake의 최소 버전은 빠른 시작 가이드에 나와 있습니다. CUDACXX 환경 변수가 시스템에 설치된 CUDA 툴킷의 NVCC를 가리키는지 확인하세요.
$ 내보내기 CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc CUTLASS 프로젝트 내에 빌드 디렉터리를 만든 다음 CMake를 실행합니다. 기본적으로 CUTLASS는 CUDA 아키텍처 버전 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9 및 9.0용 커널을 빌드합니다. 컴파일 시간을 줄이려면 CMake 구성 설정 CUTLASS_NVCC_ARCHS 를 변경하여 CUTLASS를 빌드할 아키텍처를 지정할 수 있습니다.
$ mkdir 빌드 && cd 빌드 $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # NVIDIA의 Ampere 아키텍처용으로 컴파일합니다.
build/ 디렉터리에서 make를 사용하여 대상 test_unit 빌드하여 CUTLASS 단위 테스트를 컴파일하고 실행합니다.
단위 테스트는 CUTLASS의 최상위 네임스페이스를 미러링하는 여러 바이너리로 구성되며 make의 -j 명령줄 인수를 통해 병렬로 실행될 수 있습니다.
$ make test_unit -j ... ... ... [----------] 글로벌 테스트 환경 분해 [==========] 57개의 테스트 사례에서 946개의 테스트가 실행되었습니다. (총 10812ms) [ 합격 ] 946번의 테스트를 통과했습니다.
모든 테스트는 지원되는 플랫폼에서 통과해야 하지만 정확한 테스트 수는 시간에 따라 달라질 수 있습니다.
CUTLASS는 유틸리티, 도구, 예제, 단위 테스트와 함께 헤더 전용 라이브러리로 구성됩니다. Doxygen 문서는 CUTLASS 프로젝트에 정의된 파일, 클래스 및 템플릿 개념의 전체 목록을 제공합니다.
소스 코드 구성에 대한 자세한 설명은 CUTLASS 문서에서 찾을 수 있지만 몇 가지 주요 구성 요소가 아래에 요약되어 있습니다.
include/ # client applications should target this directory in their build's include paths cutlass/ # CUDA Templates for Linear Algebra Subroutines and Solvers - headers only arch/ # direct exposure of architecture features (including instruction-level GEMMs) conv/ # code specialized for convolution epilogue/ # code specialized for the epilogue of gemm/convolution gemm/ # code specialized for general matrix product computations layout/ # layout definitions for matrices, tensors, and other mathematical objects in memory platform/ # CUDA-capable Standard Library components reduction/ # bandwidth-limited reduction kernels that do not fit the "gemm" model thread/ # simt code that can be performed within a CUDA thread transform/ # code specialized for layout, type, and domain transformations * # core vocabulary types, containers, and basic numeric operations cute/ # CuTe Layout, layout algebra, MMA/Copy atoms, tiled MMA/Copy algorithm/ # Definitions of core operations such as copy, gemm, and operations on cute::tuples arch/ # Bare bones PTX wrapper structs for copy and math instructions atom/ # Meta-information either link to or built from arch/ operators mma_atom.hpp # cute::Mma_Atom and cute::TiledMma copy_atom.hpp # cute::Copy_Atom and cute::TiledCopy *sm*.hpp # Arch specific meta-information for copy and math operations * # Core library types such as Shape, Stride, Layout, Tensor, and associated operations
CUTLASS SDK 예제는 CUTLASS 템플릿을 적용하여 기본 계산을 구현합니다.
tools/ library/ # CUTLASS Instance Library - contains instantiations of all supported CUTLASS templates include/ cutlass/ library/ profiler/ # CUTLASS Profiler - command-line utility for executing operations in the # CUTLASS Library util/ # CUTLASS Utilities - contains numerous helper classes for include/ # manging tensors in device memory, reference cutlass/ # implementations for GEMM, random initialization util/ # of tensors, and I/O.
test/unit/ 디렉토리는 Core API 구성 요소의 기본 사용법을 보여주는 Google Test로 구현된 단위 테스트와 CUTLASS GEMM 계산의 전체 테스트로 구성됩니다.
단위 테스트 빌드 및 실행에 대한 지침은 빠른 시작 가이드에 설명되어 있습니다.
tools/profiler/ 디렉토리에는 각 GEMM 커널을 시작하기 위한 명령줄 유틸리티가 포함되어 있습니다. 다음과 같이 구축할 수 있습니다.
$ make cutlass_profiler -j16
기본적으로 각 데이터 유형, 수학 명령어 및 레이아웃에 대해 하나의 타일 크기만 인스턴스화됩니다. 모두 인스턴스화하려면 빈 build/ 디렉터리에서 CMake를 실행할 때 다음 환경 변수를 설정하세요. 주의하세요. 이로 인해 수만 개의 커널이 생성되고 빌드 시간이 길어집니다. 이로 인해 바이너리 크기가 커지고 일부 플랫폼에서는 링커가 라이브러리 빌드에 실패하게 됩니다. 따라서 아래 하위 섹션에 설명된 대로 커널의 하위 집합만 생성하는 것이 좋습니다.
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=모두 ... $ make cutlass_profiler -j16
하나의 커널 또는 작은 커널 세트만 컴파일하려면 와일드카드 문자가 포함된 쉼표로 구분된 커널 이름 목록을 사용하여 커널 세트를 줄일 수 있습니다. 다음 예에서는 NVIDIA Ampere 및 Turing 아키텍처에 대해 정확히 하나 또는 커널 하위 집합을 구축하는 방법을 보여줍니다.
NVIDIA Ampere 및 Turing 아키텍처를 대상으로 하는 FP32 누적 및 FP16 입력을 사용하여 Tensor Core GEMM 커널의 하위 집합을 컴파일하려면 아래 cmake 명령줄을 사용하세요.
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 ... $ make cutlass_profiler -j16
Tensor Core GEMM 커널의 하위 집합을 프로파일링하기 위한 명령줄 예시는 다음과 같습니다.
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
=============================
문제 ID: 1
제공자: 커틀라스
OperationKind: gemm
작업: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
상태: 성공
확인: 켜짐
처분: 합격
reference_device: 통과됨
cuBLAS: 합격
인수: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:column --B=f16:row --C=f32:column --alpha=1
--beta=0 --split_k_slices=1 --batch_count=1 --op_class=tensorop --accum=f32 --cta_m=256 --cta_n=128
--cta_k=32 --stages=2 --warps_m=4 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=8 --min_cc=75
--max_cc=1024
바이트: 118489088바이트
플롭: 115992428544 플롭
런타임: 1.55948ms
메모리: 70.7616GiB/s
수학: 74378.8 GFLOP/s
=============================
...NVIDIA Ampere 및 Turing 아키텍처를 대상으로 하는 하나의 SGEMM 커널을 컴파일하려면 아래 cmake 명령줄을 사용하십시오.
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 ... $ make cutlass_profiler -j16
단일 SGEMM CUDA 커널을 프로파일링하기 위한 명령줄의 예는 다음과 같습니다.
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
=============================
문제 ID: 1
제공자: 커틀라스
OperationKind: gemm
작업: cutlass_simt_sgemm_128x128_8x2_nn_align1
상태: 성공
확인: 켜짐
처분: 합격
cuBLAS: 합격
인수: --m=3456 --n=4096 --k=4096 --A=f32:column --B=f32:column --C=f32:column --alpha=1 --beta=0 -- 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
바이트: 180355072바이트
플롭: 115992428544 플롭
런타임: 6.73655ms
메모리: 24.934GiB/s
수학: 17218.4 GFLOP/s
=============================NVIDIA Ampere 및 Turing 아키텍처를 대상으로 하는 FP32 누적 및 FP16 입력을 사용하여 순방향 전파(fprop)를 구현하는 Tensor 코어 컨볼루션 커널의 하위 집합을 컴파일하려면 아래 cmake 명령줄을 사용하세요.
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 ... $ make cutlass_profiler -j16
Tensor Core 컨볼루션 커널의 하위 집합을 프로파일링하기 위한 명령줄의 예는 다음과 같습니다.
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*fprop_optimized_f16 --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3
...
=============================
문제 ID: 1
제공자: 커틀라스
OperationKind: 전환2d
작업: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
상태: 성공
확인: 켜짐
처분: 합격
reference_device: 통과됨
인수: --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 --활성화=f16:nhwc --필터=f16:nhwc --출력=f32:nhwc
--conv_mode=교차 --iterator_algorithm=최적화 --alpha=1 --beta=0 --split_k_mode=직렬 --split_k_slices=1
--eq_gemm_provider=none --op_class=tensorop --accum=f32 --cta_m=128 --cta_n=128 --cta_k=32 --stages=5
--warps_m=2 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=16 --min_cc=80 --max_cc=1024
바이트: 1130659840바이트
플롭: 118482796544 플롭
런타임: 0.711496ms
메모리: 1479.99GiB/s
수학: 166526 GFLOP/s
=============================
...NVIDIA Ampere 및 Turing 아키텍처를 대상으로 F32 누적 및 FP32 입력을 사용하여 순방향 전파(fprop)를 구현하는 하나의 CUDA 코어 컨볼루션 커널을 컴파일하고 실행하려면 아래 cmake 명령줄을 사용하세요.
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc ... $ make cutlass_profiler -j16
하나의 CUDA Core 컨볼루션 커널을 프로파일링하기 위한 명령줄 예시:
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3
=============================
문제 ID: 1
제공자: 커틀라스
OperationKind: 전환2d
작업: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
상태: 성공
확인: 켜짐
처분: 합격
reference_device: 통과됨
인수: --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 --활성화=f32:nhwc --필터=f32:nhwc --출력=f32:nhwc
--conv_mode=교차 --iterator_algorithm=최적화 --alpha=1 --beta=0 --split_k_mode=직렬 --split_k_slices=1
--eq_gemm_provider=none --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4
--warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024
바이트: 2055798784바이트
플롭: 118482796544 플롭
런타임: 7.34266ms
메모리: 260.752GiB/s
수학: 16136.2 GFLOP/s
=============================CUTLASS 커널을 선택적으로 컴파일하는 방법에 대한 더 많은 CMake 예제를 보려면 링크를 따르십시오.
GEMM CMake 예
암시적 GEMM 컨볼루션 CMake 예
CUTLASS 프로파일러에 대한 자세한 내용은 여기에 설명되어 있습니다.
CUTLASS는 NVIDIA Corporation에서 3절 "새" BSD 라이센스에 따라 오픈 소스 소프트웨어로 출시되었습니다.
CUTLASS 개발자 및 기여자의 공식 목록은 CONTRIBUTORS에서 확인할 수 있습니다.
저작권 (c) 2017 - 2024 NVIDIA 기업 및 계열사. 모든 권리 보유. SPDX-라이센스-식별자: BSD-3-Clause
Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: 1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. 2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. 3. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission. THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
