بهینه سازی عملکرد ضرب در ماتریس ورودی مخلوط – وبلاگ تحقیقاتی گوگل

ارسال شده توسط Manish Gupta، مهندس نرم افزار کارکنان، Google Research

فناوری‌های مبتنی بر هوش مصنوعی با پتانسیل افزایش دسترسی ما به دانش و افزایش بهره‌وری کلی ما، خود را در تار و پود کارهای روزمره ما می‌پیوندند. ستون فقرات این برنامه ها در مدل های زبان بزرگ (LLM) نهفته است. LLM ها حافظه فشرده هستند و معمولاً به شتاب دهنده های سخت افزاری تخصصی نیاز دارند تا ده ها اگزافلاپس قدرت محاسباتی را به طور موثر ارائه کنند. این پست وبلاگ نشان می‌دهد که چگونه می‌توانیم با استفاده مؤثرتر از حافظه شروع به پرداختن به چالش‌های محاسباتی کنیم.

بخش عمده ای از حافظه و محاسبات یک LLM توسط وزن ها در عملیات ضرب ماتریس مصرف می شود. با استفاده از باریکتر انواع داده ها مصرف حافظه را کاهش می دهد. به عنوان مثال، ذخیره وزن ها در نوع داده 8 بیتی عدد صحیح (یعنی U8 یا S8)، ردپای حافظه را 4× نسبت به تک دقیق (F32) و 2× نسبت به نیمه دقیق (F16) یا bfloat16 (BF16) کاهش می دهد. ). علاوه بر این، کار قبلی نشان داده است که LLM ضرب ماتریس را با مدل‌سازی می‌کند وزنه ها در S8 و ورودی در F16 (حفظ دقت بالاتر ورودی کاربر) روشی موثر برای افزایش کارایی با مبادلات قابل قبول در دقت است. این تکنیک به عنوان شناخته شده است کوانتیزاسیون فقط وزنی و نیاز به اجرای کارآمد ضرب ماتریس با ورودی های مختلطبه عنوان مثال، ورودی نیمه دقیق ضرب شده با عدد صحیح 8 بیتی. شتاب‌دهنده‌های سخت‌افزاری، از جمله پردازنده‌های گرافیکی، مجموعه‌ای ثابت از انواع داده‌ها را پشتیبانی می‌کنند، و بنابراین، ضرب ماتریس ورودی مختلط به تبدیل‌های نرم‌افزاری برای نقشه‌برداری به عملیات سخت‌افزاری نیاز دارد.

برای این منظور، در این وبلاگ بر روی نقشه‌برداری ضرب ماتریس ورودی مختلط در معماری NVIDIA Ampere تمرکز می‌کنیم. ما تکنیک‌های نرم‌افزاری را ارائه می‌کنیم که تبدیل نوع داده و انطباق طرح‌بندی را برای نگاشت ضرب ماتریس ورودی مختلط بر روی انواع داده‌ها و طرح‌بندی‌های سخت‌افزاری ارائه می‌کنیم. نتایج ما نشان می‌دهد که سربار کار اضافی در نرم‌افزار حداقل است و عملکردی نزدیک به حداکثر قابلیت‌های سخت‌افزاری را ممکن می‌سازد. تکنیک های نرم افزاری که در اینجا توضیح داده شده است در مخزن منبع باز NVIDIA/CUTLASS منتشر شده است.

ردپای حافظه برای مدل LLM پارامتر 175B با فرمت‌های مختلف انواع داده.

عملیات ماتریس ضرب-انباشت

شتاب‌دهنده‌های سخت‌افزار هوش مصنوعی مدرن مانند TPU Google و GPU NVIDIA با هدف قرار دادن Tensor Cores، که عناصر پردازش تخصصی برای تسریع عملیات ماتریس هستند، به‌ویژه برای بارهای کاری هوش مصنوعی، ماتریس‌ها را به‌طور طبیعی در سخت‌افزار ضرب می‌کنند. در این وبلاگ، ما بر هسته های تانسور آمپر NVIDIA تمرکز می کنیم که این هسته ها را ارائه می دهند ماتریس-ضرب-انباشت (mma) عمل. برای بقیه مطالب وبلاگ ارجاع به mma برای هسته های تانسور آمپر است. انواع داده های پشتیبانی شده، شکل ها و طرح داده های دو ماتریس ورودی (که عملوند نامیده می شوند) برای mma عملیات در سخت افزار ثابت شده است. این بدان معنی است که ضرب ماتریس با انواع داده های مختلف و اشکال بزرگتر در نرم افزار با کاشی کردن مشکل بر روی انواع داده ها، اشکال و طرح بندی های پشتیبانی شده از سخت افزار پیاده سازی می شود.

