Runtime Particle Physics Kernel Compilation via FunGT IR

[juanchuletas]

Goal

Enable users to define custom particle force laws through the FunGT GUI using a
restricted domain-specific IR, compiled at runtime through an MLIR pipeline to
SPIR-V and loaded via SYCL's kernel_compiler extension, eliminating the need to
recompile FunGT for each behavior modification while guaranteeing GPU memory safety.

Background

Currently, particle behaviors are defined at compile-time via template
instantiation. For interactive experimentation and live demos, users need the
ability to modify particle physics in real-time through the GUI.

The naive approach of compiling raw SYCL C++ strings at runtime via SYCL-RTC
was evaluated and rejected as the primary path: LLM-generated SYCL can produce
GPU hangs, out-of-bounds memory access, and undefined behavior that crashes the
driver. Instead, FunGT defines a restricted IR where these failure modes are
impossible at the language level. The MLIR pipeline compiles this IR to SPIR-V.
SYCL-RTC loads the SPIR-V binary. The GPU never executes unvalidated code.

FunGT IR

The language the user or LLM generates. Four constructs exist. Nothing else:

force.attract:
    r = distance(self.pos, other.pos)
    f = G * self.mass * other.mass / (r * r)
    guard: r > 5.0

The vocabulary is intentionally Turing-incomplete. No pointers. No loops. No
memory allocation. No function calls outside the provided primitives. An LLM
cannot generate a GPU hang in this language because the constructs required to
produce one do not exist.

Technical Approach

User FunGT IR text
      ↓
  fungt-opt  (MLIR pipeline)
      ↓  fungt dialect → arith/math → gpu dialect → spirv dialect
  SPIR-V binary
      ↓
  sycl::create_kernel_bundle_from_source()
      ↓
  iGPU executes kernel

Requirements:

• fungt-dialect build (separate tool, links against LLVM/MLIR)
• DPC++ compiler available for SYCL runtime loading
• Device must support ext_oneapi_can_compile(source_language::spirv)
• Intel GPUs via Level Zero/OpenCL, CPUs via OpenCL

Workflow:

1. User types force law in GUI text field in FunGT IR syntax
2. FunGT calls fungt-opt as subprocess with IR string as input
3. fungt-opt runs lowering pipeline: FunGT ops → arith/math → gpu → spirv
4. fungt-opt serializes SPIR-V binary to temp file or stdout
5. FunGT reads SPIR-V bytes
6. SYCL loads binary via create_kernel_bundle_from_source
7. Kernel executes on particle data
8. Particles respond to new force law immediately

MLIR Dialect

Four ops define the entire language:

fungt.distance — Euclidean distance between two 3D points. Returns f32.

fungt.guard — Takes an i1 condition. If false, force contribution is zero.
Lowers to scf.if wrapping the remainder of the force body.

fungt.scalar_mul — Multiplies two f32 scalars. Returns f32.

fungt.force — Container op wrapping a force kernel body. Lowering pass
wraps this in a gpu.launch with the particle loop scaffold. The loop indexing,
memory coalescing, and work-group sizing are written once in the lowering pass
and are not visible to the user or the LLM.

Implementation Plan

Phase 1: FunGT IR Compiler Pipeline

• Bootstrap fungt-dialect from MLIR standalone example
• Define fungt.distance, fungt.guard, fungt.scalar_mul in TableGen
• Define fungt.force op with region body
• Write lowering pass: FunGT ops → arith and math dialect ops
• Write lowering pass: → gpu.launch with particle loop scaffold
• Write lowering pass: gpu dialect → spirv dialect
• SPIR-V binary serialization via mlir-translate
• Verify SPIR-V loads and executes correctly in SYCL on Intel iGPU
• Verify on NVIDIA via SYCL-CUDA backend

Phase 2: GUI Integration

• Add multi-line text editor widget for FunGT IR input
• Add Compile and Apply button
• Display fungt-opt output and errors in GUI console
• Show compilation status indicator (compiling, success, failed)
• Stream fungt-opt stderr to GUI in real time during compilation

Phase 3: Kernel Management

• Implement kernel caching (map: source hash → compiled SPIR-V binary)
• Handle compilation failures gracefully without crashing or hanging GPU
• Auto re-prompt LLM with compiler error appended on failure
• Add hot-reload: detect IR changes and recompile automatically

Phase 4: LLM Integration

• LLM API call takes natural language physics description
• LLM generates FunGT IR (not SYCL, not C++)
• Validated FunGT IR fed into fungt-opt pipeline
• Compilation errors fed back to LLM for self-correction
• Cache successful kernels keyed by natural language prompt hash

Phase 5: User Experience

• Provide example force snippets in FunGT IR (gravity, orbital, repulsion)
• Show performance metrics (fungt-opt compile time, kernel execution time)
• Comparison mode: show FunGT IR alongside generated SPIR-V op count

Backends

Both GPU backends are supported through the same SPIR-V binary:

Intel iGPU (Arc/Iris Xe): SPIR-V loaded via Level Zero or OpenCL backend.
Primary demo target. No discrete GPU required.

NVIDIA: SPIR-V loaded via SYCL-CUDA backend. Same binary, same pipeline.
Backend switch exposed in GUI for live demo.

Safety Guarantees

The FunGT IR type system eliminates an entire class of GPU errors at the
language level:

• No pointer arithmetic — memory access patterns are fixed in the lowering pass
• No unbounded loops — the particle loop scaffold is in the lowering pass, not
  the user IR
• No work-group synchronization — barriers are not expressible in FunGT IR
• No raw memory allocation — fungt.force operates on scalar arguments only
• guard is the only conditional — no arbitrary branching

A kernel generated from valid FunGT IR cannot produce a GPU hang, a
segmentation fault, or an out-of-bounds memory access. Invalidity is caught at
fungt-opt parse time, not at GPU driver level.

Open Questions

1. fungt-opt as subprocess vs linked MLIR libraries directly inside FunGT
2. SPIR-V temp file on disk vs pipe to FunGT process stdin
3. Minimum MLIR op set needed before Phase 1 is demo-ready
4. How to handle fungt-opt not found at runtime (bundled vs system install)
5. Persistent SPIR-V cache location and invalidation strategy