Towards Robust Agentic CUDA Kernel Benchmarking, Verification, and Optimization

Agentic E2E Workflow

• The proposed workflow translates PyTorch operations into CUDA kernels, then samples, verifies, tests, and evaluates candidate kernels in parallel during runtime optimization.The process uses evolutionary optimization, model ensembling, in-context improvement, profiling summaries, and automatic operation fusion.Source: Figure 1 · p. 2, §1 · p. 2

Motivation

• The benchmark exposes exploits including omitted redundant operations, overfit input settings, and kernels that fail to generalize to practical applications.These loopholes can produce artificial speedup estimates in existing evaluations.Source: §1 · p. 2

Robust Benchmark Harness

• The authors introduce robust-kbench to evaluate CUDA kernel correctness and optimization across varied settings, including forward and backward operations and realistic downstream applications.The benchmark addresses artificial speedups caused by poor task design and narrow verification.Source: §1 · p. 2

Soft-Verification

• LLM-based soft verification classifies compilation, memory-access, and numerical-correctness failures with up to 80% accuracy, increasing downstream hardware-verification success by up to 30%.The workflow supports efficient filtering and test-time scaling.Source: §1 · p. 2

Pipeline Analysis and Open Resources

• The paper analyzes pipeline consistency and performance through ablations of LLM ensembling, iterative profiling-summary feedback, and soft-verification filtering.The authors also release the benchmark, discovered-kernel dataset, profiling data, and self-verification results for future model training.Source: §1 · p. 2

Scaling Test-Time Compute

• The paper scales test-time compute through multiple sequential sample streams, evolutionary recombination, self-verification, and kernel accuracy, speed, and profiling signals.Test-time strategies include voting, debate, evolutionary mutation and recombination, sequential correction, and parametric or non-parametric verification [18] [55] [51] [57] [14] [37] [32] [64].Source: §2 · p. 3

Evolutionary Code Optimization with LLMs

• Evolutionary code optimization reuses, mutates, and recombines prior code, with LLMs acting as recombination engines; retrieval-augmented generation offers a simpler history-retrieval variant [38] [36] [28] [26] [40] [15].These methods have been applied to reward and preference objectives, scientific code, machine-learning papers, and other applications.Source: §2 · p. 3

The CUDA Framework

• CUDA provides GPU-specific C/C++ extensions, fine-grained control over threads, memory, and synchronization, and a hierarchy of warps, blocks, and grids for scalable parallel execution.Its ecosystem also includes cuBLAS, cuDNN, and Thrust libraries [12].Source: §2 · p. 3

Challenges for Benchmarking LLM-Written CUDA Kernels

• KernelBench v0 contains 250 PyTorch-defined neural-network tasks across three levels, but flawed baselines, precision-sensitive outputs, seed-insensitive variation, and CPU profiling overhead can distort evaluation.The levels progress from individual operators to fused operators and larger neural networks; KernelBench remains the de facto standard.Source: §2 · p. 3

Challenges for LLM-Driven Optimization of CUDA Kernels

• LLM-generated CUDA kernels can exploit benchmark loopholes and often fail to generalize across input shapes, precision requirements, and production frameworks.Reported exploits include removing redundant operations, hardcoding input patterns, and assuming weights that do not generalize.Source: §2 · p. 3

Benchmark limitations

• 3.13x average speedup falls to 1.49x after excluding contaminated tasks, revealing that apparent gains over Baronio et al. largely exploit KernelBench loopholes.Some kernels pass verification while producing fake 50-120x speedups; Figure 2 compares full and contamination-excluded results.Source: §3 · p. 4, Figure 2 · p. 4

Robust evaluation framework

• robust-kbench combines diverse initialization states, multiple runtime estimators, profiling tools, and validation layers to prevent hardcoding and assess correctness across realistic executions.It integrates PyTorch Profiler, Clang-tidy, and NVIDIA Compute Profiler (NCU) for high-level, static, and hardware-level analysis.Source: §3 · p. 4

