Skip to main content

tang_expr/
codegen.rs

1//! Multi-dialect compute shader code generation.
2//!
3//! Generates compute kernels in WGSL, MSL (Metal Shading Language),
4//! CUDA C, and plain C from the same expression graph. The expression
5//! syntax is nearly identical across dialects — only kernel boilerplate
6//! and a few edge cases (select, literal suffixes) differ.
7
8use std::fmt::Write;
9
10use crate::graph::ExprGraph;
11use crate::node::{ExprId, Node};
12
13/// Target shader dialect.
14#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
15pub enum Dialect {
16    /// WebGPU Shading Language.
17    Wgsl,
18    /// Metal Shading Language.
19    Msl,
20    /// CUDA C (compiled via NVRTC).
21    Cuda,
22    /// Plain C (for CPU fallback / verification).
23    C,
24}
25
26/// A generated compute kernel.
27pub struct ComputeKernel {
28    /// Complete kernel source code.
29    pub source: String,
30    /// Number of input values per work item.
31    pub n_inputs: usize,
32    /// Number of output values per work item.
33    pub n_outputs: usize,
34    /// Workgroup / threadgroup / block size.
35    pub workgroup_size: u32,
36    /// Which dialect this kernel was generated for.
37    pub dialect: Dialect,
38    /// Entry point function name.
39    pub entry_point: &'static str,
40}
41
42impl ExprGraph {
43    /// Generate a compute kernel for the given dialect.
44    ///
45    /// Each work item reads `n_inputs` f32 values and writes `outputs.len()` f32 values.
46    /// Shared subexpressions are computed once per thread.
47    pub fn to_kernel(
48        &self,
49        outputs: &[ExprId],
50        n_inputs: usize,
51        dialect: Dialect,
52    ) -> ComputeKernel {
53        let workgroup_size = 256u32;
54        let n_outputs = outputs.len();
55        let live = self.live_set(outputs);
56        let max_id = if live.is_empty() {
57            0
58        } else {
59            *live.iter().max().unwrap()
60        };
61
62        let mut src = String::with_capacity(2048);
63
64        match dialect {
65            Dialect::Wgsl => emit_wgsl(
66                &mut src,
67                self,
68                outputs,
69                n_inputs,
70                n_outputs,
71                &live,
72                max_id,
73                workgroup_size,
74            ),
75            Dialect::Msl => emit_msl(&mut src, self, outputs, n_inputs, n_outputs, &live, max_id),
76            Dialect::Cuda => emit_cuda(&mut src, self, outputs, n_inputs, n_outputs, &live, max_id),
77            Dialect::C => emit_c(&mut src, self, outputs, n_inputs, n_outputs, &live, max_id),
78        }
79
80        ComputeKernel {
81            source: src,
82            n_inputs,
83            n_outputs,
84            workgroup_size,
85            dialect,
86            entry_point: "k0",
87        }
88    }
89}
90
91// ---------------------------------------------------------------------------
92// Shared body emission
93// ---------------------------------------------------------------------------
94
95/// Emit SSA evaluation lines shared by all dialects.
96///
97/// `decl` is the variable declaration prefix: `"let"` for WGSL, `"float"` for MSL/CUDA/C.
98/// `lit_suffix` is appended to float literals: `""` for WGSL, `"f"` for MSL/CUDA/C.
99fn emit_body(
100    src: &mut String,
101    graph: &ExprGraph,
102    outputs: &[ExprId],
103    n_inputs: usize,
104    n_outputs: usize,
105    live: &std::collections::HashSet<usize>,
106    max_id: usize,
107    indent: &str,
108    decl: &str,
109    lit_suffix: &str,
110    thread_id: &str,
111    dialect: Dialect,
112) {
113    // Load inputs
114    if n_inputs > 0 {
115        let base = format!("{thread_id} * {n_inputs}u");
116        let base_wgsl = format!("{thread_id} * {n_inputs}");
117        for i in 0..n_inputs {
118            match dialect {
119                Dialect::Wgsl => {
120                    writeln!(src, "{indent}{decl} x{i} = inputs[{base} + {i}u];").unwrap();
121                }
122                Dialect::Msl => {
123                    writeln!(src, "{indent}{decl} x{i} = inputs[{base_wgsl} + {i}];").unwrap();
124                }
125                Dialect::Cuda => {
126                    writeln!(src, "{indent}{decl} x{i} = in{i}[{thread_id}];").unwrap();
127                }
128                Dialect::C => {
129                    writeln!(src, "{indent}{decl} x{i} = inputs[i * {n_inputs} + {i}];").unwrap();
130                }
131            }
132        }
133        writeln!(src).unwrap();
134    }
135
136    // Evaluate in topological order (SSA form)
137    for i in 0..=max_id {
138        if !live.contains(&i) {
139            continue;
140        }
141        let node = graph.node(ExprId(i as u32));
142        match node {
143            Node::Var(_) | Node::Lit(_) => continue,
144            _ => {}
145        }
146        let rhs = expr_str(graph, node, lit_suffix, dialect);
147        writeln!(src, "{indent}{decl} t{i} = {rhs};").unwrap();
148    }
149    writeln!(src).unwrap();
150
151    // Store outputs
152    if n_outputs > 0 {
153        for (k, out) in outputs.iter().enumerate() {
154            let val = ref_str(graph, *out, lit_suffix);
155            match dialect {
156                Dialect::Wgsl => {
157                    let base = format!("{thread_id} * {n_outputs}u");
158                    writeln!(src, "{indent}outputs[{base} + {k}u] = {val};").unwrap();
159                }
160                Dialect::Msl => {
161                    let base = format!("{thread_id} * {n_outputs}");
162                    writeln!(src, "{indent}outputs[{base} + {k}] = {val};").unwrap();
163                }
164                Dialect::Cuda => {
165                    let base = format!("{thread_id} * {n_outputs}");
166                    writeln!(src, "{indent}outputs[{base} + {k}] = {val};").unwrap();
167                }
168                Dialect::C => {
169                    writeln!(src, "{indent}outputs[i * {n_outputs} + {k}] = {val};").unwrap();
170                }
171            }
172        }
173    }
174}
175
176/// Generate expression string for a node.
177fn expr_str(graph: &ExprGraph, node: Node, suffix: &str, dialect: Dialect) -> String {
178    match node {
179        Node::Var(n) => format!("x{n}"),
180        Node::Lit(bits) => format_literal(f64::from_bits(bits), suffix),
181        Node::Add(a, b) => format!(
182            "({} + {})",
183            ref_str(graph, a, suffix),
184            ref_str(graph, b, suffix)
185        ),
186        Node::Mul(a, b) => format!(
187            "({} * {})",
188            ref_str(graph, a, suffix),
189            ref_str(graph, b, suffix)
190        ),
191        Node::Neg(a) => format!("(-{})", ref_str(graph, a, suffix)),
192        Node::Recip(a) => format!("(1.0{suffix} / {})", ref_str(graph, a, suffix)),
193        Node::Sqrt(a) => format!("sqrt({})", ref_str(graph, a, suffix)),
194        Node::Sin(a) => format!("sin({})", ref_str(graph, a, suffix)),
195        Node::Atan2(y, x) => format!(
196            "atan2({}, {})",
197            ref_str(graph, y, suffix),
198            ref_str(graph, x, suffix)
199        ),
200        Node::Exp2(a) => format!("exp2({})", ref_str(graph, a, suffix)),
201        Node::Log2(a) => format!("log2({})", ref_str(graph, a, suffix)),
202        Node::Select(c, a, b) => {
203            match dialect {
204                Dialect::Wgsl => {
205                    // WGSL: select(false_val, true_val, cond)
206                    format!(
207                        "select({}, {}, {} > 0.0)",
208                        ref_str(graph, b, suffix),
209                        ref_str(graph, a, suffix),
210                        ref_str(graph, c, suffix),
211                    )
212                }
213                _ => {
214                    // MSL/CUDA/C: ternary
215                    format!(
216                        "({} > 0.0{suffix} ? {} : {})",
217                        ref_str(graph, c, suffix),
218                        ref_str(graph, a, suffix),
219                        ref_str(graph, b, suffix),
220                    )
221                }
222            }
223        }
224    }
225}
226
227/// Reference a node inline: Var → x{n}, Lit → literal, others → t{index}.
228fn ref_str(graph: &ExprGraph, id: ExprId, suffix: &str) -> String {
229    match graph.node(id) {
230        Node::Var(n) => format!("x{n}"),
231        Node::Lit(bits) => format_literal(f64::from_bits(bits), suffix),
232        _ => format!("t{}", id.0),
233    }
234}
235
236/// Format f64 as a float literal with optional suffix.
237fn format_literal(v: f64, suffix: &str) -> String {
238    let base = if v == 0.0 {
239        "0.0".to_string()
240    } else if v == 1.0 {
241        "1.0".to_string()
242    } else if v == -1.0 {
243        "-1.0".to_string()
244    } else if v == 2.0 {
245        "2.0".to_string()
246    } else {
247        let s = format!("{v}");
248        if s.contains('.') || s.contains('e') || s.contains('E') {
249            s
250        } else {
251            format!("{s}.0")
252        }
253    };
254    format!("{base}{suffix}")
255}
256
257// ---------------------------------------------------------------------------
258// WGSL dialect
259// ---------------------------------------------------------------------------
260
261fn emit_wgsl(
262    src: &mut String,
263    graph: &ExprGraph,
264    outputs: &[ExprId],
265    n_inputs: usize,
266    n_outputs: usize,
267    live: &std::collections::HashSet<usize>,
268    max_id: usize,
269    workgroup_size: u32,
270) {
271    writeln!(src, "// Auto-generated by tang-expr").unwrap();
272    writeln!(src).unwrap();
273
274    // Params struct
275    writeln!(src, "struct Params {{").unwrap();
276    writeln!(src, "    count: u32,").unwrap();
277    writeln!(src, "    _pad1: u32,").unwrap();
278    writeln!(src, "    _pad2: u32,").unwrap();
279    writeln!(src, "    _pad3: u32,").unwrap();
280    writeln!(src, "}}").unwrap();
281    writeln!(src).unwrap();
282
283    // Bindings
284    writeln!(
285        src,
286        "@group(0) @binding(0) var<storage, read> inputs: array<f32>;"
287    )
288    .unwrap();
289    writeln!(
290        src,
291        "@group(0) @binding(1) var<storage, read_write> outputs: array<f32>;"
292    )
293    .unwrap();
294    writeln!(src, "@group(0) @binding(2) var<uniform> params: Params;").unwrap();
295    writeln!(src).unwrap();
296
297    // Entry point
298    writeln!(src, "@compute @workgroup_size({workgroup_size})").unwrap();
299    writeln!(
300        src,
301        "fn k0(@builtin(global_invocation_id) gid: vec3<u32>) {{"
302    )
303    .unwrap();
304    writeln!(src, "    let idx = gid.x;").unwrap();
305    writeln!(src, "    if (idx >= params.count) {{ return; }}").unwrap();
306    writeln!(src).unwrap();
307
308    emit_body(
309        src,
310        graph,
311        outputs,
312        n_inputs,
313        n_outputs,
314        live,
315        max_id,
316        "    ",
317        "let",
318        "",
319        "idx",
320        Dialect::Wgsl,
321    );
322
323    writeln!(src, "}}").unwrap();
324}
325
326// ---------------------------------------------------------------------------
327// MSL dialect
328// ---------------------------------------------------------------------------
329
330fn emit_msl(
331    src: &mut String,
332    graph: &ExprGraph,
333    outputs: &[ExprId],
334    n_inputs: usize,
335    n_outputs: usize,
336    live: &std::collections::HashSet<usize>,
337    max_id: usize,
338) {
339    writeln!(src, "// Auto-generated by tang-expr").unwrap();
340    writeln!(src, "#include <metal_stdlib>").unwrap();
341    writeln!(src, "using namespace metal;").unwrap();
342    writeln!(src).unwrap();
343
344    write!(src, "kernel void k0(").unwrap();
345    writeln!(src, "device const float* inputs [[buffer(0)]],").unwrap();
346    writeln!(src, "    device float* outputs [[buffer(1)]],").unwrap();
347    writeln!(src, "    device const uint& count [[buffer(2)]],").unwrap();
348    writeln!(src, "    uint gid [[thread_position_in_grid]]) {{").unwrap();
349    writeln!(src, "    if (gid >= count) {{ return; }}").unwrap();
350    writeln!(src).unwrap();
351
352    emit_body(
353        src,
354        graph,
355        outputs,
356        n_inputs,
357        n_outputs,
358        live,
359        max_id,
360        "    ",
361        "float",
362        "f",
363        "gid",
364        Dialect::Msl,
365    );
366
367    writeln!(src, "}}").unwrap();
368}
369
370// ---------------------------------------------------------------------------
371// CUDA dialect
372// ---------------------------------------------------------------------------
373
374fn emit_cuda(
375    src: &mut String,
376    graph: &ExprGraph,
377    outputs: &[ExprId],
378    n_inputs: usize,
379    n_outputs: usize,
380    live: &std::collections::HashSet<usize>,
381    max_id: usize,
382) {
383    writeln!(src, "// Auto-generated by tang-expr").unwrap();
384    // NVRTC provides all math builtins (expf, sqrtf, fmaxf, etc.) — no #include needed.
385    writeln!(src).unwrap();
386
387    // Separate input pointers: in0, in1, ... instead of one interleaved buffer
388    write!(src, "extern \"C\" __global__ void k0(").unwrap();
389    for i in 0..n_inputs {
390        writeln!(src, "const float* __restrict__ in{i},").unwrap();
391        write!(src, "    ").unwrap();
392    }
393    writeln!(src, "float* __restrict__ outputs,").unwrap();
394    writeln!(src, "    const unsigned int count) {{").unwrap();
395    writeln!(
396        src,
397        "    unsigned int gid = blockIdx.x * blockDim.x + threadIdx.x;"
398    )
399    .unwrap();
400    writeln!(src, "    if (gid >= count) {{ return; }}").unwrap();
401    writeln!(src).unwrap();
402
403    emit_body(
404        src,
405        graph,
406        outputs,
407        n_inputs,
408        n_outputs,
409        live,
410        max_id,
411        "    ",
412        "float",
413        "f",
414        "gid",
415        Dialect::Cuda,
416    );
417
418    writeln!(src, "}}").unwrap();
419}
420
421// ---------------------------------------------------------------------------
422// C dialect
423// ---------------------------------------------------------------------------
424
425fn emit_c(
426    src: &mut String,
427    graph: &ExprGraph,
428    outputs: &[ExprId],
429    n_inputs: usize,
430    n_outputs: usize,
431    live: &std::collections::HashSet<usize>,
432    max_id: usize,
433) {
434    writeln!(src, "// Auto-generated by tang-expr").unwrap();
435    writeln!(src, "#include <math.h>").unwrap();
436    writeln!(src).unwrap();
437
438    write!(src, "void k0(").unwrap();
439    writeln!(src, "const float* inputs,").unwrap();
440    writeln!(src, "    float* outputs,").unwrap();
441    writeln!(src, "    int count) {{").unwrap();
442    writeln!(src, "    for (int i = 0; i < count; i++) {{").unwrap();
443
444    emit_body(
445        src,
446        graph,
447        outputs,
448        n_inputs,
449        n_outputs,
450        live,
451        max_id,
452        "        ",
453        "float",
454        "f",
455        "i",
456        Dialect::C,
457    );
458
459    writeln!(src, "    }}").unwrap();
460    writeln!(src, "}}").unwrap();
461}
462
463#[cfg(test)]
464mod tests {
465    use super::*;
466    use crate::graph::ExprGraph;
467
468    #[test]
469    fn wgsl_matches_original() {
470        // Verify that to_kernel(Wgsl) produces output matching to_wgsl()
471        let mut g = ExprGraph::new();
472        let x = g.var(0);
473        let y = g.var(1);
474        let xx = g.mul(x, x);
475        let yy = g.mul(y, y);
476        let sum = g.add(xx, yy);
477        let dist = g.sqrt(sum);
478
479        let old = g.to_wgsl(&[dist], 2);
480        let new = g.to_kernel(&[dist], 2, Dialect::Wgsl);
481
482        assert_eq!(new.n_inputs, old.n_inputs);
483        assert_eq!(new.n_outputs, old.n_outputs);
484        assert_eq!(new.workgroup_size, old.workgroup_size);
485        // Both should contain core elements
486        assert!(new.source.contains("@compute"));
487        assert!(new.source.contains("sqrt("));
488    }
489
490    #[test]
491    fn msl_entry_point() {
492        let mut g = ExprGraph::new();
493        let x = g.var(0);
494        let y = g.var(1);
495        let sum = g.add(x, y);
496        let kernel = g.to_kernel(&[sum], 2, Dialect::Msl);
497        assert!(kernel.source.contains("kernel void k0("));
498        assert!(kernel.source.contains("thread_position_in_grid"));
499        assert!(kernel.source.contains("#include <metal_stdlib>"));
500        assert_eq!(kernel.entry_point, "k0");
501    }
502
503    #[test]
504    fn cuda_entry_point() {
505        let mut g = ExprGraph::new();
506        let x = g.var(0);
507        let y = g.var(1);
508        let prod = g.mul(x, y);
509        let kernel = g.to_kernel(&[prod], 2, Dialect::Cuda);
510        assert!(kernel.source.contains("extern \"C\" __global__ void k0("));
511        assert!(kernel
512            .source
513            .contains("blockIdx.x * blockDim.x + threadIdx.x"));
514        // Separate input pointers instead of interleaved buffer
515        assert!(kernel.source.contains("const float* __restrict__ in0,"));
516        assert!(kernel.source.contains("const float* __restrict__ in1,"));
517        assert!(!kernel.source.contains("inputs["));
518        // Loads from separate pointers
519        assert!(kernel.source.contains("in0[gid]"));
520        assert!(kernel.source.contains("in1[gid]"));
521    }
522
523    #[test]
524    fn c_loop() {
525        let mut g = ExprGraph::new();
526        let x = g.var(0);
527        let s = g.sin(x);
528        let kernel = g.to_kernel(&[s], 1, Dialect::C);
529        assert!(kernel.source.contains("for (int i = 0; i < count; i++)"));
530        assert!(kernel.source.contains("sin("));
531    }
532
533    #[test]
534    fn msl_select_ternary() {
535        let mut g = ExprGraph::new();
536        let x = g.var(0);
537        let a = g.lit(3.0);
538        let b = g.lit(7.0);
539        let s = g.select(x, a, b);
540        let kernel = g.to_kernel(&[s], 1, Dialect::Msl);
541        // MSL should use ternary, not select()
542        assert!(kernel.source.contains("?"));
543        assert!(!kernel.source.contains("select("));
544    }
545
546    #[test]
547    fn wgsl_select_builtin() {
548        let mut g = ExprGraph::new();
549        let x = g.var(0);
550        let a = g.lit(3.0);
551        let b = g.lit(7.0);
552        let s = g.select(x, a, b);
553        let kernel = g.to_kernel(&[s], 1, Dialect::Wgsl);
554        assert!(kernel.source.contains("select("));
555    }
556
557    #[test]
558    #[allow(clippy::approx_constant)]
559    fn msl_literal_suffix() {
560        let mut g = ExprGraph::new();
561        let x = g.var(0);
562        let c = g.lit(3.14);
563        let prod = g.mul(x, c);
564        let kernel = g.to_kernel(&[prod], 1, Dialect::Msl);
565        assert!(kernel.source.contains("3.14f"));
566    }
567
568    #[test]
569    fn multiple_outputs_all_dialects() {
570        let mut g = ExprGraph::new();
571        let x = g.var(0);
572        let y = g.var(1);
573        let sum = g.add(x, y);
574        let prod = g.mul(x, y);
575
576        for dialect in [Dialect::Wgsl, Dialect::Msl, Dialect::Cuda, Dialect::C] {
577            let kernel = g.to_kernel(&[sum, prod], 2, dialect);
578            assert_eq!(kernel.n_outputs, 2);
579            assert_eq!(kernel.n_inputs, 2);
580        }
581    }
582
583    #[test]
584    fn full_pipeline_all_dialects() {
585        let mut g = ExprGraph::new();
586        let x = g.var(0);
587        let xx = g.mul(x, x);
588        let dx = g.diff(xx, 0);
589        let dx = g.simplify(dx);
590
591        for dialect in [Dialect::Wgsl, Dialect::Msl, Dialect::Cuda, Dialect::C] {
592            let kernel = g.to_kernel(&[xx, dx], 1, dialect);
593            assert_eq!(kernel.n_outputs, 2);
594            assert!(!kernel.source.is_empty());
595        }
596    }
597}