هسته تنسور mma عملیات با تعیین دو ماتریس ورودی تعریف می شود (به عنوان مثال، آ & ب، نشان داده شده در زیر) برای تولید یک ماتریس نتیجه، سی. این mma عملیات به طور بومی از دقت مخلوط پشتیبانی می کند. هسته های تانسوری با دقت ترکیبی اجازه مخلوط کردن ورودی (آ و ب) نوع داده با نتیجه (سی) نوع داده. متقابلا، ورودی مختلط ضرب ماتریس شامل مخلوط کردن انواع داده های ورودی است و توسط سخت افزار پشتیبانی نمی شود، بنابراین باید در نرم افزار پیاده سازی شود.

عملکرد Tensor Core M-by-N-by-K روی ماتریس ورودی A از M-by-K و ماتریس B از K-by-N ماتریس خروجی C از M-by-N را تولید می کند.

چالش های ضرب ماتریس ورودی مخلوط

برای ساده‌تر کردن بحث، به یک مثال خاص از ضرب ماتریس ورودی مختلط محدود می‌کنیم: F16 برای ورودی کاربر و U8 برای وزن‌های مدل (نوشته شده به صورت F16 * U8). تکنیک های توضیح داده شده در اینجا برای ترکیب های مختلفی از انواع داده های ورودی مختلط کار می کنند.

یک برنامه نویس GPU می تواند به سلسله مراتبی از حافظه از جمله حافظه جهانی، حافظه مشترک و رجیسترها دسترسی داشته باشد که به ترتیب کاهش ظرفیت اما افزایش سرعت مرتب شده اند. هسته تانسور آمپر NVIDIA mma عملیات ماتریس های ورودی را از ثبات ها مصرف می کند. علاوه بر این، ماتریس‌های ورودی و خروجی برای انطباق با طرح‌بندی داده‌ها در یک گروه از 32 رشته به نام پیچ و تاب. نوع داده پشتیبانی شده و چیدمان در یک تار برای یک ثابت هستند mma عملیات، بنابراین برای اجرای موثر ضرب ورودی مختلط، حل چالش‌های تبدیل نوع داده و انطباق طرح‌بندی در نرم‌افزار ضروری است.

تبدیل نوع داده

این mma عملیات به دو ماتریس ورودی با یک نوع داده نیاز دارد. بنابراین، ضرب ماتریس ورودی مختلط، که در آن یکی از عملوندها در U8 در حافظه جهانی و دیگری در F16 ذخیره می شود، نیاز به تبدیل نوع داده از U8 به F16 دارد. این تبدیل، دو عملوند را به F16 می آورد و از آن نقشه برداری می کند ورودی مختلط ضرب ماتریس به سخت افزار پشتیبانی می شود دقت مخلوط هسته های تانسور با توجه به تعداد زیاد وزن ها، تعداد زیادی از این عملیات وجود دارد، و تکنیک های ما نشان می دهد که چگونه می توان تأخیر آنها را کاهش داد و عملکرد را بهبود بخشید.

انطباق طرح

این mma عملیات همچنین مستلزم چیدمان دو ماتریس ورودی، در رجیسترهای یک Warp است تا با مشخصات سخت افزاری مطابقت داشته باشد. طرح بندی برای ماتریس ورودی ب نوع داده U8 در ضرب ماتریس ورودی مختلط (F16 * U8) باید با نوع داده تبدیل شده F16 مطابقت داشته باشد. به این می گویند مطابقت چیدمان و باید در نرم افزار به دست آید.