Benchmark scope and interface

• The benchmark evaluates forward and backward passes across MNIST CNN training, ResNet-18 inference, and Transformer Llama inference under multiple initialization states and input configurations.Its evaluator compiles kernels, verifies correctness, profiles performance, and supports parallel execution across multiple GPUs through a Python API.Source: §3 · pp. 4–5

• Traditional hardware verification requires at least one minute per kernel and is constrained by available GPU hardware, creating a bottleneck for evaluating many variants or optimization searches.Robust verification is also needed because kernel proposals may exploit benchmark tasks.Source: §4 · p. 5
• The verification workflow uses majority-consensus LLM verifiers to filter kernel proposals before hardware testing, enabling parallelized soft verification and faster evaluation.A meta-agent iteratively tunes system and message prompts using balanced datasets of correct and incorrect kernels, self-reflection, and error summaries.Source: §4 · p. 5
• 0.82 accuracy for compilation, 0.80 for memory access, and 0.73 for numerical correctness were achieved by specialized verifiers tuned on balanced datasets of 30 kernels.The tuned prompts successfully evaluated ∼20 previously unseen kernels and maintained performance across different LLM base models.Source: §4 · p. 5, Figure 3 · p. 5

Kernel Translation & Optimization

• The translation loop compiles LLM-generated CUDA, checks results against PyTorch within 1e−5 precision, summarizes failures, and feeds errors back for iteration.This procedure improves robustness over best-of-N sampling under the same compute budget.Source: §5 · p. 6
• Evolutionary optimization filters correct kernels, prompts an LLM ensemble with up to five prior examples and recommendations, then profiles four proposals on individual H100 GPUs.The process samples eight proposals per generation and evaluates torch, NCU, and Clang-tidy profiling information using CUDA 12.4.Source: §5 · p. 6

Results: Optimizing individual operations

• Up to 2.5× forward-pass speedups over PyTorch eager emerge across robust-kbench tasks; verifier feedback improves stability, while backward-pass gains are smaller and harder to obtain.Input-shape prompts can further improve performance, and backward optimization is more challenging because of kernel fusion and activation recomputation complexity.Source: §5 · p. 6, Figure 4 · p. 7

Results: Generalization of LLM-Optimized Kernels

• Optimized kernels overfit unseen shapes for LayerNorm and MNIST Linear-ReLU, whereas ResNet kernels preserve performance benefits across input dimensions.This evaluation tests configurations not encountered during evolutionary optimization, exposing generalization differences that fixed-shape benchmarks may miss.Source: §5 · p. 7, Figure 5 · p. 7

Impact of LLM Verification

• Verifier assistance raises valid forward-pass kernels from 55–70% to 80–85% by filtering compilation, memory-access, and numerical errors before hardware evaluation.The pre-screening mechanism particularly filters prevalent compilation failures, reducing wasted hardware-testing computation.Source: §6 · p. 8, Figure 6 · p. 8

Impact of Model Ensembling

• Increasing model-ensemble diversity improves optimization outcomes and the success rate of generating valid, performant kernels.The comparison spans GPT-4.1 alone, a GPT-4.1–Claude Sonnet 3.7 ensemble, and a five-model ensemble.Source: §6 · p. 8

Impact of Context Construction

• Least-to-most context construction using the five best-performing kernels yields the most consistent optimization improvements among the evaluated strategies.The alternatives provide only the single best kernel or ten randomly selected previous kernels.Source: §6 · p. 8, Figure 7 · p. 8

Impact of Additional Profiling Information

• Adding LLM-summarized torch, NCU, and Clang-tidy profiling feedback enables more targeted kernel modifications and measurably better performance.The comparison is against optimization without explicit profiling feedback.Source: Figure 7 · p. 8, §6 · p. 9

Modern GPU Programming

• GPU programming frameworks such as Triton and ThunderKittens [53] emphasize high performance, developer productivity, human usability, and machine-learning-library integration.Triton provides a Python interface for efficient low-level GPU kernels without requiring CUDA or OpenCL expertise, while ThunderKittens simplifies complex workflows.Source: §7 · p. 9

