بهینه‌سازی هسته‌های GPU با Codex؛ دستیابی به سرعت ۲۳۲ برابر در محاسبات QR

تصور کنید کدی که در حالت عادی چند ثانیه زمان می‌برد، ناگهان در چند میلی‌ثانیه اجرا شود. این دقیقاً همان اتفاقی است که در یک رقابت بهینه‌سازی سخت‌افزاری رخ داد: دستیابی به سرعت ۲۳۲ برابر نسبت به استانداردهای فعلی.

به نقل از گزارش‌های مسابقه GPU Mode، یک پژوهشگر توانست با استفاده از Codex (مدل زبانی بزرگ OpenAI) در یک چرخه «مهندسی حلقه»، هسته‌های تجزیه QR را به شکلی بهینه کند که رتبه ۱۲ از میان ۱۸۳ شرکت‌کننده را در ۸ جولای ۲۰۲۶ کسب کند. این دستاورد در جریان سری مسابقات «هسته‌های جبر خطی در عصر پژوهش» و با همکاری Core Automation به دست آمد.

این نتیجه در حالی منتشر می‌شود که صنعت از مدل‌های ساده پرسش-پاسخ به سمت گردش‌های کاری عامل‌محور (Agentic) — شبیه به داشتن یک دستیار متخصص که نه تنها کد می‌نویسد، بلکه خودش آزمایش می‌کند و اشتباهاتش را اصلاح می‌کند — حرکت می‌کند. در حالی که اکثر توسعه‌دهندگان از مدل‌های زبانی بزرگ (LLM) برای نوشتن تکه‌های کوچک کد استفاده می‌کنند، این رویکرد با هوش مصنوعی به عنوان یک عامل پژوهشی برخورد می‌کند که قادر به بنچ‌مارکینگ تکرار شونده، پروفایلینگ و تست فرضیات است. همان‌طور که در تحلیل‌های قبلی ما درباره امنیت مدل‌های بازمتن اشاره کردیم، قدرت واقعی هوش مصنوعی زمانی آشکار می‌شود که از تولید متن ساده به سمت مدیریت فرآیندهای پیچیده برود. در اینجا، نقش انسان از «کدنویس» به «هدایت‌کننده» تغییر کرده است؛ کسی که بر استراتژی سطح بالا تمرکز می‌کند تا سینتکس کد. این پژوهشگر طی ۱۴ روز، بیش از ۱,۵۰۰ نسخه مختلف از کد را برای ثبت در جدول امتیازات ارسال کرد.

چالش فنی: تجزیه QR

هدف این مسابقه، پیاده‌سازی تجزیه QR به روش Householder برای دسته‌هایی از ماتریس‌های مربع (Batched Square Compact-Householder QR Factorization) بود. ورودی‌ها شامل دسته‌ای از ماتریس‌های CUDA با دقت FP32 و ابعاد batch x n x n بودند. هدف نهایی، بازگرداندن یک نمایش فشرده هاوس‌هولدر بود که دقیقاً با خروجی تابع torch.geqrf(A) مطابقت داشته باشد: یک ماتریس H (که در آن مثلث بالایی R و مثلث پایینی بردارهای هاوس‌هولدر را ذخیره می‌کند) و یک بردار tau از ضرایب بازتاب‌دهنده.

برای تأیید صحت، سیستم بررسی (Checker) ابتدا ماتریس Q را با استفاده از torch.linalg.householder_product(H, tau) بازسازی می‌کرد، سپس R را به عنوان مثلث بالایی H در نظر می‌گرفت و در نهایت بررسی می‌کرد که آیا A ≈ QR، QᵀQ ≈ I و QᵀA ≈ R برقرار است یا خیر. معیار رتبه‌بندی، میانگین هندسی زمان اجرا برای ابعاد مختلف و حالت‌های مختلف شرطی‌سازی (Conditioning) بود، به ویژه برای ماتریس‌های مربع با ابعاد ۵۱۲، ۱۰۲۴، ۲۰۴۸ و ۴۰۹۶. اگرچه استفاده داخلی از دقت‌های FP16، FP8 یا NVFP4 مجاز بود، اما فاکتورهای نهایی باید استانداردهای بررسی دقت FP32 را پاس می‌کردند.

ریاضیات بازتاب‌های هاوس‌هولدر

