## Documents
ユーザーランドのeBPFランタイムであるbpftimeを開発しているeunomiaから、GPU関連のテレメトリーのゼロ計装方法に関するドキュメントがいくつかでていた。どれも丁寧にサーベイがしてあって、知らないことも結構あった。
- [[eBPF × AILLMs The Convergence of System Observability and Artificial Intelligence - eunomia]]
- AI x eBPFの既存文献が網羅されている
- [[eBPF Tutorial Tracing CUDA GPU Operations - eunomia]]
- [[CUDA]]を[[uprobes]]でトレースする方法
- [[eBPF Tutorial by Example Monitoring GPU Driver Activity with Kernel Tracepoints - eunomia]]
- GPUドライバをtracepointでトレースする方法
- [[Write and Run eBPF on GPU with bpftime - eunomia]]
- GPUカーネルの中でeBPFプログラムを実行する。[[2025__HCDS__eGPU - Extending eBPF Programmability and Observability to GPUs|eGPU]]の話。
| 観点 | (A) eBPF×uprobes：CUDA APIトレーシング | (B) eBPF×kernel tracepoints：GPUドライバ監視 | (C) bpftime：GPU上でeBPF実行 |
| ------------ | -------------------------------------------------------------------------------- | ------------------------------------------------------------------------------------ | ------------------------------------------------------------- |
| レイヤ | **ユーザー空間**（CUDAランタイム） | **カーネル空間**（DRM/各GPUドライバ） | **GPUデバイス内部**（カーネル関数） |
| フック方式 | **uprobes/uretprobes**（`libcudart.so`） | **tracepoints**（`gpu_scheduler` 等）, **kprobes**（NVIDIA `nvidia_*`） | **PTX/SPIR-V注入**で**CUDA/ROCmカーネル**入口/出口ほかにeBPF |
| 観測できる事象 | `cudaMalloc/Free`, `cudaMemcpy(kind)`, `cudaLaunchKernel`, stream/event操作、デバイス選択 | ジョブ実行開始/完了、依存待ち、GEM/BO割当・bind・shrink、コマンド送信、割り込み、**Xid** | **ワープ/スレッドの時刻**, メモリアクセス、分岐/同期、占有率の要因など**カーネル内挙動** |
| 粒度/可視性 | **APIレベル**の“投入側”の意図と引数、戻り値、同期点 | **スケジューラ/ドライバ**の状態遷移とボトルネック | **命令/ワープ〜スレッド**級の“実行そのもの” |
| 非同期への強さ | API投入は追えるが、**キュー滞留や実体時間は隠れやすい** | 実行開始/完了/待ちが時系列で掴め、**非同期の滞留が見える** | **すべてのフェーズをGPU内基準で刻める**（非同期でも分解可能） |
| 相関のしやすさ | プロセス/関数/引数で**アプリ意図**と結び付けやすい | **GPUがいつ走り出し/止まったか**をOS視点で時刻同期可 | **GPU内部⇄CPU/OS**を**同じBPFマップ/リング**で統合も可能 |
| ベンダー/適用範囲 | CUDAランタイム依存（主に**NVIDIA**） | **DRM系はベンダー非依存**（Intel/AMD/Nouveau）。NVIDIAは**専用kprobe + Xid TP** | **NVIDIA/AMD（ROCm）**対応設計（実験段階/一部制約） |
| アプリ改変 | **不要** | **不要**（カーネルのトレースポイント/kprobe有効化のみ） | **不要**（実行時パッチ）だがビルド/環境要件は重め |
| 代表イベント/スクリプト | `cudaMalloc/Memcpy/LaunchKernel`のENTER/EXIT | `drm_run_job / drm_sched_process_job / ...`、`nvidia_unlocked_ioctl`、`nvidia_dev_xid` | `ATTACH_CUDA_PROBE/RETPROBE`、`__memcapture` |
| 取得データ | サイズ/方向、関数エラー、PID/comm、時刻 | ジョブID/リング/キュー深さ、フェンス、BO移動、割込、Xid | グローバルタイマ、blockIdx/threadIdx、**GPU専用BPFマップ**（per-thread配列/リング） |
| オーバーヘッド | **API呼び出しあたり≈+~2µs**（計測例） | **イベント時のみ極小**（ns級。継続監視向き） | **低〜中**（設計上低オーバーヘッドだが**実験段階**） |
| 強み | アプリの“意図”と引数が分かる/導入容易 | 本番適用しやすい/ベンダー非依存（DRM）/スケジューリング可視 | **GPU内部の真因に踏み込める**/条件収集/CPU側BPFと統合 |
| 限界 | **GPU内部は不可視**/非同期だと時間分解が難しい | **カーネル外/内部は見えるが**、**カーネル“中身”は不可視** | 実験的/環境依存/安全検証・運用ノウハウがまだ少ない |
| 使いどころ | **誰が何をいつ投げたか**、API失敗/サイズ/方向、同期点の洗い出し | **実行キュー滞留/待ち/VRAM圧**、Xid障害、リング偏りの把握 | **ワープ分岐/メモリアクセス/占有率**など**根因診断**と高度な最適化 |