VectorWare کد SIMD زبان Rust را مستقیماً روی GPU اجرا کرد

تصور کنید یک روتین ریاضی پیچیده را یک‌بار در Rust بنویسید و آن را بدون تغییر حتی در یک خط کد، هم روی لپ‌تاپ x86 و هم روی یک GPU انویدیا اجرا کنید. این رویایی که سال‌ها برای برنامه‌نویسان سیستم دور از دسترس بود، اکنون با دستاوردی از VectorWare به واقعیت تبدیل شده است.

بر اساس گزارش فنی مورخ ۱۰ اوت ۲۰۲۶، این تیم موفق شده است SIMD قابل‌حمل (Portable SIMD) را در سطح core::simd روی GPU پیاده کند. برای سال‌ها، توسعه‌دهندگان با یک انتخاب سخت مواجه بودند: یا کدهای اسکالر (Scalar) می‌نوشتند که قابل‌حمل اما کند بود، یا کدهای SIMD با سرعت خیره‌کننده می‌نوشتند که تنها روی یک معماری خاص CPU کار می‌کرد. از نظر تاریخی، نوشتن SIMD در Rust به معنای استفاده از دستورات اختصاصی هر سازنده (Vendor Intrinsics) در core::arch بود؛ برای مثال استفاده از _mm256_add_ps برای معماری x86-64 یا vaddq_f32 برای Arm. این دستورات به یک مجموعه دستورالعمل (Instruction Set) واحد گره خورده‌اند، به این معنی که هر برنامه‌ای که بخواهد چندین معماری را هدف قرار دهد، نیاز به پیاده‌سازی‌های مجزا برای هر یک داشت.

قابلیت SIMD قابل‌حمل در Rust برای حل این مشکل در CPUها طراحی شد. این کار از طریق ارائه یک نوع جنریک Simd<T, N> صورت گرفت — که نشان‌دهنده برداری از N المان از نوع T است — و کامپایلر آن را به دستورات خاص سخت‌افزار هدف تبدیل می‌کند. پیشرفت VectorWare در این است که با GPU نه به عنوان یک موجود بیگانه و مجزا، بلکه به عنوان یک سخت‌افزار برداری دیگر می‌نگرد که این انتزاع می‌تواند آن را هدف قرار دهد.

سازوکار: تبدیل SIMT به SIMD

اساس این پیاده‌سازی بر این مشاهده استوار است که مدل SIMT (Single Instruction, Multiple Thread) در کارت‌های گرافیک انویدیا، در واقع یک واحد برداری بسیار عریض است. در مدل SIMT، یک «وارپ» (Warp) یک دستور صادر می‌کند و هر یک از ۳۲ مسیر (Lane) آن، آن دستور را روی داده‌های خود اجرا می‌کنند. این دقیقاً همان تعریف SIMD است: یک دستور که روی داده‌های متعدد عمل می‌کند.

• نگاشت مستقیم: یک نوع Simd<i16, 32> در Rust به‌طور یک‌به‌یک روی یک وارپ ۳۲-مسیره GPU قرار می‌گیرد. در این حالت، هر مسیر دقیقاً یک المان i16 را دریافت می‌کند.
• تبدیل دستورات: جمع دو بردار از این نوع، به یک دستور واحد add.s16 در سطح وارپ روی GPU کامپایل می‌شود. این دقیقاً مشابه روشی است که همان کد Rust در یک CPU از نوع x86 به دستور vpaddw تبدیل می‌شد.
• استقلال از کتابخانه استاندارد: چون SIMD قابل‌حمل در core قرار دارد و نه در std (کتابخانه استاندارد)، نیازی به پشتیبانی کامل از کتابخانه استاندارد که VectorWare پیش‌تر برای GPU فراهم کرده بود، ندارد.

سلسله‌مراتب موازی‌سازی

این نگاشت، سلسله‌مراتب گسترده‌تری از موازی‌سازی را که توسط VectorWare توسعه یافته، تکمیل می‌کند. پیش از این، آن‌ها std::thread را به یک وارپ GPU نگاشت کرده بودند که اجازه می‌داد بسیاری از رشته‌های هم‌زمان روی GPU اجرا شوند. با این حال، آن رویکرد از مسیرهای موازی (Parallel Lanes) درون هر رشته یا وارپ استفاده نمی‌کرد.

در CPU، انتزاع برای موازی‌سازی درون یک رشته، همان SIMD است. یک دستور واحد روی چندین المان داده که در یک واحد برداری بسته‌بندی شده‌اند، عمل می‌کند. در حالی که کد اسکالر دو عدد را جمع می‌کند، یک جمع SIMD دو بردار (مثلاً شامل هشت مقدار f32) را می‌گیرد و هشت مجموع را به‌طور هم‌زمان تولید می‌کند. این موازی‌سازی داده‌ها در داخل یک رشته واحد و در سطحی پایین‌تر از جایی که سیستم‌عامل زمان‌بندی را مدیریت می‌کند، وجود دارد.

