在 TileLang 中T.Kernel、T.Parallel和T.Pipelined是构建层次化 GPU/NPU 并行与计算重叠的三大核心控制原语。它们分别对应Grid 级网格分配、Block 内线程级并行映射以及片上软件流水线掩盖访存延迟。1. T.KernelGPU 线程块网格入口Grid LaunchT.Kernel用于定义 Kernel 的执行网格Grid以及每个 Thread Block 的线程数量。它明确了算子在硬件最高层级的并行粒度。核心作用声明 GPU Grid 的维度绑定 Block ID并指定每个 Block 内分配的线程数threads。典型语法withT.Kernel(grid_x,grid_y,threads128)as(bx,by):# bx, by 相当于 CUDA 中的 blockIdx.x, blockIdx.y# 此处编写当前 Block 负责处理的 Tile 逻辑物理映射直接映射到 CUDA 的grid, threadsLaunch 配置。在with T.Kernel作用域内所有的片上资源如T.alloc_shared都是以 Block 为单位独立开辟的。2. T.ParallelTile 内部的数据并行映射Parallel LoopT.Parallel用来表示 Block 内部针对数据 Tile 的并行循环。它取代了传统 CUDA 中手写threadIdx.x索引偏移的繁琐逻辑。核心作用向编译器声明“该循环中的所有迭代互相独立可以并发执行”。TileLang 编译器会自动将循环变量均匀映射到当前 Block 内分配的线程Threads/Warps上。典型语法# 将 [Tile_M, Tile_N] 的元素初始化编译器自动分配 Thread 去并行写 0fori,jinT.Parallel(Tile_M,Tile_N):Accumulator_fragment[i,j]0.0物理映射编译器根据T.Kernel中定义的threads数量自动对T.Parallel的迭代空间Iteration Space做 Vectorization向量化和 Thread Index Binding线程索引绑定。3. T.Pipelined软件流水线重叠Hide Memory LatencyT.Pipelined是 TileLang 榨干硬件性能的核心利器用于在主 Loop如 GEMM 中的 K 轴循环中开启多级软件流水线Software Pipelining。核心作用自动将“从 Global Memory 搬运数据到 Shared MemoryT.copy”与“在 Tensor Core 上做计算T.gemm”在时间轴上异步重叠起来利用异步拷贝指令如 CUDAcp.async/ TMA掩盖高延迟的内存访问。典型语法# num_stages3 表示开辟 3 级多重缓冲区 (Triple Buffering)forkinT.Pipelined(K_blocks,num_stages3):T.copy(A[bx*Tile_M,k*Tile_K],A_shared)T.copy(B[k*Tile_K,by*Tile_N],B_shared)T.gemm(A_shared,B_shared,C_fragment)物理映射编译器会自动生成双/多缓冲区Double/Triple Buffer将上一步的 GEMM 计算与下一步的 Async Copy 安排在同一个时钟周期内重叠运行。三者协作逻辑对照原语控制层级对应 CUDA / 硬件概念开发者解决的核心问题T.KernelGrid 级blockIdx,blockDim决定整体任务怎么拆给不同的 SM / BlockT.ParallelBlock/Tile 级threadIdx, Thread 协同避免手写 Thread 索引推导自动分配并行工作T.PipelinedInstruction 级cp.async, 多重缓冲区 (Double Buffer)隐藏 Global Memory 访存延迟让 Tensor Core 保持满载
在 TileLang 中T.Kernel、T.Parallel和T.Pipelined是构建层次化 GPU/NPU 并行与计算重叠的三大核心控制原语。它们分别对应Grid 级网格分配、Block 内线程级并行映射以及片上软件流水线掩盖访存延迟。1. T.KernelGPU 线程块网格入口Grid LaunchT.Kernel用于定义 Kernel 的执行网格Grid以及每个 Thread Block 的线程数量。它明确了算子在硬件最高层级的并行粒度。核心作用声明 GPU Grid 的维度绑定 Block ID并指定每个 Block 内分配的线程数threads。典型语法withT.Kernel(grid_x,grid_y,threads128)as(bx,by):# bx, by 相当于 CUDA 中的 blockIdx.x, blockIdx.y# 此处编写当前 Block 负责处理的 Tile 逻辑物理映射直接映射到 CUDA 的grid, threadsLaunch 配置。在with T.Kernel作用域内所有的片上资源如T.alloc_shared都是以 Block 为单位独立开辟的。2. T.ParallelTile 内部的数据并行映射Parallel LoopT.Parallel用来表示 Block 内部针对数据 Tile 的并行循环。它取代了传统 CUDA 中手写threadIdx.x索引偏移的繁琐逻辑。核心作用向编译器声明“该循环中的所有迭代互相独立可以并发执行”。TileLang 编译器会自动将循环变量均匀映射到当前 Block 内分配的线程Threads/Warps上。典型语法# 将 [Tile_M, Tile_N] 的元素初始化编译器自动分配 Thread 去并行写 0fori,jinT.Parallel(Tile_M,Tile_N):Accumulator_fragment[i,j]0.0物理映射编译器根据T.Kernel中定义的threads数量自动对T.Parallel的迭代空间Iteration Space做 Vectorization向量化和 Thread Index Binding线程索引绑定。3. T.Pipelined软件流水线重叠Hide Memory LatencyT.Pipelined是 TileLang 榨干硬件性能的核心利器用于在主 Loop如 GEMM 中的 K 轴循环中开启多级软件流水线Software Pipelining。核心作用自动将“从 Global Memory 搬运数据到 Shared MemoryT.copy”与“在 Tensor Core 上做计算T.gemm”在时间轴上异步重叠起来利用异步拷贝指令如 CUDAcp.async/ TMA掩盖高延迟的内存访问。典型语法# num_stages3 表示开辟 3 级多重缓冲区 (Triple Buffering)forkinT.Pipelined(K_blocks,num_stages3):T.copy(A[bx*Tile_M,k*Tile_K],A_shared)T.copy(B[k*Tile_K,by*Tile_N],B_shared)T.gemm(A_shared,B_shared,C_fragment)物理映射编译器会自动生成双/多缓冲区Double/Triple Buffer将上一步的 GEMM 计算与下一步的 Async Copy 安排在同一个时钟周期内重叠运行。三者协作逻辑对照原语控制层级对应 CUDA / 硬件概念开发者解决的核心问题T.KernelGrid 级blockIdx,blockDim决定整体任务怎么拆给不同的 SM / BlockT.ParallelBlock/Tile 级threadIdx, Thread 协同避免手写 Thread 索引推导自动分配并行工作T.PipelinedInstruction 级cp.async, 多重缓冲区 (Double Buffer)隐藏 Global Memory 访存延迟让 Tensor Core 保持满载
