# cuda-knl-launch-tracing

**Repository Path**: zeasa/cuda-knl-launch-tracing

## Basic Information

- **Project Name**: cuda-knl-launch-tracing
- **Description**: No description available
- **Primary Language**: C++
- **License**: Apache-2.0
- **Default Branch**: master
- **Homepage**: None
- **GVP Project**: No

## Statistics

- **Stars**: 2
- **Forks**: 1
- **Created**: 2026-01-29
- **Last Updated**: 2026-03-23

## Categories & Tags

**Categories**: Uncategorized

**Tags**: None

## README

# CUDA Kernel Launch Latency

测量 CUDA 空内核启动延迟的 benchmark 工具。支持与外部 tracing 库（如 `libtrace_symbol.so`）配合，用于分析内核启动性能。

提供 **CUDA 原生** 和 **Python** 两种实现，选项和输出保持一致。

## 依赖

- CUDA Toolkit
- 可选：`libtrace_symbol.so`（提供 `start_tracing` / `stop_tracing` 符号，用于 tracing 时打点）

---

## 1. CUDA 版本 (kernel_launch_latency)

### 依赖

- CUDA Toolkit

### 编译

```bash
make
```

默认使用 `nvcc -O2 -arch=sm_80`。可根据 GPU 架构修改 `makefile` 中的 `-arch` 参数。

### 用法

```bash
./kernel_launch_latency [OPTIONS]
```

---

## 2. Python 版本 (kernel_launch_latency.py)

### 依赖

- Python 3
- **pycuda** 或 **cupy**（二选一，优先尝试 pycuda）

```bash
pip install pycuda
# 或
pip install cupy-cuda12x   # 按 CUDA 版本选择
```

### 用法

```bash
python kernel_launch_latency.py [OPTIONS]
```

---

## 通用选项

两种实现的命令行选项一致：

| 选项 | 长选项 | 说明 | 默认值 |
|------|--------|------|--------|
| `-i N` | `--iter N` | 测试迭代次数 | 100000 |
| `-w N` | `--warm N` | 预热迭代次数 | 1000 |
| `-s` | `--sync` | 每次启动后同步（不带则最后统一同步） | 否 |
| `-g ID` | `--gpu ID` | 使用的 GPU 设备 ID | 0 |
| `-k K` | `--kernel K` | 测试内核：0=empty，1=heavy（16 参数+约 50 条指令） | 0 |
| `-h` | `--help` | 显示帮助信息 | - |

### 内核类型

- **`-k 0` (empty)**：空内核，用于测量纯启动/入队延迟
- **`-k 1` (heavy_args)**：带 16 个标量参数及约 50 条算逻指令的内核，用于测试参数编码、PTX 解析等开销

### 同步模式

- **带 `-s` / `--sync`**：每次内核启动后调用 `cudaDeviceSynchronize()`，测量「单次启动+同步」延迟
- **不带**：循环内不同步，最后统一同步，测量内核入队（enqueue）延迟

### 示例

```bash
# 默认参数运行（CUDA / Python）
./kernel_launch_latency
python kernel_launch_latency.py

# 1000 次迭代，50 次预热
./kernel_launch_latency -i 1000 -w 50
python kernel_launch_latency.py -i 1000 -w 50

# 测量 sync 延迟（带 -s）
./kernel_launch_latency -s
python kernel_launch_latency.py -s

# 指定使用 GPU 1
./kernel_launch_latency -g 1
python kernel_launch_latency.py -g 1

# 测试 heavy_args 内核（多参数、多指令）
./kernel_launch_latency -k 1
python kernel_launch_latency.py -k 1

# 查看帮助
./kernel_launch_latency --help
python kernel_launch_latency.py --help
```

## 输出

程序初始化阶段会打印可用 GPU 列表，例如：

```
Available GPUs:
  [0] NVIDIA GeForce RTX 3080
  [1] NVIDIA GeForce RTX 3080
```

随后会打印：

- `kernel = empty` 或 `kernel = heavy_args (16 params + ~50 ops)`：当前测试内核
- `sync_in_loop = 0|1`：当前同步模式
- `Kernel launch latency (sync each): X.XXX us` 或  
  `Kernel enqueue latency (sync once): X.XXX us`：测量得到的平均延迟（微秒）