Crate ringkernel_wgpu_codegen 

WGSL code generation from Rust DSL for RingKernel.

This crate provides transpilation from a restricted Rust DSL to WGSL (WebGPU Shading Language), enabling developers to write GPU kernels in Rust that target WebGPU-compatible devices.

Overview

The transpiler supports the same DSL as ringkernel-cuda-codegen:

• Primitive types: f32, i32, u32, bool (with 64-bit emulation)
• Array slices: &[T], &mut [T] → storage buffers
• Arithmetic and comparison operators
• Let bindings and if/else expressions
• For/while/loop constructs
• Stencil intrinsics via GridPos context
• Ring kernel support with host-driven persistence emulation

Example

use ringkernel_wgpu_codegen::{transpile_stencil_kernel, StencilConfig};
use syn::parse_quote;

let func: syn::ItemFn = parse_quote! {
    fn fdtd(p: &[f32], p_prev: &mut [f32], c2: f32, pos: GridPos) {
        let curr = p[pos.idx()];
        let lap = pos.north(p) + pos.south(p) + pos.east(p) + pos.west(p) - 4.0 * curr;
        p_prev[pos.idx()] = 2.0 * curr - 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(&func, &config)?;

WGSL Limitations

Compared to CUDA, WGSL has some limitations that are handled with workarounds:

• No 64-bit atomics: Emulated using lo/hi u32 pairs
• No f64: Downcast to f32 with a warning
• No persistent kernels: Emulated with host-driven dispatch loops
• No warp operations: Mapped to subgroup operations where available
• No kernel-to-kernel messaging: K2K is not supported in WGPU

Re-exports

pub use bindings::generate_bindings;

pub use bindings::AccessMode;

pub use bindings::BindingLayout;

pub use handler::HandlerCodegenConfig;

pub use handler::HandlerParam;

pub use handler::HandlerParamKind;

pub use handler::HandlerReturnType;

pub use handler::HandlerSignature;

pub use handler::WgslContextMethod;

pub use intrinsics::IntrinsicRegistry;

pub use intrinsics::WgslIntrinsic;

pub use loops::LoopPattern;

pub use loops::RangeInfo;

pub use ring_kernel::RingKernelConfig;

pub use shared::SharedArray;

pub use shared::SharedMemoryConfig;

pub use shared::SharedMemoryDecl;

pub use shared::SharedTile;

pub use stencil::Grid;

pub use stencil::GridPos;

pub use stencil::StencilConfig;

pub use stencil::StencilLaunchConfig;

pub use transpiler::transpile_function;

pub use transpiler::WgslTranspiler;

pub use types::AddressSpace;

pub use types::TypeMapper;

pub use types::WgslType;

pub use u64_workarounds::U64Helpers;

pub use validation::ValidationError;

pub use validation::ValidationMode;

pub use dsl::*;

Modules

bindings

Buffer binding layout generation for WGSL.

dsl

DSL marker functions for WGSL code generation.

handler

Handler function parsing and codegen configuration.

intrinsics

Intrinsic registry for WGSL code generation.

loops

Loop transpilation for WGSL code generation.

ring_kernel

Ring kernel generation for WGSL.

shared

Shared (workgroup) memory support for WGSL code generation.

stencil

Stencil kernel support for WGSL code generation.

transpiler

Core Rust-to-WGSL transpiler.

types

Type mapping from Rust to WGSL.

u64_workarounds

64-bit integer workarounds for WGSL.

validation

DSL validation for WGSL code generation.

Enums

TranspileError

Errors that can occur during transpilation.

Functions

transpile_device_function

Transpile a Rust function to a WGSL helper function.

transpile_global_kernel

Transpile a Rust function to a WGSL @compute kernel.

transpile_ring_kernel

Transpile a Rust handler function to a WGSL ring kernel.

transpile_stencil_kernel

Transpile a Rust stencil kernel function to WGSL code.

Type Aliases

Result

Result type for transpilation operations.