logicaffeine_compile/codegen_cuda/

runtime.rs

//! Fixed templates for the emitted CUDA program. Unlike the Tenstorrent tier — a
//! baby RISC-V core that runs the whole glibc-backed C body — a GPU thread has no
//! `realloc`/`snprintf`/`strdup`. So the CUDA tier does NOT reuse the C runtime:
//! it lifts only the program's own (integer/bool) code onto the device and supplies
//! a tiny device preamble for the one thing that code touches — `Show`. Everything
//! that needs the dynamic runtime is refused upstream (see [`super::codegen_program_cuda`]).

/// Prepended to the tree-shaken, device-lifted runtime slice. Supplies only the standard headers the
/// lifted runtime needs (device `printf`, fixed-width ints, `PRId64`, `memcpy`, `isnan`/`isinf`,
/// `malloc`). The `Show` primitives (`show_i64`/`show_bool`/`show_f64`/…) are now LIFTED from the C
/// runtime itself (tree-shaken to only what the program uses), so output stays byte-identical to the
/// VM and there is no hand-maintained duplicate to drift.
pub const CUDA_DEVICE_PREAMBLE: &str = r#"// Auto-generated by LOGOS codegen_cuda — do not edit.
// A whole LOGOS program lowered to ONE CUDA kernel (launched <<<1,1>>>). Only the runtime slice the
// program references is lifted here as __device__ functions (tree-shaken from the C oracle), plus the
// program's own code; `Show` rides CUDA's native device printf, so output is byte-identical to the VM.
#include <cstdio>
#include <cstdint>
#include <cinttypes>
#include <cstring>
#include <cstdlib>
#include <cmath>

// --- lifted LOGOS runtime slice (device) ----------------------------------------
"#;

/// A device allocator shim, prepended after the includes. The C runtime grows every Seq/Map/Set with
/// `realloc`, which device code lacks (and which needs the OLD size). So every lifted allocation is
/// routed (`malloc`/`free`/`calloc`/`realloc`/`strdup` → `logos_*`) through a header-carrying malloc:
/// each block stores its byte size 16 bytes before the returned pointer, so `logos_realloc` recovers it.
/// All routed allocations share one header convention, so free/realloc are always consistent. The real
/// `malloc`/`free` are used only INSIDE these shims (the lifted code never calls them directly).
pub const CUDA_DEVICE_ALLOC_SHIM: &str = r#"// --- device allocator shim (header-carrying malloc so realloc works) ------------
__device__ static void* logos_malloc(size_t n){ size_t* p=(size_t*)malloc(n+16); if(!p) return (void*)0; p[0]=n; return (void*)((char*)p+16); }
__device__ static void logos_free(void* q){ if(q) free((void*)((char*)q-16)); }
__device__ static void* logos_realloc(void* q, size_t n){ if(!q) return logos_malloc(n); size_t old=((size_t*)((char*)q-16))[0]; void* r=logos_malloc(n); if(r){ size_t m=old<n?old:n; memcpy(r,q,m); logos_free(q);} return r; }
__device__ static void* logos_calloc(size_t c, size_t s){ size_t n=c*s; void* r=logos_malloc(n); if(r) memset(r,0,n); return r; }
__device__ static char* logos_strdup(const char* s){ size_t n=0; while(s[n]) n++; char* r=(char*)logos_malloc(n+1); if(r){ memcpy(r,s,n); r[n]=0; } return r; }
// --- lifted LOGOS runtime slice (device) ----------------------------------------
"#;

/// The checked-overflow builtins the big-integer fast path uses (`__builtin_{add,sub,mul}_overflow`),
/// which nvcc treats as host-only. Reimplemented exactly via `__int128` (device-supported): compute in
/// 128 bits, store the low 64, and report whether the value did not fit. Prepended so the lifted bignum
/// lifts verbatim.
pub const CUDA_DEVICE_INTRIN_SHIM: &str = r#"// --- device checked-overflow builtins (via __int128) ----------------------------
__device__ static int logos_add_overflow(int64_t a, int64_t b, int64_t* r){ __int128 s=(__int128)a+(__int128)b; *r=(int64_t)s; return s != (__int128)(int64_t)s; }
__device__ static int logos_sub_overflow(int64_t a, int64_t b, int64_t* r){ __int128 s=(__int128)a-(__int128)b; *r=(int64_t)s; return s != (__int128)(int64_t)s; }
__device__ static int logos_mul_overflow(int64_t a, int64_t b, int64_t* r){ __int128 s=(__int128)a*(__int128)b; *r=(int64_t)s; return s != (__int128)(int64_t)s; }
"#;

