A page target: -o file.html builds the runtime as WebAssembly, a worker a core, and a Window on its canvas

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.

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);

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) {

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<MTLDevice> dev = layer.device;
   window_pipe(dev);
   id<MTLBuffer> 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<CAMetalDrawable> 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 <emscripten/threading.h>
+#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) {

bend2/effs/window_open.c

@@ -124,9 +124,6 @@ static id<MTLDevice> 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 <file> -o <out> 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 <file> -o <out> 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 <emscripten.h>
+typedef struct { u32 w; u32 h; u32* pix; u32 cap; u32* evs; u32 got; } BendWin;
+#endif
+
+// The page's <canvas id="bend"> 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);

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 <emscripten.h>
+
+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) {