‏وقتی یه کرنل CUDA اجرا می‌شه، دقیقاً چی می‌شه؟

نقطهٔ شروع، ساده‌ترین برنامهٔ ممکنه: یه کرنل که دو تا وکتور رو جمع می‌کنه و روی یه میلیون عدد اعشاری (1

__global__ void vadd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

اولین ایستگاه PTX ه؛ یه اسمبلی مجازی و دستگاه‌ناوابسته با رجیسترهای مجازی مثل %rd1 یا %f3. نویسنده می‌گه PTX از چیزی که فکر می‌کنی پرحرف‌تره: مثلاً ساختن یه آدرس سه تا دستور می‌خواد. دلیلش اینه که پوینترهای CUDA به‌صورت پیش‌فرض «عمومی»ان و می‌تونن به حافظهٔ global، shared یا local اشاره کنن؛ پس اول cvta.to.global تضمین می‌کنه پوینتر تو پنجرهٔ global‌ه، بعد mul.wide.s32 ایندکس رو در ۴ (اندازهٔ float) ضرب و هم‌زمان از ۳۲ به ۶۴ بیت گسترش می‌ده، و آخر add.s64 اون رو به آدرس پایه اضافه می‌کنه.

بعد نوبت ptxas ه که PTX رو به SASS — اسمبلی واقعی و مخصوص همون معماری — تبدیل می‌کنه. اینجا بهینه‌سازی‌ها خودشون رو نشون می‌دن: ده‌ودوازده رجیستر مجازی به هفت رجیستر واقعی جمع می‌شن، اون توالی mul.wide + add تو یه دستور IMAD.WIDE ادغام می‌شه و تبدیل‌های cvta کلاً محو می‌شن و تو آدرس‌دهی حل می‌شن. دستور S2R هم رجیسترهای ویژهٔ سخت‌افزار (ایندکس بلاک و ایندکس نخ) رو به رجیستر معمولی کپی می‌کنه تا بشه روشون حساب کرد.

یه نکتهٔ ظریف: آرگومان‌های کرنل — پوینترهای a، b، c و اندازهٔ n — به‌علاوهٔ هندسهٔ اجرا، تو بانک ثابت ۰ می‌شینن. چرا؟ چون این یه خوندن همگانیه: همهٔ نخ‌های گرید به پوینترهای یکسان نیاز دارن و کش حافظهٔ ثابت می‌تونه هر ۳۲ لِین رو یه‌جا سرو کنه. پر کردن این بانک هم کار ساختاری به اسم QMD‌ه که درایور موقع اجرا تحویل GPU می‌ده.

فایل cubin که SASS رو نگه می‌داره، یه فایل ELF ه — همون قالبی که لینوکس برای فایل‌های اجرایی معمولی استفاده می‌کنه. ابزار fatbinary این cubin رو کنار خودِ PTX می‌ذاره و یه «فت‌باینری» می‌سازه؛ یعنی باینری نهایی هم SASS آماده داره، هم PTX. علتش سازگاری روبه‌جلوئه: اگه باینری رو ببری روی GPUی که cubin پوشش نمی‌ده، درایور همون PTX رو موقع لود به SASS تازه JIT می‌کنه. چون PTX متن ساده و حجیمه، nvcc به‌صورت پیش‌فرض فشرده‌اش می‌کنه.

حالا سمت میزبان: کامپایلر یه سازندهٔ مخفی قبل از main تزریق می‌کنه که کارش ثبت فت‌باینری تو رانتایم CUDA و ساختن یه نگاشته — از پوینتر تابع میزبان به اسم مچاله‌شدهٔ کرنل روی دستگاه. اون نحو vadd هم با یه استاب جایگزین می‌شه که آرگومان‌ها رو تو یه بافر با آفست‌های ۰، ۸، ۱۶ و ۲۴ بسته‌بندی می‌کنه — دقیقاً همون آفست‌هایی که SASS از بانک ثابت می‌خوند. بعد __cudaLaunch آدرس تابع پوچِ میزبان رو به‌عنوان کلید جست‌وجو می‌بره سراغ جدول ثبت، اسم کرنل دستگاه رو پیدا می‌کنه و می‌ره تو درایور یوزرمود .

نکتهٔ جالب اینه که تا اینجا هنوز هیچ کدی به GPU نرسیده. از CUDA 12.2 به بعد، لود شدن ماژول به‌صورت پیش‌فرض تنبل (lazy) ه: درایور آپلود کردن SASS به حافظهٔ کارت رو تا اولین باری که همون کرنل واقعاً اجرا شه عقب می‌ندازه. زیر libcuda هم درایور کرنل‌مود یعنی nvidia.ko نشسته که ارتباط باهاش از راه ioctl روی فایل‌های دستگاه انجام می‌شه.

و حرف اصلی بخش آخر: GPU مثل CPU «فراخوانی تابع» نمی‌فهمه. نه نقطهٔ ورودی داره که بپری توش، نه پشته‌ای که آرگومان روش هل بدی. کارت اون‌طرف باس PCIe نشسته و یه جریان از دستورهای درایور رو از حافظهٔ میزبان می‌خونه؛ پس تمام کاری که cuLaunchKernel از این به بعد می‌کنه اینه که یه دستور اجرای کامل رو وارد اون جریان کنه و بعد به GPU خبر بده که دستور آماده‌ست.

نکات کلیدی:

• مسیر کد: C/CUDA → PTX (مجازی و دستگاه‌ناوابسته) → SASS (اسمبلی واقعی همون معماری).
• فت‌باینری هم SASS و هم PTX رو با خودش می‌بره؛ PTX نقش نسخهٔ پشتیبان برای JIT روی GPUهای ناشناخته رو داره.
• آرگومان‌های کرنل تو بانک ثابت ۰ می‌شینن چون همهٔ نخ‌ها همزمان همون مقادیر رو می‌خونن.
• کامپایلر یه سازندهٔ مخفی و یه استاب میزبان می‌سازه تا به یه فراخوانی واقعی درایور تبدیل شه.
• از CUDA 12.2، لود ماژول پیش‌فرض تنبله؛ SASS تازه موقع اولین اجرای همون کرنل به کارت می‌ره.
• GPU فراخوانی تابع نمی‌گیره؛ فقط جریان دستورها رو از حافظهٔ میزبان روی PCIe می‌خونه.