cutlass
CUTLASS
CUTLASS 3.6.0 - ตุลาคม 2024
CUTLASS คือคอลเลกชันของเทมเพลต CUDA C++ สำหรับการใช้งานการคูณเมทริกซ์-เมทริกซ์ประสิทธิภาพสูง (GEMM) และการคำนวณที่เกี่ยวข้องในทุกระดับและทุกขนาดภายใน CUDA โดยรวมเอากลยุทธ์สำหรับการแยกย่อยแบบลำดับชั้นและการเคลื่อนไหวข้อมูลที่คล้ายคลึงกับกลยุทธ์ที่ใช้ในการนำ cuBLAS และ cuDNN ไปใช้ CUTLASS แบ่ง "ส่วนที่เคลื่อนไหว" เหล่านี้ออกเป็นส่วนประกอบซอฟต์แวร์โมดูลาร์ที่นำกลับมาใช้ใหม่ได้ ซึ่งรวบรวมโดยคลาสเทมเพลต C++ พื้นฐานสำหรับระดับต่างๆ ของลำดับชั้นของการวางแนวความคิดแบบขนานสามารถเป็นผู้เชี่ยวชาญและปรับแต่งผ่านขนาดการเรียงแบบกำหนดเอง ชนิดข้อมูล และนโยบายอัลกอริทึมอื่นๆ ความยืดหยุ่นที่เกิดขึ้นทำให้การใช้งานเป็นบล็อคส่วนประกอบภายในเคอร์เนลและแอปพลิเคชันที่กำหนดเองทำได้ง่ายขึ้น
เพื่อรองรับการใช้งานที่หลากหลาย CUTLASS ให้การสนับสนุนอย่างกว้างขวางสำหรับการคำนวณแบบผสมความแม่นยำ โดยให้การเคลื่อนไหวข้อมูลเฉพาะทางและการสะสมนามธรรมแบบทวีคูณสำหรับจุดลอยตัวแบบครึ่งความแม่นยำ (FP16), BFloat16 (BF16), Tensor Float 32 (TF32) จุดลอยตัวที่มีความแม่นยำเดี่ยว (FP32), การจำลอง FP32 ผ่านคำสั่งเทนเซอร์คอร์, ประเภทจุดลอยตัวที่มีความแม่นยำสองเท่า (FP64), ข้อมูลจำนวนเต็ม ประเภท (4b และ 8b) และประเภทข้อมูลไบนารี (1b) CUTLASS สาธิตการดำเนินการคูณเมทริกซ์แบบวาร์ปซิงโครนัสโดยกำหนดเป้าหมายไปที่ Tensor Core ที่ตั้งโปรแกรมได้และมีปริมาณงานสูง ซึ่งใช้งานโดยสถาปัตยกรรม Volta, Turing, Ampere และ Hopper ของ NVIDIA
ดูคู่มือการเริ่มต้นฉบับย่อเพื่อเริ่มต้นอย่างรวดเร็ว
ดูรายการฟังก์ชันสำหรับรายการการดำเนินการที่รองรับในแต่ละระดับของลำดับชั้นของโมเดลการดำเนินการ
CUTLASS 3.0 เปิดตัวไลบรารีหลักใหม่ CuTe เพื่ออธิบายและจัดการเทนเซอร์ของเธรดและข้อมูล CuTe คือชุดของนามธรรมเทมเพลต C++ CUDA สำหรับการกำหนดและดำเนินการบนเค้าโครงเธรดและข้อมูลหลายมิติตามลำดับชั้น CuTe มีออบเจ็กต์ Layout และ Tensor ที่จัดแพคเกจประเภท รูปร่าง พื้นที่หน่วยความจำ และเลย์เอาต์ของข้อมูลอย่างแน่นหนา ในขณะที่ดำเนินการจัดทำดัชนีที่ซับซ้อนสำหรับผู้ใช้ ซึ่งช่วยให้โปรแกรมเมอร์มุ่งเน้นไปที่คำอธิบายเชิงตรรกะของอัลกอริธึมของตนได้ ในขณะที่ CuTe จะทำบัญชีเชิงกลไกให้พวกเขา ด้วยเครื่องมือเหล่านี้ เราจึงสามารถออกแบบ นำไปใช้ และแก้ไขการดำเนินการพีชคณิตเชิงเส้นหนาแน่นทั้งหมดได้อย่างรวดเร็ว
นามธรรมหลักของ CuTe คือโครงร่างหลายมิติที่มีลำดับชั้น ซึ่งสามารถประกอบขึ้นด้วยอาร์เรย์ข้อมูลเพื่อแสดงเทนเซอร์ การแสดงเค้าโครงมีพลังมากพอที่จะแสดงเกือบทุกอย่างที่เราต้องการเพื่อใช้พีชคณิตเชิงเส้นที่มีความหนาแน่นสูงอย่างมีประสิทธิภาพ เค้าโครงยังสามารถรวมและจัดการผ่านองค์ประกอบการทำงาน ซึ่งเราสร้างชุดการดำเนินการทั่วไปขนาดใหญ่ เช่น การปูกระเบื้องและการแบ่งพาร์ติชัน
CUTLASS 3.0 ขึ้นไปใช้ CuTe ตลอดลำดับชั้น GEMM ในเทมเพลต สิ่งนี้ทำให้การออกแบบง่ายขึ้นอย่างมากและปรับปรุงความสามารถในการเขียนโค้ดและความสามารถในการอ่าน เอกสารเพิ่มเติมเฉพาะสำหรับ CuTe สามารถพบได้ในไดเร็กทอรีเอกสารประกอบเฉพาะ
นอกเหนือจาก GEMM แล้ว CUTLASS ยังใช้การบิดที่มีประสิทธิภาพสูงผ่านอัลกอริธึม GEMM โดยนัย GEMM โดยนัยคือการกำหนดการดำเนินการแบบ Convolution ในฐานะ GEMM ดังนั้นจึงใช้ประโยชน์จากไปป์ไลน์ GEMM แบบโมดูลาร์ของ CUTLASS สิ่งนี้ทำให้ CUTLASS สามารถสร้างการโน้มน้าวใจโดยการนำส่วนประกอบ GEMM ที่ได้รับการปรับปรุงประสิทธิภาพระดับสูงกลับมาใช้ใหม่
CUTLASS 3.6.0 เป็นการอัพเดตของ CUTLASS โดยเพิ่ม:
ถังมีโครงสร้างเบาบาง GEMM
FP16
เอฟพี8
INT8
TF32
การรีแฟคเตอร์ให้กับ kernel::ConvUniversal API เพื่อให้สอดคล้องกับ gemm::GemmUniversal ตอนนี้ Convolution API 3.x ไม่ถือเป็น API เบต้าอีกต่อไป
GEMM อินพุตแบบผสมที่ได้รับการปรับปรุงและการใช้งานตารางการค้นหาสำหรับโหมดเฉพาะสเกล INT4 x FP8
โหนด EVT สำหรับการเลือก Top-K และตัวอย่าง softmax และ GEMM ที่ใช้สิ่งเหล่านั้น
Programmatic Dependent Launch (PDL) ที่ใช้ประโยชน์จากฟีเจอร์ Hopper ใหม่เพื่อเร่งความเร็วเคอร์เนลสองเคอร์เนลแบบ back-to-back และเอกสารประกอบที่เกี่ยวข้อง
เครื่องมือแก้ไขข้อบกพร่องใหม่ synclog สำหรับการทิ้งกิจกรรมการซิงโครไนซ์ทั้งหมดจากภายในเคอร์เนลไปยังไฟล์ โปรดดูเอกสารประกอบของ synclog สำหรับรายละเอียด
บทส่งท้ายที่เปิดใช้งาน TMA ใหม่สำหรับ GEMM แบบกลุ่มที่นำมาซึ่งการปรับปรุงประสิทธิภาพที่สำคัญ รวมถึงการสนับสนุน EVT
บทส่งท้ายอาร์เรย์พอยน์เตอร์ที่เปิดใช้งาน SIMT
กำหนดการเคอร์เนล Ping-Pong ใหม่สำหรับ Grouped GEMM และการเพิ่มประสิทธิภาพอื่น ๆ
กลยุทธ์การสร้างอินสแตนซ์ใหม่สำหรับเคอร์เนลตัวสร้างโปรไฟล์ CUTLASS พร้อมด้วยเอกสารที่ได้รับการปรับปรุงสำหรับระดับการสร้างอินสแตนซ์ในตัวสร้างโปรไฟล์ CUTLASS
การสนับสนุนฮาร์ดแวร์ใหม่สำหรับการเปรียบเทียบและการคำนวณของ cutlass::bfloat16_t
แก้ไขการใช้ isnan บน Windows สำหรับ half_t
ข้อกำหนดขั้นต่ำ:
สถาปัตยกรรม: โวลตา
คอมไพเลอร์: ต้องรองรับอย่างน้อย C ++ 17
ชุดเครื่องมือ CUDA เวอร์ชัน: 11.4
เริ่มต้นจาก CUTLASS 3.0 CUTLASS ได้ลบการสนับสนุนต่อไปนี้:
สถาปัตยกรรม GPU ของ Maxwell และ Pascal
อูบุนตู 16.04
CUDA 10.2
ภาษา C++ เวอร์ชันต่ำกว่า 17
ดู CHANGELOG สำหรับรายการรายละเอียดของการเผยแพร่และการอัพเดต
พื้นฐาน CUTLASS มีประสิทธิภาพมาก เมื่อใช้เพื่อสร้างเคอร์เนล GEMM ทั่วทั้งอุปกรณ์ เคอร์เนลจะแสดงประสิทธิภาพสูงสุดที่เทียบได้กับ cuBLAS สำหรับการคำนวณ GEMM แบบสเกลาร์ รูปด้านบนแสดงการปรับปรุงประสิทธิภาพ CUTLASS อย่างต่อเนื่องบน NVIDIA H100 (สถาปัตยกรรม NVIDIA Hopper) ตั้งแต่ CUTLASS 3.1 CUTLASS 3.5.1 ได้รับการคอมไพล์ด้วยชุดเครื่องมือ CUDA 12.5u1 การดำเนินการ Tensor Core ดำเนินการโดยใช้คำสั่ง mma และ wgmma ของ CUDA
เมื่อใช้ Building Block ของ CUTLASS เพื่อสร้างเคอร์เนล Gemm โดยนัย (Fprop, Dgrad และ Wgrad) ทั่วทั้งอุปกรณ์ ประสิทธิภาพของ CUTLASS ยังเทียบเคียงได้กับ cuDNN เมื่อรันเลเยอร์ Resnet-50 บน NVIDIA A100 ดังแสดงในรูปด้านบน การดำเนินการ Tensor Core ดำเนินการโดยใช้คำสั่ง mma ของ CUDA
CUTLASS ต้องการคอมไพลเลอร์โฮสต์ C++17 และทำงานได้ดีที่สุดเมื่อสร้างด้วย CUDA 12.4 Toolkit นอกจากนี้ยังเข้ากันได้กับ 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 ต่อไปนี้ และคาดว่าจะมีประสิทธิภาพบน NVIDIA GPU ที่ใช้สถาปัตยกรรม Volta, Turing, Ampere, Ada และ Hopper
| จีพียู | ความสามารถในการคำนวณ CUDA | ชุดเครื่องมือ CUDA ขั้นต่ำที่ CUTLASS-3 ต้องการ |
|---|---|---|
| NVIDIA V100 เทนเซอร์คอร์ 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 เทนเซอร์คอร์ 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 เทนเซอร์คอร์ GPU | 9.0 | 11.8 |
โดยทั่วไป รหัส PTX ที่สร้างขึ้นสำหรับสถาปัตยกรรมเป้าหมายเดียวสามารถรันบนสถาปัตยกรรมในอนาคตได้ (เช่น สามารถใช้งานร่วมกับการส่งต่อได้) อย่างไรก็ตาม CUDA 12.0 ได้นำเสนอแนวคิดของ "คุณลักษณะที่เร่งด้วยสถาปัตยกรรม" ซึ่ง PTX ไม่มีการรับประกันความเข้ากันได้ล่วงหน้า คำสั่ง Hopper PTX หลายคำสั่งอยู่ภายใต้หมวดหมู่ของคุณสมบัติเร่งสถาปัตยกรรมนี้ และดังนั้นจึงต้องใช้สถาปัตยกรรมเป้าหมาย sm_90a (โปรดทราบว่ามี "a" ต่อท้าย) สำหรับรายละเอียดเพิ่มเติมเกี่ยวกับสิ่งนี้และคำแนะนำในการเร่งสถาปัตยกรรมอื่นๆ โปรดดูเอกสารประกอบ CUDA
ข้อมูลสถาปัตยกรรมเป้าหมายจะถูกส่งต่อไปยัง CUTLASS ผ่านทางแฟล็ก cmake CUTLASS_NVCC_ARCHS เพื่อเพิ่มประสิทธิภาพสูงสุดบน Hopper GH100 ผู้ใช้จะต้องสร้าง CUTLASS โดยมี 90a เป็นสถาปัตยกรรมเป้าหมาย หากผู้ใช้สร้างเคอร์เนลโดยไม่ได้ตั้งใจซึ่งใช้คุณสมบัติ SM90a (เช่นคำสั่ง Hopper Tensor Core) โดยใช้เป้าหมาย SM90 (โปรดสังเกตว่าไม่มี "a") ด้วย CUDA Toolkit 12 หรือ 11.8 เคอร์เนลคาดว่าจะล้มเหลวพร้อมกับรันไทม์ ข้อผิดพลาด.
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
โปรดดูเอกสารประกอบการทำงานสำหรับรายละเอียดว่าเคอร์เนลใดต้องการสถาปัตยกรรมเป้าหมายแบบใด
CUTLASS มีคำอธิบายอยู่ในเอกสารต่อไปนี้และเอกสารประกอบของ Doxygen ที่แนบมาด้วย
คู่มือเริ่มต้นใช้งานฉบับย่อ - สร้างและรัน CUTLASS
ฟังก์ชั่น - สรุปฟังก์ชั่นที่มีอยู่ใน CUTLASS
GEMM ที่มีประสิทธิภาพใน CUDA - อธิบายว่าเคอร์เนล GEMM สามารถนำไปใช้อย่างมีประสิทธิภาพใน CUDA ได้อย่างไร
การออกแบบ 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++
Convolution GEMM โดยนัย - อธิบายการบิด 2 มิติและ 3 มิติใน CUTLASS
Code Organisation - อธิบายองค์กรและเนื้อหาของโครงการ CUTLASS
คำศัพท์เฉพาะทาง - อธิบายคำศัพท์ที่ใช้ในโค้ด
แนวทางการเขียนโปรแกรม - แนวทางสำหรับการเขียน CUDA C++ สมัยใหม่ที่มีประสิทธิภาพ
ประเภทพื้นฐาน - อธิบายคลาส C++ พื้นฐานที่ใช้ใน CUTLASS เพื่อแสดงปริมาณตัวเลขและอาร์เรย์
เลย์เอาต์ - อธิบายเลย์เอาต์ของเมทริกซ์และเทนเซอร์ในหน่วยความจำ
ตัววนซ้ำไทล์ - อธิบายแนวคิด C++ สำหรับการวนซ้ำไทล์ของเมทริกซ์ในหน่วยความจำ
CUTLASS Profiler - แอปพลิเคชันสร้างโปรไฟล์ที่ขับเคลื่อนด้วยบรรทัดคำสั่ง
CUTLASS Utilities - เทมเพลตเพิ่มเติมที่ใช้เพื่ออำนวยความสะดวกในการพัฒนาอย่างรวดเร็ว
การเปิดเคอร์เนลขึ้นอยู่กับ - อธิบายคุณลักษณะใหม่ใน Hopper ซึ่งอนุญาตให้มีการทับซ้อนกันของเคอร์เนลในสตรีมเดียวกัน และวิธีการใช้ใน CUTLASS
นอกจากนี้เรายังได้อธิบายโครงสร้างของ GEMM ที่มีประสิทธิภาพในการพูดคุยของเราที่การประชุมเทคโนโลยี GPU ปี 2018
CUTLASS: ซอฟต์แวร์ดั้งเดิมสำหรับพีชคณิตเชิงเส้นหนาแน่นในทุกระดับและทุกระดับภายใน CUDA
การพัฒนาเคอร์เนล CUDA เพื่อผลักดัน Tensor Cores จนถึงขีดจำกัดที่แน่นอนบน NVIDIA A100
เร่งการ Convolution ด้วย Tensor Cores ใน CUTLASS
เร่งการไล่ระดับข้อมูลย้อนหลังโดยการเพิ่มการใช้งาน Tensor Core ใน CUTLASS
CUTLASS: Python API, การปรับปรุง และ NVIDIA Hopper
CUTLASS เป็นไลบรารีเทมเพลตเฉพาะส่วนหัวเท่านั้น และไม่จำเป็นต้องสร้างเพื่อใช้ในโปรเจ็กต์อื่น แอปพลิเคชันไคลเอ็นต์ควรกำหนดเป้าหมายไดเรกทอรี include/ ของ CUTLASS ในเส้นทางรวม
การทดสอบหน่วย ตัวอย่าง และยูทิลิตี้ CUTLASS สามารถสร้างได้ด้วย CMake เวอร์ชันขั้นต่ำของ CMake มีระบุไว้ในคู่มือ Quickstart ตรวจสอบให้แน่ใจว่าตัวแปรสภาพแวดล้อม CUDACXX ชี้ไปที่ NVCC ในชุดเครื่องมือ CUDA ที่ติดตั้งบนระบบของคุณ
$ ส่งออก CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc สร้างไดเร็กทอรี build ภายในโปรเจ็กต์ CUTLASS จากนั้นเรียกใช้ CMake ตามค่าเริ่มต้น CUTLASS จะสร้างเคอร์เนลสำหรับสถาปัตยกรรม CUDA เวอร์ชัน 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9 และ 9.0 เพื่อลดเวลาในการคอมไพล์ คุณสามารถระบุสถาปัตยกรรมเพื่อสร้าง CUTLASS ได้โดยการเปลี่ยนการตั้งค่าการกำหนดค่า CMake CUTLASS_NVCC_ARCHS
$ mkdir build && cd build $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # คอมไพล์สำหรับสถาปัตยกรรม Ampere ของ NVIDIA
จากไดเร็กทอรี build/ ให้คอมไพล์และรันการทดสอบหน่วย CUTLASS โดยการสร้างเป้าหมาย test_unit ด้วย make
การทดสอบหน่วยถูกจัดระเบียบเป็นไบนารีหลายตัวที่สะท้อนเนมสเปซระดับบนสุดของ CUTLASS และอาจดำเนินการแบบขนานผ่านอาร์กิวเมนต์บรรทัดคำสั่ง make's -j
$ สร้าง test_unit -j - - - [----------] สภาพแวดล้อมการทดสอบทั่วโลกรื้อลง [==========] ดำเนินการทดสอบ 946 รายการจากกรณีทดสอบ 57 กรณี (รวม 1,0812 มิลลิวินาที) [ ผ่าน ] 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/ ประกอบด้วยการทดสอบหน่วยที่ใช้กับ Google Test ซึ่งสาธิตการใช้งานพื้นฐานของส่วนประกอบ Core API และการทดสอบการคำนวณ CUTLASS GEMM ที่สมบูรณ์
คำแนะนำสำหรับการสร้างและรันการทดสอบหน่วยมีอธิบายไว้ในคู่มือ Quickstart
ไดเร็กทอรี tools/profiler/ มียูทิลิตีบรรทัดคำสั่งสำหรับการเรียกใช้เคอร์เนล GEMM แต่ละอัน สามารถสร้างได้ดังนี้:
$ ทำ cutlass_profiler -j16
ตามค่าเริ่มต้น จะมีการสร้างอินสแตนซ์ขนาดไทล์เดียวเท่านั้นสำหรับแต่ละประเภทข้อมูล คำแนะนำทางคณิตศาสตร์ และเค้าโครง หากต้องการสร้างอินสแตนซ์ทั้งหมด ให้ตั้งค่าตัวแปรสภาพแวดล้อมต่อไปนี้เมื่อเรียกใช้ CMake จากไดเร็กทอรี build/ ว่าง ระวัง สิ่งนี้ส่งผลให้เกิดเมล็ด นับหมื่น และใช้เวลาในการสร้างนาน นอกจากนี้ยังจะส่งผลให้มีขนาดไบนารี่ขนาดใหญ่และในบางแพลตฟอร์มตัวเชื่อมโยงล้มเหลวในการสร้างไลบรารี ดังนั้นจึงขอแนะนำอย่างยิ่งให้สร้างเฉพาะชุดย่อยของเมล็ดพืชตามที่แสดงไว้ในส่วนย่อยด้านล่าง
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=ทั้งหมด - $ ทำ cutlass_profiler -j16
ในการรวบรวมเคอร์เนลหนึ่งเคอร์เนลหรือชุดเคอร์เนลชุดเล็ก ๆ อย่างเคร่งครัด อาจใช้รายการชื่อเคอร์เนลที่คั่นด้วยเครื่องหมายจุลภาคพร้อมอักขระไวด์การ์ดเพื่อลดชุดเคอร์เนล ตัวอย่างต่อไปนี้แสดงการสร้างเคอร์เนลหนึ่งหรือชุดย่อยสำหรับสถาปัตยกรรม NVIDIA Ampere และ Turing:
หากต้องการรวบรวมชุดย่อยของเคอร์เนล Tensor Core GEMM ที่มีการสะสม FP32 และอินพุต FP16 ที่กำหนดเป้าหมายสถาปัตยกรรม NVIDIA Ampere และ Turing ให้ใช้บรรทัดคำสั่ง cmake ด้านล่าง:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8 - $ ทำ cutlass_profiler -j16
ตัวอย่างบรรทัดคำสั่งสำหรับการจัดทำโปรไฟล์ชุดย่อยของเคอร์เนล Tensor Core GEMM มีดังนี้:
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
-
-
รหัสปัญหา: 1
ผู้ให้บริการ: CUTLASS
OperationKind: เจมม์
การทำงาน: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
สถานะ: สำเร็จ
การยืนยัน: เปิด
นิสัย: ผ่านแล้ว
Reference_device: ผ่านแล้ว
คิวบลาส: ผ่านแล้ว
อาร์กิวเมนต์: --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.55948 มิลลิวินาที
หน่วยความจำ: 70.7616 GiB/s
คณิตศาสตร์: 74378.8 GFLOP/s
-
-หากต้องการคอมไพล์เคอร์เนล SGEMM หนึ่งตัวที่กำหนดเป้าหมายสถาปัตยกรรม NVIDIA Ampere และ Turing ให้ใช้บรรทัดคำสั่ง cmake ด้านล่าง:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1 - $ ทำ cutlass_profiler -j16
ตัวอย่างบรรทัดคำสั่งสำหรับการทำโปรไฟล์เคอร์เนล SGEMM CUDA เดี่ยวมีดังนี้:
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
-
รหัสปัญหา: 1
ผู้ให้บริการ: CUTLASS
OperationKind: เจมม์
การทำงาน: cutlass_simt_sgemm_128x128_8x2_nn_align1
สถานะ: สำเร็จ
การยืนยัน: เปิด
นิสัย: ผ่านแล้ว
คิวบลาส: ผ่านแล้ว
อาร์กิวเมนต์: --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.73655 มิลลิวินาที
หน่วยความจำ: 24.934 GiB/s
คณิตศาสตร์: 17218.4 GFLOP/s
-หากต้องการคอมไพล์ชุดย่อยของเคอร์เนล Tensor core convolution ที่ใช้การส่งต่อไปข้างหน้า (fprop) ด้วยการสะสม FP32 และอินพุต FP16 ที่กำหนดเป้าหมายสถาปัตยกรรม NVIDIA Ampere และ Turing ให้ใช้บรรทัดคำสั่ง cmake ด้านล่าง:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16 - $ ทำ cutlass_profiler -j16
ตัวอย่างบรรทัดคำสั่งสำหรับการจัดทำโปรไฟล์ชุดย่อยของเคอร์เนล Convolution 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
-
-
รหัสปัญหา: 1
ผู้ให้บริการ: CUTLASS
ประเภทการดำเนินการ: Conv2d
การทำงาน: 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 --Activation=f16:nhwc --Filter=f16:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=เพิ่มประสิทธิภาพ --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
ไบต์: 1130659840 ไบต์
ฟลอป: 118482796544 ฟลอป
รันไทม์: 0.711496 มิลลิวินาที
หน่วยความจำ: 1479.99 GiB/s
คณิตศาสตร์: 166526 GFLOP/s
-
-หากต้องการคอมไพล์และรันเคอร์เนล Convolution CUDA Core หนึ่งตัวที่ใช้การส่งต่อไปข้างหน้า (fprop) ด้วยการสะสม F32 และอินพุต FP32 ที่กำหนดเป้าหมายสถาปัตยกรรม NVIDIA Ampere และ Turing ให้ใช้บรรทัดคำสั่ง cmake ด้านล่าง:
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc - $ ทำ cutlass_profiler -j16
บรรทัดคำสั่งตัวอย่างสำหรับการสร้างโปรไฟล์เคอร์เนล 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
-
รหัสปัญหา: 1
ผู้ให้บริการ: CUTLASS
OperationKind: Conv2d
การทำงาน: 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 --Activation=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=เพิ่มประสิทธิภาพ --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
ไบต์: 2055798784 ไบต์
ฟลอป: 118482796544 ฟลอป
รันไทม์: 7.34266 มิลลิวินาที
หน่วยความจำ: 260.752 GiB/s
คณิตศาสตร์: 16136.2 GFLOP/s
-โปรดไปตามลิงก์เพื่อดูตัวอย่าง CMake เพิ่มเติมเกี่ยวกับการคอมไพล์เคอร์เนล CUTLASS แบบคัดเลือก:
ตัวอย่าง GEMM CMake
ตัวอย่าง CMake ของ GEMM Convolution โดยนัย
รายละเอียดเพิ่มเติมเกี่ยวกับ CUTLASS Profiler มีอธิบายไว้ที่นี่
CUTLASS เปิดตัวโดย NVIDIA Corporation ในรูปแบบซอฟต์แวร์โอเพ่นซอร์สภายใต้ใบอนุญาต BSD "ใหม่" 3 ข้อ
รายชื่อนักพัฒนาและผู้มีส่วนร่วมอย่างเป็นทางการของ CUTLASS มีอยู่ที่นี่: CONTRIBUTORS
ลิขสิทธิ์ (c) 2017 - 2024 NVIDIA CORPORATION และบริษัทในเครือ สงวนลิขสิทธิ์. 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.
