diff --git a/README.md b/README.md index f0101b31f..fa0f4908b 100644 --- a/README.md +++ b/README.md @@ -229,8 +229,11 @@ def add_zero(x): - Base is small: expect to write helpers other languages ship built in. - Effects are few: print, env, time, sleep, spawn, channels, files, TCP, UDP. - No TLS, HTTP library, JSON or regex for now (but you can add them as foreigns). -- Targets are C, Metal, CUDA and JavaScript; Lua, Luau and Python are planned. -- The JavaScript target runs on one core and has no graphics or audio. +- Targets are C, Metal, CUDA, JavaScript and WebAssembly; Lua, Luau and Python + are planned. +- The JavaScript target runs on one core and has no graphics or audio. The + WebAssembly target (a page) runs on every core and has a window, but no + audio, files or sockets yet. - Parallelism requires balanced calls. Flexible parallelism will be added later. - Sharing arrays with atomics across threads is experimental and needs `@unsafe`. - One GPU per program, one event loop, and no multi-machine execution yet. @@ -240,7 +243,8 @@ def add_zero(x): - We don't have as many benchmarks as we'd like yet, especially for the checker. - The compiler (not kernel) is 99% AI-written and has not been fully audited yet. - The Lean formalization and bend.ts mismatch. Early consistency bugs may occur. -- A binary needs clang 14+; ! needs 19+, Metal or CUDA 12. +- A binary needs clang 14+; ! needs 19+, Metal or CUDA 12. A page needs + Emscripten 3.1.35+ and a server with the COOP and COEP headers. - No Windows (WSL works); on Linux, Window and Audio need X11 and ALSA headers. - The hub has no names, versions, accounts or search yet. Packages are hashes. - Error messages are terse; no debugger, profiler or REPL. diff --git a/bend2/comp.ts b/bend2/comp.ts index a8b10e669..918ff49a7 100644 --- a/bend2/comp.ts +++ b/bend2/comp.ts @@ -2955,9 +2955,9 @@ function compile_tables(fl: File, entries: Seg[]): string[] { `(V)[${j}] = ${r};`).join(" ")}`, "", `#define WL_TAKE(V) ${rs.slice(0, resw).map((r, j) => `${r} = (V)[${j}];`).join(" ")}`, "", - `#define WL_SIG Env e, Stk sp, u32 seq, u32 rn, ${ws.map((w) => - "Term " + w).join(", ")}`, "", `#define WL_ALL e, sp, seq, rn, ${ws - .join(", ")}`, "", + `#define WL_SIG Corpus e_mem, DEV u64* e_alc, Stk sp, u32 seq, u32 rn, ${ws + .map((w) => "Term " + w).join(", ")}`, "", + `#define WL_ALL e.mem, e.alc, sp, seq, rn, ${ws.join(", ")}`, "", `#define WL_TABLE ${entries.map((s) => `WL_X(${s.fid})`).join(" ")}` + " WL_X(FID_EXIT)"); return defs; @@ -3445,7 +3445,8 @@ using namespace metal; #define UNLOCK(l) __atomic_store_n(&(l), 0, __ATOMIC_RELEASE) #define WL_FN static PRESERVE(preserve_none) __attribute__((noinline)) Reply #define WL_CASE(F) WL_FN WL_##F(WL_SIG) -#define WL_OPEN { WL_BANK u32 rn; +// wasm cannot tail-call with a struct +#define WL_OPEN { Env e = { e_mem, e_alc }; WL_BANK u32 rn; #define WL_JMP(F) __attribute__((musttail)) return WL_##F(WL_ALL) #define WL_DYN(F) __attribute__((musttail)) return wl_tab[F](WL_ALL) #endif @@ -4764,7 +4765,7 @@ extern "C" __global__ void bend_dev(Corpus H, u32 pass) { // pixels itself. An Image is a quadtree over 2^k x 2^k: a Qua at level // i splits its square in four (tl, tr, bl, br), a Qua under the pixels // follows tl, a Pix is 0xRRGGBB. -#if defined(__linux__) || defined(BEND_RTC) +#ifndef __METAL_VERSION__ INLINE u32 window_pix(Corpus H, Term t, u32 k, u32 x, u32 y) { for (u32 i = k; term_tag(t) == TAG_CTR;) { @@ -4826,7 +4827,12 @@ static void row_grow(Env e, Stk stk, u32 base, u32 stride, u32 want) { // Pool // ==== +#define W32 (sizeof(void*) == 4) + static void* pool_try(void* at, u64 bytes) { +#ifdef __EMSCRIPTEN__ + return memalign(16384, bytes) ?: MAP_FAILED; // mmap would memset zeroes +#endif return mmap(at, bytes, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANON | MAP_NORESERVE, -1, 0); } @@ -4840,7 +4846,7 @@ static void* pool_mmap(u64 bytes) { } static Term* pool_stack(void) { - u64 len = 1ull << 31; + u64 len = 1ull << (W32 ? 24 : 31); char* p = pool_mmap(len + 16384 + SIGSTKSZ); if (mprotect(p + len, 16384, PROT_NONE) != 0) { err_fail("stack guard failed"); @@ -5282,12 +5288,12 @@ static void cube_run(Corpus H, bool gpu) { // The cores map 8 GiB at a high base and double it in place, a hint then // a check (MAP_FIXED would replace a neighbour), so one base holds every // Loc and a run pays for the room it reaches. The banks lie past the pages -// and move up at each step. The GPU maps its whole span once. +// and move up at each step. The GPU maps its whole span once, as does wasm32. static u64 corpus_size; static void* corpus_map(u64 size) { - u64 hint = 1ull << 45; + u64 hint = W32 ? 0 : 1ull << 45; void* p = pool_try((void*)hint, size); while (p != (void*)hint && hint > size) { if (p != MAP_FAILED) { @@ -5326,7 +5332,7 @@ static bool corpus_grow(Corpus H, u64 need) { while (ok && need > a32_load(a32_at(H, H_CAP))) { u64 more = corpus_size; char* at = (char*)H + more; - void* got = io_gpu || more >= 1ull << 43 ? MAP_FAILED + void* got = io_gpu || more >= 1ull << 43 || W32 ? MAP_FAILED : pool_try(at, more); ok = got == at; if (ok) { @@ -5342,7 +5348,7 @@ static bool corpus_grow(Corpus H, u64 need) { static Corpus corpus_setup(bool gpu, long threads, u64 bytes) { io_gpu = gpu; KEEP_WORDS = gpu ? CHUNK : CAP_WORDS; - u64 dflt = gpu ? gpu_span() : 1ull << 33; + u64 dflt = gpu ? gpu_span() : 1ull << (W32 ? 30 : 33); u64 size = (gpu && bytes != 0 ? bytes : dflt) & ~16383ull; CORPUS = gpu ? gpu_map(size) : corpus_map(size); Corpus H = CORPUS; @@ -5760,7 +5766,8 @@ static void io_wait(Env e) { top = (int)a->work.word; } } - u64 len = (u64)top / 64 * 8 + 8; + // a whole fd_set at least: Emscripten's select zeroes one + u64 len = (u64)(top > 1023 ? top : 1023) / 64 * 8 + 8; u8* set[2] = { io_mem(calloc(2, len)), NULL }; set[1] = set[0] + len; io_bit(set[0], io_wake_fd[0], true); diff --git a/bend2/effs/window_close.c b/bend2/effs/window_close.c index affde2de3..fe1562647 100644 --- a/bend2/effs/window_close.c +++ b/bend2/effs/window_close.c @@ -37,6 +37,19 @@ static void window_close(intptr_t at) { free(win); } +#elif defined(__EMSCRIPTEN__) +#ifndef BendWin +#define BendWin BendWin +typedef struct { u32 w; u32 h; u32* pix; u32 cap; u32* evs; u32 got; } BendWin; +#endif + +static void window_close(intptr_t at) { + BendWin* win = (BendWin*)at; + free(win->pix); + free(win->evs); + free(win); +} + #else static void window_close(intptr_t at) { diff --git a/bend2/effs/window_frame.c b/bend2/effs/window_frame.c index f35c0bf1e..4a9ffcc48 100644 --- a/bend2/effs/window_frame.c +++ b/bend2/effs/window_frame.c @@ -3,7 +3,7 @@ // An event is five words: kind (0 key, 1 mouse, 2 move, 3 close) and // its fields; a frame answers the events pumped since the last one. -#if defined(__OBJC__) || defined(__linux__) +#if defined(__OBJC__) || defined(__linux__) || defined(__EMSCRIPTEN__) static Term window_node(Env e, const u32* ev) { static const u32 cids[3] = { CID_KEY, CID_MOUSE, CID_MOVE }; @@ -140,11 +140,9 @@ static void window_show(Env e, CAMetalLayer* layer, Term image) { id dev = layer.device; window_pipe(dev); id buf = window_corpus(e, dev); - WinArgs args = { image, layer.drawableSize.width, layer.drawableSize.height, - 0 }; - while ((1u << args.k) < args.w || (1u << args.k) < args.h) { - args.k += 1; - } + u32 w = layer.drawableSize.width; + u32 h = layer.drawableSize.height; + WinArgs args = { image, w, h, window_k(w, h) }; window_pump(); @autoreleasepool { id d = [layer nextDrawable]; @@ -308,11 +306,7 @@ static void window_fill(Env e, u32* pix, u32 w, u32 h, Term image, u32 k) { return; } #endif - for (u32 y = 0; y < h; y += 1) { - for (u32 x = 0; x < w; x += 1) { - pix[y * w + x] = window_pix(e.mem, image, k, x, y); - } - } + window_host(e.mem, image, w, h, k, pix); } // A frame waits for the next 60 Hz tick, as the Mac's display sync. @@ -329,11 +323,7 @@ static void window_pace(void) { static void window_show(Env e, BendWin* win, Term image) { u32 w = win->img->width; u32 h = win->img->height; - u32 k = 0; - while ((1u << k) < w || (1u << k) < h) { - k += 1; - } - window_fill(e, (u32*)win->img->data, w, h, image, k); + window_fill(e, (u32*)win->img->data, w, h, image, window_k(w, h)); window_pace(); XPutImage(win->dpy, win->win, DefaultGC(win->dpy, DefaultScreen(win->dpy)), win->img, 0, 0, 0, 0, w, h); @@ -350,6 +340,30 @@ static Term window_frame(Env e, intptr_t at, Term image) { return list; } +#elif defined(__EMSCRIPTEN__) +#include +#ifndef BendWin +#define BendWin BendWin +typedef struct { u32 w; u32 h; u32* pix; u32 cap; u32* evs; u32 got; } BendWin; +#endif + +static Term window_frame(Env e, intptr_t at, Term image) { + BendWin* win = (BendWin*)at; + io_sync(); + window_host(e.mem, image, win->w, win->h, window_k(win->w, win->h), win->pix); + for (u32 i = 0; i < win->w * win->h; i += 1) { + u32 c = win->pix[i]; + win->pix[i] = 0xFF000000 | c >> 16 | (c & 0xFF00) | c << 16 & 0xFF0000; + } + a32_store(&win->got, 0); + MAIN_THREAD_ASYNC_EM_ASM({ window_js_show($0, $1, $2, $3, $4, $5); }, + win->pix, win->w, win->h, win->evs, win->cap, &win->got); + while (a32_load_acq(&win->got) == 0) { + emscripten_futex_wait(&win->got, 0, 1000); + } + return window_list(e, win->evs, win->got - 1); +} + #else static Term window_frame(Env e, intptr_t at, Term image) { diff --git a/bend2/effs/window_open.c b/bend2/effs/window_open.c index 5ac73521e..777cb1eaf 100644 --- a/bend2/effs/window_open.c +++ b/bend2/effs/window_open.c @@ -124,9 +124,6 @@ static id window_dev; static u32 window_make(const char* title, u32 w, u32 h, intptr_t* out, const char** why) { - if (w < 1 || h < 1 || w > 16384 || h > 16384) { - return EINVAL; - } if (NSScreen.screens.count == 0) { *why = "Window.open: no display (build a native binary with bend -o and run it from a desktop session)"; return ENOTSUP; @@ -201,9 +198,6 @@ typedef struct { static u32 window_make(const char* title, u32 w, u32 h, intptr_t* out, const char** why) { - if (w < 1 || h < 1 || w > 16384 || h > 16384) { - return EINVAL; - } Display* dpy = XOpenDisplay(NULL); if (dpy == NULL) { *why = "Window.open: no display (build a native binary with bend -o and run it from a desktop session)"; @@ -236,6 +230,108 @@ static u32 window_make(const char* title, u32 w, u32 h, intptr_t* out, return 0; } +#elif defined(__EMSCRIPTEN__) +#ifndef BendWin +#define BendWin BendWin +#include +typedef struct { u32 w; u32 h; u32* pix; u32 cap; u32* evs; u32 got; } BendWin; +#endif + +// The page's and its listeners, on the main thread: the +// program runs on a worker. Events are the Mac's five words, its key codes +// (a character in lower case, a function key's private-use character, +// 65536 + a modifier's key code) and buttons (0 left, 1 right, 2 middle). +EM_JS(void, window_js_open, (const char* title, u32 w, u32 h), { + var c = document.getElementById("bend"); + c.width = w; + c.height = h; + Module.bendCtx = c.getContext("2d"); + Module.bendImg = new ImageData(w, h); + document.title = UTF8ToString(title); + var evs = Module.bendEvs; + if (!evs) { + evs = Module.bendEvs = []; + var put = function(k, a, b, c, d) { + if (evs.length < 5120) { + evs.push(k, a, b, c, d); + } + }; + var keys = { Escape: 27, Enter: 13, Tab: 9, Backspace: 127, + ArrowUp: 63232, ArrowDown: 63233, ArrowLeft: 63234, ArrowRight: 63235, + Insert: 63271, Delete: 63272, Home: 63273, End: 63275, PageUp: 63276, + PageDown: 63277, MetaRight: 65590, MetaLeft: 65591, ShiftLeft: 65592, + CapsLock: 65593, AltLeft: 65594, ControlLeft: 65595, ShiftRight: 65596, + AltRight: 65597, ControlRight: 65598 }; + var key = function(ev, down) { + var k = ev.key; + var f = /^F([0-9]+)$/.exec(k); + var code = keys[ev.code] || keys[k] || (f ? 63235 + Number(f[1]) + : k.length === 1 ? k.toLowerCase().codePointAt(0) : 65536 + ev.keyCode); + put(0, code, down, 0, 0); + if (!ev.metaKey && !ev.ctrlKey && !f) { + ev.preventDefault(); + } + }; + var at = function(ev) { + var r = c.getBoundingClientRect(); + var x = Math.floor((ev.clientX - r.left) * c.width / r.width); + var y = Math.floor((ev.clientY - r.top) * c.height / r.height); + return [Math.max(0, Math.min(x, c.width - 1)), + Math.max(0, Math.min(y, c.height - 1))]; + }; + var mouse = function(ev, down) { + var p = at(ev); + if (ev.button < 3) { + put(1, p[0], p[1], [0, 2, 1][ev.button], down); + } + }; + window.addEventListener("keydown", function(ev) { key(ev, 1); }); + window.addEventListener("keyup", function(ev) { key(ev, 0); }); + c.addEventListener("mousedown", function(ev) { mouse(ev, 1); }); + window.addEventListener("mouseup", function(ev) { mouse(ev, 0); }); + c.addEventListener("mousemove", function(ev) { + var p = at(ev); + put(2, p[0], p[1], 0, 0); + }); + c.addEventListener("contextmenu", function(ev) { ev.preventDefault(); }); + } + evs.length = 0; +}); + +// A frame on the display's next tick, as the Mac's display sync (a hidden +// tab has no ticks) +EM_JS(void, window_js_show, (u32* pix, u32 w, u32 h, u32* evs, u32 cap, + u32* got), { + var show = function() { + Module.bendImg.data.set(HEAPU8.subarray(pix, pix + w * h * 4)); + Module.bendCtx.putImageData(Module.bendImg, 0, 0); + Module.bendFrames = (Module.bendFrames | 0) + 1; + var q = Module.bendEvs; + var n = Math.min(q.length / 5, cap); + HEAPU32.set(q.splice(0, n * 5), evs >> 2); + Atomics.store(HEAP32, got >> 2, n + 1); + Atomics.notify(HEAP32, got >> 2); + }; + if (document.hidden) { + setTimeout(show, 16); + } else { + requestAnimationFrame(show); + } +}); + +static u32 window_make(const char* title, u32 w, u32 h, intptr_t* out, + const char** why) { + MAIN_THREAD_EM_ASM({ window_js_open($0, $1, $2); }, title, w, h); + BendWin* win = io_mem(calloc(1, sizeof *win)); + win->w = w; + win->h = h; + win->pix = io_mem(calloc((u64)w * h, 4)); + win->cap = 1024; + win->evs = io_mem(calloc(win->cap * 5, 4)); + *out = (intptr_t)win; + return 0; +} + #else static u32 window_make(const char* title, u32 w, u32 h, intptr_t* out, @@ -246,13 +342,30 @@ static u32 window_make(const char* title, u32 w, u32 h, intptr_t* out, #endif +#ifndef __METAL_VERSION__ +// the quadtree level that holds w x h, and a frame filled by the host +INLINE u32 window_k(u32 w, u32 h) { + u32 m = (w > h ? w : h) - 1; + return m ? 32 - CLZ(m) : 0; +} + +INLINE void window_host(Corpus H, Term root, u32 w, u32 h, u32 k, u32* out) { + for (u32 i = 0; i < w * h; i += 1) { + out[i] = window_pix(H, root, k, i % w, i / w); + } +} +#endif + Term window_open_run(Env e, Term* f, IoWork* w) { uint64_t n = 0; char* title = io_cstr(e, f[0], &n); intptr_t out; const char* why = NULL; + u32 wd = (u32)f[1]; + u32 ht = (u32)f[2]; u32 q = io_nul(title, n) ? EILSEQ - : window_make(title, (u32)f[1], (u32)f[2], &out, &why); + : wd < 1 || ht < 1 || wd > 16384 || ht > 16384 ? EINVAL + : window_make(title, wd, ht, &out, &why); free(title); if (q != 0) { return io_fail(e, q, why); diff --git a/bend2/effs/window_set_title.c b/bend2/effs/window_set_title.c index 91f3c7e3b..25edbd16d 100644 --- a/bend2/effs/window_set_title.c +++ b/bend2/effs/window_set_title.c @@ -36,6 +36,14 @@ static void window_set_title(intptr_t at, const char* text, u64 n) { XFlush(win->dpy); } +#elif defined(__EMSCRIPTEN__) +#include + +static void window_set_title(intptr_t at, const char* text, u64 n) { + MAIN_THREAD_EM_ASM({ document.title = UTF8ToString($0, $1); }, text, + (u32)n); +} + #else static void window_set_title(intptr_t at, const char* text, u64 n) { diff --git a/bend2/main.ts b/bend2/main.ts index bfd3fa862..57d0a28b9 100755 --- a/bend2/main.ts +++ b/bend2/main.ts @@ -35,7 +35,8 @@ const HELP = `Bend ${VERSION}: check, run, build and publish Bend programs. usage: bend [args] check the file, then run main with args - bend -o build a binary; .c emits C, .js JS + bend -o build a binary; .c emits C, .js JS, + .html a web page (WebAssembly; emcc) bend --check-only check the file and its imports; run nothing bend --publish publish the file and its imports to the hub bend -o bundle a page that imports .bend files @@ -47,6 +48,56 @@ usage: Read the guide (\`bend guide\`) before writing Bend code. `; +// PAGE is the web page a build writes beside its .js and .wasm: the canvas a +// Window draws on, a thread count to pick, its frames per second, a line per +// print. +const PAGE = ` + +NAME + + +
 threads,
+

+
+
+`;
+
 const BASE = Bend.BASE_BEND;
 
 const GUIDE = path.join(Bend.BEND_DIR, "..", "guide");
@@ -301,7 +352,11 @@ function cli_emit(book: Bend.Book, out: string): void {
     const c   = path.join(dir, path.basename(out) + ".c");
     fs.writeFileSync(c, Comp.compile_book(book));
     try {
-      cli_build(out, c);
+      if (out.endsWith(".html")) {
+        cli_build_web(out, c);
+      } else {
+        cli_build(out, c);
+      }
     } finally {
       fs.rmSync(dir, { recursive: true, force: true });
     }
@@ -373,6 +428,25 @@ function cli_build(bin: string, file: string): void {
   }
 }
 
+// cli_build_web builds the C file at `file` into the page `page` and its .js
+// and .wasm: emcc 3.1.35+ (tail calls), the program on a worker a core (the
+// pool holds 4 more: the proxied main and the IO helpers), 2 GiB of memory.
+function cli_build_web(page: string, file: string): void {
+  const base = page.slice(0, -".html".length);
+  const emcc = process.env.EMCC || "emcc";
+  const args = ["-std=gnu11", "-O3", "-pthread", "-mtail-call", file,
+    "-sPROXY_TO_PTHREAD", "-sPTHREAD_POOL_SIZE=navigator.hardwareConcurrency+4",
+    "-sINITIAL_MEMORY=2147483648", "-sENVIRONMENT=web,worker",
+    "-sEXIT_RUNTIME=1", "-o", path.resolve(base + ".js")];
+  if (child.spawnSync(emcc, args, { stdio: "inherit" }).status !== 0) {
+    throw "Error: " + emcc + " failed to build " + page
+      + " (a page needs Emscripten 3.1.35+ on PATH, or at $EMCC)";
+  }
+  const name = path.basename(base).replace(/[&<"]/g, (c) =>
+    "&#" + c.charCodeAt(0) + ";");
+  fs.writeFileSync(page, PAGE.replaceAll("NAME", name));
+}
+
 // cli_base prints the base library; with --types, its type declarations
 // (every `type`, and every law whose result is a kind); with a name, the
 // blocks declaring it or a name under it (its law, its def, its @unsafe).
diff --git a/guide/GUIDE.md b/guide/GUIDE.md
index 05e49606a..6322f7cae 100644
--- a/guide/GUIDE.md
+++ b/guide/GUIDE.md
@@ -164,7 +164,8 @@ machine without a GPU runs `!` on the CPU (still in parallel). What the lanes
 share also sets the speed: a `+` value read by every lane costs an atomic per
 read. Read `bend guide shaders` before you write a parallel app.
 
-The JavaScript target ignores all that and just runs sequentially.
+The JavaScript target ignores all that and just runs sequentially; a web
+page built with `-o file.html` runs it on every core, as a binary does.
 
 ### Arrays
 
@@ -496,6 +497,7 @@ bend file.bend            # check; run main (IO compiled; a value normalized)
 bend file.bend -o file    # compile to a native binary (clang 14+; 19+ with `!`)
 bend file.bend -o file.c  # emit the C source instead
 bend file.bend -o file.js # emit the JS source instead
+bend file.bend -o file.html # build a web page (WebAssembly on every core; emcc)
 bend page.html -o dist    # bundle a web page that imports .bend files
 ./file --threads 8        # run a native binary on 8 CPU threads
 ./file --gpu off          # run ! calls on the CPU (the GPU is on by default)
@@ -507,7 +509,12 @@ by the checker (slow for big work) and printed; a file with no `main` just
 checks. A binary that uses `!` builds its GPU program too, as `file.gpu`, which
 must stay beside it: on macOS it needs Metal, on Linux CUDA 12 at
 `/usr/local/cuda`. On Linux a program with a Window needs `libx11-dev`, one
-with Audio `libasound2-dev`. `bend guide` prints this text, `bend base` prints
+with Audio `libasound2-dev`. A page (`-o file.html`, its .js and .wasm
+beside it) is the runtime as WebAssembly on a worker per core (a `!` runs on
+them) with a Window on its canvas; it needs Emscripten 3.1.35+ and a server
+sending `Cross-Origin-Opener-Policy: same-origin` and
+`Cross-Origin-Embedder-Policy: require-corp`, since threads need cross-origin
+isolation. `bend guide` prints this text, `bend base` prints
 the Base library (`bend base Map` prints one name and everything under it), and
 `bend --help` lists the other commands.