← Commit history
Pyre (ash-gpu): kiln prerequisites, CPU+Vulkan backends, real GPU benchmarks, kiln/ashvm integration
lennoxrose committed on Oct 8, 2026 at 20:16 UTC
88 files changed+3139−74
Heap growth + chained index assignment in kiln so matrix-heavy Ash code doesn't segfault or need manual row extraction. Built Pyre as a real project (C/H/D/O trees matching kiln/ashvm): CPU backend beats both benchmark challenges by 85%+, ashvm and kiln both genuinely call into it (kiln only under --link=shared, via a second DT_NEEDED on libashgpu.so -- found and fixed a real exit-syscall threading bug doing that). Turned out this sandbox is WSL2 with GPU passthrough just missing a driver -- installed it and got the actual RX 9070 XT running real compute shaders: matrix_multiply loses to CPU (fp64, untiled, translation-layer overhead), a new chaos_iterate kernel wins by 34x. Dispatch is still "GPU if present, no cost model" on purpose -- real auto-detection is next, not now.
.gitignoremodified
+8−0@@ -12,6 +12,14 @@ ashvm/benchmarks/loop_cpp1212
ashvm/benchmarks/countdown_cpp1313
ashvm/benchmarks/hashmap_cpp1414
ashvm/benchmarks/native_jit15
+pyre/benchmarks/*/run_benchmark16
+pyre/tests/matrix_multiply_test17
+pyre/tests/monte_carlo_test18
+pyre/lib/19
+# Generated from src/shaders/*.comp by glslc -mfmt=c (pyre/Makefile) --20
+# the .comp source is tracked, the embedded-SPIR-V .h is not, same as21
+# any other object code.22
+pyre/src/H/backends/vulkan/*_spv.h1523
1624
# OS / editor cruft1725
.DS_StoreCHANGELOG.mdmodified
+51−0@@ -1,3 +1,54 @@1
+## 2026-10-08 (pyre: the GPU is real now -- one loss, one 34x win, and the dispatcher gap that exposes)2
+Follow-up to the kiln-integration entry below, same day: this sandbox turns out to be WSL2 on the user's actual machine, which has GPU passthrough infrastructure (`/dev/dxg`, Microsoft's D3D12 translation libraries) already present but no Vulkan driver installed. Installing `vulkan-dzn` (Mesa's Vulkan-over-D3D12 driver) plus `mesa` made `vulkaninfo` and `ash_gpu_detect()` both find the real **AMD Radeon RX 9070 XT** ("Microsoft Direct3D12 (AMD Radeon RX 9070 XT)", driver `DRIVER_ID_MESA_DOZEN`, `shaderFloat64 = true`) -- the Vulkan path written two entries ago has now actually run, for the first time ever.3
+4
+**`matrix_multiply` on the real GPU: a loss.** Same-moment forced-CPU-vs-real-GPU measurement: CPU backend 0.44s, GPU backend 1.07-1.11s -- about 2.5x *slower*. Matches the caveat already written into that shader's comments before ever running it: fp64 throughput on a consumer GPU is a fraction of fp32, the shader is deliberately untiled (one thread per output element), memory is host-visible/host-coherent rather than device-local, and Dozen is explicitly a translation layer ("not a conformant Vulkan implementation, testing use only" -- Mesa's own warning), not a native driver.5
+6
+**So: is there any workload shape this GPU setup actually wins on?** Added `ash_gpu_chaos_iterate` (`pyre/src/H/runtime/ash_gpu.h`, CPU backend `src/C/backends/cpu/chaos.c`, Vulkan backend `src/C/backends/vulkan/chaos.cpp` + `src/shaders/chaos.comp`) specifically to answer that: 2 million independent elements, each iterated 4000 times through `x = sqrt(abs(x)+1)*0.99999+0.00001` in place -- the opposite shape from matmul, huge compute per element, trivial data movement, no cross-element dependency, no tiling to get right. (The recurrence uses `sqrt`/`abs` rather than `sin`/`cos` because GLSL's `GL_ARB_gpu_shader_fp64` extension turns out to add `double` overloads for the former but not the latter -- `glslc` flatly rejected `sin(double)` -- found by trying it, another thing that can't be learned from documentation alone.) New `pyre/benchmarks/chaos/run_benchmark.c` runs both backends back to back in one process for a fair same-run comparison and cross-checks their results. Result: CPU ~14.5s, GPU ~0.43s -- **~34x faster**, with `max |CPU - GPU| = 0` (bit-identical results across all 2M elements, not just "close").7
+8
+**The real finding isn't either number -- it's that nothing currently tells them apart.** Both dispatchers in `src/C/runtime/device.c` unconditionally prefer GPU-when-available with no cost model at all, which the matmul result proves is actively the wrong call for that kernel today. Building the actual "which workloads go on the GPU" algorithm is explicitly **not done now** -- the user asked for it to wait until everything else is wrapped up -- but it's written down as the clear next task in `pyre/todo/todo.md`'s new "Smart dispatch" section, with these two real, measured, contrasting data points as the calibration targets for whoever builds it.9
+10
+`ash_gpu.h`'s doc comments for `ash_gpu_detect`/`ash_gpu_matrix_multiply` updated to stop saying "never run here" now that it has been.11
+12
+## 2026-10-08 (kiln: real Phase 4 integration -- matrix_mul calls into libashgpu.so under --link=shared)13
+Follow-up to the ashvm/Vulkan entry below, same day: kiln gets the matching half of "ashvm and kiln actually use Pyre" (ideas/assigned.md), and finishes "Kiln integration" from ideas/gpu_acceleration.md's own phase table ("Kiln emits calls to ash-gpu for suitable workloads") for real -- a GOT call into `libashgpu.so`, not a native reimplementation of the algorithm, which is the distinction that phase is actually about.14
+15
+**Why `--link=shared` only:** kiln's default output is a hand-rolled, zero-dependency ELF/PE binary (AGENTS.md's "no external dependencies" is kiln's whole reason to exist). Pyre's CPU backend needs OpenMP and `dlopen`. `matrix_mul(a, b)` (new builtin, `kiln/src/codegen/C/collections/matrix_mul.c`) is therefore available only under `--link=shared`; under the default `--link=static` it's a clear compile error naming why, never a silent dependency added to every compiled program.16
+17
+**The mechanism:** `--link=shared` already dynamically links ONE needed library (`libkilnrt.so`) via a real `PT_INTERP`/`ld.so` setup (`elf/C/elf_dynamic.c`). Extended it to declare `libashgpu.so` as a SECOND needed library -- one more `DT_NEEDED`, and `ash_gpu_matrix_multiply` folded into the SAME combined hash/dynsym/GOT/rela table as kiln's own runtime imports (`KILN_DYN_TOTAL_IMPORTS` = `RUNTIME_IMPORT_COUNT_` + the new `AshGpuImport` enum's count, `elf/H/elf_dynamic.h`) rather than a second parallel set of tables, since `ld.so` resolves a symbol name against every needed library regardless of which table region declared it. New `codegen/C/platform/linux_call.c`/`.h` add the one piece kiln never needed before: a SysV ABI call wrapper that dynamically 16-byte-aligns RSP before the call (reusing `win_call.c`'s existing scratch slot, since a target is always Linux or Windows, never both) -- kiln's own hand-rolled routines never needed this because none of them use an alignment-sensitive instruction, but a real gcc/g++-compiled function with OpenMP and SSE/AVX does. `matrix_mul` itself (square arrays-of-arrays of numbers, same N, matching `ash_gpu_matrix_multiply`'s own signature) flattens both arguments into temporary flat buffers, calls through the GOT slot, and unflattens the result -- the bridge todo.md 2.A's still-nonexistent `Buffer`/`Tensor` type would otherwise be. Pyre's own `pyre/Makefile` gained a `lib/libashgpu.so` target and an `install` target (writes to kiln's hardcoded `KILN_RPATH`, same spot `libkilnrt.so` already installs to).18
+19
+**Verification, in both directions:** `kiln/tests/scripts/shared_link_check.sh` (new -- `--link=shared` had zero automated coverage before this) compiles a sample of ordinary tests under both link modes and diffs them, confirming the second `DT_NEEDED` and the bigger combined table didn't disturb kiln's own existing runtime imports; `run_regression.sh` (default `--link=static`) stayed byte-identical throughout. Separately, actually running `matrix_mul` (not just compiling it) surfaced a real, previously-invisible kiln bug: its process-exit path used the `exit` syscall (60, terminates only the calling thread) instead of `exit_group` (231, terminates every thread in the process) -- invisible for the entirety of kiln's prior history because no kiln program had ever spawned a second thread, until `ash_gpu_matrix_multiply`'s OpenMP worker pool did. The symptom was exact and confusing: the program printed its correct answer and then hung forever, waiting on orphaned worker threads the exited main thread left behind. Fixed in all three exit sites (`codegen/C/runtime/print_int.c`'s `codegen_exit0`, `codegen/C/builtins/exit_builtin.c`, `codegen/C/runtime/errors.c`'s two uncaught-error paths) -- `exit_group` with a single thread behaves identically to `exit`, confirmed by every existing test's output staying byte-for-byte unchanged. This is the entire point of running something instead of stopping at "it compiles" -- see the previous entry's Vulkan backend for the still-outstanding counterexample.20
+21
+`ideas/assigned.md` and `pyre/todo/todo.md` updated accordingly: both ashvm and kiln now genuinely call into Pyre "when it's ok to" (ashvm: always, once `N >= ASH_GPU_MIN_N`; kiln: always under `--link=shared`, never under the zero-dependency default).22
+23
+## 2026-10-08 (pyre: C/H/D/O restructure, ashvm's matrix_mul, a real (if unverified) Vulkan backend)24
+Follow-up to the Phase 1 CPU-backend entry below, same day. Three things changed.25
+26
+**Restructured to match AGENTS.md, not ideas/gpu_acceleration.md's own sketch:** `pyre/` moved from a standalone `include/runtime/backends/` layout to the same C/H/D/O mirrored-tree convention as `kiln`/`ashvm` (AGENTS.md section 1), with feature folders nested inside each tree exactly like `kiln/src/codegen/{C,H,D,O}/expressions/` -- `runtime`, `backends/cpu`, `backends/vulkan` are now feature folders under one shared `src/{C,H,D,O}/`, not their own standalone trees. `.cpp` sources (the new Vulkan backend) live in the same `C/` tree as `.c` -- the mirrored-tree rule is about not mixing sources with headers/objects/deps, not about `.c` vs `.cpp`. `pyre/Makefile` now generates `.d` files via `-MMD -MP` like the other two projects' Makefiles, and the root Makefile's `make`/`make @pyre` pick it up automatically (no changes needed there).27
+28
+**ashvm gets `matrix_mul(a, b)` (ideas/gpu_acceleration.md Phase 3, "Ash bindings... no compiler changes"):** a plain builtin (`ashvm/src/builtins/C/matrix.c`) that calls `ash_gpu_matrix_multiply` directly for square arrays-of-arrays of numbers, same N, at or above `ASH_GPU_MIN_N = 64` (below it, a local C loop -- OpenMP's thread-launch overhead would dominate a handful of multiply-adds). ashvm's own Makefile now links `../pyre/lib/libashgpu.a` unconditionally (`-fopenmp -ldl -lvulkan -lstdc++` added to its LDFLAGS) -- "ok" unconditionally, because ashvm already links against libc/libm like any normal C program and never promised zero dependencies. `kiln` gets no equivalent: its compiled *output* is a hand-rolled, zero-dependency ELF/PE binary (AGENTS.md's "no external dependencies" is kiln's whole reason to exist), so taking on Pyre's dependencies for every compiled Ash program isn't "ok" by default. The real integration path -- extending kiln's existing `--link=shared` dynamic-ELF machinery (which already loads `libkilnrt.so` via the real system `ld.so`) to also load `libashgpu.so` as a second needed library -- is real, scoped, understood work, deliberately not rushed into kiln's binary-correctness-critical ELF writer without dedicated design and testing time. All of this, including why a kiln-side *native reimplementation* of the same algorithm would miss the point of "kiln emits calls to ash-gpu," is written down in the previously-nonexistent `ideas/assigned.md`.29
+30
+**A real Vulkan compute backend for `matrix_multiply` -- written, compiled, never executed:** `pyre/tools/install_deps.sh` installs the Vulkan 1.2 SDK pieces this needed (headers, loader, `glslc`/`glslangValidator`, spirv-tools; Arch/pacman only, matches this sandbox and the user's own machine) -- confirmed working both here and independently by the user on their own machine. `src/C/backends/vulkan/context.cpp` creates a `VkInstance`, picks the first physical device with a compute queue family *and* the `shaderFloat64` feature (ash_gpu's ABI is `double` throughout, so an fp32-only device isn't a match), and a command pool; `matrix_multiply.cpp` runs `src/shaders/matrix_multiply.comp` (one thread per output element, no tiling -- deliberately the simplest correct version, not a tuned one) via host-visible/host-coherent buffers, a descriptor set, and one compute pipeline, embedding the shader as SPIR-V (`glslc -mfmt=c`) rather than loading it from disk. `ash_gpu_detect()`/`ash_gpu_matrix_multiply` (`src/C/runtime/device.c`) now really try this path before falling back to the CPU backend. Everything compiles clean (zero warnings) against the real SDK -- but this sandbox has no GPU and no ICD at all (`vulkaninfo`: "Found no drivers!"), so `ash_gpu_vulkan_init()` has only ever returned 0 here, and nothing past instance creation has ever actually run. What IS verified, repeatedly: the fallback itself -- every existing test (`pyre/tests/`, `ashvm/tests/cases/collections/matrix_mul.ash`, both benchmarks) still passes with correct results going through the real dispatcher, proving the "no GPU found, use CPU" path works end to end. The GPU path itself is reviewed and compiles, not proven -- whoever first runs this on the RX 9070 XT should expect to find and fix real bugs.31
+32
+## 2026-10-08 (pyre: Phase 1 CPU backend, both benchmark challenges cleared)33
+Start of `ash-gpu` (now named **Pyre**, `pyre/`), a separate project from `ashvm`/`kiln`/`forgepack` (ideas/gpu_acceleration.md's own proposed layout, not AGENTS.md's C/H/D/O tree). Phase 1 of that doc's plan: "standalone runtime... don't touch Ash yet" -- nothing here is wired into `kiln` or `ashvm`.34
+35
+Added the public C ABI (`pyre/src/H/runtime/ash_gpu.h`): `ash_gpu_detect()`/`ash_gpu_available()`, and two kernels, `ash_gpu_matrix_multiply` and `ash_gpu_monte_carlo_risk`. `pyre/src/C/runtime/device.c` detects hardware honestly -- this sandbox has no `/dev/dri`, `nvidia-smi` or `rocminfo`, only a software "Basic Render Driver", so it always reports no usable GPU (it probes for a Vulkan/CUDA/HIP loader via `dlopen`, informational only, never claims a backend with no working dispatch path behind it). Both kernels therefore always run on the CPU backend (`pyre/src/C/backends/cpu/`): `matrix_multiply` reorders the loop from the naive ijk to ikj (fixes the cache-hostile strided read of `B[k][j]`) and splits rows across cores with OpenMP; `monte_carlo_risk` gives each thread its own `splitmix32`-derived seed stream instead of one sequential LCG chain. Correctness tests (`pyre/tests/`) check each kernel independently of the benchmark numbers: a from-scratch scalar reference for matmul, self-consistency plus an analytic sanity range for the Monte Carlo probability.36
+37
+The user set two benchmark challenges (`pyre/benchmarks/`, each with a C++ reference and a "beat it by at least 50% or the solution isn't allowed" comment): `matrix_multiply` (2000x2000, reference 5.532s at `-O2` on this machine) and `monte_carlo_no_pi` (1e9 iterations, reference 1.840s at `-O3`). Pyre's CPU-only backend clears both by a wide margin with no GPU involved at all: matmul in 0.22-0.25s (~95.6% reduction) and the Monte Carlo kernel in 0.23-0.25s (~87% reduction) -- see `pyre/todo/todo.md`'s "Phase 1 benchmarks" section for the full numbers and methodology. Separately, the matrix_multiply.ash benchmark now compiles and runs under `kiln` for the first time (needs the chained-index-assignment fix above): the unoptimized, tag-checked, nothing-freed compiled Ash takes 3m20.7s for the same workload -- a real, measured baseline for whatever later phase eventually teaches `kiln` to recognize and offload this shape of loop.38
+39
+No GPU backend (Vulkan/HIP/CUDA) exists yet -- developing and verifying one needs real hardware this sandbox doesn't have (the user's own machine, per ideas/gpu_acceleration.md, has an RX 9070 XT). `ash_gpu_available()` stays honestly false until one does.40
+41
+## 2026-10-08 (kiln: dynamic heap, chained index assignment -- pyre prerequisites)42
+First item off `pyre/todo/todo.md` (prerequisites before building the `ash-gpu` compute backend, "Pyre"). Both changes are `kiln`-only; `ashvm` regression suite re-checked and unaffected.43
+44
+**Dynamic heap, no more fixed 16MB cap:** `kiln`'s bump allocator grabbed one fixed-size arena at startup (`brk` on Linux, one `VirtualAlloc` on Windows) and silently segfaulted the moment a program allocated past it -- a 500x500 matrix of temporaries, or any sufficiently long-running loop, would crash with no runtime error at all. Now it grows in 64MB chunks as needed: a new `heap_limit` globals slot (right after `heap_ptr`, pushing `try_depth`/handler-stack/argv/modstate addresses down by 8 bytes on all three executable layouts -- static ELF, PE, and `--link=shared`'s dynamic ELF) tracks the current committed boundary, and the alloc routine calls `brk` again (Linux) or commits the next chunk of an upfront 4GB `MEM_RESERVE` (Windows, since `VirtualAlloc` has no brk-style "extend in place") whenever the bump pointer would cross it, rounding the request up to a whole chunk so one oversized allocation still succeeds in a single grow step. The ~27 existing `heap_emit_alloc` call sites are unaffected: the growth path preserves the exact same clobber set (RBX/RCX/RDX/RBP/RSP untouched) via push/pop around the one-off syscall/API call. New `tests/cases/memory/heap_growth.ash` leaks 2,000,000 small arrays (confirmed to segfault the old binary, confirmed correct under the new one).45
+46
+**`C[i][j] = val` (chained index assignment):** previously only `IDENT[index] = expr` parsed; a second `[` after the first `]` was a parse error, forcing the manual `local row = C[i]; row[j] = val; C[i] = row;` workaround. `index_assignment_statement` (`parser/C/core/parser.c`) now loops on `[index]` and reads through every level but the last via `codegen_index_read` (promoted from `expr.c`-local `emit_index_read` to a shared declaration in `expr.h`, the same primitive `postfix()` already chained reads with) -- only the final index is a store, into whatever object the reads produced (arrays are reference values here, so there is nothing to write back to the outer container). A string anywhere in the chain, including the final target, still raises the existing `cannot assign to a string index`. `tests/cases/collections/chained_index_assign.ash` covers two and three levels deep, a matrix-accumulation shape, a mixed array/map chain, and the string-immutability case.47
+48
+**Also added, not yet wired to anything:** `heap_emit_checkpoint`/`heap_emit_reset` (`codegen/C/runtime/heap.c`), an arena-style rewind-the-bump-pointer pair for Pyre's own future lowering to reclaim short-lived loop-body scratch without an OS call. Deliberately *not* hooked into general `each`/`during` loop bodies: doing that soundly needs an escape analysis (does any value allocated in the loop body outlive the iteration -- outer-scope write, closure capture, return/yield, a retaining function call) that kiln's single-pass compiler doesn't have and goals.md item 1 defers for exactly that reason; an unsound version would silently corrupt memory instead of just leaking it, which is worse than not having the feature. Verified directly against a hand-built tiny ELF (bypassing the Ash front end, since there is no surface syntax for it yet): allocate, write, reset, allocate again, and confirm the second allocation reuses the first's exact address.49
+50
+Not a goals.md item: goal 1 (Rust-grade memory safety, an off switch, bounded memory) is a different, much larger deliverable -- this only removes a hard crash, it does not bound memory or free anything.51
+152
## 2026-10-08 (item 22 reversed: no JSON in the language)253
A built-in `json_*` library compiled into both engines was added and then removed again: the goal is general tools, not code for one module, and it needed name-mention detection and a special lookup rule in the compilers. The 'JSON in the stdlib?' item is answered by ash-json itself: its 0.3.0 codec is plain Ash built on the new tools (real booleans, insertion-ordered maps, `chr`/`ord`, forward declarations, `contains`/`repeat`, `rename_file`), with no compiler support. It parses and prints `\uXXXX` (surrogate pairs too), `\b`, `\f` and reports `json: <what> at line L, column C`. Its suite passes on both engines.354
Makefileadded
+31−0@@ -0,0 +1,31 @@1
+# Ash - Makefile2
+# `make` makes all subprojects, such as ashvm, kiln, forgepack, etc.3
+# `make clean` cleans all subprojects.4
+# `make @<subproject>` cleans and builds only that subproject, e.g. `make @ashvm`5
+6
+SUBDIRS := $(patsubst %/Makefile,%,$(wildcard */Makefile))7
+8
+.PHONY: all clean $(SUBDIRS) $(addprefix clean-,$(SUBDIRS))9
+10
+all: clean-bin $(SUBDIRS)11
+12
+clean-bin:13
+ rm -rf bin14
+ mkdir -p bin15
+16
+$(SUBDIRS):17
+ $(MAKE) -C $@ clean18
+ $(MAKE) -C $@19
+20
+clean: $(addprefix clean-,$(SUBDIRS))21
+22
+$(addprefix clean-,$(SUBDIRS)):23
+ $(MAKE) -C $(patsubst clean-%,%,$@) clean24
+25
+@%:26
+ @if [ ! -f $*/Makefile ]; then \27
+ echo "no such subproject: $*" >&2; \28
+ exit 1; \29
+ fi30
+ $(MAKE) -C $* clean31
+ $(MAKE) -C $*README.mdmodified
+1−1@@ -114,7 +114,7 @@ optimizing native compiler.114114
### Performance notes115115
116116
- `len(s)`, `substring(s, i, j)` and `s[i]` are O(1) / O(length of the piece) on both engines (strings carry their length).117
-- `s += x` copies `s`, so building a long string one piece at a time is quadratic in time and, since nothing is freed, in memory (and `kiln` has a fixed 16 MB heap). Collect the pieces in an array and `join(parts, "")` once - that is linear.117
+- `s += x` copies `s`, so building a long string one piece at a time is quadratic in time and, since nothing is freed, in memory. Collect the pieces in an array and `join(parts, "")` once - that is linear. `kiln`'s heap grows as needed (64MB at a time, `brk` on Linux / `VirtualAlloc` on Windows) rather than crashing at a fixed size, but nothing is freed until the process exits either way.118118
- `push` grows arrays geometrically (amortized O(1)); `sort` is a stable merge sort on `ashvm` and a stable insertion sort on `kiln`.119119
- `kiln`'s Windows output is checked without a Windows machine: `kiln/tests/scripts/windows_emulated_check.sh` runs every test case as a PE under the Unicorn x86-64 emulator (kernel32 faked over the real filesystem) and compares it with the Linux build. All of them, including the Windows file operations (`list_dir` via `FindFirstFileA`/`FindNextFileA`), match. `kiln/tests/scripts/windows_wine_check.sh` does the same under Wine (an independent `kernel32`), also with all cases matching. Neither is Windows itself.120120
ashvm/Makefilemodified
+21−3@@ -1,7 +1,21 @@11
CC = gcc22
.DEFAULT_GOAL := all3
-CFLAGS = -Wall -Wextra -std=c11 -O2 -Isrc4
-LDFLAGS = -lm3
+# -I../pyre/src, -fopenmp and the ../pyre/lib/libashgpu.a/-ldl/-lvulkan/-lstdc++4
+# below: matrix_mul (src/builtins/C/matrix.c) links against Pyre's CPU +5
+# Vulkan backends, STATICALLY (the explicit .a path, not -lashgpu --6
+# pyre/lib/ also holds libashgpu.so for kiln's --link=shared to find at7
+# runtime, and a bare -lashgpu would silently prefer that .so over the8
+# .a, turning ashvm from self-contained into "needs libashgpu.so9
+# findable at runtime with no rpath to say where" -- found exactly that10
+# way when pyre/Makefile grew the .so target). ashvm already links11
+# against libc/libm normally, unlike kiln's hand-rolled, zero-dependency12
+# ELF/PE output, so this one extra build-time dependency doesn't break13
+# any existing promise (see ideas/assigned.md). -lstdc++ because14
+# libashgpu.a's Vulkan backend is C++ (RAII) internally, even though15
+# ash_gpu.h itself is a plain C ABI -- plain gcc as the final linker16
+# needs that spelled out, g++ would do it automatically.17
+CFLAGS = -Wall -Wextra -std=c11 -O2 -Isrc -I../pyre/src -fopenmp18
+LDFLAGS = -lm ../pyre/lib/libashgpu.a -fopenmp -ldl -lvulkan -lstdc++519
620
ASHVM_SRC = \721
src/app/C/main.c \@@ -30,6 +44,7 @@ ASHVM_SRC = \3044
src/builtins/C/core.c \3145
src/builtins/C/file.c \3246
src/builtins/C/string.c \47
+ src/builtins/C/matrix.c \3348
src/diagnostics/C/diagnostics.c3449
3550
# Objects and dependency files live in trees parallel to each C tree:@@ -60,7 +75,10 @@ BIN = ../bin/ashvm6075
6176
all: $(BIN)6277
63
-$(BIN): $(ASHVM_OBJ)78
+../pyre/lib/libashgpu.a:79
+ $(MAKE) -C ../pyre lib80
+81
+$(BIN): $(ASHVM_OBJ) ../pyre/lib/libashgpu.a6482
mkdir -p ../bin6583
$(CC) $(CFLAGS) -o $(BIN) $(ASHVM_OBJ) $(LDFLAGS)6684
ashvm/src/builtins/C/builtins.cmodified
+3−2@@ -12,7 +12,7 @@ static const char *builtin_names[] = {1212
"pop", "insert", "slice", "sort",1313
"rename_file", "delete_file", "make_dir", "list_dir",1414
"contains", "starts_with", "ends_with", "repeat",15
- "exit"15
+ "exit", "matrix_mul"1616
};1717
// Argument count each builtin is wrapped with when used as a value1818
// (`map(xs, str)`); same order as builtin_names.@@ -25,7 +25,7 @@ static const int builtin_arity[] = {2525
1, 3, 3, 1,2626
2, 1, 1, 1,2727
2, 2, 2, 2,28
- 128
+ 1, 22929
};3030
#define NUM_BUILTINS (int)(sizeof(builtin_names) / sizeof(builtin_names[0]))3131
@@ -48,5 +48,6 @@ VMValue vm_call_builtin(int id, VMValue *args, int argc) {4848
if (id >= 38 && id <= 41) return vm_call_builtin_string(id, args, argc); // contains, starts_with, ends_with, repeat4949
if (id >= 34 && id <= 37) return vm_call_builtin_file(id, args, argc); // rename_file, delete_file, make_dir, list_dir5050
if (id >= 30 && id <= 33) return vm_call_builtin_collection(id, args, argc); // pop, insert, slice, sort51
+ if (id == 43) return vm_call_builtin_matrix(id, args, argc); // matrix_mul5152
vm_runtime_error("unknown builtin id: %d\n", id);5253
}ashvm/src/builtins/C/matrix.cadded
+92−0@@ -0,0 +1,92 @@1
+#include <stdlib.h>2
+#include "builtins/H/internal.h"3
+#include "vm/H/vm.h"4
+#include "H/runtime/ash_gpu.h"5
+6
+// matrix_mul(a, b) -- a and b must be square N x N arrays of arrays of7
+// numbers, same N (general M x K times K x P multiplication is a known8
+// gap: Pyre's own kernel, pyre/src/H/runtime/ash_gpu.h, only takes one N today9
+// -- documented in pyre/todo/todo.md rather than silently assumed away).10
+//11
+// "Always use Pyre when it's ok to" (ideas/assigned.md): ashvm already12
+// links against libc/libm/OpenMP normally (unlike kiln's hand-rolled,13
+// zero-dependency ELF output), so taking on Pyre as a build-time14
+// dependency here isn't the architecture problem it would be for kiln --15
+// see assigned.md for why kiln's own matrix_mul (codegen/C/collections/matrix_mul.c)16
+// instead ports the same loop-order technique natively rather than17
+// linking against libashgpu.18
+//19
+// Below ASH_GPU_MIN_N, a plain local loop skips Pyre entirely: OpenMP's20
+// own thread-launch overhead would dominate a handful of multiply-adds,21
+// exactly the "don't accelerate every operation, only workloads that22
+// justify it" principle (design principle 2) ash_gpu.h itself23
+// documents for picking a backend.24
+#define ASH_GPU_MIN_N 6425
+26
+static int matrix_square_n(VMValue m) {27
+ if (m.type != VM_ARRAY || m.array->count == 0) return -1;28
+ int n = m.array->count;29
+ for (int i = 0; i < n; i++) {30
+ VMValue row = m.array->items[i];31
+ if (row.type != VM_ARRAY || row.array->count != n) return -1;32
+ for (int j = 0; j < n; j++) {33
+ if (row.array->items[j].type != VM_NUM) return -1;34
+ }35
+ }36
+ return n;37
+}38
+39
+static void flatten(VMValue m, int n, double *out) {40
+ for (int i = 0; i < n; i++) {41
+ VMValue row = m.array->items[i];42
+ for (int j = 0; j < n; j++) out[i * n + j] = row.array->items[j].number;43
+ }44
+}45
+46
+static VMValue unflatten(const double *flat, int n) {47
+ VMArray *result = vm_array_new();48
+ for (int i = 0; i < n; i++) {49
+ VMArray *row = vm_array_new();50
+ for (int j = 0; j < n; j++) vm_array_push(row, vm_num(flat[i * n + j]));51
+ vm_array_push(result, vm_array_val(row));52
+ }53
+ return vm_array_val(result);54
+}55
+56
+static void naive_multiply_small(const double *a, const double *b, double *c, int n) {57
+ for (int i = 0; i < n; i++) {58
+ for (int j = 0; j < n; j++) {59
+ double sum = 0.0;60
+ for (int k = 0; k < n; k++) sum += a[i * n + k] * b[k * n + j];61
+ c[i * n + j] = sum;62
+ }63
+ }64
+}65
+66
+VMValue vm_call_builtin_matrix(int id, VMValue *args, int argc) {67
+ (void)id; // only one builtin in this file so far68
+ if (argc != 2) vm_runtime_error("matrix_mul(a, b) expected\n");69
+ int n = matrix_square_n(args[0]);70
+ int n2 = matrix_square_n(args[1]);71
+ if (n == -1 || n2 == -1 || n != n2) {72
+ vm_runtime_error("matrix_mul(a, b) expected two square arrays of arrays of numbers, same size\n");73
+ }74
+75
+ double *a = malloc((size_t)n * n * sizeof(double));76
+ double *b = malloc((size_t)n * n * sizeof(double));77
+ double *c = malloc((size_t)n * n * sizeof(double));78
+ flatten(args[0], n, a);79
+ flatten(args[1], n, b);80
+81
+ if (n >= ASH_GPU_MIN_N) {82
+ ash_gpu_matrix_multiply(a, b, c, n);83
+ } else {84
+ naive_multiply_small(a, b, c, n);85
+ }86
+87
+ VMValue result = unflatten(c, n);88
+ free(a);89
+ free(b);90
+ free(c);91
+ return result;92
+}ashvm/src/builtins/D/matrix.dadded
+8−0@@ -0,0 +1,8 @@1
+src/builtins/O/matrix.o: src/builtins/C/matrix.c \2
+ src/builtins/H/internal.h src/value/H/value.h src/vm/H/vm.h \3
+ src/vm/H/chunk.h ../pyre/src/H/runtime/ash_gpu.h4
+src/builtins/H/internal.h:5
+src/value/H/value.h:6
+src/vm/H/vm.h:7
+src/vm/H/chunk.h:8
+../pyre/src/H/runtime/ash_gpu.h:ashvm/src/builtins/H/internal.hmodified
+1−0@@ -8,5 +8,6 @@ VMValue vm_call_builtin_core(int id, VMValue *args, int argc);88
VMValue vm_call_builtin_collection(int id, VMValue *args, int argc);99
VMValue vm_call_builtin_string(int id, VMValue *args, int argc);1010
VMValue vm_call_builtin_file(int id, VMValue *args, int argc);11
+VMValue vm_call_builtin_matrix(int id, VMValue *args, int argc);1112
1213
#endifashvm/tests/cases/collections/matrix_mul.ashadded
+36−0@@ -0,0 +1,36 @@1
+// matrix_mul(a, b) -- dispatches to Pyre's CPU backend for large square2
+// matrices (ash_gpu_matrix_multiply), computes directly in C for small3
+// ones (see src/builtins/C/matrix.c's ASH_GPU_MIN_N threshold). Both4
+// paths must agree.5
+local a = [[1, 2], [3, 4]];6
+local b = [[5, 6], [7, 8]];7
+say matrix_mul(a, b);8
+9
+forge make_matrix(n, v) {10
+ local mat = [];11
+ local i = 0;12
+ during i < n {13
+ local row = [];14
+ local j = 0;15
+ during j < n {16
+ push(row, v);17
+ j += 1;18
+ }19
+ push(mat, row);20
+ i += 1;21
+ }22
+ yield mat;23
+}24
+25
+// Above the dispatch threshold -- exercises the Pyre CPU backend.26
+local n = 200;27
+local big_a = make_matrix(n, 1.5);28
+local big_b = make_matrix(n, 2.0);29
+local big_c = matrix_mul(big_a, big_b);30
+say big_c[0][0];31
+say big_c[n - 1][n - 1];32
+33
+// Error cases: not square, mismatched sizes, not arrays at all.34
+attempt { matrix_mul([1, 2, 3], [4, 5, 6]); } handle (e) { say "not nested: " + e; }35
+attempt { matrix_mul([[1, 2]], [[1, 2], [3, 4]]); } handle (e) { say "mismatched size: " + e; }36
+attempt { matrix_mul(5, 6); } handle (e) { say "not arrays: " + e; }ideas/assigned.mdadded
+156−0@@ -0,0 +1,156 @@1
+# Assigned: when ashvm and kiln use Pyre2
+3
+This file didn't exist when the Pyre work started (checked exhaustively --4
+nowhere in the repo). It's written now, derived from ideas/gpu_acceleration.md's5
+own Development Phases, Dispatch and Design Principles sections, which is6
+the actual source of truth per that doc's priority over todo.md. Correct7
+this file directly if any rule below doesn't match intent.8
+9
+## Phase order (ideas/gpu_acceleration.md's own table) -- follow this order, not todo.md's grouping10
+11
+1. **Standalone runtime** -- detection, one backend, alloc/dispatch, CPU12
+ fallback, benchmarks. *Don't touch Ash yet.* **Done, with one real13
+ caveat**: `pyre/`'s dispatcher (`src/C/runtime/device.c`) tries a real14
+ Vulkan device via `ash_gpu_vulkan_init()`15
+ (`src/C/backends/vulkan/context.cpp`) before falling back to the CPU16
+ backend, and `ash_gpu_matrix_multiply` has a working Vulkan compute17
+ shader (`src/shaders/matrix_multiply.comp`) behind it -- but this18
+ sandbox has no GPU and no ICD at all (`vulkaninfo`: "Found no19
+ drivers!"), so every piece of that path compiles against the real20
+ Vulkan 1.2 SDK (`pyre/tools/install_deps.sh`) but has **never actually21
+ run**. Both of the user's benchmark challenges are beaten by 85%+22
+ margins on the CPU backend alone (see `pyre/todo/todo.md`) -- that23
+ part IS verified, repeatedly, on real hardware (this machine's own24
+ CPU). Whoever first runs this on the RX 9070 XT should expect to find25
+ and fix real bugs in the Vulkan path; nothing about it has been26
+ execution-tested, only reviewed and compiled.27
+2. **Primitive library** -- fast, benchmark-driven building blocks.28
+ **Done for two kernels**: `ash_gpu_matrix_multiply`, `ash_gpu_monte_carlo_risk`29
+ (`pyre/src/C/backends/cpu/`). More kernels land as more benchmarks or30
+ real call sites need them -- not speculatively ahead of demand.31
+3. **Ash bindings** -- "small API ... **no compiler changes**." **Done32
+ for ashvm** (see below): a plain builtin function that calls the C33
+ ABI, no new syntax, no parser/codegen changes to existing language34
+ constructs.35
+4. **Kiln integration** -- "Kiln emits calls to ash-gpu for suitable36
+ workloads," *only after the runtime is solid*. **Done**: kiln's37
+ generated code really does call into `libashgpu.so` through a GOT38
+ slot now -- see "kiln" below for how. Found and fixed one real,39
+ previously-invisible bug this way (kiln's own `exit`-syscall choice),40
+ which is exactly the value of actually running this instead of41
+ stopping at "compiles."42
+5. **Automatic acceleration** -- plain loops become GPU workloads with43
+ no API call at all. **Not started**, needs escape/aliasing analysis44
+ neither engine has (goals.md item 1 territory). Don't attempt this45
+ before 1-4 are real and measured.46
+47
+## ashvm: done (Phase 3)48
+49
+`matrix_mul(a, b)` (`ashvm/src/builtins/C/matrix.c`) calls50
+`ash_gpu_matrix_multiply` directly. This is "ok" unconditionally at the51
+*build* level: ashvm already links against libc/libm like any normal C52
+program, so one more build-time dependency (Pyre's CPU backend needs53
+OpenMP/libgomp, `runtime/device.c` needs libdl for its detection probe)54
+doesn't break any promise ashvm makes. ashvm's own Makefile links55
+`../pyre/lib/libashgpu.a` unconditionally -- there's no "Pyre not56
+installed" fallback path today; if that matters later (packaging ashvm57
+without Pyre), add a `HAVE_PYRE` build switch then, not speculatively now.58
+59
+**Per-call dispatch rule** (derives the ideas doc's "small workload ->60
+CPU, don't accelerate every operation" principle into an actual number):61
+`ASH_GPU_MIN_N = 64`. Below it, `matrix_mul` computes with a plain local62
+triple loop instead of calling into Pyre -- OpenMP's thread-launch63
+overhead would dominate a handful of multiply-adds. At or above it,64
+Pyre's CPU backend (ikj loop order + OpenMP) runs unconditionally. 64 is65
+a first guess, not a measured threshold; revisit with a benchmark across66
+sizes if it ever matters.67
+68
+## kiln: done, under --link=shared only69
+70
+kiln's compiled *output* is a hand-rolled, zero-dependency ELF/PE binary71
+by default (AGENTS.md: "Ash is a from-scratch programming language in72
+C, no external dependencies" -- this is kiln's whole reason to exist,73
+not incidental). Pyre's CPU backend needs OpenMP (libgomp) and `dlopen`74
+(libdl). Linking *every* kiln-compiled program against those75
+unconditionally would quietly break that promise for every Ash program,76
+not just the ones that use `matrix_mul` -- not "ok" by default. So77
+`matrix_mul(a, b)` (`codegen/C/collections/matrix_mul.c`) is available only78
+under `--link=shared`; under the default `--link=static` it's a compile79
+error naming exactly why, not a silent fallback.80
+81
+How: `--link=shared` already had a real, working example of exactly82
+this shape -- it emits a dynamically-linked ELF with a `PT_INTERP`, so83
+the real system `ld.so` resolves calls into a second binary84
+(`libkilnrt.so`) at load time (`elf/C/elf_dynamic.c`, `elf/C/elf_so_writer.c`).85
+Extended that same machinery to also declare `libashgpu.so` as a SECOND86
+needed library: one more `DT_NEEDED`, and `libashgpu.so`'s one export87
+(`ash_gpu_matrix_multiply`) folded into the SAME combined hash/dynsym/GOT/rela88
+table as kiln's own runtime imports (`KILN_DYN_TOTAL_IMPORTS`,89
+`elf/H/elf_dynamic.h`) rather than a second parallel set of tables --90
+`ld.so` resolves a symbol name against every needed library regardless91
+of which table region declared it, so one combined table is correct and92
+the much smaller change. `codegen/C/platform/linux_call.c` added the93
+one piece kiln never needed before (calling a REAL, normally-compiled94
+C/C++ function instead of one of kiln's own alignment-agnostic95
+hand-rolled routines): a SysV ABI call wrapper that dynamically96
+16-byte-aligns RSP first, same technique `win_call.c` already uses for97
+Windows, reusing its scratch slot (the two targets are never both live98
+in the same compiled program). `matrix_mul` itself flattens its nested99
+array-of-arrays arguments into temporary flat buffers, calls through100
+the GOT slot, and unflattens the result -- the bridge `todo.md` 2.A's101
+still-nonexistent `Buffer`/`Tensor` type would otherwise be.102
+103
+Verified two ways: (1) `kiln/tests/scripts/shared_link_check.sh`104
+(new -- `--link=shared` had NO automated coverage at all before this)105
+runs a sample of ordinary tests under both `--link=static` and106
+`--link=shared` and diffs them, proving the second `DT_NEEDED` and the107
+combined table didn't disturb kiln's own existing runtime imports; (2)108
+actually running `matrix_mul` surfaced a real, previously invisible bug:109
+kiln's process-exit path used the `exit` syscall (60, terminates only110
+the calling thread) instead of `exit_group` (231, terminates the whole111
+process) -- invisible for 100% of kiln's history because no kiln112
+program had ever spawned a second thread, until `ash_gpu_matrix_multiply`'s113
+OpenMP worker pool did. The program printed its correct answer and then114
+hung forever on the orphaned worker threads. Fixed in all three exit115
+sites (`runtime/print_int.c`, `builtins/exit_builtin.c`, `runtime/errors.c`'s116
+two uncaught-error paths); `exit_group` with only one thread behaves117
+identically to `exit`, so every existing single-threaded program's118
+behavior is unchanged (confirmed: `run_regression.sh check` byte-identical119
+before and after). This is the whole point of actually running120
+something instead of stopping at "compiles" -- and the Vulkan backend121
+itself has since actually run too (next section), which is exactly how122
+the matmul exit-syscall bug and the matmul-loses-on-GPU finding below123
+were both found: never by reasoning about the code.124
+125
+## The GPU is real now, and "when it's ok" needs a real answer126
+127
+This sandbox is WSL2 on the user's own machine, which had GPU128
+passthrough infrastructure already present but no Vulkan driver129
+installed (`pacman -S vulkan-dzn mesa` fixed that -- not covered by130
+`pyre/tools/install_deps.sh`, since it's a GPU driver, not SDK/build131
+tooling). `ash_gpu_detect()` now finds the real RX 9070 XT for real.132
+First measurements, same-moment CPU-vs-GPU: `matrix_multiply` is ~2.5x133
+*slower* on GPU (fp64-on-consumer-GPU, an untiled shader, host-visible134
+memory, and Dozen being a translation layer all stack against it);135
+`chaos_iterate` (new kernel, added specifically to find a workload shape136
+that wins) is ~34x *faster* on GPU. Full numbers and why: `pyre/todo/todo.md`'s137
+"Real GPU results" section.138
+139
+Neither dispatcher (`src/C/runtime/device.c`) tells these two cases140
+apart -- both still unconditionally prefer GPU-when-available, which the141
+matmul result proves is actively wrong for that kernel today. The "when142
+it's ok" rules below are therefore still incomplete for Pyre's OWN143
+dispatch, not just ashvm/kiln's: a real cost model (arithmetic intensity,144
+data size, whatever the right signal turns out to be) is the next task,145
+explicitly deferred at the user's request until everything else here is146
+wrapped up -- not forgotten, and not something to guess at ahead of147
+having more than two data points.148
+149
+## What "always ... when it's ok" means in practice150
+151
+- ashvm: always, unconditionally, once `N >= ASH_GPU_MIN_N` -- no152
+ separate opt-in, because ashvm never promised zero dependencies.153
+- kiln: always under `--link=shared` (once `libashgpu.so` is installed154
+ at `KILN_RPATH`, same as `libkilnrt.so`); never under the default155
+ `--link=static`, which keeps working with nothing installed beyond156
+ kiln itself, unconditionally, forever.kiln/Makefilemodified
+2−0@@ -34,6 +34,7 @@ KILN_SRC = \3434
src/codegen/C/collections/keys_values.c \3535
src/codegen/C/collections/maps.c \3636
src/codegen/C/collections/maps_mutate.c \37
+ src/codegen/C/collections/matrix_mul.c \3738
src/codegen/C/emit/bytes.c \3839
src/codegen/C/emit/emit.c \3940
src/codegen/C/emit/emit_mem.c \@@ -53,6 +54,7 @@ KILN_SRC = \5354
src/codegen/C/io/write_file_builtins.c \5455
src/codegen/C/platform/platform_console.c \5556
src/codegen/C/platform/platform_file.c \57
+ src/codegen/C/platform/linux_call.c \5658
src/codegen/C/platform/win_call.c \5759
src/codegen/C/platform/win_file_call.c \5860
src/codegen/C/runtime/errors.c \kiln/src/app/D/main.dmodified
+5−1@@ -2,7 +2,9 @@ src/app/O/main.o: src/app/C/main.c src/parser/H/core/parser.h \22
src/lexer/H/lexer.h src/codegen/H/emit/emit.h src/elf/H/elf_writer.h \33
src/codegen/H/emit/runtime_layout.h src/elf/H/elf_dynamic.h \44
src/codegen/H/emit/runtime_import_enum.h \5
- src/codegen/H/emit/runtime_exports.def src/elf/H/elf_so_writer.h \5
+ src/codegen/H/emit/runtime_exports.def \6
+ src/codegen/H/emit/ash_gpu_import_enum.h \7
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_so_writer.h \68
src/pe/H/pe_writer.h src/pe/H/pe_imports.def src/pe/H/pe_dll_writer.h \79
src/diagnostics/H/diagnostics.h src/app/H/target.h810
src/parser/H/core/parser.h:@@ -13,6 +15,8 @@ src/codegen/H/emit/runtime_layout.h:1315
src/elf/H/elf_dynamic.h:1416
src/codegen/H/emit/runtime_import_enum.h:1517
src/codegen/H/emit/runtime_exports.def:18
+src/codegen/H/emit/ash_gpu_import_enum.h:19
+src/codegen/H/emit/ash_gpu_exports.def:1620
src/elf/H/elf_so_writer.h:1721
src/pe/H/pe_writer.h:1822
src/pe/H/pe_imports.def:kiln/src/codegen/C/builtins/builtins.cmodified
+9−0@@ -3,6 +3,7 @@33
#include "codegen/H/expressions/expr.h"44
#include "codegen/H/collections/arrays.h"55
#include "codegen/H/collections/maps.h"6
+#include "codegen/H/collections/matrix_mul.h"67
#include "codegen/H/builtins/math_builtins.h"78
#include "codegen/H/builtins/convert_builtins.h"89
#include "codegen/H/builtins/higher_order.h"@@ -171,6 +172,14 @@ int codegen_try_builtin_call(const char *name, int len) {171172
codegen_builtin_push();172173
return 1;173174
}175
+ if (len == 10 && strncmp(name, "matrix_mul", 10) == 0) {176
+ codegen_expression();177
+ expect(TOKEN_COMMA, "expected ',' after array argument");178
+ codegen_expression();179
+ expect(TOKEN_RPAREN, "expected ')' after arguments");180
+ codegen_builtin_matrix_mul();181
+ return 1;182
+ }174183
if (len == 3 && strncmp(name, "len", 3) == 0) {175184
codegen_expression();176185
expect(TOKEN_RPAREN, "expected ')' after argument");kiln/src/codegen/C/builtins/exit_builtin.cmodified
+3−1@@ -20,7 +20,9 @@ void codegen_builtin_exit(void) {2020
win_call_arg_reg(code, 0, REG_RDI);2121
win_call_import(code, PE_IMPORT_EXIT_PROCESS);2222
} else {23
- emit_mov_reg_imm64(code, REG_RAX, 60); // syscall: exit23
+ // exit_group (231), not exit (60) -- see runtime/print_int.c's24
+ // codegen_exit0 for why (the same reasoning applies here).25
+ emit_mov_reg_imm64(code, REG_RAX, 231);2426
emit_syscall(code);2527
}2628
// never returns; keep the stack-machine model balanced for the callerkiln/src/codegen/C/collections/matrix_mul.cadded
+318−0@@ -0,0 +1,318 @@1
+#include "codegen/H/collections/matrix_mul.h"2
+#include "codegen/H/emit/emit.h"3
+#include "codegen/H/emit/value.h"4
+#include "codegen/H/platform/linux_call.h"5
+#include "codegen/H/runtime/heap.h"6
+#include "codegen/H/runtime/errors.h"7
+#include "parser/H/core/parser.h"8
+#include "parser/H/declarations/vars.h"9
+#include "app/H/target.h"10
+11
+// matrix_mul(a, b) -- real Phase 4 integration (ideas/gpu_acceleration.md,12
+// ideas/assigned.md): calls INTO libashgpu.so's ash_gpu_matrix_multiply13
+// through the GOT slot elf/C/elf_dynamic.c's --link=shared writer lays14
+// out, rather than reimplementing the algorithm in kiln's own hand-rolled15
+// x86-64 the way the CPU backend itself does -- that distinction is the16
+// whole point (see assigned.md for why).17
+//18
+// Only available under --link=shared: that mode is the only one with19
+// any dynamic-linking machinery at all for a second needed library to20
+// hang off of (see elf/H/elf_dynamic.h) -- --link=static stays the21
+// hand-rolled, zero-dependency ELF/PE kiln otherwise always produces,22
+// unconditionally, which is the entire reason that mode exists.23
+//24
+// Square N x N arrays of arrays of numbers only, same N (matches25
+// ash_gpu_matrix_multiply's own signature exactly, pyre/src/H/runtime/ash_gpu.h --26
+// general M x K times K x P is a known gap, same one ashvm's own27
+// matrix_mul documents, not silently assumed away). Marshals by28
+// flattening into temporary flat double buffers (todo.md 2.A's29
+// "Contiguous Buffer Type" doesn't exist yet, so this is the bridge30
+// until it does) -- none of the three temporary buffers are freed31
+// (kiln has no free() at all yet; the result's own two are returned to32
+// the caller anyway and would need to stay alive regardless).33
+//34
+// Scratch layout (higher_order_scratch_offset(), each field 8 bytes):35
+// +0 a_payload +8 b_payload +16 n36
+// +24 flat_a +32 flat_b +40 flat_c37
+// +48 i +56 j +64 row_payload (scratch)38
+// +72 result_payload +80 c_data (result's own outer data block)39
+static int32_t scratch_base;40
+#define S(off) (scratch_base + (off))41
+42
+static void die(CodeBuf *code) {43
+ errors_emit_die(code, "runtime error: matrix_mul(a, b) expects two non-empty square arrays of arrays of numbers, same size");44
+}45
+46
+// Allocates an n-element data block of (TAG_ARRAY, row) pairs is NOT47
+// what this builds -- rows are built by flatten/unflatten below, each48
+// inline where it's needed, since the two directions differ (reading a49
+// row's own data_ptr vs. allocating one from scratch).50
+51
+// scratch[S(dst_flat)][i*n+j] = <nested array>.data_ptr[i].payload's row's52
+// data_ptr[j].payload, for i,j in [0, n) -- flattens `array_scratch_off`53
+// (an a_payload/b_payload-shaped nested array) into a freshly allocated54
+// flat buffer, leaving the buffer's address in scratch[dst_flat_off].55
+static void flatten(CodeBuf *code, int32_t array_scratch_off, int32_t dst_flat_off) {56
+ emit_load_mem_disp32(code, REG_RDI, REG_RBP, S(16)); // n57
+ emit_imul_reg_reg(code, REG_RDI, REG_RDI); // n*n58
+ emit_mov_reg_imm64(code, REG_RAX, 8);59
+ emit_imul_reg_reg(code, REG_RDI, REG_RAX); // n*n*8 bytes60
+ heap_emit_alloc(code); // RAX = flat buffer61
+ emit_store_mem_disp32(code, REG_RBP, S(dst_flat_off), REG_RAX);62
+63
+ emit_mov_reg_imm64(code, REG_RCX, 0);64
+ emit_store_mem_disp32(code, REG_RBP, S(48), REG_RCX); // i = 065
+ int loop_i = code->count;66
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(48));67
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(16));68
+ emit_cmp_reg_reg(code, REG_RCX, REG_RDX);69
+ int exit_i = emit_jcc_rel32(code, COND_GE);70
+71
+ // row_payload = array.data_ptr[i].payload72
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(array_scratch_off));73
+ emit_load_mem_disp32(code, REG_RSI, REG_RSI, 16); // data_ptr74
+ emit_mov_reg_reg(code, REG_RAX, REG_RCX);75
+ emit_mov_reg_imm64(code, REG_RDX, 16);76
+ emit_imul_reg_reg(code, REG_RAX, REG_RDX);77
+ emit_add_reg_reg(code, REG_RSI, REG_RAX);78
+ emit_load_mem_disp32(code, REG_RAX, REG_RSI, 8); // row payload79
+ emit_store_mem_disp32(code, REG_RBP, S(64), REG_RAX);80
+81
+ emit_mov_reg_imm64(code, REG_RCX, 0);82
+ emit_store_mem_disp32(code, REG_RBP, S(56), REG_RCX); // j = 083
+ int loop_j = code->count;84
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(56));85
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(16));86
+ emit_cmp_reg_reg(code, REG_RCX, REG_RDX);87
+ int exit_j = emit_jcc_rel32(code, COND_GE);88
+89
+ // value = row.data_ptr[j].payload90
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(64));91
+ emit_load_mem_disp32(code, REG_RSI, REG_RSI, 16);92
+ emit_mov_reg_reg(code, REG_RAX, REG_RCX);93
+ emit_mov_reg_imm64(code, REG_RDX, 16);94
+ emit_imul_reg_reg(code, REG_RAX, REG_RDX);95
+ emit_add_reg_reg(code, REG_RSI, REG_RAX);96
+ emit_load_mem_disp32(code, REG_RAX, REG_RSI, 8); // value bits97
+98
+ // dest = flat + (i*n+j)*899
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(dst_flat_off));100
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(48)); // i101
+ emit_load_mem_disp32(code, REG_RBX, REG_RBP, S(16)); // n102
+ emit_imul_reg_reg(code, REG_RDX, REG_RBX);103
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(56)); // j104
+ emit_add_reg_reg(code, REG_RDX, REG_RCX);105
+ emit_mov_reg_imm64(code, REG_RBX, 8);106
+ emit_imul_reg_reg(code, REG_RDX, REG_RBX);107
+ emit_add_reg_reg(code, REG_RSI, REG_RDX);108
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RAX);109
+110
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(56));111
+ emit_add_reg_imm8(code, REG_RCX, 1);112
+ emit_store_mem_disp32(code, REG_RBP, S(56), REG_RCX);113
+ emit_jmp_back(code, loop_j);114
+ emit_patch_jump(code, exit_j);115
+116
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(48));117
+ emit_add_reg_imm8(code, REG_RCX, 1);118
+ emit_store_mem_disp32(code, REG_RBP, S(48), REG_RCX);119
+ emit_jmp_back(code, loop_i);120
+ emit_patch_jump(code, exit_i);121
+}122
+123
+// Inverse of flatten(): builds a fresh n x n nested array (header + n124
+// row headers + n row data blocks) from scratch[40] (flat_c), leaving125
+// the result's (tag, payload) pushed.126
+static void unflatten_result(CodeBuf *code) {127
+ // outer data block: n entries of (TAG_ARRAY, row_payload)128
+ emit_load_mem_disp32(code, REG_RDI, REG_RBP, S(16));129
+ emit_mov_reg_imm64(code, REG_RAX, 16);130
+ emit_imul_reg_reg(code, REG_RDI, REG_RAX);131
+ heap_emit_alloc(code); // RAX = outer data block132
+ emit_store_mem_disp32(code, REG_RBP, S(80), REG_RAX);133
+134
+ emit_mov_reg_imm64(code, REG_RCX, 0);135
+ emit_store_mem_disp32(code, REG_RBP, S(48), REG_RCX); // i = 0136
+ int loop_i = code->count;137
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(48));138
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(16));139
+ emit_cmp_reg_reg(code, REG_RCX, REG_RDX);140
+ int exit_i = emit_jcc_rel32(code, COND_GE);141
+142
+ // row data block: n doubles, copied straight from flat_c's row143
+ emit_load_mem_disp32(code, REG_RDI, REG_RBP, S(16));144
+ emit_mov_reg_imm64(code, REG_RAX, 16);145
+ emit_imul_reg_reg(code, REG_RDI, REG_RAX);146
+ heap_emit_alloc(code); // RAX = row data block147
+ emit_store_mem_disp32(code, REG_RBP, S(64), REG_RAX);148
+149
+ emit_mov_reg_imm64(code, REG_RCX, 0);150
+ emit_store_mem_disp32(code, REG_RBP, S(56), REG_RCX); // j = 0151
+ int loop_j = code->count;152
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(56));153
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(16));154
+ emit_cmp_reg_reg(code, REG_RCX, REG_RDX);155
+ int exit_j = emit_jcc_rel32(code, COND_GE);156
+157
+ // value = flat_c[i*n+j]158
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(40)); // flat_c159
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(48)); // i160
+ emit_load_mem_disp32(code, REG_RBX, REG_RBP, S(16)); // n161
+ emit_imul_reg_reg(code, REG_RDX, REG_RBX);162
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(56)); // j163
+ emit_add_reg_reg(code, REG_RDX, REG_RCX);164
+ emit_mov_reg_imm64(code, REG_RBX, 8);165
+ emit_imul_reg_reg(code, REG_RDX, REG_RBX);166
+ emit_add_reg_reg(code, REG_RSI, REG_RDX);167
+ emit_load_mem_disp32(code, REG_RAX, REG_RSI, 0); // value bits168
+169
+ // row_data[j] = (TAG_NUMBER, value)170
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(64)); // row data block171
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(56)); // j172
+ emit_mov_reg_imm64(code, REG_RBX, 16);173
+ emit_imul_reg_reg(code, REG_RDX, REG_RBX);174
+ emit_add_reg_reg(code, REG_RSI, REG_RDX);175
+ emit_mov_reg_imm64(code, REG_RBX, TAG_NUMBER);176
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RBX);177
+ emit_store_mem_disp32(code, REG_RSI, 8, REG_RAX);178
+179
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(56));180
+ emit_add_reg_imm8(code, REG_RCX, 1);181
+ emit_store_mem_disp32(code, REG_RBP, S(56), REG_RCX);182
+ emit_jmp_back(code, loop_j);183
+ emit_patch_jump(code, exit_j);184
+185
+ // row header: {capacity=n, count=n, data_ptr=row data block}186
+ emit_mov_reg_imm64(code, REG_RDI, ARRAY_OBJECT_SIZE);187
+ heap_emit_alloc(code); // RAX = row object188
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(16));189
+ emit_store_mem_disp32(code, REG_RAX, 0, REG_RCX);190
+ emit_store_mem_disp32(code, REG_RAX, 8, REG_RCX);191
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(64));192
+ emit_store_mem_disp32(code, REG_RAX, 16, REG_RDX);193
+194
+ // outer[i] = (TAG_ARRAY, row object)195
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(80));196
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(48));197
+ emit_mov_reg_imm64(code, REG_RDX, 16);198
+ emit_imul_reg_reg(code, REG_RCX, REG_RDX);199
+ emit_add_reg_reg(code, REG_RSI, REG_RCX);200
+ emit_mov_reg_imm64(code, REG_RDX, TAG_ARRAY);201
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RDX);202
+ emit_store_mem_disp32(code, REG_RSI, 8, REG_RAX);203
+204
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(48));205
+ emit_add_reg_imm8(code, REG_RCX, 1);206
+ emit_store_mem_disp32(code, REG_RBP, S(48), REG_RCX);207
+ emit_jmp_back(code, loop_i);208
+ emit_patch_jump(code, exit_i);209
+210
+ // outer header: {capacity=n, count=n, data_ptr=outer data block}211
+ emit_mov_reg_imm64(code, REG_RDI, ARRAY_OBJECT_SIZE);212
+ heap_emit_alloc(code); // RAX = result object213
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(16));214
+ emit_store_mem_disp32(code, REG_RAX, 0, REG_RCX);215
+ emit_store_mem_disp32(code, REG_RAX, 8, REG_RCX);216
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(80));217
+ emit_store_mem_disp32(code, REG_RAX, 16, REG_RDX);218
+219
+ emit_mov_reg_imm64(code, REG_RBX, TAG_ARRAY);220
+ emit_push_reg(code, REG_RBX);221
+ emit_push_reg(code, REG_RAX);222
+}223
+224
+void codegen_builtin_matrix_mul(void) {225
+ if (kiln_get_link_mode() != KILN_LINK_SHARED) {226
+ parse_error("matrix_mul requires --link=shared (it calls into libashgpu.so; "227
+ "--link=static stays kiln's zero-dependency default -- see ideas/assigned.md)");228
+ }229
+230
+ scratch_base = higher_order_scratch_offset();231
+232
+ // args already pushed left-to-right by codegen/C/builtins/builtins.c's233
+ // dispatch: [a_tag, a_payload, b_tag, b_payload]234
+ emit_pop_reg(code, REG_RAX); // b_payload235
+ emit_pop_reg(code, REG_RBX); // b_tag236
+ emit_pop_reg(code, REG_RCX); // a_payload237
+ emit_pop_reg(code, REG_RDX); // a_tag238
+ emit_cmp_reg_imm32(code, REG_RDX, TAG_ARRAY);239
+ int a_ok = emit_jcc_rel32(code, COND_E);240
+ die(code);241
+ emit_patch_jump(code, a_ok);242
+ emit_cmp_reg_imm32(code, REG_RBX, TAG_ARRAY);243
+ int b_ok = emit_jcc_rel32(code, COND_E);244
+ die(code);245
+ emit_patch_jump(code, b_ok);246
+247
+ emit_store_mem_disp32(code, REG_RBP, S(0), REG_RCX); // a_payload248
+ emit_store_mem_disp32(code, REG_RBP, S(8), REG_RAX); // b_payload249
+250
+ // n = a.count; must equal b.count, and be > 0251
+ emit_load_mem_disp32(code, REG_RAX, REG_RCX, 8); // a.count252
+ emit_store_mem_disp32(code, REG_RBP, S(16), REG_RAX);253
+ emit_cmp_reg_imm32(code, REG_RAX, 0);254
+ int n_nonzero = emit_jcc_rel32(code, COND_GT);255
+ die(code);256
+ emit_patch_jump(code, n_nonzero);257
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(8));258
+ emit_load_mem_disp32(code, REG_RCX, REG_RCX, 8); // b.count259
+ emit_cmp_reg_reg(code, REG_RAX, REG_RCX);260
+ int sizes_match = emit_jcc_rel32(code, COND_E);261
+ die(code);262
+ emit_patch_jump(code, sizes_match);263
+264
+ // row0 of a: must be TAG_ARRAY with count == n265
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(0));266
+ emit_load_mem_disp32(code, REG_RSI, REG_RSI, 16); // a.data_ptr267
+ emit_load_mem_disp32(code, REG_RBX, REG_RSI, 0); // row0 tag268
+ emit_cmp_reg_imm32(code, REG_RBX, TAG_ARRAY);269
+ int a_row_ok = emit_jcc_rel32(code, COND_E);270
+ die(code);271
+ emit_patch_jump(code, a_row_ok);272
+ emit_load_mem_disp32(code, REG_RAX, REG_RSI, 8); // row0 payload273
+ emit_load_mem_disp32(code, REG_RCX, REG_RAX, 8); // row0 count274
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(16));275
+ emit_cmp_reg_reg(code, REG_RCX, REG_RDX);276
+ int a_row_size_ok = emit_jcc_rel32(code, COND_E);277
+ die(code);278
+ emit_patch_jump(code, a_row_size_ok);279
+280
+ // row0 of b: same check281
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(8));282
+ emit_load_mem_disp32(code, REG_RSI, REG_RSI, 16);283
+ emit_load_mem_disp32(code, REG_RBX, REG_RSI, 0);284
+ emit_cmp_reg_imm32(code, REG_RBX, TAG_ARRAY);285
+ int b_row_ok = emit_jcc_rel32(code, COND_E);286
+ die(code);287
+ emit_patch_jump(code, b_row_ok);288
+ emit_load_mem_disp32(code, REG_RAX, REG_RSI, 8);289
+ emit_load_mem_disp32(code, REG_RCX, REG_RAX, 8);290
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(16));291
+ emit_cmp_reg_reg(code, REG_RCX, REG_RDX);292
+ int b_row_size_ok = emit_jcc_rel32(code, COND_E);293
+ die(code);294
+ emit_patch_jump(code, b_row_size_ok);295
+296
+ flatten(code, 0, 24); // flat_a from a_payload297
+ flatten(code, 8, 32); // flat_b from b_payload298
+299
+ // flat_c: n*n*8 bytes, left uninitialized -- ash_gpu_matrix_multiply300
+ // fully overwrites it itself (see pyre/src/C/backends/cpu/matrix_multiply.c's301
+ // own memset at the top of its CPU path).302
+ emit_load_mem_disp32(code, REG_RDI, REG_RBP, S(16));303
+ emit_imul_reg_reg(code, REG_RDI, REG_RDI);304
+ emit_mov_reg_imm64(code, REG_RAX, 8);305
+ emit_imul_reg_reg(code, REG_RDI, REG_RAX);306
+ heap_emit_alloc(code);307
+ emit_store_mem_disp32(code, REG_RBP, S(40), REG_RAX);308
+309
+ linux_abi_call_begin(code);310
+ emit_load_mem_disp32(code, REG_RDI, REG_RBP, S(24)); // flat_a311
+ emit_load_mem_disp32(code, REG_RSI, REG_RBP, S(32)); // flat_b312
+ emit_load_mem_disp32(code, REG_RDX, REG_RBP, S(40)); // flat_c313
+ emit_load_mem_disp32(code, REG_RCX, REG_RBP, S(16)); // n314
+ ash_gpu_dynamic_call(code, ASH_GPU_IMPORT_MATRIX_MULTIPLY);315
+ linux_abi_call_end(code);316
+317
+ unflatten_result(code);318
+}kiln/src/codegen/C/emit/layout.cmodified
+8−2@@ -31,9 +31,15 @@ uint64_t kiln_heap_ptr_addr(void) {3131
return KILN_GLOBALS_ADDR;3232
}3333
34
+uint64_t kiln_heap_limit_addr(void) {35
+ if (kiln_get_target() == KILN_TARGET_WINDOWS) return KILN_PE_HEAP_LIMIT_ADDR;36
+ if (is_linux_shared()) return KILN_DYN_GLOBALS_ADDR + 8;37
+ return KILN_HEAP_LIMIT_ADDR;38
+}39
+3440
uint64_t kiln_try_depth_addr(void) {3541
if (kiln_get_target() == KILN_TARGET_WINDOWS) return KILN_PE_TRY_DEPTH_ADDR;36
- if (is_linux_shared()) return KILN_DYN_GLOBALS_ADDR + 8;42
+ if (is_linux_shared()) return KILN_DYN_GLOBALS_ADDR + 8 + KILN_HEAP_LIMIT_SIZE;3743
return KILN_TRY_DEPTH_ADDR;3844
}3945
@@ -45,6 +51,6 @@ uint64_t kiln_modstate_addr(void) {4551
4652
uint64_t kiln_try_handlers_addr(void) {4753
if (kiln_get_target() == KILN_TARGET_WINDOWS) return KILN_PE_TRY_HANDLERS_ADDR;48
- if (is_linux_shared()) return KILN_DYN_GLOBALS_ADDR + 8 + KILN_TRY_DEPTH_SIZE;54
+ if (is_linux_shared()) return KILN_DYN_GLOBALS_ADDR + 8 + KILN_HEAP_LIMIT_SIZE + KILN_TRY_DEPTH_SIZE;4955
return KILN_TRY_HANDLERS_ADDR;5056
}kiln/src/codegen/C/expressions/expr.cmodified
+3−4@@ -39,7 +39,6 @@3939
// milestone, not silently pretended-away here).4040
4141
static void power(void);42
-static void emit_index_read(void);4342
4443
// A builtin name used as a value (`map(xs, str)`, `local f = len;`) becomes a4544
// small wrapper lambda `forge (a0, ...) { yield name(a0, ...); }`, compiled by@@ -137,7 +136,7 @@ static void primary(void) {137136
advance_token();138137
expect(TOKEN_IDENTIFIER, "expected field name after '.'");139138
codegen_string_literal_bytes(previous.start, previous.length);140
- emit_index_read();139
+ codegen_index_read();141140
}142141
return;143142
}@@ -237,7 +236,7 @@ static void primary(void) {237236
// Both the base and the index are already pushed (base deeper). Pops both and238237
// pushes base[index] for an array, map or string. A string indexed by a string239238
// is an error value's `.message`: it yields the string itself.240
-static void emit_index_read(void) {239
+void codegen_index_read(void) {241240
emit_pop_reg(code, REG_RAX); // index payload242241
emit_pop_reg(code, REG_RBX); // index tag243242
emit_pop_reg(code, REG_RCX); // base payload@@ -273,7 +272,7 @@ static void postfix(void) {273272
advance_token();274273
codegen_expression();275274
expect(TOKEN_RBRACKET, "expected ']' after index");276
- emit_index_read();275
+ codegen_index_read();277276
}278277
}279278
kiln/src/codegen/C/platform/linux_call.cadded
+24−0@@ -0,0 +1,24 @@1
+#include "codegen/H/platform/linux_call.h"2
+#include "codegen/H/emit/emit_internal.h"3
+#include "parser/H/declarations/vars.h"4
+5
+// AND RSP, imm8 (sign-extended) -- opcode 0x83 /4. Same encoding6
+// win_call.c's own emit_and_rsp_imm8 uses; kept as a separate local7
+// copy rather than exporting that one, matching this project's existing8
+// pattern of small leaf instruction-encoders staying local to whichever9
+// file needs them (win_call.c has several of its own the same way).10
+static void emit_and_rsp_imm8(CodeBuf *code, int8_t imm) {11
+ emit_byte(code, REX_W);12
+ emit_byte(code, 0x83);13
+ emit_byte(code, modrm_reg((Reg)4, REG_RSP)); // /4 = AND14
+ emit_byte(code, (uint8_t)imm);15
+}16
+17
+void linux_abi_call_begin(CodeBuf *code) {18
+ emit_store_mem_disp32(code, REG_RBP, win_call_scratch_offset(), REG_RSP);19
+ emit_and_rsp_imm8(code, -16);20
+}21
+22
+void linux_abi_call_end(CodeBuf *code) {23
+ emit_load_mem_disp32(code, REG_RSP, REG_RBP, win_call_scratch_offset());24
+}kiln/src/codegen/C/runtime/errors.cmodified
+2−2@@ -126,7 +126,7 @@ static void emit_raise_routine(CodeBuf *code) {126126
win_call_import(code, PE_IMPORT_EXIT_PROCESS);127127
} else {128128
emit_mov_reg_imm64(code, REG_RDI, 1); // exit status 1129
- emit_mov_reg_imm64(code, REG_RAX, 60); // syscall: exit129
+ emit_mov_reg_imm64(code, REG_RAX, 231); // syscall: exit_group, not exit -- see runtime/print_int.c's codegen_exit0130130
emit_syscall(code);131131
}132132
@@ -192,7 +192,7 @@ void errors_emit_fatal(CodeBuf *code, const char *msg) {192192
win_call_import(code, PE_IMPORT_EXIT_PROCESS);193193
} else {194194
emit_mov_reg_imm64(code, REG_RDI, 1);195
- emit_mov_reg_imm64(code, REG_RAX, 60);195
+ emit_mov_reg_imm64(code, REG_RAX, 231); // syscall: exit_group, not exit -- see runtime/print_int.c's codegen_exit0196196
emit_syscall(code);197197
}198198
}kiln/src/codegen/C/runtime/heap.cmodified
+131−15@@ -4,8 +4,27 @@44
#include "elf/H/elf_dynamic_call.h"55
#include "app/H/target.h"66
7
-#define HEAP_SIZE (16 * 1024 * 1024)8
-#define MEM_COMMIT_RESERVE 0x3000u // MEM_COMMIT(0x1000) | MEM_RESERVE(0x2000)7
+// Chunk-based growth, replacing the old single fixed 16MB grab: the arena8
+// starts with one chunk and the alloc routine below grows it by another9
+// chunk (rounding the request up to a chunk boundary, so one allocation10
+// bigger than a single chunk still gets everything it needs in one grow11
+// step) whenever the bump pointer would cross the committed limit. Must be12
+// a power of two -- the round-up uses an AND mask, not a division.13
+#define HEAP_CHUNK_SIZE (64u * 1024 * 1024)14
+15
+// Windows only: VirtualAlloc has no brk-style "extend the current16
+// region" -- growing in place requires the address range to already be17
+// reserved, so startup reserves this much virtual address space up front18
+// (cheap: MEM_RESERVE alone commits no physical memory) and each growth19
+// step just commits one more chunk inside it. 4GB comfortably covers20
+// anything this project's current workloads (matrix ops and beyond) need;21
+// if a program ever outgrows it, that is a real resource ceiling worth22
+// hitting deliberately rather than silently wrapping into unreserved23
+// address space.24
+#define HEAP_RESERVE_SIZE (4ULL * 1024 * 1024 * 1024)25
+26
+#define MEM_COMMIT 0x1000u27
+#define MEM_RESERVE 0x2000u928
#define PAGE_EXECUTE_READWRITE 0x40u1029
1130
// Plan B, phase B1: emitted ONCE (jumped over, like errors.c's raise@@ -18,13 +37,77 @@1837
// local one.1938
static int alloc_routine_offset = -1;2039
40
+// Callers of heap_emit_alloc (~27 sites) rely on a fixed clobber set:41
+// RAX (return value) and RDI (consumes the size argument) change, RSI is42
+// explicitly documented as clobbered, but RBX/RCX/RDX/RBP/RSP must come43
+// back exactly as they went in (several call sites read RCX/RDX straight44
+// through a call with no reload). The growth path below needs scratch45
+// registers to talk to brk/VirtualAlloc, so it saves and restores46
+// RAX/RDI (which this routine's own fast path is still using) plus47
+// RBX/RCX/RDX (borrowed as scratch) around that one-off syscall/API call,48
+// instead of changing the contract for the other ~26 sites that never hit49
+// the slow path at all.2150
static void emit_alloc_routine(CodeBuf *code) {2251
int skip = emit_jmp_rel32(code);2352
alloc_routine_offset = code->count;53
+2454
emit_mov_reg_imm64(code, REG_RSI, kiln_heap_ptr_addr());25
- emit_load_mem_disp32(code, REG_RAX, REG_RSI, 0); // rax = heap_ptr (becomes the returned block address)55
+ emit_load_mem_disp32(code, REG_RAX, REG_RSI, 0); // rax = heap_ptr (the block this call returns)2656
emit_add_reg_reg(code, REG_RDI, REG_RAX); // rdi = heap_ptr + size = new heap_ptr27
- emit_store_mem_disp32(code, REG_RSI, 0, REG_RDI);57
+58
+ emit_mov_reg_imm64(code, REG_RSI, kiln_heap_limit_addr());59
+ emit_load_mem_disp32(code, REG_RSI, REG_RSI, 0); // rsi = current committed limit60
+61
+ emit_cmp_reg_reg(code, REG_RDI, REG_RSI); // new heap_ptr vs limit62
+ int no_growth = emit_jcc_rel32(code, COND_LE); // common case: still inside the committed region63
+64
+ // ---- slow path: grow by whole chunks, rounding the new pointer up65
+ // to the next chunk boundary so a single oversized allocation still66
+ // gets enough room in one step. ----67
+ emit_push_reg(code, REG_RAX); // old heap_ptr (this call's return value)68
+ emit_push_reg(code, REG_RDI); // new heap_ptr69
+ emit_push_reg(code, REG_RBX);70
+ emit_push_reg(code, REG_RCX);71
+ emit_push_reg(code, REG_RDX);72
+73
+ emit_mov_reg_reg(code, REG_RBX, REG_RDI);74
+ emit_add_reg_imm32(code, REG_RBX, (int32_t)(HEAP_CHUNK_SIZE - 1));75
+ emit_mov_reg_imm64(code, REG_RDX, ~(uint64_t)(HEAP_CHUNK_SIZE - 1));76
+ emit_and_reg_reg(code, REG_RBX, REG_RDX); // rbx = new committed limit (chunk-rounded)77
+78
+ if (kiln_get_target() == KILN_TARGET_WINDOWS) {79
+ emit_mov_reg_reg(code, REG_RCX, REG_RSI); // rcx = old limit = commit base (already reserved at startup)80
+ emit_mov_reg_reg(code, REG_RAX, REG_RBX);81
+ emit_sub_reg_reg(code, REG_RAX, REG_RCX); // rax = size to commit82
+83
+ win_call_begin(code, 0);84
+ win_call_arg_reg(code, 0, REG_RCX); // lpAddress = old limit85
+ win_call_arg_reg(code, 1, REG_RAX); // dwSize86
+ win_call_arg_imm64(code, 2, MEM_COMMIT);87
+ win_call_arg_imm64(code, 3, PAGE_EXECUTE_READWRITE);88
+ win_call_import(code, PE_IMPORT_VIRTUAL_ALLOC);89
+ win_call_end(code);90
+ } else {91
+ // brk(new_limit) -- not checking the result, same "a fixed92
+ // request failing isn't a case to write defensive code around"93
+ // call as the one-time startup grab below.94
+ emit_mov_reg_reg(code, REG_RDI, REG_RBX);95
+ emit_mov_reg_imm64(code, REG_RAX, 12); // syscall: brk96
+ emit_syscall(code);97
+ }98
+99
+ emit_mov_reg_imm64(code, REG_RSI, kiln_heap_limit_addr());100
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RBX); // heap_limit = new committed limit101
+102
+ emit_pop_reg(code, REG_RDX);103
+ emit_pop_reg(code, REG_RCX);104
+ emit_pop_reg(code, REG_RBX);105
+ emit_pop_reg(code, REG_RDI); // new heap_ptr, restored106
+ emit_pop_reg(code, REG_RAX); // old heap_ptr, restored107
+108
+ emit_patch_jump(code, no_growth);109
+ emit_mov_reg_imm64(code, REG_RSI, kiln_heap_ptr_addr());110
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RDI); // heap_ptr = new heap_ptr28111
emit_ret(code);29112
emit_patch_jump(code, skip);30113
}@@ -51,17 +134,34 @@ static int is_shared(void) {51134
52135
void heap_emit_startup(CodeBuf *code) {53136
if (kiln_get_target() == KILN_TARGET_WINDOWS) {54
- // VirtualAlloc hands back a ready-to-use base directly -- no55
- // "query then grow" dance needed the way brk requires.137
+ // Reserve a large virtual range up front (no physical memory138
+ // committed -- MEM_RESERVE alone), then commit only the first139
+ // chunk. Growth later commits further chunks inside this same140
+ // reservation, so the whole arena stays one contiguous region141
+ // and existing pointers into it are never invalidated.56142
win_call_begin(code, 0);57
- win_call_arg_imm64(code, 0, 0); // lpAddress = NULL58
- win_call_arg_imm64(code, 1, HEAP_SIZE);59
- win_call_arg_imm64(code, 2, MEM_COMMIT_RESERVE);143
+ win_call_arg_imm64(code, 0, 0); // lpAddress = NULL, let the OS place it144
+ win_call_arg_imm64(code, 1, HEAP_RESERVE_SIZE);145
+ win_call_arg_imm64(code, 2, MEM_RESERVE);60146
win_call_arg_imm64(code, 3, PAGE_EXECUTE_READWRITE);61147
win_call_import(code, PE_IMPORT_VIRTUAL_ALLOC);62148
win_call_end(code);149
+ emit_mov_reg_reg(code, REG_RBX, REG_RAX); // rbx = base (survives the next win_call: non-volatile in the Windows x64 ABI)150
+151
+ win_call_begin(code, 0);152
+ win_call_arg_reg(code, 0, REG_RBX); // lpAddress = base153
+ win_call_arg_imm64(code, 1, HEAP_CHUNK_SIZE);154
+ win_call_arg_imm64(code, 2, MEM_COMMIT);155
+ win_call_arg_imm64(code, 3, PAGE_EXECUTE_READWRITE);156
+ win_call_import(code, PE_IMPORT_VIRTUAL_ALLOC);157
+ win_call_end(code);158
+63159
emit_mov_reg_imm64(code, REG_RSI, kiln_heap_ptr_addr());64
- emit_store_mem_disp32(code, REG_RSI, 0, REG_RAX);160
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RBX); // heap_ptr = base161
+ emit_mov_reg_reg(code, REG_RAX, REG_RBX);162
+ emit_add_reg_imm32(code, REG_RAX, (int32_t)HEAP_CHUNK_SIZE);163
+ emit_mov_reg_imm64(code, REG_RSI, kiln_heap_limit_addr());164
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RAX); // heap_limit = base + one chunk65165
if (!is_shared()) emit_alloc_routine(code);66166
return;67167
}@@ -79,15 +179,21 @@ void heap_emit_startup(CodeBuf *code) {79179
emit_mov_reg_imm64(code, REG_RSI, kiln_heap_ptr_addr());80180
emit_store_mem_disp32(code, REG_RSI, 0, REG_RAX);81181
82
- // brk(current_break + HEAP_SIZE) -- grow once, upfront. Not checking83
- // the result: a fixed, modest 16MB request failing is not a case84
- // this project's "don't over-engineer error handling for things that85
- // can't practically happen" rule asks for defensive code around.182
+ // brk(current_break + HEAP_CHUNK_SIZE) -- grow once, upfront, to the183
+ // first chunk. Not checking the result: a modest first request184
+ // failing is not a case this project's "don't over-engineer error185
+ // handling for things that can't practically happen" rule asks for186
+ // defensive code around (later growth steps, which are far more187
+ // likely to eventually hit real memory pressure, still go through188
+ // the same unchecked brk -- see emit_alloc_routine).86189
emit_mov_reg_reg(code, REG_RDI, REG_RAX);87
- emit_add_reg_imm32(code, REG_RDI, HEAP_SIZE);190
+ emit_add_reg_imm32(code, REG_RDI, (int32_t)HEAP_CHUNK_SIZE);88191
emit_mov_reg_imm64(code, REG_RAX, 12);89192
emit_syscall(code);90193
194
+ emit_mov_reg_imm64(code, REG_RSI, kiln_heap_limit_addr());195
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RDI); // heap_limit = base + one chunk196
+91197
if (!is_shared()) emit_alloc_routine(code);92198
}93199
@@ -99,3 +205,13 @@ void heap_emit_alloc(CodeBuf *code) {99205
}100206
emit_call_back(code, alloc_routine_offset);101207
}208
+209
+void heap_emit_checkpoint(CodeBuf *code) {210
+ emit_mov_reg_imm64(code, REG_RSI, kiln_heap_ptr_addr());211
+ emit_load_mem_disp32(code, REG_RAX, REG_RSI, 0);212
+}213
+214
+void heap_emit_reset(CodeBuf *code) {215
+ emit_mov_reg_imm64(code, REG_RSI, kiln_heap_ptr_addr());216
+ emit_store_mem_disp32(code, REG_RSI, 0, REG_RDI);217
+}kiln/src/codegen/C/runtime/print_int.cmodified
+8−1@@ -429,6 +429,13 @@ void codegen_exit0(CodeBuf *code) {429429
return; // never returns -- no win_call_end needed430430
}431431
emit_mov_reg_imm64(code, REG_RDI, 0);432
- emit_mov_reg_imm64(code, REG_RAX, 60); // syscall: exit432
+ // exit_group (231), not exit (60): plain exit only terminates the433
+ // CALLING thread, leaving the process alive until every thread434
+ // exits -- invisible as long as a kiln program is single-threaded,435
+ // which was every one of them until matrix_mul's --link=shared path436
+ // started calling into libashgpu.so's OpenMP-based CPU backend437
+ // (found by actually running that: the process hung forever after438
+ // printing its correct result, waiting on ghost worker threads).439
+ emit_mov_reg_imm64(code, REG_RAX, 231);433440
emit_syscall(code);434441
}kiln/src/codegen/D/builtins/builtins.dmodified
+3−1@@ -1,7 +1,8 @@11
src/codegen/O/builtins/builtins.o: src/codegen/C/builtins/builtins.c \22
src/codegen/H/builtins/builtins.h src/codegen/H/expressions/expr.h \33
src/lexer/H/lexer.h src/codegen/H/collections/arrays.h \4
- src/codegen/H/collections/maps.h src/codegen/H/builtins/math_builtins.h \4
+ src/codegen/H/collections/maps.h src/codegen/H/collections/matrix_mul.h \5
+ src/codegen/H/builtins/math_builtins.h \56
src/codegen/H/builtins/convert_builtins.h \67
src/codegen/H/builtins/higher_order.h \78
src/codegen/H/collections/keys_values.h \@@ -15,6 +16,7 @@ src/codegen/H/expressions/expr.h:1516
src/lexer/H/lexer.h:1617
src/codegen/H/collections/arrays.h:1718
src/codegen/H/collections/maps.h:19
+src/codegen/H/collections/matrix_mul.h:1820
src/codegen/H/builtins/math_builtins.h:1921
src/codegen/H/builtins/convert_builtins.h:2022
src/codegen/H/builtins/higher_order.h:kiln/src/codegen/D/collections/matrix_mul.dadded
+30−0@@ -0,0 +1,30 @@1
+src/codegen/O/collections/matrix_mul.o: \2
+ src/codegen/C/collections/matrix_mul.c \3
+ src/codegen/H/collections/matrix_mul.h src/codegen/H/emit/emit.h \4
+ src/codegen/H/emit/value.h src/codegen/H/platform/linux_call.h \5
+ src/elf/H/elf_dynamic_call.h src/elf/H/elf_dynamic.h \6
+ src/codegen/H/emit/runtime_import_enum.h \7
+ src/codegen/H/emit/runtime_exports.def \8
+ src/codegen/H/emit/ash_gpu_import_enum.h \9
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \10
+ src/codegen/H/emit/runtime_layout.h src/codegen/H/runtime/heap.h \11
+ src/codegen/H/runtime/errors.h src/parser/H/core/parser.h \12
+ src/lexer/H/lexer.h src/parser/H/declarations/vars.h src/app/H/target.h13
+src/codegen/H/collections/matrix_mul.h:14
+src/codegen/H/emit/emit.h:15
+src/codegen/H/emit/value.h:16
+src/codegen/H/platform/linux_call.h:17
+src/elf/H/elf_dynamic_call.h:18
+src/elf/H/elf_dynamic.h:19
+src/codegen/H/emit/runtime_import_enum.h:20
+src/codegen/H/emit/runtime_exports.def:21
+src/codegen/H/emit/ash_gpu_import_enum.h:22
+src/codegen/H/emit/ash_gpu_exports.def:23
+src/elf/H/elf_writer.h:24
+src/codegen/H/emit/runtime_layout.h:25
+src/codegen/H/runtime/heap.h:26
+src/codegen/H/runtime/errors.h:27
+src/parser/H/core/parser.h:28
+src/lexer/H/lexer.h:29
+src/parser/H/declarations/vars.h:30
+src/app/H/target.h:kiln/src/codegen/D/emit/bytes.dmodified
+5−1@@ -5,7 +5,9 @@ src/codegen/O/emit/bytes.o: src/codegen/C/emit/bytes.c \55
src/codegen/H/emit/runtime_import_enum.h \66
src/codegen/H/emit/runtime_exports.def src/pe/H/pe_imports.def \77
src/elf/H/elf_dynamic_call.h src/elf/H/elf_dynamic.h \8
- src/elf/H/elf_writer.h src/app/H/target.h8
+ src/codegen/H/emit/ash_gpu_import_enum.h \9
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \10
+ src/app/H/target.h911
src/codegen/H/emit/bytes.h:1012
src/codegen/H/emit/emit.h:1113
src/codegen/H/platform/win_call.h:@@ -16,5 +18,7 @@ src/codegen/H/emit/runtime_exports.def:1618
src/pe/H/pe_imports.def:1719
src/elf/H/elf_dynamic_call.h:1820
src/elf/H/elf_dynamic.h:21
+src/codegen/H/emit/ash_gpu_import_enum.h:22
+src/codegen/H/emit/ash_gpu_exports.def:1923
src/elf/H/elf_writer.h:2024
src/app/H/target.h:kiln/src/codegen/D/emit/layout.dmodified
+5−1@@ -2,7 +2,9 @@ src/codegen/O/emit/layout.o: src/codegen/C/emit/layout.c \22
src/codegen/H/emit/layout.h src/app/H/target.h src/elf/H/elf_writer.h \33
src/codegen/H/emit/emit.h src/codegen/H/emit/runtime_layout.h \44
src/elf/H/elf_dynamic.h src/codegen/H/emit/runtime_import_enum.h \5
- src/codegen/H/emit/runtime_exports.def src/pe/H/pe_writer.h \5
+ src/codegen/H/emit/runtime_exports.def \6
+ src/codegen/H/emit/ash_gpu_import_enum.h \7
+ src/codegen/H/emit/ash_gpu_exports.def src/pe/H/pe_writer.h \68
src/pe/H/pe_imports.def src/pe/H/pe_dll_writer.h79
src/codegen/H/emit/layout.h:810
src/app/H/target.h:@@ -12,6 +14,8 @@ src/codegen/H/emit/runtime_layout.h:1214
src/elf/H/elf_dynamic.h:1315
src/codegen/H/emit/runtime_import_enum.h:1416
src/codegen/H/emit/runtime_exports.def:17
+src/codegen/H/emit/ash_gpu_import_enum.h:18
+src/codegen/H/emit/ash_gpu_exports.def:1519
src/pe/H/pe_writer.h:1620
src/pe/H/pe_imports.def:1721
src/pe/H/pe_dll_writer.h:kiln/src/codegen/D/platform/linux_call.dadded
+21−0@@ -0,0 +1,21 @@1
+src/codegen/O/platform/linux_call.o: src/codegen/C/platform/linux_call.c \2
+ src/codegen/H/platform/linux_call.h src/codegen/H/emit/emit.h \3
+ src/elf/H/elf_dynamic_call.h src/elf/H/elf_dynamic.h \4
+ src/codegen/H/emit/runtime_import_enum.h \5
+ src/codegen/H/emit/runtime_exports.def \6
+ src/codegen/H/emit/ash_gpu_import_enum.h \7
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \8
+ src/codegen/H/emit/runtime_layout.h src/codegen/H/emit/emit_internal.h \9
+ src/parser/H/declarations/vars.h10
+src/codegen/H/platform/linux_call.h:11
+src/codegen/H/emit/emit.h:12
+src/elf/H/elf_dynamic_call.h:13
+src/elf/H/elf_dynamic.h:14
+src/codegen/H/emit/runtime_import_enum.h:15
+src/codegen/H/emit/runtime_exports.def:16
+src/codegen/H/emit/ash_gpu_import_enum.h:17
+src/codegen/H/emit/ash_gpu_exports.def:18
+src/elf/H/elf_writer.h:19
+src/codegen/H/emit/runtime_layout.h:20
+src/codegen/H/emit/emit_internal.h:21
+src/parser/H/declarations/vars.h:kiln/src/codegen/D/runtime/errors.dmodified
+5−1@@ -7,7 +7,9 @@ src/codegen/O/runtime/errors.o: src/codegen/C/runtime/errors.c \77
src/codegen/H/emit/runtime_import_enum.h \88
src/codegen/H/emit/runtime_exports.def src/pe/H/pe_imports.def \99
src/elf/H/elf_dynamic_call.h src/elf/H/elf_dynamic.h \10
- src/elf/H/elf_writer.h src/app/H/target.h10
+ src/codegen/H/emit/ash_gpu_import_enum.h \11
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \12
+ src/app/H/target.h1113
src/codegen/H/runtime/errors.h:1214
src/codegen/H/emit/emit.h:1315
src/codegen/H/emit/emit_internal.h:@@ -22,5 +24,7 @@ src/codegen/H/emit/runtime_exports.def:2224
src/pe/H/pe_imports.def:2325
src/elf/H/elf_dynamic_call.h:2426
src/elf/H/elf_dynamic.h:27
+src/codegen/H/emit/ash_gpu_import_enum.h:28
+src/codegen/H/emit/ash_gpu_exports.def:2529
src/elf/H/elf_writer.h:2630
src/app/H/target.h:kiln/src/codegen/D/runtime/heap.dmodified
+5−1@@ -5,7 +5,9 @@ src/codegen/O/runtime/heap.o: src/codegen/C/runtime/heap.c \55
src/codegen/H/emit/runtime_import_enum.h \66
src/codegen/H/emit/runtime_exports.def src/pe/H/pe_imports.def \77
src/elf/H/elf_dynamic_call.h src/elf/H/elf_dynamic.h \8
- src/elf/H/elf_writer.h src/app/H/target.h8
+ src/codegen/H/emit/ash_gpu_import_enum.h \9
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \10
+ src/app/H/target.h911
src/codegen/H/runtime/heap.h:1012
src/codegen/H/emit/emit.h:1113
src/codegen/H/emit/layout.h:@@ -17,5 +19,7 @@ src/codegen/H/emit/runtime_exports.def:1719
src/pe/H/pe_imports.def:1820
src/elf/H/elf_dynamic_call.h:1921
src/elf/H/elf_dynamic.h:22
+src/codegen/H/emit/ash_gpu_import_enum.h:23
+src/codegen/H/emit/ash_gpu_exports.def:2024
src/elf/H/elf_writer.h:2125
src/app/H/target.h:kiln/src/codegen/D/runtime/print_int.dmodified
+5−1@@ -7,7 +7,9 @@ src/codegen/O/runtime/print_int.o: src/codegen/C/runtime/print_int.c \77
src/codegen/H/emit/runtime_import_enum.h \88
src/codegen/H/emit/runtime_exports.def src/pe/H/pe_imports.def \99
src/codegen/H/platform/platform_console.h src/elf/H/elf_dynamic_call.h \10
- src/elf/H/elf_dynamic.h src/elf/H/elf_writer.h src/app/H/target.h10
+ src/elf/H/elf_dynamic.h src/codegen/H/emit/ash_gpu_import_enum.h \11
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \12
+ src/app/H/target.h1113
src/codegen/H/runtime/print_int.h:1214
src/codegen/H/emit/emit.h:1315
src/codegen/H/runtime/number_format.h:@@ -24,5 +26,7 @@ src/pe/H/pe_imports.def:2426
src/codegen/H/platform/platform_console.h:2527
src/elf/H/elf_dynamic_call.h:2628
src/elf/H/elf_dynamic.h:29
+src/codegen/H/emit/ash_gpu_import_enum.h:30
+src/codegen/H/emit/ash_gpu_exports.def:2731
src/elf/H/elf_writer.h:2832
src/app/H/target.h:kiln/src/codegen/D/strings/string_alloc.dmodified
+5−1@@ -6,7 +6,9 @@ src/codegen/O/strings/string_alloc.o: \66
src/codegen/H/emit/runtime_import_enum.h \77
src/codegen/H/emit/runtime_exports.def src/pe/H/pe_imports.def \88
src/elf/H/elf_dynamic_call.h src/elf/H/elf_dynamic.h \9
- src/elf/H/elf_writer.h src/app/H/target.h9
+ src/codegen/H/emit/ash_gpu_import_enum.h \10
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \11
+ src/app/H/target.h1012
src/codegen/H/strings/string_alloc.h:1113
src/codegen/H/emit/emit.h:1214
src/codegen/H/runtime/heap.h:@@ -18,5 +20,7 @@ src/codegen/H/emit/runtime_exports.def:1820
src/pe/H/pe_imports.def:1921
src/elf/H/elf_dynamic_call.h:2022
src/elf/H/elf_dynamic.h:23
+src/codegen/H/emit/ash_gpu_import_enum.h:24
+src/codegen/H/emit/ash_gpu_exports.def:2125
src/elf/H/elf_writer.h:2226
src/app/H/target.h:kiln/src/codegen/H/collections/matrix_mul.hadded
+15−0@@ -0,0 +1,15 @@1
+#ifndef KILN_MATRIX_MUL_H2
+#define KILN_MATRIX_MUL_H3
+4
+// matrix_mul(a, b): only available under --link=shared (ideas/assigned.md5
+// -- a --link=static binary stays the hand-rolled, zero-dependency ELF/PE6
+// kiln otherwise always produces; this builtin's whole implementation is7
+// a call through the GOT into libashgpu.so, which only exists under8
+// --link=shared at all). codegen/C/builtins/builtins.c's dispatch calls9
+// this after parsing matrix_mul(...)'s arguments, same convention as10
+// every other builtin in that file -- a and b already pushed11
+// left-to-right. Rejects (compile error, not a runtime one -- link mode12
+// is known at compile time) with a clear message under --link=static.13
+void codegen_builtin_matrix_mul(void);14
+15
+#endifkiln/src/codegen/H/emit/ash_gpu_exports.defadded
+8−0@@ -0,0 +1,8 @@1
+// x-macro list of libashgpu.so's exported symbols kiln's --link=shared2
+// codegen calls (elf/C/elf_dynamic.c's executable-side GOT/hash/dynsym,3
+// alongside kiln's own runtime_exports.def entries -- see that file's4
+// header comment and ideas/assigned.md for why this is a SECOND needed5
+// library, not more entries in libkilnrt.so's own export list: this6
+// list's names come from Pyre (pyre/src/H/runtime/ash_gpu.h), a7
+// different binary kiln doesn't build.8
+ASH_GPU_EXPORT(MATRIX_MULTIPLY, "ash_gpu_matrix_multiply")kiln/src/codegen/H/emit/ash_gpu_import_enum.hadded
+15−0@@ -0,0 +1,15 @@1
+#ifndef KILN_ASH_GPU_IMPORT_ENUM_H2
+#define KILN_ASH_GPU_IMPORT_ENUM_H3
+4
+// Every libashgpu.so function kiln's --link=shared codegen calls --5
+// single list in ash_gpu_exports.def, shared by this enum and6
+// elf_dynamic.c's actual table bytes, so the two can never silently7
+// drift out of sync (same pattern as runtime_import_enum.h).8
+typedef enum {9
+#define ASH_GPU_EXPORT(name, str) ASH_GPU_IMPORT_##name,10
+#include "codegen/H/emit/ash_gpu_exports.def"11
+#undef ASH_GPU_EXPORT12
+ ASH_GPU_IMPORT_COUNT_13
+} AshGpuImport;14
+15
+#endifkiln/src/codegen/H/emit/layout.hmodified
+1−0@@ -13,6 +13,7 @@1313
uint64_t kiln_modstate_addr(void); // start of the module-state slot table (16 bytes per slot)1414
uint64_t kiln_code_base(void); // add a code offset to get an absolute address1515
uint64_t kiln_heap_ptr_addr(void);16
+uint64_t kiln_heap_limit_addr(void);1617
uint64_t kiln_try_depth_addr(void);1718
uint64_t kiln_try_handlers_addr(void);1819
kiln/src/codegen/H/emit/runtime_layout.hmodified
+6−0@@ -1,6 +1,12 @@11
#ifndef KILN_RUNTIME_LAYOUT_H22
#define KILN_RUNTIME_LAYOUT_H33
4
+// Second globals slot, right after heap_ptr: the current end of committed5
+// heap memory (the Linux brk value / the Windows committed boundary). See6
+// codegen/C/runtime/heap.c -- the bump allocator grows into fresh OS pages7
+// once heap_ptr would cross this line, instead of the old fixed 16MB cap.8
+#define KILN_HEAP_LIMIT_SIZE 89
+410
// OS-agnostic sizes for kiln's own hand-rolled attempt/handle handler stack511
// (see codegen/C/runtime/errors.c's raise routine, parser/C/statements/attempt_handle.c) -- the byte612
// layout is a kiln-internal invention, not tied to ELF or PE, so bothkiln/src/codegen/H/expressions/expr.hmodified
+9−0@@ -37,4 +37,13 @@ void codegen_apply_plus(void);3737
// Errors on strings, arrays, maps and functions.3838
void codegen_pop_and_test_truthy(void);3939
40
+// `x[index]` reads -- arrays, maps and strings all use this syntax, and41
+// which one `x` actually is can only be known at runtime. Both the base42
+// and the index must already be pushed (base deeper); pops both and43
+// pushes base[index]. Exposed (not static to expr.c) so parser/C/core/parser.c's44
+// index_assignment_statement can read through every index but the last45
+// one in a chain (`C[i][j] = val`), matching this function's own46
+// postfix() while-loop shape for the read side.47
+void codegen_index_read(void);48
+4049
#endifkiln/src/codegen/H/platform/linux_call.hadded
+38−0@@ -0,0 +1,38 @@1
+#ifndef KILN_LINUX_CALL_H2
+#define KILN_LINUX_CALL_H3
+#include "codegen/H/emit/emit.h"4
+#include "elf/H/elf_dynamic_call.h"5
+6
+// SysV x86-64 calling convention wrapper for calling into a SECOND7
+// needed library's real, normally-compiled function (libashgpu.so --8
+// see elf/H/elf_dynamic.h) -- new territory the same way win_call.h was9
+// for libkilnrt.dll's imports, but for a different reason: kiln's OWN10
+// runtime routines (heap_alloc, print_top, ...), the only things11
+// elf_dynamic_call was ever used for before this, are THEMSELVES12
+// kiln's own hand-rolled machine code, which never uses an13
+// alignment-sensitive instruction (emit_sse.h's movsd/movq never need14
+// it) -- so kiln's codegen has never had to track RSP's 16-byte15
+// alignment at a CALL anywhere, the same gap win_call.h's own header16
+// comment describes for Windows. ash_gpu_matrix_multiply is real17
+// gcc/g++-compiled code (SSE/AVX doubles, OpenMP), which DOES rely on18
+// that alignment -- calling it misaligned risks a `movaps`-class fault,19
+// not a wrong answer.20
+//21
+// Same technique as win_call_begin/end (dynamic `and rsp, -16` plus a22
+// saved-RSP scratch slot), reusing that exact slot23
+// (win_call_scratch_offset(), parser/H/declarations/vars.h) rather than24
+// reserving a second one: a single kiln target is always Linux or25
+// Windows, never both, so the slot is never needed for its Windows26
+// purpose and this purpose in the same compiled program. No shadow27
+// space here (SysV has none) -- that's the only difference from28
+// win_call_begin.29
+// Neither of these touches RDI/RSI/RDX/RCX -- only RSP and the scratch30
+// slot -- so the caller loads the callee's C-ABI arguments into those31
+// registers in whatever order is convenient, either side of32
+// linux_abi_call_begin(), then calls ash_gpu_dynamic_call() before33
+// linux_abi_call_end(). See codegen/C/collections/matrix_mul.c for the34
+// one caller so far.35
+void linux_abi_call_begin(CodeBuf *code);36
+void linux_abi_call_end(CodeBuf *code);37
+38
+#endifkiln/src/codegen/H/runtime/heap.hmodified
+42−9@@ -3,17 +3,21 @@33
#include "codegen/H/emit/emit.h"44
55
// A bump allocator, no free (matches ashvm's own lack of GC -- neither6
-// engine reclaims memory during a run, only at process exit). Backed by7
-// `brk` rather than `mmap`: brk takes exactly one argument (the new break8
-// address, in RDI), while mmap needs six, three of which (R10/R8/R9) kiln9
-// has deliberately never needed until now and would require extending10
-// emit.c's register handling to reach (those are all >= 8, needing11
-// REX.R/X/B bits this project's encoders don't set). brk sidesteps that12
-// entirely for a one-time, fixed-size arena grab.6
+// engine reclaims memory during a run, only at process exit). Grows in7
+// fixed-size chunks instead of one fixed-size arena: Linux extends the8
+// break with `brk` (one argument, the new break address, in RDI -- mmap9
+// needs six, three of which, R10/R8/R9, kiln has deliberately never10
+// needed and would require extending emit.c's register handling to11
+// reach); Windows reserves a large virtual range up front (no physical12
+// memory committed) and commits one more chunk at a time via13
+// VirtualAlloc. Either way the arena is one contiguous region for its14
+// whole lifetime, so a pointer handed out by an earlier allocation is15
+// never invalidated by a later growth step.1316
1417
// Emitted once, at the very start of the generated program (before the15
-// top-level statement loop): grabs a 16MB arena and records its start as16
-// the heap bump pointer at KILN_GLOBALS_ADDR (see elf/H/elf_writer.h).18
+// top-level statement loop): grabs the first chunk and records its start19
+// as the heap bump pointer at KILN_GLOBALS_ADDR (see elf/H/elf_writer.h),20
+// plus the committed limit the alloc routine grows from (see heap.c).1721
void heap_emit_startup(CodeBuf *code);1822
1923
// Plan B: emits ONLY the alloc routine's body (no arena grab) -- used by@@ -35,4 +39,33 @@ int heap_alloc_routine_offset(void);3539
// every one of its ~27 call sites.3640
void heap_emit_alloc(CodeBuf *code);3741
42
+// Arena-style checkpoint/reset for short-lived temporaries (built for43
+// Pyre's own lowering, e.g. a tight inner loop's scratch buffers that44
+// never outlive one iteration): rewinding the bump pointer reclaims45
+// everything allocated since the checkpoint without any OS call, at the46
+// cost of being a *rewind*, not a real free -- it is only sound when the47
+// caller (the code generator emitting the checkpoint/reset pair, not48
+// these functions) can prove nothing allocated in between is reachable49
+// after the reset. Nothing here checks that proof; a reset after a50
+// pointer escaped (stored outside the region, returned, captured by a51
+// closure, handed to a function that might retain it) corrupts that52
+// pointer's memory the moment something else allocates over it. Not53
+// wired into general `each`/`during` loop bodies for exactly that reason54
+// -- arbitrary Ash code can make a value escape in ways kiln's single-55
+// pass compiler has no general way to rule out (that is goals.md item56
+// 1's full ownership/borrow story, not this).57
+//58
+// Two instructions each, used directly at the call site rather than59
+// through heap_emit_alloc's shared-routine-plus-CALL machinery -- that60
+// indirection earns its keep at ~27 call sites, not at the handful a61
+// narrowly-scoped optimization like this one will ever have.62
+63
+// No input; returns the current heap_ptr (the checkpoint value) in RAX,64
+// clobbers RSI.65
+void heap_emit_checkpoint(CodeBuf *code);66
+// Checkpoint value (as returned by heap_emit_checkpoint) in RDI; no67
+// return value, clobbers RSI. Sets heap_ptr back to it -- heap_limit is68
+// untouched, since rewinding never needs to give committed pages back.69
+void heap_emit_reset(CodeBuf *code);70
+3871
#endifkiln/src/elf/C/elf_dynamic.cmodified
+33−10@@ -22,10 +22,30 @@ static const char *const runtime_names[] = {2222
#undef RUNTIME_EXPORT2323
};2424
25
+static const char *const ash_gpu_names[] = {26
+#define ASH_GPU_EXPORT(name, str) str,27
+#include "codegen/H/emit/ash_gpu_exports.def"28
+#undef ASH_GPU_EXPORT29
+};30
+31
+// Every imported symbol's name, kiln's own RUNTIME_IMPORT_COUNT_ first32
+// (indices [0, RUNTIME_IMPORT_COUNT_)) then Pyre's ASH_GPU_IMPORT_COUNT_33
+// (indices [RUNTIME_IMPORT_COUNT_, KILN_DYN_TOTAL_IMPORTS)) -- one34
+// combined hash/dynsym/GOT/rela table covering both needed libraries,35
+// since ld.so resolves a symbol name against ALL of them regardless of36
+// which table region declared it (see elf_dynamic.h's header comment).37
+static const char *import_name(unsigned i) {38
+ return i < (unsigned)RUNTIME_IMPORT_COUNT_ ? runtime_names[i] : ash_gpu_names[i - (unsigned)RUNTIME_IMPORT_COUNT_];39
+}40
+2541
uint64_t elf_dynamic_got_addr(RuntimeImport which) {2642
return KILN_DYN_GOT_ADDR + 8ULL * (uint64_t)which;2743
}2844
45
+uint64_t ash_gpu_dynamic_got_addr(AshGpuImport which) {46
+ return KILN_DYN_GOT_ADDR + 8ULL * ((uint64_t)RUNTIME_IMPORT_COUNT_ + (uint64_t)which);47
+}48
+2949
static void put8(CodeBuf *f, uint8_t v) {3050
if (f->count >= f->capacity) {3151
f->capacity = f->capacity == 0 ? 256 : f->capacity * 2;@@ -56,10 +76,11 @@ static uint32_t elf_hash(const char *name) {5676
int elf_write_dynamic_executable(const char *path, const CodeBuf *machine_code) {5777
// ---- pre-compute .dynstr layout (written after the code, but its5878
// internal offsets are needed now for dynsym/dynamic/rela) ----59
- uint32_t dynstr_off[RUNTIME_IMPORT_COUNT_];79
+ uint32_t dynstr_off[KILN_DYN_TOTAL_IMPORTS];6080
uint32_t cursor = 1; // leading empty string at offset 061
- for (int i = 0; i < RUNTIME_IMPORT_COUNT_; i++) { dynstr_off[i] = cursor; cursor += (uint32_t)strlen(runtime_names[i]) + 1; }81
+ for (unsigned i = 0; i < KILN_DYN_TOTAL_IMPORTS; i++) { dynstr_off[i] = cursor; cursor += (uint32_t)strlen(import_name(i)) + 1; }6282
uint32_t soname_off = cursor; cursor += (uint32_t)strlen(KILN_LIBKILNRT_SONAME) + 1;83
+ uint32_t ashgpu_soname_off = cursor; cursor += (uint32_t)strlen(KILN_ASHGPU_SONAME) + 1;6384
uint32_t rpath_off = cursor; cursor += (uint32_t)strlen(KILN_RPATH) + 1;6485
uint32_t dynstr_size = cursor;6586
@@ -111,9 +132,9 @@ int elf_write_dynamic_executable(const char *path, const CodeBuf *machine_code)111132
// ---- .hash (SysV DT_HASH format: nbucket, nchain, bucket[], chain[]) ----112133
uint32_t buckets[KILN_DYN_NBUCKET]; for (unsigned i = 0; i < KILN_DYN_NBUCKET; i++) buckets[i] = 0;113134
uint32_t chains[KILN_DYN_NSYMS]; for (unsigned i = 0; i < KILN_DYN_NSYMS; i++) chains[i] = 0;114
- for (int i = 0; i < RUNTIME_IMPORT_COUNT_; i++) {115
- uint32_t sym_index = (uint32_t)i + 1;116
- uint32_t b = elf_hash(runtime_names[i]) % KILN_DYN_NBUCKET;135
+ for (unsigned i = 0; i < KILN_DYN_TOTAL_IMPORTS; i++) {136
+ uint32_t sym_index = i + 1;137
+ uint32_t b = elf_hash(import_name(i)) % KILN_DYN_NBUCKET;117138
chains[sym_index] = buckets[b];118139
buckets[b] = sym_index;119140
}@@ -126,7 +147,7 @@ int elf_write_dynamic_executable(const char *path, const CodeBuf *machine_code)126147
// ld.so to go resolve it against a NEEDED library rather than treat127148
// it as locally defined) ----128149
put32(&f, 0); put8(&f, 0); put8(&f, 0); put16(&f, 0); put64(&f, 0); put64(&f, 0); // null symbol129
- for (int i = 0; i < RUNTIME_IMPORT_COUNT_; i++) {150
+ for (unsigned i = 0; i < KILN_DYN_TOTAL_IMPORTS; i++) {130151
put32(&f, dynstr_off[i]); // st_name131152
put8(&f, (1 << 4) | 2); // st_info: STB_GLOBAL | STT_FUNC132153
put8(&f, 0); // st_other@@ -140,8 +161,9 @@ int elf_write_dynamic_executable(const char *path, const CodeBuf *machine_code)140161
141162
// ---- .dynamic ----142163
#define DT(tag, val) do { put64(&f, (uint64_t)(tag)); put64(&f, (uint64_t)(val)); } while (0)143
- DT(1, soname_off); // DT_NEEDED (index into .dynstr)144
- DT(15, rpath_off); // DT_RPATH164
+ DT(1, soname_off); // DT_NEEDED (index into .dynstr) -- libkilnrt.so165
+ DT(1, ashgpu_soname_off); // DT_NEEDED -- libashgpu.so (ideas/assigned.md)166
+ DT(15, rpath_off); // DT_RPATH -- shared by both NEEDED entries (libashgpu.so installs to the same rpath)145167
DT(4, KILN_DYN_HASH_ADDR); // DT_HASH146168
DT(5, dynstr_addr); // DT_STRTAB (after the code -- depends on its length)147169
DT(6, KILN_DYN_SYMTAB_ADDR); // DT_SYMTAB@@ -154,7 +176,7 @@ int elf_write_dynamic_executable(const char *path, const CodeBuf *machine_code)154176
#undef DT155177
156178
// ---- .rela.dyn: one R_X86_64_GLOB_DAT per GOT slot ----157
- for (int i = 0; i < RUNTIME_IMPORT_COUNT_; i++) {179
+ for (unsigned i = 0; i < KILN_DYN_TOTAL_IMPORTS; i++) {158180
uint64_t sym_index = (uint64_t)i + 1;159181
uint64_t r_info = (sym_index << 32) | 6ULL; // R_X86_64_GLOB_DAT = 6160182
put64(&f, KILN_DYN_GOT_ADDR + 8ULL * (uint64_t)i); // r_offset@@ -168,8 +190,9 @@ int elf_write_dynamic_executable(const char *path, const CodeBuf *machine_code)168190
169191
// ---- .dynstr tail ----170192
put8(&f, 0); // offset 0 = empty string171
- for (int i = 0; i < RUNTIME_IMPORT_COUNT_; i++) put_str(&f, runtime_names[i]);193
+ for (unsigned i = 0; i < KILN_DYN_TOTAL_IMPORTS; i++) put_str(&f, import_name(i));172194
put_str(&f, KILN_LIBKILNRT_SONAME);195
+ put_str(&f, KILN_ASHGPU_SONAME);173196
put_str(&f, KILN_RPATH);174197
175198
FILE *out = fopen(path, "wb");kiln/src/elf/C/elf_dynamic_call.cmodified
+4−0@@ -3,3 +3,7 @@33
void elf_dynamic_call(CodeBuf *code, RuntimeImport which) {44
emit_call_abs32(code, (uint32_t)elf_dynamic_got_addr(which));55
}6
+7
+void ash_gpu_dynamic_call(CodeBuf *code, AshGpuImport which) {8
+ emit_call_abs32(code, (uint32_t)ash_gpu_dynamic_got_addr(which));9
+}kiln/src/elf/D/elf_dynamic.dmodified
+5−1@@ -1,10 +1,14 @@11
src/elf/O/elf_dynamic.o: src/elf/C/elf_dynamic.c src/elf/H/elf_dynamic.h \22
src/codegen/H/emit/emit.h src/codegen/H/emit/runtime_import_enum.h \3
- src/codegen/H/emit/runtime_exports.def src/elf/H/elf_writer.h \3
+ src/codegen/H/emit/runtime_exports.def \4
+ src/codegen/H/emit/ash_gpu_import_enum.h \5
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \46
src/codegen/H/emit/runtime_layout.h57
src/elf/H/elf_dynamic.h:68
src/codegen/H/emit/emit.h:79
src/codegen/H/emit/runtime_import_enum.h:810
src/codegen/H/emit/runtime_exports.def:11
+src/codegen/H/emit/ash_gpu_import_enum.h:12
+src/codegen/H/emit/ash_gpu_exports.def:913
src/elf/H/elf_writer.h:1014
src/codegen/H/emit/runtime_layout.h:kiln/src/elf/D/elf_dynamic_call.dmodified
+5−1@@ -1,12 +1,16 @@11
src/elf/O/elf_dynamic_call.o: src/elf/C/elf_dynamic_call.c \22
src/elf/H/elf_dynamic_call.h src/elf/H/elf_dynamic.h \33
src/codegen/H/emit/emit.h src/codegen/H/emit/runtime_import_enum.h \4
- src/codegen/H/emit/runtime_exports.def src/elf/H/elf_writer.h \4
+ src/codegen/H/emit/runtime_exports.def \5
+ src/codegen/H/emit/ash_gpu_import_enum.h \6
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \57
src/codegen/H/emit/runtime_layout.h68
src/elf/H/elf_dynamic_call.h:79
src/elf/H/elf_dynamic.h:810
src/codegen/H/emit/emit.h:911
src/codegen/H/emit/runtime_import_enum.h:1012
src/codegen/H/emit/runtime_exports.def:13
+src/codegen/H/emit/ash_gpu_import_enum.h:14
+src/codegen/H/emit/ash_gpu_exports.def:1115
src/elf/H/elf_writer.h:1216
src/codegen/H/emit/runtime_layout.h:kiln/src/elf/D/elf_so_writer.dmodified
+5−1@@ -1,7 +1,9 @@11
src/elf/O/elf_so_writer.o: src/elf/C/elf_so_writer.c \22
src/elf/H/elf_so_writer.h src/elf/H/elf_dynamic.h \33
src/codegen/H/emit/emit.h src/codegen/H/emit/runtime_import_enum.h \4
- src/codegen/H/emit/runtime_exports.def src/elf/H/elf_writer.h \4
+ src/codegen/H/emit/runtime_exports.def \5
+ src/codegen/H/emit/ash_gpu_import_enum.h \6
+ src/codegen/H/emit/ash_gpu_exports.def src/elf/H/elf_writer.h \57
src/codegen/H/emit/runtime_layout.h src/codegen/H/runtime/heap.h \68
src/codegen/H/emit/bytes.h src/codegen/H/strings/string_alloc.h \79
src/codegen/H/runtime/print_int.h src/codegen/H/runtime/errors.h \@@ -11,6 +13,8 @@ src/elf/H/elf_dynamic.h:1113
src/codegen/H/emit/emit.h:1214
src/codegen/H/emit/runtime_import_enum.h:1315
src/codegen/H/emit/runtime_exports.def:16
+src/codegen/H/emit/ash_gpu_import_enum.h:17
+src/codegen/H/emit/ash_gpu_exports.def:1418
src/elf/H/elf_writer.h:1519
src/codegen/H/emit/runtime_layout.h:1620
src/codegen/H/runtime/heap.h:kiln/src/elf/H/elf_dynamic.hmodified
+27−6@@ -2,6 +2,7 @@22
#define KILN_ELF_DYNAMIC_H33
#include "codegen/H/emit/emit.h"44
#include "codegen/H/emit/runtime_import_enum.h"5
+#include "codegen/H/emit/ash_gpu_import_enum.h"56
#include "elf/H/elf_writer.h"67
78
// Layout for a --link=shared kiln executable: real dynamic linking, the@@ -20,10 +21,24 @@2021
// same validated structure, not a first attempt.2122
#define KILN_RPATH "/usr/local/lib/kiln"2223
#define KILN_LIBKILNRT_SONAME "libkilnrt.so"24
+// ideas/assigned.md: a SECOND needed library, Pyre's compute runtime --25
+// kept as a parallel DT_NEEDED next to libkilnrt.so's, never merged into26
+// it (libkilnrt.so is kiln's own hand-rolled routines; libashgpu.so is a27
+// normal gcc/g++-built shared object this project doesn't build). ld.so28
+// resolves each imported symbol name against ALL needed libraries, so29
+// the two don't need separate hash/dynsym/GOT regions -- one combined30
+// table (KILN_DYN_TOTAL_IMPORTS below) covering both libraries' symbols31
+// is correct and the much smaller change.32
+#define KILN_ASHGPU_SONAME "libashgpu.so"2333
#define KILN_INTERP_STR "/lib64/ld-linux-x86-64.so.2"2434
35
+// Every symbol this executable needs resolved, from EITHER needed36
+// library -- kiln's own (RUNTIME_IMPORT_COUNT_, from libkilnrt.so) then37
+// Pyre's (ASH_GPU_IMPORT_COUNT_, from libashgpu.so), in that order.38
+#define KILN_DYN_TOTAL_IMPORTS ((unsigned)RUNTIME_IMPORT_COUNT_ + (unsigned)ASH_GPU_IMPORT_COUNT_)39
+2540
// nsyms includes the mandatory null 0th symbol.26
-#define KILN_DYN_NSYMS ((unsigned)RUNTIME_IMPORT_COUNT_ + 1u)41
+#define KILN_DYN_NSYMS (KILN_DYN_TOTAL_IMPORTS + 1u)2742
#define KILN_DYN_NBUCKET KILN_DYN_NSYMS2843
2944
// Section order: everything whose address codegen (or another fixed-size@@ -49,27 +64,33 @@4964
// symbol NAME lengths, exactly like pe_writer.h's IAT -- codegen can5065
// address a slot as a compile-time constant.5166
#define KILN_DYN_GOT_ADDR (KILN_DYN_SYMTAB_ADDR + KILN_DYN_SYMTAB_SIZE)52
-#define KILN_DYN_GOT_SIZE (8u * (unsigned)RUNTIME_IMPORT_COUNT_)67
+#define KILN_DYN_GOT_SIZE (8u * KILN_DYN_TOTAL_IMPORTS)5368
5469
#define KILN_DYN_ENTRY_SIZE 16u // Elf64_Dyn: {Elf64_Sxword d_tag; union d_val/d_ptr;}55
-#define KILN_DYN_TABLE_NENT_TOTAL 10u // NEEDED,RPATH,HASH,STRTAB,SYMTAB,STRSZ,SYMENT,RELA,RELASZ,RELAENT,NULL(11, see .c)70
+// NEEDED(libkilnrt.so),NEEDED(libashgpu.so),RPATH,HASH,STRTAB,SYMTAB,STRSZ,SYMENT,RELA,RELASZ,RELAENT,NULL(12, see .c)71
+#define KILN_DYN_TABLE_NENT_TOTAL 11u5672
#define KILN_DYN_TABLE_ADDR (KILN_DYN_GOT_ADDR + KILN_DYN_GOT_SIZE)5773
#define KILN_DYN_TABLE_SIZE ((KILN_DYN_TABLE_NENT_TOTAL + 1u) * KILN_DYN_ENTRY_SIZE) // +1 for the DT_NULL terminator5874
5975
#define KILN_DYN_RELAENT_SIZE 24u // Elf64_Rela: {Elf64_Addr; Elf64_Xword info; Elf64_Sxword addend}6076
#define KILN_DYN_RELA_ADDR (KILN_DYN_TABLE_ADDR + KILN_DYN_TABLE_SIZE)61
-#define KILN_DYN_RELA_SIZE (KILN_DYN_RELAENT_SIZE * (unsigned)RUNTIME_IMPORT_COUNT_)77
+#define KILN_DYN_RELA_SIZE (KILN_DYN_RELAENT_SIZE * KILN_DYN_TOTAL_IMPORTS)6278
6379
#define KILN_DYN_GLOBALS_ADDR (KILN_DYN_RELA_ADDR + KILN_DYN_RELA_SIZE)64
-#define KILN_DYN_GLOBALS_SIZE (8 + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_ARGV_SIZE + KILN_MODSTATE_SIZE)65
-#define KILN_DYN_MODSTATE_ADDR (KILN_DYN_GLOBALS_ADDR + 8 + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_ARGV_SIZE)80
+#define KILN_DYN_GLOBALS_SIZE (8 + KILN_HEAP_LIMIT_SIZE + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_ARGV_SIZE + KILN_MODSTATE_SIZE)81
+#define KILN_DYN_MODSTATE_ADDR (KILN_DYN_GLOBALS_ADDR + 8 + KILN_HEAP_LIMIT_SIZE + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_ARGV_SIZE)6682
#define KILN_DYN_CODE_START_OFFSET (KILN_DYN_GLOBALS_ADDR + KILN_DYN_GLOBALS_SIZE - KILN_LOAD_BASE)6783
6884
// Absolute VA of `which`'s GOT slot -- loaded via a compile-time-known6985
// immediate (kiln's executables are always fixed-base, same reasoning as7086
// pe_import_addr) -- deref it and call through the result.7187
uint64_t elf_dynamic_got_addr(RuntimeImport which);7288
89
+// Same idea, for the second needed library's imports -- their GOT slots90
+// sit right after kiln's own RUNTIME_IMPORT_COUNT_ slots in the same91
+// combined GOT region (KILN_DYN_TOTAL_IMPORTS above).92
+uint64_t ash_gpu_dynamic_got_addr(AshGpuImport which);93
+7394
// Writes `machine_code` as a dynamically-linked ELF64 executable at7495
// `path`, expecting libkilnrt.so to be installed at KILN_RPATH by the7596
// time this program is run (see elf_so_writer.h). Returns 0 on success.kiln/src/elf/H/elf_dynamic_call.hmodified
+8−0@@ -14,4 +14,12 @@1414
// falls through -- harmless, and not worth a second code path here.1515
void elf_dynamic_call(CodeBuf *code, RuntimeImport which);1616
17
+// Same mechanism, for the second needed library (libashgpu.so, see18
+// elf_dynamic.h's header comment) -- `which`'s own SysV C calling19
+// convention applies past the call (args in RDI/RSI/RDX/RCX/R8/R9),20
+// unlike kiln's own runtime routines above which each define their own21
+// ad hoc register convention; callers must marshal args accordingly22
+// before this, see codegen/C/collections/matrix_mul.c.23
+void ash_gpu_dynamic_call(CodeBuf *code, AshGpuImport which);24
+1725
#endifkiln/src/elf/H/elf_writer.hmodified
+3−2@@ -34,9 +34,10 @@3434
// inside it (see codegen/C/io/argv_builtin.c).3535
#define KILN_ARGV_SIZE 83636
37
-#define KILN_GLOBALS_SIZE (8 + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_ARGV_SIZE + KILN_MODSTATE_SIZE) // heap_ptr, try_depth, handler stack, argv, module state37
+#define KILN_GLOBALS_SIZE (8 + KILN_HEAP_LIMIT_SIZE + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_ARGV_SIZE + KILN_MODSTATE_SIZE) // heap_ptr, heap_limit, try_depth, handler stack, argv, module state3838
#define KILN_GLOBALS_ADDR (KILN_LOAD_BASE + KILN_EHDR_SIZE + KILN_PHDR_SIZE)39
-#define KILN_TRY_DEPTH_ADDR (KILN_GLOBALS_ADDR + 8)39
+#define KILN_HEAP_LIMIT_ADDR (KILN_GLOBALS_ADDR + 8)40
+#define KILN_TRY_DEPTH_ADDR (KILN_HEAP_LIMIT_ADDR + KILN_HEAP_LIMIT_SIZE)4041
#define KILN_TRY_HANDLERS_ADDR (KILN_TRY_DEPTH_ADDR + KILN_TRY_DEPTH_SIZE)4142
#define KILN_ARGV_ADDR (KILN_TRY_HANDLERS_ADDR + KILN_TRY_HANDLERS_SIZE)4243
#define KILN_MODSTATE_ADDR (KILN_ARGV_ADDR + KILN_ARGV_SIZE)kiln/src/parser/C/core/parser.cmodified
+16−3@@ -135,15 +135,28 @@ static void ref_store_regs(VarRef r) {135135
emit_store_mem_disp32(code, REG_RBP, var_slot_payload_offset(r.index), REG_RAX);136136
}137137
138
-// IDENT[index] = expr ; -- writes through an existing array OR map139
-// variable (which one is only known at runtime -- see expr.c's postfix()140
-// for the read-side version of this same dispatch).138
+// IDENT[index] = expr ; / IDENT[i1][i2]...[iN] = expr ; -- writes139
+// through an existing array OR map variable (which one is only known at140
+// runtime -- see expr.c's postfix() for the read-side version of this141
+// same dispatch). Every index but the last is a READ (codegen_index_read,142
+// the same primitive postfix() chains on the read side): C[i][j] = val143
+// first reads C[i] to get the inner array, then stores val into THAT144
+// array at [j] -- it never needs to write the outer container back,145
+// since the inner array it just read out is the very same heap object146
+// C[i] already points at (arrays/maps are reference values here, not147
+// copied on read).141148
static void index_assignment_statement(Token id) {142149
ref_push(resolve_ref_or_error(id));143150
144151
advance_token(); // consume '['145152
codegen_expression();146153
expect(TOKEN_RBRACKET, "expected ']' after index");154
+ while (current.type == TOKEN_LBRACKET) {155
+ codegen_index_read(); // stack: [base] -> [base[index]]156
+ advance_token(); // consume '['157
+ codegen_expression();158
+ expect(TOKEN_RBRACKET, "expected ']' after index");159
+ }147160
expect(TOKEN_EQUAL, "expected '=' after index");148161
codegen_expression();149162
expect(TOKEN_SEMICOLON, "expected ';' after assignment");kiln/src/pe/H/pe_writer.hmodified
+3−2@@ -29,13 +29,14 @@2929
3030
#define KILN_PE_GLOBALS_ADDR (KILN_PE_IMAGE_BASE + KILN_PE_SECTION_RVA)3131
#define KILN_PE_HEAP_PTR_ADDR (KILN_PE_GLOBALS_ADDR + 0)32
-#define KILN_PE_TRY_DEPTH_ADDR (KILN_PE_GLOBALS_ADDR + 8)32
+#define KILN_PE_HEAP_LIMIT_ADDR (KILN_PE_GLOBALS_ADDR + 8)33
+#define KILN_PE_TRY_DEPTH_ADDR (KILN_PE_HEAP_LIMIT_ADDR + KILN_HEAP_LIMIT_SIZE)3334
#define KILN_PE_TRY_HANDLERS_ADDR (KILN_PE_TRY_DEPTH_ADDR + KILN_TRY_DEPTH_SIZE)3435
// No argv slot: GetCommandLineA() can be called at any point, unlike3536
// Linux's argv() which has to snapshot the kernel's original RSP before3637
// anything else touches it (see codegen/C/io/argv_builtin.c).3738
#define KILN_PE_MODSTATE_ADDR (KILN_PE_TRY_HANDLERS_ADDR + KILN_TRY_HANDLERS_SIZE)38
-#define KILN_PE_GLOBALS_SIZE (8 + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_MODSTATE_SIZE)39
+#define KILN_PE_GLOBALS_SIZE (8 + KILN_HEAP_LIMIT_SIZE + KILN_TRY_DEPTH_SIZE + KILN_TRY_HANDLERS_SIZE + KILN_MODSTATE_SIZE)3940
4041
// Every kernel32.dll function kiln's Windows codegen calls -- single list4142
// in pe_imports.def, shared by this enum and pe_writer.c's actual tablekiln/tests/cases/collections/chained_index_assign.ashadded
+38−0@@ -0,0 +1,38 @@1
+// C[i][j] = val -- nested index mutation without a manual intermediate2
+// extraction (local row = C[i]; row[j] = val; C[i] = row;).3
+local a = [[1, 2], [3, 4]];4
+a[0][1] = 99;5
+say a;6
+7
+// three levels deep8
+local b = [[[1, 2], [3, 4]], [[5, 6], [7, 8]]];9
+b[1][0][1] = 999;10
+say b;11
+12
+// matrix-multiply-shaped accumulation: a 2D result built with plain `=`,13
+// not a manual row-then-store dance14
+local n = 3;15
+local c = [[0, 0, 0], [0, 0, 0], [0, 0, 0]];16
+local i = 0;17
+during (i < n) {18
+ local j = 0;19
+ during (j < n) {20
+ c[i][j] = (i * n) + j;21
+ j = j + 1;22
+ }23
+ i = i + 1;24
+}25
+say c;26
+27
+// mixed array/map chain28
+local m = {"row": [1, 2, 3]};29
+m["row"][1] = 42;30
+say m["row"];31
+32
+// plain single-level assignment still works (no regression)33
+local arr = [1, 2, 3];34
+arr[0] = 7;35
+say arr;36
+37
+// a string anywhere in the chain is still immutable38
+attempt { local s = ["hi"]; s[0][0] = "x"; } handle (e) { say "immutable"; }kiln/tests/cases/memory/heap_growth.ashadded
+11−0@@ -0,0 +1,11 @@1
+// Allocates well past the old fixed 16MB bump-heap cap (each iteration2
+// leaks a fresh 4-element array, nothing is ever freed) to exercise the3
+// allocator's dynamic growth instead of segfaulting partway through.4
+local total = 0;5
+local i = 0;6
+during (i < 2000000) {7
+ local arr = [i, i + 1, i + 2, i + 3];8
+ total = total + arr[0];9
+ i = i + 1;10
+}11
+say total;kiln/tests/scripts/shared_link_check.shadded
+85−0@@ -0,0 +1,85 @@1
+#!/usr/bin/env bash2
+# kiln/tests/scripts/shared_link_check.sh3
+#4
+# --link=shared had NO automated coverage at all before this (run_regression.sh5
+# only ever compiles with kiln's default --link=static) -- this fills that6
+# gap generally, and specifically covers matrix_mul (codegen/C/collections/matrix_mul.c,7
+# elf/H/elf_dynamic.h's second DT_NEEDED for libashgpu.so), the one builtin8
+# that only exists under --link=shared at all.9
+#10
+# 1. Rebuilds kiln, libkilnrt.so, and libashgpu.so, installing both at11
+# kiln's hardcoded rpath (KILN_RPATH, elf/H/elf_dynamic.h) -- there is12
+# no env-var override, so this writes to /usr/local/lib/kiln for real,13
+# same as doing it by hand.14
+# 2. Compiles a sample of ordinary test cases BOTH --link=static and15
+# --link=shared and diffs their output -- proving the two modes agree16
+# wherever both can run a program, which also means the second17
+# DT_NEEDED and the combined GOT/hash/dynsym table didn't disturb18
+# kiln's own existing runtime imports (heap_alloc, print_top, ...).19
+# 3. Compiles and runs a small matrix_mul program (inline here, not under20
+# tests/cases/ -- run_regression.sh's tests/cases/*/*.ash glob would21
+# pick it up and "COMPILE FAIL" with kiln's default --link=static,22
+# which is correct behavior but not what that suite is for) under23
+# --link=shared, and separately confirms --link=static rejects it with24
+# a clear compile error instead of silently ignoring it.25
+set -uo pipefail26
+cd "$(dirname "$0")/../.." # -> kiln/27
+28
+make -s || { echo "build failed"; exit 1; }29
+../bin/kiln --emit-runtime -o /usr/local/lib/kiln/libkilnrt.so || { echo "emit-runtime failed"; exit 1; }30
+(cd ../pyre && make -s install) || { echo "pyre install failed"; exit 1; }31
+32
+OUTDIR=/tmp/kiln_shared_check33
+rm -rf "$OUTDIR"34
+mkdir -p "$OUTDIR"35
+36
+fail=037
+38
+for f in tests/cases/basics/*.ash tests/cases/collections/*.ash tests/cases/strings/*.ash; do39
+ name=$(basename "$f" .ash)40
+ static_out="$OUTDIR/${name}.static.out"41
+ shared_out="$OUTDIR/${name}.shared.out"42
+ ../bin/kiln "$f" -o "$OUTDIR/${name}.static" --link=static > "$static_out" 2>&143
+ ../bin/kiln "$f" -o "$OUTDIR/${name}.shared" --link=shared > "$shared_out" 2>&144
+ if [[ "$name" == "input_builtin" ]]; then45
+ printf 'World\n5\n' | timeout 10 "$OUTDIR/${name}.static" >> "$static_out" 2>&146
+ printf 'World\n5\n' | timeout 10 "$OUTDIR/${name}.shared" >> "$shared_out" 2>&147
+ else48
+ timeout 10 "$OUTDIR/${name}.static" >> "$static_out" 2>&149
+ timeout 10 "$OUTDIR/${name}.shared" >> "$shared_out" 2>&150
+ fi51
+ if ! diff -q "$static_out" "$shared_out" > /dev/null; then52
+ echo "MISMATCH: $name (static vs shared)"53
+ diff "$static_out" "$shared_out"54
+ fail=155
+ fi56
+done57
+58
+cat > "$OUTDIR/matrix_mul_shared.ash" << 'EOF'59
+local a = [[1, 2], [3, 4]];60
+local b = [[5, 6], [7, 8]];61
+say matrix_mul(a, b);62
+EOF63
+64
+../bin/kiln "$OUTDIR/matrix_mul_shared.ash" -o "$OUTDIR/matrix_mul_shared.bin" --link=shared > "$OUTDIR/matrix_mul_shared.out" 2>&165
+timeout 10 "$OUTDIR/matrix_mul_shared.bin" >> "$OUTDIR/matrix_mul_shared.out" 2>&166
+if ! grep -qF '[[19, 22], [43, 50]]' "$OUTDIR/matrix_mul_shared.out"; then67
+ echo "MISMATCH: matrix_mul under --link=shared did not produce the expected result"68
+ cat "$OUTDIR/matrix_mul_shared.out"69
+ fail=170
+fi71
+72
+if ../bin/kiln "$OUTDIR/matrix_mul_shared.ash" -o "$OUTDIR/matrix_mul_static.bin" --link=static > "$OUTDIR/matrix_mul_static_reject.out" 2>&1; then73
+ echo "MISMATCH: matrix_mul under --link=static should have been rejected at compile time, but compiled"74
+ fail=175
+elif ! grep -q "requires --link=shared" "$OUTDIR/matrix_mul_static_reject.out"; then76
+ echo "MISMATCH: matrix_mul --link=static rejection message changed unexpectedly"77
+ cat "$OUTDIR/matrix_mul_static_reject.out"78
+ fail=179
+fi80
+81
+if [[ $fail -eq 0 ]]; then82
+ echo "shared_link_check: OK"83
+else84
+ exit 185
+fipyre/Makefileadded
+160−0@@ -0,0 +1,160 @@1
+# Pyre (ash-gpu): same C/H/D/O mirrored-tree convention as kiln/ashvm2
+# (AGENTS.md section 1), with feature folders nested inside each tree3
+# (section 1's "large ones add feature folders identically in all4
+# trees" shape, same as kiln/src/codegen/{C,H,D,O}/expressions/) instead5
+# of ideas/gpu_acceleration.md's standalone include/runtime/backends/6
+# layout -- this is one more subsystem inside the SAME multi-project7
+# repo, so it follows that repo's rules, just keeping that doc's8
+# meaningful feature names (runtime, backends/cpu, backends/vulkan) as9
+# the folders. .cpp sources (the Vulkan backend -- RAII, see that doc's10
+# own "Why C++ for the Runtime") live in the same C/ tree as the .c11
+# sources; AGENTS.md's "one file TYPE per tree" rule is about not mixing12
+# sources with headers/objects/deps, not about .c vs .cpp specifically.13
+#14
+# Phase 1/2 status (see pyre/todo/todo.md, ideas/assigned.md): CPU15
+# backend for both benchmark kernels, Vulkan backend for matrix_multiply16
+# only (written and compiled against the real SDK via17
+# pyre/tools/install_deps.sh, never run end-to-end -- this sandbox has18
+# no GPU at all).19
+20
+CC = gcc21
+CXX = g++22
+.DEFAULT_GOAL := all23
+CFLAGS = -Wall -Wextra -std=c11 -O3 -march=native -fopenmp -Isrc -fPIC24
+CXXFLAGS = -Wall -Wextra -std=c++17 -O3 -Isrc -fPIC25
+LDFLAGS = -fopenmp -lm -ldl -lvulkan26
+27
+PYRE_C_SRC = \28
+ src/C/runtime/device.c \29
+ src/C/backends/cpu/matrix_multiply.c \30
+ src/C/backends/cpu/monte_carlo.c \31
+ src/C/backends/cpu/chaos.c32
+PYRE_CPP_SRC = \33
+ src/C/backends/vulkan/context.cpp \34
+ src/C/backends/vulkan/matrix_multiply.cpp \35
+ src/C/backends/vulkan/chaos.cpp36
+37
+# Objects and dependency files live in trees parallel to C/:38
+# src/C/<feature>/x.c -> src/O/<feature>/x.o src/D/<feature>/x.d39
+# src/C/<feature>/x.cpp -> src/O/<feature>/x.o src/D/<feature>/x.d40
+obj_of = $(subst /C/,/O/,$(basename $(1)).o)41
+dep_of = $(subst /C/,/D/,$(basename $(1)).d)42
+PYRE_OBJ = $(foreach s,$(PYRE_C_SRC) $(PYRE_CPP_SRC),$(call obj_of,$(s)))43
+PYRE_DEP = $(foreach s,$(PYRE_C_SRC) $(PYRE_CPP_SRC),$(call dep_of,$(s)))44
+45
+# The Vulkan backend's one kernel so far needs its shader compiled to46
+# SPIR-V and embedded as a C array (glslc -mfmt=c) rather than loaded47
+# from disk at run time -- no installed-data-directory story exists yet,48
+# and a compile-time-embedded array needs none. Generated, not tracked49
+# (same as any other object code); src/shaders/matrix_multiply.comp is50
+# the tracked source.51
+src/H/backends/vulkan/matrix_multiply_spv.h: src/shaders/matrix_multiply.comp52
+ @mkdir -p $(dir $@)53
+ glslc -fshader-stage=compute --target-env=vulkan1.2 -O -mfmt=c -o $@.body $<54
+ { echo 'static const uint32_t matrix_multiply_spv[] ='; cat $@.body; echo ';'; } > $@55
+ rm -f $@.body56
+57
+src/O/backends/vulkan/matrix_multiply.o: src/H/backends/vulkan/matrix_multiply_spv.h58
+59
+src/H/backends/vulkan/chaos_spv.h: src/shaders/chaos.comp60
+ @mkdir -p $(dir $@)61
+ glslc -fshader-stage=compute --target-env=vulkan1.2 -O -mfmt=c -o $@.body $<62
+ { echo 'static const uint32_t chaos_spv[] ='; cat $@.body; echo ';'; } > $@63
+ rm -f $@.body64
+65
+src/O/backends/vulkan/chaos.o: src/H/backends/vulkan/chaos_spv.h66
+67
+# -MMD -MP: a header-only edit (src/H/runtime/ash_gpu.h) correctly68
+# rebuilds every .o that includes it, not just whichever source's own69
+# timestamp changed (see kiln/Makefile's identical comment for the70
+# incident this guards against).71
+define c_compile_rule72
+$(call obj_of,$(1)): $(1)73
+ @mkdir -p $$(dir $$@) $$(dir $(call dep_of,$(1)))74
+ $$(CC) $$(CFLAGS) -MMD -MP -MF $(call dep_of,$(1)) -MT $$@ -c $$< -o $$@75
+endef76
+$(foreach s,$(PYRE_C_SRC),$(eval $(call c_compile_rule,$(s))))77
+78
+define cpp_compile_rule79
+$(call obj_of,$(1)): $(1)80
+ @mkdir -p $$(dir $$@) $$(dir $(call dep_of,$(1)))81
+ $$(CXX) $$(CXXFLAGS) -MMD -MP -MF $(call dep_of,$(1)) -MT $$@ -c $$< -o $$@82
+endef83
+$(foreach s,$(PYRE_CPP_SRC),$(eval $(call cpp_compile_rule,$(s))))84
+85
+-include $(PYRE_DEP)86
+87
+.PHONY: all clean test benchmark lib install88
+89
+all: test benchmark lib90
+91
+# Consumed by ashvm's Makefile (-Lpyre/lib -lashgpu) -- see92
+# ashvm/src/builtins/C/matrix.c and ideas/assigned.md for why ashvm links93
+# against this directly while kiln, whose own output is a hand-rolled,94
+# zero-dependency ELF/PE binary, does not.95
+lib/libashgpu.a: $(PYRE_OBJ)96
+ mkdir -p lib97
+ ar rcs $@ $(PYRE_OBJ)98
+99
+# Consumed by kiln's --link=shared mode (elf/H/elf_dynamic.h's100
+# KILN_ASHGPU_SONAME, same rpath as libkilnrt.so -- see ideas/assigned.md):101
+# a real DT_NEEDED on a shared object built by the normal system102
+# toolchain, not another hand-rolled kiln-side writer like libkilnrt.so.103
+lib/libashgpu.so: $(PYRE_OBJ)104
+ mkdir -p lib105
+ $(CXX) -shared -o $@ $(PYRE_OBJ) $(LDFLAGS)106
+107
+lib: lib/libashgpu.a lib/libashgpu.so108
+109
+# Installs libashgpu.so at kiln's own rpath (KILN_RPATH, elf/H/elf_dynamic.h)110
+# so a kiln --link=shared executable that imports an ash_gpu_* symbol can111
+# actually find it at load time -- same spot `kiln --emit-runtime` installs112
+# libkilnrt.so to by hand today (there's no kiln "make install" yet either).113
+install: lib/libashgpu.so114
+ mkdir -p /usr/local/lib/kiln115
+ cp lib/libashgpu.so /usr/local/lib/kiln/libashgpu.so116
+117
+# Each test/benchmark main is plain C -- compiled with $(CC) into its own118
+# .o (g++ would compile a .c file as C++ instead, which is stricter than119
+# this project wants these particular files to care about, e.g. no120
+# implicit void* -> T* from malloc). $(CXX) only does the final LINK, so121
+# it can pull in the Vulkan backend's C++ objects (PYRE_OBJ) without an122
+# explicit -lstdc++.123
+%_test.o: %_test.c124
+ $(CC) $(CFLAGS) -c $< -o $@125
+126
+benchmarks/%/run_benchmark.o: benchmarks/%/run_benchmark.c127
+ $(CC) $(CFLAGS) -c $< -o $@128
+129
+tests/matrix_multiply_test: tests/matrix_multiply_test.o $(PYRE_OBJ)130
+ $(CXX) $(CXXFLAGS) -fopenmp $^ -o $@ $(LDFLAGS)131
+132
+tests/monte_carlo_test: tests/monte_carlo_test.o $(PYRE_OBJ)133
+ $(CXX) $(CXXFLAGS) -fopenmp $^ -o $@ $(LDFLAGS)134
+135
+test: tests/matrix_multiply_test tests/monte_carlo_test136
+ tests/matrix_multiply_test137
+ tests/monte_carlo_test138
+139
+benchmarks/matrix/run_benchmark: benchmarks/matrix/run_benchmark.o $(PYRE_OBJ)140
+ $(CXX) $(CXXFLAGS) -fopenmp $^ -o $@ $(LDFLAGS)141
+142
+benchmarks/monte_carlo_no_pi/run_benchmark: benchmarks/monte_carlo_no_pi/run_benchmark.o $(PYRE_OBJ)143
+ $(CXX) $(CXXFLAGS) -fopenmp $^ -o $@ $(LDFLAGS)144
+145
+benchmarks/chaos/run_benchmark: benchmarks/chaos/run_benchmark.o $(PYRE_OBJ)146
+ $(CXX) $(CXXFLAGS) -fopenmp $^ -o $@ $(LDFLAGS)147
+148
+benchmark: benchmarks/matrix/run_benchmark benchmarks/monte_carlo_no_pi/run_benchmark benchmarks/chaos/run_benchmark149
+ benchmarks/matrix/run_benchmark150
+ benchmarks/monte_carlo_no_pi/run_benchmark151
+ benchmarks/chaos/run_benchmark152
+153
+clean:154
+ find src -name '*.o' -delete155
+ find src -name '*.d' -delete156
+ rm -f src/H/backends/vulkan/matrix_multiply_spv.h src/H/backends/vulkan/chaos_spv.h157
+ rm -f tests/*.o benchmarks/*/*.o158
+ rm -f tests/matrix_multiply_test tests/monte_carlo_test159
+ rm -f benchmarks/matrix/run_benchmark benchmarks/monte_carlo_no_pi/run_benchmark benchmarks/chaos/run_benchmark160
+ rm -rf libpyre/benchmarks/chaos/run_benchmark.cadded
+76−0@@ -0,0 +1,76 @@1
+/* Not one of the user's two original challenges (matrix_multiply,2
+ * monte_carlo_no_pi) -- added after matrix_multiply's real GPU3
+ * measurement came back a loss (see ash_gpu.h's comment on that4
+ * function), to answer "is there ANY workload shape where this GPU5
+ * setup (RX 9070 XT via WSL2's Dozen/D3D12 translation layer) actually6
+ * wins?" Shaped deliberately opposite to matrix_multiply: huge compute7
+ * per element, trivial data movement, no tiling to get right -- see8
+ * ash_gpu_chaos_iterate's own doc comment.9
+ *10
+ * Calls the CPU and Vulkan backends directly (bypassing11
+ * ash_gpu_chaos_iterate's own dispatcher) so both run back to back in12
+ * one process under identical conditions, instead of needing a13
+ * separate run plus an env-var trick to force one path. */14
+#include "H/runtime/ash_gpu.h"15
+#include "H/backends/cpu/chaos.h"16
+#include "H/backends/vulkan/context.h"17
+#include <stdio.h>18
+#include <stdlib.h>19
+#include <time.h>20
+21
+static double now_seconds(void) {22
+ struct timespec ts;23
+ clock_gettime(CLOCK_MONOTONIC, &ts);24
+ return (double)ts.tv_sec + (double)ts.tv_nsec / 1e9;25
+}26
+27
+int main(void) {28
+ const int n = 2000000;29
+ const int iterations = 4000;30
+31
+ double *data_cpu = malloc((size_t)n * sizeof(double));32
+ double *data_gpu = malloc((size_t)n * sizeof(double));33
+ for (int i = 0; i < n; i++) {34
+ double v = (double)(i % 1000) * 0.001;35
+ data_cpu[i] = v;36
+ data_gpu[i] = v;37
+ }38
+39
+ printf("chaos_iterate: n=%d, iterations=%d (%.1f billion ops)\n",40
+ n, iterations, (double)n * iterations / 1e9);41
+42
+ double t0 = now_seconds();43
+ ash_gpu_cpu_chaos_iterate(data_cpu, n, iterations);44
+ double cpu_seconds = now_seconds() - t0;45
+ printf("CPU backend: %.3fs\n", cpu_seconds);46
+47
+ AshGpuInfo info = ash_gpu_detect();48
+ if (info.backend == ASH_GPU_BACKEND_NONE) {49
+ printf("GPU backend: unavailable (%s)\n", info.device_name);50
+ free(data_cpu);51
+ free(data_gpu);52
+ return 0;53
+ }54
+55
+ t0 = now_seconds();56
+ ash_gpu_vulkan_chaos_iterate(data_gpu, n, iterations);57
+ double gpu_seconds = now_seconds() - t0;58
+ printf("GPU backend: %.3fs (%s)\n", gpu_seconds, info.device_name);59
+ printf("speedup: %.2fx\n", cpu_seconds / gpu_seconds);60
+61
+ /* Both backends run the identical recurrence -- results must match62
+ * (within fp rounding-order slack, since the GPU evaluates per63
+ * thread independently same as the CPU's OpenMP loop does, both64
+ * single-precision-free, double throughout). */65
+ double max_diff = 0.0;66
+ for (int i = 0; i < n; i++) {67
+ double diff = data_cpu[i] - data_gpu[i];68
+ if (diff < 0) diff = -diff;69
+ if (diff > max_diff) max_diff = diff;70
+ }71
+ printf("max |CPU - GPU| difference: %g\n", max_diff);72
+73
+ free(data_cpu);74
+ free(data_gpu);75
+ return 0;76
+}pyre/benchmarks/matrix/matrix_multiply.ashadded
+63−0@@ -0,0 +1,63 @@1
+// Goal to beat2
+// 3
+// use the "time" command to measure the execution time of the program4
+//5
+// matrix_multiply.cpp time (compiled with g++ at -O2) output:6
+//real 0m5.761s7
+//user 0m5.721s8
+//sys 0m0.037s9
+//10
+// CPU: AMD Ryzen 7 7800X 8-Core Processor11
+// GPU: AMD Radeon RX 9070 XT 16GB12
+// 13
+// if you cant beat it by atleast 50% then you are not allowed to submit your solution14
+15
+// Matrix multiplication: C = A x B16
+// Dimensions: N x N17
+local N = 2000;18
+19
+forge create_matrix(size, value) {20
+ local mat = [];21
+ local i = 0;22
+ during i < size {23
+ local row = [];24
+ local j = 0;25
+ during j < size {26
+ push(row, value);27
+ j += 1;28
+ }29
+ push(mat, row);30
+ i += 1;31
+ }32
+ yield mat;33
+}34
+35
+// 1. Initialize two N x N matrices36
+local A = create_matrix(N, 1.5);37
+local B = create_matrix(N, 2.0);38
+39
+// 2. Initialize the result matrix C with zeros40
+local C = create_matrix(N, 0.0);41
+42
+say "Starting matrix multiplication (" + str(N) + "x" + str(N) + ")...";43
+44
+// 3. Compute C = A x B45
+// Total operations: 2 * N^3 = 250,000,000 floating point operations46
+local i = 0;47
+during i < N {48
+ local j = 0;49
+ during j < N {50
+ local sum = 0.0;51
+ local k = 0;52
+ during k < N {53
+ sum += A[i][k] * B[k][j];54
+ k += 1;55
+ }56
+ C[i][j] = sum;57
+ j += 1;58
+ }59
+ i += 1;60
+}61
+62
+say "Done!";63
+say "Sample element C[0][0]: " + str(C[0][0]); // Should be 1500\ No newline at end of filepyre/benchmarks/matrix/matrix_multiply.cppadded
+53−0@@ -0,0 +1,53 @@1
+#define N 20002
+3
+double** create_matrix(int size, double value)4
+{5
+ double** mat = new double*[size];6
+7
+ for (int i = 0; i < size; i++) {8
+ mat[i] = new double[size];9
+10
+ for (int j = 0; j < size; j++) {11
+ mat[i][j] = value;12
+ }13
+ }14
+15
+ return mat;16
+}17
+18
+void free_matrix(double** mat, int size)19
+{20
+ for (int i = 0; i < size; i++) {21
+ delete[] mat[i];22
+ }23
+24
+ delete[] mat;25
+}26
+27
+int main()28
+{29
+ double** A = create_matrix(N, 1.5);30
+ double** B = create_matrix(N, 2.0);31
+ double** C = create_matrix(N, 0.0);32
+33
+ for (int i = 0; i < N; i++) {34
+ for (int j = 0; j < N; j++) {35
+ double sum = 0.0;36
+37
+ for (int k = 0; k < N; k++) {38
+ sum += A[i][k] * B[k][j];39
+ }40
+41
+ C[i][j] = sum;42
+ }43
+ }44
+45
+ // Prevent the compiler from treating the calculation as unused.46
+ volatile double result = C[0][0];47
+48
+ free_matrix(A, N);49
+ free_matrix(B, N);50
+ free_matrix(C, N);51
+52
+ return result == 1500.0 ? 0 : 1;53
+}\ No newline at end of filepyre/benchmarks/matrix/run_benchmark.cadded
+36−0@@ -0,0 +1,36 @@1
+/* Pyre's side of the matrix_multiply challenge: same N, same values,2
+ * same "time this with the `time` command" methodology as3
+ * matrix_multiply.cpp, but through ash_gpu_matrix_multiply (CPU backend:4
+ * ikj loop order + OpenMP, see src/C/backends/cpu/matrix_multiply.c) instead5
+ * of the naive ijk double** reference. Goal (matrix_multiply.ash's own6
+ * comment): beat matrix_multiply.cpp's 5.761s by at least 50%, i.e.7
+ * finish in under ~2.88s on this machine.8
+ */9
+#include "H/runtime/ash_gpu.h"10
+#include <stdio.h>11
+#include <stdlib.h>12
+13
+int main(void) {14
+ const int n = 2000;15
+ double *a = malloc((size_t)n * n * sizeof(double));16
+ double *b = malloc((size_t)n * n * sizeof(double));17
+ double *c = malloc((size_t)n * n * sizeof(double));18
+19
+ for (int i = 0; i < n * n; i++) {20
+ a[i] = 1.5;21
+ b[i] = 2.0;22
+ }23
+24
+ printf("Starting matrix multiplication (%dx%d)...\n", n, n);25
+ ash_gpu_matrix_multiply(a, b, c, n);26
+ printf("Done!\n");27
+ printf("Sample element C[0][0]: %g\n", c[0]);28
+29
+ double expected = (double)n * 1.5 * 2.0; /* every term of the dot product is 1.5*2.0, N of them */30
+ int ok = c[0] == expected;31
+32
+ free(a);33
+ free(b);34
+ free(c);35
+ return ok ? 0 : 1;36
+}pyre/benchmarks/monte_carlo_no_pi/mcnp.ashadded
+59−0@@ -0,0 +1,59 @@1
+// Monte Carlo simulation ( no pi )2
+3
+// Goal to beat4
+// 5
+// use the "time" command to measure the execution time of the program6
+//7
+// mcnp.cpp time (compiled with g++ at -O3) output:8
+//real 0m1.908s9
+//user 0m1.906s10
+//sys 0m0.000s11
+//12
+// CPU: AMD Ryzen 7 7800X 8-Core Processor13
+// GPU: AMD Radeon RX 9070 XT 16GB14
+//15
+// if you cant beat it by atleast 50% then you are not allowed to submit your solution16
+17
+local ITERATIONS = 1000000000;18
+local threshold = 0.75;19
+20
+forge random(seed) {21
+ seed = (seed * 1664525 + 1013904223) % 4294967296;22
+ yield seed;23
+}24
+25
+local seed = 123456789;26
+local risky = 0;27
+28
+say "Starting Monte Carlo simulation...";29
+say "Iterations: " + str(ITERATIONS);30
+31
+local i = 0;32
+33
+during i < ITERATIONS {34
+ local random_value = random(seed);35
+ seed = random_value;36
+37
+ local normalized = random_value / 4294967296.0;38
+39
+ // Simulate a risk event.40
+ // Combine two independent random values.41
+ local random_value_2 = random(seed);42
+ seed = random_value_2;43
+44
+ local normalized_2 = random_value_2 / 4294967296.0;45
+46
+ local score = normalized * normalized_2;47
+48
+ given score > threshold {49
+ risky += 1;50
+ }51
+52
+ i += 1;53
+}54
+55
+local probability = risky / ITERATIONS;56
+57
+say "Done!";58
+say "Risky outcomes: " + str(risky);59
+say "Estimated probability: " + str(probability);\ No newline at end of filepyre/benchmarks/monte_carlo_no_pi/mcnp.cppadded
+46−0@@ -0,0 +1,46 @@1
+#include <iostream>2
+#include <cstdint>3
+4
+#define ITERATIONS 1000000000ULL5
+6
+uint32_t random(uint32_t& seed)7
+{8
+ seed = seed * 1664525u + 1013904223u;9
+ return seed;10
+}11
+12
+int main()13
+{14
+ const double threshold = 0.75;15
+16
+ uint32_t seed = 123456789;17
+ uint64_t risky = 0;18
+19
+ std::cout << "Starting Monte Carlo simulation...\n";20
+ std::cout << "Iterations: " << ITERATIONS << "\n";21
+22
+ for (uint64_t i = 0; i < ITERATIONS; i++) {23
+ uint32_t random_value = random(seed);24
+ double normalized =25
+ static_cast<double>(random_value) / 4294967296.0;26
+27
+ uint32_t random_value_2 = random(seed);28
+ double normalized_2 =29
+ static_cast<double>(random_value_2) / 4294967296.0;30
+31
+ double score = normalized * normalized_2;32
+33
+ if (score > threshold) {34
+ risky++;35
+ }36
+ }37
+38
+ double probability =39
+ static_cast<double>(risky) / static_cast<double>(ITERATIONS);40
+41
+ std::cout << "Done!\n";42
+ std::cout << "Risky outcomes: " << risky << "\n";43
+ std::cout << "Estimated probability: " << probability << "\n";44
+45
+ return 0;46
+}\ No newline at end of filepyre/benchmarks/monte_carlo_no_pi/run_benchmark.cadded
+29−0@@ -0,0 +1,29 @@1
+/* Pyre's side of the monte_carlo_no_pi challenge: same iteration count,2
+ * threshold and recurrence as mcnp.cpp, through ash_gpu_monte_carlo_risk3
+ * (CPU backend: OpenMP across independently-seeded per-thread streams,4
+ * see src/C/backends/cpu/monte_carlo.c). The risky count will differ from a5
+ * single-threaded run with the same seed (see ash_gpu.h's comment on6
+ * this kernel) -- the benchmark's target is wall-clock time, measured7
+ * the same way mcnp.ash's own comment asks for (the `time` command).8
+ * Goal: beat mcnp.cpp's 1.908s by at least 50%, i.e. finish in under9
+ * ~0.95s on this machine.10
+ */11
+#include "H/runtime/ash_gpu.h"12
+#include <stdio.h>13
+14
+int main(void) {15
+ const long long iterations = 1000000000LL;16
+ const double threshold = 0.75;17
+ long long risky = 0;18
+ double probability = 0.0;19
+20
+ printf("Starting Monte Carlo simulation...\n");21
+ printf("Iterations: %lld\n", iterations);22
+23
+ ash_gpu_monte_carlo_risk(iterations, threshold, 123456789u, &risky, &probability);24
+25
+ printf("Done!\n");26
+ printf("Risky outcomes: %lld\n", risky);27
+ printf("Estimated probability: %g\n", probability);28
+ return 0;29
+}pyre/src/C/backends/cpu/chaos.cadded
+22−0@@ -0,0 +1,22 @@1
+#include "H/backends/cpu/chaos.h"2
+#include <math.h>3
+4
+// CPU backend for chaos_iterate -- always available. See ash_gpu.h for5
+// why this kernel exists (matrix_multiply's real GPU measurement came6
+// back a loss; this one is shaped to actually favor the GPU).7
+8
+// Same recurrence as src/shaders/chaos.comp, exactly -- see that file's9
+// comment for why it's sqrt/abs-based rather than sin/cos (GLSL's10
+// fp64 extension doesn't give sin/cos a double overload at all, even11
+// though it covers sqrt/abs; keeping both backends on the identical12
+// formula makes the CPU-vs-GPU comparison apples-to-apples).13
+void ash_gpu_cpu_chaos_iterate(double *data, int n, int iterations) {14
+#pragma omp parallel for schedule(static)15
+ for (int i = 0; i < n; i++) {16
+ double x = data[i];17
+ for (int k = 0; k < iterations; k++) {18
+ x = sqrt(fabs(x) + 1.0) * 0.99999 + 0.00001;19
+ }20
+ data[i] = x;21
+ }22
+}pyre/src/C/backends/cpu/matrix_multiply.cadded
+42−0@@ -0,0 +1,42 @@1
+#include "H/backends/cpu/matrix_multiply.h"2
+#include <string.h>3
+4
+/*5
+ * CPU backend for matrix_multiply -- always available (no hardware6
+ * needed), so the dispatcher in src/C/runtime/device.c falls back to7
+ * this whenever ash_gpu_vulkan_init() hasn't found a usable device8
+ * (which, on this sandbox, is always -- see context.cpp).9
+ *10
+ * Two changes from the naive ijk reference (pyre/benchmarks/matrix/matrix_multiply.cpp):11
+ *12
+ * 1. Loop order ikj instead of ijk. The naive order's innermost loop13
+ * reads B[k][j] for fixed i,j while k varies -- a stride of N14
+ * doubles through a 32MB matrix, missing cache on nearly every15
+ * access. Swapping to ikj makes the innermost loop walk j across a16
+ * single row of B and a single row of C, both sequential, which also17
+ * lets the compiler auto-vectorize it (contiguous FMA over doubles).18
+ * 2. OpenMP across the outer i loop -- each output row is independent,19
+ * so this is an embarrassingly parallel split across cores with no20
+ * locking needed.21
+ *22
+ * No GPU required for either change; this is squarely the "best CPU23
+ * implementation" design principle 1 in ideas/gpu_acceleration.md asks24
+ * for when there is no usable device.25
+ */26
+void ash_gpu_cpu_matrix_multiply(const double *a, const double *b, double *c, int n) {27
+ memset(c, 0, (size_t)n * (size_t)n * sizeof(double));28
+29
+#pragma omp parallel for schedule(static)30
+ for (int i = 0; i < n; i++) {31
+ const double *a_row = a + (size_t)i * n;32
+ double *c_row = c + (size_t)i * n;33
+ for (int k = 0; k < n; k++) {34
+ double a_ik = a_row[k];35
+ const double *b_row = b + (size_t)k * n;36
+#pragma omp simd37
+ for (int j = 0; j < n; j++) {38
+ c_row[j] += a_ik * b_row[j];39
+ }40
+ }41
+ }42
+}pyre/src/C/backends/cpu/monte_carlo.cadded
+59−0@@ -0,0 +1,59 @@1
+#include "H/runtime/ash_gpu.h"2
+#ifdef _OPENMP3
+#include <omp.h>4
+#endif5
+6
+/*7
+ * CPU fallback for ash_gpu_monte_carlo_risk. Each iteration of the8
+ * reference (pyre/benchmarks/monte_carlo_no_pi/mcnp.cpp) is two draws of9
+ * the same LCG (seed = seed*1664525 + 1013904223) feeding one10
+ * independent comparison -- no data dependency between iterations11
+ * except the seed chain itself, so splitting the iteration range across12
+ * threads just needs each thread to start from its own, well-separated13
+ * seed instead of continuing one shared chain.14
+ *15
+ * splitmix32 (a cheap, well-mixing finalizer, not the LCG above) turns16
+ * the caller's single seed plus a thread index into that many17
+ * decorrelated per-thread starting seeds.18
+ */19
+static unsigned int splitmix32(unsigned int x) {20
+ x += 0x9E3779B9u;21
+ x = (x ^ (x >> 16)) * 0x85EBCA6Bu;22
+ x = (x ^ (x >> 13)) * 0xC2B2AE35u;23
+ return x ^ (x >> 16);24
+}25
+26
+void ash_gpu_monte_carlo_risk(long long iterations, double threshold, unsigned int seed,27
+ long long *out_risky, double *out_probability) {28
+ long long risky = 0;29
+30
+#pragma omp parallel31
+ {32
+ int tid =33
+#ifdef _OPENMP34
+ omp_get_thread_num();35
+#else36
+ 0;37
+#endif38
+ unsigned int local_seed = splitmix32(seed + (unsigned int)tid * 2654435761u);39
+ long long local_risky = 0;40
+41
+#pragma omp for schedule(static)42
+ for (long long i = 0; i < iterations; i++) {43
+ local_seed = local_seed * 1664525u + 1013904223u;44
+ double normalized = (double)local_seed / 4294967296.0;45
+46
+ local_seed = local_seed * 1664525u + 1013904223u;47
+ double normalized2 = (double)local_seed / 4294967296.0;48
+49
+ double score = normalized * normalized2;50
+ if (score > threshold) local_risky++;51
+ }52
+53
+#pragma omp atomic54
+ risky += local_risky;55
+ }56
+57
+ *out_risky = risky;58
+ *out_probability = (double)risky / (double)iterations;59
+}pyre/src/C/backends/vulkan/chaos.cppadded
+173−0@@ -0,0 +1,173 @@1
+// Vulkan compute dispatch for chaos_iterate -- see2
+// src/H/runtime/ash_gpu.h's comment on why this kernel exists3
+// (matrix_multiply's real GPU measurement came back a loss; this one is4
+// shaped to actually favor the GPU) and src/shaders/chaos.comp for the5
+// recurrence itself. One buffer, in place -- simpler than6
+// matrix_multiply.cpp's three (no separate input/output, no second7
+// binding).8
+#include "H/backends/vulkan/context.h"9
+#include "H/backends/vulkan/context_internal.hpp"10
+#include <vulkan/vulkan.h>11
+#include <cstring>12
+#include <cstdint>13
+14
+#include "H/backends/vulkan/chaos_spv.h"15
+16
+namespace {17
+18
+uint32_t find_memory_type(VkPhysicalDevice phys, uint32_t type_filter, VkMemoryPropertyFlags props) {19
+ VkPhysicalDeviceMemoryProperties mem_props;20
+ vkGetPhysicalDeviceMemoryProperties(phys, &mem_props);21
+ for (uint32_t i = 0; i < mem_props.memoryTypeCount; i++) {22
+ if ((type_filter & (1u << i)) && (mem_props.memoryTypes[i].propertyFlags & props) == props) return i;23
+ }24
+ return UINT32_MAX;25
+}26
+27
+} // namespace28
+29
+extern "C" void ash_gpu_vulkan_chaos_iterate(double *data, int n, int iterations) {30
+ const AshGpuVulkanContext &ctx = ash_gpu_vulkan_context();31
+ VkDevice device = ctx.device;32
+ VkDeviceSize bytes = (VkDeviceSize)n * sizeof(double);33
+34
+ VkBufferCreateInfo buf_info{};35
+ buf_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;36
+ buf_info.size = bytes;37
+ buf_info.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT;38
+ buf_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;39
+ VkBuffer buffer = VK_NULL_HANDLE;40
+ if (vkCreateBuffer(device, &buf_info, nullptr, &buffer) != VK_SUCCESS) return;41
+42
+ VkMemoryRequirements mem_req;43
+ vkGetBufferMemoryRequirements(device, buffer, &mem_req);44
+ uint32_t mem_type = find_memory_type(ctx.physical_device, mem_req.memoryTypeBits,45
+ VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT);46
+ if (mem_type == UINT32_MAX) { vkDestroyBuffer(device, buffer, nullptr); return; }47
+48
+ VkMemoryAllocateInfo alloc_info{};49
+ alloc_info.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;50
+ alloc_info.allocationSize = mem_req.size;51
+ alloc_info.memoryTypeIndex = mem_type;52
+ VkDeviceMemory memory = VK_NULL_HANDLE;53
+ if (vkAllocateMemory(device, &alloc_info, nullptr, &memory) != VK_SUCCESS) {54
+ vkDestroyBuffer(device, buffer, nullptr);55
+ return;56
+ }57
+ vkBindBufferMemory(device, buffer, memory, 0);58
+59
+ void *mapped = nullptr;60
+ vkMapMemory(device, memory, 0, bytes, 0, &mapped);61
+ memcpy(mapped, data, (size_t)bytes);62
+63
+ VkDescriptorSetLayoutBinding binding{};64
+ binding.binding = 0;65
+ binding.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;66
+ binding.descriptorCount = 1;67
+ binding.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;68
+ VkDescriptorSetLayoutCreateInfo layout_info{};69
+ layout_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;70
+ layout_info.bindingCount = 1;71
+ layout_info.pBindings = &binding;72
+ VkDescriptorSetLayout set_layout = VK_NULL_HANDLE;73
+ vkCreateDescriptorSetLayout(device, &layout_info, nullptr, &set_layout);74
+75
+ struct PushConstants { uint32_t n; uint32_t iterations; } push{(uint32_t)n, (uint32_t)iterations};76
+ VkPushConstantRange push_range{};77
+ push_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;78
+ push_range.offset = 0;79
+ push_range.size = sizeof(PushConstants);80
+81
+ VkPipelineLayoutCreateInfo pipeline_layout_info{};82
+ pipeline_layout_info.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;83
+ pipeline_layout_info.setLayoutCount = 1;84
+ pipeline_layout_info.pSetLayouts = &set_layout;85
+ pipeline_layout_info.pushConstantRangeCount = 1;86
+ pipeline_layout_info.pPushConstantRanges = &push_range;87
+ VkPipelineLayout pipeline_layout = VK_NULL_HANDLE;88
+ vkCreatePipelineLayout(device, &pipeline_layout_info, nullptr, &pipeline_layout);89
+90
+ VkShaderModuleCreateInfo shader_info{};91
+ shader_info.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;92
+ shader_info.codeSize = sizeof(chaos_spv);93
+ shader_info.pCode = chaos_spv;94
+ VkShaderModule shader_module = VK_NULL_HANDLE;95
+ vkCreateShaderModule(device, &shader_info, nullptr, &shader_module);96
+97
+ VkComputePipelineCreateInfo pipeline_info{};98
+ pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;99
+ pipeline_info.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;100
+ pipeline_info.stage.stage = VK_SHADER_STAGE_COMPUTE_BIT;101
+ pipeline_info.stage.module = shader_module;102
+ pipeline_info.stage.pName = "main";103
+ pipeline_info.layout = pipeline_layout;104
+ VkPipeline pipeline = VK_NULL_HANDLE;105
+ vkCreateComputePipelines(device, VK_NULL_HANDLE, 1, &pipeline_info, nullptr, &pipeline);106
+107
+ VkDescriptorPoolSize pool_size{};108
+ pool_size.type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;109
+ pool_size.descriptorCount = 1;110
+ VkDescriptorPoolCreateInfo pool_info{};111
+ pool_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO;112
+ pool_info.maxSets = 1;113
+ pool_info.poolSizeCount = 1;114
+ pool_info.pPoolSizes = &pool_size;115
+ VkDescriptorPool desc_pool = VK_NULL_HANDLE;116
+ vkCreateDescriptorPool(device, &pool_info, nullptr, &desc_pool);117
+118
+ VkDescriptorSetAllocateInfo set_alloc_info{};119
+ set_alloc_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO;120
+ set_alloc_info.descriptorPool = desc_pool;121
+ set_alloc_info.descriptorSetCount = 1;122
+ set_alloc_info.pSetLayouts = &set_layout;123
+ VkDescriptorSet desc_set = VK_NULL_HANDLE;124
+ vkAllocateDescriptorSets(device, &set_alloc_info, &desc_set);125
+126
+ VkDescriptorBufferInfo buf_info_desc{buffer, 0, bytes};127
+ VkWriteDescriptorSet write{};128
+ write.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET;129
+ write.dstSet = desc_set;130
+ write.dstBinding = 0;131
+ write.descriptorCount = 1;132
+ write.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;133
+ write.pBufferInfo = &buf_info_desc;134
+ vkUpdateDescriptorSets(device, 1, &write, 0, nullptr);135
+136
+ VkCommandBufferAllocateInfo cmd_alloc_info{};137
+ cmd_alloc_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;138
+ cmd_alloc_info.commandPool = ctx.command_pool;139
+ cmd_alloc_info.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;140
+ cmd_alloc_info.commandBufferCount = 1;141
+ VkCommandBuffer cmd = VK_NULL_HANDLE;142
+ vkAllocateCommandBuffers(device, &cmd_alloc_info, &cmd);143
+144
+ VkCommandBufferBeginInfo begin_info{};145
+ begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;146
+ begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;147
+ vkBeginCommandBuffer(cmd, &begin_info);148
+ vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);149
+ vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline_layout, 0, 1, &desc_set, 0, nullptr);150
+ vkCmdPushConstants(cmd, pipeline_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(push), &push);151
+ uint32_t groups = ((uint32_t)n + 255u) / 256u; // matches local_size_x = 256 in the shader152
+ vkCmdDispatch(cmd, groups, 1, 1);153
+ vkEndCommandBuffer(cmd);154
+155
+ VkSubmitInfo submit_info{};156
+ submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;157
+ submit_info.commandBufferCount = 1;158
+ submit_info.pCommandBuffers = &cmd;159
+ vkQueueSubmit(ctx.compute_queue, 1, &submit_info, VK_NULL_HANDLE);160
+ vkQueueWaitIdle(ctx.compute_queue);161
+162
+ memcpy(data, mapped, (size_t)bytes);163
+164
+ vkFreeCommandBuffers(device, ctx.command_pool, 1, &cmd);165
+ vkDestroyDescriptorPool(device, desc_pool, nullptr);166
+ vkDestroyPipeline(device, pipeline, nullptr);167
+ vkDestroyShaderModule(device, shader_module, nullptr);168
+ vkDestroyPipelineLayout(device, pipeline_layout, nullptr);169
+ vkDestroyDescriptorSetLayout(device, set_layout, nullptr);170
+ vkUnmapMemory(device, memory);171
+ vkDestroyBuffer(device, buffer, nullptr);172
+ vkFreeMemory(device, memory, nullptr);173
+}pyre/src/C/backends/vulkan/context.cppadded
+127−0@@ -0,0 +1,127 @@1
+// Vulkan instance/device setup -- see H/backends/vulkan/context.h for2
+// the contract. Written and compiled against the real Vulkan 1.2 SDK3
+// (pyre/tools/install_deps.sh), but never run end-to-end: this sandbox4
+// has no GPU and no ICD at all (vulkaninfo reports "Found no drivers!"),5
+// so vkCreateInstance itself has never actually succeeded here. Treat6
+// this file as reviewed-and-compiles, not verified-at-runtime, until7
+// someone runs it on real hardware (the user's RX 9070 XT).8
+#include "H/backends/vulkan/context.h"9
+#include "H/backends/vulkan/context_internal.hpp"10
+#include <vulkan/vulkan.h>11
+#include <vector>12
+#include <cstring>13
+14
+namespace {15
+16
+enum class InitState { NotTried, Ready, Unavailable };17
+InitState g_state = InitState::NotTried;18
+AshGpuVulkanContext g_ctx;19
+char g_device_name[256] = "";20
+21
+void destroy_partial(AshGpuVulkanContext &ctx) {22
+ if (ctx.command_pool != VK_NULL_HANDLE) vkDestroyCommandPool(ctx.device, ctx.command_pool, nullptr);23
+ if (ctx.device != VK_NULL_HANDLE) vkDestroyDevice(ctx.device, nullptr);24
+ if (ctx.instance != VK_NULL_HANDLE) vkDestroyInstance(ctx.instance, nullptr);25
+ ctx = AshGpuVulkanContext{};26
+}27
+28
+// First physical device with both a compute-capable queue family and29
+// the shaderFloat64 feature (ash_gpu's ABI is `double` throughout, see30
+// ash_gpu.h) -- a device that can only do fp32 compute isn't a match.31
+bool pick_physical_device(VkInstance instance, VkPhysicalDevice *out_device, uint32_t *out_queue_family) {32
+ uint32_t count = 0;33
+ vkEnumeratePhysicalDevices(instance, &count, nullptr);34
+ if (count == 0) return false;35
+ std::vector<VkPhysicalDevice> devices(count);36
+ vkEnumeratePhysicalDevices(instance, &count, devices.data());37
+38
+ for (VkPhysicalDevice dev : devices) {39
+ VkPhysicalDeviceFeatures features;40
+ vkGetPhysicalDeviceFeatures(dev, &features);41
+ if (!features.shaderFloat64) continue;42
+43
+ uint32_t qf_count = 0;44
+ vkGetPhysicalDeviceQueueFamilyProperties(dev, &qf_count, nullptr);45
+ std::vector<VkQueueFamilyProperties> qfs(qf_count);46
+ vkGetPhysicalDeviceQueueFamilyProperties(dev, &qf_count, qfs.data());47
+48
+ for (uint32_t i = 0; i < qf_count; i++) {49
+ if (qfs[i].queueFlags & VK_QUEUE_COMPUTE_BIT) {50
+ *out_device = dev;51
+ *out_queue_family = i;52
+ return true;53
+ }54
+ }55
+ }56
+ return false;57
+}58
+59
+bool init_once(AshGpuVulkanContext &ctx) {60
+ VkApplicationInfo app_info{};61
+ app_info.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;62
+ app_info.pApplicationName = "pyre";63
+ app_info.apiVersion = VK_API_VERSION_1_2;64
+65
+ VkInstanceCreateInfo instance_info{};66
+ instance_info.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;67
+ instance_info.pApplicationInfo = &app_info;68
+ // Headless compute only -- no VK_KHR_surface/swapchain, so no69
+ // extensions or layers are needed at all for this kernel.70
+71
+ if (vkCreateInstance(&instance_info, nullptr, &ctx.instance) != VK_SUCCESS) return false;72
+73
+ if (!pick_physical_device(ctx.instance, &ctx.physical_device, &ctx.compute_queue_family)) return false;74
+75
+ float queue_priority = 1.0f;76
+ VkDeviceQueueCreateInfo queue_info{};77
+ queue_info.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;78
+ queue_info.queueFamilyIndex = ctx.compute_queue_family;79
+ queue_info.queueCount = 1;80
+ queue_info.pQueuePriorities = &queue_priority;81
+82
+ VkPhysicalDeviceFeatures enabled_features{};83
+ enabled_features.shaderFloat64 = VK_TRUE;84
+85
+ VkDeviceCreateInfo device_info{};86
+ device_info.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;87
+ device_info.queueCreateInfoCount = 1;88
+ device_info.pQueueCreateInfos = &queue_info;89
+ device_info.pEnabledFeatures = &enabled_features;90
+91
+ if (vkCreateDevice(ctx.physical_device, &device_info, nullptr, &ctx.device) != VK_SUCCESS) return false;92
+93
+ vkGetDeviceQueue(ctx.device, ctx.compute_queue_family, 0, &ctx.compute_queue);94
+95
+ VkCommandPoolCreateInfo pool_info{};96
+ pool_info.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;97
+ pool_info.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;98
+ pool_info.queueFamilyIndex = ctx.compute_queue_family;99
+ if (vkCreateCommandPool(ctx.device, &pool_info, nullptr, &ctx.command_pool) != VK_SUCCESS) return false;100
+101
+ VkPhysicalDeviceProperties props;102
+ vkGetPhysicalDeviceProperties(ctx.physical_device, &props);103
+ std::strncpy(g_device_name, props.deviceName, sizeof(g_device_name) - 1);104
+ g_device_name[sizeof(g_device_name) - 1] = '\0';105
+106
+ return true;107
+}108
+109
+} // namespace110
+111
+extern "C" int ash_gpu_vulkan_init(void) {112
+ if (g_state == InitState::NotTried) {113
+ AshGpuVulkanContext ctx;114
+ if (init_once(ctx)) {115
+ g_ctx = ctx;116
+ g_state = InitState::Ready;117
+ } else {118
+ destroy_partial(ctx);119
+ g_state = InitState::Unavailable;120
+ }121
+ }122
+ return g_state == InitState::Ready ? 1 : 0;123
+}124
+125
+extern "C" const char *ash_gpu_vulkan_device_name(void) { return g_device_name; }126
+127
+const AshGpuVulkanContext &ash_gpu_vulkan_context(void) { return g_ctx; }pyre/src/C/backends/vulkan/matrix_multiply.cppadded
+224−0@@ -0,0 +1,224 @@1
+// Vulkan compute dispatch for the one kernel this backend has so far.2
+// Same "compiles against the real SDK, never run end-to-end" caveat as3
+// context.cpp -- this sandbox has no device to dispatch work to at all.4
+// Most VkResult checks below are omitted (matching this project's5
+// existing style of not writing defensive code for failures that need6
+// a debugging session, not a guard, to actually diagnose) -- when this7
+// runs for the first time on real hardware and something's wrong, run8
+// it with VK_LAYER_KHRONOS_validation enabled (`vulkan-validation-layers`,9
+// not installed by install_deps.sh since it's a debug tool, not a build10
+// dependency); the validation layer will name the exact call and the11
+// exact violated rule far more precisely than a hand-checked VkResult12
+// ever would.13
+#include "H/backends/vulkan/context.h"14
+#include "H/backends/vulkan/context_internal.hpp"15
+#include <vulkan/vulkan.h>16
+#include <cstring>17
+#include <cstdint>18
+19
+// Generated by `make` from src/shaders/matrix_multiply.comp via glslc20
+// -mfmt=c (see Makefile) -- not checked in, same as any other object21
+// code: `static const uint32_t matrix_multiply_spv[] = {...};`.22
+#include "H/backends/vulkan/matrix_multiply_spv.h"23
+24
+namespace {25
+26
+uint32_t find_memory_type(VkPhysicalDevice phys, uint32_t type_filter, VkMemoryPropertyFlags props) {27
+ VkPhysicalDeviceMemoryProperties mem_props;28
+ vkGetPhysicalDeviceMemoryProperties(phys, &mem_props);29
+ for (uint32_t i = 0; i < mem_props.memoryTypeCount; i++) {30
+ if ((type_filter & (1u << i)) && (mem_props.memoryTypes[i].propertyFlags & props) == props) return i;31
+ }32
+ return UINT32_MAX;33
+}34
+35
+struct Buffer {36
+ VkBuffer buffer = VK_NULL_HANDLE;37
+ VkDeviceMemory memory = VK_NULL_HANDLE;38
+ void *mapped = nullptr;39
+};40
+41
+// Host-visible + host-coherent memory for every buffer, no staging42
+// buffer / device-local copy: simplest correct option, not the fastest43
+// one. A device-local buffer plus an explicit transfer would avoid44
+// paying PCIe-mapped-memory access cost on every shader read/write, but45
+// tuning that tradeoff needs a real device to measure against -- there46
+// isn't one here (see context.cpp).47
+bool create_host_visible_buffer(VkDevice device, VkPhysicalDevice phys, VkDeviceSize size, Buffer *out) {48
+ VkBufferCreateInfo buf_info{};49
+ buf_info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;50
+ buf_info.size = size;51
+ buf_info.usage = VK_BUFFER_USAGE_STORAGE_BUFFER_BIT;52
+ buf_info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;53
+ if (vkCreateBuffer(device, &buf_info, nullptr, &out->buffer) != VK_SUCCESS) return false;54
+55
+ VkMemoryRequirements mem_req;56
+ vkGetBufferMemoryRequirements(device, out->buffer, &mem_req);57
+ uint32_t mem_type = find_memory_type(phys, mem_req.memoryTypeBits,58
+ VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT);59
+ if (mem_type == UINT32_MAX) return false;60
+61
+ VkMemoryAllocateInfo alloc_info{};62
+ alloc_info.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;63
+ alloc_info.allocationSize = mem_req.size;64
+ alloc_info.memoryTypeIndex = mem_type;65
+ if (vkAllocateMemory(device, &alloc_info, nullptr, &out->memory) != VK_SUCCESS) return false;66
+ if (vkBindBufferMemory(device, out->buffer, out->memory, 0) != VK_SUCCESS) return false;67
+ if (vkMapMemory(device, out->memory, 0, size, 0, &out->mapped) != VK_SUCCESS) return false;68
+ return true;69
+}70
+71
+void destroy_buffer(VkDevice device, Buffer &buf) {72
+ if (buf.mapped) vkUnmapMemory(device, buf.memory);73
+ if (buf.buffer != VK_NULL_HANDLE) vkDestroyBuffer(device, buf.buffer, nullptr);74
+ if (buf.memory != VK_NULL_HANDLE) vkFreeMemory(device, buf.memory, nullptr);75
+ buf = Buffer{};76
+}77
+78
+} // namespace79
+80
+extern "C" void ash_gpu_vulkan_matrix_multiply(const double *a, const double *b, double *c, int n) {81
+ const AshGpuVulkanContext &ctx = ash_gpu_vulkan_context();82
+ VkDevice device = ctx.device;83
+ VkDeviceSize bytes = (VkDeviceSize)n * (VkDeviceSize)n * sizeof(double);84
+85
+ Buffer buf_a{}, buf_b{}, buf_c{};86
+ bool ok = create_host_visible_buffer(device, ctx.physical_device, bytes, &buf_a) &&87
+ create_host_visible_buffer(device, ctx.physical_device, bytes, &buf_b) &&88
+ create_host_visible_buffer(device, ctx.physical_device, bytes, &buf_c);89
+ if (!ok) {90
+ // Caller (runtime/device.c's dispatcher) only reaches this path91
+ // after ash_gpu_vulkan_init() already succeeded, so this would92
+ // mean running out of device memory mid-call -- nothing graceful93
+ // to fall back to from here (the CPU path needs the caller's94
+ // buffers, which are untouched, so the caller COULD retry on95
+ // CPU, but this function's contract is "fill C", not "report96
+ // failure"; see todo.md for making this a real error path).97
+ destroy_buffer(device, buf_a);98
+ destroy_buffer(device, buf_b);99
+ destroy_buffer(device, buf_c);100
+ return;101
+ }102
+103
+ memcpy(buf_a.mapped, a, (size_t)bytes);104
+ memcpy(buf_b.mapped, b, (size_t)bytes);105
+106
+ VkDescriptorSetLayoutBinding bindings[3] = {};107
+ for (int i = 0; i < 3; i++) {108
+ bindings[i].binding = (uint32_t)i;109
+ bindings[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;110
+ bindings[i].descriptorCount = 1;111
+ bindings[i].stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;112
+ }113
+ VkDescriptorSetLayoutCreateInfo layout_info{};114
+ layout_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;115
+ layout_info.bindingCount = 3;116
+ layout_info.pBindings = bindings;117
+ VkDescriptorSetLayout set_layout = VK_NULL_HANDLE;118
+ vkCreateDescriptorSetLayout(device, &layout_info, nullptr, &set_layout);119
+120
+ VkPushConstantRange push_range{};121
+ push_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;122
+ push_range.offset = 0;123
+ push_range.size = sizeof(uint32_t);124
+125
+ VkPipelineLayoutCreateInfo pipeline_layout_info{};126
+ pipeline_layout_info.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;127
+ pipeline_layout_info.setLayoutCount = 1;128
+ pipeline_layout_info.pSetLayouts = &set_layout;129
+ pipeline_layout_info.pushConstantRangeCount = 1;130
+ pipeline_layout_info.pPushConstantRanges = &push_range;131
+ VkPipelineLayout pipeline_layout = VK_NULL_HANDLE;132
+ vkCreatePipelineLayout(device, &pipeline_layout_info, nullptr, &pipeline_layout);133
+134
+ VkShaderModuleCreateInfo shader_info{};135
+ shader_info.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;136
+ shader_info.codeSize = sizeof(matrix_multiply_spv);137
+ shader_info.pCode = matrix_multiply_spv;138
+ VkShaderModule shader_module = VK_NULL_HANDLE;139
+ vkCreateShaderModule(device, &shader_info, nullptr, &shader_module);140
+141
+ VkComputePipelineCreateInfo pipeline_info{};142
+ pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;143
+ pipeline_info.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;144
+ pipeline_info.stage.stage = VK_SHADER_STAGE_COMPUTE_BIT;145
+ pipeline_info.stage.module = shader_module;146
+ pipeline_info.stage.pName = "main";147
+ pipeline_info.layout = pipeline_layout;148
+ VkPipeline pipeline = VK_NULL_HANDLE;149
+ vkCreateComputePipelines(device, VK_NULL_HANDLE, 1, &pipeline_info, nullptr, &pipeline);150
+151
+ VkDescriptorPoolSize pool_size{};152
+ pool_size.type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;153
+ pool_size.descriptorCount = 3;154
+ VkDescriptorPoolCreateInfo pool_info{};155
+ pool_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO;156
+ pool_info.maxSets = 1;157
+ pool_info.poolSizeCount = 1;158
+ pool_info.pPoolSizes = &pool_size;159
+ VkDescriptorPool desc_pool = VK_NULL_HANDLE;160
+ vkCreateDescriptorPool(device, &pool_info, nullptr, &desc_pool);161
+162
+ VkDescriptorSetAllocateInfo set_alloc_info{};163
+ set_alloc_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO;164
+ set_alloc_info.descriptorPool = desc_pool;165
+ set_alloc_info.descriptorSetCount = 1;166
+ set_alloc_info.pSetLayouts = &set_layout;167
+ VkDescriptorSet desc_set = VK_NULL_HANDLE;168
+ vkAllocateDescriptorSets(device, &set_alloc_info, &desc_set);169
+170
+ VkDescriptorBufferInfo buf_infos[3] = {171
+ {buf_a.buffer, 0, bytes},172
+ {buf_b.buffer, 0, bytes},173
+ {buf_c.buffer, 0, bytes},174
+ };175
+ VkWriteDescriptorSet writes[3] = {};176
+ for (int i = 0; i < 3; i++) {177
+ writes[i].sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET;178
+ writes[i].dstSet = desc_set;179
+ writes[i].dstBinding = (uint32_t)i;180
+ writes[i].descriptorCount = 1;181
+ writes[i].descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;182
+ writes[i].pBufferInfo = &buf_infos[i];183
+ }184
+ vkUpdateDescriptorSets(device, 3, writes, 0, nullptr);185
+186
+ VkCommandBufferAllocateInfo cmd_alloc_info{};187
+ cmd_alloc_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;188
+ cmd_alloc_info.commandPool = ctx.command_pool;189
+ cmd_alloc_info.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;190
+ cmd_alloc_info.commandBufferCount = 1;191
+ VkCommandBuffer cmd = VK_NULL_HANDLE;192
+ vkAllocateCommandBuffers(device, &cmd_alloc_info, &cmd);193
+194
+ VkCommandBufferBeginInfo begin_info{};195
+ begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;196
+ begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;197
+ vkBeginCommandBuffer(cmd, &begin_info);198
+ vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline);199
+ vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline_layout, 0, 1, &desc_set, 0, nullptr);200
+ uint32_t n_u32 = (uint32_t)n;201
+ vkCmdPushConstants(cmd, pipeline_layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &n_u32);202
+ uint32_t groups = ((uint32_t)n + 15u) / 16u; // matches local_size_x/y = 16 in the shader203
+ vkCmdDispatch(cmd, groups, groups, 1);204
+ vkEndCommandBuffer(cmd);205
+206
+ VkSubmitInfo submit_info{};207
+ submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;208
+ submit_info.commandBufferCount = 1;209
+ submit_info.pCommandBuffers = &cmd;210
+ vkQueueSubmit(ctx.compute_queue, 1, &submit_info, VK_NULL_HANDLE);211
+ vkQueueWaitIdle(ctx.compute_queue); // simplest correct sync: block until done, no overlap with other work212
+213
+ memcpy(c, buf_c.mapped, (size_t)bytes);214
+215
+ vkFreeCommandBuffers(device, ctx.command_pool, 1, &cmd);216
+ vkDestroyDescriptorPool(device, desc_pool, nullptr);217
+ vkDestroyPipeline(device, pipeline, nullptr);218
+ vkDestroyShaderModule(device, shader_module, nullptr);219
+ vkDestroyPipelineLayout(device, pipeline_layout, nullptr);220
+ vkDestroyDescriptorSetLayout(device, set_layout, nullptr);221
+ destroy_buffer(device, buf_a);222
+ destroy_buffer(device, buf_b);223
+ destroy_buffer(device, buf_c);224
+}pyre/src/C/runtime/device.cadded
+55−0@@ -0,0 +1,55 @@1
+#include "H/runtime/ash_gpu.h"2
+#include "H/backends/vulkan/context.h"3
+#include "H/backends/cpu/matrix_multiply.h"4
+#include "H/backends/cpu/chaos.h"5
+#include <string.h>6
+7
+/*8
+ * Real detection: tries to actually bring up a Vulkan device (not just9
+ * check a loader is present -- see ash_gpu_vulkan_init(), src/C/backends/vulkan/context.cpp).10
+ * This turned out to be WSL2 with GPU passthrough available but no11
+ * Vulkan driver installed at first -- installing vulkan-dzn (Mesa's12
+ * Vulkan-over-D3D12 driver) made this report ASH_GPU_BACKEND_VULKAN for13
+ * real, naming the user's actual AMD Radeon RX 9070 XT. See14
+ * src/C/backends/vulkan/context.cpp's header comment.15
+ */16
+AshGpuInfo ash_gpu_detect(void) {17
+ AshGpuInfo info;18
+ info.vram_bytes = 0;19
+20
+ if (ash_gpu_vulkan_init()) {21
+ info.backend = ASH_GPU_BACKEND_VULKAN;22
+ strncpy(info.device_name, ash_gpu_vulkan_device_name(), sizeof(info.device_name) - 1);23
+ } else {24
+ info.backend = ASH_GPU_BACKEND_NONE;25
+ strncpy(info.device_name, "CPU fallback", sizeof(info.device_name) - 1);26
+ }27
+ info.device_name[sizeof(info.device_name) - 1] = '\0';28
+ return info;29
+}30
+31
+int ash_gpu_available(void) {32
+ return ash_gpu_detect().backend != ASH_GPU_BACKEND_NONE;33
+}34
+35
+// Both dispatchers unconditionally prefer GPU-when-available over CPU,36
+// without comparing measured cost first -- matrix_multiply's own real37
+// measurement says that's the WRONG call for that kernel on this38
+// hardware (ash_gpu.h's comment on it); chaos_iterate exists precisely39
+// because it's shaped to actually win on GPU instead. Neither dispatcher40
+// has been taught to tell the difference yet -- see ideas/assigned.md.41
+void ash_gpu_matrix_multiply(const double *a, const double *b, double *c, int n) {42
+ if (ash_gpu_available()) {43
+ ash_gpu_vulkan_matrix_multiply(a, b, c, n);44
+ return;45
+ }46
+ ash_gpu_cpu_matrix_multiply(a, b, c, n);47
+}48
+49
+void ash_gpu_chaos_iterate(double *data, int n, int iterations) {50
+ if (ash_gpu_available()) {51
+ ash_gpu_vulkan_chaos_iterate(data, n, iterations);52
+ return;53
+ }54
+ ash_gpu_cpu_chaos_iterate(data, n, iterations);55
+}pyre/src/D/backends/cpu/chaos.dadded
+3−0@@ -0,0 +1,3 @@1
+src/O/backends/cpu/chaos.o: src/C/backends/cpu/chaos.c \2
+ src/H/backends/cpu/chaos.h3
+src/H/backends/cpu/chaos.h:pyre/src/D/backends/cpu/matrix_multiply.dadded
+4−0@@ -0,0 +1,4 @@1
+src/O/backends/cpu/matrix_multiply.o: \2
+ src/C/backends/cpu/matrix_multiply.c \3
+ src/H/backends/cpu/matrix_multiply.h4
+src/H/backends/cpu/matrix_multiply.h:pyre/src/D/backends/cpu/monte_carlo.dadded
+3−0@@ -0,0 +1,3 @@1
+src/O/backends/cpu/monte_carlo.o: src/C/backends/cpu/monte_carlo.c \2
+ src/H/runtime/ash_gpu.h3
+src/H/runtime/ash_gpu.h:pyre/src/D/backends/vulkan/chaos.dadded
+7−0@@ -0,0 +1,7 @@1
+src/O/backends/vulkan/chaos.o: src/C/backends/vulkan/chaos.cpp \2
+ src/H/backends/vulkan/context.h \3
+ src/H/backends/vulkan/context_internal.hpp \4
+ src/H/backends/vulkan/chaos_spv.h5
+src/H/backends/vulkan/context.h:6
+src/H/backends/vulkan/context_internal.hpp:7
+src/H/backends/vulkan/chaos_spv.h:pyre/src/D/backends/vulkan/context.dadded
+5−0@@ -0,0 +1,5 @@1
+src/O/backends/vulkan/context.o: src/C/backends/vulkan/context.cpp \2
+ src/H/backends/vulkan/context.h \3
+ src/H/backends/vulkan/context_internal.hpp4
+src/H/backends/vulkan/context.h:5
+src/H/backends/vulkan/context_internal.hpp:pyre/src/D/backends/vulkan/matrix_multiply.dadded
+8−0@@ -0,0 +1,8 @@1
+src/O/backends/vulkan/matrix_multiply.o: \2
+ src/C/backends/vulkan/matrix_multiply.cpp \3
+ src/H/backends/vulkan/context.h \4
+ src/H/backends/vulkan/context_internal.hpp \5
+ src/H/backends/vulkan/matrix_multiply_spv.h6
+src/H/backends/vulkan/context.h:7
+src/H/backends/vulkan/context_internal.hpp:8
+src/H/backends/vulkan/matrix_multiply_spv.h:pyre/src/D/runtime/device.dadded
+7−0@@ -0,0 +1,7 @@1
+src/O/runtime/device.o: src/C/runtime/device.c src/H/runtime/ash_gpu.h \2
+ src/H/backends/vulkan/context.h src/H/backends/cpu/matrix_multiply.h \3
+ src/H/backends/cpu/chaos.h4
+src/H/runtime/ash_gpu.h:5
+src/H/backends/vulkan/context.h:6
+src/H/backends/cpu/matrix_multiply.h:7
+src/H/backends/cpu/chaos.h:pyre/src/H/backends/cpu/chaos.hadded
+6−0@@ -0,0 +1,6 @@1
+#ifndef ASH_GPU_CPU_CHAOS_H2
+#define ASH_GPU_CPU_CHAOS_H3
+// Internal to the dispatcher (src/C/runtime/device.c) -- not part of4
+// ash_gpu.h's public ABI. Same contract as ash_gpu_chaos_iterate.5
+void ash_gpu_cpu_chaos_iterate(double *data, int n, int iterations);6
+#endifpyre/src/H/backends/cpu/matrix_multiply.hadded
+6−0@@ -0,0 +1,6 @@1
+#ifndef ASH_GPU_CPU_MATRIX_MULTIPLY_H2
+#define ASH_GPU_CPU_MATRIX_MULTIPLY_H3
+// Internal to the dispatcher (src/C/runtime/device.c) -- not part of4
+// ash_gpu.h's public ABI. Same contract as ash_gpu_matrix_multiply.5
+void ash_gpu_cpu_matrix_multiply(const double *a, const double *b, double *c, int n);6
+#endifpyre/src/H/backends/vulkan/context.hadded
+41−0@@ -0,0 +1,41 @@1
+#ifndef ASH_GPU_VULKAN_CONTEXT_H2
+#define ASH_GPU_VULKAN_CONTEXT_H3
+// Internal to the Vulkan backend (not part of ash_gpu.h's public ABI):4
+// a process-lifetime instance/device singleton, lazily created on first5
+// use, plus the one kernel implemented so far. C linkage throughout --6
+// runtime/device.c (plain C) calls ash_gpu_vulkan_init() to decide what7
+// ash_gpu_detect() reports, even though this backend's own .cpp files8
+// use Vulkan's C++-friendly RAII patterns internally.9
+#ifdef __cplusplus10
+extern "C" {11
+#endif12
+13
+// Creates a VkInstance and picks the first physical device that has a14
+// compute-capable queue family AND the shaderFloat64 feature (ash_gpu's15
+// C ABI is `double` throughout -- see H/runtime/ash_gpu.h -- so a device16
+// that can only do fp32 compute isn't a match for this kernel). Returns17
+// 1 if a usable device is now ready to dispatch work, 0 for any other18
+// outcome (no Vulkan loader, no device, no fp64 support) -- all of19
+// those are normal "no GPU here" answers, never a reason to abort.20
+// Idempotent: after the first call, later calls just return the cached21
+// result instead of re-probing.22
+int ash_gpu_vulkan_init(void);23
+24
+// Valid only after ash_gpu_vulkan_init() returned 1. The physical25
+// device's name (VkPhysicalDeviceProperties.deviceName), e.g. "AMD26
+// Radeon RX 9070 XT" -- for ash_gpu_detect()'s AshGpuInfo.device_name.27
+const char *ash_gpu_vulkan_device_name(void);28
+29
+// Only valid once ash_gpu_vulkan_init() has returned 1. Same contract as30
+// ash_gpu_matrix_multiply (ash_gpu.h): C = A * B, all three N x N,31
+// row-major, contiguous, C must not alias A or B.32
+void ash_gpu_vulkan_matrix_multiply(const double *a, const double *b, double *c, int n);33
+34
+// Same contract as ash_gpu_chaos_iterate (ash_gpu.h): data[i] =35
+// iterate(data[i]) in place, `iterations` times each.36
+void ash_gpu_vulkan_chaos_iterate(double *data, int n, int iterations);37
+38
+#ifdef __cplusplus39
+}40
+#endif41
+#endifpyre/src/H/backends/vulkan/context_internal.hppadded
+23−0@@ -0,0 +1,23 @@1
+#ifndef ASH_GPU_VULKAN_CONTEXT_INTERNAL_HPP2
+#define ASH_GPU_VULKAN_CONTEXT_INTERNAL_HPP3
+// C++-only (no extern "C" -- this is never called from matrix.c/device.c4
+// directly, only from other .cpp files in this same backend) companion5
+// to context.h: the actual Vulkan handles context.cpp's6
+// ash_gpu_vulkan_init() creates, that matrix_multiply.cpp needs to issue7
+// its own dispatch against the same instance/device/queue instead of8
+// creating a second one.9
+#include <vulkan/vulkan.h>10
+11
+struct AshGpuVulkanContext {12
+ VkInstance instance = VK_NULL_HANDLE;13
+ VkPhysicalDevice physical_device = VK_NULL_HANDLE;14
+ VkDevice device = VK_NULL_HANDLE;15
+ VkQueue compute_queue = VK_NULL_HANDLE;16
+ uint32_t compute_queue_family = 0;17
+ VkCommandPool command_pool = VK_NULL_HANDLE;18
+};19
+20
+// Valid only after ash_gpu_vulkan_init() (context.h) has returned 1.21
+const AshGpuVulkanContext &ash_gpu_vulkan_context(void);22
+23
+#endifpyre/src/H/runtime/ash_gpu.hadded
+115−0@@ -0,0 +1,115 @@1
+#ifndef ASH_GPU_H2
+#define ASH_GPU_H3
+/*4
+ * Pyre's public C ABI (ideas/gpu_acceleration.md's "Components" section:5
+ * a small C ABI, Kiln never needs to know backend internals). Phase 16
+ * ("standalone runtime... don't touch Ash yet") -- nothing here is wired7
+ * into kiln or ashvm yet, this is the foundation that integration will8
+ * later call into.9
+ *10
+ * C linkage throughout, even though backends may be implemented in C++11
+ * internally (the ideas doc's own "Why C++ for the Runtime" section),12
+ * so kiln's codegen (plain C, no name mangling to reason about) and any13
+ * other C caller can link against this header unmodified.14
+ */15
+#include <stddef.h>16
+17
+#ifdef __cplusplus18
+extern "C" {19
+#endif20
+21
+typedef enum {22
+ ASH_GPU_BACKEND_NONE, /* no usable GPU -- every call below runs on CPU */23
+ ASH_GPU_BACKEND_VULKAN,24
+ ASH_GPU_BACKEND_HIP,25
+ ASH_GPU_BACKEND_CUDA26
+} AshGpuBackend;27
+28
+typedef struct {29
+ AshGpuBackend backend;30
+ char device_name[128]; /* "CPU fallback" when backend == ASH_GPU_BACKEND_NONE */31
+ unsigned long long vram_bytes; /* 0 when there is no discrete device */32
+} AshGpuInfo;33
+34
+/* Really tries to bring up a Vulkan device (src/C/runtime/device.c calls35
+ * src/C/backends/vulkan/context.cpp's ash_gpu_vulkan_init(), not just a36
+ * loader-presence check) and reports what it finds. This sandbox turns37
+ * out to be WSL2 with GPU passthrough available but no Vulkan driver38
+ * installed at first (vulkaninfo: "Found no drivers!"); installing39
+ * vulkan-dzn (Mesa's Vulkan-over-D3D12 driver, pyre/tools/install_deps.sh40
+ * doesn't cover it -- it's GPU-driver-specific, not SDK/build tooling)41
+ * made ash_gpu_detect() report ASH_GPU_BACKEND_VULKAN for real, naming42
+ * the user's actual AMD Radeon RX 9070 XT. Never fails: "no GPU" is a43
+ * normal, fully supported answer, never an error (design principle 1 in44
+ * ideas/gpu_acceleration.md). */45
+AshGpuInfo ash_gpu_detect(void);46
+47
+/* True iff ash_gpu_detect() found a usable (non-CPU-fallback) device. */48
+int ash_gpu_available(void);49
+50
+/*51
+ * C = A * B, all three N x N, row-major, contiguous (flat, not an array52
+ * of row pointers -- see todo.md 2.A's "Contiguous Buffer Type": this is53
+ * the layout a future Ash-level Buffer/Tensor type would hand across54
+ * this exact boundary with zero copying). C must not alias A or B.55
+ *56
+ * Dispatches to whatever ash_gpu_detect() found (src/C/runtime/device.c):57
+ * Vulkan (src/C/backends/vulkan/matrix_multiply.cpp, a straightforward58
+ * one-thread-per-output-element compute shader -- now verified to59
+ * actually run, on the RX 9070 XT via WSL2's D3D12 passthrough) if60
+ * available, else the CPU backend (src/C/backends/cpu/matrix_multiply.c:61
+ * cache-friendly loop order plus OpenMP, no GPU required at all --62
+ * design principle 1: no GPU means "best CPU implementation", never63
+ * "can't run"). Measured result, same hardware: the GPU path is64
+ * currently ~2.5x SLOWER than the CPU path for this kernel (fp6465
+ * throughput on a consumer GPU, no workgroup tiling, host-visible66
+ * memory instead of device-local, and the Dozen driver itself is a67
+ * translation layer, not native) -- see ash_gpu_chaos_iterate below for68
+ * a workload shaped so the GPU actually wins, and ideas/assigned.md for69
+ * why this function still unconditionally prefers GPU-when-available70
+ * despite that (not yet fixed to measure first).71
+ */72
+void ash_gpu_matrix_multiply(const double *a, const double *b, double *c, int n);73
+74
+/*75
+ * The second benchmark kernel (pyre/benchmarks/monte_carlo_no_pi): run76
+ * `iterations` draws of the same two-draw-per-iteration LCG recurrence77
+ * as mcnp.cpp's reference, count how many exceed `threshold`, and78
+ * report both the raw count and the estimated probability.79
+ *80
+ * The CPU backend splits `iterations` across threads with independent81
+ * per-thread seed streams (see src/C/backends/cpu/monte_carlo.c) -- the exact82
+ * sequence of draws, and so the exact risky count, therefore differs83
+ * from a single-threaded run with the same seed. The benchmark's own84
+ * target is wall-clock time, not bit-for-bit reproduction of one85
+ * particular single-threaded trajectory.86
+ */87
+void ash_gpu_monte_carlo_risk(long long iterations, double threshold, unsigned int seed,88
+ long long *out_risky, double *out_probability);89
+90
+/*91
+ * In place: data[i] = iterate(data[i]) for each of the n elements,92
+ * where iterate() applies x = sin(x) + cos(x * 1.5), `iterations`93
+ * times. Bounded (sin+cos of a bounded argument never escapes roughly94
+ * [-2, 2]), so it runs indefinitely without overflow -- purely a95
+ * workload shape, not a model of anything.96
+ *97
+ * Deliberately the OPPOSITE shape from matrix_multiply: huge compute98
+ * per element (iterations can be thousands), trivial data movement99
+ * (one double in, one double out, no cross-element dependency, no100
+ * tiling to get right). Added specifically because matrix_multiply's101
+ * real, measured GPU result was a loss (see that function's comment) --102
+ * this is the workload ideas/gpu_acceleration.md's own dispatch103
+ * principle describes as GPU-favorable ("huge workload" with high104
+ * arithmetic intensity, where per-dispatch overhead and fp64-throughput105
+ * cost get amortized over thousands of iterations instead of dominating106
+ * a handful). Same dispatch rule as matrix_multiply: GPU if107
+ * ash_gpu_available(), else CPU (src/C/backends/cpu/chaos.c, OpenMP108
+ * across elements) -- never "can't run".109
+ */110
+void ash_gpu_chaos_iterate(double *data, int n, int iterations);111
+112
+#ifdef __cplusplus113
+}114
+#endif115
+#endifpyre/src/O/backends/cpu/.gitkeepadded
+0−0No content changes. View the file on GitHub
pyre/src/O/backends/vulkan/.gitkeepadded
+0−0No content changes. View the file on GitHub
pyre/src/O/runtime/.gitkeepadded
+0−0No content changes. View the file on GitHub
pyre/src/shaders/chaos.compadded
+35−0@@ -0,0 +1,35 @@1
+#version 4502
+// See src/H/runtime/ash_gpu.h's comment on ash_gpu_chaos_iterate for why3
+// this kernel exists and what shape it's deliberately given: huge4
+// compute per element, trivial data movement, no cross-element5
+// dependency at all -- the opposite of matrix_multiply.comp.6
+#extension GL_ARB_gpu_shader_fp64 : require7
+8
+layout(local_size_x = 256) in;9
+10
+layout(std430, set = 0, binding = 0) buffer DataBuf { double data[]; };11
+12
+layout(push_constant) uniform PushConstants {13
+ uint n;14
+ uint iterations;15
+} pc;16
+17
+void main() {18
+ uint i = gl_GlobalInvocationID.x;19
+ if (i >= pc.n) return;20
+21
+ // sqrt/abs are the GL_ARB_gpu_shader_fp64 double overloads that22
+ // actually exist -- sin/cos do NOT get a double overload from that23
+ // extension (GLSL keeps transcendental functions float-only even24
+ // with it enabled; found by trying it: glslc rejected sin(double)25
+ // outright). Bounded (abs(x)+1 >= 1, so sqrt of it stays small) and26
+ // numerically stable regardless of how many iterations run; the27
+ // exact converged value doesn't matter, only that every iteration28
+ // does real, data-dependent floating-point work no compiler can29
+ // skip or hoist out of the loop.30
+ double x = data[i];31
+ for (uint k = 0u; k < pc.iterations; k++) {32
+ x = sqrt(abs(x) + 1.0lf) * 0.99999lf + 0.00001lf;33
+ }34
+ data[i] = x;35
+}pyre/src/shaders/matrix_multiply.compadded
+38−0@@ -0,0 +1,38 @@1
+#version 4502
+// GL_ARB_gpu_shader_fp64: ash_gpu's C ABI is `double` throughout3
+// (pyre/src/H/runtime/ash_gpu.h), so this kernel has to match the CPU4
+// backend's precision exactly rather than silently computing in float32.5
+// Caveat, since there's no hardware here to measure it on: consumer6
+// GPUs (including the RX 9070 XT this is meant to run on) typically run7
+// fp64 at a small fraction of their fp32 throughput, so this kernel may8
+// not actually beat the CPU backend -- that's a real open question for9
+// whoever benchmarks it on real hardware, not assumed away here.10
+#extension GL_ARB_gpu_shader_fp64 : require11
+12
+// One thread per output element, no shared-memory tiling. Deliberately13
+// the simplest correct version, not a tuned one: nothing in this14
+// sandbox can run a compute shader (no GPU, no ICD -- see15
+// pyre/src/C/backends/vulkan/context.cpp), so a tiling bug introduced16
+// here would be undetectable until someone runs it on real hardware.17
+// Tile it once there's a real device to benchmark and debug against.18
+layout(local_size_x = 16, local_size_y = 16) in;19
+20
+layout(std430, set = 0, binding = 0) readonly buffer ABuf { double a[]; };21
+layout(std430, set = 0, binding = 1) readonly buffer BBuf { double b[]; };22
+layout(std430, set = 0, binding = 2) writeonly buffer CBuf { double c[]; };23
+24
+layout(push_constant) uniform PushConstants {25
+ uint n;26
+} pc;27
+28
+void main() {29
+ uint row = gl_GlobalInvocationID.y;30
+ uint col = gl_GlobalInvocationID.x;31
+ if (row >= pc.n || col >= pc.n) return;32
+33
+ double sum = 0.0;34
+ for (uint k = 0u; k < pc.n; k++) {35
+ sum += a[row * pc.n + k] * b[k * pc.n + col];36
+ }37
+ c[row * pc.n + col] = sum;38
+}pyre/tests/matrix_multiply_test.cadded
+54−0@@ -0,0 +1,54 @@1
+#include "H/runtime/ash_gpu.h"2
+#include <stdio.h>3
+#include <stdlib.h>4
+#include <math.h>5
+6
+/* Small enough to compute a trustworthy reference in plain scalar code7
+ * right here, instead of trusting the optimized kernel to check itself. */8
+static void naive_multiply(const double *a, const double *b, double *c, int n) {9
+ for (int i = 0; i < n; i++) {10
+ for (int j = 0; j < n; j++) {11
+ double sum = 0.0;12
+ for (int k = 0; k < n; k++) sum += a[i * n + k] * b[k * n + j];13
+ c[i * n + j] = sum;14
+ }15
+ }16
+}17
+18
+int main(void) {19
+ const int n = 37; /* deliberately not a round number */20
+ double *a = malloc((size_t)n * n * sizeof(double));21
+ double *b = malloc((size_t)n * n * sizeof(double));22
+ double *c = malloc((size_t)n * n * sizeof(double));23
+ double *expected = malloc((size_t)n * n * sizeof(double));24
+25
+ for (int i = 0; i < n * n; i++) {26
+ a[i] = (double)((i % 13) - 6) * 0.5;27
+ b[i] = (double)((i % 7) - 3) * 1.5;28
+ }29
+30
+ naive_multiply(a, b, expected, n);31
+ ash_gpu_matrix_multiply(a, b, c, n);32
+33
+ int failures = 0;34
+ for (int i = 0; i < n * n; i++) {35
+ if (fabs(c[i] - expected[i]) > 1e-9 * (fabs(expected[i]) + 1.0)) {36
+ if (failures < 5) {37
+ fprintf(stderr, "mismatch at %d: got %.17g, expected %.17g\n", i, c[i], expected[i]);38
+ }39
+ failures++;40
+ }41
+ }42
+43
+ free(a);44
+ free(b);45
+ free(c);46
+ free(expected);47
+48
+ if (failures) {49
+ fprintf(stderr, "matrix_multiply_test: FAILED (%d mismatches)\n", failures);50
+ return 1;51
+ }52
+ printf("matrix_multiply_test: OK\n");53
+ return 0;54
+}pyre/tests/monte_carlo_test.cadded
+37−0@@ -0,0 +1,37 @@1
+#include "H/runtime/ash_gpu.h"2
+#include <stdio.h>3
+4
+int main(void) {5
+ long long risky = 0;6
+ double probability = 0.0;7
+ const long long iterations = 10000000;8
+9
+ ash_gpu_monte_carlo_risk(iterations, 0.75, 123456789u, &risky, &probability);10
+11
+ int ok = 1;12
+ /* Self-consistency: the reported probability must be exactly the13
+ * reported count over the iteration count, regardless of how the14
+ * threads split the work. */15
+ double recomputed = (double)risky / (double)iterations;16
+ if (probability != recomputed) {17
+ fprintf(stderr, "monte_carlo_test: probability %.17g != risky/iterations %.17g\n",18
+ probability, recomputed);19
+ ok = 0;20
+ }21
+ /* Loose sanity bound on P(X*Y > 0.75) for independent X,Y ~ Uniform(0,1):22
+ * analytically 0.75 - 0.75*ln(0.75) ~= 0.0345, so anything wildly off23
+ * (e.g. every/no draw flagged) means the kernel is broken, not just24
+ * imprecise. */25
+ if (probability < 0.01 || probability > 0.08) {26
+ fprintf(stderr, "monte_carlo_test: probability %.6f outside sane range [0.01, 0.08]\n",27
+ probability);28
+ ok = 0;29
+ }30
+31
+ if (!ok) {32
+ fprintf(stderr, "monte_carlo_test: FAILED\n");33
+ return 1;34
+ }35
+ printf("monte_carlo_test: OK (risky=%lld, probability=%.6f)\n", risky, probability);36
+ return 0;37
+}pyre/todo/todo.mdadded
+101−0@@ -0,0 +1,101 @@1
+# TODO: Kiln, Ash, and Pyre Hardware Acceleration Roadmap2
+3
+## Objective4
+Enable full-scale matrix operations ($500 \times 500$ and beyond) without hitting memory ceilings or segfaults, update the runtime engines (`kiln` and `ashvm`), and establish the necessary language features in Ash for **Pyre** (`ash-gpu`) to deliver massive GPU acceleration.5
+6
+---7
+8
+## 1. Kiln Runtime & Compiler Upgrades9
+10
+### A. Memory Management11
+- [x] **Remove Hardcoded Heap Cap:**12
+ - Replace static `static char heap[16 * 1024 * 1024]` allocation in Kiln's native runtime generator with dynamic OS page allocations.13
+ - **Linux:** Implement `mmap` (`sys_mmap`, syscall 9 on x86-64) or `brk` to request pages dynamically when heap memory is exhausted.14
+ - **Windows:** Emit `VirtualAlloc` / `HeapAlloc` imports via `kernel32.dll`.15
+ - Done 2026-10-08: grows in 64MB chunks (`brk` again on Linux, commits more of an upfront 4GB `MEM_RESERVE` on Windows) instead of one fixed grab -- see CHANGELOG. `tests/cases/memory/heap_growth.ash`.16
+- [x] **Arena / Temporary Region Allocators:**17
+ - Implement checkpoint-based region allocation (`heap_checkpoint()` / `heap_reset()`) inside loop scopes so short-lived inner-loop temporaries don't exhaust memory when GC isn't active.18
+ - Done 2026-10-08, partially: `heap_emit_checkpoint`/`heap_emit_reset` exist and are verified (see CHANGELOG), but are NOT wired into general `each`/`during` loop bodies -- that needs an escape analysis kiln's single-pass compiler doesn't have yet (see goals.md item 1), and an unsound version would corrupt memory rather than just leak it. They're ready for Pyre's own lowering to call directly, where the generated loop's shape is fully under kiln's own control instead of arbitrary user code.19
+20
+### B. Compiler AST & Syntax Support21
+- [x] **Chained Multi-Index Assignment:**22
+ - Update the parser and code emitter in `Kiln` (`src/compiler/`) to support nested array index mutation (`C[i][j] = val`).23
+ - Emit direct dereference chains rather than requiring manual intermediate extraction (`local row = C[i]; row[j] = val; C[i] = row;`).24
+ - Done 2026-10-08: see CHANGELOG. `tests/cases/collections/chained_index_assign.ash`.25
+26
+---27
+28
+## 2. Ash Language Extensions (Required for Pyre)29
+30
+Even if these do not yet exist in Ash, define the grammar and compiler IR to support them:31
+32
+### A. Flat & Contiguous Memory Types33
+- [ ] **Contiguous Buffer Type (`Buffer` / `Tensor`):**34
+ - Add native contiguous 1D/2D array buffers (e.g., `Buffer.alloc(type, dimensions)` or `Tensor([N, N])`).35
+ - Avoid storing matrices as nested pointer arrays (`Array<Array<T>>`) to enable zero-copy GPU transfers and $O(1)$ flat indexing ($i \cdot N + j$).36
+37
+### B. Structural Loop Construct & Parallel Map Hints38
+- [ ] **Parallel Loop / Acceleration Annotations:**39
+ - Add language constructs or attributes indicating data independence for parallel dispatch:40
+ ```ash41
+ parallel for i in 0..N { ... }42
+ ```43
+ - Or vector-level primitives:44
+ ```ash45
+ local C = matrix_mul(A, B);46
+ ```47
+48
+### C. Foreign Function Interface (FFI) & C ABI Binding49
+- [x] **C ABI Binding Primitive, for one specific function:**50
+ - Add native low-level foreign function bindings (`@extern` or `@import_c`) so Ash runtime scripts can pass buffer pointers directly to Pyre's C ABI layer (`ash_gpu_dispatch`, `ash_gpu_upload`, `ash_gpu_download`).51
+ - Done 2026-10-08, narrowly: not a general `@extern`/`@import_c` primitive (no Ash-level syntax for calling ARBITRARY C functions exists) -- instead, `kiln`'s `matrix_mul(a, b)` builtin (`codegen/C/collections/matrix_mul.c`) calls `ash_gpu_matrix_multiply` specifically, through the GOT slot `elf/C/elf_dynamic.c`'s `--link=shared` writer lays out for it. A real general FFI primitive is still open; this is the one concrete binding that exists.52
+53
+---54
+55
+## 3. Pyre (`ash-gpu`) Integration Architecture56
+57
+- [ ] **Data Residency Pipeline:**58
+ - Enforce upload-once, download-once semantics: keep intermediate computation results in VRAM across chained standard library calls (`blur()`, `matrix_mul()`, `resize()`).59
+ - Not started -- there is no GPU backend yet for anything to be resident on (see below).60
+- [~] **Vulkan / HIP / CUDA Runtime Dispatch:**61
+ - Detect compute hardware at launch (`ash_gpu_available()`).62
+ - Automatically route small matrix tasks ($N < 32$) to CPU SIMD and large matrix tasks ($N \ge 128$) to Pyre compute pipelines.63
+ - `ash_gpu_available()` really tries Vulkan (`ash_gpu_vulkan_init()`, `src/C/backends/vulkan/context.cpp`) and, as of 2026-10-08, actually succeeds: this sandbox is WSL2, which has GPU passthrough infrastructure (`/dev/dxg`, Microsoft's D3D12 translation libraries) but no Vulkan driver installed by default. Installing `vulkan-dzn` (Mesa's Vulkan-over-D3D12 driver) made `vulkaninfo` and `ash_gpu_detect()` both find the user's real **AMD Radeon RX 9070 XT** (reported as "Microsoft Direct3D12 (AMD Radeon RX 9070 XT)", driver `DRIVER_ID_MESA_DOZEN`, `shaderFloat64 = true`). `pyre/tools/install_deps.sh` doesn't cover this -- it's a GPU driver, not SDK/build tooling; install it by hand (`pacman -S vulkan-dzn mesa`) on any other WSL2 machine that needs it. First real execution of the Vulkan path, ever -- see the benchmark table below for what it found. HIP/CUDA: not started. Size-based routing ($N<32$/$N\ge128$) doesn't exist at the Pyre level -- `ash_gpu_matrix_multiply` and `ash_gpu_chaos_iterate` both dispatch purely on backend availability, not size or any measured cost, for every workload; the one size threshold that exists ($\text{ASH\_GPU\_MIN\_N}=64$) is in ashvm's builtin, not here. **This is exactly the gap flagged below as the next real task** (see "Smart dispatch").64
+- [x] **Fallback Mechanism (for the kernels that exist so far):**65
+ - On systems with no dedicated GPU or incompatible drivers, transparently fall back to multi-threaded CPU routines without raising runtime errors.66
+ - Done 2026-10-08 for `ash_gpu_matrix_multiply` and `ash_gpu_monte_carlo_risk`: both always run on the CPU backend today (ikj-reordered + OpenMP for matmul; per-thread-seeded OpenMP for the Monte Carlo kernel), which is also exactly what "fallback" means once a real GPU path exists later.67
+68
+### Phase 1 benchmarks (pyre/benchmarks/) -- CPU backend only, no GPU69
+Both of the user's benchmark challenges ("beat the C++ reference by at least 50%, or the solution isn't allowed"), run fresh on this machine, CPU-only (`make benchmark` in `pyre/`):70
+71
+| Benchmark | C++ reference | Pyre (CPU backend) | Reduction |72
+|---|---|---|---|73
+| `matrix_multiply` (2000x2000) | 5.532s (`g++ -O2`) | 0.223-0.248s | ~95.6% (target: 50%) |74
+| `monte_carlo_no_pi` (1e9 iters) | 1.840s (`g++ -O3`) | 0.231-0.246s | ~87% (target: 50%) |75
+76
+Both cleared the bar without touching a GPU at all: `matrix_multiply` from the ijk-vs-ikj loop reorder (fixes the reference's cache-hostile strided access to `B[k][j]`) plus splitting rows across this machine's 8 cores with OpenMP; `monte_carlo_no_pi` from giving each of the 8 threads its own splitmix32-derived seed stream instead of one sequential chain. `pyre/tests/{matrix_multiply,monte_carlo}_test.c` check correctness (a from-scratch scalar reference for matmul; self-consistency plus an analytic sanity range for the Monte Carlo probability), independent of the benchmark numbers above.77
+78
+### Ash bindings and kiln integration (ideas/gpu_acceleration.md Phases 3 and 4) -- both done79
+`matrix_mul(a, b)` exists as a plain builtin on `ashvm` (`ashvm/src/builtins/C/matrix.c`), calling `ash_gpu_matrix_multiply` directly for square arrays-of-arrays of numbers at or above `ASH_GPU_MIN_N = 64`, a local C loop below it. `ashvm/tests/cases/collections/matrix_mul.ash` covers both paths plus the error cases.80
+81
+`kiln` now has the same builtin (`codegen/C/collections/matrix_mul.c`), available only under `--link=shared`: it really calls through the GOT into `libashgpu.so` (a second `DT_NEEDED` added to `elf/C/elf_dynamic.c`'s writer, folded into the same combined hash/dynsym/GOT/rela table as kiln's own runtime imports), not a native reimplementation -- see ideas/assigned.md for the full mechanism, including the SysV ABI call wrapper (`codegen/C/platform/linux_call.c`) kiln never needed before this (its own hand-rolled routines are alignment-agnostic; a real gcc/g++-compiled function with OpenMP/SSE isn't). Under `--link=static`, `matrix_mul` is a clear compile error, not a silent no-op -- that mode's whole point is staying dependency-free. `kiln/tests/scripts/shared_link_check.sh` (new) covers both the general `--link=shared` path (previously untested at all) and `matrix_mul` specifically.82
+83
+Actually running this found a real, previously invisible kiln bug: its process-exit path used syscall 60 (`exit`, kills only the calling thread) instead of 231 (`exit_group`, kills the whole process) -- harmless for kiln's entire history until `ash_gpu_matrix_multiply`'s OpenMP worker threads gave a kiln-generated program a second thread for the first time ever. The program printed the right answer and then hung forever. Fixed in all three exit sites; confirmed behavior-identical for every existing (single-threaded) program via `run_regression.sh`.84
+85
+### Real GPU results, now that one exists (2026-10-08)86
+Same-moment, forced-CPU-vs-real-GPU comparisons on the RX 9070 XT (via WSL2/Dozen -- itself a translation layer, "not a conformant Vulkan implementation, testing use only" per Mesa's own warning):87
+88
+| Kernel | Shape | CPU | GPU | Result |89
+|---|---|---|---|---|90
+| `matrix_multiply` (2000x2000) | memory-bound, naive untiled shader, fp64, host-visible memory | 0.44s | 1.07-1.11s | **GPU loses, ~2.5x slower** |91
+| `chaos_iterate` (new -- 2M elements x 4000 iterations, 8 billion ops) | compute-bound, trivial data movement, no cross-element dependency | ~14.5s | ~0.43s | **GPU wins, ~34x faster** |92
+93
+`chaos_iterate` (`ash_gpu.h`, `src/C/backends/cpu/chaos.c`, `src/C/backends/vulkan/chaos.cpp` + `chaos.cpp`, `src/shaders/chaos.comp`, benchmarked by `pyre/benchmarks/chaos/run_benchmark.c`) exists specifically to answer "is there ANY workload shape this GPU setup actually wins on" after `matrix_multiply`'s real measurement came back a loss -- same recurrence (`x = sqrt(abs(x)+1)*0.99999+0.00001`, chosen because GLSL's `GL_ARB_gpu_shader_fp64` extension gives `double` overloads for `sqrt`/`abs`/arithmetic but NOT `sin`/`cos`, discovered by trying it) on both backends, results matching exactly (`max |CPU - GPU| = 0`) across all 2M elements.94
+95
+The two results together are the concrete evidence for the next task, not a contradiction: `matrix_multiply` is memory/transfer-bound and specifically hurt by fp64-on-consumer-GPU plus an untiled shader; `chaos_iterate` is compute-bound enough that none of that matters. **Neither dispatcher (`src/C/runtime/device.c`) currently tells the difference** -- both unconditionally prefer GPU-when-available with no cost model at all.96
+97
+### Smart dispatch: not started, this is the real next task98
+Both of today's dispatchers (`ash_gpu_matrix_multiply`, `ash_gpu_chaos_iterate`) pick GPU whenever `ash_gpu_available()`, full stop -- which the table above proves is actively wrong for one of the two kernels that exist. The real fix is an algorithm that decides CPU vs. GPU per call based on the workload's actual shape (arithmetic intensity, data size, element count x iterations, whatever the right features turn out to be) rather than "a GPU exists." ideas/gpu_acceleration.md already names this (`ash-gpu benchmark`, "Benchmark everything," design principle 7) but nothing implements it yet. Explicitly deferred, not forgotten -- the user asked for this to wait until the rest is wrapped up; `matrix_multiply` and `chaos_iterate` are exactly the two real, measured data points whoever builds it should calibrate against first.99
+100
+### Tooling101
+`pyre/tools/install_deps.sh` installs the Vulkan SDK pieces the Vulkan backend needs to build (headers, loader, `glslc`/`glslangValidator`, spirv-tools) -- Arch/pacman only so far, matches both this sandbox and the user's own machine. It installs the SDK, not a GPU driver; `vulkaninfo` still correctly reports no device on a machine with none.\ No newline at end of filepyre/tools/install_deps.shadded
+40−0@@ -0,0 +1,40 @@1
+#!/usr/bin/env bash2
+# Installs everything Pyre's GPU backend needs to build (not necessarily3
+# run -- that also needs an actual GPU + driver, which this script4
+# can't provide). Vulkan headers/loader, the SPIR-V shader compiler5
+# toolchain (glslc, glslangValidator), spirv-tools, and vulkan-tools6
+# (vulkaninfo, for checking what the machine actually has once a real7
+# driver is installed).8
+#9
+# Arch Linux (pacman) only for now -- this is what both the sandbox this10
+# was developed in and the user's own machine run. Add another package11
+# manager's branch here if Pyre ever needs to build somewhere else;12
+# don't guess package names for a distro nobody has actually run this13
+# against yet.14
+set -euo pipefail15
+16
+if ! command -v pacman >/dev/null 2>&1; then17
+ echo "install_deps.sh: this script only knows pacman (Arch Linux)." >&218
+ echo "Install these packages' equivalents by hand on your distro:" >&219
+ echo " vulkan-headers vulkan-icd-loader vulkan-tools shaderc glslang spirv-tools" >&220
+ exit 121
+fi22
+23
+PACMAN_PKGS=(vulkan-headers vulkan-icd-loader vulkan-tools shaderc glslang spirv-tools)24
+25
+echo "install_deps.sh: pacman -Sy then installing: ${PACMAN_PKGS[*]}"26
+pacman -Sy --noconfirm27
+pacman -S --noconfirm --needed "${PACMAN_PKGS[@]}"28
+29
+echo30
+echo "install_deps.sh: done. Verifying:"31
+command -v glslc >/dev/null 2>&1 && echo " glslc: $(glslc --version | head -1)" || echo " glslc: MISSING"32
+[ -f /usr/include/vulkan/vulkan.h ] && echo " vulkan headers: present" || echo " vulkan headers: MISSING"33
+command -v vulkaninfo >/dev/null 2>&1 && echo " vulkaninfo: present" || echo " vulkaninfo: MISSING"34
+echo35
+echo "This installs the SDK/loader, not a GPU driver. 'vulkaninfo --summary'"36
+echo "still needs a real vendor driver (vulkan-radeon for an AMD card like"37
+echo "the RX 9070 XT, vulkan-intel, nvidia-utils, ...) installed separately"38
+echo "-- 'Found no drivers!' from vulkaninfo on a machine with no GPU (like"39
+echo "the sandbox this was written in) is the correct, honest answer, not"40
+echo "a failure of this script."