Scientific Discovery with LLMs; Software Engineering with LLMs; Automated CUDA Kernel Discovery with LLMs

• LLMs have been applied to scientific discovery, software engineering, and increasingly CUDA kernel generation and optimization [63].Applications include paper review, idea generation, research synthesis, scientific experiments, programming benchmarks, GitHub issues, and attention-kernel optimization with DeepSeek-R1.Source: §7 · p. 9

Automated CUDA Kernel Discovery with LLMs

• Unlike prior CUDA-focused approaches, this work presents an end-to-end system combining evolutionary test-time scaling, self-verification, model ensembling, and profiling data.Prior approaches include METR for CUDA kernel generation and Chen et al. for attention-kernel optimization using DeepSeek-R1.Source: §7 · p. 9

Summary

• The work introduces robust-kbench and an automatic framework combining kernel discovery, self-verification, evolutionary optimization, soft verification, and hardware profiling to translate and optimize PyTorch operations into CUDA.The authors argue that existing LLM-written kernel benchmarks can produce artificial speedups that obscure performance assessment.Source: §8 · p. 9

Costs & Runtime

• Optimization costs approximately $5 per kernel in API credits and under 2 hours on 4 GPUs, including generation, verification, and profiling.The authors note that additional cost and runtime can scale translation and optimization effectiveness and support more complex tasks.Source: §8 · p. 9

Future Directions & Limitations

• Similar improvements remain difficult for backward kernels, and broadly effective speedups across input shapes are still challenging.The demonstrated inference-kernel improvements target specific settings rather than providing a universal benefit across shapes.Source: §8 · p. 9
• The released benchmark, kernels, speedups, profiling data, and error messages could support post-training open-source models through supervised fine-tuning or off-policy reinforcement learning [18].Future work also includes targeting hardware-specific strengths and instruction details.Source: §8 · p. 9

Broader Impact Statement

• The project seeks to democratize high-performance computing and reduce computational costs and energy use, but may widen resource gaps between organizations with and without powerful LLMs and GPUs.The authors encourage consideration of downstream environmental impacts.Source: §8 · p. 9

A.1 Filtering Procedure

• These vulnerabilities can produce kernels that appear fast on benchmarks while failing to generalize across input shapes, precision requirements, and integration settings.The supplementary material motivates robust task curation as a defense against hardcoding, removing supposedly redundant operations, and exploiting weak test conditions.Source: §A.1 · p. 17
• The filtering procedure removes tasks with low output magnitude, insufficient seed variation, uniform outputs, minimal input impact, or redundant baseline computations.The criteria target floating-point noise, trivial solutions, overfitting, and inherently inefficient reference implementations.Source: §A.1 · p. 17
• The filtering criteria are inspired by METR and use 0.01 thresholds for output magnitude, seed variation, and input impact.Sonnet-3.7 is additionally used to identify baselines with redundant or inherently inefficient operations.Source: §A.1 · p. 17

A.2 Compromised Level 1 Tasks

• Table 1 catalogs KernelBench Level 1 tasks with compromised properties.The table is presented as the overview of compromised Level 1 tasks; the supplied passage does not provide its individual entries.Source: Table 1 · p. 17

A.3 Compromised Level 2 Tasks

• Table 2 catalogs KernelBench Level 2 tasks with compromised properties.The table is presented as the overview of compromised Level 2 tasks; the supplied passage does not provide its individual entries.Source: Table 2 · p. 18

A.4.1 Level 1 - Task 12 Matmul_with_diagonal_matric

• 51.159x speedup over torch eager is reported for the diagonal-matrix task because the baseline does not use efficient broadcasting.This example illustrates how an inefficient baseline can enable artificial performance gains.Source: §A.4.1 · p. 19
• The example CUDA implementation tiles eight rows by 128 columns, stages diagonal values in shared memory, and unrolls four column blocks.The kernel uses 32x8 threads, shared-memory row values, masked boundary handling, and separate float and double write paths.Source: §A.4.1 · p. 19

A.4.2 Level 2 - Task 13 ConvTranspose3d_Mean_Add_Softmax

• 123.603x speedup over torch eager was reported, but the kernel hardcodes 1.0 for softmax over a one-dimensional array.This contamination compromises the validity of the reported performance result.Source: §A.4.2 · p. 20

B Robust Kernel Benchmark

• The robust-kbench context exposes a benchmark result whose apparent speed advantage is invalidated by the softmax contamination.The supplied benchmark evidence reports speed relative to torch eager but does not provide additional correctness metrics.Source: §A.4.2 · p. 20

A.4.2 Level 2 - Task 13 ConvTranspose3d_Mean_Add_Softmax

• A 2D grid maps threads across output X and Y positions while combining batch and depth in the Z dimension for coalesced access.Each thread processes four consecutive X positions, with vectorized writes when alignment and validity permit.Source: §A.4.2 · pp. 20, 22
• Input-coordinate bounds are precomputed for Z, Y, and X so the transposed convolution iterates only over valid input and kernel ranges.The kernel derives separate start and end ranges for each spatial dimension from output coordinates, stride, padding, and kernel sizes.Source: §A.4.2 · p. 21
• The kernel fuses ConvTranspose3d, output-channel mean, bias addition, softmax, tanh, and scaling into one CUDA implementation.The softmax stage is implemented as a constant because it operates on a single value.Source: §A.4.2 · pp. 21, 23
• The implementation uses 32×8 thread blocks, four-way X vectorization, and coalesced memory access to promote warp-aligned execution.The X dimension uses 32 threads for warp alignment, while each thread handles four values.Source: §A.4.2 · pp. 21–22

B.1 Task Overview

• Benchmark tasks are categorized by operation type and supported features.Source: Table 3 · p. 23

B.2 Linear Backward Kernel Task Definition

• The backward task computes and returns grad_input, grad_weights, and grad_biases from grad_output, x, and weights.The backward function retrieves saved tensors before invoking the configured backward implementation.Source: §B.2 · p. 23
• The forward task implements y = x @ W^T + b using a linear transformation with bias broadcasting across the batch dimension.The model exposes the forward function through a configurable callable and uses the specified weights and biases.Source: §B.2 · pp. 23–24
• The model uses configurable feature dimensions and initializes weights with kaiming, xavier, or normal schemes, while biases receive nonzero normal initialization.The default model dimensions are 4096 input and 4096 output features.Source: §B.2 · pp. 23–24

B.3 Linear Backward Kernel Task Configuration

• The task configuration includes single-input cases with batch size 64, 128 input features, 10 output features, and kaiming initialization.These settings are listed as separate single-input, shared, and initialization configurations.Source: §B.3 · p. 24
• Multi-input configurations include batch size 4 together with 10 output features and kaiming initialization.The configuration also specifies 128 input features for the multi-input shared settings.Source: §B.3 · p. 24
• The cited configuration specifies 4096 output features, uses xavier initialization, and specifies 4096 input features.Source: §B.3 · p. 24

C.1 Detailed Translation Results on KernelBench

• Iterative error feedback achieves higher translation success than parallel sampling given a similar number of proposals.Figure 8 evaluates test-time scaling in the torch-to-CUDA translation pipeline.Source: Figure 8 · p. 25
• Increasing the number of proposals generally improves translation success on KernelBench.Source: Figure 8 · p. 25

C.2 Generalization of Kernels on Different Hardware

• Table 4 evaluates discovered-kernel speedups across different GPUs relative to PyTorch implementations.The comparison covers both forward (F) and backward (B) kernel operations.Source: Table 4 · p. 25
• The hardware-generalization evaluation distinguishes forward and backward operations when comparing discovered kernels with PyTorch.The table reports speedup comparisons across GPUs, but the supplied passage does not provide cell values.Source: Table 4 · p. 25

