🛌 ☃️ 👸🏻 الكتابة في Rust + CUDA C ✖️ 👨🏽‍✈️ 🎼

مرحبا بالجميع!

في هذا الدليل ، أرغب في معرفة كيفية تكوين صداقات CUDA C / C ++ و Rust. وكمثال على ذلك ، سوف نكتب برنامجًا صغيرًا في Rust لحساب المنتج القياسي للناقلات ، وسيتم حساب المنتج القياسي على GPU باستخدام CUDA C.

من يهتم القط!

كودا ج

الخطوة الأولى هي تثبيت برنامج التحويل البرمجي CUDA - nvcc. ما هي CUDA ولماذا أحتاج إلى وصفها ، يمكنني أن أقرأ عنها هنا ، على سبيل المثال. لا أستطيع إلا أن أقول أنه بمساعدتها ، يمكنك كتابة التعليمات البرمجية التي سيتم تشغيلها على بطاقات فيديو NVIDIA (المشار إليها فيما يلي باسم GPU) واستخدام كل قوتها في الحوسبة المتوازية ومعالجة الرسومات. مرة أخرى ، لا يتعلق هذا البرنامج التعليمي بكيفية كتابة التعليمات البرمجية في CUDA ، ولكن حول كيفية استخدام ميزاته من كود Rust وكتابة الحوسبة المتوازية على GPU.

لذا قم بتثبيت nvcc ومجموعة أدوات CUDA. مع هذا التعقيد ، لا ينبغي أن تنشأ تعليمات مفصلة: في خارج الموقع .

RUST + CUDA C

خلال هذا البرنامج التعليمي ، كما ذكرنا سابقًا ، سنكتب برنامجًا في Rust للبحث عن منتج العددية الخاص بمتجهين ، وستحدث عملية الحساب نفسها على وحدة معالجة الرسومات.

المنتج العددية من اثنين من المتجهات.

لنفترض أن لدينا متجهين:

a = [a_{1} ، a_{2} ، . . . a_{n}]

$a = [a_1، a_2، ... a_n]$ و

b = [b_{1} ، b_{2} ، . . . ، b_{n}]

$b = [b_1 ، b_2 ، ... ، b_n]$ ، المنتج العددية لهذه المتجهات:

a c d o t b = s u m_{i = 1}^{n} a_{i} b_{i}

$a \ cdot b = \ sum_ {i = 1} ^ {n} {a_ib_i}$

لنبدأ في إنشاء برنامجنا. علاوة على ذلك ، أفترض أن برنامج nvcc قد تم تثبيته بنجاح ، كما أن rustc والبضائع تحمل أيضًا رمز تجميع للصدأ.

أولاً ، قم بإنشاء مجلد المشروع. في مجلد المشروع ، قم بإنشاء ملف Cargo.toml ، والذي يحتوي على إرشادات لمجمع الشحنات. الملف يشبه هذا:

[package] name = "rust_cuda" #   version = "0.1.0" #   authors = ["MoonL1ght <ixav1@icloud.com>"] #    build = "build.rs" #    rust links = "cudart" #  cuda,    [dependencies] libc = "0.2" #  rust     rand = "0.5.5" #  rust      [build-dependencies] cc = "1.0" # rust

أيضًا في مجلد جذر المشروع ، قم بإنشاء ملف build.rs الذي سيحتوي على تعليمات لبناء برنامج الصدأ وتجميع كود CUDA.

أضف مجلد src إلى جذر المشروع الذي سنضع فيه ملفات التعليمات البرمجية المصدر. في مجلد src ، قم بإنشاء أربعة ملفات: main.rs - رمز البرنامج الرئيسي ، dot.cpp - الربط C ++ (المجمع لـ CUDA C) ، dot_gpu.h ، dot_gpu.cu - الملف الذي يحتوي على التعليمات البرمجية المنفذة على GPU.

المجموع لدينا مثل هيكل المشروع:

 rust-cuda/ src/ main.rs dot.cpp dot_gpu.h dot_gpu.cu Cargo.toml build.rs

في ملف build.rs ، الشيء الأكثر أهمية هو كتابة هذا:

 println!("cargo:rustc-link-search=native=/Developer/NVIDIA/CUDA-10.1/lib"); println!("cargo:rustc-link-search=/Developer/NVIDIA/CUDA-10.1/lib"); println!("cargo:rustc-env=LD_LIBRARY_PATH=/Developer/NVIDIA/CUDA-10.1/lib"); println!("cargo:rustc-link-lib=dylib=cudart");

حيث /Developer/NVIDIA/CUDA-10.1/lib هو المسار إلى الملفات القابلة للتنفيذ CUDA ، في نظام يشبه يونيكس ، يمكن العثور على هذا المسار ، على سبيل المثال ، باستخدام الأمر:

 which nvcc

بالإضافة إلى ذلك ، في ملف build.rs ، تحتاج إلى تحديد المسار إلى ملفات dot.cpp و dot_gpu.cpp:

 .files(&["./src/dot.cpp", "./src/dot_gpu.cu"])

كل رمز build.rs

 extern crate cc; fn main() { cc::Build::new() .cuda(true) .cpp(true) .flag("-cudart=shared") .files(&["./src/dot.cpp", "./src/dot_gpu.cu"]) .compile("dot.a"); println!("cargo:rustc-link-search=native=/Developer/NVIDIA/CUDA-10.1/lib"); println!("cargo:rustc-link-search=/Developer/NVIDIA/CUDA-10.1/lib"); println!("cargo:rustc-env=LD_LIBRARY_PATH=/Developer/NVIDIA/CUDA-10.1/lib"); println!("cargo:rustc-link-lib=dylib=cudart"); }

الآن يمكنك البدء في كتابة رمز البرنامج الرئيسي. في الملف main.rs ، تحتاج إلى إنشاء واجهة لوظائف C / C ++ للاتصال مباشرة من كود Rust. يمكنك قراءة المزيد حول هذا الموضوع في الوثائق الرسمية في قسم FFI .

 extern "C" { //  C        fn dot(v1: *mut c_float, v2: *mut c_float, N: size_t) -> c_float; }

للاتصال به ، تحتاج إلى استخدام كتلة التعليمات البرمجية غير الآمنة ، حيث إن الوسيطات نقوم بتمرير مؤشر قابل للتغيير إلى نوع Vec:

 unsafe { gpu_res = dot(v1.as_mut_ptr(), v2.as_mut_ptr(), VEC_SIZE); }

رمز ملف main.rs الكامل

 extern crate libc; extern crate rand; use libc::{c_float, size_t}; use rand::Rng; const VEC_SIZE: usize = 10; const MAX: f32 = 10.; const MIN: f32 = 0.; extern "C" { fn dot(v1: *mut c_float, v2: *mut c_float, N: size_t) -> c_float; } fn cpu_dot(v1: Vec<f32>, v2: Vec<f32>) -> f32 { let mut res: f32 = 0.; for i in 0..v1.len() { res += v1[i] * v2[i]; } return res; } fn main() { let mut v1: Vec<f32> = Vec::new(); let mut v2: Vec<f32> = Vec::new(); let mut gpu_res: c_float; let mut cpu_res: f32 = 0.; let mut rng = rand::thread_rng(); for _ in 0..VEC_SIZE { v1.push(rng.gen_range(MIN, MAX)); v2.push(rng.gen_range(MIN, MAX)); } println!("{:?}", v1); println!("{:?}", v2); println!("GPU computing started"); unsafe { gpu_res = dot(v1.as_mut_ptr(), v2.as_mut_ptr(), VEC_SIZE); } println!("GPU computing finished"); println!("GPU dot product result: {}", gpu_res); cpu_res = cpu_dot(v1, v2); println!("CPU dot product result: {}", cpu_res); }

نبدأ الآن في كتابة الارتباطات في C ++ ، وكذلك رمز لحساب العدد القياسي للناقلات في CUDA C.

في ملف dot.cpp ، نكتب وظيفة الربط ، ندعو هذه الوظيفة بالفعل من كود Rust:

 extern "C" { float dot(float *v1, float *v2, size_t N) { float *gpu_res; float res = 0.0; gpu_res = gpu_dot(v1, v2, N); //   GPU for (int i = 0; i < blocksPerGrid; i++) { res += gpu_res[i]; } free(gpu_res); return res; } }

رمز ملف dot.cpp الكامل

 #include <iostream> #include "dot_gpu.h" using namespace std; void display_vector(float *v, size_t N) { cout << "["; for (size_t i = 0; i < N; i++) { cout << v[i]; if (i != N - 1) { cout << ", "; } } cout << "]" << endl; } extern "C" { float dot(float *v1, float *v2, size_t N) { cout << "Calling gpu dot product" << endl; cout << "Got two vectors from rust:" << endl; display_vector(v1, N); display_vector(v2, N); float *gpu_res; float res = 0.0; gpu_res = gpu_dot(v1, v2, N); for (int i = 0; i < blocksPerGrid; i++) { res += gpu_res[i]; } free(gpu_res); return res; } }

يوجد أدناه رمز من ملف dot_gpu.cu الذي يتم فيه إجراء الحساب الرئيسي ، ولن أشرح الرمز نفسه في هذا البرنامج التعليمي ، لأنه غير مخصص لبرمجة CUDA.

dot_gpu.cu

 #include "dot_gpu.h" __global__ void dot__(float *v1, float *v2, float *res, int N) { __shared__ float cache [threadsPerBlock]; int tid = threadIdx.x + blockIdx.x * blockDim.x; int cacheIndex = threadIdx.x; float temp = 0.0; while (tid < N) { temp += v1[tid] * v2[tid]; tid += blockDim.x * gridDim.x; } cache[cacheIndex] = temp; __syncthreads(); int i = blockDim.x / 2; while (i != 0) { if (cacheIndex < i) { cache[cacheIndex] += cache[cacheIndex + i]; } __syncthreads(); i /= 2; } if (cacheIndex == 0) { res[blockIdx.x] = cache[0]; } } float * gpu_dot (float *v1, float *v2, size_t N) { float *dev_v1, *dev_v2, *dev_res, *res; res = new float[blocksPerGrid]; cudaMalloc((void**)&dev_v1, N * sizeof(float)); cudaMalloc((void**)&dev_v2, N * sizeof(float)); cudaMalloc((void**)&dev_res, blocksPerGrid * sizeof(float)); cudaMemcpy(dev_v1, v1, N * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(dev_v2, v2, N * sizeof(float), cudaMemcpyHostToDevice); dot__<<<blocksPerGrid, threadsPerBlock>>>(dev_v1, dev_v2, dev_res, (int)N); cudaMemcpy(res, dev_res, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost); cudaFree(dev_v1); cudaFree(dev_v2); cudaFree(dev_res); return res; }

كل برنامجنا الصغير مكتوب وجاهز للبناء. لتجميعها في وحدة التحكم ، اتصل بالأمر:

 cargo build

لتشغيل:

 cargo run

بعد بناء البرنامج ، سيظهر المجلد الهدف في الدليل الرئيسي للمشروع. سيتم وضع الملف القابل للتنفيذ لبرنامجنا في المجلد: ./target/debug/

علاوة على ذلك ، إذا قمنا بتشغيل ملفنا القابل للتنفيذ ، فسوف نحصل على خطأ: لم يتم تحميل مكتبة dyld. بمعنى أنه لا يمكنه العثور على المسار إلى مكتبة cuda الحيوية. لحل هذه المشكلة ، يمكنك تسجيل متغير البيئة LD_LIBRARY_PATH = path_to_CUDA_lib_directory / أو تشغيل ارتباطات رمزية في مجلد سلسلة أدوات الصدأ لـ CUDA قبل بدء تشغيل الملف القابل للتنفيذ في وحدة التحكم:

 ln -s /Developer/NVIDIA/CUDA-10.1/lib/* /Users/Alexander/.rustup/toolchains/nightly-x86_64-apple-darwin/lib

حيث /Users/Alexander/.rustup/toolchains/nightly-x86_64-apple-darwin/lib هو طريقي إلى سلسلة أدوات الصدأ المثبتة ، فقد يختلف الأمر قليلاً عنك.

عند بدء تشغيل البرنامج من خلال عملية الشحن ، لم يحدث هذا الخطأ ، لأننا سجلنا متغير البيئة LD_LIBRARY_PATH في ملف build.rs.

في النهاية

لدينا القدرة على تشغيل كود CUDA C مباشرة من كود Rust. للتحقق من ذلك ، أنشأنا برنامجًا صغيرًا ، يعمل مع المتجهات ويقوم بإجراء جميع العمليات الحسابية على وحدة معالجة الرسومات. ويمكن أيضا أن ينظر إلى رمز كامل على جيثب .

الكتابة في Rust + CUDA C

كودا ج

RUST + CUDA C

في النهاية

More articles: