# bpftime GPU Support - CUDA and ROCm eBPF Attachment [[bpftime]] 公式ドキュメント([[eunomia-bpf]]、`eunomia.dev/bpftime/documents/gpu/`)の GPU サポートページ。初出 2025-09-30、最終更新 2026-05-21。著者クレジットは [[Yusheng Zheng]]・[[Tong Yu]]・[[Yiwei Yang]] で、[[@2025__eunomia.dev__The GPU Observability Gap - Why We Need eBPF on GPU devices]] ブログ記事と同一著者陣である。ブログ記事が問題意識と概念設計を一般向けに説明するのに対し、本ページは bpftime の GPU 実装(アタッチ型・専用マップ・ヘルパー関数・ビルド手順)を開発者向けに詳述する、いわば実装リファレンスの位置づけにある。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) ## 問題設定: 同期/非同期実行と可視性の喪失 本ページはブログ記事([[@2025__eunomia.dev__The GPU Observability Gap - Why We Need eBPF on GPU devices]])と同じ問題意識を、より具体的な数値タイムラインで再提示する。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) - **同期実行**: `cudaMemcpy()` が 200μs、`cudaDeviceSynchronize()` が 115μs という壁時間は測れるが、115μs の内訳(起動オーバーヘッド~5μs、実カーネル実行~100μs、クリーンアップ~10μs)は CPU 側からは不可視。 - **非同期実行**: `cudaMemcpyAsync()`/`cudaLaunchKernel()` はキュー投入後即座に戻り、`cudaStreamSynchronize()` の 455μs という単一集計値だけが見える。最初の転送が unpinned memory で 400μs に伸びても、CPU 側は合計 655μs への増加しか観測できず、どの操作が遅延を起こしたか分からない。 - バッチ処理・永続カーネル・マルチストリーム依存・shared memory 制約を伴う LLM サービングのような現代ワークロードは、この非同期時の不可視性をさらに悪化させる。 ## Nsight Systems/Compute の限界(新規の詳細な整理) 本ページはブログ記事の 3 類型整理をさらに掘り下げ、NVIDIA Nsight Systems/Compute 固有の 4 つの限界を列挙する。ベンダー専用プロファイラを「高オーバーヘッド」と一括りにせず、具体的な設計上の制約として分解した点が新規の知見である。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) 1. **クローズドなイベントモデル**: 固定イベント集合のみで、アタッチポイントでの任意プログラマブルロジック(例:「カーネル実行が100msを超えたときだけ収集」というフィルタ述語)を書けない。 2. **本番常時運用に不適合**: カウンタ多重化・リプレイ機構を伴う特殊なプロファイリングセッションが必要で、ワークロード挙動を摂動させ、一過性の異常やレアイベントを取りこぼす。 3. **in-situ フィルタリングの欠如**: 生データを全量エクスポートしてから後処理するため、大規模非同期パイプラインで数 GB のトレースが生じ、観測状態に応じた適応的サンプリング変更ができない。 4. **NVIDIA 専用**: AMD/Intel 等ベンダー中立な展開経路がなく、`kprobes`/`uprobes`/tracepoints との統合や CPU/GPU 間でのマップ共有ができないため、「ホストのページフォールト → 遅延したカーネル起動投入 → ワープストール急増」のような因果連鎖の突き合わせが困難。 ## アーキテクチャ: CUDA Attachment Pipeline GPU サポートは `nv_attach_impl` システム(`attach/nv_attach_impl/`)上に構築され、アプリケーションプロセス内で CUDA App → bpftime Runtime → eBPF 付き GPU Kernel という計装パイプラインを形成する。Runtime はホスト側の共有メモリ(Host-GPU 通信)と GPU メモリ(IPC)の両方を介して GPU カーネルと接続する。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) ### アタッチ型(3種) `attach/nv_attach_impl/nv_attach_impl.hpp:33-34` で定義される。 - **`ATTACH_CUDA_PROBE`(ID 8)** — カーネル入口で eBPF コードを実行。 - **`ATTACH_CUDA_RETPROBE`(ID 9)** — カーネル出口で eBPF コードを実行。 - **メモリキャプチャプローブ(`__memcapture`)** — メモリアクセスパターンをキャプチャする専用プローブ型。 いずれもマングルされた C++ 関数名(例: `_Z9vectorAddPKfS0_Pf`)でターゲットカーネル関数を指定する。 ### 主要コンポーネント 1. **CUDA Runtime Hooking** — Frida ベースの動的計装で CUDA API 呼び出しを傍受。 2. **PTX Modification** — eBPF バイトコードを PTX に変換し GPU カーネルへ注入([[PTX 注入]])。 3. **Helper Trampoline** — マップ操作・タイミング・コンテキストアクセス用の GPU アクセス可能なヘルパー関数を提供。 4. **Host-GPU Communication** — ピン留め共有メモリを介した GPU→ホストの同期呼び出しを実現。 ## GPU 専用 BPF マップ bpftime は CPU 上の全マップ型をサポートするが、CPU 用 eBPF マップを GPU コードから直接使うと性能が大きく劣化するため、性能影響の小さい GPU 専用マップ型が用意されている。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) - **`BPF_MAP_TYPE_PERGPUTD_ARRAY_MAP` / `BPF_MAP_TYPE_GPU_ARRAY_MAP`** — GPU メモリ(CUDA IPC 共有メモリ)常駐の per-thread ストレージ(`max_entries × max_thread_count × value_size`)。GPU からのゼロコピーアクセスとホストへの DMA 転送、`bpf_map_lookup_elem()`/`bpf_map_update_elem()` を GPU コードから直接呼べる。実装: `runtime/src/bpf_map/gpu/nv_gpu_array_map.cpp:14-81`。 - **`BPF_MAP_TYPE_GPU_RINGBUF_MAP`(ID 1527)** — GPU メモリ上のロックフリー per-thread リングバッファ。可変長イベントレコードと低オーバーヘッドな非同期データ収集を提供し、`bpf_perf_event_output()` ヘルパーと互換。実装: `runtime/src/bpf_map/gpu/nv_gpu_ringbuf_map.cpp`。 ## GPU ヘルパー関数 `attach/nv_attach_impl/trampoline/default_trampoline.cu:331-390` で定義される GPU 専用ヘルパー(ヘルパー ID 501-506)は、`ebpf_puts`(GPU→ホストコンソール出力)、`bpf_get_globaltimer`(ナノ秒精度の GPU グローバルタイマ)、`bpf_get_block_idx`/`bpf_get_block_dim`/`bpf_get_thread_idx`(CUDA の blockIdx/blockDim/threadIdx 取得)、`bpf_gpu_membar`(`membar.sys` メモリバリア)から成る。標準 BPF ヘルパー(`bpf_map_lookup_elem`・`bpf_map_update_elem`・`bpf_trace_printk`・`bpf_get_current_pid_tgid`・`bpf_perf_event_output` 等)も GPU 上で動作するよう最適化されている(GPU array map へのファストパス、他は共有メモリ経由でホストへフォールバック)。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) ### Host-GPU 通信プロトコル ホスト対話が必要なヘルパーは、スピンロックを用いた共有メモリプロトコルとワープレベルの直列化で正当性を担保する。手順は (1) GPU スレッドがスピンロック取得、(2) 要求パラメータを共有メモリへ書き込み、(3) フラグを立ててホスト応答を待機、(4) ホストが要求を処理し完了を通知、(5) GPU が応答を読みロックを解放、という 5 段階。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) ## 実装例(bcc スタイルツール) - **kernelretsnoop** — CUDA カーネル出口にアタッチし、各 GPU スレッドが完了したナノ秒タイムスタンプを記録。ワープ内 thread 31 が他スレッドより 750ns 遅れて終わる境界条件による分岐を検出し、分岐排除後に全スレッドがナノ秒単位で揃うことを確認できる。 - **threadhist** — GPU array map で各スレッドの実行回数をヒストグラム化し、ワークロード不均衡を検出。1M 要素を 5 スレッドの grid-stride loop で処理する例で、thread 4 が他スレッドの 75% しか実行されない偏りを発見。 - **launchlate** — CPU 側 `cudaLaunchKernel()` 呼び出し時刻と GPU 側実カーネル実行開始時刻の差分を測定するカーネル起動レイテンシプロファイラ。カーネル自体は 100μs で完了するのにユーザー報告が 50ms である例で、各カーネルが前段のカーネルとメモリ転送完了を待つキュー遅延(200-500μs/カーネル)を可視化し、CUDA graphs によるバッチ化への切り替えでレイテンシが 1.2ms に低下したことを示す。 - その他: `cuda-counter`(基本 probe/retprobe)、`mem_trace`(メモリアクセスパターン解析)、`directly_run_on_gpu`(カーネルにアタッチせず GPU 上で直接 eBPF を実行)、`rocm-counter`(AMD ROCm 計装、実験的)。 いずれも動的アタッチ(preload なしで起動後に `bpftime trace` でアタッチする `threadscheduling_dynamic_hook` 等)にも対応する。(Source: [[.raw/articles/bpftime-gpu-documentation-2026-08-03.md]]) ## 性能主張とビルド - **NVBit 比で 3-10 倍高速**な計装を主張。 - **ベンダー非依存設計**(NVIDIA・AMD・Intel GPU で動作)。 - ビルドには NVIDIA CUDA Toolkit(12.x 推奨)または AMD ROCm、CMake 3.15+、LLVM 15+(PTX 生成用)、Frida-gum が必要。`cmake -Bbuild -DBPFTIME_ENABLE_CUDA_ATTACH=1 -DBPFTIME_CUDA_ROOT=/usr/local/cuda-12.6` でビルド設定する。 ## 注意点 - GPU 対応は依然**実験的**。 - ドキュメント内でも「`kprobe`/`kretprobe` という名称は仮のプレースホルダで、変更されうる」と明記されている(ブログ記事の注意点と同一)。 ## 関連 - 先行記事: [[@2025__eunomia.dev__The GPU Observability Gap - Why We Need eBPF on GPU devices]](同一著者陣による問題設定の一般向け解説) - 元論文: [[@2025__HCDS__eGPU - Extending eBPF Programmability and Observability to GPUs]] - 同一プロジェクトの入門実装: [[@2026__eunomia.dev__CUDA Events - eBPF-based CUDA API Tracing]] - 開発元: [[eunomia-bpf]] / [[bpftime]] - 著者: [[Yusheng Zheng]] / [[Tong Yu]] / [[Yiwei Yang]] - 概念: [[GPU観測性]] / [[eGPU]] / [[PTX 注入]] / [[eBPF]] / [[CUDA]] / [[CUDA API トレース]] / [[uprobe]] - エンティティ: [[NVIDIA]] / [[AMD]] / [[CUPTI]] / [[NVBit]]