|
| 1 | +--- |
| 2 | +trigger: always_on |
| 3 | +--- |
| 4 | + |
| 5 | +> **WGSL EXPERT RULE:** Before writing any shader code string, you MUST cross-reference `WGSL_REFERENCE.md` to ensure you are not using HLSL keywords (like `lerp` or `float3`) and are respecting strict type casting. |
| 6 | +
|
| 7 | +# Project Context: SpawnDev.ILGPU.WebGPU |
| 8 | +You are working on a C# library that functions as a WebGPU backend for ILGPU, running strictly inside a Blazor WebAssembly (WASM) environment. SpawnDev.ILGPU.WebGPU uses a robust and efficient WGSL code generator. |
| 9 | + |
| 10 | +**Core Stack:** |
| 11 | +- **Runtime:** Blazor WebAssembly (.NET 10). |
| 12 | +- **Interop:** `SpawnDev.BlazorJS` (Strict requirement for all JS interaction). |
| 13 | +- **Target API:** WebGPU API (accessed via `SpawnDev.BlazorJS` wrappers). |
| 14 | +- **Abstractions:** implementing `ILGPU.Runtime` interfaces (Backend, Accelerator, Buffer, etc.). |
| 15 | + |
| 16 | +# Testing |
| 17 | +- **Procedure:** Run the demo app (`SpawnDev.ILGPU.WebGPU.Demo`) in a browser and navigate to the "/tests" page. |
| 18 | +- **Do NOT use `PlaywrightTestRunner`**: It is currently reserved for CI/CD environments. |
| 19 | +- **Port Handling:** Run the demo app on a port DIFFERENT from the one in `launchSettings.json` (e.g., use 5002 if 5181 is taken) to avoid conflicts with Visual Studio. |
| 20 | +- **Process Cleanup:** `dotnet` sessions frequently hang. BEFORE starting a new test run, execute `taskkill /F /IM dotnet.exe` to ensure a clean slate. |
| 21 | +- **Verification:** Manually trigger tests on the "/tests" page and inspect results. |
| 22 | +- **Verification Safety:** ALWAYS run `dotnet build` to confirm the project compiles BEFORE attempting to run the demo app in the browser. Failed builds prevent the browser from connecting. |
| 23 | + |
| 24 | +## 🚨 CRITICAL HARD CONSTRAINTS (Violations cause deadlocks) |
| 25 | + |
| 26 | +1. **NO BLOCKING WAITS:** - NEVER use `Task.Result`, `.Wait()`, or `.GetAwaiter().GetResult()` on the main thread. |
| 27 | + - Blazor WASM on the main thread is single-threaded. Blocking the thread prevents the JS Promise callback from running, causing an instant deadlock. |
| 28 | + - **CORRECTION:** Always propagate `async/await` up the stack. If an ILGPU interface requires a synchronous return but the implementation depends on a JS Promise, you must flag this as an architectural conflict or use `SpawnDev.BlazorJS` synchronous interop methods *only if* the underlying JS API is synchronous (WebGPU is mostly async). |
| 29 | + |
| 30 | +2. **NO THREAD ASSUMPTIONS:** |
| 31 | + - Do not use `System.Threading.Thread`, `Thread.Sleep`, or assumes strictly parallel CPU execution unless explicitly working inside a `SpawnDev.BlazorJS.WebWorker`. |
| 32 | + |
| 33 | +## 🛠️ Code Style & Patterns |
| 34 | + |
| 35 | +### 1. Interop Pattern (SpawnDev.BlazorJS) |
| 36 | +Always use `SpawnDev.BlazorJS` to interact with WebGPU. Do not use standard `IJSRuntime` unless necessary. |
| 37 | +- **wrappers:** Use or create strong-typed wrappers inheriting from `JSObject` for WebGPU objects using implementations in SpawnDev.BlazorJS.JSObjects if they exist (e.g., `GPUAdapter`, `GPUDevice`, `GPUBuffer`). |
| 38 | +- **usage:** |
| 39 | + ```csharp |
| 40 | + using var navigator = BlazorJSRuntime.JS.Get<Navigator>("navigator"); |
| 41 | + using var gpu = navigator.Gpu; |
| 42 | + var adapter = await gpu.RequestAdapter(); // Async is mandatory |
| 43 | + ``` |
| 44 | + |
| 45 | +### 2. WebGPU Backend Implementation Rules |
| 46 | + |
| 47 | +#### Kernel Argument Binding |
| 48 | +- **Implicit Index:** Implicitly grouped kernels have an implicit index parameter at index 0. The backend `GenerateKernelLauncherMethod` must explicitly skip this parameter when defining the launcher signature and loading arguments (offset = 1). |
| 49 | +- **Scalar Marshaling:** WebGPU requires all buffer bindings to be `storage` or `uniform` buffers. Scalar kernel arguments (int, float) must be marshaled into **1-element storage buffers**. In WGSL, access them as `paramN[0]`. |
| 50 | +- **Buffer Flags:** All buffers passed to compute shaders **MUST** have the `GPUBufferUsage.Storage` flag. |
| 51 | + |
| 52 | +#### WGSL Generation (ArrayView Mapping) |
| 53 | +When mapping `ArrayView<T>` or `ArrayViewN<T>` in `WGSLKernelFunctionGenerator`: |
| 54 | +- **Field 0 (Ptr):** Map to the buffer reference (`¶mN`). |
| 55 | +- **Field 1 (Index/Offset):** Map to constant `0` (or `0u`). The buffer binding offset is handled by the WebGPU API `SetBindGroup` call, so the WGSL shader always sees a base-0 view. |
| 56 | +- **Field 2/3 (Length):** Map to `bitcast<i32>(arrayLength(¶mN))`. |
| 57 | + |
| 58 | +#### Compilation & Types |
| 59 | +- **Type Safety:** Use fully qualified type names for ILGPU IR types (e.g., `global::ILGPU.IR.Values.Parameter`) to avoid conflicts with reflection types. |
| 60 | +- **Value Resolution:** When handling `ValueReference` or `GetField` sources, ensure you check types correctly. Avoid relying on `.Resolve()` extension references that might be ambiguous; use direct type pattern matching. |
| 61 | + |
| 62 | +### 3. Current Debugging Context (Status as of 2026-02-04) |
| 63 | +- **Status:** All supported tests are working correctly. |
| 64 | +- **Transpiler Limitations (Throw Instruction Support):** |
| 65 | + - **CRITICAL:** The IL to WGSL transpiler **DOES NOT SUPPORT** the `Throw` instruction. |
| 66 | + - **Consequence:** If `throw` is found in the IL (e.g., explicit `throw` or implicit argument validation), the transpiler will fail and throw a compilation exception. |
| 67 | + - **Problematic Methods:** Many System.Math methods (e.g., `Math.Clamp`, `Math.Round`, `Math.Truncate`, `Math.Sign`) contain implicit `throw` checks for argument validation. |
| 68 | + - **Workarounds:** |
| 69 | + - **Math.Clamp:** Do NOT use directly. Use `Math.Min(Math.Max(val, min), max)`. |
| 70 | + - **Round/Truncate/Sign:** Avoid using these in kernels until a fix is implemented (upstream or custom Intrinsics). |
| 71 | + - **General:** Avoid any helper methods that might throw exceptions. |
| 72 | + - **Supported Intrinsics:** `Atan2`, `FusedMultiplyAdd`, `Rem`, `Min`, `Max`, `Abs`, `Pow`, `Log`, `Exp`. |
| 73 | + |
| 74 | +## 📚 Project Resources |
| 75 | +- **ILGPU Source & Docs:** Available in this workspace at `d:\users\tj\Projects\SpawnDev.ILGPU\ILGPU`. Refer to this for understanding ILGPU internals. |
| 76 | +- **Examples:** Check `SpawnDev.ILGPU.WebGPU.Demo` for working WebGPU examples. |
| 77 | +- **Resolved Issues:** |
| 78 | + - `System.NotSupportedException` (Throw) fixed via workarounds. |
| 79 | + - "Expected X, Got 0" issues in basic kernel tests appear resolved or superseded by current test suite success. |
| 80 | + |
| 81 | +#### WGSL Translation Reference (The "Rosetta Stone") |
| 82 | +*Use this mapping table as the ground truth for generating WGSL. Do not infer mappings from standard HLSL/GLSL patterns.* |
| 83 | + |
| 84 | +| C# ILGPU Concept | WGSL Implementation | Critical Notes | |
| 85 | +| :--- | :--- | :--- | |
| 86 | +| **`Index1D`** | `i32` (Cast from `u32`) | **NOT** a struct. Map from `GlobalInvocationId.x`. | |
| 87 | +| **`Index2D`** | `vec2<i32>` | **NOT** a struct. Access via `.x`, `.y`. | |
| 88 | +| **`ArrayView<T>`** | `var<storage, read_write>` | Must be declared in `@group(0)`. | |
| 89 | +| **`ArrayView2D<T>`** | **DECOMPOSE:** `arg_data`, `arg_stride` | **DO NOT** use structs. Split into **2 separate arguments**: <br>1. `var<storage> data` <br>2. `array<i32,1>` (Width/Stride). | |
| 90 | +| **Scalar Arg (`int`, `float`)** | `array<type, 1>` | **CRITICAL:** Wrapped in 1-element array. Access via `[0]`. | |
| 91 | +| **`Group.Barrier()`** | `workgroupBarrier()` | | |
| 92 | +| **`SharedMemory.Allocate`** | `var<workgroup>` | Must be declared at **module scope**, not inside `main`. | |
| 93 | +| **Launcher Logic** | `BindGroup` Creation | **CRITICAL:** For `ArrayView2D/3D`, stop field recursion. Create **2 bindings** (1: Buffer, 2: Scalar Stride). | |
| 94 | + |
| 95 | +**Translation Example (Scalar Injection):** |
| 96 | +*Input (C#):* `public static void Kernel(Index1D index, int val)` |
| 97 | +*Output (WGSL):* |
| 98 | +```wgsl |
| 99 | +@group(0) @binding(0) var<storage, read_write> val_buf : array<i32, 1>; // Scalar wrapper |
| 100 | +@compute @workgroup_size(64) |
| 101 | +fn main(@builtin(global_invocation_id) global_id : vec3<u32>) { |
| 102 | + let index = i32(global_id.x); |
| 103 | + let val = val_buf[0]; // Accessing scalar via index 0 |
| 104 | + ... |
| 105 | +} |
| 106 | +``` |
| 107 | + |
| 108 | +### Part 3: Agent Workflow & Debug Context |
| 109 | +*Paste this at the end to finish the file.* |
| 110 | + |
| 111 | +```markdown |
| 112 | +#### Agent Workflow for Transpiler Logic |
| 113 | +When asked to fix bugs in `WGSLCodeGenerator` or `WGSLKernelFunctionGenerator`: |
| 114 | +1. **Locate the AST Node:** Identify which ILGPU IR node (e.g., `Load`, `GetField`, `Atomic`) is being mishandled. |
| 115 | +2. **Consult the "Rosetta Stone":** Check the table above for the correct WGSL target. |
| 116 | +3. **Trace the Emit:** Do not just write the WGSL string. You must write the C# `StringBuilder` logic that *emits* that string. |
| 117 | +4. **Verify Scoping:** If adding a variable (especially shared memory), check if the generator is currently emitting inside `main()` or at module scope. Shared memory *must* bubble up to module scope. |
| 118 | +``` |
0 commit comments