[Skip to content](https://eunomia.dev/tutorials/47-cuda-events/#ebpf-tutorial-tracing-cuda-gpu-operations)
## eBPFチュートリアル：CUDA GPU操作のトレース
CUDAアプリケーションが実行されているとき、その裏側で何が起こっているのか疑問に思ったことはありませんか？GPUの操作はデバッグやプロファイリングが難しい場合があります。なぜなら、GPUは独自のメモリ空間を持つ別のデバイス上で動作するからです。本チュートリアルでは、CUDA API呼び出しをリアルタイムで監視できる強力なeBPFベースのトレーシングツールの構築方法をご説明します。
> 完全なソースコードはこちら： [https://github.com/eunomia-bpf/bpf-developer-tutorial/tree/main/src/47-cuda-events](https://github.com/eunomia-bpf/bpf-developer-tutorial/tree/main/src/47-cuda-events)
## CUDAとGPUトレーシングの入門
CUDA（Compute Unified Device Architecture）は、NVIDIAが開発した並列コンピューティングプラットフォームおよびプログラミングモデルであり、開発者がNVIDIA GPUを汎用計算に活用することを可能にします。CUDAアプリケーションを実行する際の一般的な処理フローは、まずホスト（CPU）がデバイス（GPU）上でメモリを割り当て、次にホストメモリからデバイスメモリへデータを転送します。その後、GPUカーネル（関数）を起動してデータを処理し、処理結果をデバイスからホストへ転送した後、最終的にデバイスメモリを解放します。
この処理プロセスの各段階では、 `cudaMalloc` によるメモリ割り当て、 `cudaMemcpy` によるデータ転送、 `cudaLaunchKernel` によるカーネル実行など、CUDA APIが呼び出されます。これらのAPI呼び出しを追跡することで、デバッグやパフォーマンス最適化に有用な情報を得られますが、これは容易な作業ではありません。GPU操作は非同期処理であるため、CPUはGPUに処理を送信した後でも待機することなく実行を継続できます。このため、従来のデバッグツールではこの非同期境界を越えてGPU内部の状態にアクセスすることが困難でした。
ここでeBPFが活躍します！uprobesを使用することで、GPUドライバに到達する前にユーザー空間のCUDAランタイムライブラリ（ `libcudart.so` ）内で実行されるCUDA APIコールをインターセプトし、重要な情報を捕捉することが可能になります。この手法により、メモリ割り当てのサイズとパターン、データ転送の方向と量、カーネル起動時のパラメータ、APIが返すエラーコードと失敗理由、および各操作の正確なタイミング情報など、詳細なインサイトを得ることができます。CPU側でこれらの呼び出しをインターセプトすることで、アプリケーションコードを変更したり専用のプロファイリングツールに依存することなく、アプリケーションのGPU使用状況を包括的に把握することが可能になります。
本チュートリアルでは主にCPU側におけるCUDA APIのトレース手法について解説します。これはアプリケーションがGPUとどのように相互作用するかを把握するマクロレベルの視点を提供します。ただし、CPU側のトレースだけでは明らかな制約があります。 `cudaLaunchKernel` のようなCUDA API関数が呼び出されると、単にGPUに対して作業要求を送信するだけだからです。カーネルの起動タイミングは把握できても、GPU内部で実際に何が起きているのかは観察できません。数千ものスレッドがどのようにメモリにアクセスしているか、実行パターンや分岐処理、同期操作などの詳細な動作状況は依然として不可視のままです。これらの詳細は、パフォーマンスのボトルネックを理解する上で極めて重要です。例えば、メモリアクセスパターンがコアレッシングアクセスの失敗を引き起こしているか、あるいはスレッドの分岐が過度に発生して実行効率が低下しているか、といった点を把握できます。
GPU操作を詳細にトレースするためには、eBPFプログラムをGPU上で直接実行させる必要があります。これはまさに『eGPU』論文と [bpftime GPUサンプルコード](https://github.com/eunomia-bpf/bpftime/tree/master/example/gpu) が探求している内容です。bpftimeはeBPFプログラムをGPUが実行可能なPTX命令に変換し、CUDAバイナリを動的に修正することで、カーネルのエントリーポイントとエグジットポイントにこれらのeBPFプログラムを挿入します。これにより、GPU内部の動作を詳細に観察することが可能になります。この手法により、開発者はブロックインデックスやスレッドインデックス、グローバルタイマーなどのGPU固有の情報にアクセスできるほか、カーネル実行時のクリティカルパス上で測定やトレースを実行することが可能になります。GPU内部のこのような可観測性は、複雑なパフォーマンス問題の診断、カーネル実行動作の理解、GPU計算の最適化において不可欠な機能であり、CPU側のトレース機能では実現できない能力です。
## 監視対象の主要CUDA関数
当社のトラッカーは、GPUコンピューティングにおける主要な操作を表すいくつかの重要なCUDA関数を監視しています。これらの関数を理解することで、トレーシング結果の解釈や、CUDAアプリケーションの問題診断が可能になります：
### メモリ管理
- **`cudaMalloc`**: GPUデバイス上にメモリを割り当てます。この関数をトレースすることで、どのタイミングでどれだけのメモリが要求され、正常に割り当てられたかを確認できます。メモリ割り当ての失敗は、CUDAアプリケーションで頻繁に発生する問題の原因となります。
```js
cudaError_t cudaMalloc(void** devPtr, size_t size);
```
- **`cudaFree`**: GPU上で以前に割り当てられたメモリを解放します。このトレース機能により、メモリリーク（解放されずに残っているメモリ）や二重解放エラーを特定できます。
```js
cudaError_t cudaFree(void* devPtr);
```
### データ転送
- **`cudaMemcpy`**: ホスト（CPU）メモリとデバイス（GPU）メモリ間、あるいはデバイスメモリ内の異なる領域間でデータをコピーします。方向パラメータ（ `kind` ）によって、データがGPUに向かって移動しているのか、GPUから移動しているのか、またはGPUメモリ内で移動しているのかが判別できます。
```js
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);
```
`kind` パラメータには以下の値が設定可能です: - `cudaMemcpyHostToDevice` （1）: CPUからGPUへのデータコピー - `cudaMemcpyDeviceToHost` （2）: GPUからCPUへのデータコピー - `cudaMemcpyDeviceToDevice` （3）: GPUメモリ内でのデータコピー
### カーネル実行
- **`cudaLaunchKernel`**: GPU上で実行するためのカーネル（関数）を起動します。ここで実際の並列計算が実行されます。この処理をトレースすることで、カーネルがいつ起動され、正常に実行されたかどうかを確認できます。
### ストリームと同期処理
CUDAでは、並行処理と非同期操作を管理するためにストリームが使用されます：
- **`cudaStreamCreate`**: 操作を順番に、かつ他のストリームと並行して実行可能な新しいストリームを作成します。
```js
cudaError_t cudaStreamCreate(cudaStream_t* pStream);
```
- **`cudaStreamSynchronize`**: ストリーム内のすべての操作が完了するまで待機します。これはパフォーマンスのボトルネックを特定する上で重要な同期ポイントとなります。
```js
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
```
### イベント
CUDAイベントはタイミング計測と同期処理に使用されます：
- **`cudaEventCreate`**: タイミング計測用のイベントオブジェクトを作成します。
```js
cudaError_t cudaEventCreate(cudaEvent_t* event);
```
- **`cudaEventRecord`**: ストリーム内にイベントを記録します。これはタイミング計測や同期処理に利用できます。
```js
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream);
```
- **`cudaEventSynchronize`**: イベントの完了を待機します。これは別の同期ポイントとして機能します。
```js
cudaError_t cudaEventSynchronize(cudaEvent_t event);
```
### デバイス管理
- **`cudaGetDevice`**: 現在使用されているデバイスを取得します。
```js
cudaError_t cudaGetDevice(int* device);
```
- **`cudaSetDevice`**: GPU実行に使用するデバイスを指定します。
```js
cudaError_t cudaSetDevice(int device);
```
これらの関数をトレースすることで、デバイス選択やメモリ割り当てからデータ転送、カーネル実行、同期処理に至るまで、GPU操作のライフサイクル全体を包括的に把握することが可能になります。これにより、ボトルネックの特定、エラーの診断、CUDAアプリケーションの動作解析が可能になります。
## アーキテクチャ概要
当社のCUDAイベントトラッカーは、主に以下の3つのコンポーネントで構成されています：
1. **ヘッダーファイル (`cuda_events.h`)**: カーネル空間とユーザー空間間の通信用データ構造を定義します
2. **eBPFプログラム（ `cuda_events.bpf.c` ）**: uprobesを使用してカーネル側でCUDA関数のフックを実装します
3. **ユーザー空間アプリケーション（ `cuda_events.c` ）**: eBPFプログラムをロードし、イベントを処理してユーザーに表示します
本ツールはeBPF uprobesを利用して、CUDAランタイムライブラリ内のCUDA API関数にアタッチします。CUDA関数が呼び出されると、eBPFプログラムが引数と戻り値を取得し、リングバッファを介してユーザー空間に送信します。
## 主要なデータ構造
本トレーサーの中核となるデータ構造は、 `cuda_events.h` で定義されている `struct event` です：
```js
struct event {
/* Common fields */
int pid; /* Process ID */
char comm[TASK_COMM_LEN]; /* Process name */
enum cuda_event_type type;/* Type of CUDA event */
/* Event-specific data (union to save space) */
union {
struct { size_t size; } mem; /* For malloc/memcpy */
struct { void *ptr; } free_data; /* For free */
struct { size_t size; int kind; } memcpy_data; /* For memcpy */
struct { void *func; } launch; /* For kernel launch */
struct { int device; } device; /* For device operations */
struct { void *handle; } handle; /* For stream/event operations */
};
bool is_return; /* True if this is from a return probe */
int ret_val; /* Return value (for return probes) */
char details[MAX_DETAILS_LEN]; /* Additional details as string */
};
```
この構造体は、さまざまな種類のCUDA操作に関する情報を効率的に取得するために設計されています。 `union` は巧妙なメモリ節約手法であり、各イベントは同時に1種類のデータのみを必要とするためです。例えば、メモリ割り当てイベントでは割り当てサイズを、解放イベントではポインタ値をそれぞれ保持する必要があります。
`cuda_event_type` 列挙型は、CUDA 操作を分類するために使用されます：
この列挙型は、メモリ管理からカーネル起動、同期処理に至るまで、トレース対象とする主要なCUDA操作を網羅しています。
## eBPFプログラムの実装
CUDA 関数にフックする eBPF プログラム（ `cuda_events.bpf.c` ）について詳しく見ていきましょう。完全なコードはリポジトリで公開されていますが、ここでは主要な部分について解説します：
まず、ユーザー空間と通信するためのリングバッファを作成します：
```js
struct {
__uint(type, BPF_MAP_TYPE_RINGBUF);
__uint(max_entries, 256 * 1024);
} rb SEC(".maps");
```
リングバッファはトレーサーシステムにおいて重要なコンポーネントです。eBPFプログラムがイベントを送信し、ユーザー空間アプリケーションがそれらを取得するための高性能なキューとして機能します。データの損失を防ぐため、バースト的なイベント処理にも対応できるよう256KBという十分なサイズを設定しています。
各CUDA操作に対して、関連するデータを収集するための補助関数を実装しています。例として `submit_malloc_event` 関数を見てみましょう：
```js
static inline int submit_malloc_event(size_t size, bool is_return, int ret_val) {
struct event *e = bpf_ringbuf_reserve(&rb, sizeof(*e), 0);
if (!e) return 0;
/* Fill common fields */
e->pid = bpf_get_current_pid_tgid() >> 32;
bpf_get_current_comm(&e->comm, sizeof(e->comm));
e->type = CUDA_EVENT_MALLOC;
e->is_return = is_return;
/* Fill event-specific data */
if (is_return) {
e->ret_val = ret_val;
} else {
e->mem.size = size;
}
bpf_ringbuf_submit(e, 0);
return 0;
}
```
この関数では、まずリングバッファ内にイベント用の領域を確保します。その後、プロセスIDやプロセス名といった共通フィールドを設定します。mallocイベントの場合、関数呼び出し時（要求サイズ取得時）または関数終了時（戻り値取得時）のいずれかの時点で、適切な値を保存します。最後に、このイベントをリングバッファに送信します。
実際のプローブは、SEC注釈を使用してCUDA関数にアタッチされています。cudaMallocの場合、以下の設定を行います：
```js
SEC("uprobe")
int BPF_KPROBE(cuda_malloc_enter, void **ptr, size_t size) {
return submit_malloc_event(size, false, 0);
}
SEC("uretprobe")
int BPF_KRETPROBE(cuda_malloc_exit, int ret) {
return submit_malloc_event(0, true, ret);
}
```
最初の関数は `cudaMalloc` が呼び出された際に実行され、要求されたメモリサイズを取得します。2番目の関数は `cudaMalloc` が終了する際に実行され、エラーコードを取得します。このパターンは、追跡対象とする各CUDA関数に対して同様に適用されます。
特に興味深いケースとして `cudaMemcpy` が挙げられます。これはホストとデバイス間でデータを転送する関数です：
```js
SEC("uprobe")
int BPF_KPROBE(cuda_memcpy_enter, void *dst, const void *src, size_t size, int kind) {
return submit_memcpy_event(size, kind, false, 0);
}
```
ここでは、サイズだけでなく「kind」パラメータも取得しています。このパラメータはデータ転送の方向（ホスト→デバイス、デバイス→ホスト、またはデバイス間）を示しており、データ移動パターンに関する貴重な情報を提供します。
## ユーザー空間アプリケーションの詳細
ユーザー空間アプリケーション（ `cuda_events.c` ）は、eBPFプログラムの読み込み、リングバッファからのイベント処理、およびその結果をユーザーフレンドリーな形式で表示する役割を担っています。
まず、プログラムはコマンドライン引数を解析して動作設定を行います：
```js
static struct env {
bool verbose;
bool print_timestamp;
char *cuda_library_path;
bool include_returns;
int target_pid;
} env = {
.print_timestamp = true,
.include_returns = true,
.cuda_library_path = NULL,
.target_pid = -1,
};
```
この構造体には、タイムスタンプの表示有無やリターンプロービングの実行有無など、アプリケーションの動作を制御する設定オプションが格納されます。デフォルト値は、実用的な初期設定として設定されています。
本プログラムでは `libbpf` ライブラリを使用してeBPFプログラムをロードし、CUDA関数にアタッチします：
```js
int attach_cuda_func(struct cuda_events_bpf *skel, const char *lib_path,
const char *func_name, struct bpf_program *prog_entry,
struct bpf_program *prog_exit) {
/* Attach entry uprobe */
if (prog_entry) {
uprobe_opts.func_name = func_name;
struct bpf_link *link = bpf_program__attach_uprobe_opts(prog_entry,
env.target_pid, lib_path, 0, &uprobe_opts);
/* Error handling... */
}
/* Attach exit uprobe */
if (prog_exit) {
/* Similar for return probe... */
}
}
```
この関数は、関数名（例："cudaMalloc"）と対応するエントリーポイント/エグジットポイント用のeBPFプログラムを受け取り、指定されたライブラリに対してこれらのプログラムをuprobeとしてアタッチします。
最も重要な機能の一つが `handle_event` 関数です。この関数はリングバッファからイベントを処理します：
```js
static int handle_event(void *ctx, void *data, size_t data_sz) {
const struct event *e = data;
struct tm *tm;
char ts[32];
char details[MAX_DETAILS_LEN];
time_t t;
/* Skip return probes if requested */
if (e->is_return && !env.include_returns)
return 0;
time(&t);
tm = localtime(&t);
strftime(ts, sizeof(ts), "%H:%M:%S", tm);
get_event_details(e, details, sizeof(details));
if (env.print_timestamp) {
printf("%-8s ", ts);
}
printf("%-16s %-7d %-20s %8s %s\n",
e->comm, e->pid,
event_type_str(e->type),
e->is_return ? "[EXIT]" : "[ENTER]",
details);
return 0;
}
```
この関数はイベント情報をフォーマットして表示します。タイムスタンプ、プロセス詳細、イベント種別、および特定のパラメータや戻り値などが含まれます。
`get_event_details` 関数は、生のイベントデータを人間が読みやすい形式に変換します：
```js
static void get_event_details(const struct event *e, char *details, size_t len) {
switch (e->type) {
case CUDA_EVENT_MALLOC:
if (!e->is_return)
snprintf(details, len, "size=%zu bytes", e->mem.size);
else
snprintf(details, len, "returned=%s", cuda_error_str(e->ret_val));
break;
/* Similar cases for other event types... */
}
}
```
この関数は各イベントタイプに応じて異なる処理を行います。例えば、mallocイベントの場合、エントリー時には要求されたメモリサイズを、エグジット時にはエラーコードを表示します。
メインのイベントループは非常に簡潔に構成されています：
```js
while (!exiting) {
err = ring_buffer__poll(rb, 100 /* timeout, ms */);
/* Error handling... */
}
```
このコードはリングバッファからイベントをポーリングし、各イベントに対して `handle_event` 関数を呼び出します。100ミリ秒のタイムアウト設定により、Ctrl+Cなどのシグナルに対してもプログラムが適切に応答できるようになっています。
当社のトレーサーの重要な機能の一つは、CUDAのエラーコードを人間が理解できるメッセージに変換することです。CUDAには「メモリ不足」といった単純なエラーコードから、「サポートされていないPTXバージョン」といった複雑なエラーコードまで、100種類以上の異なるエラーコードが存在します。
当社のツールには、数値コードを文字列説明に変換する包括的な `cuda_error_str` 関数が含まれています：
```js
static const char *cuda_error_str(int error) {
switch (error) {
case 0: return "Success";
case 1: return "InvalidValue";
case 2: return "OutOfMemory";
/* Many more error codes... */
default: return "Unknown";
}
}
```
これにより、出力結果がデバッグ作業においてはるかに有用になります。単に「エラー 2」と表示される代わりに、「OutOfMemory」という具体的なエラーメッセージが表示されるため、何が問題なのかをすぐに把握できます。
## コンパイルと実行方法
トレースツールのビルドは、付属のMakefileを使用して簡単に行えます：
```js
# Build both the tracer and the example
make
```
これにより2つのバイナリが生成されます： - `cuda_events` ：eBPFベースのCUDAトレースツール - `basic02` ：シンプルなCUDAサンプルアプリケーション
ビルドシステムは賢く、 `nvidia-smi` を使用してGPUアーキテクチャを自動検出し、適切なコンパイルオプションでCUDAコードをコンパイルします。
トレーサーの実行も非常に簡単です：
```js
# Start the tracing tool
sudo ./cuda_events -p ./basic02
# In another terminal, run the CUDA example
./basic02
```
特定のプロセスをPIDでトレースすることも可能です：
```js
# Run the CUDA example
./basic02 &
PID=$!
# Start the tracing tool with PID filtering
sudo ./cuda_events -p ./basic02 -d $PID
```
以下のサンプル出力では、各CUDA操作に関する詳細な情報が表示されています：
```js
Using CUDA library: ./basic02
TIME PROCESS PID EVENT TYPE DETAILS
17:35:41 basic02 12345 cudaMalloc [ENTER] size=4000 bytes
17:35:41 basic02 12345 cudaMalloc [EXIT] returned=Success
17:35:41 basic02 12345 cudaMalloc [ENTER] size=4000 bytes
17:35:41 basic02 12345 cudaMalloc [EXIT] returned=Success
17:35:41 basic02 12345 cudaMemcpy [ENTER] size=4000 bytes, kind=1
17:35:41 basic02 12345 cudaMemcpy [EXIT] returned=Success
17:35:41 basic02 12345 cudaLaunchKernel [ENTER] func=0x7f1234567890
17:35:41 basic02 12345 cudaLaunchKernel [EXIT] returned=Success
17:35:41 basic02 12345 cudaMemcpy [ENTER] size=4000 bytes, kind=2
17:35:41 basic02 12345 cudaMemcpy [EXIT] returned=Success
17:35:41 basic02 12345 cudaFree [ENTER] ptr=0x7f1234568000
17:35:41 basic02 12345 cudaFree [EXIT] returned=Success
17:35:41 basic02 12345 cudaFree [ENTER] ptr=0x7f1234569000
17:35:41 basic02 12345 cudaFree [EXIT] returned=Success
```
この出力結果は、典型的なCUDAアプリケーションの処理フローを示しています： 1. デバイス上でメモリを割り当てる 2. ホストからデバイスへデータをコピー（kind=1） 3. データを処理するためのカーネルを起動する 4. 処理結果をデバイスからホストへコピーバック（kind=2） 5. デバイス上のメモリを解放する
## ベンチマークツール
また、トラッカーの性能評価とCUDA API呼び出しのレイテンシ測定を行うためのベンチマークツールも提供しています。
```js
make
sudo ./cuda_events -p ./bench
./bench
```
トレースが実行されていない場合の結果は以下の通りです：
```js
Data size: 1048576 bytes (1024 KB)
Iterations: 10000
Summary (average time per operation):
-----------------------------------
cudaMalloc: 113.14 µs
cudaMemcpyH2D: 365.85 µs
cudaLaunchKernel: 7.82 µs
cudaMemcpyD2H: 393.55 µs
cudaFree: 0.00 µs
```
トレーサーがアタッチされている場合の結果は以下の通りです：
```js
Data size: 1048576 bytes (1024 KB)
Iterations: 10000
Summary (average time per operation):
-----------------------------------
cudaMalloc: 119.81 µs
cudaMemcpyH2D: 367.16 µs
cudaLaunchKernel: 8.77 µs
cudaMemcpyD2H: 383.66 µs
cudaFree: 0.00 µs
```
トレーサーは各CUDA API呼び出しに約2マイクロ秒のオーバーヘッドを追加しますが、ほとんどのケースではこの影響は無視できる程度です。オーバーヘッドをさらに削減するには、 [bpftime](https://github.com/eunomia-bpf/bpftime) ユーザー空間ランタイムを使用してeBPFプログラムを最適化する方法をお試しください。
## コマンドラインオプション
`cuda_events` ツールでは以下のオプションを使用できます：
- `-v`: デバッグ用の詳細出力を有効にします
- `-t`: タイムスタンプを表示しない
- `-r`: 関数の戻り値を表示しない（関数のエントリーポイントのみ表示）
- `-p PATH`: CUDAランタイムライブラリまたはアプリケーションのパスを指定します
- `-d PID`: 指定したプロセスIDのみをトレース対象とします
この基本的なCUDAトレースツールに慣れたら、以下の機能拡張を検討できます：
1. より多くのCUDA API関数に対応する機能を追加
2. パフォーマンスのボトルネックを分析するためのタイミング情報を追加します
3. 関連する操作間の相関関係を実装します（例：mallocとfreeの対応関係の特定）
4. CUDA操作の可視化データを作成し、分析を容易にします
5. OpenCLやROCmなど、他のGPUフレームワークへの対応機能を追加
CUDAのサンプルコードとチュートリアルの詳細については、リポジトリ [https://github.com/eunomia-bpf/basic-cuda-tutorial](https://github.com/eunomia-bpf/basic-cuda-tutorial) をご覧ください。
本チュートリアルのコードは [https://github.com/eunomia-bpf/bpf-developer-tutorial/tree/main/src/47-cuda-events](https://github.com/eunomia-bpf/bpf-developer-tutorial/tree/main/src/47-cuda-events) で公開されています。
## 参考文献
- CUDAプログラミングガイド: [https://docs.nvidia.com/cuda/cuda-c-programming-guide/](https://docs.nvidia.com/cuda/cuda-c-programming-guide/)
- NVIDIA CUDAランタイムAPI: [https://docs.nvidia.com/cuda/cuda-runtime-api/](https://docs.nvidia.com/cuda/cuda-runtime-api/)
- libbpfドキュメント: [https://libbpf.readthedocs.io/](https://libbpf.readthedocs.io/)
- Linux uprobesドキュメント: [https://www.kernel.org/doc/Documentation/trace/uprobetracer.txt](https://www.kernel.org/doc/Documentation/trace/uprobetracer.txt)
- eGPU: GPU上でのeBPF: [https://dl.acm.org/doi/10.1145/3723851.3726984](https://dl.acm.org/doi/10.1145/3723851.3726984)
- bpftime GPUサンプルコード: [https://github.com/eunomia-bpf/bpftime/tree/master/example/gpu](https://github.com/eunomia-bpf/bpftime/tree/master/example/gpu)
eBPFについてさらに詳しく学びたい方は、私たちのチュートリアルリポジトリ [https://github.com/eunomia-bpf/bpf-developer-tutorial](https://github.com/eunomia-bpf/bpf-developer-tutorial) をご覧いただくか、公式ウェブサイト [https://eunomia.dev/tutorials/](https://eunomia.dev/tutorials/) をご覧ください。