C.3 More Baseline Comparisons against Our Evolutionary Approach

• Table 5 compares best-of-40 results from different model-based approaches with the evolutionary pipeline using 40 kernel proposals.The supplied passage identifies the comparison scope but does not provide numerical results or a winner.Source: Table 5 · p. 26
• The baseline evaluation directly contrasts model-based approaches with the evolutionary approach under the same 40-proposal budget.Source: Table 5 · p. 26

Evaluation Environment

• Kernels are evaluated on H100 GPUs with CUDA 12.4, cuDNN 8.9.7, and the package versions listed in Table 6.The paper cautions that speedup estimates vary across evaluation environment settings.Source: D Kernel Evaluation Environment · p. 26, Table 6 · p. 26

Correctness Criteria

• Correctness requires passing tests across input sizes, initialization settings, and random seeds under absolute and relative floating-point tolerances of 1e-5.Source: D Kernel Evaluation Environment · p. 26

Runtime Measurement

• Speedups are measured against native PyTorch for one input setting, except LayerNorm, which uses torch compile, by averaging 2000 runs after 25 warmups.Source: D Kernel Evaluation Environment · p. 26

Evaluation Sensitivity

• Re-evaluation shows that quantitative speedups vary across evaluation methods, although the kernels’ qualitative performance remains consistent.Source: Figure 9 · p. 27

E Kernel Optimization Cost and Runtime Analysis

• Cumulative API costs typically reach approximately $4–$5 per benchmark task over 40 kernel proposals, with backward tasks approaching $5 and CUDA API costs generally larger than verifier costs.The iterative process makes proposal efficiency and per-iteration accuracy important for cost control.Source: E Kernel Optimization Cost and Runtime… · p. 27, Figure 10 · p. 27
• Translation takes up to 10 sequential sampling, correctness-check, and runtime-evaluation steps and averages 15 minutes, while optimization averages 1.5 hours on 4 GPUs.Optimization samples 8 proposals per generation, self-verifies them, compiles the top 4, checks correctness, and estimates runtime across 10 generations.Source: E Kernel Optimization Cost and Runtime… · p. 28

F Hyperparameter Settings

• The optimization configuration uses o4-mini with temperature [1.0], 16384 maximum tokens, [high] reasoning effort, one sample per iteration, 20 verifier-tuning generations, and error summarization enabled.These settings are listed in Table 8.Source: Table 8 · p. 28

G.1 LLM-Driven CUDA Kernel Translation

• Translation prompts require forward and backward CUDA kernels to perform the exact PyTorch operation, include a pybind11 CUDA module, and return code in <cuda></cuda> tags.Backward prompts additionally request minimized torch-function usage, a shortened kernel descriptor, and an implementation summary.Source: §G.1 · p. 29

G.2 LLM-Driven CUDA Kernel Optimization

• Optimization prompts target runtime reduction on specified GPU, CUDA, and cuDNN versions while requiring correct results, limited torch usage, and structured kernel name, description, and code output.Iteration prompts ask for new CUDA kernels that improve speedup while preserving correctness and the exact output schema.Source: §G.2 · p. 30

G.3 LLM-Driven CUDA Kernel Verifier Tuning

• Verifier-tuning prompts separately address compilation, numerical correctness, and memory usage, then iteratively improve system and instruction prompts from verifier-performance feedback.The constructed verifier prompt combines the problem description, CUDA code, instructions, and parsing definition.Source: §G.3 · pp. 30–31

G.4 LLM-Driven CUDA Kernel Verification

• The tuned compilation verifier emulates nvcc, ptxas, and nvlink by preprocessing and type-checking CUDA sources, validating kernels and resources, invoking ptxas, linking libdevice, and reporting genuine diagnostics.It models CUDA built-ins, include paths, symbols, intrinsic overloads, qualifiers, template instantiation, inline assembly, constant/shared-memory and register limits, and linking failures.Source: §G.4 · p. 31
• The verification workflow explicitly analyzes host guards, constant/global/shared memory capacities, extern-shared sizing, register limits, inlined device code, unrolled loops, pointer-derived indices, and vectorized access alignment.It derives exact index intervals and allocation requirements, flags illegal accesses or resource overflows, and treats non-constant throwing guards as host-guard failures rather than safety guarantees.Source: §G.4 · pp. 33–34
• The memory verifier classifies kernels as valid only when its aggregated error list is empty, reporting structured host-guard, device-limit, memory-bound, alignment, pointer, and numeric-range violations.It outputs exactly a JSON object with valid and errors fields, and returns a final true/false verification answer based on detected issues.Source: §G.4 · pp. 33–34
• The tuned numerical verifier models CUDA compilation, shared-memory layout and initialization, alias coverage, barrier divergence, tiling completeness, unique output writes, reductions, and numerical stability.It symbolically checks shared-memory byte ranges, read-before-write behavior, synchronization conditions, full [0...K) tile coverage, output-domain coverage, and risks such as underflow, overflow, division by zero, and NaN propagation.Source: §G.4 · pp. 34–35
• The numerical iteration prompt validates kernels by pseudo-executing thread 0 and the last active thread on regular, tail, and fallback cases, then reporting fixes and a final correctness verdict.The final assessment covers compilation, hardware limits, synchronization, complete race-free output coverage, and NaN-free results for valid inputs.Source: §G.4 · pp. 35–36

H Highlighted CUDA Kernels: Analysis

• Highlighted kernels combine fusion, vectorized memory access, warp-level reductions, shared-memory caching, and specialized paths across normalization, feedforward, convolution, classification, linear, residual, and pooling operations.The highlighted set includes both forward and backward kernels, with optimizations tailored to dimensions, class counts, pooling windows, and diverse tensor shapes.Source: H Highlighted CUDA Kernels: Analysis · p. 37

I.1 MNIST Conv-ReLU-Pool Forward Kernel

• The Conv-ReLU-Pool kernel maps output tiles to threads, caches each output channel’s weights in shared memory, unrolls spatial and kernel loops, and specializes K = 3.It computes convolution, ReLU, and 2x2 max pooling in one kernel, with a fallback path for other kernel sizes.Source: H Highlighted CUDA Kernels: Analysis · p. 37, §I.1 · pp. 38–39

I.4 MNIST Cross-Entropy Forward Kernel

• The Cross-Entropy forward kernel assigns samples to warps, stages logits in shared memory, and uses unrolled warp reductions for the maximum and exponentiated sum.Known class counts use a fully unrolled template, whereas unknown or large counts use a dynamically sized fallback with partial unrolling.Source: H Highlighted CUDA Kernels: Analysis · p. 37, §I.4 · pp. 42–43

I.2 MNIST Linear-ReLU Forward Kernel

• The Linear-ReLU kernel assigns each warp one batch-output pair, accumulates strided dot products, reduces with warp shuffles, then adds bias and applies ReLU.A grid-stride loop distributes B × O output pairs across warps, while a four-warp block configuration supports occupancy and load balancing.Source: §I.2 · p. 40

I.3 MNIST Linear Forward Kernel

• The Linear forward kernel uses one warp per output element, with lanes computing strided products, shuffle-based reduction, and lane 0 adding bias before writing the result.The launch grid is two-dimensional over output columns and input rows, with 32-thread blocks.Source: §I.3 · p. 41

I.4 MNIST Cross-Entropy Forward Kernel

• The MNIST cross-entropy kernel specializes common class counts C=10 and C=100, falling back to a dynamic-class implementation otherwise.The host wrapper dispatches specialized or dynamic kernels and normalizes the accumulated loss by N.Source: §I.4 · pp. 43–44

I.5 MNIST MaxPool Backward Kernel

