From 26062e1f98d4f955d04ba9060e043af47eca7030 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Wed, 29 May 2024 08:06:07 +0200
Subject: [PATCH 01/12] WIP

---
 Cargo.lock                 |   2 +-
 examples/demo_io/main.bend |  10 +-
 flake.lock                 | 115 ++++++++++++++++++++
 flake.nix                  |  26 +++++
 io.hvm                     | 212 +++++++++++++++++++++++++++++++++++++
 src/hvm.c                  | 124 +++++++++++++++++-----
 6 files changed, 461 insertions(+), 28 deletions(-)
 create mode 100644 flake.lock
 create mode 100644 flake.nix
 create mode 100644 io.hvm

diff --git a/Cargo.lock b/Cargo.lock
index 73e147ee..3d600240 100644
--- a/Cargo.lock
+++ b/Cargo.lock
@@ -163,7 +163,7 @@ checksum = "809e18805660d7b6b2e2b9f316a5099521b5998d5cba4dda11b5157a21aaef03"
 
 [[package]]
 name = "hvm"
-version = "2.0.16"
+version = "2.0.17"
 dependencies = [
  "TSPL",
  "cc",
diff --git a/examples/demo_io/main.bend b/examples/demo_io/main.bend
index cfb34bd4..05434f0c 100644
--- a/examples/demo_io/main.bend
+++ b/examples/demo_io/main.bend
@@ -23,6 +23,12 @@ def call_io(func, argm):
 
 def main:
   do IO_T:
-    x <- call_io("PUT_TEXT", "hello\n")
-    y <- call_io("PUT_TEXT", "world\n")
+    * <- call_io("PUT_TEXT", "what is your name?\n")
+    n <- call_io("GET_LINE", "")
+    c <- call_io("PUT_LINE", "pick a number 0-9?\n")
+    * <- call_io("PUT_TEXT", "hello, '")
+    * <- call_io("PUT_TEXT", n)
+    * <- call_io("PUT_TEXT", "', you picked '")
+    * <- call_io("PUT_TEXT", c)
+    * <- call_io("PUT_TEXT", "'! Nice choice!")
     return 42
diff --git a/flake.lock b/flake.lock
new file mode 100644
index 00000000..424b8522
--- /dev/null
+++ b/flake.lock
@@ -0,0 +1,115 @@
+{
+  "nodes": {
+    "fenix": {
+      "inputs": {
+        "nixpkgs": "nixpkgs",
+        "rust-analyzer-src": "rust-analyzer-src"
+      },
+      "locked": {
+        "lastModified": 1714544767,
+        "narHash": "sha256-kF1bX+YFMedf1g0PAJYwGUkzh22JmULtj8Rm4IXAQKs=",
+        "owner": "nix-community",
+        "repo": "fenix",
+        "rev": "73124e1356bde9411b163d636b39fe4804b7ca45",
+        "type": "github"
+      },
+      "original": {
+        "owner": "nix-community",
+        "ref": "monthly",
+        "repo": "fenix",
+        "type": "github"
+      }
+    },
+    "flake-utils": {
+      "inputs": {
+        "systems": "systems"
+      },
+      "locked": {
+        "lastModified": 1710146030,
+        "narHash": "sha256-SZ5L6eA7HJ/nmkzGG7/ISclqe6oZdOZTNoesiInkXPQ=",
+        "owner": "numtide",
+        "repo": "flake-utils",
+        "rev": "b1d9ab70662946ef0850d488da1c9019f3a9752a",
+        "type": "github"
+      },
+      "original": {
+        "owner": "numtide",
+        "repo": "flake-utils",
+        "type": "github"
+      }
+    },
+    "nixpkgs": {
+      "locked": {
+        "lastModified": 1714253743,
+        "narHash": "sha256-mdTQw2XlariysyScCv2tTE45QSU9v/ezLcHJ22f0Nxc=",
+        "owner": "nixos",
+        "repo": "nixpkgs",
+        "rev": "58a1abdbae3217ca6b702f03d3b35125d88a2994",
+        "type": "github"
+      },
+      "original": {
+        "owner": "nixos",
+        "ref": "nixos-unstable",
+        "repo": "nixpkgs",
+        "type": "github"
+      }
+    },
+    "nixpkgs_2": {
+      "locked": {
+        "lastModified": 1716312448,
+        "narHash": "sha256-PH3w5av8d+TdwCkiWN4UPBTxrD9MpxIQPDVWctlomVo=",
+        "owner": "NixOS",
+        "repo": "nixpkgs",
+        "rev": "e381a1288138aceda0ac63db32c7be545b446921",
+        "type": "github"
+      },
+      "original": {
+        "owner": "NixOS",
+        "ref": "nixpkgs-unstable",
+        "repo": "nixpkgs",
+        "type": "github"
+      }
+    },
+    "root": {
+      "inputs": {
+        "fenix": "fenix",
+        "flake-utils": "flake-utils",
+        "nixpkgs": "nixpkgs_2"
+      }
+    },
+    "rust-analyzer-src": {
+      "flake": false,
+      "locked": {
+        "lastModified": 1714501997,
+        "narHash": "sha256-g31zfxwUFzkPgX0Q8sZLcrqGmOxwjEZ/iqJjNx4fEGo=",
+        "owner": "rust-lang",
+        "repo": "rust-analyzer",
+        "rev": "49e502b277a8126a9ad10c802d1aaa3ef1a280ef",
+        "type": "github"
+      },
+      "original": {
+        "owner": "rust-lang",
+        "ref": "nightly",
+        "repo": "rust-analyzer",
+        "type": "github"
+      }
+    },
+    "systems": {
+      "locked": {
+        "lastModified": 1681028828,
+        "narHash": "sha256-Vy1rq5AaRuLzOxct8nz4T6wlgyUR7zLU309k9mBC768=",
+        "owner": "nix-systems",
+        "repo": "default",
+        "rev": "da67096a3b9bf56a91d16901293e51ba5b49a27e",
+        "type": "github"
+      },
+      "original": {
+        "owner": "nix-systems",
+        "repo": "default",
+        "type": "github"
+      }
+    }
+  },
+  "root": "root",
+  "version": 7
+}
diff --git a/flake.nix b/flake.nix
new file mode 100644
index 00000000..cab61933
--- /dev/null
+++ b/flake.nix
@@ -0,0 +1,26 @@
+{
+  description = "hvm2";
+  inputs.nixpkgs.url = "github:NixOS/nixpkgs/nixpkgs-unstable";
+  inputs.flake-utils.url = "github:numtide/flake-utils";
+  inputs.fenix.url = "github:nix-community/fenix/monthly";
+
+  outputs = { self, nixpkgs, flake-utils, fenix }:
+    flake-utils.lib.eachDefaultSystem (system:
+      let
+        pkgs = nixpkgs.legacyPackages.${system};
+
+        hoc-upload = pkgs.writeShellScriptBin "hoc-upload" ''
+          scp src/hvm.cu hoc@192.168.194.38:~/enricozb/hvm/src/hvm.cu
+        '';
+      in {
+        devShells.default = pkgs.mkShell {
+          packages = with fenix.packages.${system}; [
+            minimal.rustc
+            minimal.cargo
+            pkgs.clippy
+
+            hoc-upload
+          ];
+        };
+      });
+}
diff --git a/io.hvm b/io.hvm
new file mode 100644
index 00000000..9bcef7d0
--- /dev/null
+++ b/io.hvm
@@ -0,0 +1,212 @@
+@IO_T/Call = (a (b (c (d ((1 (a (b (c (d e))))) e)))))
+
+@IO_T/Done = (a (b ((0 (a (b c))) c)))
+
+@IO_T/MAGIC = (13683217 16719857)
+
+@IO_T/bind = ((@IO_T/bind__C2 a) a)
+
+@IO_T/bind__C0 = (* (a ((a b) b)))
+
+@IO_T/bind__C1 = (* (* (a (b ((c d) (e g))))))
+  & @IO_T/Call ~ (@IO_T/MAGIC (a (b ((c f) g))))
+  & @IO_T/bind ~ (d (e f))
+
+@IO_T/bind__C2 = (?((@IO_T/bind__C0 @IO_T/bind__C1) a) a)
+
+@IO_T/wrap = a
+  & @IO_T/Done ~ (@IO_T/MAGIC a)
+
+@String/Cons = (a (b ((1 (a (b c))) c)))
+
+@String/Nil = ((0 a) a)
+
+@call_io = (a (b c))
+  & @IO_T/Call ~ (@IO_T/MAGIC (a (b (@call_io__C0 c))))
+
+@call_io__C0 = a
+  & @IO_T/Done ~ (@IO_T/MAGIC a)
+
+@main = a
+  & @IO_T/bind ~ (@main__C10 (@main__C9 a))
+
+@main__C0 = x
+  & @call_io ~ (h (w x))
+  & @String/Cons ~ (80 (g h))
+  & @String/Cons ~ (85 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (84 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (88 (a b))
+  & @String/Cons ~ (84 (@String/Nil a))
+  & @String/Cons ~ (39 (v w))
+  & @String/Cons ~ (33 (u v))
+  & @String/Cons ~ (32 (t u))
+  & @String/Cons ~ (78 (s t))
+  & @String/Cons ~ (105 (r s))
+  & @String/Cons ~ (99 (q r))
+  & @String/Cons ~ (101 (p q))
+  & @String/Cons ~ (32 (o p))
+  & @String/Cons ~ (99 (n o))
+  & @String/Cons ~ (104 (m n))
+  & @String/Cons ~ (111 (l m))
+  & @String/Cons ~ (105 (k l))
+  & @String/Cons ~ (99 (j k))
+  & @String/Cons ~ (101 (i j))
+  & @String/Cons ~ (33 (@String/Nil i))
+
+@main__C1 = h
+  & @String/Cons ~ (80 (g h))
+  & @String/Cons ~ (85 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (84 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (88 (a b))
+  & @String/Cons ~ (84 (@String/Nil a))
+
+@main__C10 = bb
+  & @call_io ~ (h (ab bb))
+  & @String/Cons ~ (80 (g h))
+  & @String/Cons ~ (85 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (84 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (88 (a b))
+  & @String/Cons ~ (84 (@String/Nil a))
+  & @String/Cons ~ (119 (z ab))
+  & @String/Cons ~ (104 (y z))
+  & @String/Cons ~ (97 (x y))
+  & @String/Cons ~ (116 (w x))
+  & @String/Cons ~ (32 (v w))
+  & @String/Cons ~ (105 (u v))
+  & @String/Cons ~ (115 (t u))
+  & @String/Cons ~ (32 (s t))
+  & @String/Cons ~ (121 (r s))
+  & @String/Cons ~ (111 (q r))
+  & @String/Cons ~ (117 (p q))
+  & @String/Cons ~ (114 (o p))
+  & @String/Cons ~ (32 (n o))
+  & @String/Cons ~ (110 (m n))
+  & @String/Cons ~ (97 (l m))
+  & @String/Cons ~ (109 (k l))
+  & @String/Cons ~ (101 (j k))
+  & @String/Cons ~ (63 (i j))
+  & @String/Cons ~ (10 (@String/Nil i))
+
+@main__C2 = (* a)
+  & @IO_T/bind ~ (@main__C0 ((* 42) a))
+
+@main__C3 = x
+  & @call_io ~ (h (w x))
+  & @String/Cons ~ (80 (g h))
+  & @String/Cons ~ (85 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (84 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (88 (a b))
+  & @String/Cons ~ (84 (@String/Nil a))
+  & @String/Cons ~ (39 (v w))
+  & @String/Cons ~ (44 (u v))
+  & @String/Cons ~ (32 (t u))
+  & @String/Cons ~ (121 (s t))
+  & @String/Cons ~ (111 (r s))
+  & @String/Cons ~ (117 (q r))
+  & @String/Cons ~ (32 (p q))
+  & @String/Cons ~ (112 (o p))
+  & @String/Cons ~ (105 (n o))
+  & @String/Cons ~ (99 (m n))
+  & @String/Cons ~ (107 (l m))
+  & @String/Cons ~ (101 (k l))
+  & @String/Cons ~ (100 (j k))
+  & @String/Cons ~ (32 (i j))
+  & @String/Cons ~ (39 (@String/Nil i))
+
+@main__C4 = h
+  & @String/Cons ~ (80 (g h))
+  & @String/Cons ~ (85 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (84 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (88 (a b))
+  & @String/Cons ~ (84 (@String/Nil a))
+
+@main__C5 = q
+  & @call_io ~ (h (p q))
+  & @String/Cons ~ (80 (g h))
+  & @String/Cons ~ (85 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (84 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (88 (a b))
+  & @String/Cons ~ (84 (@String/Nil a))
+  & @String/Cons ~ (104 (o p))
+  & @String/Cons ~ (101 (n o))
+  & @String/Cons ~ (108 (m n))
+  & @String/Cons ~ (108 (l m))
+  & @String/Cons ~ (111 (k l))
+  & @String/Cons ~ (44 (j k))
+  & @String/Cons ~ (32 (i j))
+  & @String/Cons ~ (39 (@String/Nil i))
+
+@main__C6 = bb
+  & @call_io ~ (h (ab bb))
+  & @String/Cons ~ (71 (g h))
+  & @String/Cons ~ (69 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (67 (c d))
+  & @String/Cons ~ (72 (b c))
+  & @String/Cons ~ (65 (a b))
+  & @String/Cons ~ (82 (@String/Nil a))
+  & @String/Cons ~ (112 (z ab))
+  & @String/Cons ~ (105 (y z))
+  & @String/Cons ~ (99 (x y))
+  & @String/Cons ~ (107 (w x))
+  & @String/Cons ~ (32 (v w))
+  & @String/Cons ~ (97 (u v))
+  & @String/Cons ~ (32 (t u))
+  & @String/Cons ~ (110 (s t))
+  & @String/Cons ~ (117 (r s))
+  & @String/Cons ~ (109 (q r))
+  & @String/Cons ~ (98 (p q))
+  & @String/Cons ~ (101 (o p))
+  & @String/Cons ~ (114 (n o))
+  & @String/Cons ~ (32 (m n))
+  & @String/Cons ~ (48 (l m))
+  & @String/Cons ~ (45 (k l))
+  & @String/Cons ~ (57 (j k))
+  & @String/Cons ~ (63 (i j))
+  & @String/Cons ~ (10 (@String/Nil i))
+
+@main__C7 = (a i)
+  & @IO_T/bind ~ (@main__C6 ((c h) i))
+  & @IO_T/bind ~ (@main__C5 ((* g) h))
+  & @IO_T/bind ~ (b ((* f) g))
+  & @call_io ~ (@main__C4 (a b))
+  & @IO_T/bind ~ (@main__C3 ((* e) f))
+  & @IO_T/bind ~ (d (@main__C2 e))
+  & @call_io ~ (@main__C1 (c d))
+
+@main__C8 = i
+  & @call_io ~ (h (@String/Nil i))
+  & @String/Cons ~ (71 (g h))
+  & @String/Cons ~ (69 (f g))
+  & @String/Cons ~ (84 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (76 (c d))
+  & @String/Cons ~ (73 (b c))
+  & @String/Cons ~ (78 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+
+@main__C9 = (* a)
+  & @IO_T/bind ~ (@main__C8 (@main__C7 a))
+
+@test-skip = 1
+
+
diff --git a/src/hvm.c b/src/hvm.c
index a1bb0217..2dab33c2 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -599,7 +599,7 @@ static inline void node_create(Net* net, u32 loc, Pair val) {
   atomic_store_explicit(&net->node_buf[loc], val, memory_order_relaxed);
 }
 
-// Stores a var on global. Returns old.
+// Stores a var on global.
 static inline void vars_create(Net* net, u32 var, Port val) {
   atomic_store_explicit(&net->vars_buf[var], val, memory_order_relaxed);
 }
@@ -619,11 +619,6 @@ static inline void node_store(Net* net, u32 loc, Pair val) {
   atomic_store_explicit(&net->node_buf[loc], val, memory_order_relaxed);
 }
 
-// Stores a var on global. Returns old.
-static inline void vars_store(Net* net, u32 var, Port val) {
-  atomic_store_explicit(&net->vars_buf[var], val, memory_order_relaxed);
-}
-
 // Exchanges a node on global by a value. Returns old.
 static inline Pair node_exchange(Net* net, u32 loc, Pair val) {
   return atomic_exchange_explicit(&net->node_buf[loc], val, memory_order_relaxed);
@@ -1231,7 +1226,7 @@ Port expand(Net* net, Book* book, Port port) {
     normalize(net, book);
     got = peek(net, vars_load(net, get_val(ROOT)));
   }
-  vars_store(net, get_val(ROOT), old);
+  vars_create(net, get_val(ROOT), old);
   return got;
 }
 
@@ -1240,7 +1235,7 @@ Port expand(Net* net, Book* book, Port port) {
 
 // Reads back a λ-Encoded constructor from device to host.
 // Encoding: λt ((((t TAG) arg0) arg1) ...)
-Ctr read_ctr(Net* net, Book* book, Port port) {
+Ctr port_to_ctr(Net* net, Book* book, Port port) {
   Ctr ctr;
   ctr.tag = -1;
   ctr.args_len = 0;
@@ -1272,13 +1267,13 @@ Ctr read_ctr(Net* net, Book* book, Port port) {
   return ctr;
 }
 
-// Reads back a UTF-32 (truncated to 24 bits) string.
+// Converts a UTF-32 (truncated to 24 bits) string to a Port.
 // Since unicode scalars can fit in 21 bits, HVM's u24
 // integers can contain any unicode scalar value.
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Str read_str(Net* net, Book* book, Port port) {
+Str port_to_str(Net* net, Book* book, Port port) {
   // Result
   Str str;
   str.text_len = 0;
@@ -1291,7 +1286,7 @@ Str read_str(Net* net, Book* book, Port port) {
     //printf("reading str %s\n", show_port(peek(net, port)).x);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = read_ctr(net, book, peek(net, port));
+    Ctr ctr = port_to_ctr(net, book, peek(net, port));
 
     //printf("reading tag %d | len %d\n", ctr.tag, ctr.args_len);
 
@@ -1319,6 +1314,58 @@ Str read_str(Net* net, Book* book, Port port) {
   return str;
 }
 
+/// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
+Port nil_port(Net* net) {
+  if (!get_resources(net, tm[0], 0, 2, 1)) {
+    printf("failed to get resources... (nil)\n");
+    return new_port(ERA, 0);
+  }
+
+  vars_create(net, tm[0]->vloc[0], NONE);
+  Port var = new_port(VAR, tm[0]->vloc[0]);
+
+  node_create(net, tm[0]->nloc[0], new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
+  node_create(net, tm[0]->nloc[1], new_pair(new_port(CON, tm[0]->nloc[0]), var));
+
+  return new_port(CON, tm[0]->nloc[1]);
+}
+
+/// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
+Port cons_port(Net* net, Port head, Port tail) {
+  if (!get_resources(net, tm[0], 0, 4, 1)) {
+    printf("failed to get resources... (cons)\n");
+    return new_port(ERA, 0);
+  }
+
+  vars_create(net, tm[0]->vloc[0], NONE);
+  Port var = new_port(VAR, tm[0]->vloc[0]);
+
+  node_create(net, tm[0]->nloc[0], new_pair(tail, var));
+  node_create(net, tm[0]->nloc[1], new_pair(head, new_port(CON, tm[0]->nloc[0])));
+  node_create(net, tm[0]->nloc[2], new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, tm[0]->nloc[1])));
+  node_create(net, tm[0]->nloc[3], new_pair(new_port(CON, tm[0]->nloc[2]), var));
+
+  return new_port(CON, tm[0]->nloc[3]);
+}
+
+// Converts a UTF-32 (truncated to 24 bits) string to a Port.
+// Since unicode scalars can fit in 21 bits, HVM's u24
+// integers can contain any unicode scalar value.
+// Encoding:
+// - λt (t NIL)
+// - λt (((t CONS) head) tail)
+Port str_to_port(Net* net, Str *str) {
+  Port port = nil_port(net);
+
+  u32 len = str->text_len;
+  for (u32 i = 0; i < len; i++) {
+    Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
+    port = cons_port(net, chr, port);
+  }
+
+  return port;
+}
+
 // Reads back an image.
 // Encoding: (<tree>,<tree>) | #RRGGBB
 void read_img(Net* net, Port port, u32 width, u32 height, u32* buffer) {
@@ -1372,16 +1419,42 @@ void read_img(Net* net, Port port, u32 width, u32 height, u32* buffer) {
 // Primitive IO Fns
 // -----------------
 
-// IO: GetText
-Port io_get_text(Net* net, Book* book, Port argm) {
-  printf("TODO\n");
-  return new_port(ERA, 0);
+// IO: GetChar
+// Reads a single char from stdin.
+Port io_get_char(Net* net, Book* book, Port argm) {
+  /// Read a string.
+  Str str;
+
+  str.text_buf[0] = fgetc(stdin);
+  str.text_buf[1] = 0;
+  str.text_len = 1;
+
+  return str_to_port(net, &str);
+}
+
+// IO: GetLine
+// Reads from stdin at most 255 characters or until a newline is seen.
+Port io_get_line(Net* net, Book* book, Port argm) {
+  /// Read a string.
+  Str str;
+
+  fgets(str.text_buf, sizeof(str.text_buf), stdin);
+  str.text_len = strlen(str.text_buf);
+
+  // Strip any trailing newline.
+  if (str.text_len > 0 && str.text_buf[str.text_len - 1] == '\n') {
+    str.text_buf[str.text_len] = 0;
+    str.text_len--;
+  }
+
+  // Convert it to a port.
+  return str_to_port(net, &str);
 }
 
 // IO: PutText
 Port io_put_text(Net* net, Book* book, Port argm) {
   // Converts argument to C string
-  Str str = read_str(net, book, argm);
+  Str str = port_to_str(net, book, argm);
   // Prints it
   printf("%s", str.text_buf);
   // Returns result (in this case, just an eraser)
@@ -1554,7 +1627,7 @@ void do_run_io(Net* net, Book* book, Port port) {
     normalize(net, book);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = read_ctr(net, book, peek(net, port));
+    Ctr ctr = port_to_ctr(net, book, peek(net, port));
 
     // Checks if IO Magic Number is a CON
     if (get_tag(ctr.args_buf[0]) != CON) {
@@ -1570,7 +1643,7 @@ void do_run_io(Net* net, Book* book, Port port) {
 
     switch (ctr.tag) {
       case IO_CALL: {
-        Str  func = read_str(net, book, ctr.args_buf[1]);
+        Str  func = port_to_str(net, book, ctr.args_buf[1]);
         Port argm = ctr.args_buf[2];
         Port cont = ctr.args_buf[3];
         u32  lps  = 0;
@@ -1607,13 +1680,14 @@ void do_run_io(Net* net, Book* book, Port port) {
 
 // TODO: initialize ffns_len with the builtin ffns
 void book_init(Book* book) {
-  book->ffns_len = 6;
-  book->ffns_buf[0] = (FFn){"GET_TEXT", io_get_text};
-  book->ffns_buf[1] = (FFn){"PUT_TEXT", io_put_text};
-  book->ffns_buf[2] = (FFn){"GET_FILE", io_get_file};
-  book->ffns_buf[3] = (FFn){"PUT_FILE", io_put_file};
-  book->ffns_buf[4] = (FFn){"GET_TIME", io_get_time};
-  book->ffns_buf[5] = (FFn){"PUT_TIME", io_put_time};
+  book->ffns_len = 7;
+  book->ffns_buf[0] = (FFn){"GET_CHAR", io_get_char};
+  book->ffns_buf[1] = (FFn){"GET_LINE", io_get_line};
+  book->ffns_buf[2] = (FFn){"PUT_TEXT", io_put_text};
+  book->ffns_buf[3] = (FFn){"GET_FILE", io_get_file};
+  book->ffns_buf[4] = (FFn){"PUT_FILE", io_put_file};
+  book->ffns_buf[5] = (FFn){"GET_TIME", io_get_time};
+  book->ffns_buf[6] = (FFn){"PUT_TIME", io_put_time};
 }
 
 void book_load(Book* book, u32* buf) {

From c5d497a22ccb9c9fc5a15d01583d09cc6e88bb68 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Thu, 30 May 2024 11:46:17 +0200
Subject: [PATCH 02/12] C io

---
 src/hvm.c | 151 +++++++++++++++++++++++++++++++++++++++++++-----------
 1 file changed, 120 insertions(+), 31 deletions(-)

diff --git a/src/hvm.c b/src/hvm.c
index 2dab33c2..1461d3aa 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -1419,26 +1419,66 @@ void read_img(Net* net, Port port, u32 width, u32 height, u32* buffer) {
 // Primitive IO Fns
 // -----------------
 
-// IO: GetChar
-// Reads a single char from stdin.
-Port io_get_char(Net* net, Book* book, Port argm) {
+// Open file pointers. Indices into this array
+// are used as "file descriptors".
+// Indices 0 1 and 2 are reserved.
+// - 0 -> stdin
+// - 1 -> stdout
+// - 2 -> stderr
+static FILE* FILE_POINTERS[256];
+
+FILE* port_to_file(Port port) {
+  if (get_tag(port) != NUM) {
+    fprintf(stderr, "non-num where file descriptor was expected: %i\n", get_tag(port));
+    return NULL;
+  }
+
+  u32 idx = get_u24(get_val(port));
+
+  if (idx == 0) return stdin;
+  if (idx == 1) return stdout;
+  if (idx == 2) return stderr;
+
+  FILE* fp = FILE_POINTERS[idx];
+  if (fp == NULL) {
+    fprintf(stderr, "invalid file descriptor\n");
+    return NULL;
+  }
+
+  return fp;
+}
+
+// Reads a single char from `argm`.
+Port io_read_char(Net* net, Book* book, Port argm) {
+  FILE* fp = port_to_file(peek(net, argm));
+  if (fp == NULL) {
+    return new_port(ERA, 0);
+  }
+
   /// Read a string.
   Str str;
 
-  str.text_buf[0] = fgetc(stdin);
+  str.text_buf[0] = fgetc(fp);
   str.text_buf[1] = 0;
   str.text_len = 1;
 
   return str_to_port(net, &str);
 }
 
-// IO: GetLine
-// Reads from stdin at most 255 characters or until a newline is seen.
-Port io_get_line(Net* net, Book* book, Port argm) {
+// Reads from `argm` at most 255 characters or until a newline is seen.
+Port io_read_line(Net* net, Book* book, Port argm) {
+  FILE* fp = port_to_file(peek(net, argm));
+  if (fp == NULL) {
+    fprintf(stderr, "io_read_line: invalid file descriptor\n");
+    return new_port(ERA, 0);
+  }
+
   /// Read a string.
   Str str;
 
-  fgets(str.text_buf, sizeof(str.text_buf), stdin);
+  if (fgets(str.text_buf, sizeof(str.text_buf), fp) == NULL) {
+    fprintf(stderr, "io_read_line: failed to read\n");
+  }
   str.text_len = strlen(str.text_buf);
 
   // Strip any trailing newline.
@@ -1451,29 +1491,77 @@ Port io_get_line(Net* net, Book* book, Port argm) {
   return str_to_port(net, &str);
 }
 
-// IO: PutText
-Port io_put_text(Net* net, Book* book, Port argm) {
-  // Converts argument to C string
-  Str str = port_to_str(net, book, argm);
-  // Prints it
-  printf("%s", str.text_buf);
-  // Returns result (in this case, just an eraser)
+// Opens a file with the provided mode.
+// `argm` is a tuple (CON node) of the
+// file name and mode as strings.
+Port io_open_file(Net* net, Book* book, Port argm) {
+  if (get_tag(peek(net, argm)) != CON) {
+    fprintf(stderr, "io_open_file: expected tuple\n");
+    return new_port(ERA, 0);
+  }
+
+  Pair args = node_load(net, get_val(argm));
+  Str name = port_to_str(net, book, get_fst(args));
+  Str mode = port_to_str(net, book, get_snd(args));
+
+  for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
+    if (FILE_POINTERS[fd] == NULL) {
+      FILE_POINTERS[fd] = fopen(name.text_buf, mode.text_buf);
+      return new_port(NUM, new_u24(fd));
+    }
+  }
+
+  fprintf(stderr, "io_open_file: too many open files\n");
+
   return new_port(ERA, 0);
 }
 
-// IO: GetFile
-Port io_get_file(Net* net, Book* book, Port argm) {
-  printf("TODO\n");
+// Closes a file, reclaiming the file descriptor.
+Port io_close_file(Net* net, Book* book, Port argm) {
+  FILE* fp = port_to_file(peek(net, argm));
+  if (fp == NULL) {
+    fprintf(stderr, "io_close_file: failed to close\n");
+    return new_port(ERA, 0);
+  }
+
+  int err = fclose(fp) != 0;
+  if (err != 0) {
+    fprintf(stderr, "io_close_file: failed to close: %i\n", err);
+    return new_port(ERA, 0);
+  }
+
+  FILE_POINTERS[get_u24(get_val(argm))] = NULL;
+
   return new_port(ERA, 0);
 }
 
-// IO: PutFile
-Port io_put_file(Net* net, Book* book, Port argm) {
-  printf("TODO\n");
+// Writes a string to a file.
+// `argm` is a tuple (CON node) of the
+// file descriptor and string to write.
+Port io_write(Net* net, Book* book, Port argm) {
+  if (get_tag(peek(net, argm)) != CON) {
+    fprintf(stderr, "io_write: expected tuple, but got %u, port: %u\n", get_tag(peek(net, argm)), argm);
+    return new_port(ERA, 0);
+  }
+
+  Pair args = node_load(net, get_val(argm));
+  FILE* fp = port_to_file(peek(net, get_fst(args)));
+  Str str = port_to_str(net, book, get_snd(args));
+
+  if (fp == NULL) {
+    fprintf(stderr, "io_write: invalid file descriptor\n");
+    return new_port(ERA, 0);
+  }
+
+  if (fputs(str.text_buf, fp) == EOF) {
+    fprintf(stderr, "io_write: failed to write\n");
+  }
+
   return new_port(ERA, 0);
 }
 
-// IO: GetTime
+// Returns the current time as a tuple of the high
+// and low 24 bits of a 48-bit nanosecond timestamp.
 Port io_get_time(Net* net, Book* book, Port argm) {
   // Get the current time in nanoseconds
   u64 time_ns = time64();
@@ -1488,9 +1576,10 @@ Port io_get_time(Net* net, Book* book, Port argm) {
   return new_port(CON, loc);
 }
 
-// IO: PutTime
-// NOTE: changing this name will corrupt the timeline. You've been warned.
-Port io_put_time(Net* net, Book* book, Port argm) {
+// Sleeps.
+// `argm` is a tuple (CON node) of the high and low
+// 24 bits for a 48-bit duration in nanoseconds.
+Port io_sleep(Net* net, Book* book, Port argm) {
   // Get the sleep duration node
   Pair dur_node = node_load(net, get_val(argm));
   // Get the high and low 24-bit parts of the duration
@@ -1681,13 +1770,13 @@ void do_run_io(Net* net, Book* book, Port port) {
 // TODO: initialize ffns_len with the builtin ffns
 void book_init(Book* book) {
   book->ffns_len = 7;
-  book->ffns_buf[0] = (FFn){"GET_CHAR", io_get_char};
-  book->ffns_buf[1] = (FFn){"GET_LINE", io_get_line};
-  book->ffns_buf[2] = (FFn){"PUT_TEXT", io_put_text};
-  book->ffns_buf[3] = (FFn){"GET_FILE", io_get_file};
-  book->ffns_buf[4] = (FFn){"PUT_FILE", io_put_file};
+  book->ffns_buf[0] = (FFn){"READ_CHAR", io_read_char};
+  book->ffns_buf[1] = (FFn){"READ_LINE", io_read_line};
+  book->ffns_buf[2] = (FFn){"OPEN_FILE", io_open_file};
+  book->ffns_buf[3] = (FFn){"CLOSE_FILE", io_close_file};
+  book->ffns_buf[4] = (FFn){"WRITE", io_write};
   book->ffns_buf[5] = (FFn){"GET_TIME", io_get_time};
-  book->ffns_buf[6] = (FFn){"PUT_TIME", io_put_time};
+  book->ffns_buf[6] = (FFn){"SLEEP", io_sleep};
 }
 
 void book_load(Book* book, u32* buf) {

From ab6d6cf3c59621855233917224134f51c8ba8cbf Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Thu, 30 May 2024 13:51:48 +0200
Subject: [PATCH 03/12] comments + fprintf

---
 src/hvm.c | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/src/hvm.c b/src/hvm.c
index 1461d3aa..1bc6f77e 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -1317,7 +1317,7 @@ Str port_to_str(Net* net, Book* book, Port port) {
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
 Port nil_port(Net* net) {
   if (!get_resources(net, tm[0], 0, 2, 1)) {
-    printf("failed to get resources... (nil)\n");
+    fprintf(stderr, "nil_port: failed to get resources\n");
     return new_port(ERA, 0);
   }
 
@@ -1333,7 +1333,7 @@ Port nil_port(Net* net) {
 /// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
 Port cons_port(Net* net, Port head, Port tail) {
   if (!get_resources(net, tm[0], 0, 4, 1)) {
-    printf("failed to get resources... (cons)\n");
+    fprintf(stderr, "cons_port: failed to get resources\n");
     return new_port(ERA, 0);
   }
 
@@ -1427,6 +1427,7 @@ void read_img(Net* net, Port port, u32 width, u32 height, u32* buffer) {
 // - 2 -> stderr
 static FILE* FILE_POINTERS[256];
 
+// Converts a NUM port (file descriptor) to file pointer.
 FILE* port_to_file(Port port) {
   if (get_tag(port) != NUM) {
     fprintf(stderr, "non-num where file descriptor was expected: %i\n", get_tag(port));

From 7f178a6e3453c176cbac8fd26fe7b7fd19e4895a Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Thu, 30 May 2024 13:52:36 +0200
Subject: [PATCH 04/12] cuda attempt

---
 src/hvm.cu | 241 +++++++++++++++++++++++++++++++++++++++++++++--------
 1 file changed, 208 insertions(+), 33 deletions(-)

diff --git a/src/hvm.cu b/src/hvm.cu
index 5a9b821c..3eddc9b7 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -1839,7 +1839,7 @@ Port gnet_make_node(GNet* gnet, Tag tag, Port fst, Port snd) {
 
 // Reads back a λ-Encoded constructor from device to host.
 // Encoding: λt ((((t TAG) arg0) arg1) ...)
-Ctr gnet_read_ctr(GNet* gnet, Port port) {
+Ctr port_to_ctr(GNet* gnet, Port port) {
   Ctr ctr;
   ctr.tag = -1;
   ctr.args_len = 0;
@@ -1877,7 +1877,7 @@ Ctr gnet_read_ctr(GNet* gnet, Port port) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Str gnet_read_str(GNet* gnet, Port port) {
+Str port_to_str(GNet* gnet, Port port) {
   // Result
   Str str;
   str.text_len = 0;
@@ -1888,7 +1888,7 @@ Str gnet_read_str(GNet* gnet, Port port) {
     gnet_normalize(gnet);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = gnet_read_ctr(gnet, gnet_peek(gnet, port));
+    Ctr ctr = port_to_ctr(gnet, gnet_peek(gnet, port));
 
     // Reads string layer
     switch (ctr.tag) {
@@ -1914,38 +1914,211 @@ Str gnet_read_str(GNet* gnet, Port port) {
   return str;
 }
 
+/// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
+Port nil_port(GNet* gnet) {
+  TM tm;
+  Net net = vnet_new(gnet, NULL, gnet->turn);
+
+  if (!get_resources(net, tm[0], 0, 2, 1)) {
+    fprintf(stderr, "nil_port: failed to get resources\n");
+    return new_port(ERA, 0);
+  }
+
+  vars_create(net, tm[0]->vloc[0], NONE);
+  Port var = new_port(VAR, tm[0]->vloc[0]);
+
+  node_create(net, tm[0]->nloc[0], new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
+  node_create(net, tm[0]->nloc[1], new_pair(new_port(CON, tm[0]->nloc[0]), var));
+
+  return new_port(CON, tm[0]->nloc[1]);
+}
+
+/// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
+Port cons_port(Net* net, Port head, Port tail) {
+  TM tm;
+  Net net = vnet_new(gnet, NULL, gnet->turn);
+
+  if (!get_resources(net, tm[0], 0, 4, 1)) {
+    fprintf(stderr, "cons_port: failed to get resources\n");
+    return new_port(ERA, 0);
+  }
+
+  vars_create(net, tm[0]->vloc[0], NONE);
+  Port var = new_port(VAR, tm[0]->vloc[0]);
+
+  node_create(net, tm[0]->nloc[0], new_pair(tail, var));
+  node_create(net, tm[0]->nloc[1], new_pair(head, new_port(CON, tm[0]->nloc[0])));
+  node_create(net, tm[0]->nloc[2], new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, tm[0]->nloc[1])));
+  node_create(net, tm[0]->nloc[3], new_pair(new_port(CON, tm[0]->nloc[2]), var));
+
+  return new_port(CON, tm[0]->nloc[3]);
+}
+
+// Converts a UTF-32 (truncated to 24 bits) string to a Port.
+// Since unicode scalars can fit in 21 bits, HVM's u24
+// integers can contain any unicode scalar value.
+// Encoding:
+// - λt (t NIL)
+// - λt (((t CONS) head) tail)
+Port str_to_port(Net* net, Str *str) {
+  Port port = nil_port(net);
+
+  u32 len = str->text_len;
+  for (u32 i = 0; i < len; i++) {
+    Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
+    port = cons_port(net, chr, port);
+  }
+
+  return port;
+}
+
 // Primitive IO Fns
 // -----------------
 
-// IO: GetText
-Port io_get_text(GNet* gnet, Port argm) {
-  printf("TODO\n");
-  return new_port(ERA, 0);
+// Open file pointers. Indices into this array
+// are used as "file descriptors".
+// Indices 0 1 and 2 are reserved.
+// - 0 -> stdin
+// - 1 -> stdout
+// - 2 -> stderr
+static FILE* FILE_POINTERS[256];
+
+// Converts a NUM port (file descriptor) to file pointer.
+FILE* port_to_file(Port port) {
+  if (get_tag(port) != NUM) {
+    fprintf(stderr, "non-num where file descriptor was expected: %i\n", get_tag(port));
+    return NULL;
+  }
+
+  u32 idx = get_u24(get_val(port));
+
+  if (idx == 0) return stdin;
+  if (idx == 1) return stdout;
+  if (idx == 2) return stderr;
+
+  FILE* fp = FILE_POINTERS[idx];
+  if (fp == NULL) {
+    fprintf(stderr, "invalid file descriptor\n");
+    return NULL;
+  }
+
+  return fp;
 }
 
-// IO: PutText
-Port io_put_text(GNet* gnet, Port argm) {
-  // Converts argument to C string
-  Str str = gnet_read_str(gnet, argm);
-  // Prints it
-  printf("%s", str.text_buf);
-  // Returns result (in this case, just an eraser)
+// Reads a single char from `argm`.
+Port io_read_char(GNet* gnet, Book* book, Port argm) {
+  FILE* fp = port_to_file(gnet_peek(gnet, argm));
+  if (fp == NULL) {
+    return new_port(ERA, 0);
+  }
+
+  /// Read a string.
+  Str str;
+
+  str.text_buf[0] = fgetc(fp);
+  str.text_buf[1] = 0;
+  str.text_len = 1;
+
+  return str_to_port(gnet, &str);
+}
+
+// Reads from `argm` at most 255 characters or until a newline is seen.
+Port io_read_line(GNet* gnet, Book* book, Port argm) {
+  FILE* fp = port_to_file(gnet_peek(gnet, argm));
+  if (fp == NULL) {
+    fprintf(stderr, "io_read_line: invalid file descriptor\n");
+    return new_port(ERA, 0);
+  }
+
+  /// Read a string.
+  Str str;
+
+  if (fgets(str.text_buf, sizeof(str.text_buf), fp) == NULL) {
+    fprintf(stderr, "io_read_line: failed to read\n");
+  }
+  str.text_len = strlen(str.text_buf);
+
+  // Strip any trailing newline.
+  if (str.text_len > 0 && str.text_buf[str.text_len - 1] == '\n') {
+    str.text_buf[str.text_len] = 0;
+    str.text_len--;
+  }
+
+  // Convert it to a port.
+  return str_to_port(gnet, &str);
+}
+
+// Opens a file with the provided mode.
+// `argm` is a tuple (CON node) of the
+// file name and mode as strings.
+Port io_open_file(GNet* gnet, Book* book, Port argm) {
+  if (get_tag(gnet_peek(gnet, argm)) != CON) {
+    fprintf(stderr, "io_open_file: expected tuple\n");
+    return new_port(ERA, 0);
+  }
+
+  Pair args = gnet_node_load(gnet, get_val(argm));
+  Str name = port_to_str(gnet, book, get_fst(args));
+  Str mode = port_to_str(gnet, book, get_snd(args));
+
+  for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
+    if (FILE_POINTERS[fd] == NULL) {
+      FILE_POINTERS[fd] = fopen(name.text_buf, mode.text_buf);
+      return new_port(NUM, new_u24(fd));
+    }
+  }
+
+  fprintf(stderr, "io_open_file: too many open files\n");
+
   return new_port(ERA, 0);
 }
 
-// IO: GetFile
-Port io_get_file(GNet* gnet, Port argm) {
-  printf("TODO\n");
+// Closes a file, reclaiming the file descriptor.
+Port io_close_file(GNet* gnet, Book* book, Port argm) {
+  FILE* fp = port_to_file(gnet_peek(gnet, argm));
+  if (fp == NULL) {
+    fprintf(stderr, "io_close_file: failed to close\n");
+    return new_port(ERA, 0);
+  }
+
+  int err = fclose(fp) != 0;
+  if (err != 0) {
+    fprintf(stderr, "io_close_file: failed to close: %i\n", err);
+    return new_port(ERA, 0);
+  }
+
+  FILE_POINTERS[get_u24(get_val(argm))] = NULL;
+
   return new_port(ERA, 0);
 }
 
-// IO: PutFile
-Port io_put_file(GNet* gnet, Port argm) {
-  printf("TODO\n");
+// Writes a string to a file.
+// `argm` is a tuple (CON node) of the
+// file descriptor and string to write.
+Port io_write(GNet* gnet, Book* book, Port argm) {
+  if (get_tag(gnet_peek(gnet, argm)) != CON) {
+    fprintf(stderr, "io_write: expected tuple, but got %u, port: %u\n", get_tag(peek(net, argm)), argm);
+    return new_port(ERA, 0);
+  }
+
+  Pair args = gnet_node_load(gnet, get_val(argm));
+  FILE* fp = port_to_file(gnet_peek(gnet, get_fst(args)));
+  Str str = port_to_str(gnet, book, get_snd(args));
+
+  if (fp == NULL) {
+    fprintf(stderr, "io_write: invalid file descriptor\n");
+    return new_port(ERA, 0);
+  }
+
+  if (fputs(str.text_buf, fp) == EOF) {
+    fprintf(stderr, "io_write: failed to write\n");
+  }
+
   return new_port(ERA, 0);
 }
 
-// IO: GetTime
+// Returns the current time as a tuple of the high
+// and low 24 bits of a 48-bit nanosecond timestamp.
 Port io_get_time(GNet* gnet, Port argm) {
   // Get the current time in nanoseconds
   u64 time_ns = time64();
@@ -1956,9 +2129,10 @@ Port io_get_time(GNet* gnet, Port argm) {
   return gnet_make_node(gnet, CON, new_port(NUM, new_u24(time_hi)), new_port(NUM, new_u24(time_lo)));
 }
 
-// IO: PutTime
-// NOTE: changing this name will corrupt the timeline. You've been warned.
-Port io_put_time(GNet* gnet, Port argm) {
+// Sleeps.
+// `argm` is a tuple (CON node) of the high and low
+// 24 bits for a 48-bit duration in nanoseconds.
+Port io_sleep(GNet* gnet, Port argm) {
   // Get the sleep duration node
   Pair dur_node = gnet_node_load(gnet, get_val(argm));
   // Get the high and low 24-bit parts of the duration
@@ -1986,7 +2160,7 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
     gnet_normalize(gnet);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = gnet_read_ctr(gnet, gnet_peek(gnet, port));
+    Ctr ctr = port_to_ctr(gnet, gnet_peek(gnet, port));
 
     // Checks if IO Magic Number is a CON
     if (get_tag(ctr.args_buf[0]) != CON) {
@@ -2002,7 +2176,7 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
 
     switch (ctr.tag) {
       case IO_CALL: {
-        Str  func = gnet_read_str(gnet, ctr.args_buf[1]);
+        Str  func = port_to_str(gnet, ctr.args_buf[1]);
         FFn* ffn  = NULL;
         // FIXME: optimize this linear search
         for (u32 fid = 0; fid < book->ffns_len; ++fid) {
@@ -2039,13 +2213,14 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
 
 // TODO: initialize ffns_len with the builtin ffns
 void book_init(Book* book) {
-  book->ffns_len = 6;
-  book->ffns_buf[0] = (FFn){"GET_TEXT", io_get_text};
-  book->ffns_buf[0] = (FFn){"PUT_TEXT", io_put_text};
-  book->ffns_buf[2] = (FFn){"GET_FILE", io_get_file};
-  book->ffns_buf[3] = (FFn){"PUT_FILE", io_put_file};
-  book->ffns_buf[4] = (FFn){"GET_TIME", io_get_time};
-  book->ffns_buf[5] = (FFn){"PUT_TIME", io_put_time};
+  book->ffns_len = 7;
+  book->ffns_buf[0] = (FFn){"READ_CHAR", io_read_char};
+  book->ffns_buf[1] = (FFn){"READ_LINE", io_read_line};
+  book->ffns_buf[2] = (FFn){"OPEN_FILE", io_open_file};
+  book->ffns_buf[3] = (FFn){"CLOSE_FILE", io_close_file};
+  book->ffns_buf[4] = (FFn){"WRITE", io_write};
+  book->ffns_buf[5] = (FFn){"GET_TIME", io_get_time};
+  book->ffns_buf[6] = (FFn){"SLEEP", io_sleep};
 }
 
 void book_load(Book* book, u32* buf) {

From dea5d054b88b7df7459ddfb41458f58100c9dca1 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Thu, 30 May 2024 16:37:12 +0200
Subject: [PATCH 05/12] WIP - cuda hanging sometimes...

---
 src/hvm.c  |   2 +-
 src/hvm.cu | 166 ++++++++++++++++++++++++++++++++---------------------
 2 files changed, 103 insertions(+), 65 deletions(-)

diff --git a/src/hvm.c b/src/hvm.c
index 1bc6f77e..33611422 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -1541,7 +1541,7 @@ Port io_close_file(Net* net, Book* book, Port argm) {
 // file descriptor and string to write.
 Port io_write(Net* net, Book* book, Port argm) {
   if (get_tag(peek(net, argm)) != CON) {
-    fprintf(stderr, "io_write: expected tuple, but got %u, port: %u\n", get_tag(peek(net, argm)), argm);
+    fprintf(stderr, "io_write: expected tuple, but got %u\n", get_tag(peek(net, argm)));
     return new_port(ERA, 0);
   }
 
diff --git a/src/hvm.cu b/src/hvm.cu
index 3eddc9b7..3efb4b5e 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -4,6 +4,7 @@
 
 #include <stdint.h>
 #include <stdio.h>
+#include <stdlib.h>
 
 // Integers
 // --------
@@ -1057,12 +1058,23 @@ __device__ void link_pair(Net* net, TM* tm, Pair AB) {
 
 // Gets the necessary resources for an interaction.
 __device__ bool get_resources(Net* net, TM* tm, u8 need_rbag, u8 need_node, u8 need_vars) {
+  // printf("get resources %u %u %u", need_rbag, need_node, need_vars);
+  // printf("tm: %p", tm);
+  // printf("", tm->page);
+
   u32 got_rbag = min(RLEN - tm->rbag.lo_end, RLEN - tm->rbag.hi_end);
   u32 got_node;
   u32 got_vars;
   if (tm->mode != WORK) {
+    // printf("get resources (not work): %u %u %u", need_rbag, need_node, need_vars);
+
     got_node = g_node_alloc(net, tm, need_node);
+
+    // printf("got node");
+
     got_vars = g_vars_alloc(net, tm, need_vars);
+
+    // printf("got vars");
   } else {
     got_node = l_node_alloc(net, tm, need_node);
     got_vars = l_vars_alloc(net, tm, need_vars);
@@ -1444,6 +1456,65 @@ __global__ void boot_redex(GNet* gnet, Pair redex) {
   }
 }
 
+/// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
+__device__ Port nil_port(Net* net, TM* tm) {
+  if (!get_resources(net, tm, 0, 2, 1)) {
+    printf("nil_port: failed to get resources\n");
+    return new_port(ERA, 0);
+  }
+
+  vars_create(net, tm->vloc[0], NONE);
+  Port var = new_port(VAR, tm->vloc[0]);
+
+  node_create(net, tm->nloc[0], new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
+  node_create(net, tm->nloc[1], new_pair(new_port(CON, tm->nloc[0]), var));
+
+  return new_port(CON, tm->nloc[1]);
+}
+
+/// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
+__device__ Port cons_port(Net* net, TM* tm, Port head, Port tail) {
+  if (!get_resources(net, tm, 0, 4, 1)) {
+    return new_port(ERA, 0);
+  }
+
+  vars_create(net, tm->vloc[0], NONE);
+  Port var = new_port(VAR, tm->vloc[0]);
+
+  node_create(net, tm->nloc[0], new_pair(tail, var));
+  node_create(net, tm->nloc[1], new_pair(head, new_port(CON, tm->nloc[0])));
+  node_create(net, tm->nloc[2], new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, tm->nloc[1])));
+  node_create(net, tm->nloc[3], new_pair(new_port(CON, tm->nloc[2]), var));
+
+  return new_port(CON, tm->nloc[3]);
+}
+
+// Converts a UTF-32 (truncated to 24 bits) string to a Port.
+// Since unicode scalars can fit in 21 bits, HVM's u24
+// integers can contain any unicode scalar value.
+// Encoding:
+// - λt (t NIL)
+// - λt (((t CONS) head) tail)
+__device__ Port str_to_port(Net* net, TM* tm, Str *str) {
+  Port port = nil_port(net, tm);
+
+  u32 len = str->text_len;
+  for (u32 i = 0; i < len; i++) {
+    Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
+    port = cons_port(net, tm, chr, port);
+  }
+
+  return port;
+}
+
+__global__ void make_str_port(GNet* gnet, Str *str, Port* ret) {
+  if (GID() == 0) {
+    TM tm;
+    Net net = vnet_new(gnet, NULL, gnet->turn);
+    *ret = str_to_port(&net, &tm, str);
+  }
+}
+
 // Creates a node.
 __global__ void make_node(GNet* gnet, Tag tag, Port fst, Port snd, Port* ret) {
   if (GID() == 0) {
@@ -1839,7 +1910,7 @@ Port gnet_make_node(GNet* gnet, Tag tag, Port fst, Port snd) {
 
 // Reads back a λ-Encoded constructor from device to host.
 // Encoding: λt ((((t TAG) arg0) arg1) ...)
-Ctr port_to_ctr(GNet* gnet, Port port) {
+Ctr gnet_port_to_ctr(GNet* gnet, Port port) {
   Ctr ctr;
   ctr.tag = -1;
   ctr.args_len = 0;
@@ -1877,7 +1948,7 @@ Ctr port_to_ctr(GNet* gnet, Port port) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Str port_to_str(GNet* gnet, Port port) {
+Str gnet_port_to_str(GNet* gnet, Port port) {
   // Result
   Str str;
   str.text_len = 0;
@@ -1888,7 +1959,7 @@ Str port_to_str(GNet* gnet, Port port) {
     gnet_normalize(gnet);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = port_to_ctr(gnet, gnet_peek(gnet, port));
+    Ctr ctr = gnet_port_to_ctr(gnet, gnet_peek(gnet, port));
 
     // Reads string layer
     switch (ctr.tag) {
@@ -1914,62 +1985,27 @@ Str port_to_str(GNet* gnet, Port port) {
   return str;
 }
 
-/// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
-Port nil_port(GNet* gnet) {
-  TM tm;
-  Net net = vnet_new(gnet, NULL, gnet->turn);
-
-  if (!get_resources(net, tm[0], 0, 2, 1)) {
-    fprintf(stderr, "nil_port: failed to get resources\n");
-    return new_port(ERA, 0);
-  }
-
-  vars_create(net, tm[0]->vloc[0], NONE);
-  Port var = new_port(VAR, tm[0]->vloc[0]);
-
-  node_create(net, tm[0]->nloc[0], new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
-  node_create(net, tm[0]->nloc[1], new_pair(new_port(CON, tm[0]->nloc[0]), var));
-
-  return new_port(CON, tm[0]->nloc[1]);
-}
-
-/// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
-Port cons_port(Net* net, Port head, Port tail) {
-  TM tm;
-  Net net = vnet_new(gnet, NULL, gnet->turn);
-
-  if (!get_resources(net, tm[0], 0, 4, 1)) {
-    fprintf(stderr, "cons_port: failed to get resources\n");
-    return new_port(ERA, 0);
-  }
-
-  vars_create(net, tm[0]->vloc[0], NONE);
-  Port var = new_port(VAR, tm[0]->vloc[0]);
-
-  node_create(net, tm[0]->nloc[0], new_pair(tail, var));
-  node_create(net, tm[0]->nloc[1], new_pair(head, new_port(CON, tm[0]->nloc[0])));
-  node_create(net, tm[0]->nloc[2], new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, tm[0]->nloc[1])));
-  node_create(net, tm[0]->nloc[3], new_pair(new_port(CON, tm[0]->nloc[2]), var));
-
-  return new_port(CON, tm[0]->nloc[3]);
-}
-
 // Converts a UTF-32 (truncated to 24 bits) string to a Port.
 // Since unicode scalars can fit in 21 bits, HVM's u24
 // integers can contain any unicode scalar value.
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Port str_to_port(Net* net, Str *str) {
-  Port port = nil_port(net);
+Port gnet_make_str(GNet* gnet, Str *str) {
+  Port* d_ret;
+  cudaMalloc(&d_ret, sizeof(Port));
 
-  u32 len = str->text_len;
-  for (u32 i = 0; i < len; i++) {
-    Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
-    port = cons_port(net, chr, port);
-  }
+  Str* cu_str;
+  cudaMalloc(&cu_str, sizeof(Str));
+  cudaMemcpy(cu_str, str, sizeof(Str), cudaMemcpyHostToDevice);
 
-  return port;
+  make_str_port<<<1,1>>>(gnet, cu_str, d_ret);
+
+  Port ret;
+  cudaMemcpy(&ret, d_ret, sizeof(Port), cudaMemcpyDeviceToHost);
+  cudaFree(d_ret);
+
+  return ret;
 }
 
 // Primitive IO Fns
@@ -2006,7 +2042,7 @@ FILE* port_to_file(Port port) {
 }
 
 // Reads a single char from `argm`.
-Port io_read_char(GNet* gnet, Book* book, Port argm) {
+Port io_read_char(GNet* gnet, Port argm) {
   FILE* fp = port_to_file(gnet_peek(gnet, argm));
   if (fp == NULL) {
     return new_port(ERA, 0);
@@ -2019,11 +2055,11 @@ Port io_read_char(GNet* gnet, Book* book, Port argm) {
   str.text_buf[1] = 0;
   str.text_len = 1;
 
-  return str_to_port(gnet, &str);
+  return gnet_make_str(gnet, &str);
 }
 
 // Reads from `argm` at most 255 characters or until a newline is seen.
-Port io_read_line(GNet* gnet, Book* book, Port argm) {
+Port io_read_line(GNet* gnet, Port argm) {
   FILE* fp = port_to_file(gnet_peek(gnet, argm));
   if (fp == NULL) {
     fprintf(stderr, "io_read_line: invalid file descriptor\n");
@@ -2044,22 +2080,24 @@ Port io_read_line(GNet* gnet, Book* book, Port argm) {
     str.text_len--;
   }
 
+  printf("read this string: '%s'\n", str.text_buf);
+
   // Convert it to a port.
-  return str_to_port(gnet, &str);
+  return gnet_make_str(gnet, &str);
 }
 
 // Opens a file with the provided mode.
 // `argm` is a tuple (CON node) of the
 // file name and mode as strings.
-Port io_open_file(GNet* gnet, Book* book, Port argm) {
+Port io_open_file(GNet* gnet, Port argm) {
   if (get_tag(gnet_peek(gnet, argm)) != CON) {
     fprintf(stderr, "io_open_file: expected tuple\n");
     return new_port(ERA, 0);
   }
 
   Pair args = gnet_node_load(gnet, get_val(argm));
-  Str name = port_to_str(gnet, book, get_fst(args));
-  Str mode = port_to_str(gnet, book, get_snd(args));
+  Str name = gnet_port_to_str(gnet, get_fst(args));
+  Str mode = gnet_port_to_str(gnet, get_snd(args));
 
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
@@ -2074,7 +2112,7 @@ Port io_open_file(GNet* gnet, Book* book, Port argm) {
 }
 
 // Closes a file, reclaiming the file descriptor.
-Port io_close_file(GNet* gnet, Book* book, Port argm) {
+Port io_close_file(GNet* gnet, Port argm) {
   FILE* fp = port_to_file(gnet_peek(gnet, argm));
   if (fp == NULL) {
     fprintf(stderr, "io_close_file: failed to close\n");
@@ -2095,15 +2133,15 @@ Port io_close_file(GNet* gnet, Book* book, Port argm) {
 // Writes a string to a file.
 // `argm` is a tuple (CON node) of the
 // file descriptor and string to write.
-Port io_write(GNet* gnet, Book* book, Port argm) {
+Port io_write(GNet* gnet, Port argm) {
   if (get_tag(gnet_peek(gnet, argm)) != CON) {
-    fprintf(stderr, "io_write: expected tuple, but got %u, port: %u\n", get_tag(peek(net, argm)), argm);
+    fprintf(stderr, "io_write: expected tuple, but got %u", get_tag(gnet_peek(gnet, argm)));
     return new_port(ERA, 0);
   }
 
   Pair args = gnet_node_load(gnet, get_val(argm));
   FILE* fp = port_to_file(gnet_peek(gnet, get_fst(args)));
-  Str str = port_to_str(gnet, book, get_snd(args));
+  Str str = gnet_port_to_str(gnet, get_snd(args));
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
@@ -2160,7 +2198,7 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
     gnet_normalize(gnet);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = port_to_ctr(gnet, gnet_peek(gnet, port));
+    Ctr ctr = gnet_port_to_ctr(gnet, gnet_peek(gnet, port));
 
     // Checks if IO Magic Number is a CON
     if (get_tag(ctr.args_buf[0]) != CON) {
@@ -2176,7 +2214,7 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
 
     switch (ctr.tag) {
       case IO_CALL: {
-        Str  func = port_to_str(gnet, ctr.args_buf[1]);
+        Str  func = gnet_port_to_str(gnet, ctr.args_buf[1]);
         FFn* ffn  = NULL;
         // FIXME: optimize this linear search
         for (u32 fid = 0; fid < book->ffns_len; ++fid) {

From 137b1f20cfdff97690265d1f56fa510363d5326e Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 07:36:12 +0200
Subject: [PATCH 06/12] allocate resources up-front when converting strings

---
 src/hvm.c  | 63 ++++++++++++++++++++++++++++++-----------------
 src/hvm.cu | 72 +++++++++++++++++++++++++++++++-----------------------
 2 files changed, 82 insertions(+), 53 deletions(-)

diff --git a/src/hvm.c b/src/hvm.c
index 33611422..a7c9e382 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -1315,37 +1315,46 @@ Str port_to_str(Net* net, Book* book, Port port) {
 }
 
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
+/// Should only be called within `str_to_port`, as a previous call
+/// to `get_resources` is expected.
 Port nil_port(Net* net) {
-  if (!get_resources(net, tm[0], 0, 2, 1)) {
-    fprintf(stderr, "nil_port: failed to get resources\n");
-    return new_port(ERA, 0);
-  }
+  u32 v1 = tm[0]->vloc[0];
+
+  u32 n1 = tm[0]->nloc[0];
+  u32 n2 = tm[0]->nloc[1];
 
-  vars_create(net, tm[0]->vloc[0], NONE);
-  Port var = new_port(VAR, tm[0]->vloc[0]);
+  vars_create(net, v1, NONE);
+  Port var = new_port(VAR, v1);
 
-  node_create(net, tm[0]->nloc[0], new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
-  node_create(net, tm[0]->nloc[1], new_pair(new_port(CON, tm[0]->nloc[0]), var));
+  node_create(net, n1, new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
+  node_create(net, n2, new_pair(new_port(CON, n1), var));
 
-  return new_port(CON, tm[0]->nloc[1]);
+  return new_port(CON, n2);
 }
 
 /// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
-Port cons_port(Net* net, Port head, Port tail) {
-  if (!get_resources(net, tm[0], 0, 4, 1)) {
-    fprintf(stderr, "cons_port: failed to get resources\n");
-    return new_port(ERA, 0);
-  }
+/// Should only be called within `str_to_port`, as a previous call
+/// to `get_resources` is expected.
+/// The `char_idx` parameter is used to offset the vloc and nloc
+/// allocations, otherwise they would conflict with each other on
+/// subsequent calls.
+Port cons_port(Net* net, Port head, Port tail, u32 char_idx) {
+  u32 v1 = tm[0]->vloc[1 + char_idx];
+
+  u32 n1 = tm[0]->nloc[2 + char_idx * 4 + 0];
+  u32 n2 = tm[0]->nloc[2 + char_idx * 4 + 1];
+  u32 n3 = tm[0]->nloc[2 + char_idx * 4 + 2];
+  u32 n4 = tm[0]->nloc[2 + char_idx * 4 + 3];
 
-  vars_create(net, tm[0]->vloc[0], NONE);
-  Port var = new_port(VAR, tm[0]->vloc[0]);
+  vars_create(net, v1, NONE);
+  Port var = new_port(VAR, v1);
 
-  node_create(net, tm[0]->nloc[0], new_pair(tail, var));
-  node_create(net, tm[0]->nloc[1], new_pair(head, new_port(CON, tm[0]->nloc[0])));
-  node_create(net, tm[0]->nloc[2], new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, tm[0]->nloc[1])));
-  node_create(net, tm[0]->nloc[3], new_pair(new_port(CON, tm[0]->nloc[2]), var));
+  node_create(net, n1, new_pair(tail, var));
+  node_create(net, n2, new_pair(head, new_port(CON, n1)));
+  node_create(net, n3, new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, n2)));
+  node_create(net, n4, new_pair(new_port(CON, n3), var));
 
-  return new_port(CON, tm[0]->nloc[3]);
+  return new_port(CON, n4);
 }
 
 // Converts a UTF-32 (truncated to 24 bits) string to a Port.
@@ -1355,12 +1364,20 @@ Port cons_port(Net* net, Port head, Port tail) {
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
 Port str_to_port(Net* net, Str *str) {
+  // Allocate all resources up front:
+  // - NIL needs  2 nodes & 1 var
+  // - CONS needs 4 nodes & 1 var
+  u32 len = str->text_len;
+  if (!get_resources(net, tm[0], 0, 2 + 4 * len, 1 + len)) {
+    printf("str_to_port: failed to get resources\n");
+    return new_port(ERA, 0);
+  }
+
   Port port = nil_port(net);
 
-  u32 len = str->text_len;
   for (u32 i = 0; i < len; i++) {
     Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
-    port = cons_port(net, chr, port);
+    port = cons_port(net, chr, port, i);
   }
 
   return port;
diff --git a/src/hvm.cu b/src/hvm.cu
index 3efb4b5e..7e4b857e 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -705,8 +705,8 @@ __device__ Net vnet_new(GNet* gnet, void* smem, u32 turn) {
   Net net;
   net.l_node_dif   = 0;
   net.l_vars_dif   = 0;
-  net.l_node_buf   = ((LNet*)smem)->node_buf;
-  net.l_vars_buf   = ((LNet*)smem)->vars_buf;
+  net.l_node_buf   = smem == NULL ? net.l_node_buf : ((LNet*)smem)->node_buf;
+  net.l_vars_buf   = smem == NULL ? net.l_vars_buf : ((LNet*)smem)->vars_buf;
   net.g_rbag_use_A = turn % 2 == 0 ? &gnet->rbag_use_A : &gnet->rbag_use_B;
   net.g_rbag_use_B = turn % 2 == 0 ? &gnet->rbag_use_B : &gnet->rbag_use_A;
   net.g_rbag_buf_A = turn % 2 == 0 ? gnet->rbag_buf_A : gnet->rbag_buf_B;
@@ -1058,10 +1058,6 @@ __device__ void link_pair(Net* net, TM* tm, Pair AB) {
 
 // Gets the necessary resources for an interaction.
 __device__ bool get_resources(Net* net, TM* tm, u8 need_rbag, u8 need_node, u8 need_vars) {
-  // printf("get resources %u %u %u", need_rbag, need_node, need_vars);
-  // printf("tm: %p", tm);
-  // printf("", tm->page);
-
   u32 got_rbag = min(RLEN - tm->rbag.lo_end, RLEN - tm->rbag.hi_end);
   u32 got_node;
   u32 got_vars;
@@ -1457,36 +1453,46 @@ __global__ void boot_redex(GNet* gnet, Pair redex) {
 }
 
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
+/// Should only be called within `str_to_port`, as a previous call
+/// to `get_resources` is expected.
 __device__ Port nil_port(Net* net, TM* tm) {
-  if (!get_resources(net, tm, 0, 2, 1)) {
-    printf("nil_port: failed to get resources\n");
-    return new_port(ERA, 0);
-  }
+  u32 v1 = tm->vloc[0];
 
-  vars_create(net, tm->vloc[0], NONE);
-  Port var = new_port(VAR, tm->vloc[0]);
+  u32 n1 = tm->nloc[0];
+  u32 n2 = tm->nloc[1];
 
-  node_create(net, tm->nloc[0], new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
-  node_create(net, tm->nloc[1], new_pair(new_port(CON, tm->nloc[0]), var));
+  vars_create(net, v1, NONE);
+  Port var = new_port(VAR, v1);
 
-  return new_port(CON, tm->nloc[1]);
+  node_create(net, n1, new_pair(new_port(NUM, new_u24(LIST_NIL)), var));
+  node_create(net, n2, new_pair(new_port(CON, n1), var));
+
+  return new_port(CON, n2);
 }
 
 /// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
-__device__ Port cons_port(Net* net, TM* tm, Port head, Port tail) {
-  if (!get_resources(net, tm, 0, 4, 1)) {
-    return new_port(ERA, 0);
-  }
+/// Should only be called within `str_to_port`, as a previous call
+/// to `get_resources` is expected.
+/// The `char_idx` parameter is used to offset the vloc and nloc
+/// allocations, otherwise they would conflict with each other on
+/// subsequent calls.
+__device__ Port cons_port(Net* net, TM* tm, Port head, Port tail, u32 char_idx) {
+  u32 v1 = tm->vloc[1 + char_idx];
 
-  vars_create(net, tm->vloc[0], NONE);
-  Port var = new_port(VAR, tm->vloc[0]);
+  u32 n1 = tm->nloc[2 + char_idx * 4 + 0];
+  u32 n2 = tm->nloc[2 + char_idx * 4 + 1];
+  u32 n3 = tm->nloc[2 + char_idx * 4 + 2];
+  u32 n4 = tm->nloc[2 + char_idx * 4 + 3];
 
-  node_create(net, tm->nloc[0], new_pair(tail, var));
-  node_create(net, tm->nloc[1], new_pair(head, new_port(CON, tm->nloc[0])));
-  node_create(net, tm->nloc[2], new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, tm->nloc[1])));
-  node_create(net, tm->nloc[3], new_pair(new_port(CON, tm->nloc[2]), var));
+  vars_create(net, v1, NONE);
+  Port var = new_port(VAR, v1);
 
-  return new_port(CON, tm->nloc[3]);
+  node_create(net, n1, new_pair(tail, var));
+  node_create(net, n2, new_pair(head, new_port(CON, n1)));
+  node_create(net, n3, new_pair(new_port(NUM, new_u24(LIST_CONS)), new_port(CON, n2)));
+  node_create(net, n4, new_pair(new_port(CON, n3), var));
+
+  return new_port(CON, n4);
 }
 
 // Converts a UTF-32 (truncated to 24 bits) string to a Port.
@@ -1496,12 +1502,20 @@ __device__ Port cons_port(Net* net, TM* tm, Port head, Port tail) {
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
 __device__ Port str_to_port(Net* net, TM* tm, Str *str) {
+  // Allocate all resources up front:
+  // - NIL needs  2 nodes & 1 var
+  // - CONS needs 4 nodes & 1 var
+  u32 len = str->text_len;
+  if (!get_resources(net, tm, 0, 2 + 4 * len, 1 + len)) {
+    printf("str_to_port: failed to get resources\n");
+    return new_port(ERA, 0);
+  }
+
   Port port = nil_port(net, tm);
 
-  u32 len = str->text_len;
   for (u32 i = 0; i < len; i++) {
     Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
-    port = cons_port(net, tm, chr, port);
+    port = cons_port(net, tm, chr, port, i);
   }
 
   return port;
@@ -2080,8 +2094,6 @@ Port io_read_line(GNet* gnet, Port argm) {
     str.text_len--;
   }
 
-  printf("read this string: '%s'\n", str.text_buf);
-
   // Convert it to a port.
   return gnet_make_str(gnet, &str);
 }

From 899340610a7930404d15724e02d97c2a26a48c31 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 11:05:46 +0200
Subject: [PATCH 07/12] add io tests, move snapshots

---
 tests/programs/io/read_and_print.hvm          | 117 ++++++++++++++++++
 tests/programs/{ => numerics}/f24.hvm         |   0
 tests/programs/{ => numerics}/i24.hvm         |   0
 tests/programs/{ => numerics}/u24.hvm         |   0
 tests/run.rs                                  |  91 ++++++++++----
 ....snap => run__file@numerics__f24.hvm.snap} |   0
 ....snap => run__file@numerics__i24.hvm.snap} |   0
 ....snap => run__file@numerics__u24.hvm.snap} |   0
 .../run__io_file@io__read_and_print.hvm.snap  |   8 ++
 9 files changed, 192 insertions(+), 24 deletions(-)
 create mode 100644 tests/programs/io/read_and_print.hvm
 rename tests/programs/{ => numerics}/f24.hvm (100%)
 rename tests/programs/{ => numerics}/i24.hvm (100%)
 rename tests/programs/{ => numerics}/u24.hvm (100%)
 rename tests/snapshots/{run__file@f24.hvm.snap => run__file@numerics__f24.hvm.snap} (100%)
 rename tests/snapshots/{run__file@i24.hvm.snap => run__file@numerics__i24.hvm.snap} (100%)
 rename tests/snapshots/{run__file@u24.hvm.snap => run__file@numerics__u24.hvm.snap} (100%)
 create mode 100644 tests/snapshots/run__io_file@io__read_and_print.hvm.snap

diff --git a/tests/programs/io/read_and_print.hvm b/tests/programs/io/read_and_print.hvm
new file mode 100644
index 00000000..d60a16b5
--- /dev/null
+++ b/tests/programs/io/read_and_print.hvm
@@ -0,0 +1,117 @@
+@IO_T/Call = (a (b (c (d ((1 (a (b (c (d e))))) e)))))
+
+@IO_T/Done = (a (b ((0 (a (b c))) c)))
+
+@IO_T/MAGIC = (13683217 16719857)
+
+@IO_T/bind = ((@IO_T/bind__C2 a) a)
+
+@IO_T/bind__C0 = (* (a ((a b) b)))
+
+@IO_T/bind__C1 = (* (* (a (b ((c d) (e g))))))
+  & @IO_T/Call ~ (@IO_T/MAGIC (a (b ((c f) g))))
+  & @IO_T/bind ~ (d (e f))
+
+@IO_T/bind__C2 = (?((@IO_T/bind__C0 @IO_T/bind__C1) a) a)
+
+@IO_T/wrap = a
+  & @IO_T/Done ~ (@IO_T/MAGIC a)
+
+@String/Cons = (a (b ((1 (a (b c))) c)))
+
+@String/Nil = ((0 a) a)
+
+@call_io = (a (b c))
+  & @IO_T/Call ~ (@IO_T/MAGIC (a (b (@call_io__C0 c))))
+
+@call_io__C0 = a
+  & @IO_T/Done ~ (@IO_T/MAGIC a)
+
+@main = a
+  & @IO_T/bind ~ (@main__C7 (@main__C6 a))
+
+@main__C0 = i
+  & @call_io ~ (e ((1 h) i))
+  & @String/Cons ~ (87 (d e))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (73 (b c))
+  & @String/Cons ~ (84 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+  & @String/Cons ~ (39 (g h))
+  & @String/Cons ~ (33 (f g))
+  & @String/Cons ~ (10 (@String/Nil f))
+
+@main__C1 = e
+  & @String/Cons ~ (87 (d e))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (73 (b c))
+  & @String/Cons ~ (84 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+
+@main__C2 = (* a)
+  & @IO_T/bind ~ (@main__C0 ((* 42) a))
+
+@main__C3 = n
+  & @call_io ~ (e ((1 m) n))
+  & @String/Cons ~ (87 (d e))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (73 (b c))
+  & @String/Cons ~ (84 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+  & @String/Cons ~ (72 (l m))
+  & @String/Cons ~ (101 (k l))
+  & @String/Cons ~ (108 (j k))
+  & @String/Cons ~ (108 (i j))
+  & @String/Cons ~ (111 (h i))
+  & @String/Cons ~ (44 (g h))
+  & @String/Cons ~ (32 (f g))
+  & @String/Cons ~ (39 (@String/Nil f))
+
+@main__C4 = (a d)
+  & @IO_T/bind ~ (@main__C3 ((* c) d))
+  & @IO_T/bind ~ (b (@main__C2 c))
+  & @call_io ~ (@main__C1 ((1 a) b))
+
+@main__C5 = j
+  & @call_io ~ (i (0 j))
+  & @String/Cons ~ (82 (h i))
+  & @String/Cons ~ (69 (g h))
+  & @String/Cons ~ (65 (f g))
+  & @String/Cons ~ (68 (e f))
+  & @String/Cons ~ (95 (d e))
+  & @String/Cons ~ (76 (c d))
+  & @String/Cons ~ (73 (b c))
+  & @String/Cons ~ (78 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+
+@main__C6 = (* a)
+  & @IO_T/bind ~ (@main__C5 (@main__C4 a))
+
+@main__C7 = y
+  & @call_io ~ (e ((1 x) y))
+  & @String/Cons ~ (87 (d e))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (73 (b c))
+  & @String/Cons ~ (84 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+  & @String/Cons ~ (87 (w x))
+  & @String/Cons ~ (104 (v w))
+  & @String/Cons ~ (97 (u v))
+  & @String/Cons ~ (116 (t u))
+  & @String/Cons ~ (32 (s t))
+  & @String/Cons ~ (105 (r s))
+  & @String/Cons ~ (115 (q r))
+  & @String/Cons ~ (32 (p q))
+  & @String/Cons ~ (121 (o p))
+  & @String/Cons ~ (111 (n o))
+  & @String/Cons ~ (117 (m n))
+  & @String/Cons ~ (114 (l m))
+  & @String/Cons ~ (32 (k l))
+  & @String/Cons ~ (110 (j k))
+  & @String/Cons ~ (97 (i j))
+  & @String/Cons ~ (109 (h i))
+  & @String/Cons ~ (101 (g h))
+  & @String/Cons ~ (63 (f g))
+  & @String/Cons ~ (10 (@String/Nil f))
+
+@test-io = 1
diff --git a/tests/programs/f24.hvm b/tests/programs/numerics/f24.hvm
similarity index 100%
rename from tests/programs/f24.hvm
rename to tests/programs/numerics/f24.hvm
diff --git a/tests/programs/i24.hvm b/tests/programs/numerics/i24.hvm
similarity index 100%
rename from tests/programs/i24.hvm
rename to tests/programs/numerics/i24.hvm
diff --git a/tests/programs/u24.hvm b/tests/programs/numerics/u24.hvm
similarity index 100%
rename from tests/programs/u24.hvm
rename to tests/programs/numerics/u24.hvm
diff --git a/tests/run.rs b/tests/run.rs
index 743c0add..2f881ad9 100644
--- a/tests/run.rs
+++ b/tests/run.rs
@@ -1,10 +1,16 @@
+use std::{
+  collections::HashMap,
+  error::Error,
+  ffi::OsStr,
+  fs,
+  io::{Read, Write},
+  path::{Path, PathBuf},
+  process::{Command, Stdio},
+};
 
 use hvm::ast::Tree;
 use insta::assert_snapshot;
 use TSPL::Parser;
-use std::{
-  collections::HashMap, error::Error, ffi::OsStr, fs, io::Read, path::{Path, PathBuf}, process::{Command, Stdio}
-};
 
 #[test]
 fn test_run_programs() {
@@ -25,33 +31,65 @@ fn manifest_relative(sub: &str) -> PathBuf {
 }
 
 fn test_file(path: &Path) {
-  if fs::read_to_string(path).unwrap().contains("@test-skip = 1") {
+  let contents = fs::read_to_string(path).unwrap();
+  if contents.contains("@test-skip = 1") {
     println!("skipping {path:?}");
     return;
   }
+  if contents.contains("@test-io = 1") {
+    test_io_file(path);
+    return;
+  }
+
   println!("testing {path:?}...");
-  let rust_output = execute_hvm(&["run".as_ref(), path.as_os_str()]).unwrap();
+  let rust_output = execute_hvm(&["run".as_ref(), path.as_os_str()], false).unwrap();
   assert_snapshot!(rust_output);
+
   println!("  testing {path:?}, C...");
-  let c_output = execute_hvm(&["run-c".as_ref(), path.as_os_str()]).unwrap();
+  let c_output = execute_hvm(&["run-c".as_ref(), path.as_os_str()], false).unwrap();
   assert_eq!(c_output, rust_output, "{path:?}: C output does not match rust output");
+
   if cfg!(feature = "cuda") {
     println!("  testing {path:?}, CUDA...");
-    let cuda_output = execute_hvm(&["run-cu".as_ref(), path.as_os_str()]).unwrap();
-    assert_eq!(cuda_output, rust_output, "{path:?}: CUDA output does not match rust output");
+    let cuda_output = execute_hvm(&["run-cu".as_ref(), path.as_os_str()], false).unwrap();
+    assert_eq!(
+      cuda_output, rust_output,
+      "{path:?}: CUDA output does not match rust output"
+    );
+  }
+}
+
+fn test_io_file(path: &Path) {
+  println!("  testing (io) {path:?}, C...");
+  let c_output = execute_hvm(&["run-c".as_ref(), path.as_os_str()], true).unwrap();
+  assert_snapshot!(c_output);
+
+  if cfg!(feature = "cuda") {
+    println!("  testing (io) {path:?}, CUDA...");
+    let cuda_output = execute_hvm(&["run-cu".as_ref(), path.as_os_str()], true).unwrap();
+    assert_eq!(cuda_output, c_output, "{path:?}: CUDA output does not match C output");
   }
 }
 
-fn execute_hvm(args: &[&OsStr]) -> Result<String, Box<dyn Error>> {
+fn execute_hvm(args: &[&OsStr], send_io: bool) -> Result<String, Box<dyn Error>> {
   // Spawn the command
-  let mut child =
-    Command::new(env!("CARGO_BIN_EXE_hvm")).args(args).stdout(Stdio::piped()).stderr(Stdio::piped()).spawn()?;
+  let mut child = Command::new(env!("CARGO_BIN_EXE_hvm"))
+    .args(args)
+    .stdin(Stdio::piped())
+    .stdout(Stdio::piped())
+    .stderr(Stdio::piped())
+    .spawn()?;
 
   // Capture the output of the command
   let mut stdout = child.stdout.take().ok_or("Couldn't capture stdout!")?;
   let mut stderr = child.stderr.take().ok_or("Couldn't capture stderr!")?;
 
   // Wait for the command to finish and get the exit status
+  if send_io {
+    let mut stdin = child.stdin.take().ok_or("Couldn't capture stdin!")?;
+    stdin.write_all(b"io from the tests\n")?;
+    drop(stdin);
+  }
   let status = child.wait()?;
 
   // Read the output
@@ -62,22 +100,27 @@ fn execute_hvm(args: &[&OsStr]) -> Result<String, Box<dyn Error>> {
   Ok(if !status.success() {
     format!("exited with code {status}:\n{output}")
   } else {
-    parse_output(&output).unwrap_or_else(|err| {
-      panic!("error parsing output:\n{err}\n\n{output}")
-    })
+    parse_output(&output).unwrap_or_else(|err| panic!("error parsing output:\n{err}\n\n{output}"))
   })
 }
 
 fn parse_output(output: &str) -> Result<String, String> {
-  let mut parser = hvm::ast::CoreParser::new(output);
-  parser.consume("Result:")?;
-  let mut tree = parser.parse_tree()?;
-  normalize_vars(&mut tree, &mut HashMap::new());
-  // TODO: include iteration count in snapshot once consistent
-  // parser.consume("- ITRS:")?;
-  // let itrs = parser.parse_u64()?;
-  // Ok(format!("Result: {}\n- ITRS: {}", tree.show(), itrs))
-  Ok(format!("Result: {}", tree.show()))
+  let mut lines = Vec::new();
+
+  for line in output.lines() {
+    if line.starts_with("Result:") {
+      let mut parser = hvm::ast::CoreParser::new(line);
+      parser.consume("Result:")?;
+      let mut tree = parser.parse_tree()?;
+      normalize_vars(&mut tree, &mut HashMap::new());
+      lines.push(format!("Result: {}", tree.show()));
+    } else if !line.starts_with("- ITRS:") && !line.starts_with("- TIME:") && !line.starts_with("- MIPS:") {
+      // TODO: include iteration count in snapshot once consistent
+      lines.push(line.to_string())
+    }
+  }
+
+  Ok(lines.join("\n"))
 }
 
 fn normalize_vars(tree: &mut Tree, vars: &mut HashMap<String, usize>) {
@@ -85,7 +128,7 @@ fn normalize_vars(tree: &mut Tree, vars: &mut HashMap<String, usize>) {
     Tree::Var { nam } => {
       let next_var = vars.len();
       *nam = format!("x{}", vars.entry(std::mem::take(nam)).or_insert(next_var));
-    },
+    }
     Tree::Era | Tree::Ref { .. } | Tree::Num { .. } => {}
     Tree::Con { fst, snd } | Tree::Dup { fst, snd } | Tree::Opr { fst, snd } | Tree::Swi { fst, snd } => {
       normalize_vars(fst, vars);
diff --git a/tests/snapshots/run__file@f24.hvm.snap b/tests/snapshots/run__file@numerics__f24.hvm.snap
similarity index 100%
rename from tests/snapshots/run__file@f24.hvm.snap
rename to tests/snapshots/run__file@numerics__f24.hvm.snap
diff --git a/tests/snapshots/run__file@i24.hvm.snap b/tests/snapshots/run__file@numerics__i24.hvm.snap
similarity index 100%
rename from tests/snapshots/run__file@i24.hvm.snap
rename to tests/snapshots/run__file@numerics__i24.hvm.snap
diff --git a/tests/snapshots/run__file@u24.hvm.snap b/tests/snapshots/run__file@numerics__u24.hvm.snap
similarity index 100%
rename from tests/snapshots/run__file@u24.hvm.snap
rename to tests/snapshots/run__file@numerics__u24.hvm.snap
diff --git a/tests/snapshots/run__io_file@io__read_and_print.hvm.snap b/tests/snapshots/run__io_file@io__read_and_print.hvm.snap
new file mode 100644
index 00000000..e00390ac
--- /dev/null
+++ b/tests/snapshots/run__io_file@io__read_and_print.hvm.snap
@@ -0,0 +1,8 @@
+---
+source: tests/run.rs
+expression: c_output
+input_file: tests/programs/io/read_and_print.hvm
+---
+What is your name?
+Hello, 'io from the tests'!
+Result: 42

From efc21a55863097150d7453deb0d61b95d339ba92 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 11:12:04 +0200
Subject: [PATCH 08/12] ignore LEAK output

---
 tests/run.rs | 6 +++++-
 1 file changed, 5 insertions(+), 1 deletion(-)

diff --git a/tests/run.rs b/tests/run.rs
index 2f881ad9..745a4882 100644
--- a/tests/run.rs
+++ b/tests/run.rs
@@ -114,7 +114,11 @@ fn parse_output(output: &str) -> Result<String, String> {
       let mut tree = parser.parse_tree()?;
       normalize_vars(&mut tree, &mut HashMap::new());
       lines.push(format!("Result: {}", tree.show()));
-    } else if !line.starts_with("- ITRS:") && !line.starts_with("- TIME:") && !line.starts_with("- MIPS:") {
+    } else if !line.starts_with("- ITRS:")
+      && !line.starts_with("- TIME:")
+      && !line.starts_with("- MIPS:")
+      && !line.starts_with("- LEAK:")
+    {
       // TODO: include iteration count in snapshot once consistent
       lines.push(line.to_string())
     }

From 45c5395e71b24a04048f655aa9b5ff2c324f0dc2 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 11:30:37 +0200
Subject: [PATCH 09/12] comments, remove flake

---
 examples/demo_io/main.bend | 14 ++++++--------
 flake.nix                  | 26 --------------------------
 src/hvm.cu                 |  7 -------
 3 files changed, 6 insertions(+), 41 deletions(-)
 delete mode 100644 flake.nix

diff --git a/examples/demo_io/main.bend b/examples/demo_io/main.bend
index 05434f0c..5f7ccbd6 100644
--- a/examples/demo_io/main.bend
+++ b/examples/demo_io/main.bend
@@ -23,12 +23,10 @@ def call_io(func, argm):
 
 def main:
   do IO_T:
-    * <- call_io("PUT_TEXT", "what is your name?\n")
-    n <- call_io("GET_LINE", "")
-    c <- call_io("PUT_LINE", "pick a number 0-9?\n")
-    * <- call_io("PUT_TEXT", "hello, '")
-    * <- call_io("PUT_TEXT", n)
-    * <- call_io("PUT_TEXT", "', you picked '")
-    * <- call_io("PUT_TEXT", c)
-    * <- call_io("PUT_TEXT", "'! Nice choice!")
+    * <- call_io("WRITE", (1, "What is your name?\n"))
+    name <- call_io("READ_LINE", 0)
+    * <- call_io("WRITE", (1, "Hello, '"))
+    * <- call_io("WRITE", (1, name))
+    * <- call_io("WRITE", (1, "'!\n"))
+
     return 42
diff --git a/flake.nix b/flake.nix
deleted file mode 100644
index cab61933..00000000
--- a/flake.nix
+++ /dev/null
@@ -1,26 +0,0 @@
-{
-  description = "hvm2";
-  inputs.nixpkgs.url = "github:NixOS/nixpkgs/nixpkgs-unstable";
-  inputs.flake-utils.url = "github:numtide/flake-utils";
-  inputs.fenix.url = "github:nix-community/fenix/monthly";
-
-  outputs = { self, nixpkgs, flake-utils, fenix }:
-    flake-utils.lib.eachDefaultSystem (system:
-      let
-        pkgs = nixpkgs.legacyPackages.${system};
-
-        hoc-upload = pkgs.writeShellScriptBin "hoc-upload" ''
-          scp src/hvm.cu hoc@192.168.194.38:~/enricozb/hvm/src/hvm.cu
-        '';
-      in {
-        devShells.default = pkgs.mkShell {
-          packages = with fenix.packages.${system}; [
-            minimal.rustc
-            minimal.cargo
-            pkgs.clippy
-
-            hoc-upload
-          ];
-        };
-      });
-}
diff --git a/src/hvm.cu b/src/hvm.cu
index 7e4b857e..324c7d74 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -1062,15 +1062,8 @@ __device__ bool get_resources(Net* net, TM* tm, u8 need_rbag, u8 need_node, u8 n
   u32 got_node;
   u32 got_vars;
   if (tm->mode != WORK) {
-    // printf("get resources (not work): %u %u %u", need_rbag, need_node, need_vars);
-
     got_node = g_node_alloc(net, tm, need_node);
-
-    // printf("got node");
-
     got_vars = g_vars_alloc(net, tm, need_vars);
-
-    // printf("got vars");
   } else {
     got_node = l_node_alloc(net, tm, need_node);
     got_vars = l_vars_alloc(net, tm, need_vars);

From 1ca51f0bc48cef236577e9038bed583fce68e057 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 11:31:12 +0200
Subject: [PATCH 10/12] remove io.hvm

---
 io.hvm | 212 ---------------------------------------------------------
 1 file changed, 212 deletions(-)
 delete mode 100644 io.hvm

diff --git a/io.hvm b/io.hvm
deleted file mode 100644
index 9bcef7d0..00000000
--- a/io.hvm
+++ /dev/null
@@ -1,212 +0,0 @@
-@IO_T/Call = (a (b (c (d ((1 (a (b (c (d e))))) e)))))
-
-@IO_T/Done = (a (b ((0 (a (b c))) c)))
-
-@IO_T/MAGIC = (13683217 16719857)
-
-@IO_T/bind = ((@IO_T/bind__C2 a) a)
-
-@IO_T/bind__C0 = (* (a ((a b) b)))
-
-@IO_T/bind__C1 = (* (* (a (b ((c d) (e g))))))
-  & @IO_T/Call ~ (@IO_T/MAGIC (a (b ((c f) g))))
-  & @IO_T/bind ~ (d (e f))
-
-@IO_T/bind__C2 = (?((@IO_T/bind__C0 @IO_T/bind__C1) a) a)
-
-@IO_T/wrap = a
-  & @IO_T/Done ~ (@IO_T/MAGIC a)
-
-@String/Cons = (a (b ((1 (a (b c))) c)))
-
-@String/Nil = ((0 a) a)
-
-@call_io = (a (b c))
-  & @IO_T/Call ~ (@IO_T/MAGIC (a (b (@call_io__C0 c))))
-
-@call_io__C0 = a
-  & @IO_T/Done ~ (@IO_T/MAGIC a)
-
-@main = a
-  & @IO_T/bind ~ (@main__C10 (@main__C9 a))
-
-@main__C0 = x
-  & @call_io ~ (h (w x))
-  & @String/Cons ~ (80 (g h))
-  & @String/Cons ~ (85 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (84 (c d))
-  & @String/Cons ~ (69 (b c))
-  & @String/Cons ~ (88 (a b))
-  & @String/Cons ~ (84 (@String/Nil a))
-  & @String/Cons ~ (39 (v w))
-  & @String/Cons ~ (33 (u v))
-  & @String/Cons ~ (32 (t u))
-  & @String/Cons ~ (78 (s t))
-  & @String/Cons ~ (105 (r s))
-  & @String/Cons ~ (99 (q r))
-  & @String/Cons ~ (101 (p q))
-  & @String/Cons ~ (32 (o p))
-  & @String/Cons ~ (99 (n o))
-  & @String/Cons ~ (104 (m n))
-  & @String/Cons ~ (111 (l m))
-  & @String/Cons ~ (105 (k l))
-  & @String/Cons ~ (99 (j k))
-  & @String/Cons ~ (101 (i j))
-  & @String/Cons ~ (33 (@String/Nil i))
-
-@main__C1 = h
-  & @String/Cons ~ (80 (g h))
-  & @String/Cons ~ (85 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (84 (c d))
-  & @String/Cons ~ (69 (b c))
-  & @String/Cons ~ (88 (a b))
-  & @String/Cons ~ (84 (@String/Nil a))
-
-@main__C10 = bb
-  & @call_io ~ (h (ab bb))
-  & @String/Cons ~ (80 (g h))
-  & @String/Cons ~ (85 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (84 (c d))
-  & @String/Cons ~ (69 (b c))
-  & @String/Cons ~ (88 (a b))
-  & @String/Cons ~ (84 (@String/Nil a))
-  & @String/Cons ~ (119 (z ab))
-  & @String/Cons ~ (104 (y z))
-  & @String/Cons ~ (97 (x y))
-  & @String/Cons ~ (116 (w x))
-  & @String/Cons ~ (32 (v w))
-  & @String/Cons ~ (105 (u v))
-  & @String/Cons ~ (115 (t u))
-  & @String/Cons ~ (32 (s t))
-  & @String/Cons ~ (121 (r s))
-  & @String/Cons ~ (111 (q r))
-  & @String/Cons ~ (117 (p q))
-  & @String/Cons ~ (114 (o p))
-  & @String/Cons ~ (32 (n o))
-  & @String/Cons ~ (110 (m n))
-  & @String/Cons ~ (97 (l m))
-  & @String/Cons ~ (109 (k l))
-  & @String/Cons ~ (101 (j k))
-  & @String/Cons ~ (63 (i j))
-  & @String/Cons ~ (10 (@String/Nil i))
-
-@main__C2 = (* a)
-  & @IO_T/bind ~ (@main__C0 ((* 42) a))
-
-@main__C3 = x
-  & @call_io ~ (h (w x))
-  & @String/Cons ~ (80 (g h))
-  & @String/Cons ~ (85 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (84 (c d))
-  & @String/Cons ~ (69 (b c))
-  & @String/Cons ~ (88 (a b))
-  & @String/Cons ~ (84 (@String/Nil a))
-  & @String/Cons ~ (39 (v w))
-  & @String/Cons ~ (44 (u v))
-  & @String/Cons ~ (32 (t u))
-  & @String/Cons ~ (121 (s t))
-  & @String/Cons ~ (111 (r s))
-  & @String/Cons ~ (117 (q r))
-  & @String/Cons ~ (32 (p q))
-  & @String/Cons ~ (112 (o p))
-  & @String/Cons ~ (105 (n o))
-  & @String/Cons ~ (99 (m n))
-  & @String/Cons ~ (107 (l m))
-  & @String/Cons ~ (101 (k l))
-  & @String/Cons ~ (100 (j k))
-  & @String/Cons ~ (32 (i j))
-  & @String/Cons ~ (39 (@String/Nil i))
-
-@main__C4 = h
-  & @String/Cons ~ (80 (g h))
-  & @String/Cons ~ (85 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (84 (c d))
-  & @String/Cons ~ (69 (b c))
-  & @String/Cons ~ (88 (a b))
-  & @String/Cons ~ (84 (@String/Nil a))
-
-@main__C5 = q
-  & @call_io ~ (h (p q))
-  & @String/Cons ~ (80 (g h))
-  & @String/Cons ~ (85 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (84 (c d))
-  & @String/Cons ~ (69 (b c))
-  & @String/Cons ~ (88 (a b))
-  & @String/Cons ~ (84 (@String/Nil a))
-  & @String/Cons ~ (104 (o p))
-  & @String/Cons ~ (101 (n o))
-  & @String/Cons ~ (108 (m n))
-  & @String/Cons ~ (108 (l m))
-  & @String/Cons ~ (111 (k l))
-  & @String/Cons ~ (44 (j k))
-  & @String/Cons ~ (32 (i j))
-  & @String/Cons ~ (39 (@String/Nil i))
-
-@main__C6 = bb
-  & @call_io ~ (h (ab bb))
-  & @String/Cons ~ (71 (g h))
-  & @String/Cons ~ (69 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (67 (c d))
-  & @String/Cons ~ (72 (b c))
-  & @String/Cons ~ (65 (a b))
-  & @String/Cons ~ (82 (@String/Nil a))
-  & @String/Cons ~ (112 (z ab))
-  & @String/Cons ~ (105 (y z))
-  & @String/Cons ~ (99 (x y))
-  & @String/Cons ~ (107 (w x))
-  & @String/Cons ~ (32 (v w))
-  & @String/Cons ~ (97 (u v))
-  & @String/Cons ~ (32 (t u))
-  & @String/Cons ~ (110 (s t))
-  & @String/Cons ~ (117 (r s))
-  & @String/Cons ~ (109 (q r))
-  & @String/Cons ~ (98 (p q))
-  & @String/Cons ~ (101 (o p))
-  & @String/Cons ~ (114 (n o))
-  & @String/Cons ~ (32 (m n))
-  & @String/Cons ~ (48 (l m))
-  & @String/Cons ~ (45 (k l))
-  & @String/Cons ~ (57 (j k))
-  & @String/Cons ~ (63 (i j))
-  & @String/Cons ~ (10 (@String/Nil i))
-
-@main__C7 = (a i)
-  & @IO_T/bind ~ (@main__C6 ((c h) i))
-  & @IO_T/bind ~ (@main__C5 ((* g) h))
-  & @IO_T/bind ~ (b ((* f) g))
-  & @call_io ~ (@main__C4 (a b))
-  & @IO_T/bind ~ (@main__C3 ((* e) f))
-  & @IO_T/bind ~ (d (@main__C2 e))
-  & @call_io ~ (@main__C1 (c d))
-
-@main__C8 = i
-  & @call_io ~ (h (@String/Nil i))
-  & @String/Cons ~ (71 (g h))
-  & @String/Cons ~ (69 (f g))
-  & @String/Cons ~ (84 (e f))
-  & @String/Cons ~ (95 (d e))
-  & @String/Cons ~ (76 (c d))
-  & @String/Cons ~ (73 (b c))
-  & @String/Cons ~ (78 (a b))
-  & @String/Cons ~ (69 (@String/Nil a))
-
-@main__C9 = (* a)
-  & @IO_T/bind ~ (@main__C8 (@main__C7 a))
-
-@test-skip = 1
-
-

From fdfd9a8c7ad2606b8896f34bb714b75440879238 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 11:31:31 +0200
Subject: [PATCH 11/12] remove flake.lock

---
 flake.lock | 115 -----------------------------------------------------
 1 file changed, 115 deletions(-)
 delete mode 100644 flake.lock

diff --git a/flake.lock b/flake.lock
deleted file mode 100644
index 424b8522..00000000
--- a/flake.lock
+++ /dev/null
@@ -1,115 +0,0 @@
-{
-  "nodes": {
-    "fenix": {
-      "inputs": {
-        "nixpkgs": "nixpkgs",
-        "rust-analyzer-src": "rust-analyzer-src"
-      },
-      "locked": {
-        "lastModified": 1714544767,
-        "narHash": "sha256-kF1bX+YFMedf1g0PAJYwGUkzh22JmULtj8Rm4IXAQKs=",
-        "owner": "nix-community",
-        "repo": "fenix",
-        "rev": "73124e1356bde9411b163d636b39fe4804b7ca45",
-        "type": "github"
-      },
-      "original": {
-        "owner": "nix-community",
-        "ref": "monthly",
-        "repo": "fenix",
-        "type": "github"
-      }
-    },
-    "flake-utils": {
-      "inputs": {
-        "systems": "systems"
-      },
-      "locked": {
-        "lastModified": 1710146030,
-        "narHash": "sha256-SZ5L6eA7HJ/nmkzGG7/ISclqe6oZdOZTNoesiInkXPQ=",
-        "owner": "numtide",
-        "repo": "flake-utils",
-        "rev": "b1d9ab70662946ef0850d488da1c9019f3a9752a",
-        "type": "github"
-      },
-      "original": {
-        "owner": "numtide",
-        "repo": "flake-utils",
-        "type": "github"
-      }
-    },
-    "nixpkgs": {
-      "locked": {
-        "lastModified": 1714253743,
-        "narHash": "sha256-mdTQw2XlariysyScCv2tTE45QSU9v/ezLcHJ22f0Nxc=",
-        "owner": "nixos",
-        "repo": "nixpkgs",
-        "rev": "58a1abdbae3217ca6b702f03d3b35125d88a2994",
-        "type": "github"
-      },
-      "original": {
-        "owner": "nixos",
-        "ref": "nixos-unstable",
-        "repo": "nixpkgs",
-        "type": "github"
-      }
-    },
-    "nixpkgs_2": {
-      "locked": {
-        "lastModified": 1716312448,
-        "narHash": "sha256-PH3w5av8d+TdwCkiWN4UPBTxrD9MpxIQPDVWctlomVo=",
-        "owner": "NixOS",
-        "repo": "nixpkgs",
-        "rev": "e381a1288138aceda0ac63db32c7be545b446921",
-        "type": "github"
-      },
-      "original": {
-        "owner": "NixOS",
-        "ref": "nixpkgs-unstable",
-        "repo": "nixpkgs",
-        "type": "github"
-      }
-    },
-    "root": {
-      "inputs": {
-        "fenix": "fenix",
-        "flake-utils": "flake-utils",
-        "nixpkgs": "nixpkgs_2"
-      }
-    },
-    "rust-analyzer-src": {
-      "flake": false,
-      "locked": {
-        "lastModified": 1714501997,
-        "narHash": "sha256-g31zfxwUFzkPgX0Q8sZLcrqGmOxwjEZ/iqJjNx4fEGo=",
-        "owner": "rust-lang",
-        "repo": "rust-analyzer",
-        "rev": "49e502b277a8126a9ad10c802d1aaa3ef1a280ef",
-        "type": "github"
-      },
-      "original": {
-        "owner": "rust-lang",
-        "ref": "nightly",
-        "repo": "rust-analyzer",
-        "type": "github"
-      }
-    },
-    "systems": {
-      "locked": {
-        "lastModified": 1681028828,
-        "narHash": "sha256-Vy1rq5AaRuLzOxct8nz4T6wlgyUR7zLU309k9mBC768=",
-        "owner": "nix-systems",
-        "repo": "default",
-        "rev": "da67096a3b9bf56a91d16901293e51ba5b49a27e",
-        "type": "github"
-      },
-      "original": {
-        "owner": "nix-systems",
-        "repo": "default",
-        "type": "github"
-      }
-    }
-  },
-  "root": "root",
-  "version": 7
-}

From ea8e3ef050a65ed596d75d9c101f5b3b62291801 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 31 May 2024 18:05:53 +0200
Subject: [PATCH 12/12] readback & inject

---
 src/hvm.c  | 46 +++++++++++++++++++++++-----------------------
 src/hvm.cu | 36 ++++++++++++++++++------------------
 2 files changed, 41 insertions(+), 41 deletions(-)

diff --git a/src/hvm.c b/src/hvm.c
index a7c9e382..b76d27d7 100644
--- a/src/hvm.c
+++ b/src/hvm.c
@@ -1235,7 +1235,7 @@ Port expand(Net* net, Book* book, Port port) {
 
 // Reads back a λ-Encoded constructor from device to host.
 // Encoding: λt ((((t TAG) arg0) arg1) ...)
-Ctr port_to_ctr(Net* net, Book* book, Port port) {
+Ctr readback_ctr(Net* net, Book* book, Port port) {
   Ctr ctr;
   ctr.tag = -1;
   ctr.args_len = 0;
@@ -1273,7 +1273,7 @@ Ctr port_to_ctr(Net* net, Book* book, Port port) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Str port_to_str(Net* net, Book* book, Port port) {
+Str readback_str(Net* net, Book* book, Port port) {
   // Result
   Str str;
   str.text_len = 0;
@@ -1286,7 +1286,7 @@ Str port_to_str(Net* net, Book* book, Port port) {
     //printf("reading str %s\n", show_port(peek(net, port)).x);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = port_to_ctr(net, book, peek(net, port));
+    Ctr ctr = readback_ctr(net, book, peek(net, port));
 
     //printf("reading tag %d | len %d\n", ctr.tag, ctr.args_len);
 
@@ -1315,9 +1315,9 @@ Str port_to_str(Net* net, Book* book, Port port) {
 }
 
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
-/// Should only be called within `str_to_port`, as a previous call
+/// Should only be called within `inject_str`, as a previous call
 /// to `get_resources` is expected.
-Port nil_port(Net* net) {
+Port inject_nil(Net* net) {
   u32 v1 = tm[0]->vloc[0];
 
   u32 n1 = tm[0]->nloc[0];
@@ -1333,12 +1333,12 @@ Port nil_port(Net* net) {
 }
 
 /// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
-/// Should only be called within `str_to_port`, as a previous call
+/// Should only be called within `inject_str`, as a previous call
 /// to `get_resources` is expected.
 /// The `char_idx` parameter is used to offset the vloc and nloc
 /// allocations, otherwise they would conflict with each other on
 /// subsequent calls.
-Port cons_port(Net* net, Port head, Port tail, u32 char_idx) {
+Port inject_cons(Net* net, Port head, Port tail, u32 char_idx) {
   u32 v1 = tm[0]->vloc[1 + char_idx];
 
   u32 n1 = tm[0]->nloc[2 + char_idx * 4 + 0];
@@ -1363,21 +1363,21 @@ Port cons_port(Net* net, Port head, Port tail, u32 char_idx) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Port str_to_port(Net* net, Str *str) {
+Port inject_str(Net* net, Str *str) {
   // Allocate all resources up front:
   // - NIL needs  2 nodes & 1 var
   // - CONS needs 4 nodes & 1 var
   u32 len = str->text_len;
   if (!get_resources(net, tm[0], 0, 2 + 4 * len, 1 + len)) {
-    printf("str_to_port: failed to get resources\n");
+    printf("inject_str: failed to get resources\n");
     return new_port(ERA, 0);
   }
 
-  Port port = nil_port(net);
+  Port port = inject_nil(net);
 
   for (u32 i = 0; i < len; i++) {
     Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
-    port = cons_port(net, chr, port, i);
+    port = inject_cons(net, chr, port, i);
   }
 
   return port;
@@ -1445,7 +1445,7 @@ void read_img(Net* net, Port port, u32 width, u32 height, u32* buffer) {
 static FILE* FILE_POINTERS[256];
 
 // Converts a NUM port (file descriptor) to file pointer.
-FILE* port_to_file(Port port) {
+FILE* readback_file(Port port) {
   if (get_tag(port) != NUM) {
     fprintf(stderr, "non-num where file descriptor was expected: %i\n", get_tag(port));
     return NULL;
@@ -1468,7 +1468,7 @@ FILE* port_to_file(Port port) {
 
 // Reads a single char from `argm`.
 Port io_read_char(Net* net, Book* book, Port argm) {
-  FILE* fp = port_to_file(peek(net, argm));
+  FILE* fp = readback_file(peek(net, argm));
   if (fp == NULL) {
     return new_port(ERA, 0);
   }
@@ -1480,12 +1480,12 @@ Port io_read_char(Net* net, Book* book, Port argm) {
   str.text_buf[1] = 0;
   str.text_len = 1;
 
-  return str_to_port(net, &str);
+  return inject_str(net, &str);
 }
 
 // Reads from `argm` at most 255 characters or until a newline is seen.
 Port io_read_line(Net* net, Book* book, Port argm) {
-  FILE* fp = port_to_file(peek(net, argm));
+  FILE* fp = readback_file(peek(net, argm));
   if (fp == NULL) {
     fprintf(stderr, "io_read_line: invalid file descriptor\n");
     return new_port(ERA, 0);
@@ -1506,7 +1506,7 @@ Port io_read_line(Net* net, Book* book, Port argm) {
   }
 
   // Convert it to a port.
-  return str_to_port(net, &str);
+  return inject_str(net, &str);
 }
 
 // Opens a file with the provided mode.
@@ -1519,8 +1519,8 @@ Port io_open_file(Net* net, Book* book, Port argm) {
   }
 
   Pair args = node_load(net, get_val(argm));
-  Str name = port_to_str(net, book, get_fst(args));
-  Str mode = port_to_str(net, book, get_snd(args));
+  Str name = readback_str(net, book, get_fst(args));
+  Str mode = readback_str(net, book, get_snd(args));
 
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
@@ -1536,7 +1536,7 @@ Port io_open_file(Net* net, Book* book, Port argm) {
 
 // Closes a file, reclaiming the file descriptor.
 Port io_close_file(Net* net, Book* book, Port argm) {
-  FILE* fp = port_to_file(peek(net, argm));
+  FILE* fp = readback_file(peek(net, argm));
   if (fp == NULL) {
     fprintf(stderr, "io_close_file: failed to close\n");
     return new_port(ERA, 0);
@@ -1563,8 +1563,8 @@ Port io_write(Net* net, Book* book, Port argm) {
   }
 
   Pair args = node_load(net, get_val(argm));
-  FILE* fp = port_to_file(peek(net, get_fst(args)));
-  Str str = port_to_str(net, book, get_snd(args));
+  FILE* fp = readback_file(peek(net, get_fst(args)));
+  Str str = readback_str(net, book, get_snd(args));
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
@@ -1734,7 +1734,7 @@ void do_run_io(Net* net, Book* book, Port port) {
     normalize(net, book);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = port_to_ctr(net, book, peek(net, port));
+    Ctr ctr = readback_ctr(net, book, peek(net, port));
 
     // Checks if IO Magic Number is a CON
     if (get_tag(ctr.args_buf[0]) != CON) {
@@ -1750,7 +1750,7 @@ void do_run_io(Net* net, Book* book, Port port) {
 
     switch (ctr.tag) {
       case IO_CALL: {
-        Str  func = port_to_str(net, book, ctr.args_buf[1]);
+        Str  func = readback_str(net, book, ctr.args_buf[1]);
         Port argm = ctr.args_buf[2];
         Port cont = ctr.args_buf[3];
         u32  lps  = 0;
diff --git a/src/hvm.cu b/src/hvm.cu
index 324c7d74..78573d6c 100644
--- a/src/hvm.cu
+++ b/src/hvm.cu
@@ -1446,7 +1446,7 @@ __global__ void boot_redex(GNet* gnet, Pair redex) {
 }
 
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
-/// Should only be called within `str_to_port`, as a previous call
+/// Should only be called within `inject_str`, as a previous call
 /// to `get_resources` is expected.
 __device__ Port nil_port(Net* net, TM* tm) {
   u32 v1 = tm->vloc[0];
@@ -1464,7 +1464,7 @@ __device__ Port nil_port(Net* net, TM* tm) {
 }
 
 /// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
-/// Should only be called within `str_to_port`, as a previous call
+/// Should only be called within `inject_str`, as a previous call
 /// to `get_resources` is expected.
 /// The `char_idx` parameter is used to offset the vloc and nloc
 /// allocations, otherwise they would conflict with each other on
@@ -1494,13 +1494,13 @@ __device__ Port cons_port(Net* net, TM* tm, Port head, Port tail, u32 char_idx)
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-__device__ Port str_to_port(Net* net, TM* tm, Str *str) {
+__device__ Port inject_str(Net* net, TM* tm, Str *str) {
   // Allocate all resources up front:
   // - NIL needs  2 nodes & 1 var
   // - CONS needs 4 nodes & 1 var
   u32 len = str->text_len;
   if (!get_resources(net, tm, 0, 2 + 4 * len, 1 + len)) {
-    printf("str_to_port: failed to get resources\n");
+    printf("inject_str: failed to get resources\n");
     return new_port(ERA, 0);
   }
 
@@ -1518,7 +1518,7 @@ __global__ void make_str_port(GNet* gnet, Str *str, Port* ret) {
   if (GID() == 0) {
     TM tm;
     Net net = vnet_new(gnet, NULL, gnet->turn);
-    *ret = str_to_port(&net, &tm, str);
+    *ret = inject_str(&net, &tm, str);
   }
 }
 
@@ -1917,7 +1917,7 @@ Port gnet_make_node(GNet* gnet, Tag tag, Port fst, Port snd) {
 
 // Reads back a λ-Encoded constructor from device to host.
 // Encoding: λt ((((t TAG) arg0) arg1) ...)
-Ctr gnet_port_to_ctr(GNet* gnet, Port port) {
+Ctr gnet_readback_ctr(GNet* gnet, Port port) {
   Ctr ctr;
   ctr.tag = -1;
   ctr.args_len = 0;
@@ -1955,7 +1955,7 @@ Ctr gnet_port_to_ctr(GNet* gnet, Port port) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Str gnet_port_to_str(GNet* gnet, Port port) {
+Str gnet_inject_str(GNet* gnet, Port port) {
   // Result
   Str str;
   str.text_len = 0;
@@ -1966,7 +1966,7 @@ Str gnet_port_to_str(GNet* gnet, Port port) {
     gnet_normalize(gnet);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = gnet_port_to_ctr(gnet, gnet_peek(gnet, port));
+    Ctr ctr = gnet_readback_ctr(gnet, gnet_peek(gnet, port));
 
     // Reads string layer
     switch (ctr.tag) {
@@ -2027,7 +2027,7 @@ Port gnet_make_str(GNet* gnet, Str *str) {
 static FILE* FILE_POINTERS[256];
 
 // Converts a NUM port (file descriptor) to file pointer.
-FILE* port_to_file(Port port) {
+FILE* readback_file(Port port) {
   if (get_tag(port) != NUM) {
     fprintf(stderr, "non-num where file descriptor was expected: %i\n", get_tag(port));
     return NULL;
@@ -2050,7 +2050,7 @@ FILE* port_to_file(Port port) {
 
 // Reads a single char from `argm`.
 Port io_read_char(GNet* gnet, Port argm) {
-  FILE* fp = port_to_file(gnet_peek(gnet, argm));
+  FILE* fp = readback_file(gnet_peek(gnet, argm));
   if (fp == NULL) {
     return new_port(ERA, 0);
   }
@@ -2067,7 +2067,7 @@ Port io_read_char(GNet* gnet, Port argm) {
 
 // Reads from `argm` at most 255 characters or until a newline is seen.
 Port io_read_line(GNet* gnet, Port argm) {
-  FILE* fp = port_to_file(gnet_peek(gnet, argm));
+  FILE* fp = readback_file(gnet_peek(gnet, argm));
   if (fp == NULL) {
     fprintf(stderr, "io_read_line: invalid file descriptor\n");
     return new_port(ERA, 0);
@@ -2101,8 +2101,8 @@ Port io_open_file(GNet* gnet, Port argm) {
   }
 
   Pair args = gnet_node_load(gnet, get_val(argm));
-  Str name = gnet_port_to_str(gnet, get_fst(args));
-  Str mode = gnet_port_to_str(gnet, get_snd(args));
+  Str name = gnet_inject_str(gnet, get_fst(args));
+  Str mode = gnet_inject_str(gnet, get_snd(args));
 
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
@@ -2118,7 +2118,7 @@ Port io_open_file(GNet* gnet, Port argm) {
 
 // Closes a file, reclaiming the file descriptor.
 Port io_close_file(GNet* gnet, Port argm) {
-  FILE* fp = port_to_file(gnet_peek(gnet, argm));
+  FILE* fp = readback_file(gnet_peek(gnet, argm));
   if (fp == NULL) {
     fprintf(stderr, "io_close_file: failed to close\n");
     return new_port(ERA, 0);
@@ -2145,8 +2145,8 @@ Port io_write(GNet* gnet, Port argm) {
   }
 
   Pair args = gnet_node_load(gnet, get_val(argm));
-  FILE* fp = port_to_file(gnet_peek(gnet, get_fst(args)));
-  Str str = gnet_port_to_str(gnet, get_snd(args));
+  FILE* fp = readback_file(gnet_peek(gnet, get_fst(args)));
+  Str str = gnet_inject_str(gnet, get_snd(args));
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
@@ -2203,7 +2203,7 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
     gnet_normalize(gnet);
 
     // Reads the λ-Encoded Ctr
-    Ctr ctr = gnet_port_to_ctr(gnet, gnet_peek(gnet, port));
+    Ctr ctr = gnet_readback_ctr(gnet, gnet_peek(gnet, port));
 
     // Checks if IO Magic Number is a CON
     if (get_tag(ctr.args_buf[0]) != CON) {
@@ -2219,7 +2219,7 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
 
     switch (ctr.tag) {
       case IO_CALL: {
-        Str  func = gnet_port_to_str(gnet, ctr.args_buf[1]);
+        Str  func = gnet_inject_str(gnet, ctr.args_buf[1]);
         FFn* ffn  = NULL;
         // FIXME: optimize this linear search
         for (u32 fid = 0; fid < book->ffns_len; ++fid) {