🎗️ 🎛️ 😳 Écriture en rouille + CUDA C 🎮 🚍 👶

Bonjour à tous!

Dans ce guide, je veux dire comment se faire des amis CUDA C / C ++ et Rust. Et à titre d'exemple, nous écrirons un petit programme en Rust pour calculer le produit scalaire des vecteurs, le calcul du produit scalaire sera effectué sur le GPU en utilisant CUDA C.

Qui se soucie du chat!

CUDA C

La première étape consiste à installer le compilateur CUDA - nvcc. Qu'est-ce que CUDA et pourquoi je dois le décrire, je peux le lire ici, par exemple. Je peux seulement dire qu'avec son aide, vous pouvez écrire du code qui fonctionnera sur les cartes vidéo NVIDIA (ci-après dénommé le GPU) et utiliser toute leur puissance pour le calcul parallèle et le traitement graphique. Encore une fois, ce didacticiel n'est pas sur la façon d'écrire du code dans CUDA, mais sur la façon d'utiliser ses avantages à partir du code Rust et d'écrire le calcul parallèle sur le GPU.

Installez donc nvcc et la boîte à outils CUDA. Avec cette complexité, des instructions détaillées ne devraient pas apparaître: sur le site .

ROUILLE + CUDA C

Au cours de ce tutoriel, comme mentionné précédemment, nous écrirons un programme dans Rust pour trouver le produit scalaire de deux vecteurs, le processus de calcul lui-même se produira sur le GPU.

Le produit scalaire de deux vecteurs.

Supposons que nous ayons deux vecteurs:

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

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

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

$b = [b_1, b_2, ..., b_n]$ , le produit scalaire de ces vecteurs:

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

Commençons à créer notre programme. De plus, je suppose que nvcc est installé avec succès, rustc et cargo représentent également la compilation du code de rouille.

Créez d'abord le dossier du projet. Dans le dossier du projet, créez le fichier Cargo.toml, qui contient des instructions pour le collecteur de marchandises. Le fichier ressemble à ceci:

[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

Toujours dans le dossier racine du projet, créez le fichier build.rs qui contiendra des instructions pour construire le programme rust et compiler le code CUDA C.

Ajoutez le dossier src à la racine du projet dans lequel nous placerons les fichiers de code source. Dans le dossier src, créez quatre fichiers: main.rs - le code du programme principal, dot.cpp - liaison C ++ (wrapper pour CUDA C), dot_gpu.h, dot_gpu.cu - le fichier qui contient le code exécuté sur le GPU.

Au total, nous avons une telle structure de projet:

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

Dans le fichier build.rs, la chose la plus importante est d'écrire ceci:

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

où /Developer/NVIDIA/CUDA-10.1/lib est le chemin d'accès aux fichiers exécutables CUDA, dans un système de type Unix, ce chemin peut être trouvé, par exemple, avec la commande:

 which nvcc

De plus, dans le fichier build.rs, vous devez spécifier le chemin d'accès aux fichiers dot.cpp et dot_gpu.cpp:

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

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

Vous pouvez maintenant commencer à écrire le code du programme principal. Dans le fichier main.rs, vous devez créer une interface pour les fonctions C / C ++ pour appeler directement à partir du code Rust. Vous pouvez en savoir plus à ce sujet dans la documentation officielle de la section FFI .

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

Pour l'appeler, vous devez utiliser le bloc de code dangereux, comme arguments, nous passons un pointeur mutable au type Vec:

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

Code de fichier main.rs complet

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

Nous commençons maintenant à écrire les liaisons en C ++, ainsi que le code pour calculer le produit scalaire des vecteurs dans CUDA C.

Dans le fichier dot.cpp, nous écrivons la fonction de liaison, nous appelons en fait cette fonction à partir du code 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; } }

Code de fichier dot.cpp complet

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

Ce qui suit est le code du fichier dot_gpu.cu dans lequel le calcul principal est effectué, je n'expliquerai pas le code lui-même dans ce tutoriel, car il n'est pas dédié à la programmation 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; }

Tout notre petit programme est écrit et prêt à être construit. Pour l'assembler dans la console, appelez la commande:

 cargo build

Pour exécuter:

 cargo run

Après avoir construit le programme, le dossier cible apparaîtra dans le répertoire principal du projet. Le fichier exécutable de notre programme sera situé dans le dossier: ./target/debug/

De plus, si nous exécutons simplement notre fichier exécutable, nous obtenons une erreur: la bibliothèque dyld n'est pas chargée. Autrement dit, il ne peut pas trouver le chemin vers la bibliothèque dynamique cuda. Pour résoudre ce problème, vous pouvez enregistrer la variable d'environnement LD_LIBRARY_PATH = path_to_CUDA_lib_directory / ou exécuter des liens symboliques dans le dossier rust toolchain pour CUDA avant de démarrer le fichier exécutable dans la console:

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

où /Users/Alexander/.rustup/toolchains/nightly-x86_64-apple-darwin/lib est mon chemin vers la chaîne d'outils rouille installée, cela peut différer légèrement pour vous.

Lors du démarrage du programme via l'exécution du chargement, une telle erreur ne s'est pas produite, car nous avons enregistré la variable d'environnement LD_LIBRARY_PATH dans le fichier build.rs.

En fin de compte

Nous avons la possibilité d'exécuter du code CUDA C directement à partir du code Rust. Afin de vérifier cela, nous avons créé un petit programme, il fonctionne avec des vecteurs et effectue tous les calculs sur le GPU. Le code complet peut également être consulté sur github .

Écriture en rouille + CUDA C

CUDA C

ROUILLE + CUDA C

En fin de compte

More articles: