Jitify

一个单头文件的 C++ 库，用于简化 CUDA 运行时编译（NVRTC）的使用。

设计动机

将 NVRTC 集成到现有的或模板化的 CUDA 代码中可能会比较棘手。Jitify 的目标是通过提供一个简单、高层次的接口来隐藏这些复杂性，从而简化这一过程。

快速示例

const char* program_source = "my_program\n"
    "template<int N, typename T>\n"
    "__global__\n"
    "void my_kernel(T* data) {\n"
    "    T data0 = data[0];\n"
    "    for( int i=0; i<N-1; ++i ) {\n"
    "        data[0] *= data0;\n"
    "    }\n"
    "}\n";
static jitify::JitCache kernel_cache;
jitify::Program program = kernel_cache.program(program_source);
// ...设置数据等
dim3 grid(1);
dim3 block(1);
using jitify::reflection::type_of;
program.kernel("my_kernel")
       .instantiate(3, type_of(*data))
       .configure(grid, block)
       .launch(data);

特性

Jitify 提供并处理以下功能：

• 所有 NVRTC 和 CUDA 驱动程序 API 调用
• 简单的内核实例化和启动语法
• 缓存已编译的内核
• 从字符串、文件或嵌入在可执行文件中的源代码加载
• 忽略运行时编译源中的主机代码
• 跳过不必要的头文件
• 支持 JIT 安全的标准库头文件（如 float.h、stdint.h 等）
• 处理内核名称修饰问题
• 将内核模板参数反射为字符串
• 针对当前设备的计算能力进行专门编译
• 链接预编译的 PTX/CUBIN/FATBIN/对象文件或库
• 支持 CUDA 7.0、7.5、8.0、9.x、10.x 版本，适用于 Linux 和 Windows 平台
• 方便的 parallel_for 函数及 lambda 支持
• *新增* jitify::experimental API 提供序列化功能，以实现 用户管理的哈希与缓存

使用 Jitify 和 NVRTC 可以做到：

• 快速移植现有代码 以使用 CUDA 运行时编译
• 大幅减少代码量 并缩短离线编译时间
• 提升内核性能，通过在运行时嵌入常量和自动调优来实现

如何构建

Jitify 只是一个单头文件：

#include <jitify.hpp>

编译时需添加：-pthread（如果 JITIFY_THREAD_SAFE 定义为 0，则无需此选项）

链接时需添加：-lcuda -lcudart -lnvrtc

此外，还包含一个名为 stringify 的小工具，用于将文本文件转换为 C 字符串字面量，从而方便地将 JIT 编译的源代码集成到构建流程中。

运行测试

可以通过以下命令运行测试：

$ make test

该命令会自动下载并构建 GoogleTest 库，而 GoogleTest 需要系统上安装 CMake。

文档

示例

请参阅 jitify_example.cpp，其中包含一些使用该库的示例。Makefile 也展示了如何使用提供的 stringify 工具。

GTC 2017 演讲——Ben Barsdell 和 Kate Clark

API 文档

可以通过运行以下命令生成 Doxygen 文档：

$ make doc

生成的 HTML 和 LaTeX 文件将放置在 doc/ 子目录中。

许可证

BSD-3-Clause

作者

Ben Barsdell（NVIDIA，邮箱：bbarsdell at nvidia dot com）

Kate Clark（NVIDIA，邮箱：mclark at nvidia dot com）

Jitify 快速上手指南

Jitify 是一个单头文件的 C++ 库，旨在简化 CUDA 运行时编译（NVRTC）的使用。它通过高级接口隐藏了底层复杂性，支持模板内核的即时编译、缓存及自动处理名称修饰，帮助开发者快速移植代码并减少离线编译时间。

环境准备

在使用 Jitify 之前，请确保您的开发环境满足以下要求：

