Crate ringkernel_cuda_codegen 

CUDA code generation from Rust DSL for RingKernel stencil kernels.

This crate provides transpilation from a restricted Rust DSL to CUDA C code, enabling developers to write GPU kernels in Rust without directly writing CUDA.

Overview

The transpiler supports a subset of Rust focused on stencil/grid operations:

• Primitive types: f32, f64, i32, u32, i64, u64, bool
• Array slices: &[T], &mut [T]
• Arithmetic and comparison operators
• Let bindings and if/else expressions
• Stencil intrinsics via GridPos context

Example

use ringkernel_cuda_codegen::{transpile_stencil_kernel, StencilConfig};

let rust_code = r#"
    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 {
    id: "fdtd".to_string(),
    grid: Grid::Grid2D,
    tile_size: (16, 16),
    halo: 1,
};

let cuda_code = transpile_stencil_kernel(rust_code, &config)?;

Re-exports

pub use handler::generate_cuda_struct;

pub use handler::generate_message_deser;

pub use handler::generate_response_ser;

pub use handler::ContextMethod;

pub use handler::HandlerCodegenConfig;

pub use handler::HandlerParam;

pub use handler::HandlerParamKind;

pub use handler::HandlerReturnType;

pub use handler::HandlerSignature;

pub use handler::MessageTypeInfo;

pub use handler::MessageTypeRegistry;

pub use loops::LoopPattern;

pub use loops::RangeInfo;

pub use persistent_fdtd::generate_persistent_fdtd_kernel;

pub use persistent_fdtd::PersistentFdtdConfig;

pub use reduction_intrinsics::generate_inline_block_reduce;

pub use reduction_intrinsics::generate_inline_grid_reduce;

pub use reduction_intrinsics::generate_inline_reduce_and_broadcast;

pub use reduction_intrinsics::generate_reduction_helpers;

pub use reduction_intrinsics::transpile_reduction_call;

pub use reduction_intrinsics::ReductionCodegenConfig;

pub use reduction_intrinsics::ReductionIntrinsic;

pub use reduction_intrinsics::ReductionOp as CodegenReductionOp;

pub use ring_kernel::generate_control_block_struct;

pub use ring_kernel::generate_hlc_struct;

pub use ring_kernel::generate_k2k_structs;

pub use ring_kernel::KernelReductionConfig;

pub use ring_kernel::RingKernelConfig;

pub use shared::SharedArray;

pub use shared::SharedMemoryConfig;

pub use shared::SharedMemoryDecl;

pub use shared::SharedTile;

Modules

dsl

Rust DSL functions for writing CUDA kernels.

handler

Handler function integration for ring kernel transpilation.

loops

Loop transpilation helpers for CUDA code generation.

persistent_fdtd

Persistent FDTD kernel code generation.

reduction_intrinsics

Reduction Intrinsics for CUDA Code Generation

ring_kernel

Ring kernel configuration and code generation for persistent actor kernels.

shared

Shared memory support for CUDA code generation.

Structs

CudaTranspiler

CUDA code transpiler.

GridPos

Context for grid position within a stencil kernel.

IntrinsicRegistry

Registry for mapping Rust function names to GPU intrinsics.

SharedVarInfo

Information about a shared memory variable.

StencilConfig

Configuration for a stencil kernel.

StencilLaunchConfig

Launch configuration for stencil kernels.

TypeMapper

Type mapper for Rust to CUDA conversions.

Enums

CudaType

CUDA type representation.

GpuIntrinsic

GPU intrinsic operations.

Grid

Grid dimensionality for stencil kernels.

RingKernelIntrinsic

Ring kernel intrinsics for persistent actor kernels.

RingKernelParamKind

Known ring kernel parameter types.

StencilIntrinsic

Stencil-specific intrinsics for neighbor access.

TranspileError

Errors that can occur during transpilation.

ValidationError

Validation errors for DSL constraint violations.

ValidationMode

Validation mode for different kernel types.

Functions

get_slice_element_type

Extract the inner element type from a slice reference.

is_control_block_type

Check if a type is the ControlBlock type.

is_mutable_reference

Check if a type is a mutable reference.

is_ring_context_type

Check if a type is the RingContext type.

is_simple_assignment

Check if a statement might need special handling.

ring_kernel_type_mapper

Create a TypeMapper with ring kernel types pre-registered.

transpile_device_function

Transpile a Rust function to a CUDA __device__ function.

transpile_function

Transpile a function to CUDA without stencil configuration.

transpile_global_kernel

Transpile a Rust function to a CUDA __global__ kernel.

transpile_ring_kernel

Transpile a Rust handler function to a persistent ring kernel.

transpile_stencil_kernel

Transpile a Rust stencil kernel function to CUDA C code.

validate_function

Validate that a function conforms to the stencil kernel DSL.

validate_function_with_mode

Validate a function with a specific validation mode.

validate_stencil_signature

Validate function signature for stencil kernels.

Type Aliases

Result

Result type for transpilation operations.