/// `floor`/`ceil`/`round` of an INTEGER argument. The C tier emits `(int64_t)round(x)` whatever `x`'s type — C
/// converts an integer to `double` through the prototype, so `round(1000000LL)` is `1000000.0`. CUDA C++ resolves
/// the same call to the host-only integral template (`constexpr __host__`), which device code cannot call. These
/// non-template exact matches win overload resolution and keep the C meaning: an integer rounds to itself.
pub const CUDA_DEVICE_INTEGRAL_ROUNDING: &str = r#"// --- floor/ceil/round of an integer argument (C's double conversion) ------------
__host__ __device__ static inline double floor(long long x){ return (double)x; }
__host__ __device__ static inline double floor(long x){ return (double)x; }
__host__ __device__ static inline double floor(int x){ return (double)x; }
__host__ __device__ static inline double ceil(long long x){ return (double)x; }
__host__ __device__ static inline double ceil(long x){ return (double)x; }
__host__ __device__ static inline double ceil(int x){ return (double)x; }
__host__ __device__ static inline double round(long long x){ return (double)x; }
__host__ __device__ static inline double round(long x){ return (double)x; }
__host__ __device__ static inline double round(int x){ return (double)x; }
"#;

/// The device RUNTIME-ERROR PATH. A fail site in the C runtime is `fprintf(stderr, msg); exit(1);` (a
/// non-exhaustive `Inspect`, a div-by-zero guard, a bignum parse error, an explicit `FailWith`, …). On a
/// GPU thread none of `fprintf`/`stderr`/`exit`/`abort` exist — but the stderr diagnostic is NOT part of
/// the ≡VM comparison (stdout + exit status are), so the fail path lowers to: DROP the stderr message
/// (`logos_fprintf` is a no-op) and turn `exit(nonzero)`/`abort()` into a device `trap` — the kernel
/// faults, `cudaDeviceSynchronize` returns non-success, and the host orchestrator exits nonzero, exactly
/// mirroring the VM's `error.is_some()`. These names are routed into the LIFTED `__device__` runtime only
/// (`fprintf`→`logos_fprintf`, `stderr`→`logos_stderr_sink`, `exit`→`logos_exit`, `abort`→`logos_abort`,
/// `fflush`→`logos_fflush`, `stdout`→`logos_stdout_sink`), so the HOST orchestrator keeps real libc. For a program whose fail path is UNREACHABLE (e.g. an
/// `Inspect` whose match is exhaustive at runtime) the trap is dead code that merely has to compile.
pub const CUDA_DEVICE_FAULT_SHIM: &str = r#"// --- device runtime-error path (stderr elided; nonzero exit / abort -> trap) ----
__device__ static int logos_stderr_sink = 0;
__device__ static int logos_stdout_sink = 0;
__device__ static int logos_fprintf(int sink, const char *fmt, ...) { (void)sink; (void)fmt; return 0; }
__device__ static int logos_fflush(int sink) { (void)sink; return 0; }
__device__ static void logos_exit(int code) { if (code != 0) { asm("trap;"); } }
__device__ static void logos_abort(void) { asm("trap;"); }
"#;

/// The device CLOCK globals for `today`/`now`. The C runtime reads the injected fixed clock via
/// `getenv("LOGOS_FIXED_DAYS"/"LOGOS_FIXED_NANOS")` + `time()` — neither exists on a GPU thread. Instead
/// the HOST orchestrator (a CPU) reads that same env (or `time()`) and `cudaMemcpyToSymbol`s it into these
/// globals BEFORE the launch (see [`CUDA_HOST_TEMPLATE`]); the device `logos_today`/`logos_now` overrides
/// (see [`super::device_lift::device_override`]) simply return them. Non-`static` so the host symbol
/// reference resolves. Always emitted (0 for a non-temporal program, which never reads them).
pub const CUDA_DEVICE_CLOCK: &str = r#"// --- device clock (host injects the fixed instant before launch) ----------------
__device__ long long logos_dev_fixed_days = 0;
__device__ long long logos_dev_fixed_nanos = 0;
"#;

/// The CUDA tier's prologue to the stepped scheduler ([`crate::codegen_c::runtime::STEPPED_SCHEDULER`]), under
/// which every concurrent program runs on the GPU as the host C tier schedules it — the same sources, placed
/// where the host C carries them: device linkage, the device allocator shim, no failure escape (a runtime
/// helper's failure traps the run; a failing statement fails its task through the core), the run's seed from
/// a device global the host fills from `LOGOS_SEED` before the launch, and a failed run's message in a device
/// buffer the host prints after it, with exit status 1.
pub const CUDA_DEVICE_STEPPED_PROLOGUE: &str = r#"// --- the stepped scheduler (device) ---------------------------------------------
#define LOGOS_FN __device__ static
#define LOGOS_VAR __device__ static
#define LOGOS_API __device__ static
#define logos_alloc(n) logos_calloc(1, (n))
#define logos_release(p) logos_free(p)
typedef int (*logos_step_fn)(void *);
#define LOGOS_TASK_TIER logos_step_fn step; void *frame;
#define LOGOS_STEP_NO_ESCAPE
__device__ unsigned long long logos_dev_seed = 0;
#define LOGOS_RUN_SEED logos_dev_seed
__device__ int logos_dev_failed = 0;
__device__ char logos_dev_failure[256];
__device__ static void logos_dev_fail(const char *message) {
    int i = 0;
    for (; message[i] && i < 255; i++) logos_dev_failure[i] = message[i];
    logos_dev_failure[i] = 0;
    logos_dev_failed = 1;
}
#define LOGOS_RUN_FAILS(message) logos_dev_fail(message)
"#;

/// The host's side of a concurrent program's run, each spliced at its line of [`CUDA_HOST_TEMPLATE`] (and the
/// line dropped for a program with no tasks): the seed reader, the seed copied into the device before the
/// launch, and a failed run reported after it.
pub const CUDA_HOST_SEED_READER: &str = include_str!("../codegen_c/host_seed.c");
pub const CUDA_HOST_RUN_SEED: &str =
    "    { unsigned long long seed = logos_env_seed(0); cudaMemcpyToSymbol(logos_dev_seed, &seed, sizeof seed); }\n";
pub const CUDA_HOST_RUN_FAILED: &str = "    { int failed = 0; cudaMemcpyFromSymbol(&failed, logos_dev_failed, sizeof failed);\n      \
     if (failed) { char message[256]; cudaMemcpyFromSymbol(message, logos_dev_failure, sizeof message); \
     fprintf(stderr, \"%s\\n\", message); return 1; } }\n";

/// The `__global__` entry that runs the whole program on one thread, plus the host
/// orchestrator: launch, sync (flushes device printf), and map any device fault to a
/// nonzero exit so error behaviour matches the VM. Substituted after the lifted code.
pub const CUDA_HOST_TEMPLATE: &str = r#"
__global__ void logos_kernel() { (void) logos_main(0, (char **)0); }

// --- host orchestrator ----------------------------------------------------------
#include <cstdlib>
#include <ctime>
#include <cstring>
/*LOGOS_HOST_SEED_READER*/
int main(int argc, char **argv) {
    (void)argc; (void)argv;
    // The device runtime uses recursion + a fixed-limb bignum on the stack (the shortest-float
    // formatter) and device malloc (string/heap builders); raise both limits above their small
    // defaults so a whole program runs on one thread.
    cudaDeviceSetLimit(cudaLimitStackSize, (size_t)64 * 1024);
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, (size_t)256 * 1024 * 1024);
    // Inject the deterministic fixed clock the device `today`/`now` read, mirroring codegen_c's
    // getenv(LOGOS_FIXED_*)-or-time() shim on the HOST, then copy it into the device globals. Always set;
    // a non-temporal program simply never reads them.
    { const char *e; long long v;
      e = getenv("LOGOS_FIXED_DAYS");  v = e ? atoll(e) : (long long)(time(NULL) / 86400);
      cudaMemcpyToSymbol(logos_dev_fixed_days, &v, sizeof(v));
      e = getenv("LOGOS_FIXED_NANOS"); v = e ? atoll(e) : (long long)time(NULL) * 1000000000LL;
      cudaMemcpyToSymbol(logos_dev_fixed_nanos, &v, sizeof(v)); }
/*LOGOS_ARGV_MARSHAL*/
/*LOGOS_HOST_RUN_SEED*/
    logos_kernel<<<1, 1>>>();
    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) { fprintf(stderr, "kernel launch: %s\n", cudaGetErrorString(err)); return 1; }
    err = cudaDeviceSynchronize();
    if (err != cudaSuccess) { fprintf(stderr, "kernel exec: %s\n", cudaGetErrorString(err)); return 1; }
/*LOGOS_HOST_RUN_FAILED*/
    return 0;
}
"#;

/// The single canonical name of the emitted CUDA source (referenced by the Makefile).
pub const KERNEL_FILE_NAME: &str = "logos_cuda.cu";

/// Makefile glue: builds the single-file CUDA program with `nvcc` for the Ampere
/// (`sm_86`, the RTX 3060 class) target. `{KERNEL_FILE}` is substituted with the
/// emitted source name.
pub const CUDA_MAKEFILE_TEMPLATE: &str = r#"# Auto-generated by LOGOS codegen_cuda.
# Builds the single-file CUDA program with nvcc. Override ARCH for other GPUs.
ARCH ?= sm_86
NVCC ?= nvcc
logos_cuda: {KERNEL_FILE}
	$(NVCC) -O2 -arch=$(ARCH) -fmad=false $< -o $@
.PHONY: clean
clean:
	rm -f logos_cuda
"#;