WGSL Code Generation

The ringkernel-wgpu-codegen crate provides a Rust-to-WGSL transpiler that lets you write GPU kernels in a Rust DSL and generate equivalent WGSL (WebGPU Shading Language) code.

Overview

The transpiler provides full feature parity with ringkernel-cuda-codegen, supporting the same three kernel types:

1. Global Kernels - Standard @compute entry points
2. Stencil Kernels - Tile-based kernels with GridPos abstraction
3. Ring Kernels - Persistent actor kernels (host-driven emulation)

The same Rust DSL code works with both CUDA and WGSL backends—just change the import.

Global Kernels

Generic compute shaders without stencil-specific patterns. Use DSL functions for thread/workgroup indices.

use ringkernel_wgpu_codegen::transpile_global_kernel;
use syn::parse_quote;

let kernel: syn::ItemFn = parse_quote! {
    fn saxpy(x: &[f32], y: &mut [f32], a: f32, n: i32) {
        let idx = block_idx_x() * block_dim_x() + thread_idx_x();
        if idx >= n { return; }
        y[idx as usize] = a * x[idx as usize] + y[idx as usize];
    }
};

let wgsl_code = transpile_global_kernel(&kernel)?;

Generated WGSL:

@group(0) @binding(0) var<storage, read> x: array<f32>;
@group(0) @binding(1) var<storage, read_write> y: array<f32>;
@group(0) @binding(2) var<storage, read> a: f32;
@group(0) @binding(3) var<storage, read> n: i32;

@compute @workgroup_size(256, 1, 1)
fn saxpy(
    @builtin(local_invocation_id) local_invocation_id: vec3<u32>,
    @builtin(workgroup_id) workgroup_id: vec3<u32>,
    @builtin(global_invocation_id) global_invocation_id: vec3<u32>,
    @builtin(num_workgroups) num_workgroups: vec3<u32>
) {
    var idx: i32 = i32(workgroup_id.x) * i32(256u) + i32(local_invocation_id.x);
    if (idx >= n) { return; }
    y[idx] = a * x[idx] + y[idx];
}

Stencil Kernels

For stencil computations (e.g., FDTD, convolutions), use the GridPos abstraction:

use ringkernel_wgpu_codegen::{transpile_stencil_kernel, StencilConfig};

let kernel: syn::ItemFn = parse_quote! {
    fn fdtd(p: &[f32], p_prev: &mut [f32], c2: f32, pos: GridPos) {
        let lap = pos.north(p) + pos.south(p) + pos.east(p) + pos.west(p)
                  - 4.0 * p[pos.idx()];
        p_prev[pos.idx()] = 2.0 * p[pos.idx()] - p_prev[pos.idx()] + c2 * lap;
    }
};

let config = StencilConfig::new("fdtd")
    .with_tile_size(16, 16)
    .with_halo(1);

let wgsl_code = transpile_stencil_kernel(&kernel, &config)?;

Stencil Intrinsics:

• pos.idx() - Current cell linear index
• pos.north(buf) - Access cell above (buf[idx - buffer_width])
• pos.south(buf) - Access cell below (buf[idx + buffer_width])
• pos.east(buf) - Access cell to the right (buf[idx + 1])
• pos.west(buf) - Access cell to the left (buf[idx - 1])
• pos.at(buf, dx, dy) - Access cell at relative offset

Ring Kernels

Persistent GPU kernels that process messages in a loop.

Important: WebGPU does not support true persistent kernels. Ring kernels are emulated using host-driven dispatch loops—the host re-dispatches the kernel until termination.

use ringkernel_wgpu_codegen::{transpile_ring_kernel, RingKernelConfig};

let handler: syn::ItemFn = parse_quote! {
    fn process(ctx: &RingContext, msg: &Request) -> Response {
        let tid = ctx.global_thread_id();
        ctx.sync_threads();
        let result = msg.value * 2.0;
        Response { value: result, id: tid as u64 }
    }
};

