[CUDA] 動的リンクライブラリの作り方

  • 2022年8月9日
  • 2022年9月17日
  • CUDA
  • 29View
  • 0件

はじめに

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 関数のサンプルです。
要素が3つのベクトルを 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