From 55a42b922a5ed73315457a94f86665c10af8091e Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Mon, 10 Jun 2024 14:43:36 +0200
Subject: [PATCH 01/11] add io_seek

---
 src/run.c | 189 +++++++++++++++++++++++++++++++++++++++++++-----------
 1 file changed, 150 insertions(+), 39 deletions(-)

diff --git a/src/run.c b/src/run.c
index d71687f5..27eac33c 100644
--- a/src/run.c
+++ b/src/run.c
@@ -15,6 +15,14 @@ typedef struct Str {
   char text_buf[256];
 } Str;
 
+// Readback: λ-Encoded list of bytes
+typedef struct Bytes {
+  u32  len;
+  char *buf;
+} Bytes;
+
+#define MAX_BYTES 256
+
 // IO Magic Number
 #define IO_MAGIC_0 0xD0CA11
 #define IO_MAGIC_1 0xFF1FF1
@@ -64,7 +72,7 @@ Ctr readback_ctr(Net* net, Book* book, Port port) {
   return ctr;
 }
 
-// Converts a UTF-32 (truncated to 24 bits) string to a Port.
+// Converts a Port into a UTF-32 (truncated to 24 bits) string.
 // Since unicode scalars can fit in 21 bits, HVM's u24
 // integers can contain any unicode scalar value.
 // Encoding:
@@ -80,13 +88,9 @@ Str readback_str(Net* net, Book* book, Port port) {
     // Normalizes the net
     normalize(net, book);
 
-    //printf("reading str %s\n", show_port(peek(net, port)).x);
-
     // Reads the λ-Encoded Ctr
     Ctr ctr = readback_ctr(net, book, peek(net, port));
 
-    //printf("reading tag %d | len %d\n", ctr.tag, ctr.args_len);
-
     // Reads string layer
     switch (ctr.tag) {
       case LIST_NIL: {
@@ -95,8 +99,7 @@ Str readback_str(Net* net, Book* book, Port port) {
       case LIST_CONS: {
         if (ctr.args_len != 2) break;
         if (get_tag(ctr.args_buf[0]) != NUM) break;
-        if (str.text_len >= 256) { printf("ERROR: for now, HVM can only readback strings of length <256."); break; }
-        //printf("reading chr %d\n", get_u24(get_val(ctr.args_buf[0])));
+        if (str.text_len > 256) { printf("ERROR: for now, HVM can only readback strings of length <256."); break; }
         str.text_buf[str.text_len++] = get_u24(get_val(ctr.args_buf[0]));
         boot_redex(net, new_pair(ctr.args_buf[1], ROOT));
         port = ROOT;
@@ -111,6 +114,45 @@ Str readback_str(Net* net, Book* book, Port port) {
   return str;
 }
 
+// Converts a Port into a list of bytes.
+// Encoding:
+// - λt (t NIL)
+// - λt (((t CONS) head) tail)
+Bytes readback_bytes(Net* net, Book* book, Port port) {
+  // Result
+  Bytes bytes;
+  bytes.buf = malloc(sizeof(char) * MAX_BYTES);
+  bytes.len = 0;
+
+  // Readback loop
+  while (TRUE) {
+    // Normalizes the net
+    normalize(net, book);
+
+    // Reads the λ-Encoded Ctr
+    Ctr ctr = readback_ctr(net, book, peek(net, port));
+
+    // Reads string layer
+    switch (ctr.tag) {
+      case LIST_NIL: {
+        break;
+      }
+      case LIST_CONS: {
+        if (ctr.args_len != 2) break;
+        if (get_tag(ctr.args_buf[0]) != NUM) break;
+        if (bytes.len >= MAX_BYTES) { printf("ERROR: for now, HVM can only readback list of bytes of length <=%u.", MAX_BYTES); break; }
+        bytes.buf[bytes.len++] = get_u24(get_val(ctr.args_buf[0]));
+        boot_redex(net, new_pair(ctr.args_buf[1], ROOT));
+        port = ROOT;
+        continue;
+      }
+    }
+    break;
+  }
+
+  return bytes;
+}
+
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
 /// Should only be called within `inject_str`, as a previous call
 /// to `get_resources` is expected.
@@ -180,6 +222,30 @@ Port inject_str(Net* net, Str *str) {
   return port;
 }
 
+// Converts a list of bytes to a Port.
+// Encoding:
+// - λt (t NIL)
+// - λt (((t CONS) head) tail)
+Port inject_bytes(Net* net, Bytes *bytes) {
+  // Allocate all resources up front:
+  // - NIL needs  2 nodes & 1 var
+  // - CONS needs 4 nodes & 1 var
+  u32 len = bytes->len;
+  if (!get_resources(net, tm[0], 0, 2 + 4 * len, 1 + len)) {
+    printf("inject_str: failed to get resources\n");
+    return new_port(ERA, 0);
+  }
+
+  Port port = inject_nil(net);
+
+  for (u32 i = 0; i < len; i++) {
+    Port byte = new_port(NUM, new_u24(bytes->buf[len - i - 1]));
+    port = inject_cons(net, byte, port, i);
+  }
+
+  return port;
+}
+
 // Primitive IO Fns
 // -----------------
 
@@ -213,47 +279,38 @@ FILE* readback_file(Port port) {
   return fp;
 }
 
-// Reads a single char from `argm`.
-Port io_read_char(Net* net, Book* book, Port argm) {
-  FILE* fp = readback_file(peek(net, argm));
-  if (fp == NULL) {
+// Reads from a file a specified number of bytes.
+// `argm` is a tuple of (file_descriptor, num_bytes).
+Port io_read(Net* net, Book* book, Port argm) {
+  if (get_tag(peek(net, argm)) != CON) {
+    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(peek(net, argm)));
     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 inject_str(net, &str);
-}
+  Pair args = node_load(net, get_val(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 = readback_file(peek(net, argm));
+  FILE* fp = readback_file(peek(net, get_fst(args)));
   if (fp == NULL) {
     fprintf(stderr, "io_read_line: invalid file descriptor\n");
     return new_port(ERA, 0);
   }
 
+  u32 num_bytes = get_u24(get_val(peek(net, get_snd(args))));
+
   /// Read a string.
-  Str str;
+  Bytes bytes;
+  bytes.buf = malloc(sizeof(char) * num_bytes);
+  bytes.len = fread(bytes.buf, sizeof(char), num_bytes, fp);
 
-  if (fgets(str.text_buf, sizeof(str.text_buf), fp) == NULL) {
+  if ((bytes.len != num_bytes) && ferror(fp)) {
     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 inject_str(net, &str);
+  Port ret = inject_bytes(net, &bytes);
+  free(bytes.buf);
+
+  return ret;
 }
 
 // Opens a file with the provided mode.
@@ -269,6 +326,8 @@ Port io_open_file(Net* net, Book* book, Port argm) {
   Str name = readback_str(net, book, get_fst(args));
   Str mode = readback_str(net, book, get_snd(args));
 
+  printf("opening file '%s' with mode '%s'\n", name.text_buf, mode.text_buf);
+
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
       FILE_POINTERS[fd] = fopen(name.text_buf, mode.text_buf);
@@ -311,17 +370,69 @@ Port io_write(Net* net, Book* book, Port argm) {
 
   Pair args = node_load(net, get_val(argm));
   FILE* fp = readback_file(peek(net, get_fst(args)));
-  Str str = readback_str(net, book, get_snd(args));
+  Bytes bytes = readback_bytes(net, book, get_snd(args));
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
+    free(bytes.buf);
+
     return new_port(ERA, 0);
   }
 
-  if (fputs(str.text_buf, fp) == EOF) {
+  if (fwrite(bytes.buf, sizeof(char), bytes.len, fp) != bytes.len) {
     fprintf(stderr, "io_write: failed to write\n");
   }
 
+  free(bytes.buf);
+
+  return new_port(ERA, 0);
+}
+
+// Seeks to a position in a file.
+// `argm` is a 3-tuple (CON fd (CON offset whence)), where
+// - fd is a file descriptor
+// - offset is a signed byte offset
+// - whence is what that offset is relative to:
+//    - 0 (SEEK_SET): beginning of file
+//    - 1 (SEEK_CUR): current position of the file pointer
+//    - 2 (SEEK_END): end of the file
+Port io_seek(Net* net, Book* book, Port argm) {
+  if (get_tag(peek(net, argm)) != CON) {
+    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(peek(net, argm)));
+    return new_port(ERA, 0);
+  }
+
+  Pair args1 = node_load(net, get_val(argm));
+  if (get_tag(peek(net, get_snd(args1))) != CON) {
+    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(peek(net, get_snd(args1))));
+    return new_port(ERA, 0);
+  }
+
+  Pair args2 = node_load(net, get_val(peek(net, get_snd(args1))));
+
+  FILE* fp = readback_file(peek(net, get_fst(args1)));
+  if (fp == NULL) {
+    fprintf(stderr, "io_write: invalid file descriptor\n");
+    return new_port(ERA, 0);
+  }
+
+  i32 offset = get_i24(get_val(peek(net, get_fst(args2))));
+  u32 whence = get_i24(get_val(peek(net, get_snd(args2))));
+
+  int cwhence;
+  switch (whence) {
+    case 0: cwhence = SEEK_SET; break;
+    case 1: cwhence = SEEK_CUR; break;
+    case 2: cwhence = SEEK_END; break;
+    default:
+      fprintf(stderr, "io_seek: invalid whence\n");
+      return new_port(ERA, 0);
+  }
+
+  if (fseek(fp, offset, cwhence) != 0) {
+    fprintf(stderr, "io_seek: failed to seek\n");
+  }
+
   return new_port(ERA, 0);
 }
 
@@ -365,11 +476,11 @@ Port io_sleep(Net* net, Book* book, Port argm) {
 // -----------
 
 void book_init(Book* book) {
-  book->ffns_buf[book->ffns_len++] = (FFn){"READ_CHAR", io_read_char};
-  book->ffns_buf[book->ffns_len++] = (FFn){"READ_LINE", io_read_line};
-  book->ffns_buf[book->ffns_len++] = (FFn){"OPEN_FILE", io_open_file};
-  book->ffns_buf[book->ffns_len++] = (FFn){"CLOSE_FILE", io_close_file};
+  book->ffns_buf[book->ffns_len++] = (FFn){"READ", io_read};
+  book->ffns_buf[book->ffns_len++] = (FFn){"OPEN", io_open_file};
+  book->ffns_buf[book->ffns_len++] = (FFn){"CLOSE", io_close_file};
   book->ffns_buf[book->ffns_len++] = (FFn){"WRITE", io_write};
+  book->ffns_buf[book->ffns_len++] = (FFn){"SEEK", io_seek};
   book->ffns_buf[book->ffns_len++] = (FFn){"GET_TIME", io_get_time};
   book->ffns_buf[book->ffns_len++] = (FFn){"SLEEP", io_sleep};
 }

From 1121eaa2aad70b08361fb246831f9c15b488c37b Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Mon, 10 Jun 2024 21:57:43 +0200
Subject: [PATCH 02/11] io_seec on cuda (broken :/)

---
 src/run.c  |  64 +++++----------
 src/run.cu | 231 +++++++++++++++++++++++++++++++++++++++--------------
 2 files changed, 192 insertions(+), 103 deletions(-)

diff --git a/src/run.c b/src/run.c
index 27eac33c..34d7b127 100644
--- a/src/run.c
+++ b/src/run.c
@@ -99,7 +99,7 @@ Str readback_str(Net* net, Book* book, Port port) {
       case LIST_CONS: {
         if (ctr.args_len != 2) break;
         if (get_tag(ctr.args_buf[0]) != NUM) break;
-        if (str.text_len > 256) { printf("ERROR: for now, HVM can only readback strings of length <256."); break; }
+        if (str.text_len >= 256) { printf("ERROR: for now, HVM can only readback strings of length <256."); break; }
         str.text_buf[str.text_len++] = get_u24(get_val(ctr.args_buf[0]));
         boot_redex(net, new_pair(ctr.args_buf[1], ROOT));
         port = ROOT;
@@ -121,7 +121,7 @@ Str readback_str(Net* net, Book* book, Port port) {
 Bytes readback_bytes(Net* net, Book* book, Port port) {
   // Result
   Bytes bytes;
-  bytes.buf = malloc(sizeof(char) * MAX_BYTES);
+  bytes.buf = (char*) malloc(sizeof(char) * MAX_BYTES);
   bytes.len = 0;
 
   // Readback loop
@@ -140,7 +140,7 @@ Bytes readback_bytes(Net* net, Book* book, Port port) {
       case LIST_CONS: {
         if (ctr.args_len != 2) break;
         if (get_tag(ctr.args_buf[0]) != NUM) break;
-        if (bytes.len >= MAX_BYTES) { printf("ERROR: for now, HVM can only readback list of bytes of length <=%u.", MAX_BYTES); break; }
+        if (bytes.len >= MAX_BYTES) { printf("ERROR: for now, HVM can only readback list of bytes of length <%u.", MAX_BYTES); break; }
         bytes.buf[bytes.len++] = get_u24(get_val(ctr.args_buf[0]));
         boot_redex(net, new_pair(ctr.args_buf[1], ROOT));
         port = ROOT;
@@ -154,7 +154,7 @@ Bytes readback_bytes(Net* net, Book* book, Port port) {
 }
 
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
-/// Should only be called within `inject_str`, as a previous call
+/// Should only be called within `inject_bytes`, as a previous call
 /// to `get_resources` is expected.
 Port inject_nil(Net* net) {
   u32 v1 = tm[0]->vloc[0];
@@ -172,7 +172,7 @@ Port inject_nil(Net* net) {
 }
 
 /// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
-/// Should only be called within `inject_str`, as a previous call
+/// Should only be called within `inject_bytes`, 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
@@ -196,32 +196,6 @@ Port inject_cons(Net* net, Port head, Port tail, u32 char_idx) {
   return new_port(CON, n4);
 }
 
-// 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 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("inject_str: failed to get resources\n");
-    return new_port(ERA, 0);
-  }
-
-  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 = inject_cons(net, chr, port, i);
-  }
-
-  return port;
-}
-
 // Converts a list of bytes to a Port.
 // Encoding:
 // - λt (t NIL)
@@ -232,7 +206,7 @@ Port inject_bytes(Net* net, Bytes *bytes) {
   // - CONS needs 4 nodes & 1 var
   u32 len = bytes->len;
   if (!get_resources(net, tm[0], 0, 2 + 4 * len, 1 + len)) {
-    printf("inject_str: failed to get resources\n");
+    printf("inject_bytes: failed to get resources\n");
     return new_port(ERA, 0);
   }
 
@@ -291,7 +265,7 @@ Port io_read(Net* net, Book* book, Port argm) {
 
   FILE* fp = readback_file(peek(net, get_fst(args)));
   if (fp == NULL) {
-    fprintf(stderr, "io_read_line: invalid file descriptor\n");
+    fprintf(stderr, "io_read: invalid file descriptor\n");
     return new_port(ERA, 0);
   }
 
@@ -299,11 +273,11 @@ Port io_read(Net* net, Book* book, Port argm) {
 
   /// Read a string.
   Bytes bytes;
-  bytes.buf = malloc(sizeof(char) * num_bytes);
+  bytes.buf = (char*) malloc(sizeof(char) * num_bytes);
   bytes.len = fread(bytes.buf, sizeof(char), num_bytes, fp);
 
   if ((bytes.len != num_bytes) && ferror(fp)) {
-    fprintf(stderr, "io_read_line: failed to read\n");
+    fprintf(stderr, "io_read: failed to read\n");
   }
 
   // Convert it to a port.
@@ -316,9 +290,9 @@ Port io_read(Net* net, Book* book, Port argm) {
 // 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) {
+Port io_open(Net* net, Book* book, Port argm) {
   if (get_tag(peek(net, argm)) != CON) {
-    fprintf(stderr, "io_open_file: expected tuple\n");
+    fprintf(stderr, "io_open: expected tuple\n");
     return new_port(ERA, 0);
   }
 
@@ -335,22 +309,22 @@ Port io_open_file(Net* net, Book* book, Port argm) {
     }
   }
 
-  fprintf(stderr, "io_open_file: too many open files\n");
+  fprintf(stderr, "io_open: too many open files\n");
 
   return new_port(ERA, 0);
 }
 
 // Closes a file, reclaiming the file descriptor.
-Port io_close_file(Net* net, Book* book, Port argm) {
+Port io_close(Net* net, Book* book, Port argm) {
   FILE* fp = readback_file(peek(net, argm));
   if (fp == NULL) {
-    fprintf(stderr, "io_close_file: failed to close\n");
+    fprintf(stderr, "io_close: 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);
+    fprintf(stderr, "io_close: failed to close: %i\n", err);
     return new_port(ERA, 0);
   }
 
@@ -359,9 +333,9 @@ Port io_close_file(Net* net, Book* book, Port argm) {
   return new_port(ERA, 0);
 }
 
-// Writes a string to a file.
+// Writes a list of bytes to a file.
 // `argm` is a tuple (CON node) of the
-// file descriptor and string to write.
+// file descriptor and list of bytes 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\n", get_tag(peek(net, argm)));
@@ -477,8 +451,8 @@ Port io_sleep(Net* net, Book* book, Port argm) {
 
 void book_init(Book* book) {
   book->ffns_buf[book->ffns_len++] = (FFn){"READ", io_read};
-  book->ffns_buf[book->ffns_len++] = (FFn){"OPEN", io_open_file};
-  book->ffns_buf[book->ffns_len++] = (FFn){"CLOSE", io_close_file};
+  book->ffns_buf[book->ffns_len++] = (FFn){"OPEN", io_open};
+  book->ffns_buf[book->ffns_len++] = (FFn){"CLOSE", io_close};
   book->ffns_buf[book->ffns_len++] = (FFn){"WRITE", io_write};
   book->ffns_buf[book->ffns_len++] = (FFn){"SEEK", io_seek};
   book->ffns_buf[book->ffns_len++] = (FFn){"GET_TIME", io_get_time};
diff --git a/src/run.cu b/src/run.cu
index 8893e263..3b0f5796 100644
--- a/src/run.cu
+++ b/src/run.cu
@@ -15,6 +15,14 @@ struct Str {
   char text_buf[256];
 };
 
+// Readback: λ-Encoded list of bytes
+typedef struct Bytes {
+  u32  len;
+  char *buf;
+} Bytes;
+
+#define MAX_BYTES 256
+
 // IO Magic Number
 #define IO_MAGIC_0 0xD0CA11
 #define IO_MAGIC_1 0xFF1FF1
@@ -91,7 +99,7 @@ Str gnet_readback_str(GNet* gnet, Port port) {
       case LIST_CONS: {
         if (ctr.args_len != 2) break;
         if (get_tag(ctr.args_buf[0]) != NUM) break;
-        if (str.text_len >= 256) { printf("ERROR: for now, HVM can only readback strings of length <256."); break; }
+        if (str.text_len > 256) { printf("ERROR: for now, HVM can only readback strings of length <256."); break; }
 
         str.text_buf[str.text_len++] = get_u24(get_val(ctr.args_buf[0]));
         gnet_boot_redex(gnet, new_pair(ctr.args_buf[1], ROOT));
@@ -107,8 +115,50 @@ Str gnet_readback_str(GNet* gnet, Port port) {
   return str;
 }
 
+// Converts a Port into a list of bytes.
+// Encoding:
+// - λt (t NIL)
+// - λt (((t CONS) head) tail)
+Bytes gnet_readback_bytes(GNet* gnet, Port port) {
+  // Result
+  Bytes bytes;
+  bytes.buf = (char*) malloc(sizeof(char) * MAX_BYTES);
+  bytes.len = 0;
+
+  // Readback loop
+  while (TRUE) {
+    // Normalizes the net
+    gnet_normalize(gnet);
+
+    // Reads the λ-Encoded Ctr
+    Ctr ctr = gnet_readback_ctr(gnet, gnet_peek(gnet, port));
+
+    // Reads string layer
+    switch (ctr.tag) {
+      case LIST_NIL: {
+        break;
+      }
+      case LIST_CONS: {
+        if (ctr.args_len != 2) break;
+        if (get_tag(ctr.args_buf[0]) != NUM) break;
+        if (bytes.len >= MAX_BYTES) { printf("ERROR: for now, HVM can only readback list of bytes of length <%u.", MAX_BYTES); break; }
+
+        bytes.buf[bytes.len++] = get_u24(get_val(ctr.args_buf[0]));
+        gnet_boot_redex(gnet, new_pair(ctr.args_buf[1], ROOT));
+        port = ROOT;
+        continue;
+      }
+    }
+    break;
+  }
+
+  bytes.buf[bytes.len] = '\0';
+
+  return bytes;
+}
+
 /// Returns a λ-Encoded Ctr for a NIL: λt (t NIL)
-/// Should only be called within `inject_str`, as a previous call
+/// Should only be called within `inject_bytes`, as a previous call
 /// to `get_resources` is expected.
 __device__ Port inject_nil(Net* net, TM* tm) {
   u32 v1 = tm->vloc[0];
@@ -126,7 +176,7 @@ __device__ Port inject_nil(Net* net, TM* tm) {
 }
 
 /// Returns a λ-Encoded Ctr for a CONS: λt (((t CONS) head) tail)
-/// Should only be called within `inject_str`, as a previous call
+/// Should only be called within `inject_bytes`, 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
@@ -150,19 +200,17 @@ __device__ Port inject_cons(Net* net, TM* tm, Port head, Port tail, u32 char_idx
   return new_port(CON, n4);
 }
 
-// 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.
+// Converts a list of bytes to a Port.
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-__device__ Port inject_str(Net* net, TM* tm, Str *str) {
+__device__ Port inject_bytes(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("inject_str: failed to get resources\n");
+    printf("inject_bytes: failed to get resources\n");
     return new_port(ERA, 0);
   }
 
@@ -176,29 +224,51 @@ __device__ Port inject_str(Net* net, TM* tm, Str *str) {
   return port;
 }
 
-__global__ void make_str_port(GNet* gnet, Str *str, Port* ret) {
+// Converts a list of bytes to a Port.
+// Encoding:
+// - λt (t NIL)
+// - λt (((t CONS) head) tail)
+__device__ Port inject_bytes(Net* net, TM* tm, Bytes *bytes) {
+  // Allocate all resources up front:
+  // - NIL needs  2 nodes & 1 var
+  // - CONS needs 4 nodes & 1 var
+  u32 len = bytes->len;
+  if (!get_resources(net, tm, 0, 2 + 4 * len, 1 + len)) {
+    printf("inject_bytes: failed to get resources\n");
+    return new_port(ERA, 0);
+  }
+
+  Port port = inject_nil(net, tm);
+
+  for (u32 i = 0; i < len; i++) {
+    Port byte = new_port(NUM, new_u24(bytes->buf[len - i - 1]));
+    port = inject_cons(net, tm, byte, port, i);
+  }
+
+  return port;
+}
+
+__global__ void make_bytes_port(GNet* gnet, Bytes *bytes, Port* ret) {
   if (GID() == 0) {
     TM tm;
     Net net = vnet_new(gnet, NULL, gnet->turn);
-    *ret = inject_str(&net, &tm, str);
+    *ret = inject_bytes(&net, &tm, bytes);
   }
 }
 
-// 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.
+// Converts a list of bytes to a Port.
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Port gnet_inject_str(GNet* gnet, Str *str) {
+Port gnet_inject_bytes(GNet* gnet, Bytes *bytes) {
   Port* d_ret;
   cudaMalloc(&d_ret, sizeof(Port));
 
-  Str* cu_str;
-  cudaMalloc(&cu_str, sizeof(Str));
-  cudaMemcpy(cu_str, str, sizeof(Str), cudaMemcpyHostToDevice);
+  Bytes* cu_bytes;
+  cudaMalloc(&cu_bytes, sizeof(Str));
+  cudaMemcpy(cu_bytes, bytes, sizeof(Str), cudaMemcpyHostToDevice);
 
-  make_str_port<<<1,1>>>(gnet, cu_str, d_ret);
+  make_bytes_port<<<1,1>>>(gnet, cu_bytes, d_ret);
 
   Port ret;
   cudaMemcpy(&ret, d_ret, sizeof(Port), cudaMemcpyDeviceToHost);
@@ -240,55 +310,46 @@ FILE* readback_file(Port port) {
   return fp;
 }
 
-// Reads a single char from `argm`.
-Port io_read_char(GNet* gnet, Port argm) {
-  FILE* fp = readback_file(gnet_peek(gnet, argm));
-  if (fp == NULL) {
+// Reads from a file a specified number of bytes.
+// `argm` is a tuple of (file_descriptor, num_bytes).
+Port io_read(GNet* gnet, Port argm) {
+  if (get_tag(gnet_peek(gnet, argm)) != CON) {
+    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(gnet_peek(gnet, argm)));
     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 gnet_inject_str(gnet, &str);
-}
+  Pair args = gnet_node_load(gnet, get_val(argm));
 
-// Reads from `argm` at most 255 characters or until a newline is seen.
-Port io_read_line(GNet* gnet, Port 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);
   }
 
-  /// Read a string.
-  Str str;
+  u32 num_bytes = get_u24(get_val(gnet_peek(gnet, get_snd(args))));
 
-  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);
+  /// Read a string.
+  Bytes bytes;
+  bytes.buf = (char*) malloc(sizeof(char) * num_bytes);
+  bytes.len = fread(bytes.buf, sizeof(char), num_bytes, fp);
 
-  // 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--;
+  if ((bytes.len != num_bytes) && ferror(fp)) {
+    fprintf(stderr, "io_read: failed to read\n");
   }
 
   // Convert it to a port.
-  return gnet_inject_str(gnet, &str);
+  Port ret = gnet_inject_bytes(gnet, &bytes);
+  free(bytes.buf);
+
+  return ret;
 }
 
 // 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, Port argm) {
+Port io_open(GNet* gnet, Port argm) {
   if (get_tag(gnet_peek(gnet, argm)) != CON) {
-    fprintf(stderr, "io_open_file: expected tuple\n");
+    fprintf(stderr, "io_open: expected tuple\n");
     return new_port(ERA, 0);
   }
 
@@ -296,6 +357,8 @@ Port io_open_file(GNet* gnet, Port argm) {
   Str name = gnet_readback_str(gnet, get_fst(args));
   Str mode = gnet_readback_str(gnet, get_snd(args));
 
+  printf("opening file '%s' with mode '%s'\n", name.text_buf, mode.text_buf);
+
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
       FILE_POINTERS[fd] = fopen(name.text_buf, mode.text_buf);
@@ -303,22 +366,22 @@ Port io_open_file(GNet* gnet, Port argm) {
     }
   }
 
-  fprintf(stderr, "io_open_file: too many open files\n");
+  fprintf(stderr, "io_open: too many open files\n");
 
   return new_port(ERA, 0);
 }
 
 // Closes a file, reclaiming the file descriptor.
-Port io_close_file(GNet* gnet, Port argm) {
+Port io_close(GNet* gnet, Port argm) {
   FILE* fp = readback_file(gnet_peek(gnet, argm));
   if (fp == NULL) {
-    fprintf(stderr, "io_close_file: failed to close\n");
+    fprintf(stderr, "io_close: 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);
+    fprintf(stderr, "io_close: failed to close: %i\n", err);
     return new_port(ERA, 0);
   }
 
@@ -327,9 +390,9 @@ Port io_close_file(GNet* gnet, Port argm) {
   return new_port(ERA, 0);
 }
 
-// Writes a string to a file.
+// Writes a list of bytes to a file.
 // `argm` is a tuple (CON node) of the
-// file descriptor and string to write.
+// file descriptor and list of bytes to write.
 Port io_write(GNet* gnet, Port argm) {
   if (get_tag(gnet_peek(gnet, argm)) != CON) {
     fprintf(stderr, "io_write: expected tuple, but got %u", get_tag(gnet_peek(gnet, argm)));
@@ -338,17 +401,69 @@ Port io_write(GNet* gnet, Port argm) {
 
   Pair args = gnet_node_load(gnet, get_val(argm));
   FILE* fp = readback_file(gnet_peek(gnet, get_fst(args)));
-  Str str = gnet_readback_str(gnet, get_snd(args));
+  Bytes bytes = gnet_readback_bytes(gnet, get_snd(args));
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
+    free(bytes.buf);
+
     return new_port(ERA, 0);
   }
 
-  if (fputs(str.text_buf, fp) == EOF) {
+  if (fwrite(bytes.buf, sizeof(char), bytes.len, fp) != bytes.len) {
     fprintf(stderr, "io_write: failed to write\n");
   }
 
+  free(bytes.buf);
+
+  return new_port(ERA, 0);
+}
+
+// Seeks to a position in a file.
+// `argm` is a 3-tuple (CON fd (CON offset whence)), where
+// - fd is a file descriptor
+// - offset is a signed byte offset
+// - whence is what that offset is relative to:
+//    - 0 (SEEK_SET): beginning of file
+//    - 1 (SEEK_CUR): current position of the file pointer
+//    - 2 (SEEK_END): end of the file
+Port io_seek(GNet* gnet, Port argm) {
+  if (get_tag(gnet_peek(gnet, argm)) != CON) {
+    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(gnet_peek(gnet, argm)));
+    return new_port(ERA, 0);
+  }
+
+  Pair args1 = gnet_node_load(gnet, get_val(argm));
+  if (get_tag(gnet_peek(gnet, get_snd(args1))) != CON) {
+    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(gnet_peek(gnet, get_snd(args1))));
+    return new_port(ERA, 0);
+  }
+
+  Pair args2 = gnet_node_load(gnet, get_val(gnet_peek(gnet, get_snd(args1))));
+
+  FILE* fp = readback_file(gnet_peek(gnet, get_fst(args1)));
+  if (fp == NULL) {
+    fprintf(stderr, "io_write: invalid file descriptor\n");
+    return new_port(ERA, 0);
+  }
+
+  i32 offset = get_i24(get_val(gnet_peek(gnet, get_fst(args2))));
+  u32 whence = get_i24(get_val(gnet_peek(gnet, get_snd(args2))));
+
+  int cwhence;
+  switch (whence) {
+    case 0: cwhence = SEEK_SET; break;
+    case 1: cwhence = SEEK_CUR; break;
+    case 2: cwhence = SEEK_END; break;
+    default:
+      fprintf(stderr, "io_seek: invalid whence\n");
+      return new_port(ERA, 0);
+  }
+
+  if (fseek(fp, offset, cwhence) != 0) {
+    fprintf(stderr, "io_seek: failed to seek\n");
+  }
+
   return new_port(ERA, 0);
 }
 
@@ -385,11 +500,11 @@ Port io_sleep(GNet* gnet, Port argm) {
 }
 
 void book_init(Book* book) {
-  book->ffns_buf[book->ffns_len++] = (FFn){"READ_CHAR", io_read_char};
-  book->ffns_buf[book->ffns_len++] = (FFn){"READ_LINE", io_read_line};
-  book->ffns_buf[book->ffns_len++] = (FFn){"OPEN_FILE", io_open_file};
-  book->ffns_buf[book->ffns_len++] = (FFn){"CLOSE_FILE", io_close_file};
+  book->ffns_buf[book->ffns_len++] = (FFn){"READ", io_read};
+  book->ffns_buf[book->ffns_len++] = (FFn){"OPEN", io_open};
+  book->ffns_buf[book->ffns_len++] = (FFn){"CLOSE", io_close};
   book->ffns_buf[book->ffns_len++] = (FFn){"WRITE", io_write};
+  book->ffns_buf[book->ffns_len++] = (FFn){"SEEK", io_seek};
   book->ffns_buf[book->ffns_len++] = (FFn){"GET_TIME", io_get_time};
   book->ffns_buf[book->ffns_len++] = (FFn){"SLEEP", io_sleep};
 

From b092fce82a67a5694ea7a2e43e9fb5f109c347d7 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Tue, 11 Jun 2024 07:41:44 +0200
Subject: [PATCH 03/11] fix cuda alloc bug

---
 src/run.c  | 27 ++++++++++++---------------
 src/run.cu | 54 ++++++++++++++----------------------------------------
 2 files changed, 26 insertions(+), 55 deletions(-)

diff --git a/src/run.c b/src/run.c
index 34d7b127..c745cd18 100644
--- a/src/run.c
+++ b/src/run.c
@@ -206,7 +206,7 @@ Port inject_bytes(Net* net, Bytes *bytes) {
   // - CONS needs 4 nodes & 1 var
   u32 len = bytes->len;
   if (!get_resources(net, tm[0], 0, 2 + 4 * len, 1 + len)) {
-    printf("inject_bytes: failed to get resources\n");
+    fprintf(stderr, "inject_bytes: failed to get resources\n");
     return new_port(ERA, 0);
   }
 
@@ -278,12 +278,13 @@ Port io_read(Net* net, Book* book, Port argm) {
 
   if ((bytes.len != num_bytes) && ferror(fp)) {
     fprintf(stderr, "io_read: failed to read\n");
+    free(bytes.buf);
+    return new_port(ERA, 0);
   }
 
   // Convert it to a port.
   Port ret = inject_bytes(net, &bytes);
   free(bytes.buf);
-
   return ret;
 }
 
@@ -300,8 +301,6 @@ Port io_open(Net* net, Book* book, Port argm) {
   Str name = readback_str(net, book, get_fst(args));
   Str mode = readback_str(net, book, get_snd(args));
 
-  printf("opening file '%s' with mode '%s'\n", name.text_buf, mode.text_buf);
-
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
       FILE_POINTERS[fd] = fopen(name.text_buf, mode.text_buf);
@@ -310,7 +309,6 @@ Port io_open(Net* net, Book* book, Port argm) {
   }
 
   fprintf(stderr, "io_open: too many open files\n");
-
   return new_port(ERA, 0);
 }
 
@@ -329,7 +327,6 @@ Port io_close(Net* net, Book* book, Port argm) {
   }
 
   FILE_POINTERS[get_u24(get_val(argm))] = NULL;
-
   return new_port(ERA, 0);
 }
 
@@ -349,16 +346,15 @@ Port io_write(Net* net, Book* book, Port argm) {
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
     free(bytes.buf);
-
     return new_port(ERA, 0);
   }
 
+
   if (fwrite(bytes.buf, sizeof(char), bytes.len, fp) != bytes.len) {
     fprintf(stderr, "io_write: failed to write\n");
   }
 
   free(bytes.buf);
-
   return new_port(ERA, 0);
 }
 
@@ -489,11 +485,6 @@ void do_run_io(Net* net, Book* book, Port port) {
     switch (ctr.tag) {
       case IO_CALL: {
         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;
-        u32  loc  = node_alloc_1(net, tm[0], &lps);
-        Port ret  = new_port(ERA, 0);
         FFn* ffn  = NULL;
         // FIXME: optimize this linear search
         for (u32 fid = 0; fid < book->ffns_len; ++fid) {
@@ -503,16 +494,22 @@ void do_run_io(Net* net, Book* book, Port port) {
           }
         }
         if (ffn == NULL) {
+          fprintf(stderr, "Unknown IO func '%s'\n", func.text_buf);
           break;
         }
-        ret = ffn->func(net, book, argm);
+
+        Port argm = ctr.args_buf[2];
+        Port cont = ctr.args_buf[3];
+        Port ret = ffn->func(net, book, argm);
+
+        u32 lps = 0;
+        u32 loc = node_alloc_1(net, tm[0], &lps);
         node_create(net, loc, new_pair(ret, ROOT));
         boot_redex(net, new_pair(new_port(CON, loc), cont));
         port = ROOT;
         continue;
       }
       case IO_DONE: {
-        printf("DONE\n");
         break;
       }
     }
diff --git a/src/run.cu b/src/run.cu
index 3b0f5796..5c8c9e8a 100644
--- a/src/run.cu
+++ b/src/run.cu
@@ -200,30 +200,6 @@ __device__ Port inject_cons(Net* net, TM* tm, Port head, Port tail, u32 char_idx
   return new_port(CON, n4);
 }
 
-// Converts a list of bytes to a Port.
-// Encoding:
-// - λt (t NIL)
-// - λt (((t CONS) head) tail)
-__device__ Port inject_bytes(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("inject_bytes: failed to get resources\n");
-    return new_port(ERA, 0);
-  }
-
-  Port port = inject_nil(net, tm);
-
-  for (u32 i = 0; i < len; i++) {
-    Port chr = new_port(NUM, new_u24(str->text_buf[len - i - 1]));
-    port = inject_cons(net, tm, chr, port, i);
-  }
-
-  return port;
-}
-
 // Converts a list of bytes to a Port.
 // Encoding:
 // - λt (t NIL)
@@ -248,11 +224,11 @@ __device__ Port inject_bytes(Net* net, TM* tm, Bytes *bytes) {
   return port;
 }
 
-__global__ void make_bytes_port(GNet* gnet, Bytes *bytes, Port* ret) {
+__global__ void make_bytes_port(GNet* gnet, Bytes bytes, Port* ret) {
   if (GID() == 0) {
     TM tm;
     Net net = vnet_new(gnet, NULL, gnet->turn);
-    *ret = inject_bytes(&net, &tm, bytes);
+    *ret = inject_bytes(&net, &tm, &bytes);
   }
 }
 
@@ -264,15 +240,18 @@ Port gnet_inject_bytes(GNet* gnet, Bytes *bytes) {
   Port* d_ret;
   cudaMalloc(&d_ret, sizeof(Port));
 
-  Bytes* cu_bytes;
-  cudaMalloc(&cu_bytes, sizeof(Str));
-  cudaMemcpy(cu_bytes, bytes, sizeof(Str), cudaMemcpyHostToDevice);
+  Bytes cu_bytes;
+  cu_bytes.len = bytes->len;
+
+  cudaMalloc(&cu_bytes.buf, sizeof(char) * cu_bytes.len);
+  cudaMemcpy(cu_bytes.buf, bytes->buf, sizeof(char) * cu_bytes.len, cudaMemcpyHostToDevice);
 
   make_bytes_port<<<1,1>>>(gnet, cu_bytes, d_ret);
 
   Port ret;
   cudaMemcpy(&ret, d_ret, sizeof(Port), cudaMemcpyDeviceToHost);
   cudaFree(d_ret);
+  cudaFree(cu_bytes.buf);
 
   return ret;
 }
@@ -291,7 +270,7 @@ static FILE* FILE_POINTERS[256];
 // Converts a NUM port (file descriptor) to file pointer.
 FILE* readback_file(Port port) {
   if (get_tag(port) != NUM) {
-    fprintf(stderr, "non-num where file descriptor was expected: %i\n", get_tag(port));
+    fprintf(stderr, "non-num where file descriptor was expected: %s\n", show_port(port).x);
     return NULL;
   }
 
@@ -320,9 +299,9 @@ Port io_read(GNet* gnet, Port argm) {
 
   Pair args = gnet_node_load(gnet, get_val(argm));
 
-  FILE* fp = readback_file(gnet_peek(gnet, argm));
+  FILE* fp = readback_file(gnet_peek(gnet, get_fst(args)));
   if (fp == NULL) {
-    fprintf(stderr, "io_read_line: invalid file descriptor\n");
+    fprintf(stderr, "io_read: invalid file descriptor\n");
     return new_port(ERA, 0);
   }
 
@@ -335,12 +314,13 @@ Port io_read(GNet* gnet, Port argm) {
 
   if ((bytes.len != num_bytes) && ferror(fp)) {
     fprintf(stderr, "io_read: failed to read\n");
+    free(bytes.buf);
+    return new_port(ERA, 0);
   }
 
   // Convert it to a port.
   Port ret = gnet_inject_bytes(gnet, &bytes);
   free(bytes.buf);
-
   return ret;
 }
 
@@ -357,8 +337,6 @@ Port io_open(GNet* gnet, Port argm) {
   Str name = gnet_readback_str(gnet, get_fst(args));
   Str mode = gnet_readback_str(gnet, get_snd(args));
 
-  printf("opening file '%s' with mode '%s'\n", name.text_buf, mode.text_buf);
-
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
       FILE_POINTERS[fd] = fopen(name.text_buf, mode.text_buf);
@@ -367,7 +345,6 @@ Port io_open(GNet* gnet, Port argm) {
   }
 
   fprintf(stderr, "io_open: too many open files\n");
-
   return new_port(ERA, 0);
 }
 
@@ -406,7 +383,6 @@ Port io_write(GNet* gnet, Port argm) {
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
     free(bytes.buf);
-
     return new_port(ERA, 0);
   }
 
@@ -415,7 +391,6 @@ Port io_write(GNet* gnet, Port argm) {
   }
 
   free(bytes.buf);
-
   return new_port(ERA, 0);
 }
 
@@ -550,7 +525,7 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
           }
         }
         if (ffn == NULL) {
-          printf("FOUND NOTHING when looking for %s\n", func.text_buf);
+          fprintf(stderr, "Unknown IO func '%s'\n", func.text_buf);
           break;
         }
 
@@ -564,7 +539,6 @@ void do_run_io(GNet* gnet, Book* book, Port port) {
         continue;
       }
       case IO_DONE: {
-        printf("DONE\n");
         break;
       }
     }

From e9aff7e3792e28b5af1beb9f5877f72c6186679e Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Tue, 11 Jun 2024 07:51:58 +0200
Subject: [PATCH 04/11] add seek to tests

---
 examples/demo_io/main.bend                    |  36 ++-
 examples/demo_io/main.hvm                     | 219 ++++++++++++------
 tests/programs/io/read_and_print.hvm          | 210 +++++++++++++----
 .../run__io_file@demo_io__main.hvm.snap       |   4 +-
 .../run__io_file@io__read_and_print.hvm.snap  |   2 +-
 5 files changed, 348 insertions(+), 123 deletions(-)

diff --git a/examples/demo_io/main.bend b/examples/demo_io/main.bend
index 5d982f50..e06c913d 100644
--- a/examples/demo_io/main.bend
+++ b/examples/demo_io/main.bend
@@ -1,4 +1,3 @@
-# run io on tests
 test-io = 1
 
 type IO_T:
@@ -18,12 +17,35 @@ def IO_T/bind(a, b):
 def call_io(func, argm):
   return IO_T/Call(IO_T/MAGIC, func, argm, lambda x: IO_T/Done(IO_T/MAGIC, x))
 
-def main:
+def read_input():
   with IO_T:
-    * <- call_io("WRITE", (1, "Hi! What's your name?\n"))
-    name <- call_io("READ_LINE", 0)
-    * <- call_io("WRITE", (1, "Your name is '"))
-    * <- call_io("WRITE", (1, name))
-    * <- call_io("WRITE", (1, "'!\n"))
+    * <- call_io("WRITE", (1, "What is your name?\n"))
+    return call_io("READ", (0, 10))
+
+def write_to_file():
+  with IO_T:
+    fp <- call_io("OPEN", ("testing.txt", "w"))
+    input <- read_input()
+    * <- call_io("WRITE", (fp, input))
+    * <- call_io("WRITE", (fp, "\n"))
+
+    return call_io("CLOSE", fp)
+
+def read_from_file():
+  with IO_T:
+    fp <- call_io("OPEN", ("testing.txt", "r"))
+    bytes <- call_io("READ", (fp, 5))
+    * <- call_io("WRITE", (1, bytes))
+    * <- call_io("SEEK", (fp, 2, 0))
+    bytes <- call_io("READ", (fp, 5))
+    * <- call_io("WRITE", (1, bytes))
+
+    return call_io("CLOSE", fp)
+
+def main():
+  with IO_T:
+    * <- write_to_file()
+    * <- read_from_file()
+    * <- call_io("WRITE", (1, "\n"))
 
     return 42
diff --git a/examples/demo_io/main.hvm b/examples/demo_io/main.hvm
index 53dcf595..bcdc54cf 100644
--- a/examples/demo_io/main.hvm
+++ b/examples/demo_io/main.hvm
@@ -33,43 +33,125 @@
 @call_io__C0 = a
   & @IO_T/Done ~ (@IO_T/MAGIC a)
 
-@main = a
-  & @IO_T/bind ~ (@main__C12 (@main__C11 a))
+@main = m
+  & @IO_T/bind ~ (@write_to_file ((((* k) l) l) m))
+  & @IO_T/bind ~ (@read_from_file ((((* i) j) j) k))
+  & @IO_T/bind ~ (g ((((* 42) h) h) i))
+  & @call_io ~ (e ((1 f) g))
+  & @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 ~ (10 (@String/Nil f))
 
-@main__C0 = (((* 42) a) a)
+@read_from_file = a
+  & @IO_T/bind ~ (@read_from_file__C11 (@read_from_file__C10 a))
 
-@main__C1 = i
-  & @call_io ~ (e ((1 h) i))
+@read_from_file__C0 = (f (* g))
+  & @call_io ~ (e (f g))
+  & @String/Cons ~ (67 (d e))
+  & @String/Cons ~ (76 (c d))
+  & @String/Cons ~ (79 (b c))
+  & @String/Cons ~ (83 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+
+@read_from_file__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))
-  & @String/Cons ~ (39 (g h))
-  & @String/Cons ~ (33 (f g))
-  & @String/Cons ~ (10 (@String/Nil f))
 
-@main__C10 = (* a)
-  & @IO_T/bind ~ (@main__C9 (@main__C8 a))
+@read_from_file__C10 = ((@read_from_file__C9 a) a)
+
+@read_from_file__C11 = q
+  & @call_io ~ (d ((o p) q))
+  & @String/Cons ~ (79 (c d))
+  & @String/Cons ~ (80 (b c))
+  & @String/Cons ~ (69 (a b))
+  & @String/Cons ~ (78 (@String/Nil a))
+  & @String/Cons ~ (116 (n o))
+  & @String/Cons ~ (101 (m n))
+  & @String/Cons ~ (115 (l m))
+  & @String/Cons ~ (116 (k l))
+  & @String/Cons ~ (105 (j k))
+  & @String/Cons ~ (110 (i j))
+  & @String/Cons ~ (103 (h i))
+  & @String/Cons ~ (46 (g h))
+  & @String/Cons ~ (116 (f g))
+  & @String/Cons ~ (120 (e f))
+  & @String/Cons ~ (116 (@String/Nil e))
+  & @String/Cons ~ (114 (@String/Nil p))
+
+@read_from_file__C2 = (c (a e))
+  & @IO_T/bind ~ (b (((@read_from_file__C0 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C1 ((1 a) b))
+
+@read_from_file__C3 = d
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (65 (a b))
+  & @String/Cons ~ (68 (@String/Nil a))
 
-@main__C11 = ((@main__C10 a) a)
+@read_from_file__C4 = ({a c} (* e))
+  & @IO_T/bind ~ (b (((@read_from_file__C2 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C3 ((a 5) b))
 
-@main__C12 = bb
-  & @call_io ~ (e ((1 ab) bb))
+@read_from_file__C5 = d
+  & @String/Cons ~ (83 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (69 (a b))
+  & @String/Cons ~ (75 (@String/Nil a))
+
+@read_from_file__C6 = ({a c} (* e))
+  & @IO_T/bind ~ (b (((@read_from_file__C4 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C5 ((a (2 0)) b))
+
+@read_from_file__C7 = 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))
-  & @String/Cons ~ (72 (z ab))
-  & @String/Cons ~ (105 (y z))
-  & @String/Cons ~ (33 (x y))
-  & @String/Cons ~ (32 (w x))
-  & @String/Cons ~ (87 (v w))
-  & @String/Cons ~ (104 (u v))
-  & @String/Cons ~ (97 (t u))
-  & @String/Cons ~ (116 (s t))
-  & @String/Cons ~ (39 (r s))
+
+@read_from_file__C8 = (c (a e))
+  & @IO_T/bind ~ (b (((@read_from_file__C6 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C7 ((1 a) b))
+
+@read_from_file__C9 = ({e g} i)
+  & @IO_T/bind ~ (f (((@read_from_file__C8 (g h)) h) i))
+  & @call_io ~ (d ((e 5) f))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (65 (a b))
+  & @String/Cons ~ (68 (@String/Nil a))
+
+@read_input = a
+  & @IO_T/bind ~ (@read_input__C2 (@read_input__C1 a))
+
+@read_input__C0 = (* e)
+  & @call_io ~ (d ((0 10) e))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (65 (a b))
+  & @String/Cons ~ (68 (@String/Nil a))
+
+@read_input__C1 = ((@read_input__C0 a) a)
+
+@read_input__C2 = 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))
@@ -84,63 +166,68 @@
   & @String/Cons ~ (63 (f g))
   & @String/Cons ~ (10 (@String/Nil f))
 
-@main__C2 = (* a)
-  & @IO_T/bind ~ (@main__C1 (@main__C0 a))
+@test-io = 1
+
+@undefer = (((a a) b) b)
+
+@write_to_file = a
+  & @IO_T/bind ~ (@write_to_file__C8 (@write_to_file__C7 a))
+
+@write_to_file__C0 = (f (* g))
+  & @call_io ~ (e (f g))
+  & @String/Cons ~ (67 (d e))
+  & @String/Cons ~ (76 (c d))
+  & @String/Cons ~ (79 (b c))
+  & @String/Cons ~ (83 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+
+@write_to_file__C1 = a
+  & @String/Cons ~ (10 (@String/Nil a))
 
-@main__C3 = e
+@write_to_file__C2 = 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__C4 = ((@main__C2 a) a)
+@write_to_file__C3 = ({a c} (* e))
+  & @IO_T/bind ~ (b (((@write_to_file__C0 (c d)) d) e))
+  & @call_io ~ (@write_to_file__C2 ((a @write_to_file__C1) b))
 
-@main__C5 = (a (* c))
-  & @IO_T/bind ~ (b (@main__C4 c))
-  & @call_io ~ (@main__C3 ((1 a) b))
-
-@main__C6 = t
-  & @call_io ~ (e ((1 s) t))
+@write_to_file__C4 = 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))
-  & @String/Cons ~ (89 (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 ~ (32 (i j))
-  & @String/Cons ~ (105 (h i))
-  & @String/Cons ~ (115 (g h))
-  & @String/Cons ~ (32 (f g))
-  & @String/Cons ~ (39 (@String/Nil f))
-
-@main__C7 = (a c)
-  & @IO_T/bind ~ (@main__C6 (((@main__C5 (a b)) b) c))
-
-@main__C8 = ((@main__C7 a) a)
-
-@main__C9 = 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))
 
-@test-io = 1
-
-@undefer = (((a a) b) b)
+@write_to_file__C5 = ({a d} (b f))
+  & @IO_T/bind ~ (c (((@write_to_file__C3 (d e)) e) f))
+  & @call_io ~ (@write_to_file__C4 ((a b) c))
+
+@write_to_file__C6 = (a c)
+  & @IO_T/bind ~ (@read_input (((@write_to_file__C5 (a b)) b) c))
+
+@write_to_file__C7 = ((@write_to_file__C6 a) a)
+
+@write_to_file__C8 = q
+  & @call_io ~ (d ((o p) q))
+  & @String/Cons ~ (79 (c d))
+  & @String/Cons ~ (80 (b c))
+  & @String/Cons ~ (69 (a b))
+  & @String/Cons ~ (78 (@String/Nil a))
+  & @String/Cons ~ (116 (n o))
+  & @String/Cons ~ (101 (m n))
+  & @String/Cons ~ (115 (l m))
+  & @String/Cons ~ (116 (k l))
+  & @String/Cons ~ (105 (j k))
+  & @String/Cons ~ (110 (i j))
+  & @String/Cons ~ (103 (h i))
+  & @String/Cons ~ (46 (g h))
+  & @String/Cons ~ (116 (f g))
+  & @String/Cons ~ (120 (e f))
+  & @String/Cons ~ (116 (@String/Nil e))
+  & @String/Cons ~ (119 (@String/Nil p))
 
 
diff --git a/tests/programs/io/read_and_print.hvm b/tests/programs/io/read_and_print.hvm
index d60a16b5..bcdc54cf 100644
--- a/tests/programs/io/read_and_print.hvm
+++ b/tests/programs/io/read_and_print.hvm
@@ -1,12 +1,17 @@
-@IO_T/Call = (a (b (c (d ((1 (a (b (c (d e))))) e)))))
+@IO_T/Call = (a (b (c (d ((@IO_T/Call/tag (a (b (c (d e))))) e)))))
 
-@IO_T/Done = (a (b ((0 (a (b c))) c)))
+@IO_T/Call/tag = 1
+
+@IO_T/Done = (a (b ((@IO_T/Done/tag (a (b c))) c)))
+
+@IO_T/Done/tag = 0
 
 @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__C0 = (* (b (a c)))
+  & @undefer ~ (a (b c))
 
 @IO_T/bind__C1 = (* (* (a (b ((c d) (e g))))))
   & @IO_T/Call ~ (@IO_T/MAGIC (a (b ((c f) g))))
@@ -14,12 +19,13 @@
 
 @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 ((@String/Cons/tag (a (b c))) c)))
 
-@String/Cons = (a (b ((1 (a (b c))) c)))
+@String/Cons/tag = 1
 
-@String/Nil = ((0 a) a)
+@String/Nil = ((@String/Nil/tag a) a)
+
+@String/Nil/tag = 0
 
 @call_io = (a (b c))
   & @IO_T/Call ~ (@IO_T/MAGIC (a (b (@call_io__C0 c))))
@@ -27,67 +33,113 @@
 @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))
+@main = m
+  & @IO_T/bind ~ (@write_to_file ((((* k) l) l) m))
+  & @IO_T/bind ~ (@read_from_file ((((* i) j) j) k))
+  & @IO_T/bind ~ (g ((((* 42) h) h) i))
+  & @call_io ~ (e ((1 f) g))
   & @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
+@read_from_file = a
+  & @IO_T/bind ~ (@read_from_file__C11 (@read_from_file__C10 a))
+
+@read_from_file__C0 = (f (* g))
+  & @call_io ~ (e (f g))
+  & @String/Cons ~ (67 (d e))
+  & @String/Cons ~ (76 (c d))
+  & @String/Cons ~ (79 (b c))
+  & @String/Cons ~ (83 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+
+@read_from_file__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))
+@read_from_file__C10 = ((@read_from_file__C9 a) a)
+
+@read_from_file__C11 = q
+  & @call_io ~ (d ((o p) q))
+  & @String/Cons ~ (79 (c d))
+  & @String/Cons ~ (80 (b c))
+  & @String/Cons ~ (69 (a b))
+  & @String/Cons ~ (78 (@String/Nil a))
+  & @String/Cons ~ (116 (n o))
+  & @String/Cons ~ (101 (m n))
+  & @String/Cons ~ (115 (l m))
+  & @String/Cons ~ (116 (k l))
+  & @String/Cons ~ (105 (j k))
+  & @String/Cons ~ (110 (i j))
+  & @String/Cons ~ (103 (h i))
+  & @String/Cons ~ (46 (g h))
+  & @String/Cons ~ (116 (f g))
+  & @String/Cons ~ (120 (e f))
+  & @String/Cons ~ (116 (@String/Nil e))
+  & @String/Cons ~ (114 (@String/Nil p))
+
+@read_from_file__C2 = (c (a e))
+  & @IO_T/bind ~ (b (((@read_from_file__C0 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C1 ((1 a) b))
+
+@read_from_file__C3 = d
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (65 (a b))
+  & @String/Cons ~ (68 (@String/Nil a))
+
+@read_from_file__C4 = ({a c} (* e))
+  & @IO_T/bind ~ (b (((@read_from_file__C2 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C3 ((a 5) b))
 
-@main__C3 = n
-  & @call_io ~ (e ((1 m) n))
+@read_from_file__C5 = d
+  & @String/Cons ~ (83 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (69 (a b))
+  & @String/Cons ~ (75 (@String/Nil a))
+
+@read_from_file__C6 = ({a c} (* e))
+  & @IO_T/bind ~ (b (((@read_from_file__C4 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C5 ((a (2 0)) b))
+
+@read_from_file__C7 = 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))
-  & @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))
+@read_from_file__C8 = (c (a e))
+  & @IO_T/bind ~ (b (((@read_from_file__C6 (c d)) d) e))
+  & @call_io ~ (@read_from_file__C7 ((1 a) b))
+
+@read_from_file__C9 = ({e g} i)
+  & @IO_T/bind ~ (f (((@read_from_file__C8 (g h)) h) i))
+  & @call_io ~ (d ((e 5) f))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (65 (a b))
+  & @String/Cons ~ (68 (@String/Nil a))
+
+@read_input = a
+  & @IO_T/bind ~ (@read_input__C2 (@read_input__C1 a))
+
+@read_input__C0 = (* e)
+  & @call_io ~ (d ((0 10) e))
+  & @String/Cons ~ (82 (c d))
+  & @String/Cons ~ (69 (b c))
+  & @String/Cons ~ (65 (a b))
+  & @String/Cons ~ (68 (@String/Nil a))
+
+@read_input__C1 = ((@read_input__C0 a) a)
 
-@main__C7 = y
+@read_input__C2 = y
   & @call_io ~ (e ((1 x) y))
   & @String/Cons ~ (87 (d e))
   & @String/Cons ~ (82 (c d))
@@ -115,3 +167,67 @@
   & @String/Cons ~ (10 (@String/Nil f))
 
 @test-io = 1
+
+@undefer = (((a a) b) b)
+
+@write_to_file = a
+  & @IO_T/bind ~ (@write_to_file__C8 (@write_to_file__C7 a))
+
+@write_to_file__C0 = (f (* g))
+  & @call_io ~ (e (f g))
+  & @String/Cons ~ (67 (d e))
+  & @String/Cons ~ (76 (c d))
+  & @String/Cons ~ (79 (b c))
+  & @String/Cons ~ (83 (a b))
+  & @String/Cons ~ (69 (@String/Nil a))
+
+@write_to_file__C1 = a
+  & @String/Cons ~ (10 (@String/Nil a))
+
+@write_to_file__C2 = 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))
+
+@write_to_file__C3 = ({a c} (* e))
+  & @IO_T/bind ~ (b (((@write_to_file__C0 (c d)) d) e))
+  & @call_io ~ (@write_to_file__C2 ((a @write_to_file__C1) b))
+
+@write_to_file__C4 = 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))
+
+@write_to_file__C5 = ({a d} (b f))
+  & @IO_T/bind ~ (c (((@write_to_file__C3 (d e)) e) f))
+  & @call_io ~ (@write_to_file__C4 ((a b) c))
+
+@write_to_file__C6 = (a c)
+  & @IO_T/bind ~ (@read_input (((@write_to_file__C5 (a b)) b) c))
+
+@write_to_file__C7 = ((@write_to_file__C6 a) a)
+
+@write_to_file__C8 = q
+  & @call_io ~ (d ((o p) q))
+  & @String/Cons ~ (79 (c d))
+  & @String/Cons ~ (80 (b c))
+  & @String/Cons ~ (69 (a b))
+  & @String/Cons ~ (78 (@String/Nil a))
+  & @String/Cons ~ (116 (n o))
+  & @String/Cons ~ (101 (m n))
+  & @String/Cons ~ (115 (l m))
+  & @String/Cons ~ (116 (k l))
+  & @String/Cons ~ (105 (j k))
+  & @String/Cons ~ (110 (i j))
+  & @String/Cons ~ (103 (h i))
+  & @String/Cons ~ (46 (g h))
+  & @String/Cons ~ (116 (f g))
+  & @String/Cons ~ (120 (e f))
+  & @String/Cons ~ (116 (@String/Nil e))
+  & @String/Cons ~ (119 (@String/Nil p))
+
+
diff --git a/tests/snapshots/run__io_file@demo_io__main.hvm.snap b/tests/snapshots/run__io_file@demo_io__main.hvm.snap
index f7f42b28..8503fe52 100644
--- a/tests/snapshots/run__io_file@demo_io__main.hvm.snap
+++ b/tests/snapshots/run__io_file@demo_io__main.hvm.snap
@@ -3,6 +3,6 @@ source: tests/run.rs
 expression: c_output
 input_file: examples/demo_io/main.hvm
 ---
-Hi! What's your name?
-Your name is 'io from the tests'!
+What is your name?
+io fr from
 Result: 42
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
index e00390ac..8b7aabc9 100644
--- a/tests/snapshots/run__io_file@io__read_and_print.hvm.snap
+++ b/tests/snapshots/run__io_file@io__read_and_print.hvm.snap
@@ -4,5 +4,5 @@ expression: c_output
 input_file: tests/programs/io/read_and_print.hvm
 ---
 What is your name?
-Hello, 'io from the tests'!
+io fr from
 Result: 42

From c8df99d7401aa1dd87525d8090198af2f44d23e3 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Tue, 11 Jun 2024 07:54:57 +0200
Subject: [PATCH 05/11] bench fix attempt

---
 .github/workflows/bench.yml | 1 -
 1 file changed, 1 deletion(-)

diff --git a/.github/workflows/bench.yml b/.github/workflows/bench.yml
index c0933053..316a678d 100644
--- a/.github/workflows/bench.yml
+++ b/.github/workflows/bench.yml
@@ -13,7 +13,6 @@ jobs:
     timeout-minutes: 10
     steps:
       - uses: actions/checkout@v3
-      - uses: dsherret/rust-toolchain-file@v1
       - name: compare perf
         run: |
           git fetch origin main

From 3652957f48d3f6374266200a7b23b79ea15fda95 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Tue, 11 Jun 2024 19:16:16 +0200
Subject: [PATCH 06/11] fix io_sleep hanging on var nodes

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

diff --git a/src/run.c b/src/run.c
index c745cd18..601bc1bd 100644
--- a/src/run.c
+++ b/src/run.c
@@ -429,8 +429,8 @@ 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
-  u32 dur_hi = get_u24(get_val(get_fst(dur_node)));
-  u32 dur_lo = get_u24(get_val(get_snd(dur_node)));
+  u32 dur_hi = get_u24(get_val(peek(net, get_fst(dur_node))));
+  u32 dur_lo = get_u24(get_val(peek(net, get_snd(dur_node))));
   // Combine into a 48-bit duration in nanoseconds
   u64 dur_ns = (((u64)dur_hi) << 24) | dur_lo;
   // Sleep for the specified duration

From 1401b24a83f5a979e5d6009d53b2d5769c0a4584 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 14 Jun 2024 17:40:35 +0200
Subject: [PATCH 07/11] add expand to io args

---
 examples/demo_io/main.bend | 14 +++++++++++-
 examples/demo_io/main.hvm  | 30 ++++++++++++++++---------
 src/run.c                  | 46 +++++++++++++++++++-------------------
 3 files changed, 55 insertions(+), 35 deletions(-)

diff --git a/examples/demo_io/main.bend b/examples/demo_io/main.bend
index e06c913d..b9c22fab 100644
--- a/examples/demo_io/main.bend
+++ b/examples/demo_io/main.bend
@@ -42,10 +42,22 @@ def read_from_file():
 
     return call_io("CLOSE", fp)
 
+def write:
+  return "WRITE"
+
+def one:
+  return 1
+
+def newline:
+  return "\n"
+
+def one_newline_pair:
+  return (one, newline)
+
 def main():
   with IO_T:
     * <- write_to_file()
     * <- read_from_file()
-    * <- call_io("WRITE", (1, "\n"))
+    * <- call_io(write, one_newline_pair)
 
     return 42
diff --git a/examples/demo_io/main.hvm b/examples/demo_io/main.hvm
index bcdc54cf..f5f3da3f 100644
--- a/examples/demo_io/main.hvm
+++ b/examples/demo_io/main.hvm
@@ -33,17 +33,18 @@
 @call_io__C0 = a
   & @IO_T/Done ~ (@IO_T/MAGIC a)
 
-@main = m
-  & @IO_T/bind ~ (@write_to_file ((((* k) l) l) m))
-  & @IO_T/bind ~ (@read_from_file ((((* i) j) j) k))
-  & @IO_T/bind ~ (g ((((* 42) h) h) i))
-  & @call_io ~ (e ((1 f) g))
-  & @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 ~ (10 (@String/Nil f))
+@main = g
+  & @IO_T/bind ~ (@write_to_file ((((* e) f) f) g))
+  & @IO_T/bind ~ (@read_from_file ((((* c) d) d) e))
+  & @IO_T/bind ~ (a ((((* 42) b) b) c))
+  & @call_io ~ (@write (@one_newline_pair a))
+
+@newline = a
+  & @String/Cons ~ (10 (@String/Nil a))
+
+@one = 1
+
+@one_newline_pair = (@one @newline)
 
 @read_from_file = a
   & @IO_T/bind ~ (@read_from_file__C11 (@read_from_file__C10 a))
@@ -170,6 +171,13 @@
 
 @undefer = (((a a) b) b)
 
+@write = 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))
+
 @write_to_file = a
   & @IO_T/bind ~ (@write_to_file__C8 (@write_to_file__C7 a))
 
diff --git a/src/run.c b/src/run.c
index 601bc1bd..5fea52a7 100644
--- a/src/run.c
+++ b/src/run.c
@@ -256,21 +256,21 @@ FILE* readback_file(Port port) {
 // Reads from a file a specified number of bytes.
 // `argm` is a tuple of (file_descriptor, num_bytes).
 Port io_read(Net* net, Book* book, Port argm) {
-  if (get_tag(peek(net, argm)) != CON) {
-    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(peek(net, argm)));
+  if (get_tag(argm) != CON) {
+    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(argm));
     return new_port(ERA, 0);
   }
 
   Pair args = node_load(net, get_val(argm));
+  Port file = expand(net, book, get_fst(args));
+  u32 num_bytes = get_u24(get_val(expand(net, book, get_snd(args))));
 
-  FILE* fp = readback_file(peek(net, get_fst(args)));
+  FILE* fp = readback_file(file);
   if (fp == NULL) {
     fprintf(stderr, "io_read: invalid file descriptor\n");
     return new_port(ERA, 0);
   }
 
-  u32 num_bytes = get_u24(get_val(peek(net, get_snd(args))));
-
   /// Read a string.
   Bytes bytes;
   bytes.buf = (char*) malloc(sizeof(char) * num_bytes);
@@ -292,14 +292,14 @@ Port io_read(Net* net, Book* book, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file name and mode as strings.
 Port io_open(Net* net, Book* book, Port argm) {
-  if (get_tag(peek(net, argm)) != CON) {
-    fprintf(stderr, "io_open: expected tuple\n");
+  if (get_tag(argm) != CON) {
+    fprintf(stderr, "io_open: expected tuple, but got %u\n", get_tag(argm));
     return new_port(ERA, 0);
   }
 
   Pair args = node_load(net, get_val(argm));
-  Str name = readback_str(net, book, get_fst(args));
-  Str mode = readback_str(net, book, get_snd(args));
+  Str name = readback_str(net, book, expand(net, book, get_fst(args)));
+  Str mode = readback_str(net, book, expand(net, book, get_snd(args)));
 
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
@@ -314,7 +314,7 @@ Port io_open(Net* net, Book* book, Port argm) {
 
 // Closes a file, reclaiming the file descriptor.
 Port io_close(Net* net, Book* book, Port argm) {
-  FILE* fp = readback_file(peek(net, argm));
+  FILE* fp = readback_file(argm);
   if (fp == NULL) {
     fprintf(stderr, "io_close: failed to close\n");
     return new_port(ERA, 0);
@@ -334,13 +334,13 @@ Port io_close(Net* net, Book* book, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file descriptor and list of bytes 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\n", get_tag(peek(net, argm)));
+  if (get_tag(argm) != CON) {
+    fprintf(stderr, "io_write: expected tuple, but got %u\n", get_tag(argm));
     return new_port(ERA, 0);
   }
 
   Pair args = node_load(net, get_val(argm));
-  FILE* fp = readback_file(peek(net, get_fst(args)));
+  FILE* fp = readback_file(expand(net, book, get_fst(args)));
   Bytes bytes = readback_bytes(net, book, get_snd(args));
 
   if (fp == NULL) {
@@ -367,28 +367,28 @@ Port io_write(Net* net, Book* book, Port argm) {
 //    - 1 (SEEK_CUR): current position of the file pointer
 //    - 2 (SEEK_END): end of the file
 Port io_seek(Net* net, Book* book, Port argm) {
-  if (get_tag(peek(net, argm)) != CON) {
-    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(peek(net, argm)));
+  if (get_tag(argm) != CON) {
+    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(argm));
     return new_port(ERA, 0);
   }
 
   Pair args1 = node_load(net, get_val(argm));
-  if (get_tag(peek(net, get_snd(args1))) != CON) {
+  Port argm2 = expand(net, book, get_snd(args1));
+  if (get_tag(argm2) != CON) {
     fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(peek(net, get_snd(args1))));
     return new_port(ERA, 0);
   }
+  Pair args2 = node_load(net, get_val(argm2));
 
-  Pair args2 = node_load(net, get_val(peek(net, get_snd(args1))));
+  FILE* fp = readback_file(expand(net, book, get_fst(args1)));
+  i32 offset = get_i24(get_val(expand(net, book, get_fst(args2))));
+  u32 whence = get_i24(get_val(expand(net, book, get_snd(args2))));
 
-  FILE* fp = readback_file(peek(net, get_fst(args1)));
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
     return new_port(ERA, 0);
   }
 
-  i32 offset = get_i24(get_val(peek(net, get_fst(args2))));
-  u32 whence = get_i24(get_val(peek(net, get_snd(args2))));
-
   int cwhence;
   switch (whence) {
     case 0: cwhence = SEEK_SET; break;
@@ -429,8 +429,8 @@ 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
-  u32 dur_hi = get_u24(get_val(peek(net, get_fst(dur_node))));
-  u32 dur_lo = get_u24(get_val(peek(net, get_snd(dur_node))));
+  u32 dur_hi = get_u24(get_val(expand(net, book, get_fst(dur_node))));
+  u32 dur_lo = get_u24(get_val(expand(net, book, get_snd(dur_node))));
   // Combine into a 48-bit duration in nanoseconds
   u64 dur_ns = (((u64)dur_hi) << 24) | dur_lo;
   // Sleep for the specified duration

From 1eaf5fe7fc36357b94e5e9ee4c5fdfb2daac7125 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Fri, 14 Jun 2024 17:50:44 +0200
Subject: [PATCH 08/11] cuda impl

---
 src/run.c  |  3 +--
 src/run.cu | 47 +++++++++++++++++++++++------------------------
 2 files changed, 24 insertions(+), 26 deletions(-)

diff --git a/src/run.c b/src/run.c
index 5fea52a7..3329a90a 100644
--- a/src/run.c
+++ b/src/run.c
@@ -349,7 +349,6 @@ Port io_write(Net* net, Book* book, Port argm) {
     return new_port(ERA, 0);
   }
 
-
   if (fwrite(bytes.buf, sizeof(char), bytes.len, fp) != bytes.len) {
     fprintf(stderr, "io_write: failed to write\n");
   }
@@ -375,7 +374,7 @@ Port io_seek(Net* net, Book* book, Port argm) {
   Pair args1 = node_load(net, get_val(argm));
   Port argm2 = expand(net, book, get_snd(args1));
   if (get_tag(argm2) != CON) {
-    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(peek(net, get_snd(args1))));
+    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(argm2));
     return new_port(ERA, 0);
   }
   Pair args2 = node_load(net, get_val(argm2));
diff --git a/src/run.cu b/src/run.cu
index 5c8c9e8a..9a814035 100644
--- a/src/run.cu
+++ b/src/run.cu
@@ -292,21 +292,21 @@ FILE* readback_file(Port port) {
 // Reads from a file a specified number of bytes.
 // `argm` is a tuple of (file_descriptor, num_bytes).
 Port io_read(GNet* gnet, Port argm) {
-  if (get_tag(gnet_peek(gnet, argm)) != CON) {
-    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(gnet_peek(gnet, argm)));
+  if (get_tag(argm) != CON) {
+    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(argm));
     return new_port(ERA, 0);
   }
 
   Pair args = gnet_node_load(gnet, get_val(argm));
+  Port file = gnet_expand(gnet, get_fst(args));
+  u32 num_bytes = get_u24(get_val(gnet_expand(gnet, get_snd(args))));
 
-  FILE* fp = readback_file(gnet_peek(gnet, get_fst(args)));
+  FILE* fp = readback_file(file);
   if (fp == NULL) {
     fprintf(stderr, "io_read: invalid file descriptor\n");
     return new_port(ERA, 0);
   }
 
-  u32 num_bytes = get_u24(get_val(gnet_peek(gnet, get_snd(args))));
-
   /// Read a string.
   Bytes bytes;
   bytes.buf = (char*) malloc(sizeof(char) * num_bytes);
@@ -328,14 +328,14 @@ Port io_read(GNet* gnet, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file name and mode as strings.
 Port io_open(GNet* gnet, Port argm) {
-  if (get_tag(gnet_peek(gnet, argm)) != CON) {
+  if (get_tag(argm) != CON) {
     fprintf(stderr, "io_open: expected tuple\n");
     return new_port(ERA, 0);
   }
 
   Pair args = gnet_node_load(gnet, get_val(argm));
-  Str name = gnet_readback_str(gnet, get_fst(args));
-  Str mode = gnet_readback_str(gnet, get_snd(args));
+  Str name = gnet_readback_str(gnet, gnet_expand(gnet, get_fst(args)));
+  Str mode = gnet_readback_str(gnet, gnet_expand(gnet, get_snd(args)));
 
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
@@ -350,7 +350,7 @@ Port io_open(GNet* gnet, Port argm) {
 
 // Closes a file, reclaiming the file descriptor.
 Port io_close(GNet* gnet, Port argm) {
-  FILE* fp = readback_file(gnet_peek(gnet, argm));
+  FILE* fp = readback_file(argm);
   if (fp == NULL) {
     fprintf(stderr, "io_close: failed to close\n");
     return new_port(ERA, 0);
@@ -363,7 +363,6 @@ Port io_close(GNet* gnet, Port argm) {
   }
 
   FILE_POINTERS[get_u24(get_val(argm))] = NULL;
-
   return new_port(ERA, 0);
 }
 
@@ -371,13 +370,13 @@ Port io_close(GNet* gnet, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file descriptor and list of bytes to write.
 Port io_write(GNet* gnet, Port argm) {
-  if (get_tag(gnet_peek(gnet, argm)) != CON) {
-    fprintf(stderr, "io_write: expected tuple, but got %u", get_tag(gnet_peek(gnet, argm)));
+  if (get_tag(argm) != CON) {
+    fprintf(stderr, "io_write: expected tuple, but got %u\n", get_tag(argm));
     return new_port(ERA, 0);
   }
 
   Pair args = gnet_node_load(gnet, get_val(argm));
-  FILE* fp = readback_file(gnet_peek(gnet, get_fst(args)));
+  FILE* fp = readback_file(gnet_expand(gnet, get_fst(args)));
   Bytes bytes = gnet_readback_bytes(gnet, get_snd(args));
 
   if (fp == NULL) {
@@ -403,28 +402,28 @@ Port io_write(GNet* gnet, Port argm) {
 //    - 1 (SEEK_CUR): current position of the file pointer
 //    - 2 (SEEK_END): end of the file
 Port io_seek(GNet* gnet, Port argm) {
-  if (get_tag(gnet_peek(gnet, argm)) != CON) {
-    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(gnet_peek(gnet, argm)));
+  if (get_tag(argm) != CON) {
+    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(argm));
     return new_port(ERA, 0);
   }
 
   Pair args1 = gnet_node_load(gnet, get_val(argm));
-  if (get_tag(gnet_peek(gnet, get_snd(args1))) != CON) {
-    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(gnet_peek(gnet, get_snd(args1))));
+  Port argm2 = gnet_expand(gnet, get_snd(args1));
+  if (get_tag(argm2) != CON) {
+    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(argm2));
     return new_port(ERA, 0);
   }
+  Pair args2 = gnet_node_load(gnet, get_val(argm2));
 
-  Pair args2 = gnet_node_load(gnet, get_val(gnet_peek(gnet, get_snd(args1))));
+  FILE* fp = readback_file(gnet_expand(gnet, get_fst(args1)));
+  i32 offset = get_i24(get_val(gnet_expand(gnet, get_fst(args2))));
+  u32 whence = get_i24(get_val(gnet_expand(gnet, get_snd(args2))));
 
-  FILE* fp = readback_file(gnet_peek(gnet, get_fst(args1)));
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
     return new_port(ERA, 0);
   }
 
-  i32 offset = get_i24(get_val(gnet_peek(gnet, get_fst(args2))));
-  u32 whence = get_i24(get_val(gnet_peek(gnet, get_snd(args2))));
-
   int cwhence;
   switch (whence) {
     case 0: cwhence = SEEK_SET; break;
@@ -461,8 +460,8 @@ 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
-  u32 dur_hi = get_u24(get_val(get_fst(dur_node)));
-  u32 dur_lo = get_u24(get_val(get_snd(dur_node)));
+  u32 dur_hi = get_u24(get_val(gnet_expand(gnet, get_fst(dur_node))));
+  u32 dur_lo = get_u24(get_val(gnet_expand(gnet, get_snd(dur_node))));
   // Combine into a 48-bit duration in nanoseconds
   u64 dur_ns = (((u64)dur_hi) << 24) | dur_lo;
   // Sleep for the specified duration

From 4ee1bf4e50436fb188c1cc09de2a66713c926949 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Mon, 17 Jun 2024 11:07:06 +0200
Subject: [PATCH 09/11] readback_tup run.c

---
 src/run.c | 86 +++++++++++++++++++++++++++++++++++--------------------
 1 file changed, 55 insertions(+), 31 deletions(-)

diff --git a/src/run.c b/src/run.c
index 3329a90a..3dd8a129 100644
--- a/src/run.c
+++ b/src/run.c
@@ -7,6 +7,12 @@ typedef struct Ctr {
   Port args_buf[16];
 } Ctr;
 
+// Readback: Tuples
+typedef struct Tup {
+  u32  elem_len;
+  Port elem_buf[8];
+} Tup;
+
 // Readback: λ-Encoded Str (UTF-32)
 // FIXME: this is actually ASCII :|
 // FIXME: remove len limit
@@ -72,6 +78,26 @@ Ctr readback_ctr(Net* net, Book* book, Port port) {
   return ctr;
 }
 
+// Reads back a tuple of at most `size` elements. Tuples are
+// (right-nested con nodes) (CON 1 (CON 2 (CON 3 (...))))
+// The provided `port` should be `expanded` before calling.
+Tup readback_tup(Net* net, Book* book, Port port, u32 size) {
+  Tup tup;
+  tup.elem_len = 0;
+
+  // Loads remaining arguments
+  while (get_tag(port) == CON && (tup.elem_len + 1 < size)) {
+    Pair node = node_load(net, get_val(port));
+    tup.elem_buf[tup.elem_len++] = expand(net, book, get_fst(node));
+
+    port = expand(net, book, get_snd(node));
+  }
+
+  tup.elem_buf[tup.elem_len++] = port;
+
+  return tup;
+}
+
 // Converts a Port into a UTF-32 (truncated to 24 bits) string.
 // Since unicode scalars can fit in 21 bits, HVM's u24
 // integers can contain any unicode scalar value.
@@ -256,16 +282,15 @@ FILE* readback_file(Port port) {
 // Reads from a file a specified number of bytes.
 // `argm` is a tuple of (file_descriptor, num_bytes).
 Port io_read(Net* net, Book* book, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(argm));
+  Tup tup = readback_tup(net, book, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_read: expected 2-tuple\n");
     return new_port(ERA, 0);
   }
 
-  Pair args = node_load(net, get_val(argm));
-  Port file = expand(net, book, get_fst(args));
-  u32 num_bytes = get_u24(get_val(expand(net, book, get_snd(args))));
+  FILE* fp = readback_file(tup.elem_buf[0]);
+  u32 num_bytes = get_u24(get_val(tup.elem_buf[1]));
 
-  FILE* fp = readback_file(file);
   if (fp == NULL) {
     fprintf(stderr, "io_read: invalid file descriptor\n");
     return new_port(ERA, 0);
@@ -292,14 +317,14 @@ Port io_read(Net* net, Book* book, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file name and mode as strings.
 Port io_open(Net* net, Book* book, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_open: expected tuple, but got %u\n", get_tag(argm));
+  Tup tup = readback_tup(net, book, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_open: expected 2-tuple\n");
     return new_port(ERA, 0);
   }
 
-  Pair args = node_load(net, get_val(argm));
-  Str name = readback_str(net, book, expand(net, book, get_fst(args)));
-  Str mode = readback_str(net, book, expand(net, book, get_snd(args)));
+  Str name = readback_str(net, book, tup.elem_buf[0]);
+  Str mode = readback_str(net, book, tup.elem_buf[1]);
 
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
@@ -334,14 +359,14 @@ Port io_close(Net* net, Book* book, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file descriptor and list of bytes to write.
 Port io_write(Net* net, Book* book, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_write: expected tuple, but got %u\n", get_tag(argm));
+  Tup tup = readback_tup(net, book, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_write: expected 2-tuple\n");
     return new_port(ERA, 0);
   }
 
-  Pair args = node_load(net, get_val(argm));
-  FILE* fp = readback_file(expand(net, book, get_fst(args)));
-  Bytes bytes = readback_bytes(net, book, get_snd(args));
+  FILE* fp = readback_file(expand(net, book, tup.elem_buf[0]));
+  Bytes bytes = readback_bytes(net, book, tup.elem_buf[1]);
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
@@ -366,22 +391,15 @@ Port io_write(Net* net, Book* book, Port argm) {
 //    - 1 (SEEK_CUR): current position of the file pointer
 //    - 2 (SEEK_END): end of the file
 Port io_seek(Net* net, Book* book, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(argm));
-    return new_port(ERA, 0);
-  }
-
-  Pair args1 = node_load(net, get_val(argm));
-  Port argm2 = expand(net, book, get_snd(args1));
-  if (get_tag(argm2) != CON) {
-    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(argm2));
+  Tup tup = readback_tup(net, book, argm, 3);
+  if (tup.elem_len != 3) {
+    fprintf(stderr, "io_seek: expected 3-tuple\n");
     return new_port(ERA, 0);
   }
-  Pair args2 = node_load(net, get_val(argm2));
 
-  FILE* fp = readback_file(expand(net, book, get_fst(args1)));
-  i32 offset = get_i24(get_val(expand(net, book, get_fst(args2))));
-  u32 whence = get_i24(get_val(expand(net, book, get_snd(args2))));
+  FILE* fp = readback_file(tup.elem_buf[0]);
+  i32 offset = get_i24(get_val(tup.elem_buf[1]));
+  u32 whence = get_i24(get_val(tup.elem_buf[2]));
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
@@ -425,11 +443,17 @@ Port io_get_time(Net* net, Book* book, Port argm) {
 // `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) {
+  Tup tup = readback_tup(net, book, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_sleep: expected 3-tuple\n");
+    return new_port(ERA, 0);
+  }
+
   // 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
-  u32 dur_hi = get_u24(get_val(expand(net, book, get_fst(dur_node))));
-  u32 dur_lo = get_u24(get_val(expand(net, book, get_snd(dur_node))));
+  u32 dur_hi = get_u24(get_val(tup.elem_buf[0]));
+  u32 dur_lo = get_u24(get_val(tup.elem_buf[1]));
   // Combine into a 48-bit duration in nanoseconds
   u64 dur_ns = (((u64)dur_hi) << 24) | dur_lo;
   // Sleep for the specified duration

From 9d6d0f61b18532b86d62bb6fba8f0b7b38472846 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Mon, 17 Jun 2024 11:15:50 +0200
Subject: [PATCH 10/11] readback_tup run.cu

---
 src/run.c  |  2 +-
 src/run.cu | 87 +++++++++++++++++++++++++++++++++++-------------------
 2 files changed, 57 insertions(+), 32 deletions(-)

diff --git a/src/run.c b/src/run.c
index 3dd8a129..28304508 100644
--- a/src/run.c
+++ b/src/run.c
@@ -365,7 +365,7 @@ Port io_write(Net* net, Book* book, Port argm) {
     return new_port(ERA, 0);
   }
 
-  FILE* fp = readback_file(expand(net, book, tup.elem_buf[0]));
+  FILE* fp = readback_file(tup.elem_buf[0]);
   Bytes bytes = readback_bytes(net, book, tup.elem_buf[1]);
 
   if (fp == NULL) {
diff --git a/src/run.cu b/src/run.cu
index 9a814035..d8a7c000 100644
--- a/src/run.cu
+++ b/src/run.cu
@@ -7,6 +7,12 @@ struct Ctr {
   Port args_buf[16];
 };
 
+// Readback: Tuples
+struct Tup {
+  u32  elem_len;
+  Port elem_buf[8];
+};
+
 // Readback: λ-Encoded Str (UTF-32)
 // FIXME: this is actually ASCII :|
 // FIXME: remove len limit
@@ -72,6 +78,27 @@ Ctr gnet_readback_ctr(GNet* gnet, Port port) {
   return ctr;
 }
 
+// Reads back a tuple of at most `size` elements. Tuples are
+// (right-nested con nodes) (CON 1 (CON 2 (CON 3 (...))))
+// The provided `port` should be `expanded` before calling.
+Tup gnet_readback_tup(GNet* gnet, Port port, u32 size) {
+  Tup tup;
+  tup.elem_len = 0;
+
+  // Loads remaining arguments
+  while (get_tag(port) == CON && (tup.elem_len + 1 < size)) {
+    Pair node = gnet_node_load(gnet, get_val(port));
+    tup.elem_buf[tup.elem_len++] = gnet_expand(gnet, get_fst(node));
+
+    port = gnet_expand(gnet, get_snd(node));
+  }
+
+  tup.elem_buf[tup.elem_len++] = port;
+
+  return tup;
+}
+
+
 // Reads back a UTF-32 (truncated to 24 bits) string.
 // Since unicode scalars can fit in 21 bits, HVM's u24
 // integers can contain any unicode scalar value.
@@ -292,16 +319,15 @@ FILE* readback_file(Port port) {
 // Reads from a file a specified number of bytes.
 // `argm` is a tuple of (file_descriptor, num_bytes).
 Port io_read(GNet* gnet, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_read: expected tuple, but got %u\n", get_tag(argm));
+  Tup tup = gnet_readback_tup(gnet, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_read: expected 2-tuple\n");
     return new_port(ERA, 0);
   }
 
-  Pair args = gnet_node_load(gnet, get_val(argm));
-  Port file = gnet_expand(gnet, get_fst(args));
-  u32 num_bytes = get_u24(get_val(gnet_expand(gnet, get_snd(args))));
+  FILE* fp = readback_file(tup.elem_buf[0]);
+  u32 num_bytes = get_u24(get_val(tup.elem_buf[1]));
 
-  FILE* fp = readback_file(file);
   if (fp == NULL) {
     fprintf(stderr, "io_read: invalid file descriptor\n");
     return new_port(ERA, 0);
@@ -328,14 +354,14 @@ Port io_read(GNet* gnet, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file name and mode as strings.
 Port io_open(GNet* gnet, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_open: expected tuple\n");
+  Tup tup = gnet_readback_tup(gnet, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_open: expected 2-tuple\n");
     return new_port(ERA, 0);
   }
 
-  Pair args = gnet_node_load(gnet, get_val(argm));
-  Str name = gnet_readback_str(gnet, gnet_expand(gnet, get_fst(args)));
-  Str mode = gnet_readback_str(gnet, gnet_expand(gnet, get_snd(args)));
+  Str name = gnet_readback_str(gnet, tup.elem_buf[0]);
+  Str mode = gnet_readback_str(gnet, tup.elem_buf[1]);
 
   for (u32 fd = 3; fd < sizeof(FILE_POINTERS); fd++) {
     if (FILE_POINTERS[fd] == NULL) {
@@ -370,14 +396,14 @@ Port io_close(GNet* gnet, Port argm) {
 // `argm` is a tuple (CON node) of the
 // file descriptor and list of bytes to write.
 Port io_write(GNet* gnet, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_write: expected tuple, but got %u\n", get_tag(argm));
+  Tup tup = gnet_readback_tup(gnet, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_write: expected 2-tuple\n");
     return new_port(ERA, 0);
   }
 
-  Pair args = gnet_node_load(gnet, get_val(argm));
-  FILE* fp = readback_file(gnet_expand(gnet, get_fst(args)));
-  Bytes bytes = gnet_readback_bytes(gnet, get_snd(args));
+  FILE* fp = readback_file(tup.elem_buf[0]);
+  Bytes bytes = gnet_readback_bytes(gnet, tup.elem_buf[1]);
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
@@ -402,22 +428,15 @@ Port io_write(GNet* gnet, Port argm) {
 //    - 1 (SEEK_CUR): current position of the file pointer
 //    - 2 (SEEK_END): end of the file
 Port io_seek(GNet* gnet, Port argm) {
-  if (get_tag(argm) != CON) {
-    fprintf(stderr, "io_seek: expected first tuple, but got %u\n", get_tag(argm));
-    return new_port(ERA, 0);
-  }
-
-  Pair args1 = gnet_node_load(gnet, get_val(argm));
-  Port argm2 = gnet_expand(gnet, get_snd(args1));
-  if (get_tag(argm2) != CON) {
-    fprintf(stderr, "io_seek: expected second tuple, but got %u\n", get_tag(argm2));
+  Tup tup = gnet_readback_tup(gnet, argm, 3);
+  if (tup.elem_len != 3) {
+    fprintf(stderr, "io_seek: expected 3-tuple\n");
     return new_port(ERA, 0);
   }
-  Pair args2 = gnet_node_load(gnet, get_val(argm2));
 
-  FILE* fp = readback_file(gnet_expand(gnet, get_fst(args1)));
-  i32 offset = get_i24(get_val(gnet_expand(gnet, get_fst(args2))));
-  u32 whence = get_i24(get_val(gnet_expand(gnet, get_snd(args2))));
+  FILE* fp = readback_file(tup.elem_buf[0]);
+  i32 offset = get_i24(get_val(tup.elem_buf[1]));
+  u32 whence = get_i24(get_val(tup.elem_buf[2]));
 
   if (fp == NULL) {
     fprintf(stderr, "io_write: invalid file descriptor\n");
@@ -457,11 +476,17 @@ Port io_get_time(GNet* gnet, Port argm) {
 // `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) {
+  Tup tup = gnet_readback_tup(gnet, argm, 2);
+  if (tup.elem_len != 2) {
+    fprintf(stderr, "io_sleep: expected 3-tuple\n");
+    return new_port(ERA, 0);
+  }
+
   // 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
-  u32 dur_hi = get_u24(get_val(gnet_expand(gnet, get_fst(dur_node))));
-  u32 dur_lo = get_u24(get_val(gnet_expand(gnet, get_snd(dur_node))));
+  u32 dur_hi = get_u24(get_val(tup.elem_buf[0]));
+  u32 dur_lo = get_u24(get_val(tup.elem_buf[1]));
   // Combine into a 48-bit duration in nanoseconds
   u64 dur_ns = (((u64)dur_hi) << 24) | dur_lo;
   // Sleep for the specified duration

From 2b4ce85907ec346349c25dae31824588866a3583 Mon Sep 17 00:00:00 2001
From: Enrico Zandomeni Borba <enricozb@gmail.com>
Date: Mon, 17 Jun 2024 20:24:08 +0200
Subject: [PATCH 11/11] typo

---
 src/run.c | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/run.c b/src/run.c
index 28304508..dc1eb187 100644
--- a/src/run.c
+++ b/src/run.c
@@ -445,7 +445,7 @@ Port io_get_time(Net* net, Book* book, Port argm) {
 Port io_sleep(Net* net, Book* book, Port argm) {
   Tup tup = readback_tup(net, book, argm, 2);
   if (tup.elem_len != 2) {
-    fprintf(stderr, "io_sleep: expected 3-tuple\n");
+    fprintf(stderr, "io_sleep: expected 2-tuple\n");
     return new_port(ERA, 0);
   }