با اضافه کردن core::simd ، VectorWare یک مدل سازگار در تمامی پلتفرم‌ها ایجاد کرده است. در CPU، یک رشته شامل مسیرهای SIMD است. در GPU، std::thread در واقع یک وارپ است که مسیرهای سخت‌افزاری آن همان نقش را ایفا می‌کنند. در هر دو محیط، core::simd محرک این مسیرهاست.

اجرای عملیات پیچیده

تیم VectorWare به محاسبات ساده بسنده نکرد. آن‌ها تمام مجموعه عملیات SIMD را به دستورات اولیه GPU نگاشت کردند و اجازه دادند روتین‌های پیچیده‌ای مانند relu_dot به‌طور بومی روی دستگاه اجرا شوند. این روتین ویژگی‌های کلیدی را به چالش می‌کشد: محاسبات عنصر-به-عنصر، مقایسه‌ای که یک ماسک مسیر تولید می‌کند، یک انتخاب (Select) که توسط آن ماسک هدایت می‌شود و یک کاهش افقی (Horizontal Reduction) در عرض مسیرها.

• عملیات عنصر-به-عنصر: جمع، ضرب و مقایسه‌ها مستقیماً به اپراتورهای بومی مسیر-محور GPU نگاشت می‌شوند. برای مثال، یک ضرب با عرض ۳۲، سی و دو حاصل‌ضرب را به‌طور هم‌زمان محاسبه می‌کند. این‌ها از پیاده‌سازی‌های معمولی Trait در Rust روی Simd (مانند Add) می‌آیند.
• کاهش (Reductions): عملیاتی مانند reduce_sum و reduce_max هر مسیر را در یک اسکالر واحد ترکیب می‌کنند. این عملیات از دستورات Shuffle وارپ در GPU برای تبادل و ترکیب مقادیر بین مسیرها استفاده می‌کنند تا نتیجه اسکالر یکسانی در هر مسیر تولید شود.
• جابجایی‌های بین-مسیری (Cross-lane Shuffles): دستورات simd_swizzle! و چرخش‌ها (Rotates)، المان‌ها را بین مسیرها جابجا می‌کنند. این‌ها روی همان دستورات اولیه Shuffle وارپ نگاشت می‌شوند که مسیرهای GPU را در تبادل داده‌ها کارآمد می‌کنند.
• ماسک‌گذاری (Masking): انواع Mask<T, N> برای هر مسیر SIMD یک پیش‌شرط (Predicate) فراهم می‌کنند. Mask::select یک انتخاب در هر مسیر وارپ انجام می‌دهد، در حالی که پرس‌وجوهای افقی مانند any و all از دستورات Vote و Ballot در GPU استفاده می‌کنند.

داده‌های یکنواخت در مقابل متغیر (Uniform vs Varying)

این پیاده‌سازی از تمایزی بهره می‌برد که مشابه زبان‌های موازی-داده مانند ISPC است. در مدل VectorWare، مقادیر اسکالر در کد اطراف — مانند یک مقدار ثابت یا شمارنده حلقه — به‌طور یکسان توسط هر مسیر محاسبه می‌شوند. این‌ها مقادیر «یکنواخت» (Uniform) هستند و به‌سادگی در عرض وارپ تکثیر می‌شوند.

در مقابل، یک Simd<f32, 32> داده‌ای «متغیر» (Varying) است. این تمایز به‌طور طبیعی از سیستم تایپ خود Rust حاصل می‌شود: یک f32 ساده، یکنواخت است، در حالی که نوع Simd متغیر است.

مدیریت تفاوت‌های سخت‌افزاری

یک چالش قابل توجه این است که عرض SIMD در CPUها متغیر است (از ۱ تا ۶۴)، در حالی که GPUها عرض‌های ثابتی دارند: ۳۲ مسیر در انویدیا و ۳۲ یا ۶۴ مسیر در AMD. نگاشت تنها زمانی یک‌به‌یک است که N با عرض سخت‌افزار مطابقت داشته باشد. یک N کوچک‌تر باعث می‌شود برخی مسیرها بیکار بمانند، در حالی که یک N بزرگ‌تر مستلزم آن است که مسیرها بیش از یک المان را پردازش کنند.

برای مدیریت این موضوع، VectorWare وارپ را به عنوان یک «ماشین» کوچک با مجموعه‌ای ثابت از دستورات اولیه برای جابجایی و ترکیب داده‌ها در نظر می‌گیرد. آن‌ها به این ماشین یک نمایش میانی (IR) می‌دهند که در سیستم تایپ Rust با استفاده از موارد زیر کدگذاری شده است:

• ابزارهای سیستم تایپ: Generics، Const Generics و Trait Bounds.
• عملیات تایپ‌شده: Ballotها، Shuffleها، Reductionها، Scanها، Gatherها، Scatterها، Atomics و Strip Mining برای بردارهایی که عریض‌تر از وارپ هستند.
• ایمنی: از آنجایی که عملوندها، شکل اجرا و ظرفیت، همگی تایپ‌شده هستند، بسیاری از برنامه‌های نامعتبر اصلاً قابل ساخت نیستند.

این IR نیازی به مفسر روی GPU ندارد؛ بلکه مستقیماً به دستورات تبدیل می‌شود و هیچ هزینه اضافی نسبت به کد PTX دست‌نویس ندارد. برای تست، VectorWare یک مفسر مرجع — یک «Miri برای برنامه‌نویسی مسیر-وارپ» — ساخت تا کدهای GPU را شبیه‌سازی کرده و تست‌های تفاضلی (Differential Testing) انجام دهد.

موازنه و محدودیت‌ها

با وجود افزایش کارایی، این رویکرد محدودیت‌هایی دارد. SIMD قابل‌حمل همچنان یک ویژگی ناپایدار در Rust است که به کامپایلر Nightly و پرچم #! [feature(portable_simd)] نیاز دارد. ساختار آن ممکن است پیش از تثبیت تغییر کند.

کارایی تنها زمانی «بدون هزینه» (Zero-cost) است که عرض بردار کاملاً با مسیرهای وارپ مطابقت داشته باشد. بردارهای باریک‌تر سخت‌افزار را بیکار می‌گذارند، در حالی که بردارهای عریض‌تر هر عملیات را به چندین دستور تبدیل می‌کنند. علاوه بر این، هر عملیات بین-مسیری کارآمد نیست؛ در حالی که الگوهای Shuffle پشتیبانی‌شده ارزان هستند، جایگشت‌های دلخواه (Arbitrary Permutations) ممکن است به چندین دستور یا عبور از حافظه مشترک (Shared Memory) نیاز داشته باشند.

علاوه بر این، عملیات افقی مانند Reductionها و all/any به عنوان نقاط همگام‌سازی در داخل وارپ عمل می‌کنند. این موضوع محدود می‌کند که زمان‌بند (Scheduler) تا چه حد می‌تواند کارها را هم‌پوشانی کند. VectorWare همچنین اشاره کرد که آن‌ها مجبور شدند کامپایلر را تغییر دهند تا اطمینان حاصل کنند که این انتزاع هنگام تعامل با سایر ویژگی‌های Rust، سالم (Sound) باقی می‌ماند. از آنجایی که این یک قلمرو ناشناخته است، آن‌ها هنوز مطمئن نیستند که تمام موارد پوشش داده شده باشند.

تحلیل: تغییر پارادایم GPU

این توسعه، این فرض بنیادی را که برنامه‌نویسی GPU نیازمند یک طرز فکر یا زبان «خارجی» است، تغییر می‌دهد. با تبدیل GPU به یک پلتفرم معمولی Rust، VectorWare در حال کاهش سد ورود برای محاسبات با کارایی بالا (HPC) است.

برای توسعه‌دهنده، این بدان معناست که Borrow Checker، Lifetimes و ایمنی تایپ اکنون به عملیات برداری GPU تعمیم می‌یابند. شما دیگر مجبور نیستید برای دستیابی به حداکثر کارایی، از محیط امن Rust خارج شده و به «غرب وحشی» دستورات CUDA وارد شوید. اثر ثانویه این است که کتابخانه‌های موجود Rust که از SIMD قابل‌حمل استفاده می‌کنند، اکنون کاندیداهای فوری برای شتاب‌دهی GPU بدون نیاز به بازنویسی هستند.

مسیر پیش رو

تیم VectorWare اکنون به دنبال ترکیب SIMD با کارهای قبلی خود روی رشته‌های GPU و Async است. هدف آن‌ها ایجاد سلسله‌مراتبی است که در آن رشته‌ها کار را بین وارپ‌ها پخش کنند، core::simd داده‌ها را در مسیرها توزیع کند و Async هم‌زمانی بین آن‌ها را مدیریت نماید.

آن‌ها همچنین در حال بررسی امکان نگاشت SIMD با شکل ماتریسی روی Tensor Coreهای GPU و برداری‌سازی خودکار (Auto-vectorizing) حلقه‌های اسکالر Rust هستند تا کدها بدون نیاز به نوشتن صریح با core::simd ، موازی‌سازی در سطح وارپ را به دست آورند. مهندسان VectorWare به عنوان اعضای تیم کامپایلر Rust، قصد دارند بخش زیادی از این منطق را مستقیماً به درون خود کامپایلر منتقل کنند.

در حالی که کار فعلی NVIDIA را هدف قرار داده است، IR طراحی شده مستقل از معماری است. Wavefrontهای AMD و Subgroupهای Vulkan دستورات و معناشناسی مشابهی را ارائه می‌دهند و آن‌ها را به اهداف طبیعی در آینده تبدیل می‌کنند. VectorWare تأکید می‌کند که اگرچه در نهایت از زبان‌ها و محیط‌های اجرای متعددی پشتیبانی خواهند کرد، اما Rust به‌طور منحصر‌به‌فردی برای ساخت برنامه‌های بومی و قابل اعتماد GPU مناسب است.