RustでOpenCL(ベクトル加算編)

はじめに

概要

• RustでOpenCLを使ってGPU上で10個の要素のベクトルの加算を行いました。
• ベクトル加算以外の他の処理を行う場合でも、殆どの部分を使いまわし出来ると思います。
• APIの詳細はこちらをご覧ください。

ブログラム

Githubに公開しています。

環境

• OS: Ubuntu 20.04
• CPU: Intel Core i9 9900K
• GPU: NVIDIA RTX 2070 Super

実装

Cargo.tomlファイルの編集

Cargo.tomlファイルにopencl3 = "0.6"を追加します。

Cargo.toml

[dependencies]
opencl3 = "0.6"

ベクトル加算のカーネル

下記がGPUで実行されるベクトル加算のカーネルコードです。カーネルコードはOpenCL C形式ですのでホスト側をCやC++、C#で実装する場合と変わりありません。

vadd.cl

__kernel void vadd(__global int* a, __global int* b, __global int* c) {
  int i = get_global_id(0);

  c[i] = a[i] + b[i];
}

ホスト側

続いて、ホスト側のコードです。プログラム全文はGithubに公開しています。

プラットフォームの取得

// get platforms
let platforms = opencl3::platform::get_platforms().expect(
  "Failed to get platforms."
);
let platform = platforms[PLATFORM_ID];

プラットフォーム（実行環境）を取得します。実行環境は各ベンダ毎に提供されています（Intel、AMD、NVIDIA、Xilinx等）。
上記では、opencl3::platform::get_platforms()でプラットフォームの一覧を配列として受け取り、platforms[PLATFORM_ID]でプラットフォームを一つ選択しています（今回の例ではNVIDIA）。

デバイスの取得

let devices = platform.get_devices(opencl3::device::CL_DEVICE_TYPE_ALL).expect(
  "Failed to get devices."
);
let device = opencl3::device::Device::new(devices[DEVICE_ID]);

プラットフォームを一つ選択したので、次にデバイスを選択します。
今回はプラットフォームとしてNVIDIAを選んだのでこのプラットフォームで実行できるデバイスの一覧がplatform.get_devices(opencl3::device::CL_DEVICE_TYPE_ALL)で取得できます（NVIDIA GeForce RTX 2070 SUPER、NVIDIA GeForce RTX 3080等）。
opencl3::device::Device::new(devices[DEVICE_ID])でデバイスを一つ選択しています（今回の例ではNVIDIA GeForce RTX 2070 SUPER）。

コンテキストの作成

// create context
let context = opencl3::context::Context::from_device(&device).expect(
  "Failed to create context."
);

デバイスを一つ選択したので、次にコンテキストを作成します。コンテキストはホストコード上でのデバイスのような感じで、ホストコードからデバイスを操作するにはコンテキストを通じて行います。このコンテキストの作成を経てデバイスのセットアップが完了します。
opencl3::context::Context::from_device(&device)でコンテキストを作成します。

カーネルプログラムのビルド

let source = std::fs::read_to_string("./src/vadd.cl").unwrap();

// create program
let mut program = opencl3::program::Program::create_from_source(&context, &source).expect(
  "Failed to create program."
);
  
// build program
program.build(&devices, "").expect(
  "Failed to build program."
);

次に、下記の手順でカーネルプログラムのビルドを行います。

• std::fs::read_to_string("./src/vadd.cl")でカーネルプログラムを読み込みます。
• 読み込んだソースコードからopencl3::program::Program::create_from_source(&context, &source)でプログラムを作成します。
• program.build(&devices, "")でビルドを行います。

program.build(&devices, "")の第1引数ではビルド対象のデバイスを配列で渡すことが出来ます。また、第2引数でビルドオプションを指定できます。今回は指定しないので空文字にしています。

カーネルの作成

// create kernel
let kernel = opencl3::kernel::Kernel::create(&program, "vadd").expect(
  "Failed to create kernel."
);

ビルドしたプログラムを使ってカーネルを作成します。カーネルはデバイスで実行する関数のホストコード上の表現のようなものです。opencl3::kernel::Kernel::create(&program, "vadd")で作成でき、第2引数で使用するカーネル関数を選択できます（今回はvadd関数）。

コマンドキューの作成

// create command queue
let command_queue = opencl3::command_queue::CommandQueue::create(&context, device.id(), opencl3::command_queue::CL_QUEUE_PROFILING_ENABLE).expect(
  "Failed to create command queue."
);

