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