در بازتاب هاوس‌هولدر، هدف این است که یک بردار (مثلاً یک ستون از ماتریس A) در یک مرحله روی یک محور هدف قرار گیرد. برای مثال، در یک ماتریس ۳x۳، اولین بازتاب‌دهنده ستون اول (۱۲، ۶، -۴) را به (-۱۴، ۰، ۰) منتقل می‌کند. مقدار -۱۴ تبدیل به R11 می‌شود و درایه‌های زیر آن صفر می‌گردند.

• مکانیزم آینه: بردار هاوس‌هولدر v مانند یک آینه عمل می‌کند. بازتاب بر اساس فرمول ℋx = x - τv(vᵀx) محاسبه می‌شود، که در آن τ = 2 / (vᵀv). آینه در واقع هر چیزی است که بر v عمود است و بازتاب در امتداد v حرکت می‌کند. بخش عمود بر بردار v، در واقع سایه x روی v است که به صورت (vᵀxvᵀv) / vᵀv از کپی‌های v محاسبه می‌شود.
• فرمت فشرده: برای صرفه‌جویی در حافظه، تابع geqrf از فضاهای خالی زیر قطر اصلی ماتریس H برای ذخیره دنباله بردارهای هاوس‌هولدر v_j استفاده می‌کند (در حالی که عدد ۱ ابتدایی به صورت ضمنی در نظر گرفته می‌شود). این کار اجازه می‌دهد تنها یک ماتریس و دو محموله (H و tau) کل تجزیه را نمایش دهند.
• گلوگاه ترتیبی: مشکل اصلی این است که تجزیه QR استاندارد هاوس‌هولدر ذاتاً ترتیبی (Serial) است. بازتاب‌دهنده j+1 تا زمانی که بازتاب‌دهنده j روی بلوک باقی‌مانده اعمال نشود، نمی‌تواند ساخته شود. این کارهای ماتریس-برداری ترتیبی در مسیرهای کند برداری اجرا می‌شوند و باعث می‌شوند هسته‌های تنسور (Tensor Cores) در GPU بیکار بمانند.

الگوریتم هاوس‌هولدر بلوکی

برای حل این گلوگاه، پژوهشگر از «الگوریتم هاوس‌هولدر بلوکی» استفاده کرد. در این روش، کارهای ترتیبی به یک پنل باریک از b ستون (مثلاً ۳۲ یا ۶۴ ستون) محدود می‌شوند.

به جای اعمال تک‌تک بازتاب‌ها، آن‌ها در یک به‌روزرسانی تک‌رتبه به نام نمایش WY فشرده می‌شوند. این کار بازتاب‌ها را به فرم I - VTVᵀ در می‌آورد، که در آن V ماتریسی از بردارهای هاوس‌هولدر و T یک ماتریس مثلثی بالایی کوچک b x b است. این تغییر، به‌روزرسانی بلوک‌های باقی‌مانده را به سه عملیات ضرب ماتریس بزرگ (GEMM) متوالی تبدیل می‌کند:

1. W = VᵀAtrail
2. Z = TᵀW
3. Atrail ← Atrail - VZ

این تغییر تضمین می‌کند که بخش عمده محاسبات به شکل GEMM باشد، که زبان مادری هسته‌های تنسور است. سپس پنل روی بلوک باقی‌مانده می‌لغزد و این فرآیند با کوچک شدن بلوک تکرار می‌شود.

گردش کار پژوهش خودکار

این سیستم بر پایه یک حلقه بازخورد تنگ بین Codex (نسخه GPT-5.5) و Claude بنا شده بود. ابزار اصلی این ارتباط، رابط کاربری popcorn CLI بود که به عامل اجازه می‌داد کد را تست کند، بنچ‌مارک بگیرد و نتایج را مستقیماً به جدول امتیازات بفرستد. برگزارکنندگان مسابقه اجازه ارسال نامحدود کد را دادند، به شرطی که برای جلوگیری از Time-out شدن صف‌ها، فاصله زمانی بین ارسال‌ها رعایت شود. در یک مرحله، حجم ارسال‌ها به قدری زیاد بود که اعتبار حساب Modal پژوهشگر تمام شد.

اجزای کلیدی این سیستم عبارت بودند از:

• AGENTS.md: فایلی حاوی دستورالعمل‌های ارشد که دیسیپلین عامل را تعریف می‌کرد. این فایل شامل «دیسیپلین جست‌وجوی پرتویی» (Beam Search) بود تا از حذف زودهنگام خانواده‌های ایده‌های امیدوارکننده جلوگیری شود و همچنین «دیسیپلین ارسال» را تعریف می‌کرد تا اطمینان حاصل شود که پس از هر پیشرفت بزرگ، پروفایلینگ انجام شود. همچنین صراحتاً به عامل دستور می‌داد که برای کارهای سخت از زیر-عامل‌ها (Sub-agents) استفاده کند و Time-outها را به عنوان نتایج غیرقطعی تلقی کند.
• log.md: یک دفترچه یادداشت برای ثبت تمام ارسال‌ها، وضعیت آن‌ها (پذیرفته/رد شده) و زمان‌های اجرا برای هر ابعاد ماتریس، تا مدل ایده‌های شکست‌خورده را تکرار نکند. این فایل پس از رسیدن به رکورد ۳,۰۰۰ میکروثانیه حیاتی شد.
• Modal: پلتفرمی ابری برای پروفایلینگ و اجرای ابزارهای nsys و NCU (NVIDIA Compute Utility) جهت شناسایی گلوگاه‌های سخت‌افزاری. پژوهشگر دسترسی مستقیم Codex به پروفایلینگ Modal را فراهم کرد تا پارامترها را سریع‌تر بررسی کند.

عبور از نقاط بهینه محلی

وقتی سرعت اجرا به مرز ۳,۰۰۰ میکروثانیه رسید، مدل در «بهینه محلی» (Local Maxima) گیر کرد؛ یعنی تغییرات کوچک در پارامترها بدون پیشرفت چشمگیر تکرار می‌شدند. برای خروج از این وضعیت، استراتژی «پرتوی کاندیداها» (Beam-of-candidates) به کار گرفته شد.

به جای اینکه هر آزمایش جدید باید بلافاصله از بهترین رکورد فعلی سریع‌تر باشد، مدل ۳ تا ۵ «خانواده ایده‌ای» فعال را به طور هم‌زمان دنبال می‌کرد. این کار اجازه می‌داد تغییرات ساختاری که شاید در ابتدا کندتر باشند، اما در بلندمدت پس از چند تکرار به جهشی بزرگ منجر شوند، بررسی شوند. پژوهشگر در این میان نقش «سس مخفی» را داشت و مدل را به سمت ایده‌های جسورانه‌تر سوق می‌داد، مخصوصاً زمانی که مدل بیش از حد محافظه‌کار می‌شد.

برای مدیریت این تنوع، دستورات خاصی به Codex داده می‌شد:

• /goal: برای تعیین اهداف کمی و قابل دستیابی (مثلاً: «فقط از Triton یا CUDA استفاده کن و رکورد n=512 را بزن. cuSolver را کاملاً حذف کن»).
• /btw یا /side: برای پرسیدن سوالات جانبی درباره پیشرفت کار یا موفقیت‌ها بدون متوقف کردن حلقه اصلی بهینه‌سازی. سوالاتی مانند «پسر، داریم می‌بریم؟»، «در حال حاضر روی چه ایده‌هایی کار می‌کنی؟» و «آیا می‌توانیم این را به ابعاد دیگر منتقل کنیم؟».

مسیر رسیدن به سرعت ۲۳۲ برابر

تکامل کد از مسیرهای وابسته به کتابخانه به سمت اسمبلی سفارشی پیش رفت. زمان اجرای اولیه در torch.geqrf حدود ۴۱۹,۰۰۰ میکروثانیه بود که در نهایت به ۱,۸۰۵ میکروثانیه رسید.

مراحل تکامل هسته:

• خط‌مبنا: استفاده کامل از torch.geqrf در همه جا (بیش از ۱۰۸.۸ هزار میکروثانیه).
• تغییر الگوریتمی: پیاده‌سازی WY QR بلوکی برای ابعاد n=512 (۱۰۸.۸ هزار میکروثانیه).
• تعمیم: استفاده از مسیر بلوکی برای تمام ابعاد؛ برای n=32 از full-QR با به‌روزرسانی‌های LARFB16 استفاده شد (۱۰.۲ هزار میکروثانیه).
• تخصص‌یافتگی: استفاده از پنل‌های Triton و هسته‌های گروهی WY panel16/32 (۴.۳ هزار میکروثانیه).
• چرخش الگوریتمی: استفاده از Cholesky-ORHR برای ابعاد ۴۰۹۶ با استفاده از Gram-Schmidt و بازسازی بازتاب‌دهنده‌ها (۴ هزار میکروثانیه).
• کاهش سربار: استفاده از CUDA graph replay برای حذف سربار لانچ (۳.۴ هزار میکروثانیه).
• بهینه‌سازی حافظه: ادغام چیدمان V/T در اسمبلی برای حذف کپی‌های اسلایس و متغیرهای موقت (۲.۷۵ هزار میکروثانیه).
• بهینه‌سازی دنباله: استفاده از پنل‌های split16 و tail-Gram برای پرش از WY کامل در نزدیکی انتهای ماتریس (۲.۵ هزار میکروثانیه).
• سخت‌افزاری کردن: تخصص‌یافتگی برای ابعاد ثابت با ردیف‌های سخت‌افزاری (Hardcoded) و ادغام کاهش‌ها (۲ هزار میکروثانیه).
• معماری نهایی: ترکیب سوپرپنل‌ها، بسته‌های سفارشی Cholesky V256/T256 و بازگشت مستقیم H (۱.۸ هزار میکروثانیه).

تحلیل: گذار به «سلیقه پژوهشی»

این تجربه نشان می‌دهد که مرز بعدی کدنویسی با هوش مصنوعی، «تولید کد بهتر» نیست، بلکه «سلیقه پژوهشی» (Research Taste) است. توانایی مدل در تصمیم‌گیری درباره اینکه کدام آزمایش را بر اساس بازخورد پروفایلر اجرا کند، بسیار ارزشمندتر از توانایی نوشتن یک حلقه ساده است. کلید پیشرفت در پرسیدن سوالات درست بود؛ سوالاتی مثل «چه سیگنالی را می‌توانیم از زمان اجرا به سیگنال استاتیک تبدیل کنیم؟» یا «آیا ادغام کاهش‌ها (Reduction Fusions) ممکن است؟»

برای توسعه‌دهندگان، این یعنی مزیت رقابتی به کسانی می‌رسد که بتوانند حلقه‌های اعتبارسنجی مستحکم بسازند. ارزش انسان اکنون در «هدایت» است؛ ارائه تخصص دامنه برای اینکه به عامل بگوید به دنبال «ادغام کاهش‌ها» یا «سعه دقت» (Precision Slack) بگردد، مخصوصاً زمانی که مدل به یک پلاتو (سطح ثابت) می‌رسد. پژوهشگر اشاره کرد که در حالی که Codex بسیار سرسخت بود، Claude اغلب با بهانه‌هایی مبنی بر اینکه تمام بهینه‌سازی‌ها تمام شده است، تسلیم می‌شد.

درس‌های آموخته شده و بهبودهای آینده

بررسی ۱۰ رتبه اول مسابقه نشان داد که برخی فرصت‌ها از دست رفته بود:

• توزیع داده‌ها: برندگان از تشخیص‌دهنده‌های داده برای بهره‌برداری از موارد کم‌رتبه (Low-rank) که حاوی صفرهای زیادی هستند، استفاده کرده بودند.
• حذف کتابخانه: موفق‌ترین راهکارها، حل مثلثی PyTorch را با معکوس‌های مثلثی سفارشی جایگزین کرده بودند.
• مدیریت دقت: نگه داشتن ماتریس باقی‌مانده در حالت FP16 به جای جابجایی مکرر بین فرمت‌ها، سرعت را افزایش می‌داد.
• بهره‌برداری سخت‌افزاری: پژوهشگر نتوانست از دستورات tcgen05 برای بهره‌برداری کامل از هسته‌های تنسور B200 استفاده کند.
• استراتژی: استراتژی «پرتوی کاندیداها» باید از همان ابتدا اجرا می‌شد تا از بهینه محلی‌های اولیه جلوگیری شود.

گام بعدی شما

• اگر به بهینه‌سازی سطح پایین علاقه دارید، سری مسابقات GPU Mode را دنبال کنید.
• الگوی «مشاور-مجری» (Advisor-Executor) را امتحان کنید: از یک مدل استدلالی قوی مثل Claude Opus برای استراتژی و از یک مدل سریع‌تر برای اجرای کد استفاده کنید.
• ابزارهای پروفایلینگ انویدیا مانند NCU را یاد بگیرید تا بتوانید بازخوردهای دقیق‌تری به مدل‌های زبانی بدهید.

اما داستان سخت‌افزاری این تحول حتی شگفت‌انگیزتر است — به تحلیل ما درباره تراشه‌های Blackwell مراجعه کنید.