続いて、コマンドキューを作成します。コマンドキューはホスト側とデバイス側をつなぐパイプのようなものです。使用するデバイス（コンテキスト）を指定してキューを作成します。第1引数でコンテキストを指定し、第2引数でデバイスIDを指定します。デバイスIDは第1引数のコンテキストを作成したときのデバイスIDである必要があります。

バッファの作成

GPUで実行するためにGPU上でメモリの確保を行います。

// create buffer
let mut a_buffer = opencl3::memory::Buffer::<opencl3::types::cl_int>::create(&context, opencl3::memory::CL_MEM_READ_ONLY, 10, std::ptr::null_mut()).expect(
  "Failed to create buffer a."
);
let mut b_buffer = opencl3::memory::Buffer::<opencl3::types::cl_int>::create(&context, opencl3::memory::CL_MEM_READ_ONLY, 10, std::ptr::null_mut()).expect(
  "Failed to create buffer b."
);
let mut c_buffer = opencl3::memory::Buffer::<opencl3::types::cl_int>::create(&context, opencl3::memory::CL_MEM_WRITE_ONLY, 10, std::ptr::null_mut()).expect(
  "Failed to create buffer c."
);

opencl3::memory::Buffer::createでメモリの確保を行います。今回は加算する10要素のベクトル2つと出力用の10要素のベクトル1つを使用しますのでそれぞれのメモリを確保します。

• 第1引数でコンテキストを指定します。
• 第2引数でフラグを指定します。c[i] = a[i] + b[i]の加算を行うので、a, bにはopencl3::memory::CL_MEM_READ_ONLY、cにはopencl3::memory::CL_MEM_WRITE_ONLYを指定してます。
• 第3引数で要素数を指定します。
• 第4引数は確保済のバッファを指定します。今回はここで始めて作成しているのでstd::ptr::null_mut()を与えています。

バッファへの書き込み

GPU上のメモリの確保が終わったので次に書き込みを行います。

let a: [i32; 10] = [1, 2, 3, 4, 5, 6, 7, 8, 9, 10];
let b: [i32; 10] = [1, 2, 3, 4, 5, 6, 7, 8, 9, 10];
let mut c: [i32; 10] = [0; 10];

// write buffer
command_queue.enqueue_write_buffer(&mut a_buffer, opencl3::types::CL_TRUE, 0, &a[..], &[]).expect(
  "Failed to write buffer a."
);
command_queue.enqueue_write_buffer(&mut b_buffer, opencl3::types::CL_TRUE, 0, &b[..], &[]).expect(
  "Failed to write buffer b."
);

• 第1引数でバッファを指定し、第4引数で書き込むデータを指定します。
• 第2引数は同期的に書き込むか非同期的に書き込むかを選択できます。ここでは同期的に行います。
• 第3引数はバッファに書き込むオフセットを指定できます。
• 第5引数では書き込みを実行する前に終了しておくべきイベントをリストとして与えることが出来ます。今回は無いので空のリストを渡しています。

カーネル実行

// execute kernel
let kernel_event = opencl3::kernel::ExecuteKernel::new(&kernel)
  .set_arg(&a_buffer)
  .set_arg(&b_buffer)
  .set_arg(&c_buffer)
  .set_global_work_size(10)
  .set_local_work_size(1)
  .set_global_work_offset(0)
  .enqueue_nd_range(&command_queue)
  .expect("Failed to enqueue kernel execution.");

opencl3::kernel::ExecuteKernelのenqueue_nd_rangeでカーネルを実行します。
set_argでカーネルの引数をそれぞれバッファと結びつけ、set_global_work_size(10)でワークサイズを指定します。これによって、カーネルコード内のget_global_id(0)が0から9まで回ります。
enqueue_nd_rangeでカーネルを実行します。この式は非同期なのでここを抜けてもカーネルの実行が完了している保証はありません。

バッファからの読み込み

// read buffer
command_queue.enqueue_read_buffer(&mut c_buffer, opencl3::types::CL_TRUE, 0, &mut c[..], &[kernel_event.get()]).expect(
  "Failed to read buffer c."
);

ベクトル加算の結果をバッファから読み込みます。引数は書き込みの場合と同様です。
第5引数については、バッファの読み込みの前にカーネルの実行が完了している必要があるのでカーネルイベントを渡しています。

実行

$ cargo run
Using platform: NVIDIA CUDA
Using device: NVIDIA GeForce RTX 2070 SUPER
c[0] = 2
c[1] = 4
c[2] = 6
c[3] = 8
c[4] = 10
c[5] = 12
c[6] = 14
c[7] = 16
c[8] = 18
c[9] = 20