🈹 🆗 🖌️ Escrita em Rust + CUDA C 🍷 🧑🏽 👨‍🚒

Olá pessoal!

Neste guia, quero dizer como fazer amigos CUDA C / C ++ e Rust. E como exemplo, escreveremos um pequeno programa no Rust para calcular o produto escalar de vetores, o cálculo do produto escalar será realizado na GPU usando CUDA C.

Quem se importa com o gato!

CUDA C

O primeiro passo é instalar o compilador CUDA - nvcc. O que é CUDA e por que preciso descrevê-lo, posso ler sobre isso aqui, por exemplo. Só posso dizer que, com sua ajuda, você pode escrever o código que será executado nas placas de vídeo da NVIDIA (doravante denominada GPU) e usar todo o seu poder para computação paralela e processamento gráfico. Mais uma vez, este tutorial não é sobre como escrever código no CUDA, mas sobre como usar suas vantagens do código Rust e escrever computação paralela na GPU.

Portanto, instale o nvcc e o CUDA Toolkit. Com essa complexidade, instruções detalhadas não devem surgir: fora do local .

RUST + CUDA C

Durante este tutorial, como mencionado anteriormente, escreveremos um programa no Rust para encontrar o produto escalar de dois vetores; o próprio processo de cálculo ocorrerá na GPU.

O produto escalar de dois vetores.

Suponha que temos dois vetores:

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

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

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

$b = [b_1, b_2, ..., b_n]$ , o produto escalar desses vetores:

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

Vamos começar a criar nosso programa. Além disso, suponho que o nvcc seja instalado com sucesso, rustc e cargo também representem a compilação do código de ferrugem.

Primeiro, crie a pasta do projeto. Na pasta do projeto, crie o arquivo Cargo.toml, que contém instruções para o coletor de carga. O arquivo fica assim:

[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

Também na pasta raiz do projeto, crie o arquivo build.rs que conterá instruções para criar o programa rust e compilar o código CUDA C.

Adicione a pasta src à raiz do projeto na qual colocaremos os arquivos de código-fonte. Na pasta src, crie quatro arquivos: main.rs - o código do programa principal, dot.cpp - ligação C ++ (wrapper para CUDA C), dot_gpu.h, dot_gpu.cu - o arquivo que contém o código executado na GPU.

Total, temos uma estrutura de projeto:

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

No arquivo build.rs, o mais importante é escrever o seguinte:

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

onde /Developer/NVIDIA/CUDA-10.1/lib é o caminho para os arquivos executáveis da CUDA, em um sistema unix, esse caminho pode ser encontrado, por exemplo, com o comando:

 which nvcc

Além disso, no arquivo build.rs, você precisa especificar o caminho para os arquivos dot.cpp e dot_gpu.cpp:

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

Todo o código 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"); }

Agora você pode começar a escrever o código principal do programa. No arquivo main.rs, você precisa criar uma interface para funções C / C ++ para chamar diretamente do código Rust. Você pode ler mais sobre isso na documentação oficial na seção FFI .

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

Para chamá-lo, você precisa usar o bloco de código não seguro, pois argumentos passamos um ponteiro mutável para o tipo Vec:

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

Código completo do arquivo 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); }

Agora começamos a escrever as ligações em C ++, bem como o código para calcular o produto escalar de vetores em CUDA C.

No arquivo dot.cpp, escrevemos a função de ligação, na verdade chamamos essa função do código 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; } }

Código completo do arquivo 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; } }

A seguir está o código do arquivo dot_gpu.cu no qual o cálculo principal é realizado, não explicarei o código neste tutorial, pois ele não é dedicado à programação 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; }

Todo o nosso pequeno programa está escrito e pronto para ser construído. Para montá-lo no console, chame o comando

 cargo build

Para executar:

 cargo run

Após criar o programa, a pasta de destino aparecerá no diretório principal do projeto. O arquivo executável do nosso programa estará localizado na pasta: ./target/debug/

Além disso, se apenas rodarmos nosso arquivo executável, obteremos um erro: biblioteca dyld não carregada. Ou seja, ele não consegue encontrar o caminho para a biblioteca dinâmica cuda. Para resolver esse problema, você pode registrar a variável de ambiente LD_LIBRARY_PATH = path_to_CUDA_lib_directory / ou executar links simbólicos na pasta rust toolchain para CUDA antes de iniciar o arquivo executável no console:

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

onde /Users/Alexander/.rustup/toolchains/nightly-x86_64-apple-darwin/lib é o meu caminho para a cadeia de ferramentas ferrugem instalada, ela pode ser um pouco diferente para você.

Ao iniciar o programa através da execução de carga, esse erro não ocorreu, porque registramos a variável de ambiente LD_LIBRARY_PATH no arquivo build.rs.

No final

Temos a capacidade de executar o código CUDA C diretamente do código Rust. Para verificar isso, criamos um pequeno programa, que trabalha com vetores e realiza todos os cálculos na GPU. O código completo também pode ser visualizado no github .

Escrita em Rust + CUDA C

CUDA C

RUST + CUDA C

No final

More articles: