MultiKernelBench: A Multi-Platform Benchmark for Kernel Generation

1 Introduction

• Existing benchmarks focus mainly on Nvidia, use coarse kernel categories, and omit or undersample important operations such as optimizers.The coarse categories can group tasks with sharply different correctness rates: DeepSeek-V3 scored 0% on CUDA convolution and over 80% on activation.Source: §1 · p. 1
• MultiKernelBench covers Nvidia GPUs, Huawei NPUs, and Google TPUs with 285 tasks across 14 functional categories.It adds missing categories and augments existing ones, with a modular backend abstraction separating platform-specific logic from the core evaluation pipeline.Source: §1 · pp. 1–2
• The paper evaluates seven diverse LLMs using compilation, correctness, and speedup metrics, and compares prompting strategies including category-aware one-shot examples.The default prompt uses the add task as its one-shot example; the category-aware strategy selects an example from the target task’s category.Source: §1 · p. 2
• Category-aware one-shot prompting improves performance, particularly on platforms with limited representation in LLM training data.The evaluation also analyzes platform sensitivity, category-wise difficulty, failures, and cases where models discover optimizations such as kernel fusion and exploiting data sparsity.Source: §1 · p. 2

2.1 Hardware for Deep Learning Computing

• Nvidia GPUs use many arithmetic units and Tensor Cores for mixed-precision matrix operations, while Google TPUs use interconnected multiply-accumulate units for matrix multiplication.The text notes that TPUs are less suited to frequent branching or element-wise operations.Source: §2.1 · p. 2
• Huawei NPUs combine scalar, vector, and cube units for logical control, element-wise operations, and matrix computation, respectively.The three unit types can operate in parallel.Source: §2.1 · p. 2

2.2 Deep Learning Kernels

• DL kernels implement computational primitives, and efficient implementations require attention to memory hierarchy, parallelism, and hardware-specific features.Examples include matrix multiplication, convolution, and activation functions.Source: §2.2 · p. 2
• CUDA exposes GPU threads and memory hierarchy, AscendC provides access to NPU compute units and explicit data movement, and Pallas offers a concise Python interface for TPU kernels.The Pallas example uses JAX transformations and pallas_call; the figure compares vector-addition kernels across the three platforms.Source: §2.2 · p. 2, Figure 1 · p. 3

3.1 Benchmark Overview

• Each task pairs a reference PyTorch module and input tensors with a platform-specific prompt; generated kernel and module code is built, then checked for correctness and performance against the reference.The framework’s workflow is illustrated in Figure 2.Source: §3.1 · p. 3, Figure 2 · p. 4

3.2 Input Prompt and Target Output

• Prompts combine platform-specific instructions and a one-shot example with the target PyTorch module and its input tensors.The example demonstrates platform conventions and required output format; expected outputs differ across CUDA, Pallas, and AscendC.Source: §3.2 · p. 3

3.3 Kernel Tasks

• The taxonomy groups tasks by functional semantics and yields 14 categories, including new Broadcast, Resize, and Optimizer categories and an expanded Indexing category.New tasks are selected from kernel libraries, refined, and manually verified; Table 1 summarizes categories, representative tasks, and task counts.Source: §3.3 · pp. 3–4, Table 1 · p. 4

3.4 Platform-Specific Compilation

• CUDA and Pallas outputs are extracted and executed as Python code, while AscendC components are organized by glue scripts and passed through its compilation pipeline.AscendC requires a more involved build process because it lacks automatic utilities to compile and bind its components directly from Python.Source: §3.4 · p. 4

3.5 Evaluation Approach

• Evaluation measures compilation success, randomized-test correctness, and performance; correctness uses five random inputs with atol and rtol both set to 1e-2.CUDA and AscendC timings use synchronization markers, while Pallas timing uses debugging tools; experiments found the methods stable with controlled variance.Source: §3.5 · pp. 4–5

3.6 Extensible Backend Abstraction for Multi-Platform Evaluation

• A unified Backend interface encapsulates device setup, compilation, execution, correctness and performance checks, and resource cleanup; new backends can be registered without changing core evaluation logic.The framework loads the backend at runtime based on the specified platform name.Source: §3.6 · p. 5

4.1 Methodology

