[Skip to content](https://eunomia.dev/en/bpftime/documents/gpu/#write-and-run-ebpf-on-gpu-with-bpftime) ## bpftime を使用して GPU 上で eBPF を記述・実行 > 鄭宇生=、于通=、楊一威= bpftimeはCUDA/ROCmアタッチメント実装を通じてGPUサポートを提供しており、これによりeBPFプログラムを **NVIDIAおよびAMD製GPUのカーネル内部で直接実行** することが可能になります。これにより、eBPFのプログラム可能性、監視機能、カスタマイズ機能をGPUコンピューティングワークロードに適用可能となり、ソースコードを変更せずにGPUアプリケーションのリアルタイムプロファイリング、デバッグ、実行時拡張を実現できます。 > **注意：** GPUサポートは現在実験段階です。ご質問やご提案がございましたら、 [問題を報告](https://github.com/eunomia-bpf/bpftime/issues) いただくか、 [お問い合わせ](https://eunomia.dev/en/bpftime/documents/gpu/) ください。 ## 課題：GPUの可観測性に関する課題 GPUは現在、機械学習、科学技術計算、高性能コンピューティングワークロードにおいて主流のアクセラレータとなっていますが、SIMT（Single Instruction, Multiple Thread）実行モデルを採用しているため、可観測性と拡張性において重大な課題が生じています。最新のGPUでは、数千ものスレッドがワーフ（通常32スレッド）に編成され、ストリーミングマルチプロセッサ（SM）上で同期的に実行されます。カーネルはホストから非同期的に起動されます。これらのスレッドは、高速だが容量制限のあるスレッドごとのレジスタから、スレッドブロック内の共有メモリ/LDS、L1/L2キャッシュ、そして低速ながら大容量のデバイスメモリへと至る複雑な多層メモリ階層を移動します。その一方で、プリエンプション機能が限られているため、カーネルの実行を中断・検査・拡張することが困難という課題を抱えています。この複雑なアーキテクチャ構造により、ワーフダイブ（分岐処理による処理分岐）、メモリのコアレッシングパターン、バンク競合、占有率の変動など、スループットに直接影響を与える多様な性能特性が生じます。しかし、これらの挙動は従来の監視ツールではほぼ把握できない状態にあります。カーネルストール、メモリボトルネック、非効率的な同期処理、あるいはSM（ストリーミング・マルチプロセッサ）の非効率的な使用といった問題を理解し最適化するためには、GPU内部の深いレベルで実行フロー、メモリアクセスパターン、ワープ間連携などの詳細な可視化が必要です。さらに、動的にカスタムロジックを注入できる能力も不可欠です。これらの機能こそが、既存のツールでは柔軟かつプログラム可能な形で提供することが困難な要素なのです。 既存のGPUトレースおよびプロファイリングツールは、大きく2つのカテゴリに分類され、それぞれに重大な制約があります。第一に、多くのトレーシングツールはCPU-GPU境界でのみ動作し、CUDA/ROCmのユーザー空間ライブラリ呼び出し（例えばlibcuda.soに対するLD\_PRELOADフックを使用）や、システムコール層でカーネルドライバをインストゥルメントする方法を採用しています。このアプローチでは、カーネルの起動、メモリ転送、APIの実行タイミングといったホスト側のイベントは捕捉できますが、GPUをブラックボックスとして扱うため、カーネル実行時の内部動作や特定のワープ動作・メモリストールとの関連性を可視化することはできません。また、デバイス内部の実行状況に応じて動作を動的に変更する機能も提供されません。第二に、GPUベンダー固有のプロファイラ（NVIDIAのCUPTI、Nsight Compute、IntelのGTPin、AMDのROCProfiler、NVBitやNeutrinoなどの研究用ツール）はデバイス側の計測機能を提供しており、ハードウェア性能カウンタやワープレベルのトレース、命令レベルのメトリクスなどを収集することが可能です。ただし、これらのツールはLinuxカーネルの監視機能や拡張機能スタックとは独立した環境で動作するため、相互連携が図られていません。これらのツールでは、GPUイベントとCPU側のeBPFプローブ（kprobes、uprobes、tracepoints）を相関分析することはできません。また、データ収集には個別のパイプラインと分析ワークフローが必要となり、多くの場合、詳細なプロファイリングでは10～100倍の処理速度低下を引き起こすという重大なオーバーヘッドが発生します。さらに、これらのツールには、再コンパイルやサービス中断なしに本番環境で収集するデータの内容や処理方法をカスタマイズできる、制御プレーンからの動的なプログラム可能性と制御機能が欠けています。これがeBPFを強力なツールにしている本質的な利点です。 ### タイムライン可視化の限界：観測可能な範囲と不可能な範囲 一般的なデバッグシナリオを考えてみましょう。「CUDAアプリケーションの実行に500msかかっていますが、その時間がどこに使われているのか分かりません。メモリ転送なのか、カーネル実行なのか、それともAPIオーバーヘッドなのか？」この回答は、アプリケーションが同期型CUDA APIを使用しているか非同期型APIを使用しているかによって大きく異なります。これは、CPU側からの観測可能性に根本的な制約があることを示しています。 #### 同期実行モード：CPU側ツールで確認できる範囲と限界 同期モードでは、CUDA API呼び出しはGPUが各処理を完了するまで待機するため、CPUとGPUのタイムラインが密接に連動します。典型的なワークフローとして、デバイスメモリの確保、ホストからGPUへのデータ転送（ホスト→デバイス）、カーネル実行、完了待機といった一連の処理が挙げられます。CPU側のプロファイラは、各ブロック型API呼び出しの実時間を測定できるため、高レベルな傾向分析に有用な情報を提供します。例えば、 `cudaMemcpy()` の実行時間が200μ秒であるのに対し、カーネル処理完了を待つ `cudaDeviceSynchronize()` の時間がわずか115μ秒しかない場合、開発者はデータ転送が計算処理よりも圧倒的に時間を消費していることを即座に把握できます。これは、PCIe帯域幅のボトルネックが原因である可能性が高く、固定メモリの使用、バッチサイズの拡大、あるいは非同期転送の採用などによって改善が見込めることを示唆しています。 ```js CPU Timeline (what traditional tools see): ─────────────────────────────────────────────────────────────────────────── cudaMalloc() cudaMemcpy() cudaLaunchKernel() cudaDeviceSync() ──────●─────────●──────────────────●──────────────────●──────────────────── ~1μs wait 200μs wait returns 115μs wait (blocked) (H→D transfer) immediately (kernel done) GPU Timeline (actual execution with hidden phases): ─────────────────────────────────────────────────────────────────────────── ◄─Alloc─►◄────H→D DMA────►◄──Launch──►◄──Kernel Exec──►◄─Cleanup─► │ ~1μs │ 200μs │ 5μs │ 100μs │ ~10μs │ ──────┴───────┴──────────────────┴──────────┴─────────────────┴─────────┴── (SM busy) (SM idle) ``` しかし、開発者が「カーネル同期に115μ秒かかっているのに、なぜカーネル処理が遅いのか？ その原因は起動オーバーヘッドなのか、メモリストールなのか、ワープ分岐なのか、またはSMの利用率が低いためなのか？」と疑問を抱いた場合、CPU側のツールでは根本的な限界に突き当たります。115μ秒という同期時間は、GPU内部で発生する複数の隠れた処理フェーズを一括した不透明な集計値です。これには、カーネル起動時のオーバーヘッド（ストリーミングマルチプロセッサ上で作業を開始するまで約5μ秒）、実際のカーネル実行時間（ストリーミングマルチプロセッサ上での計算処理約100μ秒）、および後処理時間（パイプラインのクリアとリソース解放に約10μ秒）が含まれます。これらは前述のGPUタイムライン図に示されている通りです。 同期API呼び出しのタイミングが完璧であっても、CPU側のツールでは、カーネル性能の低下が以下のどの要因によるものかを区別できません： (1) 過剰な起動オーバーヘッド（例：小規模なカーネル起動が多すぎる場合） (2) 実行時間100μs内の計算効率の悪さ（例：分岐によるワープの30%しかアクティブになっていない場合） (3) メモリアクセスパターンによる処理遅延（例：非結合型グローバルメモリロードの場合） (4) ストリーミングマルチプロセッサの未活用状態（例：利用可能なSMの50%しか稼働していない場合）これらの情報を取得するには、カーネル実行時にGPU内部からのみアクセス可能な、ワープレレベルの実行状況、メモリトランザクション統計、およびスレッドごとの動作状況を把握する必要があります。カーネルの実行中にどのような処理が行われているのかを理解するためには、起動時と終了時だけでなく、より詳細な粒度でプログラム可能なGPU内部の視点が必要です。 #### 非同期実行：時間的デカップリングが可視化機能を無効化する仕組み 最新のCUDAアプリケーションでは、CPU処理とGPU実行をオーバーラップさせることでハードウェアの利用効率を最大化するため、非同期API（ `cudaMemcpyAsync()` やストリームを使用した `cudaLaunchKernel()` など）が採用されています。これは時間的非同期処理の概念を導入するもので、API呼び出しはストリームへの作業投入後すぐに戻り、CPUは作業を継続できる一方、GPUはバックグラウンドで順次処理を実行します。これにより、同期処理モードでCPU側ツールが有していた可観測性が損なわれることになります。 同じワークフローを非同期モードで実行した場合を考えてみましょう。開発者はホストからデバイスへのデータ転送（200μ秒）、カーネル起動処理（実行時間100μ秒）、デバイスからホストへのデータ転送（150μ秒）をキューに登録し、その後CPU側の処理を継続します。最終的に `cudaStreamSynchronize()` を呼び出して、すべてのGPU操作が完了するのを待ちます。CPU側から見ると、すべてのエンキュー操作はマイクロ秒単位で即座に完了し、最終的な同期ポイントでのみ処理がブロックされ、合計455μsの時間が報告されます（GPU上で順次実行された200μsのホスト→デバイス転送、5μsのカーネル起動、100μsのデバイス→ホスト転送の合計）。 ```js CPU Timeline (what traditional tools see): ───────────────────────────────────────────────────────────────────────────────── cudaMallocAsync() cudaMemcpyAsync() cudaLaunchKernel() cudaMemcpyAsync() cudaStreamSync() ●─────●─────●─────●─────────────────────────────────────────────────────────────────●──── 1μs 1μs 1μs 1μs CPU continues doing other work... 455μs wait (alloc)(H→D)(kernel)(D→H) (all done) GPU Timeline (actual execution - sequential in stream): ───────────────────────────────────────────────────────────────────────────────── ◄─Alloc─►◄───────H→D DMA────────►◄─Launch─►◄────Kernel Exec────►◄────D→H DMA────► │ ~1μs │ 200μs │ 5μs │ 100μs │ 150μs │ ┴────────┴────────────────────────┴─────────┴────────────────────┴────────────────┴───── ↑ ↑ ↑ CPU already moved on GPU still working Sync returns ``` 同期実行モードでは、個々のAPI呼び出しの所要時間を測定することで、データ転送と計算処理のどちらがボトルネックになっているかを特定できました。非同期モードでは、この機能が完全に失われてしまいます。すべてのタイミング情報が同期ポイントで単一の455μsの合計値に集約されてしまうためです。CPU側からでは、「ボトルネックがメモリ転送なのかカーネル実行なのか」という疑問に答えることができなくなります。もし最初のデータ転送がアンピンドメモリの影響で通常の2倍の時間（400μs）かかった場合、すべての後続操作も200μs遅延することになります。開発者が認識できるのは、全体の処理時間が455μsから655μsに増加したという事実だけで、どの操作が遅延を引き起こしたのか、いつ発生したのか、あるいはその影響が後続操作に伝播したのかどうかといった詳細な情報は一切得られません。 非同期実行では、同期モードでは見えなかったGPU内部の詳細（ワープ分岐、メモリストール、SM使用率など）が隠蔽されるだけでなく、CPUツールが従来提供していた粗粒度のフェーズタイミング情報も失われてしまいます。開発者は基本的なトリアージすら行えなくなるという問題が生じます。メモリ転送処理、カーネルロジック、API使用パターンのいずれに最適化の重点を置くべきかを判断できません。デバッグのために低速な同期実行に戻すか（非同期実行の利点を損なう）、あるいは再コンパイルが必要な手動計測ポイントを追加するかの二者択一を強いられるからです。 LLMサービングなどの最新GPUアプリケーションでは、高度な最適化技術がさらに状況を複雑にしています。バッチ処理戦略は複数の操作をパイプラインのように組み合わせてスループットを最大化しますが、その一方で、どの個別の操作がボトルネックになっているかを特定することを困難にします。永続カーネルはGPU上に常駐し、複数のワークバッチを処理するため、起動オーバーヘッドは発生しませんが、処理フェーズの境界が不明瞭になるという問題が生じます。複数ストリーム間の複雑な依存関係を伴うマルチストリーム実行では、各ストリームの処理が予測不能に交錯する複雑な実行グラフが形成されます。スレッドブロックごとの共有メモリ使用量は占有率を制限し、同時に実行可能なワープ数を制限するため、カーネル構成によって微妙なリソース競合が発生します。これらの最適化はスループットを大幅に向上させますが、もともとCPU側からの観測・デバッグが困難だった非同期実行モデルをさらに不透明なものにしてしまいます。 ただし、GPU内部の動作を完全に可視化できたとしても、メモリストールサイクルが高い原因を特定することはできません。これはワープダイバージェンスによる非結合アクセスが原因なのか、それともホストスレッドの再スケジュール遅延が非同期メモリコピーを妨げているのか、あるいは複数のRDMA操作が同時に実行されることでPCIe帯域が混雑しているためなのか、といった問題です。あるいは、カーネルによるオンデマンドページ移動の遅延が原因でしょうか？同様に、SMの利用率が低い場合でも、デバイス単体のメトリクスでは、グリッドサイズが小さすぎるのか、ユーザー空間のミューテックスによって起動処理がシリアル化されているのか、ECCエラーや電源イベント後にドライバが制限をかけているのかを区別することができません。 特に本番環境では、テールレイテンシが突発的に急上昇するケースが問題となります。これはGPUキャッシュの影響なのか、ホスト側のプロデューサースレッドに対するcgroupによるスロットリングなのか、それとも他のコンテナが大規模なDMA転送を実行しているためなのか――こうした原因の特定が困難です。デバイス専用のツール群は「GPU上で何が起こったか」は報告できますが、「なぜこの異種混合システムで今このタイミングでそれが起こったのか」までは説明できません。ユーザー空間の動作（スレッド、システムコール、メモリ割り当て）との時間同期された相関関係（ページフォルト、スケジューラのコンテキストスイッチ、ブロックI/O）、ドライバの判断（ストリーム依存関係の解決、メモリ登録、電源状態遷移）がなければ、エンジニアは以下のような非効率な試行錯誤を繰り返す必要があります：GPUの症状を観察 → ホスト側の原因を推測 → その場しのぎの計測機能を組み込んで再ビルド → このプロセスを繰り返す。このフィードバックループはコストが高く、遅延に敏感な環境では実現が困難な場合が多いです。 > **重要な洞察：** GPUの効果的な監視性と拡張性を実現するためには、ヘテロジニアスコンピューティングスタックの複数層にまたがる統合ソリューションが必要です。具体的には、CUDA APIを呼び出すユーザー空間アプリケーションから、デバイスリソースを管理するOSカーネルドライバ、そしてGPUハードウェア上で実行されるデバイスコードに至るまで、一貫したアプローチが求められます。従来のツールはこれらの各層に分散しており、CPU-GPU境界やGPUカーネル内部といった個別のレベルでの可視性は提供できるものの、あるレベルでの判断やイベントが別のレベルでのパフォーマンスや動作にどのように影響するかを理解するために必要な層間の相関関係が欠けています。 #### 現行のツール：Nsight Systems/Computeとその限界 NVIDIAのNsightスイートは、異なる領域にまたがる可視性の課題を解決しようとしています。 **Nsight Systems** は、CPUスレッド、CUDA API呼び出し、カーネル起動、メモリ転送などを統合的に表示するシステム全体のタイムラインを提供します。 **Nsight Compute** は、CUPTIカウンターとリプレイベースの分析を通じて、カーネルレベルの詳細なマイクロアーキテクチャメトリクスを提供するとともに、最適化のためのガイド付きルールも提供します。これらのツールは、カーネル起動とCPUスレッドスケジューリングの相関関係を可視化できるほか、各カーネルのストール要因やメモリ使用状況に関する詳細な情報を提供します。 ただし、Nsightやその他ベンダー提供ツールは、統一的なeBPFアプローチと比較して根本的な制約を抱えています。第一に、Nsightは固定されたイベントセットとアタッチメントポイントでの任意のプログラム可能なロジックを持たない閉鎖的なイベントモデルを採用しています。つまり、アプリケーションを再コンパイルせずにカスタム計測機能を動的に追加したり、「カーネル実行時間が100msを超えた場合のみデータを収集する」といったフィルタリング条件を定義することはできません。第二に、これらのツールはカウンター多重化やリプレイ機構によってワークロードの動作を意図的に変更する特別なプロファイリングセッションを必要とします。このため、本番環境で常時稼働する継続的なテレメトリ収集には適しておらず、リプレイベースの収集方式では一時的な異常値や稀なイベントを捕捉できないという問題があります。第三に、Nsightにはインシトゥフィルタリング機能がないため、すべての生データをエクスポート後に後処理する必要があり、大規模な非同期パイプラインからは状態観測に基づいてサンプリングロジックを動的に変更できないにもかかわらず、数GB規模のトレースデータが生成されてしまいます。第四に、システム統合の機能が著しく制限されています。Nsightでは、再起動なしで永続的なカーネルに動的プローブをアタッチすることができず、Linux eBPFインフラストラクチャ（kprobes/uprobes/tracepoints）との統合機能も欠いています。さらに、CPUとGPUの両方で計測を行う場合、データ構造（マップ）を共有できないため、ページフォルト（ホスト側）→遅延起動キューへの追加→ワープストールスパイクといった因果関係の連鎖を追跡することが極めて困難です。これらのツールはNVIDIA製品専用であり、AMDやIntel製、あるいはヘテロジニアスシステムにおける他のアクセラレータ製品へのベンダー中立的な展開方法が明確ではありません。実際には、開発者は大規模なトレースデータを扱う場合、根本原因の特定に時間がかかり、プロファイリング時のオーバーヘッドでは再現しない本番環境の問題を見逃してしまうことがあります。また、GPUイベントを既存の本番環境監視スタック（perf、bpftrace、カスタムeBPFエージェントなど）と相関させるには、複雑なモード切り替えを行って特別な「プロファイリングセッション」を起動する必要があり、非効率な作業を強いられています。 #### ギャップを埋める：eBPFによるCPUとGPUの統合的な監視・拡張性 bpftimeはeBPFプログラムをGPUカーネル内で直接実行することで、スタック全体にわたるプログラム可能な統合型の監視機能と拡張性を実現します。セッションベースのプロファイラとは異なり、常に稼働するプロダクション環境向けモニタリングを実現します。プローブの動的ロード/アンロードやデバイス側の条件フィルタリング機能により、オーバーヘッドを最小限に抑えられます。非同期モードの可観測性を回復し、各処理フェーズごとにタイムスタンプを付与します（CPU→GPU：T+200μs、カーネル処理：T+205μs、GPU→CPU：T+455μs）。さらに、ナノ秒単位の粒度でワープ実行状況やメモリアクセスパターンなどのGPU内部情報を可視化し、従来の個別プロファイラのような重いオーバーヘッドなしにCPUとGPUのイベントを相関分析できます。 このアーキテクチャでは、CPUとGPUのプローブを統一された制御プレーンにおける対等な要素として扱います。共有されるBPFマップとリングバッファによって直接的なデータ交換が可能となり、動的な計測機能は再コンパイルやシステム再起動を必要としません。さらに、既存のeBPFインフラストラクチャ（perf、bpftrace、カスタムエージェントなど）との連携においても、モード切り替えなしでシームレスに統合できます。開発者は、uprobesを使用してCPU側のCUDA API呼び出しを、kprobesを使用してカーネルドライバとの相互作用を、そしてCUDAプローブを使用してGPU側のカーネル実行を、すべて同じeBPFツールチェーンで同時にトレースできます。さらに、ホスト-デバイス間の境界を越えてイベントを相関させることも可能です。これまで答えられなかった疑問にも回答可能になります：「T+50μ秒時点のCPUシステムコール遅延が、T+150μ秒時点のGPUカーネルストールの原因だったのか？」「高いワープ分岐率を示すカーネルを起動しているCPUスレッドはどれか？」といった疑問です。このレイヤー横断的な可視性により、ユーザー空間アプリケーションロジックからカーネルドライバ、GPUハードウェア動作に至るまで、本番環境の監視ワークフローを離れることなく、システム全体にわたる根本原因分析が可能になります。 ## 解決策：bpftimeによるGPU上でのeBPFの実現 **bpftimeのアプローチ** は、このギャップを埋めるため、eBPFのプログラム可能性とカスタマイズモデルをGPU実行環境に直接拡張することで、eBPFプログラムをアプリケーションワークロードと並行してGPUカーネル内でネイティブに実行可能にします。本システムでは、CPU側のkprobes/uprobesと同等の柔軟性を備えた、GPU側向けの包括的なアタッチポイントセットを定義しています。開発者は、CUDA/ROCmデバイス関数のエントリーポイントとエグジットポイント（関数プローブに相当）、スレッドブロックのライフサイクルイベント（ブロック開始/終了）、同期プリミティブ（バリア、アトミック操作）、メモリ操作（ロード/ストア/転送）、およびストリーム/イベント操作を計測対象として設定できます。制限付きC言語で記述されたeBPFプログラムはLLVMを介してコンパイルされ、デバイスネイティブなバイトコード形式（NVIDIA GPUの場合はPTX（Parallel Thread Executionアセンブリ）、AMD/Intelの場合はSPIR-V）に変換されます。これらのプログラムは、ソースコードの修正や再コンパイルを必要とせず、バイナリインストゥルメンテーションによって実行時に対象カーネルに動的に注入されます。本ランタイムは、GPU上で完全なeBPF実行環境を提供します。その主な機能は以下の通りです： （1）SIMT環境における実行範囲の制限とメモリ安全性を保証する安全検証機構 （2）スレッド/ブロック/グリッドコンテキストへのアクセス、タイミング計測、同期処理、フォーマット出力などのためのGPU対応ヘルパー関数群 （3）高スループットなスレッド単位データ収集用のGPUメモリ常駐型BPFマップタイプ（GPU配列マップ）と、イベントストリーミング用のGPUリングバッファマップ （4）共有メモリとスピンロックを利用したホスト-GPU間通信プロトコルにより、必要に応じてホスト側ヘルパー関数を安全に呼び出すことが可能このアーキテクチャは、ナノ秒単位の精度でワープごとのタイミング情報、メモリアクセスパターン、制御フロー分岐などの詳細なテレメトリを収集できるだけでなく、実行時の状況に応じてカーネルの動作を動的に調整したり、カスタム拡張機能や最適化を実装したり、GPUの監視機能を既存のCPU側eBPFプログラムと統合して単一の分析パイプラインを構築したりすることが可能です。しかも、これらの機能を本番環境でも許容可能なオーバーヘッドで実現しています。これにより以下の機能が可能になります： - **計測ツールNVBitなどと比較して3～10倍高速なパフォーマンス** を実現 - **ベンダー非依存の設計** により、NVIDIA、AMD、Intel製GPUで共通して使用可能 - **統合された監視機能と制御機能** をLinuxカーネルのeBPFプログラム（kprobes、uprobes）で実現 - **ワープ単位または命令レベルでの詳細なプロファイリングと実行時カスタマイズ** - **GPUカーネルメモリの適応的最適化** とSM間のプログラム可能なスケジューリング - **動的拡張機能** により、再コンパイルなしでGPUワークロードを拡張可能 - **GPUの計算能力を活用したeBPFアプリケーションの高速化** 本アーキテクチャは4つの主要な目標を達成するように設計されています：(1) ユーザー空間、カーネル、異なるベンダーの複数CPU/GPUコンテキスト間でシームレスに動作する統一的なeBPFベースのインターフェースを提供すること、(2) ソースコードの修正や再コンパイルを必要とせずに動的な実行時計測を可能にすること、(3) GPUハードウェアの制約とSIMT実行モデルの範囲内で安全かつ効率的な実行を実現することです。（4）依存関係が少なく導入が容易：既存のCUDA/ROCm/OpenGLランタイム上に構築されており、カスタムカーネルドライバやファームウェアの変更、記録・再生システムのような大規模なランタイムを必要としません。 ## アーキテクチャ ### CUDA アタッチメント・パイプライン GPUサポートは `nv_attach_impl` システム（ `attach/nv_attach_impl/` ディレクトリ）上に構築されており、以下の計測パイプラインを実装しています： ## 使用例 以下に、bccスタイルのツールを使用したGPU向けeBPFの機能を、いくつかの具体例を通じてご紹介します： CUDAカーネルの終了ポイントにアタッチし、各GPUスレッドが実行を完了した正確なナノ秒単位のタイムスタンプを記録します。これにより、従来のプロファイラでは検出できないスレッドの分岐、メモリアクセスパターン、ワープスケジューリングに関する問題を明らかにできます。 > ここで使用されている `kprobe` / `kretprobe` という名称は適切ではありません。単なる仮の名称として使用しているだけです。この部分は適切な名称に変更すべきでしょう。 **使用例**: カーネルの動作が想定より遅いことに気づいた場合、 `kernelretsnoop` を使用すると、各ワープ内のスレッド31が他のスレッド0～30よりも750ナノ秒遅れて実行を終了していることが判明します。これは分岐条件によるスレッドの分岐現象を示しています。この問題を解決するために分岐処理をリファクタリングした結果、すべてのスレッドがほぼナノ秒単位で同時に処理を完了するようになりました。 ```js // eBPF program runs on GPU at kernel exit SEC("kretprobe/_Z9vectorAddPKfS0_Pf") int ret__cuda() { u64 tid_x, tid_y, tid_z; bpf_get_thread_idx(&tid_x, &tid_y, &tid_z); // Which thread am I? u64 ts = bpf_get_globaltimer(); // When did I finish? // Write to ringbuffer for userspace analysis bpf_perf_event_output(ctx, &events, 0, &data, sizeof(data)); } ``` ### threadhist - スレッド実行回数ヒストグラム GPU配列マップを使用して、各スレッドが実行される回数をカウントします。一部のスレッドが他のスレッドよりも大幅に多くの処理を行っているワークロードの不均衡を検出し、GPUの計算リソースが無駄になっている状況を明らかにします。 **使用例**: グリッドストライドループで100万要素を5つのスレッドで処理する場合を考えます。均等な負荷分散を期待しますが、 `threadhist` の結果ではスレッド4の実行頻度がスレッド0～3の75%しかありません。境界要素の処理が不均等に分割されたため、スレッド4が待機状態となり、他のスレッドだけが作業を行っている状況です。分布を調整することで、実行負荷を均等化できます。 ```js // eBPF program runs on GPU at kernel exit SEC("kretprobe/_Z9vectorAddPKfS0_Pf") int ret__cuda() { u64 tid_x, tid_y, tid_z; bpf_get_thread_idx(&tid_x, &tid_y, &tid_z); // Per-thread counter in GPU array map u64 *count = bpf_map_lookup_elem(&thread_counts, &tid_x); if (count) { __atomic_add_fetch(count, 1, __ATOMIC_SEQ_CST); // Thread N executed once more } } ``` CPU上の `cudaLaunchKernel()` 呼び出しからGPU上での実際のカーネル実行までの時間を測定します。これにより、高速に動作するはずのカーネルが実運用環境で遅くなる原因となる、隠れたキュー遅延やストリーム依存関係、スケジューリングオーバーヘッドを明らかにします。 **使用例**: カーネルの実行時間は100μsですが、ユーザーからは50msの遅延が報告されています。 `launchlate` を使用すると、各カーネルの起動遅延が200～500μsであることが判明します。これは各カーネルが前のカーネルの実行完了とメモリ転送の終了を待つためです。実際の合計時間は5msであり、1msではありません。CUDAグラフモードに切り替え、すべてのカーネル起動をバッチ処理すると、遅延が1.2msに短縮されます。 ```js BPF_MAP_DEF(BPF_MAP_TYPE_ARRAY, launch_time); // CPU-side uprobe captures launch time SEC("uprobe/app:cudaLaunchKernel") int uprobe_launch(struct pt_regs *ctx) { u64 ts_cpu = bpf_ktime_get_ns(); // When did CPU request launch? bpf_map_update_elem(&launch_time, &key, &ts_cpu, BPF_ANY); } // GPU-side kprobe captures execution start SEC("kprobe/_Z9vectorAddPKfS0_Pf") int kprobe_exec() { u64 ts_gpu = bpf_get_globaltimer(); // When did GPU actually start? u64 *ts_cpu = bpf_map_lookup_elem(&launch_time, &key); u64 latency = ts_gpu - *ts_cpu; // How long did kernel wait in queue? u32 bin = get_hist_bin(latency); // Update histogram... } ``` ### その他の使用例 - **[cuda-counter](https://github.com/eunomia-bpf/bpftime/tree/master/example/gpu/cuda-counter)**: 基本的なプローブ/リターンプローブ機能とタイミング測定 - **[mem\_trace](https://github.com/eunomia-bpf/bpftime/tree/master/example/gpu/mem_trace)**: メモリアクセスパターンの追跡と解析 - **[directly\_run\_on\_gpu](https://github.com/eunomia-bpf/bpftime/tree/master/example/gpu/directly_run_on_gpu)**: カーネルにアタッチすることなく、直接GPU上でeBPFプログラムを実行します - **[rocm-counter](https://github.com/eunomia-bpf/bpftime/tree/master/example/gpu/rocm-counter)**: AMD ROCm GPU向けの計測機能（実験段階） ### 主要構成要素 1. **CUDAランタイムフック**: Fridaベースの動的計測技術を用いてCUDA API呼び出しをインターセプトします 2. **PTX 改変処理**: eBPF バイトコードを PTX（並列スレッド実行）アセンブリに変換し、GPU カーネルに注入します 3. **ヘルパー・トランポリン**: マップ操作、タイミング計測、コンテキストアクセス用のGPUから利用可能な補助関数を提供します 4. **ホスト-GPU間通信**: ピン留めされた共有メモリを介してGPUからホストへの同期呼び出しを可能にします ### アタッチタイプ bpftimeではGPUカーネルに対して以下の3種類のアタッチ方式をサポートしています（ `attach/nv_attach_impl/nv_attach_impl.hpp:33-34` で定義されています）： - **`ATTACH_CUDA_PROBE` (8)** - カーネルエントリポイントでeBPFコードを実行します - **`ATTACH_CUDA_RETPROBE` (9)** - カーネル終了時にeBPFコードを実行します - **メモリアクセスプロービング（ `__memcapture` ）** - メモリアクセスパターンをキャプチャするための特殊なプロービングタイプ すべてのタイプで、ターゲットとするカーネル関数を名前で指定できます（例：C++の名前修飾後の関数名 `_Z9vectorAddPKfS0_Pf` など）。 ## GPU専用BPFマップ bpftimeには、GPU操作向けに最適化された専用のマップタイプが含まれています： ### BPF\_MAP\_TYPE\_PERGPUTD\_ARRAY\_MAP（1502） **スレッドごとに独立したストレージ** を備えたGPU常駐型配列マップで、高性能なデータ収集を実現します。 主な特徴： - データはGPUメモリ（CUDA IPC共有メモリ）に直接格納されます - 各スレッドごとに独立したストレージ領域が割り当てられます（ `max_entries × max_thread_count × value_size` ） - GPUからのゼロコピーアクセスが可能で、データをホスト側にDMA転送できます - GPUコード内で `bpf_map_lookup_elem()` および `bpf_map_update_elem()` 関数を使用できます 実装箇所: `runtime/src/bpf_map/gpu/nv_gpu_array_map.cpp:14-81` ### BPF\_MAP\_TYPE\_GPU\_RINGBUF\_MAP（1527） **スレッドごとのイベントストリーミング** を効率的にホスト側に転送するためのGPUリングバッファマップ。 主な特徴： - GPUメモリ上でスレッドごとにロックフリーなリングバッファを実現 - メタデータ付きの可変サイズイベントレコードをサポート - 低オーバーヘッドで非同期的なデータ収集が可能 - `bpf_perf_event_output()` ヘルパー関数と互換性あり 実装コード: `runtime/src/bpf_map/gpu/nv_gpu_ringbuf_map.cpp` ## GPU用ヘルパー関数 bpftime では、CUDA カーネルから利用可能な GPU 専用の eBPF ヘルパー関数を提供しています（ `attach/nv_attach_impl/trampoline/default_trampoline.cu:331-390` 参照）： ### コアGPUヘルパー関数 | ヘルパーID | 関数シグネチャ | 説明 | | --- | --- | --- | | **501** | `ebpf_puts(const char *str) ` | GPUカーネルからホストコンソールへ文字列を出力します | | **502** | `bpf_get_globaltimer(void) ` | GPUのグローバルタイマー値を読み取る（ナノ秒単位の精度） | | **503** | `bpf_get_block_idx(u64 *x, u64 *y, u64 *z) ` | CUDAブロックインデックスを取得します（blockIdx） | | **504** | `bpf_get_block_dim(u64 *x, u64 *y, u64 *z) ` | CUDAブロックの次元情報を取得します（blockDim） | | **505** | `bpf_get_thread_idx(u64 *x, u64 *y, u64 *z) ` | CUDAスレッドインデックスを取得します（threadIdx） | | **506** | `bpf_gpu_membar(void) ` | GPUメモリバリアを実行します（ `membar.sys` ） | ### 標準 BPF ヘルパー関数（GPU 対応版） 以下の標準的なeBPFヘルパー関数は、GPU上で動作するように特別な最適化が施されています： - **`bpf_map_lookup_elem()`** (1): GPU配列マップ用の高速パス。その他のマップタイプの場合はホスト側で処理されます - **`bpf_map_update_elem()`** (2): GPU配列マップ用の高速パス。その他のマップタイプの場合はホスト側で処理されます - **`bpf_map_delete_elem()`** (3): 共有メモリを介したホスト側呼び出し - **`bpf_trace_printk()`** (6): ホスト側コンソールへのフォーマット出力 - **`bpf_get_current_pid_tgid()`** (14): ホストプロセスのPID/TIDを取得します - **`bpf_perf_event_output()`** (25): GPUリングバッファマップ向けに最適化された機能 ### ホスト-GPU間通信プロトコル ホストとの相互作用を必要とするヘルパー関数の場合、bpftimeはスピンロックとワープレベルのシリアル化を採用した共有メモリプロトコルを採用しており、これにより処理の正確性が保証されます。このプロトコルの具体的な手順は以下の通りです： 1. GPUスレッドがスピンロックを取得します 2. リクエストパラメータを共有メモリに書き込みます 3. フラグを設定し、ホストからの応答を待機します 4. ホスト側でリクエストを処理し、処理完了を通知します 5. GPUが応答を読み取り、ロックを解放します ## GPUサポートを有効にしたビルド方法 ### 必要条件 - **NVIDIA CUDA Toolkit** （バージョン12.x推奨）または **AMD ROCm** - **CMake** 3.15 以降 - **LLVM** 15以降（PTX生成用） - 実行時フック用の **Frida-gum** ### ビルド構成 ```js # For NVIDIA CUDA cmake -Bbuild \ -DBPFTIME_ENABLE_CUDA_ATTACH=1 \ -DBPFTIME_CUDA_ROOT=/usr/local/cuda-12.6 \ -DCMAKE_BUILD_TYPE=Release make -j$(nproc) ``` ## 参考文献 1. [bpftime OSDI '25 論文](https://www.usenix.org/conference/osdi25/presentation/zheng-yusheng) 2. [CUDAランタイムAPI](https://docs.nvidia.com/cuda/cuda-runtime-api/) 3. [PTX命令セットアーキテクチャ（ISA）](https://docs.nvidia.com/cuda/parallel-thread-execution/) 4. [eBPF ドキュメント](https://ebpf.io/) 5. [eGPU：GPUへのeBPFプログラム可能性と可観測性の拡張](https://dl.acm.org/doi/10.1145/3723851.3726984) 引用文献: ```js @inproceedings{yang2025egpu, title={eGPU: Extending eBPF Programmability and Observability to GPUs}, author={Yang, Yiwei and Yu, Tong and Zheng, Yusheng and Quinn, Andrew}, booktitle={Proceedings of the 4th Workshop on Heterogeneous Composable and Disaggregated Systems}, pages={73--79}, year={2025} } ``` ご質問やご意見がございましたら、 [GitHub](https://github.com/eunomia-bpf/bpftime) で問題を報告するか、 [お問い合わせ](https://eunomia.dev/en/bpftime/documents/gpu/) ください。 [共有先:](https://x.com/intent/tweet?text=Write%20and%20Run%20eBPF%20on%20GPU%20with%20bpftime%0A&url=https://eunomia.dev/en/bpftime/documents/gpu/)[共有先:](https://www.facebook.com/sharer/sharer.php?u=https://eunomia.dev/en/bpftime/documents/gpu/)