From 8322bad5bcbc3f1ccd056ad6e868f0b8af0eb96b Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 09:46:01 +0200
Subject: [PATCH 01/17] WIP - hvm.rs numeric casts

---
 src/ast.rs                       | 19 ++++++++++++++++
 src/hvm.rs                       | 38 ++++++++++++++++++++++++++++----
 tests/programs/numeric-casts.hvm | 34 ++++++++++++++++++++++++++++
 3 files changed, 87 insertions(+), 4 deletions(-)
 create mode 100644 tests/programs/numeric-casts.hvm

diff --git a/src/ast.rs b/src/ast.rs
index 23d8cb39..dedb462d 100644
--- a/src/ast.rs
+++ b/src/ast.rs
@@ -43,8 +43,22 @@ impl<'i> CoreParser<'i> {
   pub fn parse_numb_sym(&mut self) -> Result<Numb, String> {
     self.consume("[")?;
 
+      // numeric casts
+    if let Some(cast) = match () {
+      _ if self.try_consume("to_u24") => Some(hvm::TY_U24),
+      _ if self.try_consume("to_i24") => Some(hvm::TY_I24),
+      _ if self.try_consume("to_f24") => Some(hvm::TY_F24),
+      _ => None
+    } {
+      // Closes symbol bracket, as casts can't be partially applied
+      self.consume("]")?;
+
+      return Ok(Numb(hvm::Numb::new_sym(cast).0));
+    }
+
     // Parses the symbol
     let op = hvm::Numb::new_sym(match () {
+      // numeric operations
       _ if self.try_consume("+")   => hvm::OP_ADD,
       _ if self.try_consume("-")   => hvm::OP_SUB,
       _ if self.try_consume(":-")  => hvm::FP_SUB,
@@ -224,6 +238,11 @@ impl Numb {
     let numb = hvm::Numb(self.0);
     match numb.get_typ() {
       hvm::TY_SYM => match numb.get_sym() as hvm::Tag {
+        // casts
+        hvm::TY_U24 => "[to_u24]".to_string(),
+        hvm::TY_I24 => "[to_i24]".to_string(),
+        hvm::TY_F24 => "[to_f24]".to_string(),
+        // operations
         hvm::OP_ADD => "[+]".to_string(),
         hvm::OP_SUB => "[-]".to_string(),
         hvm::FP_SUB => "[:-]".to_string(),
diff --git a/src/hvm.rs b/src/hvm.rs
index 5d8cd9f2..6add82d1 100644
--- a/src/hvm.rs
+++ b/src/hvm.rs
@@ -266,16 +266,40 @@ impl Numb {
   }
 
   // Gets the numeric type.
-
   pub fn get_typ(&self) -> Tag {
-    return (self.0 & 0x1F) as Tag;
+    (self.0 & 0x1F) as Tag
+  }
+
+  pub fn is_num(&self) -> bool {
+    self.get_typ() >= TY_U24 && self.get_typ() <= TY_F24
   }
 
-  // Flip flag.
+  pub fn is_cast(&self) -> bool {
+    self.get_typ() == TY_SYM && self.get_sym() >= TY_U24 && self.get_sym() <= TY_F24
+  }
 
   // Partial application.
   pub fn partial(a: Self, b: Self) -> Self {
-    return Numb((b.0 & !0x1F) | a.get_sym() as u32);
+    Numb((b.0 & !0x1F) | a.get_sym() as u32)
+  }
+
+  // Cast a number to another type.
+  pub fn cast(a: Self, b: Self) -> Self {
+    match (a.get_sym(), b.get_typ()) {
+      (TY_U24, TY_U24) => b,
+      (TY_U24, TY_I24) => Self::new_u24(b.get_i24() as u32),
+      (TY_U24, TY_F24) => Self::new_u24(b.get_f24() as u32),
+
+      (TY_I24, TY_U24) => Self::new_i24(b.get_u24() as i32),
+      (TY_I24, TY_I24) => b,
+      (TY_I24, TY_F24) => Self::new_i24(b.get_f24() as i32),
+
+      (TY_F24, TY_U24) => Self::new_f24(b.get_u24() as f32),
+      (TY_F24, TY_I24) => Self::new_f24(b.get_i24() as f32),
+      (TY_F24, TY_F24) => b,
+      // invalid cast
+      (_, _) => Self::new_u24(0),
+    }
   }
 
   pub fn operate(a: Self, b: Self) -> Self {
@@ -285,6 +309,12 @@ impl Numb {
     if at == TY_SYM && bt == TY_SYM {
       return Numb::new_u24(0);
     }
+    if a.is_cast() && b.is_num() {
+      return Numb::cast(a, b);
+    }
+    if b.is_cast() && a.is_num() {
+      return Numb::cast(b, a);
+    }
     if at == TY_SYM && bt != TY_SYM {
       return Numb::partial(a, b);
     }
diff --git a/tests/programs/numeric-casts.hvm b/tests/programs/numeric-casts.hvm
new file mode 100644
index 00000000..d02de646
--- /dev/null
+++ b/tests/programs/numeric-casts.hvm
@@ -0,0 +1,34 @@
+@main = x & @tu0 ~ (* x)
+
+// casting to u24
+@tu0 = (* {n x}) & @tu1 ~ (* x) &  2.15 ~ $([to_u24] n) // 2
+@tu1 = (* {n x}) & @tu2 ~ (* x) & -1.15 ~ $([to_u24] n) // 0
+@tu2 = (* {n x}) & @tu3 ~ (* x) &  0.15 ~ $([to_u24] n) // 0
+@tu3 = (* {n x}) & @tu4 ~ (* x) & -1234 ~ $([to_u24] n) // 16775982 (reinterpreting bits)
+@tu4 = (* {n x}) & @tu5 ~ (* x) & +1234 ~ $([to_u24] n) // 1234
+@tu5 = (* {n x}) & @tu6 ~ (* x) &  4321 ~ $([to_u24] n) // 4321
+@tu6 = (* {n x}) & @tu7 ~ (* x) &  +inf ~ $([to_u24] n) // 16777215
+@tu7 = (* {n x}) & @tu8 ~ (* x) &  -inf ~ $([to_u24] n) // 0
+@tu8 = (* {n x}) & @ti0 ~ (* x) &  +NaN ~ $([to_u24] n) // 0
+
+// casting to i24
+@ti0 = (* {n x}) & @ti1 ~ (* x) &  2.15    ~ $([to_i24] n) // +2
+@ti1 = (* {n x}) & @ti2 ~ (* x) & -1.15    ~ $([to_i24] n) // -1
+@ti2 = (* {n x}) & @ti3 ~ (* x) &  0.15    ~ $([to_i24] n) // +0
+@ti3 = (* {n x}) & @ti4 ~ (* x) & -1234    ~ $([to_i24] n) // -1234
+@ti4 = (* {n x}) & @ti5 ~ (* x) & +1234    ~ $([to_i24] n) // +1234
+@ti5 = (* {n x}) & @ti6 ~ (* x) &  4321    ~ $([to_i24] n) // +4321
+@ti6 = (* {n x}) & @ti7 ~ (* x) & 16775982 ~ $([to_i24] n) // -1234 (reinterpreting bits)
+@ti7 = (* {n x}) & @ti8 ~ (* x) &  +inf    ~ $([to_i24] n) // -1 ??
+@ti8 = (* {n x}) & @ti9 ~ (* x) &  -inf    ~ $([to_i24] n) // +0 ??
+@ti9 = (* {n x}) & @tf0 ~ (* x) &  +NaN    ~ $([to_i24] n) // +0
+
+// casting to f24
+@tf0 = (* {n x}) & @tf1 ~ (* x) &  2.15    ~ $([to_f24] n) //  2.15
+@tf1 = (* {n x}) & @tf2 ~ (* x) & -1.15    ~ $([to_f24] n) // -1.15
+@tf2 = (* {n x}) & @tf3 ~ (* x) &  0.15    ~ $([to_f24] n) //  0.15
+@tf3 = (* {n x}) & @tf4 ~ (* x) & -1234    ~ $([to_f24] n) // -1234.0
+@tf4 = (* {n x}) & @tf5 ~ (* x) & +1234    ~ $([to_f24] n) // +1234.0
+@tf5 = (* {n x}) & @t   ~ (* x) & 16775982 ~ $([to_f24] n) // 16775936.0
+
+@t = *

From cb584fff5118b0e679666f83ff9def1c4c8df91c Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 12:22:36 +0200
Subject: [PATCH 02/17] numeric casts in c

---
 src/ast.rs |  2 +-
 src/hvm.c  | 35 ++++++++++++++++++++++++++++++++++-
 2 files changed, 35 insertions(+), 2 deletions(-)

diff --git a/src/ast.rs b/src/ast.rs
index dedb462d..03cf3206 100644
--- a/src/ast.rs
+++ b/src/ast.rs
@@ -283,7 +283,7 @@ impl Numb {
         } else if val.is_nan() {
           format!("+NaN")
         } else {
-          format!("{:?}", val)
+          format!("{:.7e}", val)
         }
       }
       _ => {
diff --git a/src/hvm.c b/src/hvm.c
index 28b7fafa..68d51ec4 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -21,6 +21,8 @@ typedef uint16_t u16;
 typedef  int32_t i32;
 typedef uint32_t u32;
 typedef uint64_t u64;
+typedef    float f32;
+typedef   double f64;
 
 typedef _Atomic(u8) a8;
 typedef _Atomic(u16) a16;
@@ -429,11 +431,36 @@ static inline Tag get_typ(Numb word) {
   return word & 0x1F;
 }
 
+static inline bool is_num(Numb word) {
+  return get_typ(word) >= TY_U24 && get_typ(word) <= TY_F24;
+}
+
+static inline bool is_cast(Numb word) {
+  return get_typ(word) == TY_SYM && get_sym(word) >= TY_U24 && get_sym(word) <= TY_F24;
+}
+
 // Partial application
 static inline Numb partial(Numb a, Numb b) {
   return (b & ~0x1F) | get_sym(a);
 }
 
+// Cast a number to another type.
+static inline Numb cast(Numb a, Numb b) {
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_U24) return b;
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_I24) return new_u24((u32) get_i24(b));
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_F24) return new_u24((u32) get_f24(b));
+
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_U24) return new_i24((i32) get_u24(b));
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_I24) return b;
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_F24) return new_i24((i32) get_f24(b));
+
+  if (get_sym(a) == TY_F24 && get_typ(b) == TY_U24) return new_f24((f32) get_u24(b));
+  if (get_sym(a) == TY_F24 && get_typ(b) == TY_I24) return new_f24((f32) get_i24(b));
+  if (get_sym(a) == TY_F24 && get_typ(b) == TY_F24) return b;
+
+  return new_u24(0);
+}
+
 // Operate function
 static inline Numb operate(Numb a, Numb b) {
   Tag at = get_typ(a);
@@ -441,6 +468,12 @@ static inline Numb operate(Numb a, Numb b) {
   if (at == TY_SYM && bt == TY_SYM) {
     return new_u24(0);
   }
+  if (is_cast(a) && is_num(b)) {
+    return cast(a, b);
+  }
+  if (is_cast(b) && is_num(a)) {
+    return cast(b, a);
+  }
   if (at == TY_SYM && bt != TY_SYM) {
     return partial(a, b);
   }
@@ -1776,7 +1809,7 @@ void pretty_print_numb(Numb word) {
       } else if (isnan(get_f24(word))) {
         printf("+NaN");
       } else {
-        printf("%f", get_f24(word));
+        printf("%.7e", get_f24(word));
       }
       break;
     }

From a838ef0f4600157b49bda3078fd1fde30a16a60f Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 12:24:51 +0200
Subject: [PATCH 03/17] numeric casts in cuda

---
 src/hvm.cu | 33 ++++++++++++++++++++++++++++++++-
 1 file changed, 32 insertions(+), 1 deletion(-)

diff --git a/src/hvm.cu b/src/hvm.cu
index ed8cc2ed..e5da83c0 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -540,6 +540,31 @@ __device__ __host__ inline Tag get_typ(Numb word) {
   return word & 0x1F;
 }
 
+__device__ __host__ inline bool is_num(Numb word) {
+  return get_typ(word) >= TY_U24 && get_typ(word) <= TY_F24;
+}
+
+__device__ __host__ inline bool is_cast(Numb word) {
+  return get_typ(word) == TY_SYM && get_sym(word) >= TY_U24 && get_sym(word) <= TY_F24;
+}
+
+// Cast a number to another type.
+__device__ __host__ inline Numb cast(Numb a, Numb b) {
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_U24) return b;
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_I24) return new_u24((u32) get_i24(b));
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_F24) return new_u24((u32) get_f24(b));
+
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_U24) return new_i24((i32) get_u24(b));
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_I24) return b;
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_F24) return new_i24((i32) get_f24(b));
+
+  if (get_sym(a) == TY_F24 && get_typ(b) == TY_U24) return new_f24((f32) get_u24(b));
+  if (get_sym(a) == TY_F24 && get_typ(b) == TY_I24) return new_f24((f32) get_i24(b));
+  if (get_sym(a) == TY_F24 && get_typ(b) == TY_F24) return b;
+
+  return new_u24(0);
+}
+
 // Partial application
 __device__ __host__ inline Numb partial(Numb a, Numb b) {
   return (b & ~0x1F) | get_sym(a);
@@ -552,6 +577,12 @@ __device__ __host__ inline Numb operate(Numb a, Numb b) {
   if (at == TY_SYM && bt == TY_SYM) {
     return new_u24(0);
   }
+  if (is_cast(a) && is_num(b)) {
+    return cast(a, b);
+  }
+  if (is_cast(b) && is_num(a)) {
+    return cast(b, a);
+  }
   if (at == TY_SYM && bt != TY_SYM) {
     return partial(a, b);
   }
@@ -2226,7 +2257,7 @@ __device__ void pretty_print_numb(Numb word) {
       } else if (isnan(get_f24(word))) {
         printf("+NaN");
       } else {
-        printf("%f", get_f24(word));
+        printf("%.7e", get_f24(word));
       }
       break;
     }

From 59da92348331c5772736ab9042bb598b71c55a20 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 12:46:16 +0200
Subject: [PATCH 04/17] change test file to not include UB cases

remove out of range casts as in C these are UB
---
 tests/programs/numeric-casts.hvm | 49 ++++++++++++++------------------
 1 file changed, 22 insertions(+), 27 deletions(-)

diff --git a/tests/programs/numeric-casts.hvm b/tests/programs/numeric-casts.hvm
index d02de646..e515c703 100644
--- a/tests/programs/numeric-casts.hvm
+++ b/tests/programs/numeric-casts.hvm
@@ -1,34 +1,29 @@
 @main = x & @tu0 ~ (* x)
 
-// casting to u24
-@tu0 = (* {n x}) & @tu1 ~ (* x) &  2.15 ~ $([to_u24] n) // 2
-@tu1 = (* {n x}) & @tu2 ~ (* x) & -1.15 ~ $([to_u24] n) // 0
-@tu2 = (* {n x}) & @tu3 ~ (* x) &  0.15 ~ $([to_u24] n) // 0
-@tu3 = (* {n x}) & @tu4 ~ (* x) & -1234 ~ $([to_u24] n) // 16775982 (reinterpreting bits)
-@tu4 = (* {n x}) & @tu5 ~ (* x) & +1234 ~ $([to_u24] n) // 1234
-@tu5 = (* {n x}) & @tu6 ~ (* x) &  4321 ~ $([to_u24] n) // 4321
-@tu6 = (* {n x}) & @tu7 ~ (* x) &  +inf ~ $([to_u24] n) // 16777215
-@tu7 = (* {n x}) & @tu8 ~ (* x) &  -inf ~ $([to_u24] n) // 0
-@tu8 = (* {n x}) & @ti0 ~ (* x) &  +NaN ~ $([to_u24] n) // 0
+// casting to u24 (casts from a value of of range of u24 is UB)
+@tu0 = (* {n x}) & @tu1 ~ (* x) &  2.55 ~ $([to_u24] n) // 3
+@tu1 = (* {n x}) & @tu2 ~ (* x) &  0.15 ~ $([to_u24] n) // 0
+@tu2 = (* {n x}) & @tu3 ~ (* x) & +1234 ~ $([to_u24] n) // 1234
+@tu3 = (* {n x}) & @ti0 ~ (* x) &  4321 ~ $([to_u24] n) // 4321
 
-// casting to i24
-@ti0 = (* {n x}) & @ti1 ~ (* x) &  2.15    ~ $([to_i24] n) // +2
-@ti1 = (* {n x}) & @ti2 ~ (* x) & -1.15    ~ $([to_i24] n) // -1
-@ti2 = (* {n x}) & @ti3 ~ (* x) &  0.15    ~ $([to_i24] n) // +0
-@ti3 = (* {n x}) & @ti4 ~ (* x) & -1234    ~ $([to_i24] n) // -1234
-@ti4 = (* {n x}) & @ti5 ~ (* x) & +1234    ~ $([to_i24] n) // +1234
-@ti5 = (* {n x}) & @ti6 ~ (* x) &  4321    ~ $([to_i24] n) // +4321
-@ti6 = (* {n x}) & @ti7 ~ (* x) & 16775982 ~ $([to_i24] n) // -1234 (reinterpreting bits)
-@ti7 = (* {n x}) & @ti8 ~ (* x) &  +inf    ~ $([to_i24] n) // -1 ??
-@ti8 = (* {n x}) & @ti9 ~ (* x) &  -inf    ~ $([to_i24] n) // +0 ??
-@ti9 = (* {n x}) & @tf0 ~ (* x) &  +NaN    ~ $([to_i24] n) // +0
+// casting to i24 (casts from a value out of range of i24 is UB)
+@ti0 = (* {n x}) & @ti1 ~ (* x) &  2.55 ~ $([to_i24] n) // +3
+@ti1 = (* {n x}) & @ti2 ~ (* x) & -2.15 ~ $([to_i24] n) // -1
+@ti2 = (* {n x}) & @ti3 ~ (* x) &  0.15 ~ $([to_i24] n) // +0
+@ti3 = (* {n x}) & @ti4 ~ (* x) & -1234 ~ $([to_i24] n) // -1234
+@ti4 = (* {n x}) & @ti5 ~ (* x) & +1234 ~ $([to_i24] n) // +1234
+@ti5 = (* {n x}) & @tf0 ~ (* x) &  4321 ~ $([to_i24] n) // +4321
 
 // casting to f24
-@tf0 = (* {n x}) & @tf1 ~ (* x) &  2.15    ~ $([to_f24] n) //  2.15
-@tf1 = (* {n x}) & @tf2 ~ (* x) & -1.15    ~ $([to_f24] n) // -1.15
-@tf2 = (* {n x}) & @tf3 ~ (* x) &  0.15    ~ $([to_f24] n) //  0.15
-@tf3 = (* {n x}) & @tf4 ~ (* x) & -1234    ~ $([to_f24] n) // -1234.0
-@tf4 = (* {n x}) & @tf5 ~ (* x) & +1234    ~ $([to_f24] n) // +1234.0
-@tf5 = (* {n x}) & @t   ~ (* x) & 16775982 ~ $([to_f24] n) // 16775936.0
+@tf0 = (* {n x}) & @tf1 ~ (* x) &  +NaN    ~ $([to_f24] n) // +NaN
+@tf1 = (* {n x}) & @tf2 ~ (* x) &  +inf    ~ $([to_f24] n) // +inf
+@tf2 = (* {n x}) & @tf3 ~ (* x) &  -inf    ~ $([to_f24] n) // -inf
+@tf3 = (* {n x}) & @tf4 ~ (* x) &  2.15    ~ $([to_f24] n) //  2.15
+@tf4 = (* {n x}) & @tf5 ~ (* x) & -2.15    ~ $([to_f24] n) // -1.15
+@tf5 = (* {n x}) & @tf6 ~ (* x) &  0.15    ~ $([to_f24] n) //  0.15
+@tf6 = (* {n x}) & @tf7 ~ (* x) & -1234    ~ $([to_f24] n) // -1234.0
+@tf7 = (* {n x}) & @tf8 ~ (* x) & +1234    ~ $([to_f24] n) // +1234.0
+@tf8 = (* {n x}) & @tf9 ~ (* x) & 123456   ~ $([to_f24] n)   // 123456.0
+@tf9 = (* {n x}) & @t   ~ (* x) & 16775982 ~ $([to_f24] n) // 16775936.0
 
 @t = *

From 94975f71ab384d5aff12592e3f2e9bc04e70d952 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 12:49:03 +0200
Subject: [PATCH 05/17] typos

---
 tests/programs/numeric-casts.hvm | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/tests/programs/numeric-casts.hvm b/tests/programs/numeric-casts.hvm
index e515c703..7fb37fef 100644
--- a/tests/programs/numeric-casts.hvm
+++ b/tests/programs/numeric-casts.hvm
@@ -19,11 +19,11 @@
 @tf1 = (* {n x}) & @tf2 ~ (* x) &  +inf    ~ $([to_f24] n) // +inf
 @tf2 = (* {n x}) & @tf3 ~ (* x) &  -inf    ~ $([to_f24] n) // -inf
 @tf3 = (* {n x}) & @tf4 ~ (* x) &  2.15    ~ $([to_f24] n) //  2.15
-@tf4 = (* {n x}) & @tf5 ~ (* x) & -2.15    ~ $([to_f24] n) // -1.15
+@tf4 = (* {n x}) & @tf5 ~ (* x) & -2.15    ~ $([to_f24] n) // -2.15
 @tf5 = (* {n x}) & @tf6 ~ (* x) &  0.15    ~ $([to_f24] n) //  0.15
 @tf6 = (* {n x}) & @tf7 ~ (* x) & -1234    ~ $([to_f24] n) // -1234.0
 @tf7 = (* {n x}) & @tf8 ~ (* x) & +1234    ~ $([to_f24] n) // +1234.0
-@tf8 = (* {n x}) & @tf9 ~ (* x) & 123456   ~ $([to_f24] n)   // 123456.0
+@tf8 = (* {n x}) & @tf9 ~ (* x) & 123456   ~ $([to_f24] n) // 123456.0
 @tf9 = (* {n x}) & @t   ~ (* x) & 16775982 ~ $([to_f24] n) // 16775936.0
 
 @t = *

From 51afffd0d785bc67ac4191a122fa7b9ace28c68f Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 13:01:28 +0200
Subject: [PATCH 06/17] snapshots, fix numeric-casts.hvm comments

---
 tests/programs/numeric-casts.hvm                 | 6 +++---
 tests/snapshots/run__file@f24.hvm.snap           | 2 +-
 tests/snapshots/run__file@numeric-casts.hvm.snap | 6 ++++++
 3 files changed, 10 insertions(+), 4 deletions(-)
 create mode 100644 tests/snapshots/run__file@numeric-casts.hvm.snap

diff --git a/tests/programs/numeric-casts.hvm b/tests/programs/numeric-casts.hvm
index 7fb37fef..290863d5 100644
--- a/tests/programs/numeric-casts.hvm
+++ b/tests/programs/numeric-casts.hvm
@@ -1,14 +1,14 @@
 @main = x & @tu0 ~ (* x)
 
 // casting to u24 (casts from a value of of range of u24 is UB)
-@tu0 = (* {n x}) & @tu1 ~ (* x) &  2.55 ~ $([to_u24] n) // 3
+@tu0 = (* {n x}) & @tu1 ~ (* x) &  2.55 ~ $([to_u24] n) // 2
 @tu1 = (* {n x}) & @tu2 ~ (* x) &  0.15 ~ $([to_u24] n) // 0
 @tu2 = (* {n x}) & @tu3 ~ (* x) & +1234 ~ $([to_u24] n) // 1234
 @tu3 = (* {n x}) & @ti0 ~ (* x) &  4321 ~ $([to_u24] n) // 4321
 
 // casting to i24 (casts from a value out of range of i24 is UB)
-@ti0 = (* {n x}) & @ti1 ~ (* x) &  2.55 ~ $([to_i24] n) // +3
-@ti1 = (* {n x}) & @ti2 ~ (* x) & -2.15 ~ $([to_i24] n) // -1
+@ti0 = (* {n x}) & @ti1 ~ (* x) &  2.55 ~ $([to_i24] n) // +2
+@ti1 = (* {n x}) & @ti2 ~ (* x) & -2.15 ~ $([to_i24] n) // -2
 @ti2 = (* {n x}) & @ti3 ~ (* x) &  0.15 ~ $([to_i24] n) // +0
 @ti3 = (* {n x}) & @ti4 ~ (* x) & -1234 ~ $([to_i24] n) // -1234
 @ti4 = (* {n x}) & @ti5 ~ (* x) & +1234 ~ $([to_i24] n) // +1234
diff --git a/tests/snapshots/run__file@f24.hvm.snap b/tests/snapshots/run__file@f24.hvm.snap
index f528dbd4..99bc3bad 100644
--- a/tests/snapshots/run__file@f24.hvm.snap
+++ b/tests/snapshots/run__file@f24.hvm.snap
@@ -3,4 +3,4 @@ source: tests/run.rs
 expression: rust_output
 input_file: tests/programs/f24.hvm
 ---
-Result: {+inf {-inf {+NaN {2.5 {-1.5 {1.1499939 {0.25 {0.5 {0 {1 {1 {0 {0 {0 {0 {+NaN {+inf {-inf {1.019989 {0.1000061 {0.1000061 {-0.1000061 {-0.1000061 *}}}}}}}}}}}}}}}}}}}}}}}
+Result: {+inf {-inf {+NaN {2.5000000e0 {-1.5000000e0 {1.1499939e0 {2.5000000e-1 {5.0000000e-1 {0 {1 {1 {0 {0 {0 {0 {+NaN {+inf {-inf {1.0199890e0 {1.0000610e-1 {1.0000610e-1 {-1.0000610e-1 {-1.0000610e-1 *}}}}}}}}}}}}}}}}}}}}}}}
diff --git a/tests/snapshots/run__file@numeric-casts.hvm.snap b/tests/snapshots/run__file@numeric-casts.hvm.snap
new file mode 100644
index 00000000..b6a80202
--- /dev/null
+++ b/tests/snapshots/run__file@numeric-casts.hvm.snap
@@ -0,0 +1,6 @@
+---
+source: tests/run.rs
+expression: rust_output
+input_file: tests/programs/numeric-casts.hvm
+---
+Result: {2 {0 {1234 {4321 {+2 {-2 {+0 {-1234 {+1234 {+4321 {+NaN {+inf {-inf {2.1500244e0 {-2.1500244e0 {1.5000153e-1 {-1.2340000e3 {1.2340000e3 {1.2345600e5 {1.6775936e7 *}}}}}}}}}}}}}}}}}}}}

From eccb19eb34ec421202c47389dc3c0f273e080ba4 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 13:05:49 +0200
Subject: [PATCH 07/17] merge main, snapshot review

---
 tests/snapshots/run__file@empty.hvm.snap | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/tests/snapshots/run__file@empty.hvm.snap b/tests/snapshots/run__file@empty.hvm.snap
index b5b4735e..953bfae9 100644
--- a/tests/snapshots/run__file@empty.hvm.snap
+++ b/tests/snapshots/run__file@empty.hvm.snap
@@ -4,6 +4,6 @@ expression: rust_output
 input_file: tests/programs/empty.hvm
 ---
 exit status: 101
-thread 'main' panicked at src/ast.rs:508:41:
+thread 'main' panicked at src/ast.rs:527:41:
 missing `@main` definition
 note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace

From 3853c6902d6cd62aaefb2dde45c2e6d5b423a135 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 13:51:40 +0200
Subject: [PATCH 08/17] comments

---
 src/ast.rs | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/src/ast.rs b/src/ast.rs
index 538d354e..17276093 100644
--- a/src/ast.rs
+++ b/src/ast.rs
@@ -43,14 +43,14 @@ impl<'i> CoreParser<'i> {
   pub fn parse_numb_sym(&mut self) -> Result<Numb, String> {
     self.consume("[")?;
 
-      // numeric casts
+    // numeric casts
     if let Some(cast) = match () {
       _ if self.try_consume("to_u24") => Some(hvm::TY_U24),
       _ if self.try_consume("to_i24") => Some(hvm::TY_I24),
       _ if self.try_consume("to_f24") => Some(hvm::TY_F24),
       _ => None
     } {
-      // Closes symbol bracket, as casts can't be partially applied
+      // Casts can't be partially applied, so nothing should follow.
       self.consume("]")?;
 
       return Ok(Numb(hvm::Numb::new_sym(cast).0));

From baf7e2388a16147189df92439ac237b6a9128de9 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Thu, 30 May 2024 07:22:13 +0200
Subject: [PATCH 09/17] define behavior for all numeric casts

---
 src/hvm.c                                     | 42 ++++++++++++++++---
 src/hvm.cu                                    | 42 ++++++++++++++++---
 src/hvm.rs                                    |  4 +-
 tests/programs/numeric-casts.hvm              | 35 ++++++++++------
 .../run__file@numeric-casts.hvm.snap          |  2 +-
 5 files changed, 100 insertions(+), 25 deletions(-)

diff --git a/src/hvm.c b/src/hvm.c
index 68d51ec4..1562acac 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -280,10 +280,15 @@ static inline void swap(Port *a, Port *b) {
   Port x = *a; *a = *b; *b = x;
 }
 
-u32 min(u32 a, u32 b) {
+inline u32 min(u32 a, u32 b) {
   return (a < b) ? a : b;
 }
 
+inline f32 clamp(f32 x, f32 min, f32 max) {
+  const f32 t = x < min ? min : x;
+  return (t > max) ? max : t;
+}
+
 // A simple spin-wait barrier using atomic operations
 a64 a_reached = 0; // number of threads that reached the current barrier
 a64 a_barrier = 0; // number of barriers passed during this program
@@ -445,14 +450,41 @@ static inline Numb partial(Numb a, Numb b) {
 }
 
 // Cast a number to another type.
+// The semantics are meant to spiritually resemble rust's numeric casts:
+// - i24 <-> u24: is just reinterpretation of bits
+// - f24  -> i24,
+//   f24  -> u24: casts to the "closest" integer representing this float,
+//                saturating if out of range and 0 if NaN
+// - i24  -> f24,
+//   u24  -> f24: casts to the "closest" float representing this integer.
 static inline Numb cast(Numb a, Numb b) {
   if (get_sym(a) == TY_U24 && get_typ(b) == TY_U24) return b;
-  if (get_sym(a) == TY_U24 && get_typ(b) == TY_I24) return new_u24((u32) get_i24(b));
-  if (get_sym(a) == TY_U24 && get_typ(b) == TY_F24) return new_u24((u32) get_f24(b));
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_I24) {
+    // reinterpret bits
+    i32 val = get_i24(b);
+    return new_u24(*(u32*) &val);
+  }
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_F24) {
+    f32 val = get_f24(b);
+    if (isnan(val)) {
+      return new_u24(0);
+    }
+    return new_u24((u32) clamp(val, 0.0, 16777215));
+  }
 
-  if (get_sym(a) == TY_I24 && get_typ(b) == TY_U24) return new_i24((i32) get_u24(b));
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_U24) {
+    // reinterpret bits
+    u32 val = get_u24(b);
+    return new_i24(*(i32*) &val);
+  }
   if (get_sym(a) == TY_I24 && get_typ(b) == TY_I24) return b;
-  if (get_sym(a) == TY_I24 && get_typ(b) == TY_F24) return new_i24((i32) get_f24(b));
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_F24) {
+    f32 val = get_f24(b);
+    if (isnan(val)) {
+      return new_i24(0);
+    }
+    return new_i24((i32) clamp(val, -8388608.0, 8388607.0));
+  }
 
   if (get_sym(a) == TY_F24 && get_typ(b) == TY_U24) return new_f24((f32) get_u24(b));
   if (get_sym(a) == TY_F24 && get_typ(b) == TY_I24) return new_f24((f32) get_i24(b));
diff --git a/src/hvm.cu b/src/hvm.cu
index e5da83c0..92eea98d 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -261,6 +261,11 @@ __global__ void print_heatmap(GNet* gnet, u32 turn);
 // Utils
 // -----
 
+inline f32 clamp(f32 x, f32 min, f32 max) {
+  const f32 t = x < min ? min : x;
+  return (t > max) ? max : t;
+}
+
 // TODO: write a time64() function that returns the time as fast as possible as a u64
 static inline u64 time64() {
   struct timespec ts;
@@ -549,14 +554,41 @@ __device__ __host__ inline bool is_cast(Numb word) {
 }
 
 // Cast a number to another type.
-__device__ __host__ inline Numb cast(Numb a, Numb b) {
+// The semantics are meant to spiritually resemble rust's numeric casts:
+// - i24 <-> u24: is just reinterpretation of bits
+// - f24  -> i24,
+//   f24  -> u24: casts to the "closest" integer representing this float,
+//                saturating if out of range and 0 if NaN
+// - i24  -> f24,
+//   u24  -> f24: casts to the "closest" float representing this integer.
+static inline Numb cast(Numb a, Numb b) {
   if (get_sym(a) == TY_U24 && get_typ(b) == TY_U24) return b;
-  if (get_sym(a) == TY_U24 && get_typ(b) == TY_I24) return new_u24((u32) get_i24(b));
-  if (get_sym(a) == TY_U24 && get_typ(b) == TY_F24) return new_u24((u32) get_f24(b));
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_I24) {
+    // reinterpret bits
+    i32 val = get_i24(b);
+    return new_u24(*(u32*) &val);
+  }
+  if (get_sym(a) == TY_U24 && get_typ(b) == TY_F24) {
+    f32 val = get_f24(b);
+    if (isnan(val)) {
+      return new_u24(0);
+    }
+    return new_u24((u32) clamp(val, 0.0, 16777215));
+  }
 
-  if (get_sym(a) == TY_I24 && get_typ(b) == TY_U24) return new_i24((i32) get_u24(b));
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_U24) {
+    // reinterpret bits
+    u32 val = get_u24(b);
+    return new_i24(*(i32*) &val);
+  }
   if (get_sym(a) == TY_I24 && get_typ(b) == TY_I24) return b;
-  if (get_sym(a) == TY_I24 && get_typ(b) == TY_F24) return new_i24((i32) get_f24(b));
+  if (get_sym(a) == TY_I24 && get_typ(b) == TY_F24) {
+    f32 val = get_f24(b);
+    if (isnan(val)) {
+      return new_i24(0);
+    }
+    return new_i24((i32) clamp(val, -8388608.0, 8388607.0));
+  }
 
   if (get_sym(a) == TY_F24 && get_typ(b) == TY_U24) return new_f24((f32) get_u24(b));
   if (get_sym(a) == TY_F24 && get_typ(b) == TY_I24) return new_f24((f32) get_i24(b));
diff --git a/src/hvm.rs b/src/hvm.rs
index 6add82d1..07333886 100644
--- a/src/hvm.rs
+++ b/src/hvm.rs
@@ -288,11 +288,11 @@ impl Numb {
     match (a.get_sym(), b.get_typ()) {
       (TY_U24, TY_U24) => b,
       (TY_U24, TY_I24) => Self::new_u24(b.get_i24() as u32),
-      (TY_U24, TY_F24) => Self::new_u24(b.get_f24() as u32),
+      (TY_U24, TY_F24) => Self::new_u24(b.get_f24().clamp(0.0, 16777215.0) as u32),
 
       (TY_I24, TY_U24) => Self::new_i24(b.get_u24() as i32),
       (TY_I24, TY_I24) => b,
-      (TY_I24, TY_F24) => Self::new_i24(b.get_f24() as i32),
+      (TY_I24, TY_F24) => Self::new_i24(b.get_f24().clamp(-8388608.0, 8388607.0) as i32),
 
       (TY_F24, TY_U24) => Self::new_f24(b.get_u24() as f32),
       (TY_F24, TY_I24) => Self::new_f24(b.get_i24() as f32),
diff --git a/tests/programs/numeric-casts.hvm b/tests/programs/numeric-casts.hvm
index 290863d5..d01a1ebf 100644
--- a/tests/programs/numeric-casts.hvm
+++ b/tests/programs/numeric-casts.hvm
@@ -1,18 +1,29 @@
 @main = x & @tu0 ~ (* x)
 
-// casting to u24 (casts from a value of of range of u24 is UB)
-@tu0 = (* {n x}) & @tu1 ~ (* x) &  2.55 ~ $([to_u24] n) // 2
-@tu1 = (* {n x}) & @tu2 ~ (* x) &  0.15 ~ $([to_u24] n) // 0
-@tu2 = (* {n x}) & @tu3 ~ (* x) & +1234 ~ $([to_u24] n) // 1234
-@tu3 = (* {n x}) & @ti0 ~ (* x) &  4321 ~ $([to_u24] n) // 4321
+// casting to u24
+@tu0 = (* {n x}) & @tu1 ~ (* x) &  0          ~ $([to_u24] n) // 0
+@tu1 = (* {n x}) & @tu2 ~ (* x) &  1234       ~ $([to_u24] n) // 1234
+@tu2 = (* {n x}) & @tu3 ~ (* x) & +4321       ~ $([to_u24] n) // 4321
+@tu3 = (* {n x}) & @tu4 ~ (* x) & -5678       ~ $([to_u24] n) // 16771538 (reinterprets bits)
+@tu4 = (* {n x}) & @tu5 ~ (* x) &  2.8        ~ $([to_u24] n) // 2 (rounds to zero)
+@tu5 = (* {n x}) & @tu6 ~ (* x) & -12.5       ~ $([to_u24] n) // 0 (saturates)
+@tu6 = (* {n x}) & @tu7 ~ (* x) &  16777216.0 ~ $([to_u24] n) // 16777215 (saturates)
+@tu7 = (* {n x}) & @tu8 ~ (* x) & +inf        ~ $([to_u24] n) // 16777215 (saturates)
+@tu8 = (* {n x}) & @tu9 ~ (* x) & -inf        ~ $([to_u24] n) // 0 (saturates)
+@tu9 = (* {n x}) & @ti0 ~ (* x) & +NaN        ~ $([to_u24] n) // 0
 
-// casting to i24 (casts from a value out of range of i24 is UB)
-@ti0 = (* {n x}) & @ti1 ~ (* x) &  2.55 ~ $([to_i24] n) // +2
-@ti1 = (* {n x}) & @ti2 ~ (* x) & -2.15 ~ $([to_i24] n) // -2
-@ti2 = (* {n x}) & @ti3 ~ (* x) &  0.15 ~ $([to_i24] n) // +0
-@ti3 = (* {n x}) & @ti4 ~ (* x) & -1234 ~ $([to_i24] n) // -1234
-@ti4 = (* {n x}) & @ti5 ~ (* x) & +1234 ~ $([to_i24] n) // +1234
-@ti5 = (* {n x}) & @tf0 ~ (* x) &  4321 ~ $([to_i24] n) // +4321
+// casting to i24
+@ti0 = (* {n x}) & @ti1 ~ (* x) &  0          ~ $([to_i24] n) // +0
+@ti1 = (* {n x}) & @ti2 ~ (* x) &  1234       ~ $([to_i24] n) // +1234
+@ti2 = (* {n x}) & @ti3 ~ (* x) & +4321       ~ $([to_i24] n) // +4321
+@ti3 = (* {n x}) & @ti4 ~ (* x) & -5678       ~ $([to_i24] n) // -5678
+@ti4 = (* {n x}) & @ti5 ~ (* x) &  2.8        ~ $([to_i24] n) // +2   (rounds to zero)
+@ti5 = (* {n x}) & @ti6 ~ (* x) & -12.7       ~ $([to_i24] n) // -12 (rounds to zero)
+@ti6 = (* {n x}) & @ti7 ~ (* x) &  8388610.0  ~ $([to_i24] n) // +8388607 (saturates)
+@ti7 = (* {n x}) & @ti8 ~ (* x) & -8388610.0  ~ $([to_i24] n) // -8388608 (saturates)
+@ti8 = (* {n x}) & @ti9 ~ (* x) & +inf        ~ $([to_i24] n) // +8388607 (saturates)
+@ti9 = (* {n x}) & @ti10 ~ (* x) & -inf        ~ $([to_i24] n) // -8388608 (saturates)
+@ti10 = (* {n x}) & @tf0 ~ (* x) & +NaN        ~ $([to_i24] n) // +0
 
 // casting to f24
 @tf0 = (* {n x}) & @tf1 ~ (* x) &  +NaN    ~ $([to_f24] n) // +NaN
diff --git a/tests/snapshots/run__file@numeric-casts.hvm.snap b/tests/snapshots/run__file@numeric-casts.hvm.snap
index b6a80202..df75977d 100644
--- a/tests/snapshots/run__file@numeric-casts.hvm.snap
+++ b/tests/snapshots/run__file@numeric-casts.hvm.snap
@@ -3,4 +3,4 @@ source: tests/run.rs
 expression: rust_output
 input_file: tests/programs/numeric-casts.hvm
 ---
-Result: {2 {0 {1234 {4321 {+2 {-2 {+0 {-1234 {+1234 {+4321 {+NaN {+inf {-inf {2.1500244e0 {-2.1500244e0 {1.5000153e-1 {-1.2340000e3 {1.2340000e3 {1.2345600e5 {1.6775936e7 *}}}}}}}}}}}}}}}}}}}}
+Result: {0 {1234 {4321 {16771538 {2 {0 {16777215 {16777215 {0 {0 {+0 {+1234 {+4321 {-5678 {+2 {-12 {+8388607 {-8388608 {+8388607 {-8388608 {+0 {+NaN {+inf {-inf {2.1500244e0 {-2.1500244e0 {1.5000153e-1 {-1.2340000e3 {1.2340000e3 {1.2345600e5 {1.6775936e7 *}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}

From a51bcf2f10d78757f9d397d577d2460ebe90a74a Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Thu, 30 May 2024 07:23:19 +0200
Subject: [PATCH 10/17] comment

---
 src/hvm.rs | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/src/hvm.rs b/src/hvm.rs
index 07333886..c8f0c404 100644
--- a/src/hvm.rs
+++ b/src/hvm.rs
@@ -284,6 +284,13 @@ impl Numb {
   }
 
   // Cast a number to another type.
+  // The semantics are meant to spiritually resemble rust's numeric casts:
+  // - i24 <-> u24: is just reinterpretation of bits
+  // - f24  -> i24,
+  //   f24  -> u24: casts to the "closest" integer representing this float,
+  //                saturating if out of range and 0 if NaN
+  // - i24  -> f24,
+  //   u24  -> f24: casts to the "closest" float representing this integer.
   pub fn cast(a: Self, b: Self) -> Self {
     match (a.get_sym(), b.get_typ()) {
       (TY_U24, TY_U24) => b,

From 934d4a06602ad5fbf9305e99911bfbec385395f5 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Thu, 30 May 2024 07:50:22 +0200
Subject: [PATCH 11/17] fix cuda

---
 src/hvm.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/src/hvm.cu b/src/hvm.cu
index 92eea98d..c9bf18f8 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -261,7 +261,7 @@ __global__ void print_heatmap(GNet* gnet, u32 turn);
 // Utils
 // -----
 
-inline f32 clamp(f32 x, f32 min, f32 max) {
+__device__ __host__ f32 clamp(f32 x, f32 min, f32 max) {
   const f32 t = x < min ? min : x;
   return (t > max) ? max : t;
 }
@@ -561,7 +561,7 @@ __device__ __host__ inline bool is_cast(Numb word) {
 //                saturating if out of range and 0 if NaN
 // - i24  -> f24,
 //   u24  -> f24: casts to the "closest" float representing this integer.
-static inline Numb cast(Numb a, Numb b) {
+__device__ __host__ inline Numb cast(Numb a, Numb b) {
   if (get_sym(a) == TY_U24 && get_typ(b) == TY_U24) return b;
   if (get_sym(a) == TY_U24 && get_typ(b) == TY_I24) {
     // reinterpret bits

From ed5936beb6519d17ea1d6ec558639e3f5fd96321 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 06:22:54 +0200
Subject: [PATCH 12/17] constants, clamp ints on rust impl

---
 src/ast.rs | 2 +-
 src/hvm.c  | 8 ++++++--
 src/hvm.rs | 8 ++++++--
 3 files changed, 13 insertions(+), 5 deletions(-)

diff --git a/src/ast.rs b/src/ast.rs
index 17276093..378ece56 100644
--- a/src/ast.rs
+++ b/src/ast.rs
@@ -283,7 +283,7 @@ impl Numb {
         } else if val.is_nan() {
           format!("+NaN")
         } else {
-          format!("{:.7e}", val)
+          format!("{:?}", val)
         }
       }
       _ => {
diff --git a/src/hvm.c b/src/hvm.c
index 1562acac..3a4cda49 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -77,6 +77,10 @@ typedef u32 Numb; // Numb ::= 29-bit (rounded up to u32)
 #define SWIT 0x7
 
 // Numbers
+static const f32 U24_MAX = (f32) (1 << 24) - 1;
+static const f32 U24_MIN = 0.0;
+static const f32 I24_MAX = (f32) (1 << 23) - 1;
+static const f32 I24_MIN = (f32) (i32) ((-1u) << 23);
 #define TY_SYM 0x00
 #define TY_U24 0x01
 #define TY_I24 0x02
@@ -469,7 +473,7 @@ static inline Numb cast(Numb a, Numb b) {
     if (isnan(val)) {
       return new_u24(0);
     }
-    return new_u24((u32) clamp(val, 0.0, 16777215));
+    return new_u24((u32) clamp(val, U24_MIN, U24_MAX));
   }
 
   if (get_sym(a) == TY_I24 && get_typ(b) == TY_U24) {
@@ -483,7 +487,7 @@ static inline Numb cast(Numb a, Numb b) {
     if (isnan(val)) {
       return new_i24(0);
     }
-    return new_i24((i32) clamp(val, -8388608.0, 8388607.0));
+    return new_i24((i32) clamp(val, I24_MIN, I24_MAX));
   }
 
   if (get_sym(a) == TY_F24 && get_typ(b) == TY_U24) return new_f24((f32) get_u24(b));
diff --git a/src/hvm.rs b/src/hvm.rs
index c8f0c404..b494822d 100644
--- a/src/hvm.rs
+++ b/src/hvm.rs
@@ -25,6 +25,10 @@ pub struct APair(pub AtomicU64);
 
 // Number
 pub struct Numb(pub Val);
+const U24_MAX : u32 = 1 << 24 - 1;
+const U24_MIN : u32 = 0;
+const I24_MAX : i32 = 1 << 23 - 1;
+const I24_MIN : i32 = -1 << 23;
 
 // Tags
 pub const VAR : Tag = 0x0; // variable
@@ -295,11 +299,11 @@ impl Numb {
     match (a.get_sym(), b.get_typ()) {
       (TY_U24, TY_U24) => b,
       (TY_U24, TY_I24) => Self::new_u24(b.get_i24() as u32),
-      (TY_U24, TY_F24) => Self::new_u24(b.get_f24().clamp(0.0, 16777215.0) as u32),
+      (TY_U24, TY_F24) => Self::new_u24((b.get_f24() as u32).clamp(U24_MIN, U24_MAX)),
 
       (TY_I24, TY_U24) => Self::new_i24(b.get_u24() as i32),
       (TY_I24, TY_I24) => b,
-      (TY_I24, TY_F24) => Self::new_i24(b.get_f24().clamp(-8388608.0, 8388607.0) as i32),
+      (TY_I24, TY_F24) => Self::new_i24((b.get_f24() as i32).clamp(I24_MIN, I24_MAX)),
 
       (TY_F24, TY_U24) => Self::new_f24(b.get_u24() as f32),
       (TY_F24, TY_I24) => Self::new_f24(b.get_i24() as f32),

From 55e0ce83e54fd61acb3e386ac29a4810cc38bbdc Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 06:26:43 +0200
Subject: [PATCH 13/17] cargo insta review

---
 tests/snapshots/run__file@f24.hvm.snap           | 2 +-
 tests/snapshots/run__file@numeric-casts.hvm.snap | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/tests/snapshots/run__file@f24.hvm.snap b/tests/snapshots/run__file@f24.hvm.snap
index 99bc3bad..f528dbd4 100644
--- a/tests/snapshots/run__file@f24.hvm.snap
+++ b/tests/snapshots/run__file@f24.hvm.snap
@@ -3,4 +3,4 @@ source: tests/run.rs
 expression: rust_output
 input_file: tests/programs/f24.hvm
 ---
-Result: {+inf {-inf {+NaN {2.5000000e0 {-1.5000000e0 {1.1499939e0 {2.5000000e-1 {5.0000000e-1 {0 {1 {1 {0 {0 {0 {0 {+NaN {+inf {-inf {1.0199890e0 {1.0000610e-1 {1.0000610e-1 {-1.0000610e-1 {-1.0000610e-1 *}}}}}}}}}}}}}}}}}}}}}}}
+Result: {+inf {-inf {+NaN {2.5 {-1.5 {1.1499939 {0.25 {0.5 {0 {1 {1 {0 {0 {0 {0 {+NaN {+inf {-inf {1.019989 {0.1000061 {0.1000061 {-0.1000061 {-0.1000061 *}}}}}}}}}}}}}}}}}}}}}}}
diff --git a/tests/snapshots/run__file@numeric-casts.hvm.snap b/tests/snapshots/run__file@numeric-casts.hvm.snap
index df75977d..0255c0ef 100644
--- a/tests/snapshots/run__file@numeric-casts.hvm.snap
+++ b/tests/snapshots/run__file@numeric-casts.hvm.snap
@@ -3,4 +3,4 @@ source: tests/run.rs
 expression: rust_output
 input_file: tests/programs/numeric-casts.hvm
 ---
-Result: {0 {1234 {4321 {16771538 {2 {0 {16777215 {16777215 {0 {0 {+0 {+1234 {+4321 {-5678 {+2 {-12 {+8388607 {-8388608 {+8388607 {-8388608 {+0 {+NaN {+inf {-inf {2.1500244e0 {-2.1500244e0 {1.5000153e-1 {-1.2340000e3 {1.2340000e3 {1.2345600e5 {1.6775936e7 *}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}
+Result: {0 {1234 {4321 {16771538 {2 {0 {8388608 {8388608 {0 {0 {+0 {+1234 {+4321 {-5678 {+2 {-12 {+4194304 {-8388608 {+4194304 {-8388608 {+0 {+NaN {+inf {-inf {2.1500244 {-2.1500244 {0.15000153 {-1234.0 {1234.0 {123456.0 {16775936.0 *}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}

From 53099d0ada4d58dec6cbf18be90f521a7621c88d Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 06:33:09 +0200
Subject: [PATCH 14/17] shift precedence

---
 src/hvm.rs | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/src/hvm.rs b/src/hvm.rs
index b494822d..d7007aad 100644
--- a/src/hvm.rs
+++ b/src/hvm.rs
@@ -25,10 +25,10 @@ pub struct APair(pub AtomicU64);
 
 // Number
 pub struct Numb(pub Val);
-const U24_MAX : u32 = 1 << 24 - 1;
+const U24_MAX : u32 = (1 << 24) - 1;
 const U24_MIN : u32 = 0;
-const I24_MAX : i32 = 1 << 23 - 1;
-const I24_MIN : i32 = -1 << 23;
+const I24_MAX : i32 = (1 << 23) - 1;
+const I24_MIN : i32 = (-1) << 23;
 
 // Tags
 pub const VAR : Tag = 0x0; // variable

From c34c4e20918f7cba553370c7ec766935d6848375 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 06:37:51 +0200
Subject: [PATCH 15/17] snapshots

---
 tests/snapshots/run__file@numeric-casts.hvm.snap | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/tests/snapshots/run__file@numeric-casts.hvm.snap b/tests/snapshots/run__file@numeric-casts.hvm.snap
index 0255c0ef..30ff9a5b 100644
--- a/tests/snapshots/run__file@numeric-casts.hvm.snap
+++ b/tests/snapshots/run__file@numeric-casts.hvm.snap
@@ -3,4 +3,4 @@ source: tests/run.rs
 expression: rust_output
 input_file: tests/programs/numeric-casts.hvm
 ---
-Result: {0 {1234 {4321 {16771538 {2 {0 {8388608 {8388608 {0 {0 {+0 {+1234 {+4321 {-5678 {+2 {-12 {+4194304 {-8388608 {+4194304 {-8388608 {+0 {+NaN {+inf {-inf {2.1500244 {-2.1500244 {0.15000153 {-1234.0 {1234.0 {123456.0 {16775936.0 *}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}
+Result: {0 {1234 {4321 {16771538 {2 {0 {16777215 {16777215 {0 {0 {+0 {+1234 {+4321 {-5678 {+2 {-12 {+8388607 {-8388608 {+8388607 {-8388608 {+0 {+NaN {+inf {-inf {2.1500244 {-2.1500244 {0.15000153 {-1234.0 {1234.0 {123456.0 {16775936.0 *}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}

From f217b2088c9e86cc98f2d9284014068a79048736 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Mon, 3 Jun 2024 06:51:07 +0200
Subject: [PATCH 16/17] add printing of casts to c and cu, remove to_ prefix

---
 src/ast.rs                       | 12 +++---
 src/hvm.c                        |  5 +++
 src/hvm.cu                       |  5 +++
 tests/programs/numeric-casts.hvm | 67 +++++++++++++++++---------------
 4 files changed, 52 insertions(+), 37 deletions(-)

diff --git a/src/ast.rs b/src/ast.rs
index 378ece56..2d6c0370 100644
--- a/src/ast.rs
+++ b/src/ast.rs
@@ -45,9 +45,9 @@ impl<'i> CoreParser<'i> {
 
     // numeric casts
     if let Some(cast) = match () {
-      _ if self.try_consume("to_u24") => Some(hvm::TY_U24),
-      _ if self.try_consume("to_i24") => Some(hvm::TY_I24),
-      _ if self.try_consume("to_f24") => Some(hvm::TY_F24),
+      _ if self.try_consume("u24") => Some(hvm::TY_U24),
+      _ if self.try_consume("i24") => Some(hvm::TY_I24),
+      _ if self.try_consume("f24") => Some(hvm::TY_F24),
       _ => None
     } {
       // Casts can't be partially applied, so nothing should follow.
@@ -239,9 +239,9 @@ impl Numb {
     match numb.get_typ() {
       hvm::TY_SYM => match numb.get_sym() as hvm::Tag {
         // casts
-        hvm::TY_U24 => "[to_u24]".to_string(),
-        hvm::TY_I24 => "[to_i24]".to_string(),
-        hvm::TY_F24 => "[to_f24]".to_string(),
+        hvm::TY_U24 => "[u24]".to_string(),
+        hvm::TY_I24 => "[i24]".to_string(),
+        hvm::TY_F24 => "[f24]".to_string(),
         // operations
         hvm::OP_ADD => "[+]".to_string(),
         hvm::OP_SUB => "[-]".to_string(),
diff --git a/src/hvm.c b/src/hvm.c
index 3a4cda49..fa9caa28 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -1804,6 +1804,11 @@ void pretty_print_numb(Numb word) {
   switch (get_typ(word)) {
     case TY_SYM: {
       switch (get_sym(word)) {
+        // types
+        case TY_U24: printf("[u24]"); break;
+        case TY_I24: printf("[i24]"); break;
+        case TY_F24: printf("[f24]"); break;
+        // operations
         case OP_ADD: printf("[+]"); break;
         case OP_SUB: printf("[-]"); break;
         case FP_SUB: printf("[:-]"); break;
diff --git a/src/hvm.cu b/src/hvm.cu
index c9bf18f8..b824e7be 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -2248,6 +2248,11 @@ __device__ void pretty_print_numb(Numb word) {
   switch (get_typ(word)) {
     case TY_SYM: {
       switch (get_sym(word)) {
+        // types
+        case TY_U24: printf("[u24]"); break;
+        case TY_I24: printf("[i24]"); break;
+        case TY_F24: printf("[f24]"); break;
+        // operations
         case OP_ADD: printf("[+]"); break;
         case OP_SUB: printf("[-]"); break;
         case FP_SUB: printf("[:-]"); break;
diff --git a/tests/programs/numeric-casts.hvm b/tests/programs/numeric-casts.hvm
index d01a1ebf..16df5e65 100644
--- a/tests/programs/numeric-casts.hvm
+++ b/tests/programs/numeric-casts.hvm
@@ -1,40 +1,45 @@
 @main = x & @tu0 ~ (* x)
 
 // casting to u24
-@tu0 = (* {n x}) & @tu1 ~ (* x) &  0          ~ $([to_u24] n) // 0
-@tu1 = (* {n x}) & @tu2 ~ (* x) &  1234       ~ $([to_u24] n) // 1234
-@tu2 = (* {n x}) & @tu3 ~ (* x) & +4321       ~ $([to_u24] n) // 4321
-@tu3 = (* {n x}) & @tu4 ~ (* x) & -5678       ~ $([to_u24] n) // 16771538 (reinterprets bits)
-@tu4 = (* {n x}) & @tu5 ~ (* x) &  2.8        ~ $([to_u24] n) // 2 (rounds to zero)
-@tu5 = (* {n x}) & @tu6 ~ (* x) & -12.5       ~ $([to_u24] n) // 0 (saturates)
-@tu6 = (* {n x}) & @tu7 ~ (* x) &  16777216.0 ~ $([to_u24] n) // 16777215 (saturates)
-@tu7 = (* {n x}) & @tu8 ~ (* x) & +inf        ~ $([to_u24] n) // 16777215 (saturates)
-@tu8 = (* {n x}) & @tu9 ~ (* x) & -inf        ~ $([to_u24] n) // 0 (saturates)
-@tu9 = (* {n x}) & @ti0 ~ (* x) & +NaN        ~ $([to_u24] n) // 0
+@tu0 = (* {n x}) & @tu1 ~ (* x) &  0          ~ $([u24] n) // 0
+@tu1 = (* {n x}) & @tu2 ~ (* x) &  1234       ~ $([u24] n) // 1234
+@tu2 = (* {n x}) & @tu3 ~ (* x) & +4321       ~ $([u24] n) // 4321
+@tu3 = (* {n x}) & @tu4 ~ (* x) & -5678       ~ $([u24] n) // 16771538 (reinterprets bits)
+@tu4 = (* {n x}) & @tu5 ~ (* x) &  2.8        ~ $([u24] n) // 2 (rounds to zero)
+@tu5 = (* {n x}) & @tu6 ~ (* x) & -12.5       ~ $([u24] n) // 0 (saturates)
+@tu6 = (* {n x}) & @tu7 ~ (* x) &  16777216.0 ~ $([u24] n) // 16777215 (saturates)
+@tu7 = (* {n x}) & @tu8 ~ (* x) & +inf        ~ $([u24] n) // 16777215 (saturates)
+@tu8 = (* {n x}) & @tu9 ~ (* x) & -inf        ~ $([u24] n) // 0 (saturates)
+@tu9 = (* {n x}) & @ti0 ~ (* x) & +NaN        ~ $([u24] n) // 0
 
 // casting to i24
-@ti0 = (* {n x}) & @ti1 ~ (* x) &  0          ~ $([to_i24] n) // +0
-@ti1 = (* {n x}) & @ti2 ~ (* x) &  1234       ~ $([to_i24] n) // +1234
-@ti2 = (* {n x}) & @ti3 ~ (* x) & +4321       ~ $([to_i24] n) // +4321
-@ti3 = (* {n x}) & @ti4 ~ (* x) & -5678       ~ $([to_i24] n) // -5678
-@ti4 = (* {n x}) & @ti5 ~ (* x) &  2.8        ~ $([to_i24] n) // +2   (rounds to zero)
-@ti5 = (* {n x}) & @ti6 ~ (* x) & -12.7       ~ $([to_i24] n) // -12 (rounds to zero)
-@ti6 = (* {n x}) & @ti7 ~ (* x) &  8388610.0  ~ $([to_i24] n) // +8388607 (saturates)
-@ti7 = (* {n x}) & @ti8 ~ (* x) & -8388610.0  ~ $([to_i24] n) // -8388608 (saturates)
-@ti8 = (* {n x}) & @ti9 ~ (* x) & +inf        ~ $([to_i24] n) // +8388607 (saturates)
-@ti9 = (* {n x}) & @ti10 ~ (* x) & -inf        ~ $([to_i24] n) // -8388608 (saturates)
-@ti10 = (* {n x}) & @tf0 ~ (* x) & +NaN        ~ $([to_i24] n) // +0
+@ti0 = (* {n x}) & @ti1 ~ (* x) &  0          ~ $([i24] n) // +0
+@ti1 = (* {n x}) & @ti2 ~ (* x) &  1234       ~ $([i24] n) // +1234
+@ti2 = (* {n x}) & @ti3 ~ (* x) & +4321       ~ $([i24] n) // +4321
+@ti3 = (* {n x}) & @ti4 ~ (* x) & -5678       ~ $([i24] n) // -5678
+@ti4 = (* {n x}) & @ti5 ~ (* x) &  2.8        ~ $([i24] n) // +2   (rounds to zero)
+@ti5 = (* {n x}) & @ti6 ~ (* x) & -12.7       ~ $([i24] n) // -12 (rounds to zero)
+@ti6 = (* {n x}) & @ti7 ~ (* x) &  8388610.0  ~ $([i24] n) // +8388607 (saturates)
+@ti7 = (* {n x}) & @ti8 ~ (* x) & -8388610.0  ~ $([i24] n) // -8388608 (saturates)
+@ti8 = (* {n x}) & @ti9 ~ (* x) & +inf        ~ $([i24] n) // +8388607 (saturates)
+@ti9 = (* {n x}) & @ti10 ~ (* x) & -inf        ~ $([i24] n) // -8388608 (saturates)
+@ti10 = (* {n x}) & @tf0 ~ (* x) & +NaN        ~ $([i24] n) // +0
 
 // casting to f24
-@tf0 = (* {n x}) & @tf1 ~ (* x) &  +NaN    ~ $([to_f24] n) // +NaN
-@tf1 = (* {n x}) & @tf2 ~ (* x) &  +inf    ~ $([to_f24] n) // +inf
-@tf2 = (* {n x}) & @tf3 ~ (* x) &  -inf    ~ $([to_f24] n) // -inf
-@tf3 = (* {n x}) & @tf4 ~ (* x) &  2.15    ~ $([to_f24] n) //  2.15
-@tf4 = (* {n x}) & @tf5 ~ (* x) & -2.15    ~ $([to_f24] n) // -2.15
-@tf5 = (* {n x}) & @tf6 ~ (* x) &  0.15    ~ $([to_f24] n) //  0.15
-@tf6 = (* {n x}) & @tf7 ~ (* x) & -1234    ~ $([to_f24] n) // -1234.0
-@tf7 = (* {n x}) & @tf8 ~ (* x) & +1234    ~ $([to_f24] n) // +1234.0
-@tf8 = (* {n x}) & @tf9 ~ (* x) & 123456   ~ $([to_f24] n) // 123456.0
-@tf9 = (* {n x}) & @t   ~ (* x) & 16775982 ~ $([to_f24] n) // 16775936.0
+@tf0 = (* {n x}) & @tf1 ~ (* x) &  +NaN    ~ $([f24] n) // +NaN
+@tf1 = (* {n x}) & @tf2 ~ (* x) &  +inf    ~ $([f24] n) // +inf
+@tf2 = (* {n x}) & @tf3 ~ (* x) &  -inf    ~ $([f24] n) // -inf
+@tf3 = (* {n x}) & @tf4 ~ (* x) &  2.15    ~ $([f24] n) //  2.15
+@tf4 = (* {n x}) & @tf5 ~ (* x) & -2.15    ~ $([f24] n) // -2.15
+@tf5 = (* {n x}) & @tf6 ~ (* x) &  0.15    ~ $([f24] n) //  0.15
+@tf6 = (* {n x}) & @tf7 ~ (* x) & -1234    ~ $([f24] n) // -1234.0
+@tf7 = (* {n x}) & @tf8 ~ (* x) & +1234    ~ $([f24] n) // +1234.0
+@tf8 = (* {n x}) & @tf9 ~ (* x) & 123456   ~ $([f24] n) // 123456.0
+@tf9 = (* {n x}) & @tp0 ~ (* x) & 16775982 ~ $([f24] n) // 16775936.0
+
+// printing
+@tp0 = (* {n x}) & @tp1 ~ (* x) & n ~ [u24] // [u24]
+@tp1 = (* {n x}) & @tp2 ~ (* x) & n ~ [i24] // [i24]
+@tp2 = (* {n x}) & @t   ~ (* x) & n ~ [f24] // [f24]
 
 @t = *

From f080a3ac04394f2989fffd4e0109020ae467f266 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Mon, 3 Jun 2024 06:51:46 +0200
Subject: [PATCH 17/17] snapshot

---
 tests/snapshots/run__file@numeric-casts.hvm.snap | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/tests/snapshots/run__file@numeric-casts.hvm.snap b/tests/snapshots/run__file@numeric-casts.hvm.snap
index 30ff9a5b..cc1dd365 100644
--- a/tests/snapshots/run__file@numeric-casts.hvm.snap
+++ b/tests/snapshots/run__file@numeric-casts.hvm.snap
@@ -3,4 +3,4 @@ source: tests/run.rs
 expression: rust_output
 input_file: tests/programs/numeric-casts.hvm
 ---
-Result: {0 {1234 {4321 {16771538 {2 {0 {16777215 {16777215 {0 {0 {+0 {+1234 {+4321 {-5678 {+2 {-12 {+8388607 {-8388608 {+8388607 {-8388608 {+0 {+NaN {+inf {-inf {2.1500244 {-2.1500244 {0.15000153 {-1234.0 {1234.0 {123456.0 {16775936.0 *}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}
+Result: {0 {1234 {4321 {16771538 {2 {0 {16777215 {16777215 {0 {0 {+0 {+1234 {+4321 {-5678 {+2 {-12 {+8388607 {-8388608 {+8388607 {-8388608 {+0 {+NaN {+inf {-inf {2.1500244 {-2.1500244 {0.15000153 {-1234.0 {1234.0 {123456.0 {16775936.0 {[u24] {[i24] {[f24] *}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}}