消えたCUDA関連の旧ブログ記事を復元するひとり Advent Calendar 2024の記事です。

何の話か

CUDA用の静的ライブラリを作る場合のコンパイルオプションについてです．
CUDAではコンパイルを行う時にホスト側のコードとデバイス側のコードを分離し，デバイス側のコードはgencodeで指定したアーキテクチャの種類だけカーネルイメージを作成します．
この際後方互換性のためにPTXをイメージに埋め込みます．
nvccでCUDAの静的ライブラリを作る時は，これらイメージを再配置可能にする必要があり，そのためのオプションが用意されています．
この記事はこれらのオプションを使って静的ライブラリを作る話です．

作り方

ライブラリの定義コード（test.cu）から静的ライブラリ（libtest.a）を作ることが目標です．
今回サンプルとして以下のコードを用います．

test.cu

#include <stdio.h>
#include "test.hpp"

namespace {
__global__ void hello_kernel() {
	printf("[GPU] Hello\n");
}
}

void print_hello() {
	printf("[CPU] Hello\n");
	hello_kernel<<<1, 1>>>();
	cudaDeviceSynchronize();
}

test.hpp

#ifndef __TEST_HPP__
#define __TEST_HPP__

void print_hello();

#endif

はじめに定義コード（test.cu）から普通のオブジェクトファイル（test.o）を作ります．
この際，埋め込みたいカーネルイメージの対象だけ-gencode arch=compute_XX,code=sm_XXを列挙します． 

nvcc test.cu -c -o test.o -dc -gencode arch=compute_86,code=sm_86

次に，今作ったオブジェクトファイル内のイメージを再配置可能にします．
これにはnvccの-dlinkオプションを用います．
これにより先程作ったtest.oからtest.dlink.oを得ます． 

nvcc -dlink test.o -o test.dlink.o -gencode arch=compute_86,code=sm_86

最後に，作ったtest.oとtest.dlink.oから目的のlibtest.aを作成します． 

nvcc -o libtest.a -lib test.dlink.o test.o

これだけです．
定義コードがたくさんある場合は，それぞれの.oと.*dlink.oを作って最後の手順に渡すことで一つのライブラリを作れます．

使い方

簡単に作ったライブラリを他のプログラムにリンクさせながらビルドするにはnvccを使うのがおすすめです．
例えばこんなコードをコンパイルしてみます．

main.cpp

#include "test.hpp"

int main() {
	print_hello();
}

コンパイルプションはこんな感じです． 

nvcc main.cpp -L[/path/to/libtest.a (w/o libtest.a)] -ltest

静的ライブラリは作りましたが，CUDAの公式のライブラリたちは動的ライブラリなため，このライブラリとは別途リンクさせる必要があります．
nvccであればこれを自動で行ってくれますが，gccなどの他のコンパイラを用いる場合は自分でオプションを付ける必要があります．

おわりに

対象のアーキを列挙すればするほどライブラリのファイルサイズが増加し，かつビルド時間も増加します．
不必要なアーキは列挙しないのがおすすめです．

参考