Warp API Capture（APIC）实战：CUDA 图捕获、序列化与 C++ 独立回放——以 APIC 波浪模拟为例

Warp API CaptureAPIC实战CUDA 图捕获、序列化与 C 独立回放——以 APIC 波浪模拟为例【免费下载链接】warpA Python framework for GPU-accelerated simulation, robotics, and machine learning.项目地址: https://gitcode.com/GitHub_Trending/warp/warp本篇技术指南以 Warp 仓库中的02_apic_visualization示例为主线系统讲解 Warp API CaptureAPIC的完整工作流如何在 Python 侧把一段多内核模拟帧17 次内核启动捕获成可序列化的操作流并保存为.wrp文件再由一个不依赖 Python 运行时的 C 程序加载、重建 CUDA 图并逐帧交互式回放。读完本文你将掌握wp.capture_begin(apicTrue)/wp.capture_save()的使用方式、wp_apic_*C API 的调用方法以及 CUDA Graph 一次性启动整帧的部署与性能优化思路。背景为什么需要 APICWarp 的常规wp.capture_begin()/wp.capture_end()会在设备上录制一个活 CUDA 图live graph但该图无法脱离 Python 运行时部署。APICAPI Capture扩展解决了这个问题捕获时在活 CUDA 图之外再记录一条可序列化的 APIC 操作流operation stream。wp.capture_save()把这条流与编译好的 CUDA 模块cubin一起写入磁盘C 端用wp_apic_load_graph()读回重建一个全新的 CUDA 图即可在没有 Python 的环境里反复回放同一份计算负载。从源码结构看APIC 的操作流由 C 侧直接记录apic.h 注释明确Operations are recorded directly in C via hooks in wp_cuda_launch_kernel(), wp_cpu_launch_kernel(), and the memcpy/memset functions… The C byte stream is the single source of truthPython 仅负责在每个内核启动时构建APICLaunchInfo结构。因此02_apic_visualization展示的Python 捕获 C 回放正是 APIC 三个消费路径中最典型的部署形态。示例一览文件与产物示例位于 warp/examples/cpp/02_apic_visualization包含以下文件capture_wave.pyPython 脚本捕获整帧波浪模拟图并保存main.cuC 程序加载图、更新输入、GLFW/OpenGL 交互可视化Makefile / CMakeLists.txt两套构建系统Make 面向 UnixCMake 跨平台捕获脚本运行后自动生成generated/目录其中generated/wave_sim.wrp序列化的 APIC 图表示generated/wave_sim_modules/编译好的 CUDA 模块cubins示例演示的三项核心能力独立部署C 应用运行时不需要 Python图重建APIC 操作流重建出的 CUDA 图包含17 次内核启动1 次位移 16 次波动方程求解迭代动态参数命名输入named inputs可在每一帧直接更新无需重建图低开销回放C 每帧只调用一次cudaGraphLaunch()完成整个模拟。环境准备与快速开始前置要求示例在 Warp 仓库内运行需要Python 3.10CUDA Toolkit 12.0含nvcc编译器NVIDIA GPUGLFW 库用于创建 OpenGL 窗口Warp 原生库随 Warp 一起构建构建系统GNU MakeUnix/Linux或 CMake 3.20跨平台注意macOS 不支持macOS 上无 CUDAMakefile 中对 Darwin 直接$(error)拒绝构建克隆并构建 Warpgit clone https://github.com/NVIDIA/warp.git cd warp # 构建并安装二选一 uv run build_lib.py # 方式 A推荐 # 或python build_lib.py pip install -e . # 方式 B cd warp/examples/cpp/02_apic_visualization使用 Make 构建运行Unix/Linuxmake # 构建全部必要时自动捕获图 ./02_apic_visualization # 运行Make 的默认目标链值得说明02_apic_visualization可执行文件依赖generated/wave_sim.wrp与 glad OpenGL 加载器通过 Makefile 中的$(GRAPH_SENTINEL)和$(GLAD_HEADER)规则自动生成也就是说make会自动先跑capture_wave.py无需手动捕获。Make 目标一览目标作用make默认构建全部必要时自动捕获图make cpp仅构建 C 程序快速迭代假定图已存在make capture仅生成 APIC 图产物不构建 Cmake clean仅删除可执行文件make distclean删除可执行文件与generated/目录Make 构建中两个可覆盖的环境变量WARP_NATIVE_DIR指向warp/native/与WARP_BIN_DIR指向含libwarp.so的warp/bin/默认值均为相对仓库根目录的相对路径。此外CUDA_ARCH可覆盖设备端update_vertices内核的架构默认sm_80。GLFW 头文件在 Make 路径下是系统依赖构建前会先用宿主编译器探测GLFW/glfw3.h缺失时给出安装提示。使用 CMake 构建跨平台python capture_wave.py # 第 1 步捕获图 cmake -B build -DCMAKE_BUILD_TYPERelease # 第 2 步配置 cmake --build build --config Release # 第 3 步构建 ./build/02_apic_visualization # 第 4 步运行Ninja / Unix MakefilesWindows 上使用 Visual Studio 生成器时运行.\build\Release\02_apic_visualization.exeCMake 路径与 Make 路径的差异详见 CMakeLists.txtGLFW 与 GLAD 通过FetchContent自动拉取GLFW 3.3.8、glad v2.0.8不需要系统安装 GLFW头文件下是否有 OpenGL 是示例是否构建的开关find_package(OpenGL QUIET)失败时直接跳过该示例适合无显示服务器的 CIWarp 库以IMPORTED目标warp::runtime链接避免 Linux 上libwarp.so无libname.so前缀导致-lwarp链接失败默认 CUDA 架构为 75Turing可用-DCMAKE_CUDA_ARCHITECTURES覆盖Windows 下构建后自动把warp.dll复制到输出目录。无头冒烟模式Headless Smoke示例接受单参数--smoke不打开 GLFW 窗口只做加载图、查询参数、连续回放 10 次的冒烟检查。CTest 以apic_visualization_smoke注册该模式PASS_REGULAR_EXPRESSION断言输出smoke OK (10 graph launches)防止冒烟分支静默退出 0 的回归因此它可以在没有显示服务器的 CI 主机上运行./02_apic_visualization --smoke # 退出码 0输出 smoke OK (10 graph launches)工作原理解析阶段 1Python捕获多内核整帧图核心思路是把一整个模拟帧多个子步的全部内核启动装进一张图而不是只捕获单个内核。capture_wave.py 中定义了三个 Warp 组件sample_height带边界钳制的采样函数wave_displace内核读取mouse_pos参数数组对落在半径radius内的格点施加平滑衰减的位移falloff 1.0 - dist_sq / radius²产生点击出涟漪的效果wave_solve内核五点模板计算拉普拉斯算子用显式有限差分积分波动方程h_new 2*h - h0 dt²·(k_speed·laplacian − k_damp·(h − h0))其中k_speed60.0是波速平方、k_damp1000.0是阻尼系数、dt 1/60/substeps。捕获主流程模拟常量可从命令行覆盖--width、--height、--substeps、--outputgrid0 wp.zeros(width * height, dtypefloat, devicedevice) grid1 wp.zeros(width * height, dtypefloat, devicedevice) mouse_pos wp.zeros(1, dtypewp.vec2, devicedevice) # 把 ALL 子步捕获进一张图 wp.capture_begin(devicedevice, apicTrue) for s in range(substeps): # 例如 16 次迭代 if s 0: # 第一个子步施加鼠标位移 wp.launch(wave_displace, dimwidth * height, inputs[grid0, grid1, mouse_pos, ...]) # 每个子步积分波动方程 wp.launch(wave_solve, dimwidth * height, inputs[grid0, grid1, ...]) # 交换缓冲区 grid0, grid1 grid1, grid0 graph wp.capture_end(devicedevice) # 带命名绑定保存 wp.capture_save( graph, generated/wave_sim, inputs{heights: grid1, heights_prev: grid0, mouse_pos: mouse_pos}, outputs{heights_out: grid1, heights_prev_out: grid0}, )关键点apicTrue是 APIC 捕获的开关必须在capture_begin时传入波动方程需要两个缓冲区当前与上一帧用于求速度每子步后交换16 子步后grid1持有最新的h(t16)grid0持有h(t15)因此保存时heights绑定grid1、heights_prev绑定grid0inputs/outputs中的命名绑定会成为 C 端动态更新的入口该操作流共描述17 次内核启动1 次 displacement 16 次 solve。APIC 操作流的底层结构可参考 warp/_src/apic/types.py操作类型枚举APIC_OP_KERNEL_LAUNCH、APIC_OP_MEMCPY_H2D/D2H/D2D、APIC_OP_MEMSET、APIC_OP_ALLOC、APIC_OP_IF/WHILE、各种排序/扫描/规约算子等与标量类型枚举每个内核参数被描述为一个APICLaunchParamRecord参数切片[value_offset, value_offsetvalue_size)加重定位计数指针字段由APICLaunchPtrLocation重定位条目在回放时按三种 kind 修补APIC_RELOC_DATA_PTR按 region 解析指针、APIC_RELOC_HANDLE句柄重映射、APIC_RELOC_NULL显式零指针。在 warp/_src/apic/capture.py 中可以看到对应的 Python 侧实现APICCapture负责内存区域注册track_array把数组的基地址分配注册为 region、APICLaunchInfo构造以及句柄字段的递归查找_find_handle_offsets用于wp.handle字段在加载后重映射 Mesh/Volume/BVH 对象。注意_walk_value_pointers中有明确的限制APIC 捕获尚不支持wp.fixedarray与wp.fabricarray作为启动参数会抛NotImplementedError编写可捕获内核时应避开这些类型。阶段 2C加载图并可视化main.cu 的初始化流程--graph参数可指定图路径默认generated/wave_simcuInit(0)初始化驱动 APIcuDevicePrimaryCtxRetain获取主 CUDA 上下文wp_init(nullptr)初始化 Warp 运行时wp_apic_load_graph(context, path, 0)加载图0 APIC_DEVICE_CUDA遍历wp_apic_get_num_params/wp_apic_get_param_name/wp_apic_get_param_size打印全部命名参数wp_apic_get_cuda_graph_exec(graph)构建 CUDA 图可执行句柄首次调用时构建。主循环每帧只做一次图启动#include aot.h // Warp AOT 工具含 CUDA #include warp.h // Warp C API #include apic.h // APIC 图加载与执行 APICGraph* graph wp_apic_load_graph(context, generated/wave_sim, 0); // 0 CUDA int n_params wp_apic_get_num_params(graph); for (int i 0; i n_params; i) { const char* name wp_apic_get_param_name(graph, i); size_t size wp_apic_get_param_size(graph, name); printf( %s: %zu bytes\n, name, size); } cudaGraphExec_t exec (cudaGraphExec_t)wp_apic_get_cuda_graph_exec(graph); while (!glfwWindowShouldClose(window)) { // 从 GLFW 输入更新鼠标位置 float mouse_grid[2] { ... }; cudaMemcpyAsync(d_mouse_pos, mouse_grid, 2 * sizeof(float), cudaMemcpyHostToDevice, stream); // 同时设置两个波状态缓冲与交互输入 wp_apic_set_param(graph, heights, d_heights[current_buffer], heights_size); wp_apic_set_param(graph, heights_prev, d_heights[1 - current_buffer], heights_size); wp_apic_set_param(graph, mouse_pos, d_mouse_pos, 2 * sizeof(float)); // 执行图——一次调用跑完全部 17 个内核 cudaGraphLaunch(exec, stream); cudaStreamSynchronize(stream); // 读取图输出前必须同步 // 读取两个输出缓冲然后为下一帧交换 wp_apic_get_param(graph, heights_out, d_heights[1 - current_buffer], heights_size); wp_apic_get_param(graph, heights_prev_out, d_heights[current_buffer], heights_size); current_buffer 1 - current_buffer; // OpenGL 渲染... } wp_apic_destroy_graph(graph);交互参数通过wp_apic_set_param每帧原位更新heights、heights_prev、mouse_pos三个命名输入图结构完全不变因此无需重建。结果经wp_apic_get_param读回两个输出缓冲再交换角色作为下一帧输入形成连续模拟。渲染侧的要点见 main.cu一个自定义 CUDA 内核update_vertices把高度场写入网格顶点X/Z 由格点坐标乘以GRID_SCALE0.1得到Y 为高度值再cudaMemcpyAsync拷回主机更新 VBO着色器通过dFdx/dFdy从世界坐标求法线做简单的漫反射 环境光着色蓝色水面头部注释特别提醒包含顺序GLAD 必须先于其他 GL 头文件且aot.h含 CUDA须在 GLAD 之后以避免类型冲突窗口标题实时显示 FPS。APIC C API 参考以下为示例实际使用的 APIC C API 全集签名与 apic.h 保持一致// 从 .wrp 文件加载图device_type: APIC_DEVICE_CUDA0, APIC_DEVICE_CPU1 APICGraph* wp_apic_load_graph(void* context, const char* path, int device_type); // 设置/获取命名参数成功返回 true失败返回 false bool wp_apic_set_param(APICGraph* graph, const char* name, const void* data, size_t size); bool wp_apic_get_param(APICGraph* graph, const char* name, void* data, size_t size); void* wp_apic_get_param_ptr(APICGraph* graph, const char* name); // 获取 CUDA 图句柄 void* wp_apic_get_cuda_graph(APICGraph* graph); void* wp_apic_get_cuda_graph_exec(APICGraph* graph); // 首次调用时创建可执行图 // 查询参数 int wp_apic_get_num_params(APICGraph* graph); const char* wp_apic_get_param_name(APICGraph* graph, int index); size_t wp_apic_get_param_size(APICGraph* graph, const char* name); // 清理 void wp_apic_destroy_graph(APICGraph* graph);wp_apic_get_cuda_graph_exec返回的句柄可直接强转为cudaGraphExec_t用于cudaGraphLaunch。除参数外加载后的图还暴露内核元数据查询接口wp_apic_get_num_kernels、wp_apic_get_kernel_key等CPU 回放加载.o模块时使用。APIC 的整体消费路径在 apic.h 中有明确注释wp_apic_state_save()写.wrp文件、wp_apic_cpu_replay_state()执行活捕获、wp_apic_cpu_replay_graph()/wp_apic_get_cuda_graph()执行已加载的.wrp图CPU / CUDA。交互控制左键点击在鼠标位置制造波纹右键拖动旋转相机俯仰角钳制在 ±89°滚轮缩放距离钳制在 5.0–50.0ESC退出相机/交互状态在 main.cu 中由 GLFW 回调维护mouse_button_callback、cursor_pos_callback、scroll_callback、framebuffer_size_callback。鼠标到网格的映射是简化实现注释明确说明完整实现应使用射线拾取左键按下时把屏幕坐标线性映射到[0, GRID_WIDTH] × [0, GRID_HEIGHT]未按下时置为网格外坐标(-1000, -1000)。部署与性能APIC 让被捕获的工作负载可以脱离 Python 独立部署。性能收益来自 CUDA 图回放本身的启动效率与在 C 中逐个发出 17 次 CUDA 启动相比重建后的图每帧只需一次cudaGraphLaunch()避免了反复的主机端启动开销——这对包含大量小内核的工作负载尤其有价值。从实现看图的执行还复用了 Warp 序列化时保存的编译模块cubins重建图无需重新编译内核。故障排查Failed to load graph先确认图已捕获python capture_wave.pyCould not find Warp library设置库路径export LD_LIBRARY_PATH/path/to/warp/bin:$LD_LIBRARY_PATH # Linux # 或 cmake -DWARP_BIN_DIR/path/to/warp/bin ... # CMakeGLFW not found安装 GLFW 开发包sudo apt install libglfw3-dev # Ubuntu/Debian sudo dnf install glfw-devel # Fedora brew install glfw # macOS注意macOS 仍不支持 CUDA继续深入仓库示例完整源码capture_wave.py、main.cu构建脚本Makefile、CMakeLists.txtAPIC 底层实现warp/_src/apic/capture.py区域注册、APICLaunchInfo构造、重定位、warp/_src/apic/types.py操作类型/重定位 kind 定义C 接口与消费路径warp/native/apic.h、warp/native/apic_types.h相关测试可关注 warp/tests 下与 graph、compile、aot 相关的用例验证捕获/回放行为【免费下载链接】warpA Python framework for GPU-accelerated simulation, robotics, and machine learning.项目地址: https://gitcode.com/GitHub_Trending/warp/warp创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考