ringkernel-wgpu-codegen 0.4.0

WGSL code generation from Rust DSL for RingKernel
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108

# ringkernel-wgpu-codegen

Rust-to-WGSL transpiler for RingKernel GPU kernels.

## Overview

This crate provides transpilation from a restricted Rust DSL to WGSL (WebGPU Shading Language). It shares the same DSL as `ringkernel-cuda-codegen`, allowing the same kernel code to target both CUDA and WebGPU.

## Installation

```toml
[dependencies]
ringkernel-wgpu-codegen = "0.2"
syn = { version = "2.0", features = ["full"] }
```

## Global Kernels

```rust
use ringkernel_wgpu_codegen::transpile_global_kernel;
use syn::parse_quote;

let func: 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(&func)?;
```

## Stencil Kernels

```rust
use ringkernel_wgpu_codegen::{transpile_stencil_kernel, StencilConfig};

let func: syn::ItemFn = parse_quote! {
    fn heat(temp: &[f32], temp_new: &mut [f32], alpha: f32, pos: GridPos) {
        let t = temp[pos.idx()];
        let neighbors = pos.north(temp) + pos.south(temp)
                      + pos.east(temp) + pos.west(temp);
        temp_new[pos.idx()] = t + alpha * (neighbors - 4.0 * t);
    }
};

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

let wgsl_code = transpile_stencil_kernel(&func, &config)?;
```

## Ring Kernels

Ring kernels in WGPU are emulated using host-driven dispatch loops:

```rust
use ringkernel_wgpu_codegen::{transpile_ring_kernel, RingKernelConfig};

let handler: syn::ItemFn = parse_quote! {
    fn process(value: f32) -> f32 {
        value * 2.0
    }
};

let config = RingKernelConfig::new("processor")
    .with_workgroup_size(256)
    .with_hlc(true);

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

## WGSL Limitations

The transpiler handles these WebGPU limitations:

| CUDA Feature | WGSL Workaround |
|--------------|-----------------|
| 64-bit atomics | Lo/hi u32 pairs with manual carry |
| `f64` | Downcast to `f32` (with warning) |
| Persistent kernels | Host-driven dispatch loop |
| Warp operations | Subgroup operations (where available) |
| K2K messaging | Not supported |

## Type Mapping

| Rust Type | WGSL Type |
|-----------|-----------|
| `f32` | `f32` |
| `i32` | `i32` |
| `u32` | `u32` |
| `bool` | `bool` |
| `&[T]` | `array<T>` (storage, read) |
| `&mut [T]` | `array<T>` (storage, read_write) |

## Testing

```bash
cargo test -p ringkernel-wgpu-codegen
```

The crate includes 50 tests covering types, intrinsics, and transpilation.

## License

Apache-2.0