Snooping on your GPU - Using eBPF to Build Zero-instrumentation CUDA Monitoring

GPUprobeはLinuxの[[eBPF]]技術を活用し、CUDAアプリケーションをコード変更なしで監視します - 再コンパイルもプロファイリング用コードの挿入も不要で、単にアタッチして実行するだけで動作します。 [リポジトリはこちら](https://github.com/GPUprobe/gpuprobe-daemon) をご覧ください！ ## はじめに私がソフトウェアエンジニアとして初めて仕事を始めた時、最初の課題はGoで書かれたサービスで発生したメモリリークを修正することでした。これによりサービスが頻繁にクラッシュして再起動を繰り返す状態になっていました。問題の本質は、GoコードがCgoを介してCライブラリを呼び出していたため、ガベージコレクション機能を備えたGoランタイムの保護範囲から外れ、生の `malloc()` および `free()` メモリ管理関数を直接使用する状態になっていたことです。このため、通常のGoプロファイリングツールは全く役に立たなくなっていました。そこで私は[[Valgrind]]やその他のツールを試しましたが、Goバイナリではこれが全く機能しませんでした。このツールはメモリアクセスをすべて監視するため、処理速度が大幅に低下し、バイナリを起動することすらままならず、ましてやデバッグなど不可能でした。私は必死に別の解決策を探し始めました... ここで登場するのが [BCC-memleak](https://github.com/iovisor/bcc/blob/master/tools/memleak.py) です。これはeBPFベースのツールで、実行中のプロセスにアタッチすることでメモリリークを検出できます。必要なのは、デバッグフラグ付きでリークしているライブラリをコンパイルし、memleakを対象プロセスにアタッチするだけ。これだけで問題箇所を特定できます。30分以内に、メモリリークを引き起こしている正確な関数呼び出しを特定し、1回の呼び出しあたりのメモリリーク量を計測し、1行の修正コードを含むPRを作成することができました。数ヶ月後、サイドプロジェクトとして取り組める可能性のあるアイデアについて考えていた時、BCC Memleakを使用した経験とその時間節約効果を思い出しました。GPUメモリの割り当てにもこのようなツールが適用できるのではないか、また他にもどんな面白いことが実現できるのかと想像を膨らませました。こうして誕生したのが **GPUprobe** です。CUDA向けのeBPFベースの可観測性ツールです。現在、このツールはメモリ割り当てパターン、メモリリーク、カーネル起動パターンなどの情報を提供しており、今後さらに多くの機能が追加される予定です。 ## GPUモニタリングに欠けているもの GPUは高価なハードウェアです。場合によっては「非常に」高価な場合もあります。GPUユーザーとして、可能な限り最大限の性能を引き出すことは、自分自身と財布に対する義務と言えるでしょう。さらに、一般的なシステムソフトウェアと同様、デバッグと障害検出は容易ではありません。CUDAの補助関数は通常 `cudaError_t` 型の値を返しますが、これは単なる列挙型値に過ぎません。当然ながらコード内では常にエラー処理を行うべきですが、実行時の可観測性を確保するためには、以下のようなエラーチェック用のステートメントをコード中に散りばめる必要があり、何か問題が発生した際には標準出力を必死に確認しなければなりません。 ``` cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { printf("CUDA error: %s\n", cudaGetErrorString(err)); } ``` 当然ながら、GPUの監視に関しては既に様々なソリューションが存在しており、それぞれに長所と短所があります。 ## NVIDIA NSight Systems NSight SystemsはNVIDIAが提供する主要なGPUプロファイリングツールです。開発時のプロファイリングと最適化において非常に強力な機能を備えており、CUPTI（CUDAプロファイリングツールインターフェース）を使用してGPU使用状況に関する詳細なメトリクスを収集します。ただし、そのワークフローは継続的なモニタリングとは根本的に異なります： - プロファイリングセッションを明示的に開始する必要があります - プロファイリングを行うと、多くの場合2～10倍のパフォーマンス低下という重大なオーバーヘッドが生じます - データ分析はプログラムの実行完了後に行われます - 継続的な本番環境モニタリングを目的としたツールではありません。これは監視ツールというよりも、GPU専用のデバッガーに近い機能を持っています。開発段階でのボトルネック特定には優れていますが、継続的な本番環境モニタリングには適していません。 ## DCGM（Data Center GPU Manager） DCGM（Data Center GPU Manager）はNVIDIAが提供するデータセンター向けGPU監視ソリューションです。システムレベルのメトリクス収集に優れており、以下のような指標を取得できます： - GPU使用率 - メモリ使用量（総使用量/空き容量） - 温度および消費電力 - ハードウェアの健康状態ただし、DCGMは高レベルの監視機能を提供するため、アプリケーション固有の詳細な情報は取得できません： - プロセスごとのメモリ割り当てパターンに関する可視性が欠如しています - 個々のCUDAカーネル起動を追跡できません - メモリリークの検出機能が限定的 - APIレベルの動作に関する詳細な情報は取得できません ## GPUprobeの位置付け GPUprobeは、GPU監視における特定のギャップを埋めるツールです。アプリケーションレベルで軽量かつ継続的な監視を実現し、以下の機能を提供します： 1. ほぼオーバーヘッドのないランタイム監視を実現： - 既存のシステムコールに乗っかる形でeBPF uprobesを使用します - コードへの変更は一切不要 - アプリケーション性能への影響が最小限（ベンチマークテストでは4%未満） 2. アプリケーションレベルの詳細な分析： - プロセスごとのメモリ割り当て状況と潜在的なメモリリークを追跡します - CUDAカーネルの起動パターンを監視 - 実際の関数名と呼び出しパターンを確認可能 - 本番環境でAPIレベルの問題をデバッグ可能 3. 最新の可観測性統合機能： - Grafanaダッシュボード用のPrometheusメトリクスをエクスポートします - 本番環境でも使用可能な継続的なモニタリング機能 - 既存の監視システムに容易に統合可能 GPUprobeは、eBPF uprobesを利用してLinuxカーネルレベルでCUDAランタイムAPIコールを監視することでこの機能を実現しています。この独自のアプローチにより、アプリケーションコードを変更することなく、またパフォーマンスへの影響を最小限に抑えながら、詳細なメトリクスを収集することが可能です。NSightの詳細だがリソース消費の大きいプロファイリングと、DCGMの高レベルなシステム監視の中間的な位置づけと考えてください。 ## eBPF とその魅力についてここではeBPFで実現できるすべての機能について詳しく説明することはしませんが、より詳細な情報については [eBPF公式サイト](https://ebpf.io/) をご覧ください。以下に公式サイトからの抜粋を掲載します： > eBPF（Extended Berkeley Packet Filter）は、Linuxカーネルを起源とする革新的な技術で、オペレーティングシステムカーネルなどの特権環境でサンドボックス化されたプログラムを実行できます。これは、カーネルソースコードを変更したりカーネルモジュールをロードしたりすることなく、カーネルの機能を安全かつ効率的に拡張するために使用されます。従来、オペレーティングシステムはカーネルがシステム全体を監視・制御できる特権的な立場にあるため、可観測性、セキュリティ、ネットワーク機能を実装する理想的な場所とされてきました。... ## アップローブ（uprobes） eBPFの最も強力な機能の一つが、uprobeを介してユーザー空間プログラムにアタッチできる点です。uprobeを、実行中のプログラム内の任意の関数に貼り付け可能な顕微鏡レベルのブレークポイントと考えてください。この関数が呼び出されると、eBPFプログラムが通知を受け取り、呼び出しに関するデータの検査や収集を行うことができます。これらの機能が特に強力な理由は、対象プログラム自体を変更する必要がなく、従来の計測手法と比較してオーバーヘッドが極めて少ない点にあります。 GPUprobeでは、uprobesが中核技術として機能します。これらはCUDAランタイムAPI（ `libcudart.so` ）に直接アタッチされ、 `cudaMalloc()` 、 `cudaFree()` 、 `cudaLaunchKernel()` などの関数呼び出しによってトリガーされます。プログラムから関数が呼び出されると、私たちのeBPFプログラムがそれをインターセプトし、必要なデータを収集して、プログラム側に知られることなく監視パイプラインに送信します。 ## 事例研究：メモリリーク検出ツールの実装 CUDAメモリリークをリアルタイムで検出するツールの実装方法について詳しく説明します。大まかに言うと、私たちは各プロセスごとにCUDAメモリマップを保持しています。メモリは `cudaMalloc()` 関数の呼び出しによって割り当てられ、対応する `cudaFree()` 関数の呼び出しによって解放されます。これらの関数の署名について、初めてご覧になる方のために説明します： ``` // allocate \`size\` bytes on device, device address is copied to \`*devPtr\` cudaError_t cudaMalloc (void** devPtr, size_t size) /// free an allocation at device address \`devPtr\` cudaError_t cudaFree (void* devPtr) ``` eBPFを使用してこの監視を実現する基本的な考え方を、Python風の擬似コードで説明します。 ``` class MemoryMaps: ... memory_maps = MemoryMaps() def uprobe_cuda_malloc(devPtr, size): pid = get_pid() memory_maps.make_entry(pid, *devPtr, size) def uprobe_cuda_free(devPtr): pid = get_pid() memory_maps.free_entry(pid, devPtr) def process_exits(pid): memory_maps.free_all(pid) ``` 実にシンプルな仕組みです！ *ほぼ*... まず、メモリマップをカーネル領域に配置することは避けたいと考えています。メトリクスをエクスポートする際には、常にデータ構造を参照し、どのメモリ割り当てが未処理のまま残っているかを確認し、さらに終了したプロセスのデータ構造をクリーンアップする必要があります。一方、uprobeはこれらのデータ構造へのアクセスを競合します。CUDAランタイム関数によってトリガーされた場合でも、ロック待ちが発生すると、アプリケーションの処理速度が低下してしまいます。そこで私たちは、「イベント駆動型」システムを採用することにしました。このシステムはeBPFキューを使用して実装されており、uprobesからプッシュされたイベントをユーザー空間プログラムが消費します。このキューには、 `cudaMalloc()` や `cudaFree()` の呼び出しに関して把握したい重要な情報が格納されます。 ``` /** * Wraps the arguments passed to \`cudaMalloc\` or \`cudaFree\`, and return code, * and some metadata */ struct memleak_event { __u64 start; __u64 end; void *device_addr; __u64 size; __u32 pid; int32 ret; enum memleak_event_t event_type; }; /** * Queue of memleak events that are updated from eBPF space, then dequeued * and processed from user-space by the GPUprobe daemon. */ struct { __uint(type, BPF_MAP_TYPE_QUEUE); __uint(key_size, 0); __type(value, struct memleak_event); __uint(max_entries, 1024); } memleak_events_queue SEC(".maps"); ``` `cudaMalloc()` アップローブを実装する際に遭遇した課題として、 `*devPtr` にコピーされるデバイスアドレスは関数が返却された時点で初めて確定するという点が挙げられます。このため、当然ながら関数返却時にトリガーされる uretprobe *（これはアップローブと同様の機能ですが、関数の返却時に発動します）* を使用する必要があります。しかし、この方法にも限界があります。uprobesとuretprobesは `struct pt_regs *ctx` からデータを読み取るため、これはレジスタ状態のスナップショットを取得することを意味します。つまり、単にuretprobeを使用するだけでは不十分です。なぜなら、関数実行中にレジスタの内容は変化し、任意のデータを含む可能性があるからです。重要な注意点として、eBPFプログラムは相互に呼び出しを行うことができず、また `bpf.h` で公開されている許可されたヘルパー関数以外の任意の関数を呼び出すこともできません。このため、 `cudaMalloc()` のuprobeとuretprobe間でデータを共有する際には、特定のプロセスに対応する `devPtr` を保持するeBPFハッシュマップを使用します。この前提は、CUDAの公式ドキュメントでも支持されている通り、 `cudaMalloc()` がブロッキング関数であり、同一スレッドから同時に2回以上呼び出すことはできないという特性に基づいています。大まかに言えば、ロジックの流れは以下の通りです： ``` pid_to_devPtr = {} def uprobe_malloc(devPtr, size): pid = get_pid() pid_to_devPtr[pid] = devPtr def uretprobe_malloc(devPtr): pid = get_pid() # we must read from user-space to get the device address that was copied # into \`void** devPtr\`. Think of this as a pointer deref, except that we # are deferencing something in user-space device_address = read_from_user(pid_to_devPtr[pid]) ``` 実際の実装では、その他の有用なメタデータを収集することで、以下のような機能を実現しています ``` /// uprobe triggered by a call to \`cudaMalloc\` SEC("uprobe/cudaMalloc") int memleak_cuda_malloc(struct pt_regs *ctx) { struct memleak_event e = { 0 }; void **dev_ptr; u32 pid, key0 = 0; e.size = (size_t)PT_REGS_PARM2(ctx); dev_ptr = (void **)PT_REGS_PARM1(ctx); pid = (u32)bpf_get_current_pid_tgid(); e.event_type = CUDA_MALLOC; e.start = bpf_ktime_get_ns(); e.pid = pid; if (bpf_map_update_elem(&memleak_pid_to_event, &pid, &e, 0)) { return -1; } return bpf_map_update_elem(&memleak_pid_to_dev_ptr, &pid, &dev_ptr, 0); } /// uretprobe triggered when \`cudaMalloc\` returns SEC("uretprobe/cudaMalloc") int memleak_cuda_malloc_ret(struct pt_regs *ctx) { int cuda_malloc_ret; u32 pid, key0 = 0; size_t *size, *num_failures; struct memleak_event *e; void **dev_ptr; void ***map_ptr; cuda_malloc_ret = (int)PT_REGS_RC(ctx); pid = (u32)bpf_get_current_pid_tgid(); e = bpf_map_lookup_elem(&memleak_pid_to_event, &pid); if (!e) { return -1; } e->ret = cuda_malloc_ret; // lookup the value of \`devPtr\` passed to \`cudaMalloc\` by this process map_ptr = (void ***)bpf_map_lookup_elem(&memleak_pid_to_dev_ptr, &pid); if (!map_ptr) { return -1; } dev_ptr = *map_ptr; // read the value copied into \`*devPtr\` by \`cudaMalloc\` from user-space if (bpf_probe_read_user(&e->device_addr, sizeof(void *), dev_ptr)) { return -1; } e->end = bpf_ktime_get_ns(); return bpf_map_push_elem(&memleak_events_queue, e, 0); } ``` これで完了です。ユーザー空間コードは `memleak_events_queue` からデータを取得し、メモリマップを更新できるようになりました。 [![High-level architecture diagram](https://media2.dev.to/dynamic/image/width=800%2Cheight=%2Cfit=scale-down%2Cgravity=auto%2Cformat=auto/https%3A%2F%2Fdev-to-uploads.s3.amazonaws.com%2Fuploads%2Farticles%2Ftz70lyccerp28ez6541a.png)](https://media2.dev.to/dynamic/image/width=800%2Cheight=%2Cfit=scale-down%2Cgravity=auto%2Cformat=auto/https%3A%2F%2Fdev-to-uploads.s3.amazonaws.com%2Fuploads%2Farticles%2Ftz70lyccerp28ez6541a.png) `cudaFree()` のuprobeとuretprobeの実装は、先ほど説明した `cudaMalloc()` の実装と非常に類似しています。ユーザー空間で処理パイプラインを実行する利点の一つは、ユーザー空間プログラミング言語が提供する豊富なデータ構造を活用できる点です。具体的には以下のものが挙げられますが、これらに限定されません： - 動的サイズのデータ構造：eBPFマップは静的なサイズを指定するか、プログラムをアタッチする前に初期化時に明示的にサイズを設定する必要があります。 - 入れ子構造のデータ構造：各プロセスごとのメモリマップは、Bツリーマップのハッシュマップとして実装しています。具体的には、プロセスID（PID）をキーとして、そのプロセスのメモリマップをBツリーマップで管理します。このメモリマップは、CUDAデバイスアドレスの順序付き範囲と対応するサイズメタデータを保持しています。ユーザー空間処理パイプラインはRustで実装されており、 [libbpf](https://github.com/libbpf/libbpf) のRustバインディングを提供する `libbpf-rs` ライブラリを利用しています。このライブラリは、eBPFプログラムのアタッチ処理やライフサイクル管理、eBPFマップへのアクセスのための使いやすいAPIを提供しています。 uprobesによって生成されたイベントキューは、標準出力への表示時またはエクスポート時に処理されます。表示はユーザーが設定可能な固定間隔（デフォルトは5秒）でトリガーされ、メトリクスエンドポイント `:9000/metrics` へのリクエストがあるたびにエクスポートがトリガーされます（このポート番号もユーザー設定可能です）。キューの処理は比較的単純で、eBPFキューからイベントがなくなるまで順次ポップしていくだけです。なお、このキューには `cudaMalloc()` と `cudaFree()` のuprobeで生成されたイベントの両方が含まれている点に注意してください。 ``` let key: [u8; 0] = []; // key size must be zero for BPF_MAP_TYPE_QUEUE // \`lookup_and_delete\` calls. while let Ok(opt) = self .skel .skel .maps .memleak_events_queue .lookup_and_delete(&key) { let event_bytes = match opt { Some(b) => b, None => { return Ok(()); } }; let event = match MemleakEvent::from_bytes(&event_bytes) { Some(e) => e, None => { return Err(GpuprobeError::RuntimeError( "unable to construct MemleakEvent from bytes".to_string(), )); } }; // update CUDA state ``` 保持する状態情報の構造は以下の通りです： ``` pub struct MemleakState { pub memory_map: HashMap<u32, BTreeMap<u64, CudaMemoryAlloc>>, pub num_successful_mallocs: u64, pub num_failed_mallocs: u64, pub num_successful_frees: u64, pub num_failed_frees: u64, active_pids: HashSet<u32>, } ``` 追跡対象プロセスのセット（ `active_pids` ）を保持しており、観測したCUDAプログラムの生存確認を行うとともに、プロセスが終了する際には関連する状態情報を適切にクリーンアップします。 ## メモリリークとカーネル起動のリアルタイム監視 - ライブデモ！プログラムが正しく動作しており、 `cudaMalloc()` 、 `cudaFree()` 、および `cudaLaunchKernel()` イベントを正常に検知していることを実証するため、 `gpuprobe-daemon` を起動し、シンプルなCUDAバイナリを実行します。では、プログラムの出力結果にはどのような情報が含まれているのでしょうか？ ``` 2024-21-12 16:32:46 num_successful_mallocs: 3 num_failed_mallocs: 0 num_successful_frees: 0 num_failed_frees: 0 per-process memory maps: process 365159 0x0000793a44000000: 8000000 Bytes 0x0000793a48c00000: 8000000 Bytes 0x0000793a49400000: 8000000 Bytes total kernel launches: 1470 pid: 365159 0x5de98f9fba50 (_Z27optimized_convolution_part1PdS_i) -> 735 0x5de98f9fbb30 (_Z27optimized_convolution_part2PdS_i) -> 735 ============================== 2024-21-12 16:32:51 num_successful_mallocs: 3 num_failed_mallocs: 0 num_successful_frees: 2 num_failed_frees: 0 per-process memory maps: process 365159 0x0000793a44000000: 8000000 Bytes 0x0000793a48c00000: 0 Bytes 0x0000793a49400000: 0 Bytes total kernel launches: 2000 pid: 365159 0x5de98f9fba50 (_Z27optimized_convolution_part1PdS_i) -> 1000 0x5de98f9fbb30 (_Z27optimized_convolution_part2PdS_i) -> 1000 ``` なるほど！正常に動作しているようです。では、具体的にどのような情報が表示されているのでしょうか？最初の計測間隔では、プログラムが実行中であることがわかります。GPU上で必要なすべてのメモリが割り当てられており、2つのCUDAカーネルがそれぞれ735回起動されていることがわかります。まず、プログラムバイナリ内で検出されたカーネル関数のアドレスと、その関数名 *（これは関連バイナリからシンボルを解決することで取得しました）* が表示されます。次に、実際に行われたメモリ割り当ての数、GPU上の仮想アドレス、および各割り当てのサイズを確認できます。 2番目の計測間隔では、両方のカーネルがそれぞれ設定した回数（1000回）実行されていることがわかります。また、メモリの2つの領域が解放されていることも確認できます。しかし、GPUメモリを占有し続けているように見える3番目の割り当てはどうなっているのでしょうか？コードをよく見てみると、 `cudaFree(dv_intermediate)` を呼び出すのを忘れていたことがわかります！これは単なるサンプルコードですが、目的とする機能を比較的よく示していると思います。 - プログラム実行中にメモリリークを検出しました。 - 各CUDAカーネルが何回起動されたかを正確に把握できます - コードの変更は一切必要なく、GPUprobeを特定のプロセスにアタッチする必要もありませんでしたここで紹介したサンプルプログラムは、私が学部時代に取り組んだ授業課題を簡略化したバージョンです。これは私が最初にGPUprobeのテストに使用したものでした。興味深いことに、当初はメモリリークが発生するとは予想しておらず、後になってGPUprobeを使用して初めてその問題を発見しました。当時このようなツールがあれば、間違いなくメモリリークのある課題を提出することはなかったでしょう。 ## バグ情報私が遭遇したバグについて共有します。この問題については、現在も *（必ずしも積極的にではなく）* 解決を試みているところです。 CUDAカーネルのシンボル解決を行い、起動されたカーネルの名称を標準出力に表示します。まるで魔法のようですが、この機能を実現するには実に多くの試行錯誤が必要でした。これは非常に有用な機能だと思います。仮想関数アドレスとCUDAカーネルという抽象概念を結びつけるのは、通常非常に困難な作業だからです。 uprobeによって生成されたCUDAカーネル起動イベントには、カーネルを起動したプロセスのPIDが含まれています。これは非常に有用な情報です！プロセスIDが分かれば、 `/proc/[pid]` ディレクトリを参照することで、実行中のバイナリのパス情報と、その仮想ベースアドレス *（ASLR（アドレス空間配置ランダム化）により実行ごとに変化する）* を取得できます。ここから、従来のシンボル解決手法を用いて、実行時にCUDAカーネルの仮想アドレスから関数名を特定できます。これはCUDAカーネルが関数として記述されているため、バイナリの`.text` セクションに格納されているからです。これで完了です。ユーザーに対して、 `cudaLaunchKernel(const void* func, ...)` に渡されているカーネルの名前という、より人間が理解しやすい形式で出力を表示できるようになりました。大まかに言えば、シンボル解決の処理は以下のように行われます： ``` # Map from binary offset -> symbol name symbols = {} # e.g. {0x1000: "my_cuda_kernel", ...} def resolve_symbol(virtual_addr, pid): # Adjust for ASLR by subtracting the binary's base address offset = virtual_addr - get_virtual_base_offset(pid) return symbols[offset] # Get the symbol name for this offset ``` この手法はほとんどの場合において非常に有効です... *ただし例外もあります* 。イベント処理は、数秒ごとにディスプレイ表示またはエクスポートイベントが発生するたびに行われます。イベントを記録するたびに毎回 `/proc` ファイルを読み込む必要がないよう、シンボルテーブルをキャッシュしています。新しいPIDのイベントが記録された場合にのみ、この処理を実行します。問題が発生するのは、短い実行期間を持つプロセスが2つのタイムインターバルの間だけ起動・終了する場合です。このようなケースではイベントの記録が困難になります。イベントキューを処理する際、 `pid` に対応するキャッシュされたシンボルテーブルが存在するかどうかを確認します。もし特定のPIDが初めて記録された場合、当然ながらそのエントリは存在しません。そのため、 `/proc/pid` ディレクトリを参照しますが、...実際には存在しません。プロセスはすでに終了しているため、バイナリの仮想ベースオフセットや実行元バイナリの位置を把握できません。このため、シンボル解決を行うことはできません。これは少々残念な結果です。周波数ヒストグラムの記録と表示は可能ですが、カーネルのシンボル解決ができなかったため、例として示しているように単に `不明なカーネル` と表示することになります。 ``` total kernel launches: 25 pid: 365306 0x5823e39efa50 (unknown kernel) -> 10 0x5823e39efb30 (unknown kernel) -> 15 ``` シンボル解決の実装には常に複数の方法が存在します。代替案としては、システム全体でプロセスの終了を監視し、必要に応じて終了プロセスのシンボルテーブルをキャッシュする方法が考えられます。これは本プロジェクトの範囲外の問題であり、Linuxシステム上のほとんどのプロセスはCUDAランタイムAPIを使用しないため、不要なデータが大量に収集される可能性が高いと考えられます。ここで採用されているトレードオフは妥当な判断だと思います。ほとんどのCUDAジョブは、単一の表示/エクスポート間隔よりも長時間実行されるからです。少なくとも、監視対象とするCUDAジョブであればなおさらです。 ## パフォーマンスベンチマーク最後に結論に至る前に、いくつかのベンチマーク結果について考察します。まず、uprobesを使用する場合、一般的にかなりのオーバーヘッドが生じることに注意が必要です。これはコンテキストスイッチが発生するためです *（カーネルがユーザー空間からの関数呼び出しを監視するため）* 。uprobesを使用したCUDAランタイム監視においてオーバーヘッドが最小限に抑えられると期待できる理由は、以下の点にあります： - CUDAランタイムAPIはカーネル空間に存在するCUDAドライバを呼び出します。これもコンテキストスイッチを引き起こす要因となります！ - CUDAランタイムAPI関数の呼び出しは、PCIe *（またはその他のインターコネクト）* 経由でGPUとの通信を引き起こしますが、これらの関数のほとんどはブロッキング処理です。つまり、CUDAランタイムAPI関数自体が非常にリソースを消費するものであり、通常はGPUに対して実行する処理を指示するために使用されます。このAPIはスループット最適化のためにバッチ処理を前提に設計されており、低遅延のシングルスレッド性能を追求したものではありません。この手法の利点は、uprobesによって生じる追加の遅延が、CUDAランタイムAPI呼び出し自体の遅延に比べて無視できるほど小さいため、相対的なオーバーヘッドがほとんど発生しない点にあります。本記事ではベンチマークを非常にシンプルに保ち、より厳密なベンチマークは後日実施する予定です。テストはNVIDIA Quadro P520（VRAM 2048MiB）を搭載したノートPCで実施しました。私のシステムでは、 `cudaMalloc()` と `cudaFree()` のペアを5,000回実行し、各反復処理の平均レイテンシを測定しました。結果がGPU上で連続した大容量メモリブロックを割り当てる際のオーバーヘッドではなく、 `cudaMalloc()` 呼び出し自体のオーバーヘッドをより正確に反映するよう、割り当てサイズを100バイトと非常に小さく設定しました。最初の500回の実行結果はウォームアップ効果を考慮して破棄しています。 | | GPUprobe未使用時 | GPUprobe 使用時 | | --- | --- | --- | | 平均レイテンシ | 255μs | 265μs | すべての `cudaMalloc()` / `cudaFree()` 呼び出しに監視機能を実装した場合、レイテンシが約3.92%増加することが確認されました。 `cudaLaunchKernel()` 呼び出しの監視に伴うオーバーヘッドについては、ケーススタディで提示したプログラムをベンチマークしました。本プログラムでは、各反復処理で2つのCUDAカーネルを1000回ずつ実行します。このケースでは、GPUprobeを使用した場合と使用しない場合で、測定可能なパフォーマンスへの影響は確認されませんでした。これらのベンチマークから得られた主な知見は以下の通りです： - `cudaMalloc()` および `cudaFree()` への頻繁な呼び出しを監視すると、約4%のオーバーヘッドが発生します。ただしこのケースは現実的とは言えません。実際には、デバイス上でメモリブロックを一度割り当てた後、 *その後* そのメモリに対して多数の操作を行うのが一般的です。 - `cudaLaunchKernel()` の呼び出しを監視する場合、この単純なケースでは実行時オーバーヘッドはほとんど発生しません。これは当然のことです。カーネル起動を監視するための uprobe は、メモリ割り当てイベントを記録するための uprobe に比べて *はるかに単純* （中間状態を保持する必要がない）ためです。現時点での評価では、GPUprobeがCUDAアプリケーションの実行に与えるオーバーヘッドは極めて小さいと言えます。今後のベンチマーク調査では以下の内容を予定しています： - 実際の機械学習ワークロード（大規模データセット）を使用したテスト - メモリ集約型アプリケーションへの影響測定 - さまざまなGPUアーキテクチャとCUDAバージョンでのベンチマークテスト現時点では、これらの結果から、GPUprobeのオーバーヘッドは実用上問題にならない程度に小さいことが示唆されます。ローカル環境や本番環境でも十分に使用可能なレベルです。 ## 結論本稿では、 **GPUprobe** というゼロインストゥルメンテーション型のGPU監視ツールを紹介します。このツールは、eBPF技術を用いてCUDAランタイムAPIコールを検査することで、GPUの動作を監視するものです。本論文では、GPUprobeがGPU監視ツール群の中でどのような独自の位置付けにあるかを探りました。具体的には、NSightやDCGMといった既存ツールが提供していない、低オーバーヘッドでありながらアプリケーションレベルの詳細な情報を提供するという特徴について考察しています。メモリリーク検出器の実装詳細を掘り下げることで、eBPF uprobesを使用することでアプリケーションコードを変更することなく、GPUメモリ割り当てを追跡できる仕組みを明らかにしました。本プロジェクトでは、実行時におけるCUDAカーネルのシンボル解決といった興味深い技術的課題に取り組みました。ベンチマーク結果から、メモリ割り当てが頻繁に行われるワークロードにおいても、オーバーヘッドが極めて小さいことを実証しています。今後の取り組みとしては、実際の機械学習ワークロードでGPUprobeの動作テストを行い、監視ソリューションとしての位置付けをより明確にしていく予定です。特に以下の点に注目しています： - 機械学習インフラストラクチャ監視におけるユースケースの検討 - CUDAランタイムAPI関数のサポート機能をさらに追加 GPU監視技術またはeBPFにご興味のある方は、ぜひリポジトリをご覧ください！スターを付けていただけると、プロジェクトの認知度向上につながります:) [リポジトリへのリンク](https://github.com/GPUprobe/gpuprobe-daemon)