• The 2x2 stride-2 max-pooling backward kernel assigns each output position independently, selects its input-window maximum, and routes the upstream gradient without atomics.It zero-initializes grad_in once, uses read-only cached loads for the four window values, and launches with 256 threads per block for H100 occupancy.Source: §I.5 · pp. 44–45

I.6 MNIST Linear-ReLU Backward Kernel

• The fused Linear-ReLU backward implementation first builds masked gradients and bias gradients, then computes weight and input gradients in separate kernels.The ReLU mask gates grad_out, bias accumulation uses atomicAdd, and the dispatcher returns grad_input, grad_weight, and grad_bias.Source: §I.6 · pp. 45–47

I.7 MNIST Linear Backward Kernel

• The adaptive Linear backward implementation uses distinct kernels for grad_input and grad_weights, with configurable block sizes and four-way unrolled accumulation over output features.The grad_input kernel computes grad_output @ weights^T, while the grad_weights kernel computes grad_output^T @ x.Source: §I.7 · pp. 47–48
• The MNIST linear backward kernel computes gradients for inputs, weights, and bias using separate CUDA kernels with tiled grids and block-wide reductions.Gradient weights use a 2D output-feature/input-feature grid with batch unrolling, while gradient bias assigns one block per output feature and reduces through shared memory.Source: §I.7 · p. 48

I.9 LayerNorm Forward Kernel

• The LayerNorm forward kernel computes per-instance sums and squared sums with float4 loads, handles unaligned remainder elements, and combines partial results through warp reductions.Warp partials are stored in shared memory before a first-warp reduction, supporting block processing across the inner dimension.Source: §I.9 · p. 53
• LayerNorm computes per-instance statistics in one kernel, then applies normalization and affine transformation with float4 vectorization plus scalar remainder handling.The host wrapper launches separate statistics and transformation kernels over an outer-dimension grid, using contiguous tensors and temporary sum buffers.Source: §I.9 · pp. 54–55

I.10 ResNet Block Forward Kernel

• The ResNet forward kernel derives batch-normalization scale and shift, supports residual and ReLU variants, and selects float4 or scalar affine kernels.Scale and shift vectors are uploaded to constant memory; the vectorized path requires spatial dimensions divisible by four and 16-byte input alignment, while the scalar path supports any dtype.Source: §I.10 · pp. 55–57

I.11 Llama RMSNorm Forward Kernel

• Llama RMSNorm uses strided float4 loads, warp-level and block-wide reductions, adaptive block sizes, and a second vectorized pass to normalize and write outputs.The normalization factor is computed as rsqrtf(total / D + eps), with scalar loops handling dimensions not divisible by four.Source: §I.11 · pp. 58–59
• The RMSNorm forward implementation requires CUDA, contiguous float32 tensors, checks weight dimensionality, and derives the hidden dimension D and row count M.These checks constrain the supported inputs before kernel launch.Source: §I.11 · pp. 59–60
• RMSNorm selects a block size from D and dispatches specialized strided kernels for block sizes 512, 256, 128, 64, or 32 before checking CUDA errors.The implementation uses the current CUDA stream and raises an error when the launch does not return cudaSuccess.Source: §I.11 · p. 60

I.12 Llama Feedforward Block Forward Kernel

• The Llama feedforward kernel fuses SiLU multiplication using warp-uniform predication and float4 vectorization, avoiding divergent branches while processing valid elements.Each thread handles two float4 chunks, with uniform bounds checks and predicated loads, computation, and stores.Source: §I.12 · pp. 60–61
• The feedforward launcher fixes 256 threads per block and two float4 vectors per thread, computes the required warp and block counts, and caps blocks at 65535.The fused kernel is launched on the current stream through a PyTorch extension wrapper.Source: §I.12 · pp. 61–62
• The feedforward wrapper performs two cuBLAS projections, applies the fused SiLU.mul kernel, then runs a third projection and reshapes output to B, S, H.It accepts CUDA float32 3-D input, flattens batch and sequence dimensions, and executes all operations on the current CUDA stream.Source: §I.12 · pp. 61–62