اگر میخواهید از یک کاربر سادهی کتابخانههای هوش مصنوعی به یک متخصص محاسبات GPU تبدیل شوید، باید بدانید که در لایههای زیرین، هر خط کد PyTorch به صدها عملیات پیچیده حافظه تبدیل میشود. تفاوت میان یک توسعهدهنده و یک مهندس سیستم در درک همین جزئیات سطح پایین است.
به نقل از راهنمای فنی منتشر شده در dev.to در ۸ اوت ۲۰۲۶، ساخت یک آموزشدهنده برای مدل زبانی بزرگ (LLM) — شبیه کتابخانهداری که میلیاردها صفحه را خوانده و حالا با همان لحن کتابها جواب میدهد — نیازمند مدیریت سختگیرانه حافظه و پیادهسازی دقیق پسانتشار (Backpropagation) است. این راهنما هستههای GPU مورد نیاز برای تبدیل یک بلوک ترنسفورمر (Transformer) به یک مدل قابل آموزش را با استفاده از HIP (لایه قابل انتقال برای CUDA و ROCm) تشریح میکند. این رویکرد در راستای بهینهسازی زیرساختهای آموزشی است، مشابه آنچه در بررسی سازوکار همگامسازی گرههای جایگزین در PyTorch برای کاهش زمان توقف آموزش مدلها مشاهده کردیم.
بسیاری از توسعهدهندگان به انتزاعهای سطح بالا تکیه میکنند که در آنها یک خط کد، کل فرآیند گرادیان کاهشی (Gradient Descent) را مدیریت میکند. اما درک ابزارهای اولیه، همان چیزی است که تسلط بر سختافزار را به ارمغان میآورد. این رویکرد دقیقاً مشابه نیاز به کنترل معماری عمیقی است که در پوشش قبلی خود درباره نحوه مدیریت بستر محلی (Local Context) در EasyRouterAI برای پایین نگه داشتن هزینههای LLM مشاهده کردیم.
طبق مستندات این راهنما، برای اجرای محلی این سیستم، توسعهدهندگان به تسلط بر قاعده زنجیرهای (Chain Rule) و یک GPU با حداقل ۸ گیگابایت حافظه ویدیویی (VRAM) نیاز دارند. هدف این است که از مرحلهی «گذر پیشرو» عبور کنیم؛ مرحلهای که نویسنده آن را «تنها یک تولیدکننده اعداد تصادفی بسیار گرانقیمت» توصیف میکند، زیرا بدون وجود گرادیانها، هیچ یادگیری رخ نمیدهد.
برای آموزش مدل، ابتدا باید پسانتشار را برای لایههای خطی پیاده کرد. در حالی که گذر پیشرو یک ضرب ماتریسی ساده است (Y = X * W + b)، پسانتشار به دو عملیات مجزا نیاز دارد: محاسبه گرادیان نسبت به ورودی (dX) و وزنها (dW). از نظر ریاضی، اینها به صورت dX = dY * W^T و dW = X^T * dY تعریف میشوند.
بر اساس آموزش dev.to، این هدف با استفاده مجدد از hipblasSgemm یا rocBLAS و ترانهادن ماتریسها محقق میشود. در ساختاری که d_Y دارای ابعاد [B, S, D] و W دارای [D, D] است، برای محاسبه d_X از HIPBLAS_OP_N برای d_Y و HIPBLAS_OP_T برای W استفاده میشود. برای محاسبه d_W، این عملیاتها معکوس شده و از HIPBLAS_OP_T برای X و HIPBLAS_OP_N برای d_Y استفاده میشود. اگرچه گرادیانهای توجه چندسر (Multi-head Attention) به دلیل وجود سافتمکس (Softmax) و ضرب Q*K^T پیچیدهتر هستند، اما اصل ثابت است: هر ضرب ماتریسی در گذر پیشرو، با دو ضرب ماتریسی در پسانتشار متناظر است.
هستههای سفارشی برای توابع غیرخطی
لایههای خطی توسط کتابخانهها مدیریت میشوند، اما توابع غیرخطی به هستههای سفارشی GPU نیاز دارند. این راهنما پیادهسازی خاصی برای هسته پسانتشار سافتمکس ارائه میدهد. این هسته گرادیان را از طریق ضرب داخلی احتمالات پیشرو (P) و گرادیانهای ورودی (dY) محاسبه کرده و سپس فرمول dX_i = P_i * (dY_i - sum(P_j * dY_j)) را اعمال میکند. در حالی که مدیریت دستی این هستهها چالشبرانگیز است، ابزارهایی مانند TileLang امکان کامپایل کرنلهای GPU را با پایتون فراهم میکنند تا نیاز به مدیریت دستی حافظه کاهش یابد.
- مکانیزم سافتمکس: هسته
softmax_backward_kernelروی سطرها و ستونها پیمایش کرده و ابتدا یک ضرب داخلی برای سطر محاسبه میکند و سپس مقدار dX = P * (dY - dot) را مییابد. - مکانیزم LayerNorm: این لایه نیازمند محاسبه گرادیان برای ورودی X و پارامترهای مقیاس/بایاس (گاما و بتا) است.
در نرمالسازی لایهای (LayerNorm)، نویسنده پیشنهاد میکند میانگین و معکوس انحراف معیار (inv_std) از گذر پیشرو در یک بافر کوچک ذخیره شوند. این کار مانع از محاسبه مجدد این مقادیر در پسانتشار شده و چرخههای پردازشی گرانبهای GPU را ذخیره میکند. این فرآیند شامل کاهش (Reduction) روی بعد پنهان برای محاسبه d_gamma و d_beta و سپس استفاده از مقادیر استاندارد شده برای یافتن d_X است.
بهینهساز ادغامشده AdamW
بهروزرسانی وزنها برای میلیاردها پارامتر میتواند سربار شدیدی در اجرای هستهها (Kernel Launch Overhead) ایجاد کند. برای حل این مشکل، یک هسته ادغامشده برای بهینهساز AdamW پیاده شده است. به جای اجرای هستههای جداگانه برای تکانه (Momentum) و واریانس (Variance)، یک تکهسته همه چیز را در یک مسیر بهروز میکند تا پهنای باند حافظه به حداکثر برسد و سربار به حداقل برسد.
این هسته ادغامشده موارد زیر را مدیریت میکند:
- تکانه و واریانس: نگهداری دو میانگین متحرک نمایی برای هر پارامتر: m (تکانه) و v (واریانس).
- اصلاح بایاس: محاسبه
bias_correction1(1.0f - powf(beta1, step)) وbias_correction2(1.0f - powf(beta2, step)). - کاهش وزن: اعمال مستقیم کاهش وزن (Weight Decay) در هنگام بهروزرسانی با فرمول: W[idx] = W[idx] - lr * (m_hat / (sqrtf(v_hat) + eps) + weight_decay * W[idx]).
- بازنشانی گرادیان: صفر کردن اختیاری dW[idx] = 0.0f برای تکرار بعدی.
دقت ترکیبی و مدیریت حافظه
سختافزارهای مدرن مانند NVIDIA Ampere و AMD CDNA+ از سختافزار اختصاصی FP16/BF16 پشتیبانی میکنند. راهنما توضیح میدهد که آموزش باید از FP16 برای گذر پیشرو و پسرو استفاده کند تا توان عملیاتی دو برابر شود، در حالی که یک نسخه FP32 از وزنها (Master Copy) برای جلوگیری از محوشدن گرادیان (Gradient Underflow) حفظ شود.
به دلیل محدوده دینامیکی کم در FP16، از مقیاسبندی زیان (Loss Scaling) استفاده میشود. زیان قبل از پسانتشار در یک عدد بزرگ (مثلاً ۱۰۲۴) ضرب و سپس قبل از بهروزرسانی وزنها تقسیم میشود. این کار تضمین میکند که گرادیانهای کوچک به صفر تبدیل نشوند. این پیادهسازی از #include <hip/hip_fp16.h> و یک هسته cast_and_scale_gradients برای تبدیل half* dW_half به float* dW_float با استفاده از ضریب مقیاس استفاده میکند.
برای دستیابی به حداکثر دقت، نویسنده توصیه میکند از hipblasGemmEx برای انتخاب نوع محاسبات (FP32 برای تجمع یا Accumulation) و از hipblasHgemm برای عملیات FP16 استفاده شود.
نقطه بازرسی فعالسازها
حافظه، گلوگاه اصلی برای مدلهای ۷ میلیارد پارامتر به بالاست. ذخیره تمام فعالسازها برای پسانتشار اغلب غیرممکن است. راهنما تکنیک نقطه بازرسی فعالسازها (Activation Checkpointing) را معرفی میکند که در آن محاسبات در برابر حافظه اولویت مییابند و تنها ورودیهای لایههای خاص (مثلاً هر چهارمین بلوک) ذخیره میشوند. این رویکرد به مدیریت بهینه حافظه کمک میکند، مشابه روشهای پیشرفتهای که در مدیریت حافظه مجازی کتابخانه vLLM برای افزایش توان عملیاتی مدلها به کار گرفته شده است.
- استراتژی پیشرو: قبل از گذر پیشرو در یک بلوک، ورودی X در یک بافر نقطه بازرسی با استفاده از
hipMemcpyDeviceToDeviceذخیره میشود. - استراتژی پسرو: سیستم X را مجدداً بارگذاری کرده و کل گذر پیشرو آن بلوک را دوباره اجرا میکند تا فعالسازهای مورد نیاز برای پسانتشار بازسازی شوند.
این استراتژی نیاز به محاسبات را تقریباً ۳۰ تا ۴۰ درصد افزایش میدهد اما ردپای VRAM را بهشدت کاهش میدهد و برای LLMهای بزرگ اجباری است.
چرخه کامل آموزش
یک تکرار آموزشی در HIP/C++ خالص از این توالی سختگیرانه پیروی میکند:
۱. انتقال داده: کپی نامتقارن (Async) دستهها از CPU به GPU با استفاده از hipMemcpyAsync.
۲. گذر پیشرو: اجرای FP16 در ترنسفورمر.
۳. محاسبه زیان: استفاده از هسته compute_loss_kernel برای محاسبه آنتروپی متقاطع (Cross-entropy).
۴. مقیاسبندی زیان: اعمال scale_loss_kernel روی مقدار زیان.
۵. پسانتشار: پیمایش معکوس گراف برای محاسبه گرادیانهای FP16.
۶. برداشتن مقیاس: تبدیل گرادیانها به FP32 با استفاده از cast_and_scale_gradients و ضریب 1.0f/loss_scale.
۷. برش گرادیان: محاسبه اختیاری نرم L2 از طریق compute_l2_norm_kernel برای جلوگیری از انفجار گرادیان.
۸. گام بهینهساز: بهروزرسانی وزنهای FP32 توسط AdamW با اجرای هسته در (n+255)/256 بلوک.
۹. همگامسازی وزنها: تبدیل مجدد وزنهای بهروزرسانی شده FP32 به FP16 برای تکرار بعدی.
۱۰. همگامسازی نهایی: اجرای hipStreamSynchronize برای اطمینان از تکمیل تمام عملیات.
تحلیل گلوگاهها
برای بهینهسازی این هستهها، استفاده از nvprof یا Nsight Systems (nsys profile) برای انویدیا و rocprof یا OmniTrace برای AMD توصیه میشود. هدف، به حداکثر رساندن اشغال (Occupancy) و اطمینان از این است که عملیات GEMM محدود به محاسبات (Compute-bound) باشند، نه حافظه (Memory-bound).
متریکهای کلیدی برای نظارت عبارتند از:
- اشغال (Occupancy): بررسی
achieved_occupancyبرای دیدن اینکه آیا Warpها یا Wavefrontها بیکار هستند یا خیر. - پهنای باند حافظه: شناسایی هستههای محدود به حافظه مانند Softmax و LayerNorm برای پیادهسازی کاهشهای حافظه مشترک (Shared Memory Reductions).
- سربار اجرا: کاهش تعداد هستههای بسیار کوچک از طریق ادغام (Fusion).
اگر توسعهدهنده متوجه تعداد زیادی هسته کوچک شود، راهکار ادغام است؛ مثلاً ترکیب افزودن بایاس، فعالسازی و دراپاوت در یک تکهسته، سربار اجرای مکرر روی GPU را کاهش میدهد.
این پیادهسازی سطح پایین دقیقاً نشان میدهد وقتی در PyTorch دستور loss.backward() و optimizer.step() را میزنید چه اتفاقی میافتد. با دور زدن چارچوبها، امکان بهینهسازی برای سختافزارهای خاص مانند NVIDIA H100 یا AMD MI300X فراهم میشود.
این تغییر جهت به سمت توسعه هستههای سفارشی به مهندسان اجازه میدهد تا از هر قطره عملکرد سختافزار خود استفاده کنند. این روند صنعت را از آموزش مدلهای عمومی به سمت بهینهسازیهای آگاه از سختافزار (Hardware-aware) سوق میدهد.
در ادامه این مجموعه، به سراغ آموزش توزیعشده در چندین GPU خواهیم رفت و بر استراتژیهای All-Reduce، Ring-AllReduce و روشهای sharding مانند مراحل ZeRO با استفاده از NCCL و RCCL تمرکز خواهیم کرد.
گفتگو