• 操作系统：Linux 或 Windows。
• CUDA Toolkit：支持 CUDA 7.0 至 10.x 及更高版本（需包含 nvrtc 组件）。
• 编译器：支持 C++11 或更高版本的 GCC/Clang (Linux) 或 MSVC (Windows)。
• 构建工具：
  • make 和 cmake（仅运行测试或生成文档时需要）。
  • 若需运行官方测试套件，需能访问互联网以自动下载 GoogleTest。

安装步骤

Jitify 无需传统的安装过程，它是一个单头文件库。

1. 获取头文件 从 GitHub 仓库下载 jitify.hpp 文件，并将其放置在项目的包含路径（include path）中，或直接复制到项目源码目录。

   # 示例：克隆仓库并复制头文件
   git clone https://github.com/NVIDIA/jitify.git
   cp jitify/jitify.hpp /your/project/path/
   

2. 编译配置 在编译包含 Jitify 的 C++ 源文件时，需添加以下标志：

   • 包含路径：确保编译器能找到 jitify.hpp。
   • 线程支持：添加 -pthread 标志（若在代码中定义了 JITIFY_THREAD_SAFE 0 则不需要）。
   • 链接库：链接 CUDA 相关库 -lcuda -lcudart -lnvrtc。

   编译命令示例：

   g++ -o my_app my_app.cu -I. -pthread -lcuda -lcudart -lnvrtc
   

   (注：如果使用 nvcc 编译，通常也能直接识别这些依赖，但显式链接更稳妥)

基本使用

以下是最小化的使用示例，展示了如何定义 CUDA 内核字符串、实例化模板参数并启动内核。

#include <jitify.hpp>
#include <cuda_runtime.h>
#include <iostream>

// 1. 定义 CUDA 内核源代码字符串
const char* program_source = "my_program\n"
    "template<int N, typename T>\n"
    "__global__\n"
    "void my_kernel(T* data) {\n"
    "    T data0 = data[0];\n"
    "    for( int i=0; i<N-1; ++i ) {\n"
    "        data[0] *= data0;\n"
    "    }\n"
    "}\n";

int main() {
    // 2. 创建内核缓存对象
    static jitify::JitCache kernel_cache;

    // 3. 构建程序对象
    jitify::Program program = kernel_cache.program(program_source);

    // 模拟数据准备 (实际使用时请分配真实的 CUDA 内存)
    float h_data = 2.0f;
    float *d_data;
    cudaMalloc(&d_data, sizeof(float));
    cudaMemcpy(d_data, &h_data, sizeof(float), cudaMemcpyHostToDevice);

    // 4. 配置执行参数
    dim3 grid(1);
    dim3 block(1);

    // 5. 实例化模板并启动内核
    // instantiate 参数对应模板参数：<3, float>
    // type_of(*d_data) 用于自动推导类型字符串
    auto error = program.kernel("my_kernel")
           .instantiate(3, jitify::reflection::type_of(*d_data))
           .configure(grid, block)
           .launch(d_data);

    if (error) {
        std::cerr << "Kernel launch failed!" << std::endl;
        return -1;
    }

    // 等待完成并清理
    cudaDeviceSynchronize();
    cudaMemcpy(&h_data, d_data, sizeof(float), cudaMemcpyDeviceToHost);
    std::cout << "Result: " << h_data << std::endl; // 输出应为 8.0 (2^3)
    
    cudaFree(d_data);
    return 0;
}

关键点说明

• 源代码格式：内核代码必须以字符串形式提供，第一行通常是程序名（如 "my_program\n"），用于内部标识。
• 模板实例化：.instantiate() 方法接受模板参数值。使用 jitify::reflection::type_of() 可以方便地将 C++ 类型转换为 Jitify 需要的字符串格式。
• 自动缓存：JitCache 会自动管理编译后的内核，避免重复编译相同的配置。
• 头文件处理：Jitify 会自动忽略主机端代码，并支持常用的 JIT 安全标准头文件（如 float.h, stdint.h 等），无需手动处理复杂的 include 路径。