• The methodology evaluates the NVIDIA L20, Huawei Ascend 910B2, and Google TPU v2-8 with seven LLMs and reports compilation, correctness, and speedup metrics at k=1 and k=5.For k=1 the models use greedy decoding; k=5 uses stochastic sampling except for reasoning models.Source: §4.1 · pp. 5–6

4.2 RQ1: Performance in Default Setting

• Under greedy decoding, Claude-Sonnet-4 records 164 of 855 total passed kernels across platforms; sampling improves performance, yet a majority of benchmark tasks remain unsolved.The overall comparison covers 285 DL kernel tasks across CUDA, AscendC, and Pallas, with add as the one-shot example.Source: §4.2 · p. 6, Table 4 · p. 7
• Performance varies by platform: GPT-4o leads CUDA Compilation@1 at 97.5%, while Claude-Sonnet-4 leads Pallas Pass@1 at 8.4% and Pass@5 at 10.5%.DeepSeek-R1 leads CUDA Pass@1 and SpeedUp1@1 but has the lowest Pallas Pass@1, at 2.8%; weaker AscendC and Pallas results are associated with limited training-corpus coverage.Source: §4.2 · p. 6
• AscendC compilation failures indicate platform-specific syntax gaps: 74.8% contain the keyword no member named; Pallas runtime errors most commonly involve Unexpected Keyword Argument, accounting for 22.4%, and incorrect API usage.Source: §4.2 · pp. 7–8

4.3 RQ2: Category-wise Analysis

• Category performance varies substantially: CUDA Pass@1 is highest for Activation at 88.9% and lowest for Convolution at 13.7%, while Reduce and Matrix Multiply reach zero Pass@1 on AscendC and Pallas.Activation is the easiest category across all three platforms; results underscore that strong CUDA categories do not necessarily transfer to other platforms.Source: §4.3 · p. 8
• LLMs sometimes achieve speedups through task-specific sparsity exploitation or kernel fusion, which can reduce redundant computation, memory transfers, and launch overhead.For diagonal matrix multiplication, a generated kernel skips redundant work; in a five-operation Pallas task, the latter four operations are fused.Source: §4.3 · p. 8

4.4 RQ3: Category-Aware One Shot

• The category-aware one-shot study replaces the add exemplar with a representative example from the target category for five categories on each platform.Examples were collected for Activation, Loss, Matrix Multiply, Normalization, and Reduce on AscendC and Pallas.Source: §4.4 · pp. 8–9
• Category-aware examples improve AscendC compilation and correctness, including Matrix Multiply outputs where default prompting yielded none; Pallas Pass@1 gains exceed 100% and span every category.AscendC gains remain incomplete: Reduce errors include mishandled 3D inputs despite a 2D exemplar, and the new Matrix Multiply outputs do not achieve runtime speedup.Source: §4.4 · p. 10

• Added kernels were manually reviewed and their correctness verified through compilation and execution tests.Source: §5 · p. 10
• The authors identify whether findings generalize beyond the selected LLMs and hardware platforms or programming languages as a threat to validity.Source: §5 · p. 10
• Supporting Triton requires fewer than 20 lines of additional code, and the authors plan to add such platforms in future work.Source: §5 · p. 10

6.1 LLM-based Code Generation

• LLM code generation has used continual pretraining and code-focused instruction tuning, while retrieval-based prompting adds code or documentation examples for domain-specific APIs.Source: §6.1 · p. 10

6.2 Benchmarks for Code Generation

• General-purpose code-generation benchmarks progress from function-level unit-test tasks to class- and repository-level tasks involving longer-context reasoning and multi-file consistency.Source: §6.2 · p. 10
• KernelBench [29] and TritonBench [23] assess hardware-efficient kernels on NVIDIA platforms, leaving other hardware platforms outside their coverage.Source: §6.2 · p. 11
• MultiKernelBench supports multiple backends to provide a more comprehensive, platform-agnostic framework for kernel evaluation.Source: §6.2 · p. 11

• MultiKernelBench spans GPUs, NPUs, and TPUs with 285 tasks across 14 functional categories.Source: §7 · p. 11
• Experiments with seven diverse LLMs show substantial performance variability across platforms and kernel categories.Source: §7 · p. 11
• Category-aware prompting significantly improves generation performance on AscendC and Pallas.Source: §7 · p. 11