TornadoVM 7.0：TileContext 走 CUDA Tile，KernelContext 与 library tasks 同图协作

TornadoVM 7.0（2026-09-22）把 NVIDIA CUDA Tile（cuTile）接到 Java：方法首参换成 TileContext 后，按 tile 写 kernel，由 tile 编译器选线程映射与 Tensor Core 指令。文档侧原先的 Kernel API 已改称 KernelContext API；7.0.1（2026-09-24）把 CUDA backend（含 cuTile 编译路径）以及 cuBLAS、cuDNN 等 library bindings 发上 Maven Central。站内此前没有 TornadoVM 条目，下面按 GitHub Release、README、programming guide 与 tile-api 文档核对后写。

7.0 / 7.0.1 事实表

项	内容
仓库	beehive-lab/TornadoVM
版本	v7.0.0（2026-09-22）、跟进 v7.0.1（2026-09-24）
JDK	21–27；预构建 SDK 分 `jdk21` 与 `jdk22plus` 两族
许可证	Tornado-API、Examples、Annotation 等为 Apache 2.0；Tornado-Runtime、Tornado-Drivers 为 GPLv2 with Classpath Exception（README「Licenses per module」）
SDK 分发	按加速后端拆包：`opencl`、`cuda`、`metal`、`full`；GitHub Release assets 与 SDKMAN! 候选版本同名
Maven	`io.github.beehive-lab:tornado-api` / `tornado-runtime`，版本 `7.0.1-jdk21` 或 `7.0.1-jdk22plus`；7.0.1 起另有 `tornado-drivers-cuda` 与 `tornado-cublas`、`tornado-cudnn`、`tornado-cufft`、`tornado-curand`、`tornado-cusparse`、`tornado-cudf`、`tornado-cutlass`（Central 已有上述版本）
官网	tornadovm.org

7.0.0 里和选题直接相关的条目：#1083 引入 Java CUDA Tile 执行路径（TileContext API）；#1112 在 programming guide 文档化 TileContext，并把 Kernel API 章节改名为 KernelContext API；#1110 增加 cuDF library-task；#1085 增加 Method-based task API，便于运行时生成的方法也进图。半精度读写、MMA accumulator、Metal half、寄存器过重 1D kernel 等修复一并进了 7.0。recover.bailout 在 7.0 改为默认关闭（#1107）。

7.0.1 的 #1116 补上 Central 缺口：此前 Central 主要是 tornado-api、tornado-runtime、tornado-matrices，依赖 CUDA library task 或 CUDA backend 的工程（例如 jitllm）在 7.0.0 上解析失败。合并后按 JDK 族发布 tornado-drivers-cuda（含 cuTile 编译路径）和上述 library bindings；cudnn-jni / cudf-jni 等 CMake 原生 shim 仍只进 SDK，Central 上是纯 Java、运行时用 java.lang.foreign 加载 NVIDIA 库。

jdk21 构建带 --enable-preview，类文件钉在 JDK 21；jdk22plus 面向 22 及以上（含 27），换同族 JDK 通常不用重装 SDK。Maven 坐标后缀必须与所用 SDK 族一致。

Release 资产命名大致是 tornadovm-<ver>-<jdk族>-<backend>-<os>-<arch>.zip|tar.gz。Linux amd64 上常见 opencl / cuda / full；Windows 有 opencl 与 cuda；macOS aarch64 有 opencl 与 metal。选包时先定 JDK 族，再定后端，避免装了 opencl 包却去跑 TileContext（tile 只在 cuda 包 / full 包里的 CUDA 路径可用）。

TileContext、KernelContext，以及和原先路径的关系

TornadoVM 挂在现有 OpenJDK 或 GraalVM 上，把选定的 Java 方法 JIT 成 CUDA / OpenCL / Metal，自己管理主机与设备间的缓冲迁移。并行怎么写，由方法签名与注解决定，三种写法共用 TaskGraph 与 TornadoExecutionPlan，可以混在一张图里。

Loop Parallel（@Parallel） 在循环上标注并行维，运行时推断全局与局部工作大小。同一方法也能当普通顺序 Java 跑，适合先写通再上设备。

KernelContext 要求方法首参是 KernelContext。写法接近 CUDA / OpenCL：用 globalIdx / localIdx、分配 local（shared）内存、插 barrier，也可在 CUDA 上走 mma.sync 一类 Tensor Core 内建。JIT 仍走 SIMT 路径（NVIDIA 上大致是 Graal IR → PTX → NVRTC → cubin）。7.0 文档把这块从「Kernel API」统一叫成 KernelContext API，类名与用法本身是延续，主要是命名与章节对齐。

TileContext（7.0 新增，仅 CUDA backend） 要求方法首参是 TileContext。你写的是一块 tile 上的运算：view 出张量视图、partition 切块、load / store、mma / matmul、归约与 elementwise 等。API 里不出现 thread、warp、fragment。编译走 nvcc -tilecubin --tile-only，由 tile 编译器决定背后用多少线程、发哪条 Tensor Core 指令、shared memory 怎么铺。tile 边长必须是编译期常量且为 2 的幂（字面量或 static final int）；问题规模（extent）可以是运行时参数，不会特化整颗 kernel。WorkerGrid 计的是 tile block 数，不要设 local work（CUDATileScheduler 钉成 1x1x1）。

硬件门槛（官方 tile-api）：CUDA Toolkit 13.3+、驱动 R580+、算力 8.0+。缺任一条件时抛 TornadoDeviceTileNotSupported。每一条 TileContext 操作在 JVM 上也有实现，同一方法可以在 CPU 上调试；recover.bailout 默认关闭，避免编译失败时静默退回主机却看起来「算对了」。

关系上，TileContext 是第三条编译路径，并不取代 KernelContext。需要显式 local memory、barrier、按线程写算法，或从现有 CUDA/OpenCL 思路平移时，继续用 KernelContext；愿意把工作单位抬到 tile、少碰线程映射时，用 TileContext。README 与 unittests（如 TestTileChaining）强调：@Parallel、KernelContext、TileContext 与 native library task 可共享设备缓冲与同一 CUDA stream，并可 withCUDAGraph() 整图捕获后一次 launch 重放。

选型上不必追求「全部改成 tile」。很多预处理、后处理、不规则索引仍更适合 @Parallel 或 KernelContext；GEMM 类热点若已有 cuBLAS，直接 libraryTask 往往更省事。TileContext 的价值集中在：想用 Java 表达 tile 级算法，又希望由编译器去挑 Tensor Core 与 shared 布局，而不是手写 mma.sync 形状。

Java GPU kernel 与 hybrid library tasks 怎么分工

自己写的 Java kernel（三种 API 任一）适合应用特有逻辑：布局转换、融合激活、量化反量化、库里没有或要对中间张量动手的那段。vendor 库已经调得很满的稠密算子（GEMM、FFT、部分深度学习原语、RAPIDS 侧关系算子等），用 libraryTask 直接绑原生入口，和 JIT kernel 共用 TornadoVM 管的设备缓冲，中间结果可以不回主机、也不必额外做主机侧同步。

官网 hybrid 形态如下（符号与 README 一致）：

TaskGraph tg = new TaskGraph("hybrid")
    .transferToDevice(DataTransferMode.EVERY_EXECUTION, matrix, vector)
    .task("preprocess", MyKernels::preprocess, matrix)   // JIT 编译的 Java kernel
    .libraryTask("gemv", CuBlas::cublasSgemv,
            CUBLAS_OP_T, m, n, alpha, matrix, lda, vector, incx, beta, output, incy)
    .task("postprocess", MyKernels::activate, output)
    .transferToHost(DataTransferMode.EVERY_EXECUTION, output);

try (TornadoExecutionPlan plan = new TornadoExecutionPlan(tg.snapshot())) {
    plan.withCUDAGraph().execute();
}

README 当前列出的库侧入口包括 cuBLAS / cuBLASLt、cuFFT、cuDNN、cuDF；KernelContext 另可直接发 FP16 / INT8 的 mma.sync。实务上可以这样选：定制与融合前处理写 Java kernel；能对上 cuBLAS / cuDNN 等的，优先 libraryTask；希望在 tile 层表达 matmul、softmax、attention，并让编译器选 Tensor Core 指令时，用 TileContext（仍可与 libraryTask 同图）。官网首页有 SGEMM / SGEMV 在特定 GPU 与 CUDA 版本上的对比数字，那是官方演示环境的结果，本稿不转述为普适结论；换卡、换精度、换图结构后差异会很大，以本机 tornado 示例与 profiler 为准。

和「纯绑定」路线的差别也值得分清：TornadoVM 的 JIT kernel 是从 Java 字节码生成设备代码；library task 才是对已有 NVIDIA 库的绑定。两者可以同图，但责任不同：前者你要保证可并行语义与 TornadoVM 支持的语言子集，后者你要保证本机装了对应 CUDA 库且参数约定正确。

最小上手

先确认 JAVA_HOME 指向目标 JDK，再装与族、后端匹配的 SDK。SDKMAN! 示例：

sdk install tornadovm 7.0.1-jdk21-cuda
# OpenCL 默认包：sdk install tornadovm
# Apple Silicon Metal：sdk install tornadovm 7.0.1-jdk21-metal
tornado --devices
tornado --version

也可从 v7.0.1 Release 下载 zip/tar（例如 tornadovm-7.0.1-jdk21-opencl-linux-amd64.zip），解压后设 TORNADOVM_HOME 并把 $TORNADOVM_HOME/bin 加入 PATH。官网还提供 sdk install tornadovm 与 downloads 页。

Maven 侧至少需要 API 与 runtime（版本后缀跟 SDK 族走）：

<dependency>
  <groupId>io.github.beehive-lab</groupId>
  <artifactId>tornado-api</artifactId>
  <version>7.0.1-jdk21</version>
</dependency>
<dependency>
  <groupId>io.github.beehive-lab</groupId>
  <artifactId>tornado-runtime</artifactId>
  <version>7.0.1-jdk21</version>
</dependency>

仅加 API jar 不够跑设备；实际执行依赖本机驱动与对应后端的 SDK / driver 模块。编译期要用 CUDA library task 或 CUDA backend（含 cuTile）时，再引入 tornado-cublas 等与 tornado-drivers-cuda（版本同上）。跑通后可用官方示例确认：

java @$TORNADOVM_HOME/tornado-argfile \
  -cp $TORNADOVM_HOME/share/java/tornado/tornado-examples-7.0.1.jar \
  uk.ac.manchester.tornado.examples.compute.MatrixVectorRowMajor

TileContext 极简形态（摘自 programming / tile-api，略去完整 main）：

public static void gemm(TileContext tc, HalfFloatArray a, HalfFloatArray b, FloatArray c,
                        int m, int n, int k) {
    PartitionView av = tc.partition(tc.view(a, m, k), 64, 32);
    PartitionView bv = tc.partition(tc.view(b, k, n), 32, 64);
    PartitionView cv = tc.partition(tc.view(c, m, n), 64, 64);

    Tile acc = tc.zeros(DType.F32, 64, 64);
    for (int step = 0; step < k / 32; step++) {
        acc = tc.mma(av.load(tc.bidX(), step), bv.load(step, tc.bidY()), acc);
    }
    cv.store(acc, tc.bidX(), tc.bidY());
}

WorkerGrid2D worker = new WorkerGrid2D(m / 64, n / 64); // tile blocks
GridScheduler grid = new GridScheduler("s0.gemm", worker);

tornado-examples 下 uk.ac.manchester.tornado.examples.tile 有 TileMatrixMultiply、TileSoftmax、TileAttention 等可跑样例。不确定生成结果时，用 tornado --printKernel 看输出的 CUDA Tile C++（应看到 ct:: 调用，而不是手写内联 PTX）。

适用边界

OpenCL backend 覆盖 NVIDIA / AMD / Intel 等与多核 CPU；Metal 面向 Apple Silicon；TileContext 与 NVIDIA library tasks 只在 CUDA backend。Tile 额外要求 Toolkit 13.3+、R580+、CC 8.0+。数据请用 FloatArray、HalfFloatArray 等 off-heap 类型（基于 Panama MemorySegment），并在图上显式声明 transferToDevice / transferToHost。

官方 known limitations 里几条对混图有影响：tile 任务暂不支持 withBatch；@Reduce 不宜与 tile 同图（reduction 改写会改图名，导致 GridScheduler 对不上，tile 可能静默按单 block 跑）；部分 CUDA Tile C++ 表面（masked atomics、atomic_compare_exchange、permute 等）尚未暴露。调试时保持 recover.bailout 关闭，并用 --printKernel 确认设备侧确实生成了 kernel。

业务代码通常只链接 Apache 2.0 的 API 模块；打包或分发 Runtime / Drivers 时按仓库各模块的 GPLv2+CE 自行核对。TornadoVM 不替代 JVM：未放入 TaskGraph 的代码仍按普通 Java 执行；放入图中的方法则受加速器与后端能力约束（例如 Metal 上没有 TileContext，也没有 cuBLAS library task）。

若目标是多厂商可移植，优先 @Parallel / KernelContext + OpenCL（或按平台选 Metal）；若目标是 NVIDIA 上把 Java 业务与 CUDA 生态拼在一张图里，再考虑 cuda SDK、libraryTask 与 TileContext。更细的操作表见 Core Programming 与仓库 docs/source/tile-api.rst。版本号、资产名与 Maven 坐标以 v7.0.0、v7.0.1 与 Maven Central 为准。

— 感谢阅读 —

一起交流

分享你的思考，让讨论更进一步。