شکل زیر یک را نشان می دهد mma ماتریس مصرف کننده عملیات آ و ماتریس ب از ثبات ها برای تولید ماتریس سی در رجیسترها، در یک تار توزیع شده است. تهدید T0 برای نشان دادن ماتریس وزن، هایلایت و بزرگنمایی می شود ب از طریق تبدیل نوع داده می‌گذرد و برای اینکه بتواند به عملیات Tensor Core با پشتیبانی سخت‌افزار نقشه‌برداری کند، به یک انطباق طرح نیاز دارد.

استراتژی های نرم افزاری برای مقابله با چالش ها

یک تبدیل نوع داده معمولی شامل دنباله ای از عملیات روی ثبات های 32 بیتی است که در زیر نشان داده شده است. هر بلوک مستطیلی یک ثبت را نشان می دهد و متن مجاور عملیات است. کل دنباله تبدیل از 4xU8 به 2x (2xF16) را نشان می دهد. این توالی شامل تقریباً 10 عملیات است.

راه های زیادی برای دستیابی به انطباق طرح وجود دارد. دو تا از راه حل های موجود عبارتند از:

1. بارهای حافظه مشترک با پهنای بیت باریکتر: در این رویکرد، رشته ها بارهای حافظه با پهنای بیت باریکی را صادر می کنند و داده های U8 را از حافظه مشترک به رجیسترها منتقل می کنند. این منجر به دو ثبات های 32 بیتی، با هر ثبات حاوی مقادیر 2xF16 (در بالا برای ماتریس نشان داده شده است. بموضوع T0). بار حافظه مشترک باریک‌تر به انطباق طرح‌بندی مستقیماً با ثبات‌ها بدون نیاز به هیچ گونه تغییری دست می‌یابد. با این حال، از پهنای باند حافظه مشترک کامل استفاده نمی کند.
2. پیش پردازش در حافظه جهانی: یک استراتژی جایگزین شامل مرتب کردن مجدد داده ها در حافظه جهانی (یک سطح بالاتر از حافظه مشترک در سلسله مراتب حافظه) است که امکان بارگذاری حافظه مشترک گسترده تر را فراهم می کند. این رویکرد استفاده از پهنای باند حافظه مشترک را به حداکثر می‌رساند و تضمین می‌کند که داده‌ها در یک طرح منطبق مستقیماً در ثبات‌ها بارگذاری می‌شوند. اگرچه فرآیند بازآرایی را می توان قبل از استقرار LLM به صورت آفلاین اجرا کرد، و از عدم تاثیر بر عملکرد برنامه اطمینان حاصل کرد، اما یک مرحله پیش پردازش اضافی و غیر پیش پا افتاده مخصوص سخت افزار را معرفی می کند که به یک برنامه اضافی برای تنظیم مجدد داده ها نیاز دارد. NVIDIA/FasterTransformer این روش را برای مقابله موثر با چالش‌های انطباق طرح‌بندی اتخاذ می‌کند.

استراتژی های نرم افزاری بهینه شده

برای بهینه سازی بیشتر و کاهش سربار تبدیل نوع داده و انطباق طرح، ما را پیاده سازی کرده ایم FastNumericArrayConvertor و FragmentShuffler، به ترتیب.

FastNumericArrayConvertor بر روی 4xU8 در رجیسترهای 32 بیتی بدون باز کردن مقادیر جداگانه 1xU8 کار می کند. علاوه بر این، از عملیات حسابی کم‌هزینه‌تر استفاده می‌کند که تعداد دستورالعمل‌ها را کاهش می‌دهد و سرعت تبدیل را افزایش می‌دهد.

توالی تبدیل U8 به F16 در زیر نشان داده شده است. این عملیات از رجیسترهای بسته بندی شده 32b استفاده می کند و از بازکردن و بسته بندی صریح اجتناب می کند. FastNumericArrayConvertor استفاده می کند permute byte برای مرتب کردن مجدد بایت های 4xU8 به دو ثبات. علاوه بر این، FastNumericArrayConvertor از دستورالعمل های تبدیل عدد صحیح به ممیز شناور گران قیمت استفاده نمی کند و از عملیات بردار برای به دست آوردن نتایج بسته بندی شده در دو ثبات های 32 بیتی حاوی مقادیر 2x(2xF16). این FastNumericArrayConvertor برای U8-to-F16 تقریباً از شش عملیات استفاده می کند که نسبت به روش نشان داده شده در بالا 1.6× کاهش می یابد.

