夏の怪談：SYCLのプログラムを実行したら、コンテナの`/`が消えた話

あらすじ

それは、6月の蒸し暑い日の夕暮れのこと。
僕は、Podman コンテナで SYCL プログラムを書き、コンパイルしていた。
ターゲットデバイスは NVIDIA の GPU。コンパイラは intel/llvm clang。
順当に構築した環境で、ソースファイルをコンパイルし、実行した。
...
その直後、コンテナの / 配下がすべて見えなくなった。

当時のコマンド

$ /path/to/llvm/bin/clang++ -fsycl \
  -fsycl-targets=nvptx64-nvidia-cuda \
  --cuda-path=/path/to/cuda/toolkit \
  xxx.cpp -o a.out
$ LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/path/to/llvm/build/lib ./a.out

コマンドを見る限り、特別なことはしていない。
カーネルを実行しただけである。

何が起きたのだろうか...

思い当たる節が、全くなかったわけではない。
デバイスから書き込めない可能性のあるホスト側ポインタ、malloc / free、pipeやsignalなどを呼ぶいくつかの関数がデバイスコードに含まれていた可能性がある。
これらについて、一つずつ見ていく。

容疑者A：ホストポインタ

ホスト側とデバイス側のポインタを混在させちゃったかも。

SYCL 2020 仕様によると、USM でないホスト側ポインタをデバイス側で参照した場合は動作未定義となる。

Pointers and objects containing pointers may be shared. However, when a pointer is passed between SYCL devices or between the host and a SYCL device, dereferencing that pointer on the device produces undefined behavior unless the device supports USM and the pointer is an address within a USM memory region.
ポインタとポインタを含むオブジェクトは共有できます。ただし、ポインタが SYCL デバイス間またはホストと SYCL デバイス間で渡される場合、デバイスが USM をサポートし、ポインタが USM メモリ領域のアドレスでない限り、デバイス上のポインタを逆参照すると動作は未定義となります
(SYCL™ 2020 Specification, revision 11, Sec.5.4)

以下はその例である。

危険なコード：デバイスコードから通常のホストポインタへ書き込もうとしている

int *host_ptr =
    static_cast<int *>(std::malloc(size * sizeof(int)));

queue.parallel_for(
    sycl::range<1>{size},[=](sycl::id<1> index) {
        // 危険なコード：デバイスコードから通常のホストポインタへ書き込もうとしている
        host_ptr[index] = static_cast<int>(index[0]);
    }
);
:
std::free(host_ptr);

ただし、これらの処理は通常、コンパイルまたはリンクの段階で拒否されるか、カーネル停止や CUDA context のエラーで落ちるようである。
そのため、これらの問題によって、GPUカーネルがコンテナ内から見える/を直接削除したり、rootfsをアンマウントしたりする可能性は低い。
※ NVPTXカーネルはLinuxプロセスではなく、コンテナのmount namespaceを直接操作する仕組みを持たないためである。

容疑者B：デバイス側でのmalloc/free

デバイス側でホストの malloc/free を呼んでしまったかも。

CUDA には専用の cudaMalloc/cudaFree が存在する。
具体的には、デバイスメモリを確保するcudaMalloc/cudaFreeに加えて、CUDAカーネル内からデバイスヒープを操作するデバイス側malloc/freeも存在する。これらは別のメモリ管理系統であり、確保元と解放側を混同してはならない。例えば、cudaMalloc()で確保したポインタをデバイス側free()で解放することはできない。

一方、標準SYCLでは、カーネル内での動的なストレージ確保は許可されていない。したがって、SYCLカーネル内のstd::malloc()が、CUDAバックエンドによってCUDAのデバイス側malloc()へ変換されることは、標準SYCLの動作として規定されていない。

SYCL 2020 仕様では、「カーネル内でメモリ領域を確保することは許可されない」と規定されている。

