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 عالية الإنتاجية القابلة للبرمجة والتي تنفذها معماريات Volta وTuring وAmpere وHopper من NVIDIA.
راجع دليل البدء السريع للبدء بسرعة.
راجع قائمة الوظائف للحصول على قائمة العمليات المدعومة في كل مستوى من مستويات التسلسل الهرمي لنموذج التنفيذ.
قدم CUTLASS 3.0 مكتبة أساسية جديدة، CuTe، لوصف ومعالجة موترات الخيوط والبيانات. CuTe عبارة عن مجموعة من تجريدات قوالب C++ CUDA لتعريف وتشغيل تخطيطات متعددة الأبعاد هرمية للخيوط والبيانات. يوفر برنامج CuTe كائنات Layout و Tensor التي تعمل على حزم نوع البيانات وشكلها ومساحة الذاكرة وتخطيطها بشكل مضغوط، أثناء إجراء الفهرسة المعقدة للمستخدم. يتيح ذلك للمبرمجين التركيز على الأوصاف المنطقية لخوارزمياتهم بينما يقوم برنامج CuTe بمسك الدفاتر الميكانيكية لهم. باستخدام هذه الأدوات، يمكننا تصميم وتنفيذ وتعديل جميع عمليات الجبر الخطي الكثيفة بسرعة.
التجريدات الأساسية لـ CuTe عبارة عن تخطيطات متعددة الأبعاد بشكل هرمي والتي يمكن تكوينها باستخدام صفائف البيانات لتمثيل الموترات. يعد تمثيل المخططات قويًا بما يكفي لتمثيل كل ما نحتاجه تقريبًا لتنفيذ الجبر الخطي الكثيف الفعال. يمكن أيضًا دمج التخطيطات ومعالجتها من خلال التركيب الوظيفي، والذي نبني عليه مجموعة كبيرة من العمليات الشائعة مثل التبليط والتقسيم.
يعتمد CUTLASS 3.0 وما بعده برنامج CuTe عبر التسلسل الهرمي لـ GEMM في قوالبه. يؤدي هذا إلى تبسيط التصميم إلى حد كبير وتحسين قابلية تكوين التعليمات البرمجية وسهولة قراءتها. يمكن العثور على المزيد من الوثائق الخاصة بـ CuTe في دليل الوثائق المخصص لها.
بالإضافة إلى GEMMs، ينفذ CUTLASS التفافًا عالي الأداء عبر خوارزمية GEMM الضمنية. GEMM الضمنية هي صياغة عملية الالتفاف باعتبارها GEMM وبالتالي الاستفادة من خط أنابيب GEMM المعياري الخاص بـ CUTLASS. يتيح ذلك لـ CUTLASS إنشاء تلافيفات من خلال إعادة استخدام مكونات GEMM المحسنة للغاية.
CUTLASS 3.6.0 هو تحديث لـ CUTLASS يضيف:
هوبر منظم متناثر GEMM.
FP16
FP8
إنت8
TF32
إعادة بناء لـ CUTLASS 3.x convolution kernel::ConvUniversal API لجعلها متوافقة مع gemm::GemmUniversal . الآن لم تعد واجهة برمجة تطبيقات 3.x تعتبر واجهة برمجة تطبيقات تجريبية.
مدخلات مختلطة محسنة GEMM وتنفيذ جدول بحث لوضع مقياس INT4 x FP8 فقط.
عقد EVT لاختيار Top-K ومثال softmax وGEMM باستخدام تلك العقد.
الإطلاق البرمجي المعتمد (PDL) الذي يستفيد من ميزة Hopper الجديدة لتسريع نواتين متتاليتين والوثائق المقابلة لها.
أداة تصحيح جديدة، synclog، للتخلص من كافة أحداث المزامنة من داخل النواة إلى ملف. يرجى الاطلاع على وثائق Synlog للحصول على التفاصيل.
خاتمة جديدة ممكّنة لـ TMA لـ GEMM المجمعة والتي تحقق تحسينًا كبيرًا في الأداء، بالإضافة إلى دعم EVT الخاص بها.
خاتمة صفيف مؤشر تمكين SIMT.
جدول Ping-Pong kernel جديد لـ GEMM المجمعة وبعض التحسينات الأخرى.
استراتيجية جديدة لإنشاء مثيل لنواة ملف تعريف CUTLASS بالإضافة إلى توثيق محسّن لمستوى إنشاء مثيل في ملف تعريف CUTLASS.
دعم جديد للأجهزة لإجراء مقارنات وحسابات لـ cutlass::bfloat16_t
تم إصلاح استخدام isnan على نظام Windows لـ half_t .
الحد الأدنى من المتطلبات:
العمارة: فولتا
المترجم: يجب أن يدعم C++ 17 على الأقل
إصدار مجموعة أدوات CUDA: 11.4
بدءًا من الإصدار 3.0 من CUTLASS، أزال CUTLASS الدعم لما يلي:
معماريات ماكسويل وباسكال GPU
أوبونتو 16.04
كودا 10.2
إصدارات لغة C++ أقل من 17.
راجع سجل التغيير للحصول على قائمة مفصلة بالإصدارات والتحديثات.
تعتبر البدائيات CUTLASS فعالة للغاية. عند استخدامها لإنشاء نواة GEMM على مستوى الجهاز، فإنها تظهر أعلى أداء يمكن مقارنته بـ cuBLAS لحسابات GEMM العددية. يوضح الشكل أعلاه التحسينات المستمرة في أداء CUTLASS على NVIDIA H100 (بنية NVIDIA Hopper) منذ CUTLASS 3.1. تم تجميع CUTLASS 3.5.1 باستخدام مجموعة أدوات CUDA 12.5u1. يتم تنفيذ عمليات Tensor Core باستخدام تعليمات CUDA mma وwgmma.
عند استخدام كتل بناء CUTLASS لإنشاء نواة Gemm (Fprop وDgrad وWgrad) الضمنية على مستوى الجهاز، يكون أداء CUTLASS أيضًا مشابهًا لأداء cuDNN عند تشغيل طبقات Resnet-50 على NVIDIA A100 كما هو موضح في الشكل أعلاه. يتم تنفيذ عمليات Tensor Core باستخدام تعليمات CUDA's 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 | دول مجلس التعاون الخليجي 7.5.0 |
| أوبونتو 20.04 | دول مجلس التعاون الخليجي 10.3.0 |
| أوبونتو 22.04 | دول مجلس التعاون الخليجي 11.2.0 |
| أوبونتو 22.04 | رنة 10.0.0 |
| أوبونتو 22.04 | رنة 14.0.6 |
| أوبونتو 22.04 | رنة 17.0.6 |
| ويندوز 10.0 | فيجوال ستوديو 2019 v16.11.27 |
ملاحظة: لقد عرف الإصدار 8.5.0 من مجلس التعاون الخليجي تراجعات فيما يتعلق بتعبيرات الطية والمشغلين الزائدين. يوصى باستخدام الإصدار 7.5.0 من مجلس التعاون الخليجي أو (المفضل) مجلس التعاون الخليجي >= 9.
يعمل CUTLASS بنجاح على وحدات معالجة الرسوميات NVIDIA التالية، ومن المتوقع أن يكون فعالاً على وحدات معالجة الرسوميات NVIDIA المستندة إلى بنية Volta وTuring وAmpere وAda وHopper.
| GPU | القدرة على حساب CUDA | الحد الأدنى لمجموعة أدوات CUDA المطلوبة بواسطة CUTLASS-3 |
|---|---|---|
| وحدة معالجة الرسوميات NVIDIA V100 Tensor Core | 7.0 | 11.4 |
| نفيديا تيتانV | 7.0 | 11.4 |
| نفيديا جي فورس آر تي إكس 2080 تي آي، 2080، 2070 | 7.5 | 11.4 |
| نفيديا T4 | 7.5 | 11.4 |
| وحدة معالجة الرسوميات NVIDIA A100 Tensor Core | 8.0 | 11.4 |
| نفيديا A10 | 8.6 | 11.4 |
| نفيديا جي فورس آر تي إكس 3090 | 8.6 | 11.4 |
| نفيديا جي فورس آر تي إكس 4090 | 8.9 | 11.8 |
| نفيديا L40 | 8.9 | 11.8 |
| وحدة معالجة الرسوميات NVIDIA H100 Tensor Core | 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 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++
التفاف GEMM الضمني - يصف الالتفاف ثنائي الأبعاد وثلاثي الأبعاد في CUTLASS
تنظيم الكود - يصف تنظيم ومحتويات مشروع CUTLASS
المصطلحات - تصف المصطلحات المستخدمة في الكود
إرشادات البرمجة - إرشادات لكتابة CUDA C++ الحديثة بكفاءة
الأنواع الأساسية - تصف فئات C++ الأساسية المستخدمة في CUTLASS لتمثيل الكميات والمصفوفات الرقمية
التخطيطات - تصف تخطيطات المصفوفات والموترات في الذاكرة
Tile Iterators - يصف مفاهيم C++ للتكرار على مربعات المصفوفات في الذاكرة
CUTLASS Profiler - تطبيق إنشاء ملفات تعريف يعتمد على سطر الأوامر
أدوات CUTLASS - قوالب إضافية تستخدم لتسهيل التطوير السريع
إطلاق النواة التابعة - يصف ميزة جديدة في Hopper والتي تسمح بتداخل النواة التابعة في نفس الدفق، وكيفية استخدامها في CUTLASS.
لقد وصفنا أيضًا هيكل GEMM الفعال في حديثنا في مؤتمر تكنولوجيا GPU 2018.
CUTLASS: البدائيات البرمجية للجبر الخطي الكثيف على جميع المستويات والمقاييس داخل CUDA
تطوير نواة CUDA لدفع النوى الموترة إلى الحد المطلق على NVIDIA A100
تسريع الالتواء باستخدام نوى Tensor في CUTLASS
تسريع التدرج العكسي للبيانات عن طريق زيادة استخدام Tensor Core في CUTLASS
CUTLASS: واجهة برمجة تطبيقات Python والتحسينات وNVIDIA Hopper
CUTLASS هي مكتبة قوالب للترويسة فقط ولا تحتاج إلى إنشاء لاستخدامها في مشاريع أخرى. يجب أن تستهدف تطبيقات العميل دليل include/ بـ CUTLASS في مسارات التضمين الخاصة بها.
يمكن إنشاء اختبارات وحدة CUTLASS والأمثلة والأدوات المساعدة باستخدام CMake. يتوفر الحد الأدنى من إصدار CMake في دليل Quickstart. تأكد من أن متغير البيئة CUDACXX يشير إلى NVCC في مجموعة أدوات CUDA المثبتة على نظامك.
$ تصدير 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. لتقليل وقت الترجمة، يمكنك تحديد البنيات لإنشاء CUTLASS لها عن طريق تغيير إعداد تكوين CMake CUTLASS_NVCC_ARCHS .
$ mkdir build && cd build $ cmake .. -DCUTLASS_NVCC_ARCHS=80 # تجميع لبنية Ampere Architecture من NVIDIA
من الدليل build/ ، قم بتجميع وتشغيل اختبارات وحدة CUTLASS عن طريق إنشاء test_unit الهدف باستخدام make.
يتم تنظيم اختبارات الوحدة كعدة ثنائيات تعكس مساحات الأسماء ذات المستوى الأعلى لـ CUTLASS، ويمكن تنفيذها بالتوازي عبر وسيطة سطر الأوامر make's -j .
$ اصنع test_unit -j ... ... ... [----------] هدم بيئة الاختبار العالمية [==========] تم إجراء 946 اختبارًا من 57 حالة اختبار. (إجمالي 10812 مللي ثانية) [تم النجاح] 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.
تم توضيح تعليمات إنشاء اختبارات الوحدة وتشغيلها في دليل البدء السريع.
يحتوي الدليل tools/profiler/ على أداة مساعدة لسطر الأوامر لتشغيل كل نواة GEMM. يمكن بناؤها على النحو التالي:
$ اصنع Cutlass_profiler -j16
افتراضيًا، يتم إنشاء حجم تجانب واحد فقط لكل نوع بيانات وتعليمات رياضية وتخطيط. لإنشاء مثيل لكل شيء، قم بتعيين متغير البيئة التالي عند تشغيل CMake من دليل build/ فارغ. احذر، يؤدي هذا إلى عشرات الآلاف من النوى وأوقات بناء طويلة. قد يؤدي هذا أيضًا إلى حجم ثنائي كبير وفشل رابط بعض الأنظمة الأساسية في إنشاء المكتبة. لذلك، يوصى بشدة بإنشاء مجموعة فرعية فقط من النوى كما هو موضح في القسم الفرعي أدناه.
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all ... $ اصنع 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
نوع العملية: جوهرة
العملية:cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
الحالة: النجاح
التحقق: تشغيل
التصرف: مرت
مرجع_الجهاز: تم اجتيازه
كوبلاس: مرت
الوسيطات: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:عمود --B=f16:صف --C=f32:عمود --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 جيجا بايت/ثانية
الرياضيات: 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
نوع العملية: جوهرة
العملية:cutlass_simt_sgemm_128x128_8x2_nn_align1
الحالة: النجاح
التحقق: تشغيل
التصرف: مرت
كوبلاس: مرت
الوسيطات: --m=3456 --n=4096 --k=4096 --A=f32:عمود --B=f32:عمود --C=f32:عمود --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 جيجا بايت/ثانية
الرياضيات: 17218.4 GFLOP/s
======================لتجميع مجموعة فرعية من نواة الالتواء الأساسية Tensor التي تنفذ الانتشار الأمامي (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
مثال سطر الأوامر لتوصيف مجموعة فرعية من حبات الالتواء 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
الحالة: النجاح
التحقق: تشغيل
التصرف: مرت
مرجع_الجهاز: تم اجتيازه
الوسائط: --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 --Filter=f16:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=optimized --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 جيجا بايت/ثانية
الرياضيات: 166526 GFLOP/s
======================
...لتجميع وتشغيل نواة التفاف 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
مثال لسطر الأوامر لتوصيف نواة التفاف 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
نوع العملية: conv2d
العملية:cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
الحالة: النجاح
التحقق: تشغيل
التصرف: مرت
مرجع_الجهاز: تم اجتيازه
الوسائط: --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 --Filter=f32:nhwc --Output=f32:nhwc
--conv_mode=cross --iterator_algorithm=optimized --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 جيجا بايت/ثانية
الرياضيات: 16136.2 GFLOP/s
======================يرجى اتباع الروابط للحصول على المزيد من أمثلة CMake حول التجميع الانتقائي لنواة CUTLASS:
GEMM C اصنع أمثلة
أمثلة ضمنية لالتفاف GEMM CMake
مزيد من التفاصيل حول ملف تعريف CUTLASS موصوفة هنا.
تم إصدار CUTLASS بواسطة شركة NVIDIA كبرنامج مفتوح المصدر بموجب ترخيص BSD "الجديد" المكون من 3 فقرات.
القائمة الرسمية لمطوري ومساهمي CUTLASS متاحة هنا: المساهمين.
حقوق الطبع والنشر (ج) 2017 - 2024 مملوكة لشركة NVIDIA والشركات التابعة لها. جميع الحقوق محفوظة. معرف ترخيص SPDX: 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.
