cuda support

src/hvm.cuh

@@ -1,12 +1,13 @@
 #ifndef hvm_cuh_INCLUDED
 #define hvm_cuh_INCLUDED
 
+#include <math.h>
 #include <stdint.h>
-#include <stdlib.h>
 
 // Types
 // -----
 
+typedef uint8_t bool;
 typedef  uint8_t u8;
 typedef uint16_t u16;
 typedef uint32_t u32;
@@ -191,7 +192,7 @@ typedef struct Tup {
 // Reads a tuple of `size` elements from `port`.
 // Tuples are con nodes nested to the right auxilliary port,
 // For example, `(CON a (CON b (CON c)))` is a 3-tuple (a, b, c).
-extern Tup gnet_readback_tup(Net* net, Port port, u32 size);
+extern Tup gnet_readback_tup(GNet* gnet, Port port, u32 size);
 
 typedef struct Str {
   u32  text_len;
@@ -200,7 +201,7 @@ typedef struct Str {
 
 // Reads a constructor-encoded string (of length at most 255 characters),
 // into a null-terminated `Str`.
-extern Str gnet_readback_str(GNet* net, Port port);
+extern Str gnet_readback_str(GNet* gnet, Port port);
 
 typedef struct Bytes {
   u32  len;

src/hvm.h

@@ -4,10 +4,10 @@
 #include <math.h>
 #include <stdint.h>
 
-typedef uint8_t bool;
 // Types
 // -----
 
+typedef uint8_t bool;
 typedef  uint8_t  u8;
 typedef uint16_t u16;
 typedef uint32_t u32;

src/run.c

@@ -286,15 +286,15 @@ FILE* readback_file(Port port) {
 }
 
 // Converts a NUM port (dylib handle) to an opaque dylib object.
-FILE* readback_dylib(Port port) {
+void* readback_dylib(Port port) {
   if (get_tag(port) != NUM) {
     fprintf(stderr, "non-num where dylib handle was expected: %i\n", get_tag(port));
     return NULL;
   }
 
   u32 idx = get_u24(get_val(port));
 
-  FILE* dl = DYLIBS[idx];
+  void* dl = DYLIBS[idx];
   if (dl == NULL) {
     fprintf(stderr, "invalid dylib handle\n");
     return NULL;
@@ -556,7 +556,7 @@ Port io_dl_call(Net* net, Book* book, Port argm) {
 
 // Closes a loaded dylib, reclaiming the handle.
 Port io_dl_close(Net* net, Book* book, Port argm) {
-  FILE* dl = readback_dylib(argm);
+  void* dl = readback_dylib(argm);
   if (dl == NULL) {
     fprintf(stderr, "io_dl_close: invalid handle\n");
     return new_port(ERA, 0);

src/run.cu

@@ -1,3 +1,4 @@
+#include <dlfcn.h>
 #include "hvm.cu"
 
 // Readback: λ-Encoded Ctr
@@ -81,7 +82,7 @@ Ctr gnet_readback_ctr(GNet* gnet, Port port) {
 // 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) {
+extern "C" Tup gnet_readback_tup(GNet* gnet, Port port, u32 size) {
   Tup tup;
   tup.elem_len = 0;
 
@@ -105,7 +106,7 @@ Tup gnet_readback_tup(GNet* gnet, Port port, u32 size) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Str gnet_readback_str(GNet* gnet, Port port) {
+extern "C" Str gnet_readback_str(GNet* gnet, Port port) {
   // Result
   Str str;
   str.text_len = 0;
@@ -146,7 +147,7 @@ Str gnet_readback_str(GNet* gnet, Port port) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Bytes gnet_readback_bytes(GNet* gnet, Port port) {
+extern "C" Bytes gnet_readback_bytes(GNet* gnet, Port port) {
   // Result
   Bytes bytes;
   bytes.buf = (char*) malloc(sizeof(char) * MAX_BYTES);
@@ -263,7 +264,7 @@ __global__ void make_bytes_port(GNet* gnet, Bytes bytes, Port* ret) {
 // Encoding:
 // - λt (t NIL)
 // - λt (((t CONS) head) tail)
-Port gnet_inject_bytes(GNet* gnet, Bytes *bytes) {
+extern "C" Port gnet_inject_bytes(GNet* gnet, Bytes *bytes) {
   Port* d_ret;
   cudaMalloc(&d_ret, sizeof(Port));
 
@@ -294,6 +295,10 @@ Port gnet_inject_bytes(GNet* gnet, Bytes *bytes) {
 // - 2 -> stderr
 static FILE* FILE_POINTERS[256];
 
+// Open dylibs handles. Indices into this array
+// are used as opaque loadedd object "handles".
+static void* DYLIBS[256];
+
 // Converts a NUM port (file descriptor) to file pointer.
 FILE* readback_file(Port port) {
   if (get_tag(port) != NUM) {
@@ -316,6 +321,24 @@ FILE* readback_file(Port port) {
   return fp;
 }
 
+// Converts a NUM port (dylib handle) to an opaque dylib object.
+void* readback_dylib(Port port) {
+  if (get_tag(port) != NUM) {
+    fprintf(stderr, "non-num where dylib handle was expected: %i\n", get_tag(port));
+    return NULL;
+  }
+
+  u32 idx = get_u24(get_val(port));
+
+  void* dl = DYLIBS[idx];
+  if (dl == NULL) {
+    fprintf(stderr, "invalid dylib handle\n");
+    return NULL;
+  }
+
+  return dl;
+}
+
 // 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) {
@@ -515,6 +538,79 @@ Port io_sleep(GNet* gnet, Port argm) {
   return new_port(ERA, 0);
 }
 
+// Opens a dylib at the provided path.
+// `argm` is a tuple of `filename` and `lazy`.
+// `filename` is a λ-encoded string.
+// `lazy` is a `bool` indicating if functions should be lazily loaded.
+Port io_dl_open(GNet* gnet, Port argm) {
+  Tup tup = gnet_readback_tup(gnet, argm, 2);
+  Str str = gnet_readback_str(gnet, tup.elem_buf[0]);
+  u32 lazy = get_u24(get_val(tup.elem_buf[1]));
+
+  int flags = lazy ? RTLD_LAZY : RTLD_NOW;
+
+  for (u32 dl = 0; dl < sizeof(DYLIBS); dl++) {
+    if (DYLIBS[dl] == NULL) {
+      DYLIBS[dl] = dlopen(str.text_buf, flags);
+      if (DYLIBS[dl] == NULL) {
+        fprintf(stderr, "failed to open dylib '%s': %s\n", str.text_buf, dlerror());
+
+        return new_port(ERA, 0);
+      } else {
+        fprintf(stderr, "opened dylib '%s'\n", str.text_buf);
+      }
+
+      return new_port(NUM, new_u24(dl));
+    }
+  }
+
+  fprintf(stderr, "io_dl_open: too many open dylibs\n");
+  return new_port(ERA, 0);
+}
+
+// Calls a function from a loaded dylib.
+// `argm` is a 3-tuple of `dylib_handle`, `symbol`, `args`.
+// `dylib_handle` is the numeric node returned from a `DL_OPEN` call.
+// `symbol` is a λ-encoded string of the symbol name.
+// `args` is the argument to be provided to the dylib symbol.
+Port io_dl_call(GNet* gnet, Port argm) {
+  Tup tup = gnet_readback_tup(gnet, argm, 3);
+  if (tup.elem_len != 3) {
+    fprintf(stderr, "io_dl_call: expected 3-tuple\n");
+    return new_port(ERA, 0);
+  }
+
+  void* dl = readback_dylib(tup.elem_buf[0]);
+  Str symbol = gnet_readback_str(gnet, tup.elem_buf[1]);
+
+  dlerror();
+  Port (*func)(GNet*, Port) = (Port (*)(GNet*, Port)) dlsym(dl, symbol.text_buf);
+  char* error = dlerror();
+  if (error != NULL) {
+    fprintf(stderr, "io_dl_call: failed to get symbol '%s': %s\n", symbol.text_buf, error);
+  }
+
+  return func(gnet, tup.elem_buf[2]);
+}
+
+// Closes a loaded dylib, reclaiming the handle.
+Port io_dl_close(Net* net, Book* book, Port argm) {
+  void* dl = readback_dylib(argm);
+  if (dl == NULL) {
+    fprintf(stderr, "io_dl_close: invalid handle\n");
+    return new_port(ERA, 0);
+  }
+
+  int err = dlclose(dl) != 0;
+  if (err != 0) {
+    fprintf(stderr, "io_dl_close: failed to close: %i\n", err);
+    return new_port(ERA, 0);
+  }
+
+  DYLIBS[get_u24(get_val(argm))] = NULL;
+  return new_port(ERA, 0);
+}
+
 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};
@@ -524,6 +620,9 @@ void book_init(Book* book) {
   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};
+  book->ffns_buf[book->ffns_len++] = (FFn){"DL_OPEN", io_dl_open};
+  book->ffns_buf[book->ffns_len++] = (FFn){"DL_CALL", io_dl_call};
+  book->ffns_buf[book->ffns_len++] = (FFn){"DL_CLOSE", io_dl_open};
 
   cudaMemcpyToSymbol(BOOK, book, sizeof(Book));
 }