let config = RingKernelConfig::new("processor")
    .with_workgroup_size(256)
    .with_hlc(true);
    // Note: K2K is NOT supported in WGPU

let wgsl_code = transpile_ring_kernel(&handler, &config)?;

RingKernelConfig Options

K2K Limitation

Kernel-to-kernel (K2K) messaging is not supported in WGPU due to WebGPU’s execution model. Attempting to enable K2K will return an error:

let config = RingKernelConfig::new("processor")
    .with_k2k(true);  // This will error!

let result = transpile_ring_kernel(&handler, &config);
assert!(result.is_err());
// Error: "Kernel-to-kernel (K2K) messaging is not supported in WGPU"

Use host-mediated messaging instead for cross-kernel communication.

DSL Features

Control Flow

// If/else
if condition {
    do_something();
} else {
    do_other();
}

// Match (transpiles to switch/case)
match value {
    0 => handle_zero(),
    1 => handle_one(),
    _ => handle_default(),
}

// Early return
if idx >= n { return; }

Loops

// For loops (range patterns)
for i in 0..n {
    data[i] = 0.0;
}

// Inclusive range
for i in 0..=n {
    process(i);
}

// While loops
while condition {
    iterate();
}

// Infinite loops
loop {
    if done { break; }
    continue;
}

Note: Loops are forbidden in Stencil validation mode but allowed in Generic and RingKernel modes.

Shared Memory (Workgroup Variables)

use ringkernel_wgpu_codegen::{SharedTile, SharedArray};

// 2D tile: var<workgroup> tile: array<array<f32, 16>, 16>;
let tile: SharedTile<f32, 16, 16>;

// 1D array: var<workgroup> cache: array<f32, 256>;
let cache: SharedArray<f32, 256>;

Struct Literals

// Rust struct literal
let response = Response { value: result, id: tid as u64 };

// Transpiles to WGSL constructor:
// Response(result, u64(tid))

Type Mapping

GPU Intrinsics

Thread/Workgroup Indices

Synchronization

Atomics

Math Functions

Subgroup Operations (Warp Equivalents)

Note: These require the chromium_experimental_subgroups extension, which may not be available on all platforms.

64-bit Integer Emulation

WGSL 1.0 does not support 64-bit integers or atomics. The transpiler emulates them using lo/hi u32 pairs:

Type Representation

// u64 becomes vec2<u32> where:
// - x (lo) = lower 32 bits
// - y (hi) = upper 32 bits

Generated Helper Functions

The transpiler automatically includes these helpers for ring kernels:

// Read 64-bit value from atomic pair
fn read_u64(lo: ptr<storage, atomic<u32>, read_write>,
            hi: ptr<storage, atomic<u32>, read_write>) -> vec2<u32> {
    return vec2<u32>(atomicLoad(lo), atomicLoad(hi));
}

// Atomically increment 64-bit value
fn atomic_inc_u64(lo: ptr<storage, atomic<u32>, read_write>,
                  hi: ptr<storage, atomic<u32>, read_write>) {
    let old_lo = atomicAdd(lo, 1u);
    if (old_lo == 0xFFFFFFFFu) {
        atomicAdd(hi, 1u);
    }
}

// Compare two 64-bit values: returns -1, 0, or 1
fn compare_u64(a: vec2<u32>, b: vec2<u32>) -> i32 {
    if (a.y > b.y) { return 1; }
    if (a.y < b.y) { return -1; }
    if (a.x > b.x) { return 1; }
    if (a.x < b.x) { return -1; }
    return 0;
}

Validation Modes

The transpiler uses validation modes to control what constructs are allowed:

WGSL vs CUDA Differences

Testing

The crate includes 50 tests covering all features:

cargo test -p ringkernel-wgpu-codegen

Test categories:

• Type mapping (primitives, slices, 64-bit emulation)
• Expression transpilation (binary ops, method calls, indexing)
• Intrinsics (thread indices, barriers, math)
• Loops (for/while/loop patterns)
• Validation (mode-specific constraints)

Next: Migration Guide