中文内容
注：BASIC 中的 CUDA Tile 编程是一个愚人节玩笑，但它也是真实的，并且确实可以运行，展示了 CUDA 的灵活性。
CUDA 13.1 引入了 CUDA Tile，这是一种下一代基于 tile 的 GPU 编程范式，旨在让细粒度并行更加易用和灵活。它的一个关键优势是语言开放性：任何编程语言都可以面向 CUDA Tile，使开发者能够将基于 tile 的 GPU 加速带入广泛的生态系统。
为了回应各地资深开发者的巨大需求，我们发布了面向 GPU 的 cuTile BASIC，将 CUDA Tile 编程带到这门长期被忽视的语言中。
什么是 cuTile BASIC？
cuTile BASIC 是 CUDA Tile 编程模型在 BASIC 中的一种表达，构建于 CUDA Tile IR 规范之上。它使你能够使用基于 tile 的模型在 BASIC 中编写 tile kernel；对于早于多线程编程出现的 BASIC 这样的编程语言来说，这是一种自然契合。
cuTile BASIC 是 GPU 的强大能力与 BASIC 编程语言过时魅力和语法简洁性的完美结合——这是一门来自更像素化时代的优雅语言。手动给代码行编号从未如此好看，也从未运行得如此之快！
cuTile BASIC 面向谁？
BASIC 是现存最古老的编程语言之一，因此受到整整一代开发者的推崇，他们会怀念 300 波特拨号调制解调器握手的声音。对于许多这样的开发者来说，BASIC 是他们第一次接触计算机编程。
现在，那些脑海中仍深深烙着 BASIC 的开发者，首次可以把遗留应用带到 NVIDIA GPU 加速计算上。这将释放 BASIC 编程语言此前无法想象的性能和功能——让你的 Lunar Lander 以比 Artemis 任务更快的速度在月球表面飞驰。
在世界上最强大的 GPU 上运行时，重温数学课上为图形计算器编写游戏的辉煌时光。
开始设置
首先，使用 PIP 安装 cuTile BASIC：
pip install git+https://github.com/nvidia/cuda-tile.git@basic-experimental
运行 cuTile BASIC 的完整硬件和软件要求列在本文末尾（建议使用 64k 或更多 RAM）。
cuTile BASIC 示例
如果你学过 CUDA C++，可能见过经典的向量加法 kernel。CUDA C++ 中的向量加法 kernel 大致如下，它接收两个向量，并将它们按元素相加以生成第三个向量。
这是可以编写的最简单 CUDA kernel 之一：
__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
/* calculate my thread index */
int workIndex = threadIdx.x + blockIdx.x*blockDim.x;
if(workIndex < vectorLength)
{
/* perform the vector addition */
C[workIndex] = A[workIndex] + B[workIndex];
}
}
在这个 kernel 中，每个线程的工作都被明确指定，程序员在启动此 kernel 时会指定要启动的 block 和 thread 数量。
现在让我们看看用 cuTile BASIC 编写的等价代码。我们不需要指定每个线程做什么。我们只需要把数据划分成 tile，并指定这些 tile 应执行哪些数学运算。其他一切都由系统为我们处理。
cuTile BASIC 向量加法 kernel 如下所示：
10 REM Vector Add: C = A + B 20 INPUT N, A(), B() 30 DIM A(N), B(N), C(N) 40 TILE A(128), B(128), C(128) 50 LET C(BID) = A(BID) + B(BID) 60 OUTPUT C 70 END
这个示例非常基础，使用标准 BASIC，并有三点额外内容需要指出：
- 索引数组会返回一个 tile，即数组的一个子集。
- BID 是一个内置变量，用于指定 tile block 索引。
- TILE 指定数组应被划分成多大尺寸的 tile。
注意，除了加法操作外，我们无需指定任何内容。其他一切都由 cuTile BASIC 处理。
整合起来
现在我们将展示如何在 BASIC 中运行这个向量加法 kernel。直接的工作流程是，首先将 BASIC 函数编译为 cubin，然后在 GPU 上启动。为简洁起见，我们省略了无聊的 Python 主机端和封装代码，但你可以在我们的 GitHub 仓库中找到它。
如果你安装了合适版本的 CUDA Toolkit 和 Python，并已从 GitHub 下载 cuTile BASIC 仓库，就可以执行以下命令：
$ python examples/vector_add.py
[1/2] Compiling to cubin ...
Arrays: ['A', 'B', 'C'], tile_shapes={'A': [128], 'B': [128], 'C': [128]}, grid_size=8
[2/2] Launching kernel on GPU ...
Results (showing 5 samples of 1024):
C[ 0] = 0.0 (expected 0.0)
C[ 1] = 3.0 (expected 3.0)
C[ 511] = 1533.0 (expected 1533.0)
C[ 512] = 1536.0 (expected 1536.0)
C[1023] = 3069.0 (expected 3069.0)
VERIFICATION PASSED (max_diff=0.000000, 1024 elements)
如果你的输出看起来相同，恭喜，你刚刚运行了你的第一个 cuTile BASIC 程序，而且很可能也是你用 BASIC 编写的第一个任何类型的程序！Max Headroom 会为你感到骄傲。
一个 BASIC 矩阵乘法
BASIC 是一门简单的语言，只需很少几行代码就能表达常见算法。请看下面用 BASIC 编写的矩阵乘法（GEMM）kernel：
10 REM GEMM: C(M,N) = A(M,K) * B(K,N) 15 INPUT M, N, K, A(), B() 20 DIM A(M, K), B(K, N), C(M, N) 30 TILE A(128, 32), B(32, 128), C(128, 128), ACC(128, 128) 40 LET TILEM = INT(BID / INT(N / 128)) 50 LET TILEN = BID MOD INT(N / 128) 60 LET ACC = 0.0 70 FOR KI = 0 TO INT(K / 32) - 1 80 LET ACC = MMA(A(TILEM, KI), B(KI, TILEN), ACC) 90 NEXT KI 100 LET C(TILEM, TILEN) = ACC 110 OUTPUT C 120 END
在这个 kernel 中，除了标准 BASIC 语法外，TILE 指定 A、B 和 C 应如何分块，以及累加器 tile ACC 的大小。MMA 是矩阵乘加的函数调用。注意这段代码有多简单。你指定数据应如何细分为 tile，并以高层次指定算法，而在底层，CUDA Tile 会处理其他一切。
这个示例也可在 GitHub 仓库的 examples 文件夹中找到。运行它会产生以下输出：
$ python examples/gemm.py
[1/2] Compiling to cubin ...
M=512, N=512, K=512, tile_shapes={'A': [128, 32], 'B': [32, 128], 'C': [128, 128]}, grid_size=16
[2/2] Launching kernel on GPU ...
Results (showing 5 samples of 512x512 = 262144 elements):
C[0,0] = -0.1199 (expected -0.1199)
C[0,1] = -14.4456 (expected -14.4456)
C[256,0] = -15.8891 (expected -15.8891)
C[256,1] = -2.8646 (expected -2.8646)
C[511,511] = 11.4724 (expected 11.4724)
VERIFICATION PASSED (max_diff=0.000012, tol=0.005120)
如上所示的矩阵乘法是大型语言模型等人工智能工具的核心。有了 cuTile Basic，开发者现在可以从一门几乎无法想象整整 1 MB 系统内存的语言出发，探索拥有数万亿参数模型的人工智能前沿。
开发者如何获取 cuTile
要运行 cuTile BASIC 程序，你需要以下内容：
- 具有 8.x、10.x、11.x 或 12.x 计算能力的 GPU（在未来 CUDA 版本中，我们将增加对更多 GPU 架构的支持）
- NVIDIA Driver R580 或更高版本（tile 专用开发者工具支持需要 R590）
- CUDA Toolkit 13.1 或更高版本
- Python 版本 3.10 或更高
- 正文：cuTile BASIC package
开始使用
准备好所有软件后，请查看完整的 cuTile BASIC 文档，尝试 GitHub 上的所有示例程序，并从今天开始使用 cuTile BASIC 编程。享受将你的现代 AI 或科学计算代码库移植到一门具有历史关键地位的语言中的机会，同时仍能在现有最强大的硬件上运行！只是别告诉你的 Commodore 64。
任何语言中的 CUDA Tile
虽然 BASIC 可能不是开发者在想到高性能并行计算时首先想到的语言，但它是一个有启发性的演示：得益于 CUDA 软件栈的设计，CUDA Tile 几乎可以从任何编程语言中使用。通过编译为 CUDA Tile IR 格式，CUDA Tile 可以被带到几乎任何语言中……甚至 BASIC！
编者注：回想起来，要求广泛支持 CUDA Tile 编程模型的开发者也许本该说得更具体一些。敬请期待 2027 年 4 月 1 日推出的 cuTile COBOL。
标签
