👩🏽‍🤝‍👨🏿 👱 🛍️ Schreiben in Rust + CUDA C. 👋🏼 🦐 👨‍👩‍👧

Hallo allerseits!

In diesem Handbuch möchte ich erklären, wie man Freunde CUDA C / C ++ und Rust findet. Als Beispiel schreiben wir ein kleines Programm in Rust, um das Skalarprodukt von Vektoren zu berechnen. Die Berechnung des Skalarprodukts erfolgt auf der GPU mit CUDA C.

Wen interessiert die Katze?

CUDA C.

Der erste Schritt ist die Installation des CUDA-Compilers - nvcc. Was ist CUDA und warum muss ich es beschreiben? Ich kann es hier zum Beispiel lesen. Ich kann nur sagen, dass Sie mit seiner Hilfe Code schreiben können, der auf NVIDIA-Grafikkarten (im Folgenden als GPU bezeichnet) ausgeführt wird, und deren gesamte Leistung für paralleles Rechnen und Grafikverarbeitung nutzen können. In diesem Tutorial geht es nicht darum, wie man Code in CUDA schreibt, sondern wie man die Vorteile von Rust-Code nutzt und paralleles Computing auf die GPU schreibt.

Installieren Sie also nvcc und das CUDA Toolkit. Bei dieser Komplexität sollten keine detaillierten Anweisungen entstehen: vor Ort .

ROST + CUDA C.

Während dieses Tutorials werden wir, wie bereits erwähnt, ein Programm in Rust schreiben, um das Skalarprodukt zweier Vektoren zu finden. Der Berechnungsprozess selbst wird auf der GPU ausgeführt.

Das Skalarprodukt zweier Vektoren.

Angenommen, wir haben zwei Vektoren:

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

$a = [a_1, a_2, ... a_n]$ und

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

$b = [b_1, b_2, ..., b_n]$ , das Skalarprodukt dieser Vektoren:

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}$

Beginnen wir mit der Erstellung unseres Programms. Außerdem gehe ich davon aus, dass nvcc erfolgreich installiert wurde. Rustc und Fracht stehen auch für das Kompilieren von Rostcode.

Erstellen Sie zunächst den Projektordner. Erstellen Sie im Projektordner die Datei Cargo.toml, die Anweisungen für den Ladungssammler enthält. Die Datei sieht folgendermaßen aus:

[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

Erstellen Sie auch im Projektstammordner die Datei build.rs, die Anweisungen zum Erstellen des Rostprogramms und zum Kompilieren des CUDA C-Codes enthält.

Fügen Sie den Ordner src zum Projektstamm hinzu, in dem wir die Quellcodedateien ablegen. Erstellen Sie im Ordner src vier Dateien: main.rs - den Code des Hauptprogramms, dot.cpp - C ++ - Bindung (Wrapper für CUDA C), dot_gpu.h, dot_gpu.cu - die Datei, die den auf der GPU ausgeführten Code enthält.

Insgesamt haben wir eine solche Projektstruktur:

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

In der Datei build.rs ist es am wichtigsten, Folgendes zu schreiben:

 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");

Dabei ist /Developer/NVIDIA/CUDA-10.1/lib der Pfad zu den ausführbaren CUDA-Dateien. In einem Unix-ähnlichen System kann dieser Pfad beispielsweise mit dem folgenden Befehl gefunden werden:

 which nvcc

Darüber hinaus müssen Sie in der Datei build.rs den Pfad zu den Dateien dot.cpp und dot_gpu.cpp angeben:

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

Alle build.rs Code

 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"); }

Jetzt können Sie mit dem Schreiben des Hauptprogrammcodes beginnen. In der Datei main.rs müssen Sie eine Schnittstelle für C / C ++ - Funktionen erstellen, um direkt aus Rust-Code aufzurufen. Weitere Informationen hierzu finden Sie in der offiziellen Dokumentation im Abschnitt FFI .

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

Um es aufzurufen, müssen Sie den unsicheren Codeblock verwenden, da wir als Argumente einen veränderlichen Zeiger auf den Vec-Typ übergeben:

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

Vollständiger main.rs-Dateicode

 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); }

Jetzt schreiben wir die Bindungen in C ++ sowie den Code zur Berechnung des Skalarprodukts von Vektoren in CUDA C.

In die Datei dot.cpp schreiben wir die Bindungsfunktion, wir rufen diese Funktion tatsächlich aus dem Rust-Code auf:

 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; } }

Vollständiger dot.cpp-Dateicode

 #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; } }

Das Folgende ist der Code aus der Datei dot_gpu.cu, in der die Hauptberechnung durchgeführt wird. Ich werde den Code selbst in diesem Tutorial nicht erklären, da er nicht für die CUDA-Programmierung vorgesehen ist.

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; }

Alle unsere kleinen Programme sind geschrieben und bereit zum Erstellen. Rufen Sie den folgenden Befehl auf, um es in der Konsole zusammenzustellen:

 cargo build

So führen Sie aus:

 cargo run

Nach dem Erstellen des Programms wird der Zielordner im Hauptverzeichnis des Projekts angezeigt. Die ausführbare Datei unseres Programms befindet sich im Ordner: ./target/debug/

Wenn wir nur unsere ausführbare Datei ausführen, wird außerdem die folgende Fehlermeldung angezeigt: Die dyld-Bibliothek wurde nicht geladen. Das heißt, er kann den Pfad zur dynamischen cuda-Bibliothek nicht finden. Um dieses Problem zu lösen, können Sie die Umgebungsvariable LD_LIBRARY_PATH = path_to_CUDA_lib_directory / registrieren oder symbolische Links im Ordner rust rustchainchain für CUDA ausführen, bevor Sie die ausführbare Datei in der Konsole starten:

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

Wenn /Users/Alexander/.rustup/toolchains/nightly-x86_64-apple-darwin/lib mein Pfad zur installierten Rost-Toolchain ist, kann dies für Sie geringfügig abweichen.

Beim Starten des Programms über den Frachtlauf ist ein solcher Fehler nicht aufgetreten, da wir die Umgebungsvariable LD_LIBRARY_PATH in der Datei build.rs registriert haben.

Zusammenfassend

Wir haben die Möglichkeit, CUDA C-Code direkt aus Rust-Code auszuführen. Um dies zu überprüfen, haben wir ein kleines Programm erstellt, das mit Vektoren arbeitet und alle Berechnungen auf der GPU durchführt. Der vollständige Code kann auch auf github angezeigt werden .

Schreiben in Rust + CUDA C.

CUDA C.

ROST + CUDA C.

Zusammenfassend

More articles: