GPU

Kernel Paradigm triton 的 kernel 有两种写法: 传统 kernel Persistent Kernel 传统 Kernel grid = (num_tiles,) 每个 block 处理一个 tile → 结束 每个 SM 从 grid 里领一个 work item，做完就退出, CTA 调度由硬件自动完成 Persistent Kernel @triton.jit def kernel(..., n_tiles: tl.constexpr, ...): pid = tl.program_id(0) n_progs = tl.num_programs(0) # grid(0) 启动的 program 数 tile = pid while tile < n_tiles: # 处理 tile 对应的那一块工作 # ... tile += n_progs # 跳到下一个属于自己的 tile 每个 CTA 长期驻留在 SM 上，自己软件调度 work persistent kernel 常见动机： 任务总 tile 数不够多（比如小 batch、小矩阵、很多小 GEMM / MoE / grouped GEMM）:普通写法 programs 数 < SM，GPU 吃不满；persistent 让每个 SM 都有活干，并通过循环把剩余工作吃完。 ...

Background GPU Memory Model SRAM (Static RAM) Located inside the GPU core, it utilizes Registers, L1 Cache, and L2 Cache: Registers — These are tiny, ultra-fast memory locations within each GPU core. Registers store immediate values that a core is actively processing, making them the fastest type of memory. L1 Cache — This is the first-level cache inside a Streaming Multiprocessor (SM). It stores frequently accessed data to speed up calculations and reduce access to slower memory (like DRAM). L2 Cache — This is a larger, second-level cache that is shared across multiple SMs. It helps store and reuse data that might not fit in L1 cache, reducing reliance on external memory (VRAM). HBM Bigger capacity ...

Extensive Reading Author Info ‪Le Chen‬ - ‪Google Scholar‬ Haibo Chen [IPADS]: Director of Institute of Parallel and Distributed Systems Background 现有的LLM推理引擎通常只使用其中一种加速器（例如只用GPU或只用NPU），这导致了两个主要问题： 资源浪费：无法充分利用芯片上所有计算单元的算力。 性能瓶颈：单一加速器有其固有的性能短板，无法在所有场景下都达到最优性能。 Challenges Insights 设计一个能够同时、高效地利用 GPU 和 NPU 进行协同计算的 LLM 推理引擎，以最大限度地提升移动设备上的 LLM 推理速度 The NPU serves as the primary computing unit, handling the majority of computing tasks, while the GPU acts as a secondary computing unit to enhance the lower bound of NPU performance. GPU Characteristics Linear Performance: The GPU’s performance scales linearly as the tensor size increases. It transitions from being memory-bound on small tensors to compute-bound on large ones, where its performance plateaus. High-Cost Synchronization: There are two main types of synchronization overheads. Data Copy: API calls to transfer data between CPU and GPU buffers, like clEnqueueWriteBuffer, incur a fixed latency of about 400 microseconds, irrespective of the data’s size. Kernel Submission: Submitting a kernel to an active, non-empty GPU queue has a negligible overhead (10-20 microseconds). However, after a synchronization event empties the queue, submitting the next kernel incurs a much higher “startup” latency of 50-100 microseconds. ...