diff --git a/build.rs b/build.rs index aecd18ae..dd72c946 100644 --- a/build.rs +++ b/build.rs @@ -6,6 +6,7 @@ fn main() { println!("cargo:rerun-if-changed=src/hvm.c"); println!("cargo:rerun-if-changed=src/run.cu"); println!("cargo:rerun-if-changed=src/hvm.cu"); + println!("cargo:rustc-link-arg=-rdynamic"); match cc::Build::new() .file("src/run.c") diff --git a/src/hvm.cuh b/src/hvm.cuh new file mode 100644 index 00000000..d4c85f70 --- /dev/null +++ b/src/hvm.cuh @@ -0,0 +1,220 @@ +#ifndef hvm_cuh_INCLUDED +#define hvm_cuh_INCLUDED + +#include +#include + +// Types +// ----- + +typedef uint8_t bool; +typedef uint8_t u8; +typedef uint16_t u16; +typedef uint32_t u32; +typedef unsigned long long int u64; +typedef int32_t i32; +typedef float f32; +typedef double f64; + +// Local Types +typedef u8 Tag; // Tag ::= 3-bit (rounded up to u8) +typedef u32 Val; // Val ::= 29-bit (rounded up to u32) +typedef u32 Port; // Port ::= Tag + Val (fits a u32) +typedef u64 Pair; // Pair ::= Port + Port (fits a u64) + +// Numbs +typedef u32 Numb; // Numb ::= 29-bit (rounded up to u32) + +// Tags +const Tag VAR = 0x0; // variable +const Tag REF = 0x1; // reference +const Tag ERA = 0x2; // eraser +const Tag NUM = 0x3; // number +const Tag CON = 0x4; // constructor +const Tag DUP = 0x5; // duplicator +const Tag OPR = 0x6; // operator +const Tag SWI = 0x7; // switch + +// Numbers +static const f32 U24_MAX = (f32) (1 << 24) - 1; +static const f32 U24_MIN = 0.0; +static const f32 I24_MAX = (f32) (1 << 23) - 1; +static const f32 I24_MIN = (f32) (i32) ((-1u) << 23); +const Tag TY_SYM = 0x00; +const Tag TY_U24 = 0x01; +const Tag TY_I24 = 0x02; +const Tag TY_F24 = 0x03; +const Tag OP_ADD = 0x04; +const Tag OP_SUB = 0x05; +const Tag FP_SUB = 0x06; +const Tag OP_MUL = 0x07; +const Tag OP_DIV = 0x08; +const Tag FP_DIV = 0x09; +const Tag OP_REM = 0x0A; +const Tag FP_REM = 0x0B; +const Tag OP_EQ = 0x0C; +const Tag OP_NEQ = 0x0D; +const Tag OP_LT = 0x0E; +const Tag OP_GT = 0x0F; +const Tag OP_AND = 0x10; +const Tag OP_OR = 0x11; +const Tag OP_XOR = 0x12; +const Tag OP_SHL = 0x13; +const Tag FP_SHL = 0x14; +const Tag OP_SHR = 0x15; +const Tag FP_SHR = 0x16; + +typedef struct GNet GNet; + +// Debugger +// -------- + +// Port: Constructor and Getters +// ----------------------------- + +static inline Port new_port(Tag tag, Val val) { + return (val << 3) | tag; +} + +static inline Tag get_tag(Port port) { + return port & 7; +} + +static inline Val get_val(Port port) { + return port >> 3; +} + +// Pair: Constructor and Getters +// ----------------------------- + +static inline const Pair new_pair(Port fst, Port snd) { + return ((u64)snd << 32) | fst; +} + +static inline Port get_fst(Pair pair) { + return pair & 0xFFFFFFFF; +} + +static inline Port get_snd(Pair pair) { + return pair >> 32; +} + +// Utils +// ----- + +// Swaps two ports. +static inline void swap(Port *a, Port *b) { + Port x = *a; *a = *b; *b = x; +} + +static inline u32 min(u32 a, u32 b) { + return (a < b) ? a : b; +} + +static inline f32 clamp(f32 x, f32 min, f32 max) { + const f32 t = x < min ? min : x; + return (t > max) ? max : t; +} + +// Numbs +// ----- + +// Constructor and getters for SYM (operation selector) +static inline Numb new_sym(u32 val) { + return (val << 5) | TY_SYM; +} + +static inline u32 get_sym(Numb word) { + return (word >> 5); +} + +// Constructor and getters for U24 (unsigned 24-bit integer) +static inline Numb new_u24(u32 val) { + return (val << 5) | TY_U24; +} + +static inline u32 get_u24(Numb word) { + return word >> 5; +} + +// Constructor and getters for I24 (signed 24-bit integer) +static inline Numb new_i24(i32 val) { + return ((u32)val << 5) | TY_I24; +} + +static inline i32 get_i24(Numb word) { + return ((i32)word) << 3 >> 8; +} + +// Constructor and getters for F24 (24-bit float) +static inline Numb new_f24(float val) { + u32 bits = *(u32*)&val; + u32 shifted_bits = bits >> 8; + u32 lost_bits = bits & 0xFF; + // round ties to even + shifted_bits += (!isnan(val)) & ((lost_bits - ((lost_bits >> 7) & !shifted_bits)) >> 7); + // ensure NaNs don't become infinities + shifted_bits |= isnan(val); + return (shifted_bits << 5) | TY_F24; +} + +static inline float get_f24(Numb word) { + u32 bits = (word << 3) & 0xFFFFFF00; + return *(float*)&bits; +} + +static inline Tag get_typ(Numb word) { + return word & 0x1F; +} + +static inline bool is_num(Numb word) { + return get_typ(word) >= TY_U24 && get_typ(word) <= TY_F24; +} + +static inline bool is_cast(Numb word) { + return get_typ(word) == TY_SYM && get_sym(word) >= TY_U24 && get_sym(word) <= TY_F24; +} + +// Partial application +static inline Numb partial(Numb a, Numb b) { + return (b & ~0x1F) | get_sym(a); +} + +// Readback +// --------- + +// Readback: Tuples +typedef struct Tup { + u32 elem_len; + Port elem_buf[8]; +} 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(GNet* gnet, Port port, u32 size); + +typedef struct Str { + u32 text_len; + char text_buf[256]; +} Str; + +// Reads a constructor-encoded string (of length at most 255 characters), +// into a null-terminated `Str`. +extern Str gnet_readback_str(GNet* gnet, Port port); + +typedef struct Bytes { + u32 len; + char *buf; +} Bytes; + +// Reads a constructor-encoded string (of length at most 256 characters), +// into a `Bytes`. The returned `Bytes` is not null terminated. +extern Bytes gnet_readback_bytes(GNet* net, Port port); + +// Creates a construtor-encoded string of arbitrary length from the +// provided `bytes`. This string can be consumed on the HVM-side. This +// will return an `ERA` if nodes cannot be allocated. +extern Port gnet_inject_bytes(GNet* net, Bytes *bytes); + +#endif // hvm_cuh_INCLUDED diff --git a/src/hvm.h b/src/hvm.h new file mode 100644 index 00000000..3aab66b1 --- /dev/null +++ b/src/hvm.h @@ -0,0 +1,231 @@ +#ifndef hvm_h_INCLUDED +#define hvm_h_INCLUDED + +#include +#include + +// Types +// ----- + +typedef uint8_t bool; +typedef uint8_t u8; +typedef uint16_t u16; +typedef uint32_t u32; +typedef int32_t i32; +typedef uint64_t u64; +typedef float f32; +typedef double f64; + +// Local Types +typedef u8 Tag; // Tag ::= 3-bit (rounded up to u8) +typedef u32 Val; // Val ::= 29-bit (rounded up to u32) +typedef u32 Port; // Port ::= Tag + Val (fits a u32) +typedef u64 Pair; // Pair ::= Port + Port (fits a u64) + +// Numbs +typedef u32 Numb; // Numb ::= 29-bit (rounded up to u32) + +// Tags +#define VAR 0x0 // variable +#define REF 0x1 // reference +#define ERA 0x2 // eraser +#define NUM 0x3 // number +#define CON 0x4 // constructor +#define DUP 0x5 // duplicator +#define OPR 0x6 // operator +#define SWI 0x7 // switch + +// Numbers +static const f32 U24_MAX = (f32) (1 << 24) - 1; +static const f32 U24_MIN = 0.0; +static const f32 I24_MAX = (f32) (1 << 23) - 1; +static const f32 I24_MIN = (f32) (i32) ((-1u) << 23); +#define TY_SYM 0x00 +#define TY_U24 0x01 +#define TY_I24 0x02 +#define TY_F24 0x03 +#define OP_ADD 0x04 +#define OP_SUB 0x05 +#define FP_SUB 0x06 +#define OP_MUL 0x07 +#define OP_DIV 0x08 +#define FP_DIV 0x09 +#define OP_REM 0x0A +#define FP_REM 0x0B +#define OP_EQ 0x0C +#define OP_NEQ 0x0D +#define OP_LT 0x0E +#define OP_GT 0x0F +#define OP_AND 0x10 +#define OP_OR 0x11 +#define OP_XOR 0x12 +#define OP_SHL 0x13 +#define FP_SHL 0x14 +#define OP_SHR 0x15 +#define FP_SHR 0x16 + +typedef struct Net Net; +typedef struct Book Book; + +// Debugger +// -------- + +typedef struct { + char x[13]; +} Show; + +void put_u16(char* B, u16 val); +Show show_port(Port port); +void print_net(Net* net); +void pretty_print_numb(Numb word); +void pretty_print_port(Net* net, Book* book, Port port); + +// Port: Constructor and Getters +// ----------------------------- + +static inline Port new_port(Tag tag, Val val) { + return (val << 3) | tag; +} + +static inline Tag get_tag(Port port) { + return port & 7; +} + +static inline Val get_val(Port port) { + return port >> 3; +} + +// Pair: Constructor and Getters +// ----------------------------- + +static inline const Pair new_pair(Port fst, Port snd) { + return ((u64)snd << 32) | fst; +} + +static inline Port get_fst(Pair pair) { + return pair & 0xFFFFFFFF; +} + +static inline Port get_snd(Pair pair) { + return pair >> 32; +} + +// Utils +// ----- + +// Swaps two ports. +static inline void swap(Port *a, Port *b) { + Port x = *a; *a = *b; *b = x; +} + +static inline u32 min(u32 a, u32 b) { + return (a < b) ? a : b; +} + +static inline f32 clamp(f32 x, f32 min, f32 max) { + const f32 t = x < min ? min : x; + return (t > max) ? max : t; +} + +// Numbs +// ----- + +// Constructor and getters for SYM (operation selector) +static inline Numb new_sym(u32 val) { + return (val << 5) | TY_SYM; +} + +static inline u32 get_sym(Numb word) { + return (word >> 5); +} + +// Constructor and getters for U24 (unsigned 24-bit integer) +static inline Numb new_u24(u32 val) { + return (val << 5) | TY_U24; +} + +static inline u32 get_u24(Numb word) { + return word >> 5; +} + +// Constructor and getters for I24 (signed 24-bit integer) +static inline Numb new_i24(i32 val) { + return ((u32)val << 5) | TY_I24; +} + +static inline i32 get_i24(Numb word) { + return ((i32)word) << 3 >> 8; +} + +// Constructor and getters for F24 (24-bit float) +static inline Numb new_f24(float val) { + u32 bits = *(u32*)&val; + u32 shifted_bits = bits >> 8; + u32 lost_bits = bits & 0xFF; + // round ties to even + shifted_bits += (!isnan(val)) & ((lost_bits - ((lost_bits >> 7) & !shifted_bits)) >> 7); + // ensure NaNs don't become infinities + shifted_bits |= isnan(val); + return (shifted_bits << 5) | TY_F24; +} + +static inline float get_f24(Numb word) { + u32 bits = (word << 3) & 0xFFFFFF00; + return *(float*)&bits; +} + +static inline Tag get_typ(Numb word) { + return word & 0x1F; +} + +static inline bool is_num(Numb word) { + return get_typ(word) >= TY_U24 && get_typ(word) <= TY_F24; +} + +static inline bool is_cast(Numb word) { + return get_typ(word) == TY_SYM && get_sym(word) >= TY_U24 && get_sym(word) <= TY_F24; +} + +// Partial application +static inline Numb partial(Numb a, Numb b) { + return (b & ~0x1F) | get_sym(a); +} + +// Readback +// --------- + +// Readback: Tuples +typedef struct Tup { + u32 elem_len; + Port elem_buf[8]; +} 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 readback_tup(Net* net, Book* book, Port port, u32 size); + +typedef struct Str { + u32 text_len; + char text_buf[256]; +} Str; + +// Reads a constructor-encoded string (of length at most 255 characters), +// into a null-terminated `Str`. +extern Str readback_str(Net* net, Book* book, Port port); + +typedef struct Bytes { + u32 len; + char *buf; +} Bytes; + +// Reads a constructor-encoded string (of length at most 256 characters), +// into a `Bytes`. The returned `Bytes` is not null terminated. +extern Bytes readback_bytes(Net* net, Book* book, Port port); + +// Creates a construtor-encoded string of arbitrary length from the +// provided `bytes`. This string can be consumed on the HVM-side. This +// will return an `ERA` if nodes cannot be allocated. +extern Port inject_bytes(Net* net, Bytes *bytes); + +#endif // hvm_h_INCLUDED diff --git a/src/run.c b/src/run.c index 59710629..d449d9fa 100644 --- a/src/run.c +++ b/src/run.c @@ -1,3 +1,4 @@ +#include #include "hvm.c" // Readback: λ-Encoded Ctr @@ -81,7 +82,7 @@ Ctr readback_ctr(Net* net, Book* book, 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 readback_tup(Net* net, Book* book, Port port, u32 size) { +extern Tup readback_tup(Net* net, Book* book, Port port, u32 size) { Tup tup; tup.elem_len = 0; @@ -258,6 +259,10 @@ Port inject_bytes(Net* net, 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) { @@ -280,6 +285,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(Net* net, Book* book, Port argm) { @@ -483,6 +506,77 @@ Port io_sleep(Net* net, Book* book, 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(Net* net, Book* book, Port argm) { + Tup tup = readback_tup(net, book, argm, 2); + Str str = readback_str(net, book, 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); + } + + 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(Net* net, Book* book, Port argm) { + Tup tup = readback_tup(net, book, 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 = readback_str(net, book, tup.elem_buf[1]); + + dlerror(); + Port (*func)(Net*, Book*, 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(net, book, 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); +} + // Book Loader // ----------- @@ -495,6 +589,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}; } // Monadic IO Evaluator diff --git a/src/run.cu b/src/run.cu index a54a941e..78c790ae 100644 --- a/src/run.cu +++ b/src/run.cu @@ -1,3 +1,4 @@ +#include #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)); }