FastNumericArrayConvertor استفاده می کند permute bytes و محاسبات بسته بندی شده، تعداد دستورالعمل ها را در تبدیل نوع داده کاهش می دهد.

FragmentShuffler انطباق طرح را با به هم زدن داده ها به گونه ای مدیریت می کند که امکان استفاده از عملیات بارگذاری پهنای بیت گسترده تر، افزایش استفاده از پهنای باند حافظه مشترک و کاهش تعداد کل عملیات را فراهم می کند.

معماری NVIDIA Ampere یک دستورالعمل ماتریس بار ارائه می دهد (ldmatrix). این ldmatrix یک عملیات در سطح Warp است که در آن 32 رشته از یک Warp داده‌ها را از حافظه مشترک به ثبت‌های موجود در شکل و چیدمان که mma ماتریس آ و ب مصرف کردن. استفاده از ldmatrix کاهش می دهد تعداد دستورالعمل های بارگذاری و افزایش استفاده از پهنای باند حافظه از آنجا که ldmatrix دستورالعمل داده های U8 را به رجیسترها منتقل می کند، طرح بندی پس از بارگذاری با U8*U8 مطابقت دارد. mma عملیات، و نه با F16 * F16 mma عمل. اجرا کردیم FragmentShuffler برای تنظیم مجدد داده ها در ثبات ها با استفاده از shuffle (shfl.sync) عملیات برای دستیابی به انطباق طرح.

مهم‌ترین سهم این کار دستیابی به انطباق طرح‌بندی از طریق ترکیب‌های ثبت نام، اجتناب از پیش‌پردازش آفلاین در حافظه جهانی یا بارهای حافظه مشترک با پهنای بیت باریک‌تر است. علاوه بر این، ما پیاده سازی هایی برای FastNumericArrayConvertor شامل تبدیل نوع داده از U8 به F16، S8 به F16، U8 به BF16 و S8 به BF16.

نتایج عملکرد

ما عملکرد هشت نوع ورودی مختلط را اندازه‌گیری کردیم روش ما (در زیر به رنگ آبی و قرمز نشان داده شده است؛ انواع داده های ماتریس را تغییر دهید آ و ب) و دو دقت مخلوط انواع داده (به رنگ سبز نشان داده شده است) در تراشه NVIDIA A100 SXM. نتایج عملکرد در FLOPS نشان داده شده است (بالاتر بهتر است). قابل ذکر است، هشت ضرب-ماتریس اول به عملیات اضافی نسبت به دو مورد آخر نیاز دارند، زیرا انواع با دقت ترکیبی مستقیماً عملیات Tensor Core با شتاب سخت‌افزاری را هدف قرار می‌دهند و نیازی به تبدیل نوع داده و انطباق طرح‌بندی ندارند. با این حال، رویکرد ما عملکرد ضرب ماتریس ورودی مختلط را فقط کمی کمتر یا همتراز با دقت مخلوط نشان می‌دهد.

عملکرد ضرب ماتریس ورودی مختلط در تراشه NVIDIA A100 40GB SMX4 برای شکل مشکل ماتریس محاسباتی m=3456, n=4096, k=2048.

سپاسگزاریها

ما می‌خواهیم به چند نفر از افرادی اشاره کنیم که از طریق طوفان فکری فنی و بهبود پست وبلاگ کمک کرده‌اند، از جمله کوئنتین کلمبت، ژاک پینار، آلی کالپ، کالین کاسکاوال، آشیش گوندیمالا، مت والش، مارک کولودزیج، و امان بهاتیا. مایلیم از شرکای NVIDIA خود، Rawn Henry، Pradeep Ramani، Vijay Thakkar، Haicheng Wu، Andrew Kerr، Matthew Nicely و Vartika Singh تشکر کنیم.