Memory storage allocation is not allowed in kernels. All memory allocation for the device is done on the host using accessor classes or using USM as explained in Section 4.8. Consequently, the default allocation operator new overloads that allocate storage are disallowed in a SYCL kernel. The placement new operator and any user-defined overloads that do not allocate storage are permitted.
カーネルでのメモリーストレージ割り当ては許可されていません。デバイスのすべてのメモリ割り当ては、4.8 節で説明されるように、アクセサークラスや USM によりホスト上で行われます。そのため、ストレージを割り当てるデフォルトの割り当てオペレータ new のオーバーロードは、SYCL カーネルでは許可されません。プレースメント new オペレーターとストレージを割り当てないユーザー定義のオーバーロードが許可されます。
(SYCL™ 2020 Specification, revision 11, Sec.5.4)

危険なコード：デバイスカーネル内で動的メモリを扱う

queue.single_task([=]() {
    // デバイスカーネル内で動的メモリを確保しようとしている。
    int *ptr = static_cast<int *>(
        std::malloc(16 * sizeof(int)));
            :
    // 同じくデバイスカーネル内で解放しようとしている。
    std::free(ptr);
});

freeは、何でも自由にしていいという意味ではない...w
ただ、これもメモリ破壊の可能性はあるが、コンテナ rootfs 消失の直接原因とはならないようだ。

容疑者C：pipeやsignalなどのホストOS依存API

デバイスカーネル内で、Linux システムコールを呼んでたかも。

デバイス上では通常のLinuxシステムコールとして実行することはできない。

• 今回の NVPTX カーネルは Linux プロセスではない
• GPU 上に通常のファイルディスクリプタや signal 処理環境はない

危険なコード：ホスト向け関数をデバイスカーネルから呼び出そうとしている

void initialize_runtime()
{
    int file_descriptors[2];
    // Linux ファイルディスクリプタ作成
    pipe(file_descriptors);
    // プロセスのシグナルハンドラ設定
    signal(SIGTERM, SIG_IGN);
}

:
queue.single_task([=]() {
    // 危険なコード：ホスト向け関数をデバイスカーネルから呼び出そうとしている
    initialize_runtime();
});

それでも/アンマウントするには別の処理と権限が必要なため、コンテナ rootfs 消失の原因にはならない。（これらの関数は通常、SYCLのデバイスコンパイルまたはdevice linkで拒否される。仮にホスト側コードパスとして実行されたとしても、/をアンマウントするにはmount関連の操作と、通常はCAP_SYS_ADMINなどの権限が必要）

容疑者D：Podman / OverlayFS の問題

残念ながら、発生当時の Podman events、container inspect 等のログは保存していなかった。したがって、真犯人を断定することはできない。
当時の状況を振り返ると、GPUカーネルの不正アクセスがコンテナのファイルを直接消去したというより、Podmanのコンテナ停止、rootfsのマウント状態、OverlayFS、あるいは関連するストレージ層で異常が起きた可能性の方が、現象とは整合しそうである。

ただし、発生時のpodman events、podman inspect、kernel log、mount情報を保存していなかったため、rootfsが実際にアンマウントされたことまでは確認できていない。
「ファイルがすべて削除された」のではなく、ファイルを見せていたコンテナのrootfsや、そのマウント状態に問題が発生したのではないか――というのが、現在もっとも疑っている筋書きである。

終わりに

梅雨も終わり、今年も暑い夏となりました。
この一件の真相は闇の中ですが、以下の見解が得られました。

• ホストポインタの不正参照：動作未定義だが、通常はGPUカーネルやCUDA contextのエラーとして現れる
• カーネル内のmalloc/free：標準SYCLでは動的なストレージ確保自体が許可されていない
• pipe/signal：NVPTXカーネルから通常のホストOS APIとして実行できるものではない
• Podman/OverlayFSまたはコンテナ停止：現象とは比較的整合するが、ログがないため断定できない

ファイルが消えたのではなく、Podmanが見せていたrootfsや、そのマウント状態に何らかの問題が発生した可能性を疑っています。
私のような非力なエンジニアにとって、環境が潰れることほど恐ろしいものはありません...　

終わり

当時の環境

• コンパイラ：intel/llvm repository sycl branch
• nvidia/cuda 12.3.2, ubuntu24.04 (podman container)