اگر توسعهدهندهای هستید که برای رسیدن به حداکثر سرعت در GPU مجبورید ساعتها با پیچیدگیهای CUDA C++ و همگامسازی دستی در سطح Warpها دستوپنجه نرم کنید، TileLang میتواند مسیر شما را تغییر دهد. نوشتن کرنلهای با کارایی بالا معمولاً نیازمند تخصص عمیق در CUDA C++ است. TileLang این روند را تغییر میدهد و به توسعهدهندگان اجازه میدهد تا کرنلهای متمرکز بر عملکرد را با استفاده از یک زبان تخصصی (DSL) سطح بالا در پایتون تعریف کنند که از طریق TVM کامپایل میشود.
به نقل از آموزشهای منتشر شده در Marktechpost در سال ۲۰۲۴، TileLang شکاف میان کدهای سطح بالای PyTorch و عملکرد خام CUDA را پر میکند. برای سالها، این فاصله تنها توسط تعداد کمی از متخصصان که کرنلهای پیچیده مینوشتند، پل میشد. در واقع، چالشهای مدیریت دستی حافظه و پیچیدگیهای سطح پایین CUDA را میتوان در پروژههایی مانند ساخت GPT-2 از صفر با زبان C و CUDA مشاهده کرد که نشان میدهد دستیابی به کارایی بالا بدون ابزارهای انتزاعی چقدر دشوار است. در حال حاضر، اکثر توسعهدهندگان به کتابخانههای پیشساخته و صلبی مانند cuBLAS متکی هستند که اگرچه سریع هستند اما انعطافناپذیرند. TileLang در زمانی عرضه شده است که اپراتورهای سفارشی — مانند آنچه برای مکانیزمهای جدید توجه (Attention) نیاز است — به گلوگاه اصلی بهرهوری در مدلهای زبانی بزرگ (LLM) تبدیل شدهاند.
همانطور که در تحلیلهای پیشین ما درباره بهینهسازی لایههای استنتاج اشاره کردیم، حذف سربار انتقال داده بین حافظهها کلید افزایش سرعت است. TileLang دقیقاً همینجا وارد عمل میشود و سختترین بخشهای برنامهنویسی GPU را انتزاع میکند. به جای مدیریت دستی اندیسهای رشته (Thread)، توسعهدهندگان با کاشیهای حافظه مشترک (Shared-memory tiles) و قطعات رجیستری (Register fragments) کار میکنند. کامپایلر وظیفه تولید دستورات سطح پایین CUDA، مدیریت چیدمان حافظه (Memory layouts) و برداریسازی (Vectorization) را بر عهده میگیرد.
برای ایجاد یک محیط کاری، سیستم ابتدا محیط CUDA را اعتبارسنجی کرده و TileLang را نصب میکند؛ در صورتی که نسخه پایدار (Stable wheel) قابل استفاده نباشد، سیستم از کانال Nightly به عنوان جایگزین استفاده میکند. این ابزار مستقیماً با PyTorch برای مدیریت تنسورها و تأیید عددی ادغام میشود. در مرحله تنظیمات محیط، سیستم قابلیت محاسباتی (Compute Capability - CC) و نسخه SM گرافیک را شناسایی میکند تا بودجه حافظه مشترک (Smem) را تعیین کند. این مقدار معمولاً برای SMهای قدیمیتر (کمتر از ۸۰) ۴۸ کیلوبایت و برای SM 80 به بالا ۹۶ کیلوبایت در هر بلوک است.
یکی از نقاط قوت اصلی این سیستم، توانایی آن در تولید کد بهینه برای دستگاه است. هنگام پیادهسازی یک جمع برداری ساده (با استفاده از حلقه T.Parallel و T.ceildiv برای محاسبه گرید)، TileLang کد منبع CUDA را تولید میکند که از نظر عملکرد پهنای باند با PyTorch برابری میکند. این موضوع ثابت میکند که لایه انتزاعی، سربار قابلتوجهی ایجاد نمیکند. این عملکرد با استفاده از بررسی «نرم نسبی فروبنیوس» (Relative-Frobenius-norm) تأیید میشود که برای دقت fp16 بسیار معنادارتر از تلورانس مطلق است.
قلب تپنده هوش مصنوعی، ضرب ضرب ماتریسی یا GEMM است. TileLang این عملیات را از طریق یک سلسلهمراتب سختگیرانه مدیریت میکند تا دادهها را به بهینهترین شکل جابهجا کند:
- حافظه سراسری (Global Memory): نقطه شروع برای تنسورهای ورودی.
- حافظه مشترک (Shared Memory): استفاده از
T.alloc_sharedبرای کاشیبندی و کاهش دفعات مراجعه به حافظه سراسری. - قطعات رجیستری (Register Fragments): جایی که محاسبات واقعی هستههای تنسور (Tensor-core) با استفاده از
T.alloc_fragmentرخ میدهد.
در جزئیات پیادهسازی GEMM، نکات فنی زیر حائز اهمیت است:
- خطلولهسازی (Pipelining): کرنل از حلقههای
T.Pipelinedبرای همپوشانی جابهجایی داده و محاسبات استفاده میکند. تعداد مراحل (Stages) پیشفرض برای GPUهای قدیمی ۲ و برای معماریهای SM 80 به بالا ۳ است. این رویکرد مشابه استراتژیهایی است که در بهینهسازی موتور Photon برای کاهش تأخیر استنتاج به کار گرفته شد تا بهرهوری سختافزار افزایش یابد. - جابهجایی حافظه: دادهها با دستور
T.copyاز حافظه سراسری به مشترک منتقل میشوند که کامپایلر آن را به دستورات بهینهای مانندcp.asyncیاldmatrixتبدیل میکند. - محاسبات: اپراتور
T.gemmاجرای واقعی هسته تنسور را فعال میکند. کامپایلر دستورات خاص هر معماری مانندmma.syncیاwgmmaرا صادر میکند. - بهینهسازی L2: توسعهدهندگان میتوانند با استفاده از
T.use_swizzle(panel_size=10, enable=True)قابلیت L2 Rasterization را فعال کنند تا نرخ برخورد حافظه پنهان (Cache hit) بهبود یابد.
در یک بنچمارک برای ماتریس ۲۰۴۸^۳، پیادهسازی GEMM در TileLang تنها با حدود ۲۰ خط کد پایتون، درصد قابلتوجهی از عملکرد cuBLAS را به دست آورد. برای بهینهسازی بیشتر، توسعهدهندگان میتوانند «پیچهای تنظیم» (Knobs) را به صورت دستی تغییر دهند؛ مانند ابعاد کاشی (مثلاً ۶۴x۶۴x۳۲ در مقابل ۱۲۸x۱۲۸x۶۴)، تعداد رشتهها (۱۲۸ یا ۲۵۶) و تنظیمات Swizzling. طبق مستندات، بهترین زمانبندی (Schedule) کاملاً به معماری سختافزار و شکل تنسورها بستگی دارد.
یکی از کاربردیترین ویژگیهای TileLang، «ادغام اپیلوگ» (Epilogue Fusion) است. در PyTorch استاندارد، افزودن بایاس و اعمال فعالساز GELU پس از ضرب ماتریسی، نیاز به سه بار اجرای کرنل مجزا دارد. این امر باعث میشود GPU مجبور شود نتایج میانی را به حافظه پهنایباند بالا (HBM) بنویسد و دوباره آنها را بخواند. TileLang اجازه میدهد این عملیات در یک کرنل واحد ادغام شوند. با تکمیل افزودن بایاس و فعالساز GELU در حالی که دادهها هنوز در رجیسترها هستند، سیستم مگابایتها از ترافیک HBM را ذخیره میکند. برای یک ماتریس ۴۰۹۶ در ۴۰۹۶ با K=۱۰۲۴، این ادغام تقریباً ۳۲ مگابایت از خواندن/نوشتن میانی را حذف میکند (محاسبه شده به صورت ~2MN*2/2^20). فعالساز GELU با استفاده از فرمول تقریبی tanh پیاده شده است: C_local[i, j] / (1.0 + T.exp(-1.5957691216 * (C_local[i, j] + 0.044715 * C_local[i, j] * C_local[i, j] * C_local[i, j]))).
به همین ترتیب، TileLang عملیات Softmax سطر-محور را از طریق کاهش (Reduction) در سطح قطعات رجیستری مدیریت میکند. این فرآیند شامل مراحل زیر است:
- کاهش حداکثر (Max Reduction): استفاده از
T.reduce_maxبرای یافتن حداکثر سطر جهت پایداری عددی. - نماییسازی (Exponentiation): تفریق مقدار حداکثر و محاسبه
T.expدر یک حلقهT.Parallel. - کاهش مجموع (Sum Reduction): استفاده از
T.reduce_sumبرای محاسبه فاکتور نرمالسازی.
این روش تضمین میکند که فرآیند کاهش دو مرحلهای هرگز رجیسترها را ترک نکند و عملیات را به جای «محاسبه-محور»، «حافظه-محور» نگه دارد. در تستهایی با M=۸۱۹۲ و N=۱۰۲۴، این رویکرد نرخ انتقال GB/s بالایی را حفظ کرد که با PyTorch قابل مقایسه است.
در مورد پیادهسازی توجه برقآسا (FlashAttention)، TileLang اجازه میدهد تا یک کرنل Forward ادغامشده در حدود ۷۰ خط کد پایتون پیاده شود. این یک بهینهسازی حیاتی برای LLMها است که از ایجاد ماتریس کامل امتیازات توجه در حافظه سراسری جلوگیری میکند. مکانیزمهای کلیدی در این پیادهسازی عبارتند از:
- Softmax آنلاین: استفاده از حداکثرهای جاری (
m_prev,m_cur) و مجموعهای نرمالسازی (logsum) برای بهروزرسانی امتیازات در لحظه. این سیستم از یک فاکتور بازسنجیalpha = T.exp((m_prev[i] - m_cur[i]) * scale)برای تنظیم انباشتهای قبلی استفاده میکند. - GEMM کاشیبندی شده: انجام ضرب ماتریسی روی بلوکهای کوچک از تنسورهای Query، Key و Value (مثلاً block_M=64, block_N=64).
- ماسک علی (Causal Masking): پیادهسازی توجه علی با صفر کردن شرطی امتیازات با استفاده از
T.if_then_else(bx * block_M + i >= k * block_N + j, 0.0, NEG)، که در آن NEG برابر با -1.0e30 است.
این مکانیزم کامل توجه علی و ادغامشده با هسته تنسور، تنسورهایی با شکل [batch, seq_len, heads, dim] را پردازش میکند و از طریق تخصص Warp و TMA (شتابدهنده حافظه تنسور)، مسیری به سوی عملکرد در سطح FlashMLA روی GPUهای H100 فراهم میکند.
از آنجا که بهترین پیکربندی کرنل به معماری GPU و ابعاد تنسور بستگی دارد، TileLang دکوراتور @tilelang.autotune را ارائه میدهد. این ابزار جستوجوی بهینه برای زمانبندی (Schedule) را خودکار میکند. کاربران یک فضای جستوجو تعریف میکنند که شامل موارد زیر است:
- ابعاد کاشی: اندازههای بلوک M، N و K (مثلاً پیمایش M/N از ۶۴ تا ۲۵۶ و K از ۳۲ تا ۶۴).
- عمق خط لوله: تعداد مراحل برای جابهجایی غیرهمزمان دادهها (مثلاً ۲ یا ۳).
- تعداد رشتهها: تعداد رشتهها در هر بلوک (مثلاً ۱۲۸ یا ۲۵۶).
اتوتیونر پیکربندیهایی که از SMEM_CAP (بودجه حافظه مشترک) فراتر میروند را حذف میکند. سپس هر کاندید را کامپایل، بنچمارک و از نظر دقت عددی اعتبارسنجی میکند. برنده در مسیر ~/.tilelang/cache ذخیره میشود تا اجراهای بعدی آنی باشند. این حافظه پنهان را میتوان از طریق متغیر محیطی TILELANG_AUTO_TUNING_DISABLE_CACHE=1 غیرفعال کرد.
برای رفع مشکل «جعبه سیاه» بودن کرنلهای کامپایلشده، ابزارهای بازرسی دقیقی تعبیه شده است. دستور T.print اجازه چاپ کنترلشده در سمت دستگاه را میدهد که برای تأیید مقادیر رجیستر در حین اجرای کرنل مفید است. توسعهدهندگان میتوانند با get_kernel_source() کد واقعی CUDA تولید شده را مشاهده کنند و کلمات کلیدی مانند __global__ ،extern "C" ،mma ،cp.async ،__syncthreads و ldmatrix را برای تأیید صحت دستورات جستوجو کنند.
سایر ابزارهای حیاتی عبارتند از:
- پروفایلر: متد
kernel.get_profiler().do_bench()ورودیهای مصنوعی میسازد تا تأخیر خام (Raw latency) را اندازهگیری کند. - کد میزبان:
kernel.get_host_source()پوشش (Wrapper) اجرای CUDA را نشان میدهد. - تأییدها (Assertions):
T.device_assert(cond, msg)برای بررسی خطاهای زمان اجرا در GPU. - بصریسازی:
tilelang.tools.plot_layoutبرای مشاهده بصری چیدمان رجیسترها و حافظه مشترک.
TileLang مجموعهای جامع از توابع اولیه را ارائه میدهد. برای مدیریت حافظه، T.alloc_shared برای حافظه مشترک، T.alloc_fragment برای رجیسترها و T.alloc_barrier برای mbarrierهای سبک Hopper ارائه شده است. جابهجایی دادهها از طریق T.copy (برداری)، T.async_copy (برای cp.async صریح) و T.tma_copy برای انتقالهای حجیم غیرهمزمان در سختافزارهای جدید انجام میشود.
عملیات محاسباتی از T.gemm و T.gemm_sp (برای پراکندگی ساختاریافته ۲:۴) تا اسکنهای T.cumsum و T.cummax را شامل میشود. ساختارهای حلقه به سه دسته تقسیم میشوند: T.Parallel برای عملیات عنصر-به-عنصر، T.Pipelined برای خطلولهسازی نرمافزاری و T.serial برای منطق متوالی.
این گردش کار، بار بهینهسازی GPU را از کدنویسی دستی C++ به طراحی الگوریتمیک سطح بالا منتقل میکند. با خودکارسازی نگاشت کاشیها به سختافزار، TileLang تکرار سریعتر را برای پژوهشگرانی که نسل بعدی اپراتورهای بهینه AI را میسازند، ممکن میکند. برای بررسی بیشتر، توسعهدهندگان میتوانند پازلهای TileLang را در گیتهاب یا مرجع کامل API را در tilelang.com مشاهده کنند. برای کاربران پیشرفته، نمونههایی از پاسهای Backward در FlashAttention، ضرب ماتریسی W4A16 با ترفندهای LOP3 و پیادهسازیهای Decode در DeepSeek MLA در مخزن کد موجود است.
گام بعدی شما
- اگر در حال توسعه اپراتورهای سفارشی برای مدلهای LLM هستید، TileLang را جایگزین نوشتن دستی CUDA C++ کنید.
- از قابلیت Autotuning برای یافتن بهینهترین ابعاد کاشی (Tile Dimensions) متناسب با کارت گرافیک خود استفاده کنید.
- کدهای نمونه FlashAttention در گیتهاب TileLang را برای درک نحوه ادغام عملیات (Fusion) بررسی کنید.
اما داستان سختافزاری این تحول در تراشههای جدیدتر حتی شگفتانگیزتر است — به تحلیل ما درباره قابلیتهای TMA در معماری Hopper مراجعه کنید.
گفتگو