1use std::fmt::Write;
9
10use crate::graph::ExprGraph;
11use crate::node::{ExprId, Node};
12
13#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
15pub enum Dialect {
16 Wgsl,
18 Msl,
20 Cuda,
22 C,
24}
25
26pub struct ComputeKernel {
28 pub source: String,
30 pub n_inputs: usize,
32 pub n_outputs: usize,
34 pub workgroup_size: u32,
36 pub dialect: Dialect,
38 pub entry_point: &'static str,
40}
41
42impl ExprGraph {
43 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
91fn 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 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 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 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
176fn 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 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 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
227fn 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
236fn 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
257fn 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 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 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 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
326fn 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
370fn 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 writeln!(src).unwrap();
386
387 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
421fn 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 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 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 assert!(kernel.source.contains("const float* __restrict__ in0,"));
516 assert!(kernel.source.contains("const float* __restrict__ in1,"));
517 assert!(!kernel.source.contains("inputs["));
518 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 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}