مبانی سرویس‌دهی LLM: GPU

جلسه‌ی قبل گفت prefill «محدود به محاسبه» و decode «محدود به پهنای باند» است. این جلسه توضیح می‌دهد این جمله‌ها دقیقاً یعنی چه و چطور با یک حساب سرانگشتی می‌شود زمان اجرا را پیش‌بینی کرد. هدف شهود است، نه برنامه‌نویسی CUDA.

GPU یک ماشین توان عملیاتی است#

CPU مثل یک خودروی اسپرت است: برای دو مسافر سریع‌ترین است. GPU مثل ناوگان اتوبوس است: کندتر، اما هزاران نفر را جابه‌جا می‌کند. LLM دقیقاً کار اتوبوسی است: برای هر توکن همان ضرب‌های ماتریسی روی همان وزن‌ها به‌طور موازی انجام می‌شود.

warp و SIMT#

یک warp گروهی از ۳۲ نخ است که یک دستور را هم‌زمان روی داده‌های مختلف اجرا می‌کنند (SIMT: یک دستور، چند نخ):

// each thread computes one element
i = threadIdx.x;
y[i] = a * x[i] + b;

یک بار واکشی، یک بار رمزگشایی و یک تصمیم زمان‌بندی، ۳۲ مسیر را جلو می‌برد. هزینه‌اش واگرایی شاخه است:

if (i % 4 == 0)  a = f(x[i]);   // 8 of 32 threads
else             a = g(x[i]);   // 24 of 32 threads

warp هر دو مسیر را پشت سر هم اجرا می‌کند و در هر مسیر بخشی از نخ‌ها بیکارند.

پنهان کردن تأخیر با جابه‌جایی warp#

خواندن از حافظه‌ی اصلی صدها چرخه طول می‌کشد. CPU این تأخیر را با کش‌های بزرگ و پیش‌بینی پنهان می‌کند. GPU روش دیگری دارد: هر SM تا ۶۴ warp را هم‌زمان نگه می‌دارد و وقتی یکی منتظر حافظه است، بی‌هزینه به warp دیگری که آماده است می‌پرد. چون حالت همه‌ی نخ‌ها مقیم است، جابه‌جایی رایگان است.

سه مشخصه که هر GPU را محدود می‌کنند#

• محاسبه: ریاضی با چه سرعتی انجام می‌شود. موتور خودرو.
• پهنای باند: وزن‌ها و KV با چه سرعتی خوانده می‌شوند. لوله‌ی سوخت.
• ظرفیت: چند وزن و چند مدخل KV جا می‌شوند. باک.

قاعده‌ی طلایی#

kernel_time ≈ max( bytes_moved / memory_bandwidth ,  FLOPs / peak_FLOPs )

زمان اجرای هر kernel تقریباً بیشینه‌ی دو عدد است: زمانی که جابه‌جایی داده طول می‌کشد و زمانی که محاسبه طول می‌کشد. هر کدام بزرگ‌تر بود، همان گلوگاه است.

دو مثال با Llama 3 8B روی یک H200:

decode با batch برابر ۱. برای یک توکن باید حدود ۱۶ گیگابایت وزن از HBM خوانده شود:

زمان خواندن:  16 GB ÷ 4.8 TB/s ≈ 3.3 ms
زمان محاسبه:  ~16 GFLOP ÷ 989 TFLOP/s ≈ 0.016 ms

هسته‌های محاسباتی ۹۹٫۵٪ وقت بیکارند. هر وزن فقط یک بار استفاده می‌شود و ریاضی منتظر حافظه می‌ماند.

prefill با ۲۰۴۸ توکن. هر وزنی که خوانده می‌شود برای ۲۰۴۸ توکن استفاده می‌شود:

محاسبه: 2 × 8B × 2048 ≈ 33 TFLOP  ←  33 TFLOP ÷ 989 TFLOP/s ≈ 33 ms

خواندن زود تمام می‌شود و حالا ریاضی سرعت را تعیین می‌کند.

مدل roofline#

شدت حسابی یعنی چند FLOP به ازای هر بایتی که از HBM جابه‌جا می‌شود:

شدت حسابی = FLOP انجام‌شده ÷ بایت جابه‌جاشده
نقطه‌ی شکست (ridge) = محاسبه‌ی بیشینه ÷ پهنای باند
                    = 989 TFLOP/s ÷ 4.8 TB/s ≈ 206 FLOP/B   (برای H200)

  performance
  (FLOP/s)
     │                     ┌──────────────── compute roof: 989 TFLOP/s
     │                    ╱
     │                   ╱   compute-bound: prefill, large GEMM
     │                  ╱
     │                ╱
     │              ╱   memory-bound: decode, GEMV
     │            ╱     speed = bandwidth × intensity
     │          ╱
     └────────────────────┬────────────────── arithmetic intensity (FLOP/B)
                         206

سمت چپ نقطه‌ی شکست، سرعت را پهنای باند تعیین می‌کند (decode). سمت راست، سقف محاسبه تعیین می‌کند (prefill و ضرب ماتریس‌های بزرگ).

شدت حسابی عملیات رایج (b = بایت هر عنصر؛ در BF16 برابر ۲):

تمرین: H200 در decode با batch برابر ۱ چقدر مشغول است؟#

شدت حدود ۱ FLOP بر بایت است، پس:

کارایی ≈ 4.8 TB/s × 1 FLOP/B = 4.8 TFLOP/s   ←   4.8 ÷ 989 ≈ 0.5%  از توان محاسباتی

نیم درصد. decode با batch کوچک بیش از ۹۹٪ توان GPU را هدر می‌دهد. راه‌حل‌ها از خود فرمول می‌آیند:

• batch بزرگ‌تر: هر وزنی که خوانده می‌شود برای چند درخواست استفاده می‌شود و شدت بالا می‌رود
• کوانتیزه‌سازی: بایت کمتر برای هر وزن
• کم کردن ترافیک KV: GQA، MLA و فشرده‌سازی (جلسه‌ی قبل)

آیا batching در decode کافی است؟#

با batch برابر B، هر وزنی که خوانده می‌شود B بار استفاده می‌شود، پس شدت حسابی ضرب‌های وزن تقریباً B برابر می‌شود. نموداری در اسلاید شدت را برای batch از ۱ تا ۲۵۶ در BF16 رسم می‌کند: با بزرگ شدن batch، ضرب‌های وزن به نقطه‌ی شکست ۲۰۶ نزدیک می‌شوند. اما attention در decode روی 2/b می‌ماند، چون KV هر درخواست مال خودش است و فقط یک بار استفاده می‌شود. با بلند شدن کانتکست، خواندن KV بر هزینه غالب می‌شود:

Llama 3 8B، batch ۱۶، BF16 (وزن‌ها ۱۶٫۱ گیگابایت، KV حدود ۱۲۸ کیلوبایت به ازای هر توکن):

در ۳۲K، KV چهار برابر خود وزن‌هاست. برای ایجنت‌ها با کانتکست صدها هزار توکنی، decode تقریباً تماماً خواندن KV است.

def decode_step_ms(weights_gb, kv_kib_per_token, batch, context, bw_tb_s=4.8):
    """تخمین زمان یک قدم decode وقتی محدود به پهنای باند است."""
    kv_gb = batch * context * kv_kib_per_token * 1024 / 1e9
    return (weights_gb + kv_gb) / (bw_tb_s * 1e3) * 1e3      # GB ÷ (GB/ms)

for ctx in (1024, 32 * 1024, 200 * 1024):
    print(ctx, round(decode_step_ms(16.1, 128, 16, ctx), 1), "ms")
# 1024 3.8 ms | 32768 17.7 ms | 204800 92.8 ms

در خط آخر، KV به تنهایی حدود ۴۳۰ گیگابایت است و اصلاً در حافظه‌ی ۱۴۱ گیگابایتی H200 جا نمی‌شود. یعنی پیش از آنکه سرعت مسئله شود، ظرفیت مسئله می‌شود.

دقت پایین‌تر، نقطه‌ی شکست را جابه‌جا می‌کند#

در FP32 نقطه‌ی شکست خیلی پایین است: حدود ۹٫۶ در A100، ۲۰ در H100، ۱۴ در H200 و ۹٫۴ در B200. یعنی kernelهای FP32 تقریباً همیشه محدود به حافظه‌اند. در طرف دیگر، NVFP4 متراکم در Rubin حدود ۱۶۰۰ FLOP به ازای هر بایت لازم دارد تا محدود به محاسبه شود.

(اعداد جدول نقطه‌ی شکست بر حسب FLOP بر بایت‌اند.) از A100 تا Rubin، توان محاسباتی سریع‌تر از پهنای باند رشد کرده است. پس هر نسل، kernelهای بیشتری محدود به حافظه می‌شوند. بهینه‌سازی جابه‌جایی داده روزبه‌روز مهم‌تر می‌شود.

درون H200#

┌─────────────────── H200 ───────────────────┐
│  132 × SM   ┌────┐┌────┐ ... ┌────┐        │   each SM: 4 partitions, each with
│             │ SM ││ SM │     │ SM │        │   a warp scheduler, 32 FP32 cores
│             └────┘└────┘     └────┘        │   and one Tensor Core
│                                            │   registers 256 KB, shared/L1 256 KB
│                 L2: 50 MB                  │
│        HBM3e: 141 GB at 4.8 TB/s           │──── NVLink 900 GB/s to other GPUs
└────────────────────────────────────────────┘

Tensor Core#

یک دستور MMA در سطح warp حاصل D = A × B + C را روی کاشی‌هایی مثل ۱۶×۱۶ و ۱۶×۸ حساب می‌کند: ۲۰۴۸ ضرب و جمع با یک دستور. بیش از ۹۹٪ FLOPهای LLM ضرب ماتریس است و روی Tensor Core اجرا می‌شود. softmax و نرمال‌سازی روی هسته‌های معمولی FP32 اجرا می‌شوند. برای مقایسه، H200 با Tensor Core در BF16 به ۹۸۹ ترافلاپ می‌رسد و با هسته‌های FP32 معمولی فقط به ۶۷.

دلیل سرعت Tensor Core: مسیر داده‌ی تخصصی برای کل کاشی، دستور کمتر به ازای هر FLOP، استفاده‌ی دوباره از عملوندها و دقت پایین‌تر با جمع‌کردن در دقت بالاتر.

سلسله‌مراتب حافظه#

FlashAttention که در جلسه‌ی قبل دیدیم دقیقاً از همین سلسله‌مراتب استفاده می‌کند: کاشی‌ها را در حافظه‌ی مشترک نگه می‌دارد تا ماتریس بزرگ امتیازها هرگز به HBM نرود.

حافظه به‌عنوان یک لوله: قانون Little#

پهنای باند = بایت‌های در جریان ÷ تأخیر

اگر تأخیر حافظه حدود ۲۵۰ نانوثانیه باشد، برای رسیدن به ۴٫۸ ترابایت بر ثانیه باید 4.8 TB/s × 250 ns ≈ 1.2 MB داده هم‌زمان در راه باشد، یعنی حدود ۹ کیلوبایت برای هر SM. یک kernel ساده که هر نخ ۴ بایت می‌خواند فقط حدود ۴ کیلوبایت در هر SM در راه دارد و به حدود ۴۵٪ پهنای باند می‌رسد.

طول لوله را سخت‌افزار تعیین می‌کند، پس تنها راه پهنای باند بیشتر، داده‌ی بیشتر در جریان است: بارگذاری‌های ۱۶ بایتی، چند بارگذاری مستقل در هر نخ، و کپی‌های گروهی ناهمگام. هر نسل جدید این عدد را بالاتر می‌برد و ویژگی‌های جدید هر نسل هم بیشتر برای همین ساخته شده‌اند:

ابزارهای Hopper و Blackwell برای داده‌ی بیشتر در جریان (پیشرفته)#

بزرگ‌تر از یک GPU#

• مقیاس به بالا: ۸ GPU در یک سرور با NVSwitch (در H200، NVLink با ۹۰۰ GB/s برای هر GPU)، یا ۷۲ GPU در یک رک GB200 NVL72
• مقیاس به بیرون: شبکه‌ی RDMA بین سرورها (حدود ۵۰ GB/s برای هر GPU در H200)

فاصله‌ی HBM (۴٫۸ TB/s) تا NVLink (۰٫۹ TB/s) و تا شبکه (۰٫۰۵ TB/s) نشان می‌دهد چرا تقسیم مدل روی چند GPU هزینه دارد و چرا جای قرار گرفتن هر بخش مهم است. این موضوع جلسه‌ی موازی‌سازی است.

دو دهه GPU انویدیا (پیشرفته)#

سال‌های ۲۰۰۶ تا ۲۰۱۷ GPU را از گرافیک به محاسبه‌ی عمومی برد و ۲۰۱۸ تا ۲۰۲۶ آن را به سمت موتور هوش مصنوعی با دقت پایین. از G80 تا B300 توان محاسباتی بیشینه حدود ۴۳ هزار برابر شده و پهنای باند فقط حدود ۹۰ برابر. این همان دلیلی است که kernelها روزبه‌روز محدود به حافظه‌تر می‌شوند.

شتاب‌دهنده‌های دیگر، در یک نگاه#

مقایسه با B200:

شتاب‌دهنده‌هایی که بر حافظه‌ی روی تراشه (SRAM) تکیه می‌کنند:

جمع‌بندی#

1. GPU سرعت تک‌نخی را با هزاران هسته‌ی موازی معاوضه می‌کند.
2. محاسبه، پهنای باند و ظرفیت هر kernel را محدود می‌کنند.
3. نقطه‌ی شکست roofline (۲۰۶ FLOP بر بایت در H200) kernelهای محدود به حافظه را از kernelهای محدود به محاسبه جدا می‌کند.
4. batching خواندن وزن‌ها را به نقطه‌ی شکست نزدیک می‌کند، اما خواندن KV محدود به حافظه می‌ماند.
5. نسل‌های جدید نقطه‌ی شکست را به راست می‌برند، پس kernelهای بیشتری محدود به پهنای باند می‌شوند.

اتصال به بخش اول مجموعه: وقتی در جلسه‌ی ۸ گفتیم «کانتکست کوچک‌تر نگه دارید»، این فقط صرفه‌جویی در صورت‌حساب API نیست. در سطح سخت‌افزار، هر توکن اضافه در کانتکست یعنی بایت‌های بیشتری که در هر قدم decode از HBM خوانده می‌شوند.

پرسش‌هایی برای فکر کردن#

• اگر Llama 3 8B را با وزن‌های ۴ بیتی اجرا کنید، زمان یک قدم decode با batch برابر ۱ تقریباً چقدر می‌شود؟ آیا decode دیگر محدود به حافظه نیست؟
• یک ایجنت با کانتکست ۲۰۰ هزار توکنی روی Llama 3 8B و batch ۱۶: چند درصد بایت‌های هر قدم مال وزن‌هاست؟ prefix caching چه کمکی می‌کند و چه کمکی نمی‌کند؟

منابع#

• اسلاید جلسه: Efficient LLM serving basics II: intro to GPUs (10_llm_serving_basics_2a.pdf)
• CUDA C++ Programming Guide
• Williams، Waterman و Patterson، Roofline: An Insightful Visual Performance Model
• Luo و همکاران، Dissecting the NVIDIA Hopper Architecture
• NVIDIA، H200 datasheet