はじめに
CUDA のコンパイラ nvcc を利用して動的リンクライブラリを作成します。
Windows でいう DLL ですが、今回は WSL で試しているので .so ファイルを作成します。
以下の環境で試しています。
- Windows 11
- WSL Ubuntu (20.04)
- GeForce RTX 2080
- nvcc 10.1
動的リンクライブラリの作成
サンプルコード
ライブラリのサンプルとして、add_vector.cu を作成します。
単純なベクトルの足し算を計算するためのコードです。
#ifndef ADD_VECTOR_H_
#define ADD_VECTOR_H_
void add_vector(float *dest, float *lhs, float *rhs, size_t n);
#endif // ADD_VECTOR_H_
#include "add_vector.h"
__global__ void add_vector_dev(float *dest, float *lhs, float *rhs, size_t n)
{
size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n)
{
dest[idx] = lhs[idx] + rhs[idx];
}
}
void add_vector(float *dest, float *lhs, float *rhs, size_t n)
{
float *dev_dest, *dev_lhs, *dev_rhs;
size_t bytes = n * sizeof(float);
cudaMalloc((void **)&dev_dest, bytes);
cudaMalloc((void **)&dev_lhs, bytes);
cudaMalloc((void **)&dev_rhs, bytes);
cudaMemcpy(dev_lhs, lhs, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(dev_rhs, rhs, bytes, cudaMemcpyHostToDevice);
dim3 block(32);
dim3 grid((n + block.x - 1) / block.x);
add_vector_dev<<<grid, block>>>(dev_dest, dev_lhs, dev_rhs, n);
cudaMemcpy(dest, dev_dest, bytes, cudaMemcpyDeviceToHost);
cudaFree(dev_dest);
cudaFree(dev_lhs);
cudaFree(dev_rhs);
}
動的リンクライブラリ生成
ターミナルから CUDA コンパイラを実行して、動的リンクライブラリを生成します。
$ nvcc --shared -Xcompiler -fPIC add_vector.cu -o libadd_vector.so--shared は nvcc に対するオプションです。
動的リンクライブラリを生成することを nvcc に指示します。
-Xcompiler は次のオプション（ -fPIC）を内部呼び出しするコンパイラ（gcc）に渡す効果があります。
PIC は Position Independent Code の略であり、相対アドレスで配置されたオブジェクトファイルを生成します。
gcc/g++ で .so ファイルを生成するときのおまじないです。
なお、-fPIC を与えないとリンクエラーが発生してしまいます。
$ nvcc --shared add_vector.cu -o libadd_vector.so
/usr/bin/ld: /tmp/tmpxft_000014fb_00000000-10_add_vector.o: relocation R_X86_64_PC32 against symbol `_Z14add_vector_devPfS_S_m' can not be used when making a shared object; recompile with -fPIC
/usr/bin/ld: final link failed: bad value
collect2: error: ld returned 1 exit statusコマンドが正常終了すると、libadd_vector.so ファイルが生成されます。
動的リンクライブラリの利用
サンプルコード
上記で作成した libadd_vector.so は、通常の .so ファイルと同様に扱うことが可能です。
以下、main 関数のサンプルです。
要素が３つのベクトルを GPU で足し算します。
#include <cstdio>
#include <cstdlib>
#include <ctime>
#include "add_vector.h"
void initialize_list(float *list, size_t n) {
for (size_t i = 0; i < n; ++i) {
list[i] = (float)rand() / RAND_MAX;
}
}
int main() {
float *dest, *lhs, *rhs;
size_t n = 3;
size_t bytes_dest = n * sizeof(float);
size_t bytes_lhs = n * sizeof(float);
size_t bytes_rhs = n * sizeof(float);
dest = (float *)malloc(bytes_dest);
lhs = (float *)malloc(bytes_lhs);
rhs = (float *)malloc(bytes_rhs);
srand((unsigned int)time(NULL));
initialize_list(lhs, n);
initialize_list(rhs, n);
add_vector(dest, lhs, rhs, n);
for (size_t i = 0; i < n; ++i) {
printf("%.6f + %.6f = %.6f\n", lhs[i], rhs[i], dest[i]);
}
free(dest);
free(lhs);
free(rhs);
}コンパイル／リンク
ターミナルからコンパイルとリンクを行います。
CUDA のコード部分は既に .so に梱包済みのため nvcc を利用する必要はありません。
gcc や g++ でコンパイル可能です。
$ g++ main.cc -L. -ladd_vector -o main-ladd_vector オプションによって、libadd_vector.so が探索されリンクされます。
実行
LD_LIBRARY_PATH 環境変数をごまかす必要がありますが、これで main の実行ができます。
$ LD_LIBRARY_PATH=. ./main
0.416673 + 0.586987 = 1.003660
0.212794 + 0.141135 = 0.353929
0.531908 + 0.280919 = 0.812828./main だけだと、ライブラリ探索に失敗してエラーになります。
$ ./main
./main: error while loading shared libraries: libadd_vector.so: cannot open shared object file: No such file or directory