🛁 🙇🏼 👩🏿‍🏭 Rust + CUDA Cでの書き込み 👨🏼 👩🏾‍🏭 👩🏾‍🤝‍👩🏻

みなさんこんにちは！

このガイドでは、CUDA C / C ++とRustを友達にする方法を説明します。また、例として、ベクトルのスカラー積を計算するためにRustで小さなプログラムを作成します。スカラー積の計算は、CUDA Cを使用してGPUで実行されます。

猫を気にする人！

CUDA C

最初のステップは、CUDAコンパイラーnvccをインストールすることです。 CUDAとは何で、なぜそれを説明する必要があるのか、たとえばここで読むことができます。その助けを借りて、NVIDIAビデオカード（以下GPUと呼びます）で実行するコードを記述し、そのすべての能力を並列コンピューティングとグラフィックス処理に使用できるとしか言えません。繰り返しになりますが、このチュートリアルはCUDAでコードを記述する方法ではなく、Rustコードの利点を活用してGPUで並列コンピューティングを記述する方法に関するものです。

そのため、nvccとCUDA Toolkitをインストールします。この複雑さにより、詳細な指示は発生しません：オフサイトで。

Rust + CUDA C

前述のように、このチュートリアルでは、Rustで2つのベクトルのスカラー積を見つけるプログラムを作成します。計算プロセス自体はGPUで実行されます。

2つのベクトルのスカラー積。

2つのベクトルがあるとします。

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も錆コードのコンパイルを表していると思います。

まず、プロジェクトフォルダを作成します。プロジェクトフォルダーで、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

また、プロジェクトのルートフォルダーで、rustプログラムの構築とCUDA Cコードのコンパイルの手順を含むbuild.rsファイルを作成します。

ソースコードファイルを配置するプロジェクトルートにsrcフォルダーを追加します。 srcフォルダーに4つのファイルを作成します。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は、Unixライクシステム上の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ファイルでは、Rustコードから直接呼び出すためのC / C ++関数インターフェイスを作成する必要があります。これについての詳細は、 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); }

ここで、CUDA Cでベクトルのスカラー積を計算するためのコードと同様に、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 library not loaded。つまり、彼はcuda動的ライブラリへのパスを見つけることができません。この問題を解決するには、コンソールで実行可能ファイルを開始する前に、環境変数LD_LIBRARY_PATH = path_to_CUDA_lib_directory /を登録するか、CUDAのrust toolchainフォルダーでシンボリックリンクを実行します。

 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はインストール済みのrustツールチェーンへの私のパスであり、多少異なる場合があります。

環境変数LD_LIBRARY_PATHをbuild.rsファイルに登録したため、カーゴランを介してプログラムを起動したときに、このエラーは発生しませんでした。

最後に

RustコードからCUDA Cコードを直接実行することができます。これを確認するために、小さなプログラムを作成しました。ベクトルで動作し、GPUですべての計算を実行します。完全なコードはgithubでも表示できます。

Rust + CUDA Cでの書き込み

CUDA C

Rust + CUDA C

最後に

More articles: