diff --git a/.gitignore b/.gitignore index 91380077..2b9560b3 100644 --- a/.gitignore +++ b/.gitignore @@ -10,3 +10,7 @@ examples/**/main examples/**/*.c examples/**/*.cu .out.hvm + +# nix-direnv +/.direnv/ +/.envrc diff --git a/src/hvm.c b/src/hvm.c index d35b6f83..a1bb0217 100644 --- a/src/hvm.c +++ b/src/hvm.c @@ -171,7 +171,7 @@ typedef struct Ctr { Port args_buf[16]; } Ctr; -// Readback: λ-Encoded Str (UTF-16) +// Readback: λ-Encoded Str (UTF-32) // FIXME: this is actually ASCII :| // FIXME: remove len limit typedef struct Str { @@ -765,7 +765,7 @@ static inline void link(Net* net, TM* tm, Port A, Port B) { if (get_tag(A) != VAR && get_tag(B) == VAR) { Port X = A; A = B; B = X; } - + // If `A` is NODE: create the `A ~ B` redex if (get_tag(A) != VAR) { push_redex(net, tm, new_pair(A, B)); // TODO: move global ports to local @@ -806,7 +806,7 @@ static inline bool interact_link(Net* net, TM* tm, Port a, Port b) { if (!get_resources(net, tm, 1, 0, 0)) { return FALSE; } - + // Links. link_pair(net, tm, new_pair(a, b)); @@ -879,7 +879,7 @@ static inline bool interact_eras(Net* net, TM* tm, Port a, Port b) { Pair B = node_exchange(net, get_val(b), 0); Port B1 = get_fst(B); Port B2 = get_snd(B); - + //if (B == 0) printf("[%04x] ERROR2: %s\n", tid, show_port(b).x); // Links. @@ -910,7 +910,7 @@ static inline bool interact_anni(Net* net, TM* tm, Port a, Port b) { Pair B = node_take(net, get_val(b)); Port B1 = get_fst(B); Port B2 = get_snd(B); - + //if (A == 0) printf("[%04x] ERROR3: %s\n", tid, show_port(a).x); //if (B == 0) printf("[%04x] ERROR4: %s\n", tid, show_port(b).x); @@ -985,7 +985,7 @@ static inline bool interact_oper(Net* net, TM* tm, Port a, Port b) { Pair B = node_take(net, get_val(b)); Port B1 = get_fst(B); Port B2 = enter(net, get_snd(B)); - + // Performs operation. if (get_tag(B1) == NUM) { Val bv = get_val(B1); @@ -1016,7 +1016,7 @@ static inline bool interact_swit(Net* net, TM* tm, Port a, Port b) { Pair B = node_take(net, get_val(b)); Port B1 = get_fst(B); Port B2 = get_snd(B); - + // Stores new nodes. if (av == 0) { node_create(net, tm->nloc[0], new_pair(B2, new_port(ERA,0))); @@ -1124,7 +1124,7 @@ void evaluator(Net* net, TM* tm, Book* book) { // Update global idle counter if (busy) atomic_fetch_add_explicit(&net->idle, 1, memory_order_relaxed); busy = FALSE; - + //// Peeks a redex from target u32 sid = (tm->tid - 1) % TPC; u32 idx = sid*(G_RBAG_LEN/TPC) + (tm->sidx++); @@ -1149,7 +1149,7 @@ void evaluator(Net* net, TM* tm, Book* book) { //} // Stealing Everything: this will steal all redexes - + Pair got = atomic_exchange_explicit(&net->rbag_buf[idx], 0, memory_order_relaxed); if (got != 0) { //printf("[%04x] stolen one task from %04x | itrs=%d idle=%d | %s ~ %s\n", tm->tid, sid, tm->itrs, atomic_load_explicit(&net->idle, memory_order_relaxed),show_port(get_fst(got)).x, show_port(get_snd(got)).x); @@ -1159,7 +1159,7 @@ void evaluator(Net* net, TM* tm, Book* book) { //printf("[%04x] failed to steal from %04x | itrs=%d idle=%d |\n", tm->tid, sid, tm->itrs, atomic_load_explicit(&net->idle, memory_order_relaxed)); tm->sidx = 0; } - + // Chill... sched_yield(); // Halt if all threads are idle @@ -1272,7 +1272,9 @@ Ctr read_ctr(Net* net, Book* book, Port port) { return ctr; } -// Reads back a UTF-16 string. +// 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. // Encoding: // - λt (t NIL) // - λt (((t CONS) head) tail) @@ -1280,7 +1282,7 @@ Str read_str(Net* net, Book* book, Port port) { // Result Str str; str.text_len = 0; - + // Readback loop while (TRUE) { // Normalizes the net @@ -1335,7 +1337,7 @@ void read_img(Net* net, Port port, u32 width, u32 height, u32* buffer) { Port port = enter(net, rect.port); u32 lv = rect.lv; u32 x0 = rect.x0; - u32 x1 = rect.x1; + u32 x1 = rect.x1; u32 y0 = rect.y0; u32 y1 = rect.y1; if (get_tag(port) == CON) { @@ -1400,7 +1402,7 @@ Port io_put_file(Net* net, Book* book, Port argm) { // IO: GetTime Port io_get_time(Net* net, Book* book, Port argm) { - // Get the current time in nanoseconds + // Get the current time in nanoseconds u64 time_ns = time64(); // Encode the time as a 64-bit unsigned integer u32 time_hi = (u32)(time_ns >> 24) & 0xFFFFFFF; @@ -1413,12 +1415,12 @@ Port io_get_time(Net* net, Book* book, Port argm) { return new_port(CON, loc); } -// IO: PutTime +// IO: PutTime // NOTE: changing this name will corrupt the timeline. You've been warned. Port io_put_time(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 + // 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))); // Combine into a 48-bit duration in nanoseconds @@ -1625,7 +1627,7 @@ void book_load(Book* book, u32* buf) { // Gets def Def* def = &book->defs_buf[fid]; - + // Reads name memcpy(def->name, buf, 256); buf += 64; @@ -1644,7 +1646,7 @@ void book_load(Book* book, u32* buf) { // Reads rbag_buf memcpy(def->rbag_buf, buf, 8*def->rbag_len); buf += def->rbag_len * 2; - + // Reads node_buf memcpy(def->node_buf, buf, 8*def->node_len); buf += def->node_len * 2; @@ -1657,7 +1659,7 @@ void book_load(Book* book, u32* buf) { void put_u32(char* B, u32 val) { for (int i = 0; i < 8; i++, val >>= 4) { B[8-i-1] = "0123456789ABCDEF"[val & 0xF]; - } + } } Show show_port(Port port) { diff --git a/src/hvm.cu b/src/hvm.cu index 4a930396..5a9b821c 100644 --- a/src/hvm.cu +++ b/src/hvm.cu @@ -195,10 +195,20 @@ struct Def { Pair node_buf[L_NODE_LEN/TPB]; }; +typedef struct Book Book; + +// A Foreign Function +typedef struct { + char name[256]; + Port (*func)(GNet*, Port); +} FFn; + // Book of Definitions struct Book { u32 defs_len; - Def defs_buf[0x4000]; // 256 MB + Def defs_buf[0x4000]; + u32 ffns_len; + FFn ffns_buf[0x4000]; }; // Static Book @@ -211,7 +221,7 @@ struct Ctr { Port args_buf[16]; }; -// Readback: λ-Encoded Str (UTF-16) +// Readback: λ-Encoded Str (UTF-32) // FIXME: this is actually ASCII :| // FIXME: remove len limit struct Str { @@ -219,9 +229,17 @@ struct Str { char text_buf[256]; }; -// Str Type -const u32 NIL = 0; -const u32 CONS = 1; +// IO Magic Number +#define IO_MAGIC_0 0xD0CA11 +#define IO_MAGIC_1 0xFF1FF1 + +// IO Tags +#define IO_DONE 0 +#define IO_CALL 1 + +// List Type +#define LIST_NIL 0 +#define LIST_CONS 1 // Debugger // -------- @@ -243,6 +261,13 @@ __global__ void print_heatmap(GNet* gnet, u32 turn); // Utils // ----- +// TODO: write a time64() function that returns the time as fast as possible as a u64 +static inline u64 time64() { + struct timespec ts; + clock_gettime(CLOCK_MONOTONIC, &ts); + return (u64)ts.tv_sec * 1000000000ULL + (u64)ts.tv_nsec; +} + __device__ inline u32 TID() { return threadIdx.x; } @@ -303,7 +328,7 @@ __device__ __host__ Pair set_par_flag(Pair pair) { Port p1 = get_fst(pair); Port p2 = get_snd(pair); if (get_tag(p1) == REF) { - return new_pair(new_port(get_tag(p1), get_val(p1) | 0x10000000), p2); + return new_pair(new_port(get_tag(p1), get_val(p1) | 0x10000000), p2); } else { return pair; } @@ -313,7 +338,7 @@ __device__ __host__ Pair clr_par_flag(Pair pair) { Port p1 = get_fst(pair); Port p2 = get_snd(pair); if (get_tag(p1) == REF) { - return new_pair(new_port(get_tag(p1), get_val(p1) & 0xFFFFFFF), p2); + return new_pair(new_port(get_tag(p1), get_val(p1) & 0xFFFFFFF), p2); } else { return pair; } @@ -1246,7 +1271,7 @@ __device__ bool interact_oper(Net* net, TM* tm, Port a, Port b) { Pair B = node_take(net, get_val(b)); Port B1 = get_fst(B); Port B2 = enter(net, tm, get_snd(B)); - + // Performs operation. if (get_tag(B1) == NUM) { Val bv = get_val(B1); @@ -1616,7 +1641,7 @@ __global__ void evaluator(GNet* gnet) { //__syncthreads(); //} //__syncthreads(); - + //printf("[%04x] span is %d\n", TID(), span); //__syncthreads(); } @@ -1631,7 +1656,7 @@ __global__ void evaluator(GNet* gnet) { // WORK MODE // --------- - + if (tm.mode == WORK) { u32 chkt = 0; u32 chka = 1; @@ -1742,7 +1767,7 @@ void gnet_normalize(GNet* gnet) { //printf("==================================================== "); //printf("TURN: %04x | RLEN: %04x | ITRS: %012llu\n", turn, rlen, itrs); //cudaDeviceSynchronize(); - + evaluator<<>>(gnet); inbetween<<<1, 1>>>(gnet); //cudaDeviceSynchronize(); @@ -1802,16 +1827,227 @@ Port gnet_expand(GNet* gnet, Port port) { Port gnet_make_node(GNet* gnet, Tag tag, Port fst, Port snd) { Port ret; Port* d_ret; - cudaMalloc(&d_ret, sizeof(Port)); + cudaMalloc(&d_ret, sizeof(Port)); make_node<<<1,1>>>(gnet, tag, fst, snd, d_ret); cudaMemcpy(&ret, d_ret, sizeof(Port), cudaMemcpyDeviceToHost); cudaFree(d_ret); return ret; } +// Readback +// -------- + +// Reads back a λ-Encoded constructor from device to host. +// Encoding: λt ((((t TAG) arg0) arg1) ...) +Ctr gnet_read_ctr(GNet* gnet, Port port) { + Ctr ctr; + ctr.tag = -1; + ctr.args_len = 0; + + // Loads root lambda + Port lam_port = gnet_expand(gnet, port); + if (get_tag(lam_port) != CON) return ctr; + Pair lam_node = gnet_node_load(gnet, get_val(lam_port)); + + // Loads first application + Port app_port = gnet_expand(gnet, get_fst(lam_node)); + if (get_tag(app_port) != CON) return ctr; + Pair app_node = gnet_node_load(gnet, get_val(app_port)); + + // Loads first argument (as the tag) + Port arg_port = gnet_expand(gnet, get_fst(app_node)); + if (get_tag(arg_port) != NUM) return ctr; + ctr.tag = get_u24(get_val(arg_port)); + + // Loads remaining arguments + while (TRUE) { + app_port = gnet_expand(gnet, get_snd(app_node)); + if (get_tag(app_port) != CON) break; + app_node = gnet_node_load(gnet, get_val(app_port)); + arg_port = gnet_expand(gnet, get_fst(app_node)); + ctr.args_buf[ctr.args_len++] = arg_port; + } + + return ctr; +} + +// 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. +// Encoding: +// - λt (t NIL) +// - λt (((t CONS) head) tail) +Str gnet_read_str(GNet* gnet, Port port) { + // Result + Str str; + str.text_len = 0; + + // Readback loop + while (TRUE) { + // Normalizes the net + gnet_normalize(gnet); + + // Reads the λ-Encoded Ctr + Ctr ctr = gnet_read_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 (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)); + port = ROOT; + continue; + } + } + break; + } + + str.text_buf[str.text_len] = '\0'; + + return str; +} + +// Primitive IO Fns +// ----------------- + +// IO: GetText +Port io_get_text(GNet* gnet, Port argm) { + printf("TODO\n"); + return new_port(ERA, 0); +} + +// IO: PutText +Port io_put_text(GNet* gnet, Port argm) { + // Converts argument to C string + Str str = gnet_read_str(gnet, argm); + // Prints it + printf("%s", str.text_buf); + // Returns result (in this case, just an eraser) + return new_port(ERA, 0); +} + +// IO: GetFile +Port io_get_file(GNet* gnet, Port argm) { + printf("TODO\n"); + return new_port(ERA, 0); +} + +// IO: PutFile +Port io_put_file(GNet* gnet, Port argm) { + printf("TODO\n"); + return new_port(ERA, 0); +} + +// IO: GetTime +Port io_get_time(GNet* gnet, Port argm) { + // Get the current time in nanoseconds + u64 time_ns = time64(); + // Encode the time as a 64-bit unsigned integer + u32 time_hi = (u32)(time_ns >> 24) & 0xFFFFFFF; + u32 time_lo = (u32)(time_ns & 0xFFFFFFF); + // Return the encoded time + return gnet_make_node(gnet, CON, new_port(NUM, new_u24(time_hi)), new_port(NUM, new_u24(time_lo))); +} + +// IO: PutTime +// NOTE: changing this name will corrupt the timeline. You've been warned. +Port io_put_time(GNet* gnet, Port argm) { + // 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))); + // Combine into a 48-bit duration in nanoseconds + u64 dur_ns = (((u64)dur_hi) << 24) | dur_lo; + // Sleep for the specified duration + struct timespec ts; + ts.tv_sec = dur_ns / 1000000000; + ts.tv_nsec = dur_ns % 1000000000; + nanosleep(&ts, NULL); + // Return an eraser + return new_port(ERA, 0); +} + +// Monadic IO Evaluator +// --------------------- + +// Runs an IO computation. +void do_run_io(GNet* gnet, Book* book, Port port) { + // IO loop + while (TRUE) { + // Normalizes the net + gnet_normalize(gnet); + + // Reads the λ-Encoded Ctr + Ctr ctr = gnet_read_ctr(gnet, gnet_peek(gnet, port)); + + // Checks if IO Magic Number is a CON + if (get_tag(ctr.args_buf[0]) != CON) { + break; + } + + // Checks the IO Magic Number + Pair io_magic = gnet_node_load(gnet, get_val(ctr.args_buf[0])); + //printf("%08x %08x\n", get_u24(get_val(get_fst(io_magic))), get_u24(get_val(get_snd(io_magic)))); + if (get_val(get_fst(io_magic)) != new_u24(IO_MAGIC_0) || get_val(get_snd(io_magic)) != new_u24(IO_MAGIC_1)) { + break; + } + + switch (ctr.tag) { + case IO_CALL: { + Str func = gnet_read_str(gnet, ctr.args_buf[1]); + FFn* ffn = NULL; + // FIXME: optimize this linear search + for (u32 fid = 0; fid < book->ffns_len; ++fid) { + if (strcmp(func.text_buf, book->ffns_buf[fid].name) == 0) { + ffn = &book->ffns_buf[fid]; + break; + } + } + if (ffn == NULL) { + printf("FOUND NOTHING when looking for %s\n", func.text_buf); + break; + } + + Port argm = ctr.args_buf[2]; + Port cont = ctr.args_buf[3]; + Port ret = ffn->func(gnet, argm); + + Port p = gnet_make_node(gnet, CON, ret, ROOT); + gnet_boot_redex(gnet, new_pair(p, cont)); + port = ROOT; + continue; + } + case IO_DONE: { + printf("DONE\n"); + break; + } + } + break; + } +} + // Book Loader // ----------- +// TODO: initialize ffns_len with the builtin ffns +void book_init(Book* book) { + book->ffns_len = 6; + book->ffns_buf[0] = (FFn){"GET_TEXT", io_get_text}; + book->ffns_buf[0] = (FFn){"PUT_TEXT", io_put_text}; + book->ffns_buf[2] = (FFn){"GET_FILE", io_get_file}; + book->ffns_buf[3] = (FFn){"PUT_FILE", io_put_file}; + book->ffns_buf[4] = (FFn){"GET_TIME", io_get_time}; + book->ffns_buf[5] = (FFn){"PUT_TIME", io_put_time}; +} + void book_load(Book* book, u32* buf) { // Reads defs_len book->defs_len = *buf++; @@ -1825,7 +2061,7 @@ void book_load(Book* book, u32* buf) { // Gets def Def* def = &book->defs_buf[fid]; - + // Reads name memcpy(def->name, buf, 256); buf += 64; @@ -1844,7 +2080,7 @@ void book_load(Book* book, u32* buf) { // Reads rbag_buf memcpy(def->rbag_buf, buf, 8*def->rbag_len); buf += def->rbag_len * 2; - + // Reads node_buf memcpy(def->node_buf, buf, 8*def->node_len); buf += def->node_len * 2; @@ -2219,26 +2455,27 @@ __global__ void print_result(GNet* gnet) { extern "C" void hvm_cu(u32* book_buffer) { // Start the timer clock_t start = clock(); - + // Loads the Book + Book* book = (Book*)malloc(sizeof(Book)); if (book_buffer) { - Book* book = (Book*)malloc(sizeof(Book)); + book_init(book); book_load(book, (u32*)book_buffer); cudaMemcpyToSymbol(BOOK, book, sizeof(Book)); - free(book); } // Configures Shared Memory Size cudaFuncSetAttribute(evaluator, cudaFuncAttributeMaxDynamicSharedMemorySize, sizeof(LNet)); - + // Creates a new GNet GNet* gnet = gnet_create(); // Boots root redex, to expand @main - gnet_boot_redex(gnet, new_pair(new_port(REF,0), ROOT)); + gnet_boot_redex(gnet, new_pair(new_port(REF, 0), ROOT)); + + // Normalizes and runs IO + do_run_io(gnet, book, ROOT); - // Normalizes the GNet - gnet_normalize(gnet); cudaDeviceSynchronize(); // Stops the timer