diff --git a/.cargo/config.toml b/.cargo/config.toml new file mode 100644 index 00000000..559ff712 --- /dev/null +++ b/.cargo/config.toml @@ -0,0 +1,5 @@ +[build] +rustflags = ["-A", "unused"] + +[env] +MACOSX_DEPLOYMENT_TARGET = "15.0" \ No newline at end of file diff --git a/.claude/settings.local.json b/.claude/settings.local.json index d218751e..fc079cc0 100644 --- a/.claude/settings.local.json +++ b/.claude/settings.local.json @@ -3,7 +3,19 @@ "allow": [ "Bash(dir:*)", "Bash(grep:*)", - "Bash(cargo build:*)" + "Bash(cargo build:*)", + "Bash(ls:*)", + "Bash(python3:*)", + "Bash(cargo check:*)", + "Bash(cargo tree:*)", + "Bash(xargs:*)", + "Bash(cargo doc:*)", + "Bash(curl -s \"https://crates.io/api/v1/crates/winit\")", + "Bash(cargo add *)", + "Bash(sed -i '' 's/winit = \"0.29.1\"/winit = \"0.30\"/' Cargo.toml)", + "Bash(cargo fetch *)", + "Bash(ls ~/.cargo/registry/src/index.crates.io-1949cf8c6b5b557f/ | grep winit)", + "Read(//Users/alex.dixon/.cargo/registry/src/index.crates.io-1949cf8c6b5b557f/**)" ] } } diff --git a/.vscode/launch.json b/.vscode/launch.json index 9af03231..1598bec1 100644 --- a/.vscode/launch.json +++ b/.vscode/launch.json @@ -88,6 +88,50 @@ "console": "externalTerminal", "preLaunchTask": "examples" }, + { + "name": "bindful (macOS|Debug)", + "type": "lldb", + "request": "launch", + "program": "${workspaceFolder}/target/debug/examples/bindful", + "args": [], + "stopAtEntry": false, + // "cwd": "${fileDirname}", + "environment": [], + "console": "externalTerminal", + // "preLaunchTask": "examples" + }, + { + "name": "bindless (macOS|Debug)", + "type": "lldb", + "request": "launch", + "program": "${workspaceFolder}/target/debug/examples/bindless", + "args": [], + "stopAtEntry": false, + // "cwd": "${fileDirname}", + "environment": [], + "console": "externalTerminal", + // "preLaunchTask": "examples" + }, + { + "name": "imgui_demo (macOS|Debug)", + "type": "lldb", + "request": "launch", + "program": "${workspaceFolder}/target/debug/examples/imgui_demo", + "args": [], + "stopAtEntry": false, + "environment": [], + "console": "externalTerminal", + }, + { + "name": "client (macOS|Debug)", + "type": "lldb", + "request": "launch", + "program": "${workspaceFolder}/target/debug/client", + "args": [], + "stopAtEntry": false, + "environment": [], + "console": "externalTerminal", + }, { "name": "imgui_demo (Win32|Debug)", "type": "cppvsdbg", diff --git a/.vscode/settings.json b/.vscode/settings.json index b4e98008..c1ba8593 100644 --- a/.vscode/settings.json +++ b/.vscode/settings.json @@ -63,5 +63,6 @@ "xtree": "cpp", "xutility": "cpp" }, - "rust-analyzer.showUnlinkedFileNotification": false + "rust-analyzer.showUnlinkedFileNotification": false, + "rust-analyzer.checkOnSave": false } \ No newline at end of file diff --git a/Cargo.toml b/Cargo.toml index 0e2ca27b..02859c0a 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -44,13 +44,16 @@ libloading = "0.7.4" bevy_ecs.workspace = true ddsfile = "0.5.1" +[build-dependencies] +htwv = { path = "hotline-data/htwv" } + [dependencies.imgui-sys] version = "0.9.0" features = ["docking"] [target.'cfg(target_os = "macos")'.dependencies] metal = "0.28.0" -winit = "0.29.1" +winit = "0.30" objc = "0.2.4" cocoa = "0.25.0" core-graphics-types = "0.1.3" @@ -97,6 +100,10 @@ default = ["build_data", "client"] name = "bindless" crate-type = ["bin"] +[[example]] +name = "compute" +crate-type = ["bin"] + [[example]] name = "imgui_demo" crate-type = ["bin"] diff --git a/build.rs b/build.rs index 93383c2c..171216ef 100644 --- a/build.rs +++ b/build.rs @@ -1,6 +1,11 @@ +use htwv; + use std::process::Command; +#[cfg(target_os = "windows")] fn main() { + println!("cargo:rerun-if-changed=shaders"); + if std::env::var("CARGO_FEATURE_BUILD_DATA").is_ok() { let pmbuild = "hotline-data\\pmbuild.cmd"; @@ -14,3 +19,30 @@ fn main() { } } } + +#[cfg(target_os = "macos")] +fn main() { + use std::path::Path; + + println!("cargo:rerun-if-changed=shaders"); + + if std::env::var("CARGO_FEATURE_BUILD_DATA").is_ok() { + let output_dir = Path::new("target/data/shaders"); + + let pmbuild = "pmbuild"; + let status = Command::new(pmbuild) + .args(["mac-data"]) + .status() + .unwrap_or_else(|e| panic!("failed to run '{pmbuild}': {e}")); + + if !status.success() { + panic!("pmbuild mac-data failed with status: {status}"); + } + + println!("cargo:warning=Compiling shaders..."); + match htwv::compile_dir("shaders", "target/data/shaders") { + Ok(_) => println!("cargo:warning=Shader compilation succeeded"), + Err(e) => println!("cargo:warning=Shader compilation errors:\n{e}"), + } + } +} \ No newline at end of file diff --git a/client/main.rs b/client/main.rs index 70642fcc..91ed1370 100644 --- a/client/main.rs +++ b/client/main.rs @@ -1,10 +1,27 @@ use hotline_rs::*; use hotline_rs::prelude::*; +#[cfg(target_os = "macos")] +fn platform_dpi_aware() -> bool { + false +} + +#[cfg(not(target_os = "macos"))] +fn platform_dpi_aware() -> bool { + true +} + fn main() -> Result<(), hotline_rs::Error> { + std::panic::set_hook(Box::new(|info| { + let bt = std::backtrace::Backtrace::force_capture(); + eprintln!("PANIC: {info}\n{bt}"); + std::process::abort(); + })); + // create client let ctx : Client = Client::create(HotlineInfo { + dpi_aware: platform_dpi_aware(), ..Default::default() })?; diff --git a/config.jsn b/config.jsn index c23f1587..e94bc24e 100644 --- a/config.jsn +++ b/config.jsn @@ -1,9 +1,14 @@ { // configure build tools tools: { - pmfx: "hotline-data/bin/win32/pmfx/pmfx.exe" + pmfx: "py -3 hotline-data/pmfx-shader/pmfx.py" texturec: "hotline-data/bin/win32/texturec/texturec.exe" - pmfx_dev: "py -3 ../pmfx-shader/pmfx.py" + pmfx_dev: "py -3 hotline-data/pmfx-shader/pmfx.py" + } + + tools: { + pmfx_dev: "python3 hotline-data/pmfx-shader/pmfx.py" + texturec: "hotline-data/bin/macos/texturec" } tools_help: { @@ -15,7 +20,7 @@ } pmfx_dev(pmfx): {} } - + tools_update: { pmfx: { tag_name: latest @@ -121,11 +126,10 @@ // windows specific data, will also build (base) win32-data(base): { - pmfx: { + pmfx_dev: { args: [ "-shader_platform hlsl" "-shader_version 6_5" - "-i ${src_shader_dir}/" "-o ${data_dir}/shaders" "-t ${temp_dir}/shaders" @@ -137,12 +141,19 @@ "-Od" ] } + } + + mac-data(base): { pmfx_dev: { explicit: true args: [ - "-shader_platform hlsl" - "-shader_version 6_5" - "-i ${src_shader_dir}/" + "-shader_platform metal" + "-shader_version 6_0" + "-metal_version 2.3" + "-metal_sdk macosx" + "-discrete_binding 0" + "-discrete_binding 1" + "-i ${src_shader_dir}" "-o ${data_dir}/shaders" "-t ${temp_dir}/shaders" "-num_threads 1" @@ -153,21 +164,6 @@ } } - rt_shaders: { - jsn_vars: { - output_dir: "target/data/shaders" - } - shell: { - commands: [ - "hotline-data\\bin\\win32\\pmfx\\bin\\dxc\\dxc.exe -T lib_6_3 -E MyRaygenShader -Fo ${output_dir}/raygen.cso -I . shaders/raytracing_example.hlsl" - "hotline-data\\bin\\win32\\pmfx\\bin\\dxc\\dxc.exe -T lib_6_3 -E MyClosestHitShader -Fo ${output_dir}/closesthit.cso -I . shaders/raytracing_example.hlsl" - "hotline-data\\bin\\win32\\pmfx\\bin\\dxc\\dxc.exe -T lib_6_3 -E MyMissShader -Fo ${output_dir}/miss.cso -I . shaders/raytracing_example.hlsl" - "hotline-data\\bin\\win32\\pmfx\\bin\\dxc\\dxc.exe -T lib_6_3 -Fo ${output_dir}/lib.cso -I . shaders/raytracing_example.hlsl" - ] - } - } - - // win32 debug client, plugins and data win32-debug(win32-data, hotline): { copy: { diff --git a/config.toml b/config.toml deleted file mode 100644 index 678266e5..00000000 --- a/config.toml +++ /dev/null @@ -1,2 +0,0 @@ -[build] -rustflags = ["-C", "link-arg=-fuse-ld=lld"] \ No newline at end of file diff --git a/docs/binding-architecture.md b/docs/binding-architecture.md new file mode 100644 index 00000000..23a92b98 --- /dev/null +++ b/docs/binding-architecture.md @@ -0,0 +1,250 @@ +# Bindless binding architecture: D3D12 ↔ Metal + +Reference notes on how hotline maps a single bindless shader-resource model onto two very +different GPU binding APIs — D3D12 descriptor heaps and Metal argument buffers — and the two +approaches we went through to make Metal behave like D3D12. Written as background for a blog post. + +--- + +## 1. The goal + +Author shaders **once** in HLSL using bindless resource arrays, and run them on both D3D12 and +Metal with identical CPU-side code. A shader looks like this (`shaders/ecs.hlsl`): + +```hlsl +// all bindless arrays share register t1, separated by space +Texture2D textures[] : register(t1, space7); +TextureCube cubemaps[] : register(t1, space9); +Texture2DArray texture_arrays[] : register(t1, space10); +Texture3D volume_textures[] : register(t1, space11); + +StructuredBuffer draws : register(t0, space0); +StructuredBuffer materials : register(t0, space2); + +SamplerState sampler_wrap_linear : register(s1); +``` + +The CPU stores **global indices** into these arrays (e.g. a material stores the heap slot of its +albedo texture and the scene's IBL cubemap), and the shader does `textures[material.albedo_id]` +or `cubemaps[ibl_id]`. There is one global resource pool; the index is the only thing that crosses +the CPU/GPU boundary. + +--- + +## 2. Two GPU binding models + +### D3D12 — descriptor heaps +A `ID3D12DescriptorHeap` is a flat array of descriptors. The shader-visible CBV/SRV/UAV heap is +bound once; shaders index it directly (`ResourceDescriptorHeap[i]` in SM6.6, or unbounded +descriptor-table ranges pre-6.6). The **heap slot is the global index** — exactly the model above, +natively. A root signature maps each HLSL `register`/`space` to a descriptor-table range or root +constants. + +### Metal — argument buffers +Metal has no shader-visible descriptor heap that you index with an arbitrary integer (pre-Metal 3 +bindless). Instead you build an **argument buffer**: a GPU buffer whose layout is described by an +`MTLArgumentEncoder`, holding texture/buffer/sampler handles at fixed `[[id(N)]]` slots. A shader +receives argument buffers as `[[buffer(N)]]` parameters. An `MTLHeap` backs the actual resources so +they're resident; the argument buffer holds references into it. + +So the porting problem is: **make a flat, integer-indexed global pool work on top of Metal argument +buffers**, where each shader only sees the argument buffers (descriptor sets) it actually uses. + +--- + +## 3. Hotline's abstraction + +- `gfx::PipelineLayout` carries `bindings: Vec`, `push_constants`, and + `static_samplers`. Each `DescriptorBinding` records `shader_register`, `register_space`, + `binding_type` (SRV/UAV/CBV/Sampler) and `num_descriptors` (None = unbounded/bindless). +- `gfx::PipelineSlotInfo { index, count }` is the resolved location for a binding key + `(register, space, descriptor_type)`. +- A single bindless `Heap` holds all textures and all buffers. + +The two backends implement `gfx::Device::create_render_pipeline` / `set_heap` / `set_binding` +differently from here. + +--- + +## 4. D3D12 mapping (the reference) + +`src/gfx/d3d12.rs` builds a root signature from the pipeline layout. Each binding/space becomes a +descriptor-table range or root parameter; the resolved `PipelineSlotInfo.index` is the **heap +index / root slot**, used directly. There is no per-binding offset fix-up — the global index *is* +the descriptor index. (This is why `PipelineSlotInfo` carries no offset field, and why D3D12 needed +no special handling.) + +--- + +## 5. Metal mapping via SPIRV-Cross (htwv) + +Shaders are cross-compiled offline: **HLSL → DXC → SPIR-V → SPIRV-Cross → MSL** in the `htwv` crate +(`hotline-data/htwv/src/macos_impl.rs`). DXC encodes HLSL `register(tN, spaceM)` as SPIR-V +decorations `Binding = N`, `DescriptorSet = M`. htwv then **re-decorates** each resource to assign +the Metal descriptor set (`[[buffer(N)]]`) and the slot within it (`[[id(N)]]`), and feeds matching +`spvc_msl_resource_binding` entries to SPIRV-Cross. + +The runtime (`src/gfx/mtl.rs`) must assign the **same** `[[buffer(N)]]` numbering so that +`set_heap` / `set_binding` bind the heap's argument buffers to the slots the MSL expects. This +mirroring is the crux: `build_slot_lookup`, `build_stage_binders` and `build_compute_binder` in +`mtl.rs` reproduce, byte-for-byte, the grouping htwv used in codegen. + +The `Heap` keeps **two argument buffers** — one array of all textures, one of all buffer pointers +(`get_texture_argument_buffer` / `get_buffer_argument_buffer`) — each a flat array based at id 0, +indexed by global slot. Samplers live in their own argument buffer at fragment `buffer(0)`. Push +constants are *discrete* descriptor sets (no argument buffer) so they can use +`setVertexBytes`/`setFragmentBytes`. + +The open question that produced two approaches: **how do you group HLSL bindings into Metal +descriptor sets?** + +--- + +## 6. Approach 1 — group by `(kind, register)`, compensate with `sub_offset` + +Group bindings by register *kind and number only* (`t0`, `t1`, `u0`, `b0`…). All arrays sharing a +register — regardless of space — landed in **one** descriptor set, packed at consecutive ids: + +``` +descriptor set for t1: + textures -> [[id(0)]] + cubemaps -> [[id(1)]] + texture_arrays -> [[id(2)]] + volume_textures -> [[id(3)]] +``` + +**Why this was attractive:** SPIRV-Cross hard-limits argument buffers to +`kMaxArgumentBuffers = 8` (`spirv_msl.hpp`, throws *"Descriptor set index is out of range."* past +it). Packing many spaces into one register's set conserves that scarce budget — you could have up +to 8 *registers*, each holding many spaces. + +**The compensation:** with SPIRV-Cross's "unsized array hack", `array[i]` in a packed set lowers to +`arg_buffer[id + i]`. Our heap argument buffer is a flat array based at id 0, so `cubemaps[i]` +(at `id(1)`) actually reads `heap[1 + i]`. To cancel the `+1`, the CPU subtracted the binding's +`sub_offset` from the index it wrote: `index = global_index - sub_offset` (the old +`get_lookup` → `get_sub_binding_offset` path). + +**The fatal flaw:** that compensation only ran for **structured buffers routed through +`get_lookup`**. Texture indices live in *shared material/draw data* and are written as raw global +indices — materials are pipeline-agnostic, so you can't bake a pipeline-specific `sub_offset` into +them. Result: the **second and later texture array in a packed set was off by one**. `cubemaps` +(at `id(1)`) read `heap[ibl_id + 1]` and sampled a neighbouring, unrelated texture as a cube — the +classic "cube samples a flat orange, 2D samples a flat sky-blue" symptom. The first array +(`textures` at `id(0)`) worked by luck because its offset was 0. + +So Approach 1 only ever worked for shaders using a single bindless texture array. + +--- + +## 7. Approach 2 — group by `(kind, register, space)`, one set per binding (current) + +Add **space** to the grouping key. Each `(kind, register, space)` tuple becomes its own descriptor +set, so every bindless array is alone in its set at `[[id(0)]]`: + +``` +buffer(3): textures -> [[id(0)]] +buffer(4): cubemaps -> [[id(0)]] +``` + +Now `textures[i]` and `cubemaps[i]` both lower to `arg_buffer[i]` — no offset, no compensation. +The CPU writes raw global indices and they just work. `sub_offset`, `get_sub_binding_offset`, and +the `get_lookup` subtraction were all removed; with one binding per set they are permanently 0. + +**The cost we accepted:** capacity drops from "8 registers × many spaces" to roughly **8 total +binding groups per stage**, because each binding now consumes a whole descriptor set against +`kMaxArgumentBuffers = 8`. The heaviest current shader (`vs_mesh_material_indirect`) sits at +`buffer(7)` — the last legal slot. To make the ceiling visible instead of cryptic, htwv now emits a +`cargo:warning` when a stage reaches/exceeds the limit (`MAX_DESCRIPTOR_SETS`, kept in sync with +`kMaxArgumentBuffers`; bump it if a future SPIRV-Cross raises the cap). + +The two sides stay in lockstep by keying on `(kind, register, space)` in **both** +`hotline-data/htwv/src/macos_impl.rs` (codegen) and the three `mtl.rs` builders (runtime), iterating +`pipeline_layout.bindings` in the same order. + +--- + +## 8. Side-by-side + +| Concept | D3D12 | Metal (Approach 2) | +|----------------------------|----------------------------------------|-----------------------------------------------------| +| Global pool | Shader-visible descriptor heap | Two `MTLHeap`-backed argument buffers (tex / buf) | +| Index semantics | Heap slot = global index (direct) | `arg_buffer[i]`, base id 0 = global index (direct) | +| HLSL `register`/`space` | Root-sig table range / root param | One MSL descriptor set per `(kind, register, space)`| +| Per-binding offset fix-up | None (`PipelineSlotInfo` has no offset)| None (each binding alone at `[[id(0)]]`) | +| Samplers | Static samplers in root sig | Sampler argument buffer at fragment `buffer(0)` | +| Push constants | Root constants | Discrete set via `setVertex/FragmentBytes` | +| Residency | Implicit (heap is resident when set) | `use_heap` (textures) + `use_resource` (buffers) | +| Hard limit | 1M-entry heaps (effectively unbounded) | `kMaxArgumentBuffers = 8` descriptor sets | + +The design intent: **make Metal's indexing match D3D12's "the index is the index" semantics**, so +shared CPU data (material/draw indices) is correct on both backends with no per-backend fix-up. +Approach 1 broke that for textures; Approach 2 restores it, trading capacity for correctness. + +--- + +## 9. Residency: `use_heap` vs `use_resource` + +Getting the index right only solves *where* the shader looks. On Metal there's a second, separate +problem: the resource the argument buffer points at must be made **GPU-resident**, or the read +returns garbage. Pointing an argument buffer at a resource does *not* make it resident — that's +explicit, and it differs by how the resource was allocated. + +In hotline the two resource classes are allocated differently: + +- **Textures** are allocated *from* the bindless `MTLHeap` (`mtl_heap.new_texture`). One + `encoder.use_heap(&mtl_heap)` in `set_heap` makes the whole heap resident, covering every texture + the bindless argument buffer might index. +- **Structured buffers** (draw/material/light/etc.) are allocated *from the device* + (`device.new_buffer`), **not** from that heap. Two reasons they can't just live in the texture + heap: + 1. **Storage mode.** The texture heap is `Private` (GPU-only, blit-uploaded). The world buffers + are `Shared` + persistently-mapped and rewritten by the CPU every frame. A `Private` heap + can't host CPU-writable buffers, so they'd need a *separate* `Shared` heap. + 2. **Sizing.** `MTLHeap` is fixed-size and can't grow; buffer capacities are dynamic + (`reserve_world_buffers` at runtime, e.g. a 786 KB draw buffer vs a 160 B material buffer). + Sub-allocating dynamically-sized buffers from a fixed heap means over-allocating wildly or + recreating + re-encoding the heap on growth. + +So the buffers are device-allocated and only *referenced* from the heap's buffer argument buffer. +`use_heap` does nothing for them. They must be made resident explicitly with +`encoder.use_resource(buffer, Read|Write)` (mirroring `use_resource_at(..., Vertex|Fragment)` on the +render encoder) for **every** buffer in the pool — because bindless indices are resolved at runtime +in the shader, any buffer could be the one indexed. + +This produced a memorable bug. Without the `use_resource` calls, residency was left to Metal's +implicit heuristics: small/early device allocations happened to stay resident, large ones did not. +Symptoms: + +- One demo (small draw buffer) worked; another with the *same shader* but a large (786 KB) draw + buffer drew nothing — a size-dependent failure, easily mistaken for a CPU-side data difference. +- Hard-coding the world matrix to identity on the CPU **changed nothing** — the decisive clue. If + the data were the problem, identity would draw at the origin. It didn't, because the GPU was + never reading that buffer's memory: it wasn't resident. + +The fix is a few lines in `set_heap`: alongside each `use_heap`/`use_heap_at`, iterate the heap's +`buffer_slots` and `use_resource` each one. D3D12 has no analogue — a shader-visible descriptor heap +is resident for the duration it's set, so residency never surfaces as a separate step. + +--- + +## 10. Lessons / future + +- A compensation that only covers *some* index paths (buffers, not textures) is worse than none — + it hides the bug for the simple case and surfaces it only with a second array. +- The 8-set limit is now the real constraint. Options if a pipeline needs more: raise + `kMaxArgumentBuffers` in a newer SPIRV-Cross; or selectively re-pack **only buffer bindings** + (which *can* be compensated correctly through `get_lookup`, since they don't share CPU data) while + keeping textures one-set-per-binding. +- Keeping codegen (htwv) and runtime (mtl.rs) grouping identical is essential and fragile — any + change to the grouping key must be made in all four places at once. +- Indexing correctness and residency are *separate* problems on Metal. A correct bindless index + still reads garbage if the resource isn't resident — and because failure is residency-heuristic + driven, it looks data-dependent (works small, fails large), which sends you debugging the wrong + layer. "Hard-coding the data changes nothing" is the tell that it's residency, not data. + +### Key files +- `hotline-data/htwv/src/macos_impl.rs` — HLSL→MSL, descriptor-set assignment, limit warning +- `src/gfx/mtl.rs` — `build_slot_lookup`, `build_stage_binders`, `build_compute_binder`, `set_heap` +- `src/gfx/d3d12.rs` — reference root-signature mapping +- `src/pmfx.rs` — `get_lookup` / world-buffer indices +- `shaders/ecs.hlsl` — the bindless array declarations diff --git a/examples/bindful/main.rs b/examples/bindful/main.rs index 55ae1bad..6aa55812 100644 --- a/examples/bindful/main.rs +++ b/examples/bindful/main.rs @@ -153,8 +153,6 @@ fn main() -> Result<(), hotline_rs::Error> { cmdbuffer.set_scissor_rect(&scissor); cmdbuffer.set_render_pipeline(pso_pmfx); - cmdbuffer.set_heap(pso_pmfx, dev.get_shader_heap()); - let srv0 = textures[0].get_srv_index().unwrap(); let srv1 = textures[1].get_srv_index().unwrap(); let srv2 = textures[2].get_srv_index().unwrap(); diff --git a/examples/compute/main.rs b/examples/compute/main.rs new file mode 100644 index 00000000..51ad2b05 --- /dev/null +++ b/examples/compute/main.rs @@ -0,0 +1,177 @@ +use hotline_rs::{*, prelude::*}; + +use os::{App, Window}; +use gfx::{CmdBuf, Device, SwapChain, RenderPass, Texture}; + +#[repr(C)] +struct Vertex { + position: [f32; 2], + texcoord: [f32; 2], +} + +// matches `julia_constants` in shaders/julia.hlsl (5x 32-bit values) +#[repr(C)] +struct JuliaConstants { + output_index: u32, + width: u32, + height: u32, + cr: f32, + ci: f32, +} + +fn main() -> Result<(), hotline_rs::Error> { + let mut app = os_platform::App::create(os::AppInfo { + name: String::from("compute"), + window: false, + num_buffers: 0, + dpi_aware: true, + }); + + let mut dev = gfx_platform::Device::create(&gfx::DeviceInfo { + adapter_name: None, + shader_heap_size: 100, + render_target_heap_size: 100, + depth_stencil_heap_size: 100, + }); + print!("{}", dev.get_adapter_info()); + + let mut win = app.create_window(os::WindowInfo { + title: String::from("compute - animated julia set"), + rect: os::Rect { x: 100, y: 100, width: 1280, height: 720 }, + style: os::WindowStyleFlags::NONE, + parent_handle: None, + }); + + let swap_chain_info = gfx::SwapChainInfo { + num_buffers: 2, + format: gfx::Format::RGBA8n, + clear_colour: Some(gfx::ClearColour { r: 0.0, g: 0.0, b: 0.0, a: 1.0 }), + }; + let mut swap_chain = dev.create_swap_chain::(&swap_chain_info, &win)?; + let mut cmdbuffer = dev.create_cmd_buf(2); + + // fullscreen quad (NDC) with texcoords flipped so (0,0) is top-left of the image + let vertices = [ + Vertex { position: [-1.0, -1.0], texcoord: [0.0, 1.0] }, + Vertex { position: [-1.0, 1.0], texcoord: [0.0, 0.0] }, + Vertex { position: [ 1.0, 1.0], texcoord: [1.0, 0.0] }, + Vertex { position: [ 1.0, -1.0], texcoord: [1.0, 1.0] }, + ]; + let vertex_buffer = dev.create_buffer(&gfx::BufferInfo { + usage: gfx::BufferUsage::VERTEX, + cpu_access: gfx::CpuAccessFlags::NONE, + format: gfx::Format::Unknown, + stride: std::mem::size_of::(), + num_elements: 4, + initial_state: gfx::ResourceState::VertexConstantBuffer + }, Some(gfx::as_u8_slice(&vertices)))?; + + let indices: [u16; 6] = [0, 1, 2, 0, 2, 3]; + let index_buffer = dev.create_buffer(&gfx::BufferInfo { + usage: gfx::BufferUsage::INDEX, + cpu_access: gfx::CpuAccessFlags::NONE, + format: gfx::Format::R16u, + stride: std::mem::size_of::(), + num_elements: 6, + initial_state: gfx::ResourceState::IndexBuffer + }, Some(gfx::as_u8_slice(&indices)))?; + + // compute output texture - written by the julia kernel, sampled by the blit pass + let vp_rect = win.get_viewport_rect(); + let tex_width = vp_rect.width as u64; + let tex_height = vp_rect.height as u64; + let output_texture = dev.create_texture::(&gfx::TextureInfo { + format: gfx::Format::RGBA8n, + tex_type: gfx::TextureType::Texture2D, + width: tex_width, + height: tex_height, + depth: 1, + array_layers: 1, + mip_levels: 1, + samples: 1, + usage: gfx::TextureUsage::SHADER_RESOURCE | gfx::TextureUsage::UNORDERED_ACCESS, + initial_state: gfx::ResourceState::UnorderedAccess, + }, None)?; + let uav_index = output_texture.get_uav_index().unwrap() as u32; + let srv_index = output_texture.get_srv_index().unwrap() as u32; + + // load shaders and create the compute + blit pipelines via pmfx + let mut pmfx : pmfx::Pmfx = pmfx::Pmfx::create(&mut dev, 0); + pmfx.load(&hotline_rs::get_data_path("shaders/julia"))?; + pmfx.create_compute_pipeline(&dev, "julia")?; + pmfx.create_render_pipeline(&dev, "blit", swap_chain.get_backbuffer_pass())?; + + let blit_fmt = swap_chain.get_backbuffer_pass().get_format_hash(); + + let mut frame = 0u32; + while app.run() { + win.update(&mut app); + swap_chain.update::(&mut dev, &win, &mut cmdbuffer); + cmdbuffer.reset(&swap_chain); + + // animate the complex constant around a circle for a classic morphing julia set + let theta = frame as f32 * 0.0125; + let radius = 0.7885; + let constants = JuliaConstants { + output_index: uav_index, + width: tex_width as u32, + height: tex_height as u32, + cr: radius * theta.cos(), + ci: radius * theta.sin(), + }; + + // compute pass - dispatch the julia kernel into the rw texture + cmdbuffer.begin_event(0xff00ff00, "Julia Compute"); + let julia = pmfx.get_compute_pipeline("julia")?; + cmdbuffer.set_compute_pipeline(julia); + cmdbuffer.set_heap(julia, dev.get_shader_heap()); + cmdbuffer.push_compute_constants(julia, 0, 0, 5, 0, gfx::as_u8_slice(&constants)); + cmdbuffer.dispatch( + gfx::Size3 { x: (tex_width as u32 + 7) / 8, y: (tex_height as u32 + 7) / 8, z: 1 }, + gfx::Size3 { x: 8, y: 8, z: 1 } + ); + cmdbuffer.end_event(); + + // blit pass - draw the compute output to the back buffer + cmdbuffer.begin_event(0xff0000ff, "Blit Pass"); + let vp = gfx::Viewport::from(win.get_viewport_rect()); + let sc = gfx::ScissorRect::from(win.get_viewport_rect()); + + cmdbuffer.transition_barrier(&gfx::TransitionBarrier { + texture: Some(swap_chain.get_backbuffer_texture()), + buffer: None, + state_before: gfx::ResourceState::Present, + state_after: gfx::ResourceState::RenderTarget, + }); + + let blit = pmfx.get_render_pipeline_for_format("blit", blit_fmt)?; + cmdbuffer.begin_render_pass(swap_chain.get_backbuffer_pass_mut()); + cmdbuffer.set_viewport(&vp); + cmdbuffer.set_scissor_rect(&sc); + cmdbuffer.set_render_pipeline(blit); + cmdbuffer.set_heap(blit, dev.get_shader_heap()); + cmdbuffer.set_index_buffer(&index_buffer); + cmdbuffer.set_vertex_buffer(&vertex_buffer, 0); + let srv = [srv_index, 0, 0, 0]; + cmdbuffer.push_render_constants(blit, 0, 0, 4, 0, gfx::as_u8_slice(&srv)); + cmdbuffer.draw_indexed_instanced(6, 1, 0, 0, 0); + cmdbuffer.end_render_pass(); + + cmdbuffer.transition_barrier(&gfx::TransitionBarrier { + texture: Some(swap_chain.get_backbuffer_texture()), + buffer: None, + state_before: gfx::ResourceState::RenderTarget, + state_after: gfx::ResourceState::Present, + }); + cmdbuffer.end_event(); + + cmdbuffer.close()?; + dev.execute(&cmdbuffer); + swap_chain.swap(&mut dev); + frame += 1; + } + + swap_chain.wait_for_last_frame(); + dev.cleanup_dropped_resources(&swap_chain); + Ok(()) +} diff --git a/hotline-data b/hotline-data index d8f117c4..44e76c60 160000 --- a/hotline-data +++ b/hotline-data @@ -1 +1 @@ -Subproject commit d8f117c42d9e9960800eb288cae30231ebcaaf65 +Subproject commit 44e76c6066234a06824cb94d3ab46b2d741ba20e diff --git a/plugins/ecs_examples/src/bindless_material_ibl.rs b/plugins/ecs_examples/src/bindless_material_ibl.rs index 6f4d45d7..a6c91a99 100644 --- a/plugins/ecs_examples/src/bindless_material_ibl.rs +++ b/plugins/ecs_examples/src/bindless_material_ibl.rs @@ -1,9 +1,6 @@ /// /// Bindless Material IBL /// -/// 10 PBR materials × 4 mesh types arranged in a curated 10×4 grid, -/// lit by image-based lighting (cubemap + BRDF LUT) instead of point lights. -/// use crate::prelude::*; @@ -104,7 +101,7 @@ pub fn setup_bindless_material_ibl( let x = -x_offset + col as f32 * cell; let z = -z_offset + row as f32 * cell; - + let pos = vec3f(x, 0.0, z); commands.spawn(( diff --git a/plugins/ecs_examples/src/claude.rs b/plugins/ecs_examples/src/claude.rs index c654f2c5..a6319a56 100644 --- a/plugins/ecs_examples/src/claude.rs +++ b/plugins/ecs_examples/src/claude.rs @@ -1,6 +1,3 @@ -// currently windows only because here we need a concrete gfx and os implementation -#![cfg(target_os = "windows")] - use crate::prelude::*; /// diff --git a/plugins/ecs_examples/src/draw.rs b/plugins/ecs_examples/src/draw.rs index 1b4e3cde..ae91ba7f 100644 --- a/plugins/ecs_examples/src/draw.rs +++ b/plugins/ecs_examples/src/draw.rs @@ -1,6 +1,6 @@ /// /// Draw -/// +/// use crate::prelude::*; diff --git a/plugins/ecs_examples/src/draw_indexed.rs b/plugins/ecs_examples/src/draw_indexed.rs index 74af6363..8d91fa0a 100644 --- a/plugins/ecs_examples/src/draw_indexed.rs +++ b/plugins/ecs_examples/src/draw_indexed.rs @@ -2,7 +2,7 @@ /// Draw Indexed /// -use crate::prelude::*; +use crate::prelude::*; /// Sets up a single cube mesh to test draw indexed call with a single enity #[no_mangle] @@ -17,7 +17,7 @@ pub fn draw_indexed(client: &mut Client) } } -/// Set's up a single cube mesh. The draw all is made with `draw_meshes` in `draw.rs` +/// Set's up a single cube mesh. The draw all is made with `draw_meshes` in `draw.rs` #[export_update_fn] pub fn setup_draw_indexed( mut device: ResMut, @@ -44,7 +44,7 @@ pub fn draw_meshes_indexed( view: &pmfx::View, cmd_buf: &mut ::CmdBuf, mesh_draw_query: Query<(&WorldMatrix, &MeshComponent)>) -> Result<(), hotline_rs::Error> { - + let fmt = view.pass.get_format_hash(); let pipeline = pmfx.get_render_pipeline_for_format(&view.view_pipeline, fmt)?; let camera = pmfx.get_camera_constants(&view.camera)?; diff --git a/plugins/ecs_examples/src/draw_indirect.rs b/plugins/ecs_examples/src/draw_indirect.rs index 9c1d3479..a787fc19 100644 --- a/plugins/ecs_examples/src/draw_indirect.rs +++ b/plugins/ecs_examples/src/draw_indirect.rs @@ -22,13 +22,13 @@ pub fn draw_indirect(client: &mut Client pub fn setup_draw_indirect( mut device: ResMut, mut commands: Commands) -> Result<(), hotline_rs::Error> { - + let scalar_scale = 10.0; let scale = Mat34f::from_scale(splat3f(scalar_scale)); // draw indirect let tri = hotline_rs::primitives::create_triangle_mesh(&mut device.0); - let pos = Mat34f::from_translation(vec3f(-scalar_scale, scalar_scale, 0.0)); + let pos = Mat34f::from_translation(vec3f(-scalar_scale, scalar_scale, 0.0)); let args = gfx::DrawArguments { vertex_count_per_instance: 3, @@ -50,7 +50,7 @@ pub fn setup_draw_indirect( vec![gfx::IndirectArgument{ argument_type: gfx::IndirectArgumentType::Draw, arguments: None - }], + }], None ).unwrap(); @@ -65,7 +65,7 @@ pub fn setup_draw_indirect( // draw indexed indirect let teapot = hotline_rs::primitives::create_teapot_mesh(&mut device.0, 8); - let pos = Mat34f::from_translation(vec3f(scalar_scale, scalar_scale, 0.0)); + let pos = Mat34f::from_translation(vec3f(scalar_scale, scalar_scale, 0.0)); let args = gfx::DrawIndexedArguments { index_count_per_instance: teapot.num_indices, @@ -88,7 +88,7 @@ pub fn setup_draw_indirect( vec![gfx::IndirectArgument{ argument_type: gfx::IndirectArgumentType::DrawIndexed, arguments: None - }], + }], None ).unwrap(); @@ -110,9 +110,9 @@ pub fn draw_meshes_indirect( pmfx: &Res, view: &pmfx::View, cmd_buf: &mut ::CmdBuf, - mesh_draw_indirect_query: Query<(&WorldMatrix, &MeshComponent, &CommandSignatureComponent, &BufferComponent)>) + mesh_draw_indirect_query: Query<(&WorldMatrix, &MeshComponent, &CommandSignatureComponent, &BufferComponent)>) -> Result<(), hotline_rs::Error> { - + let fmt = view.pass.get_format_hash(); let pipeline = pmfx.get_render_pipeline_for_format(&view.view_pipeline, fmt)?; let camera = pmfx.get_camera_constants(&view.camera)?; @@ -126,11 +126,11 @@ pub fn draw_meshes_indirect( cmd_buf.set_vertex_buffer(&mesh.0.vb, 0); cmd_buf.execute_indirect( - &command.0, - 1, - &args.0, - 0, - None, + &command.0, + 1, + &args.0, + 0, + None, 0 ); } diff --git a/plugins/ecs_examples/src/draw_cbuffer_instanced.rs b/plugins/ecs_examples/src/draw_structured_buffer_instanced.rs similarity index 69% rename from plugins/ecs_examples/src/draw_cbuffer_instanced.rs rename to plugins/ecs_examples/src/draw_structured_buffer_instanced.rs index ccb98770..66ee7f74 100644 --- a/plugins/ecs_examples/src/draw_cbuffer_instanced.rs +++ b/plugins/ecs_examples/src/draw_structured_buffer_instanced.rs @@ -1,28 +1,28 @@ /// -/// Draw cbuffer Instanced +/// Draw Structured Buffer Instanced /// -use crate::prelude::*; +use crate::prelude::*; -/// Creates a instance batch, where the `InstanceBatch` parent will update a cbuffer containing -/// the cbuffer is created in a separate heap and the matrices and indexed into using the instance id system value semantic +/// Creates an instance batch where the per-instance world matrices live in a `StructuredBuffer` +/// bound as an SRV and indexed by `SV_InstanceID` in the vertex shader. #[no_mangle] -pub fn draw_cbuffer_instanced(client: &mut Client) -> ScheduleInfo { +pub fn draw_structured_buffer_instanced(client: &mut Client) -> ScheduleInfo { client.pmfx.load(&hotline_rs::get_data_path("shaders/ecs_examples").as_str()).unwrap(); ScheduleInfo { setup: systems![ - "setup_draw_cbuffer_instanced" + "setup_draw_structured_buffer_instanced" ], update: systems![ "rotate_meshes", "batch_world_matrix_instances" ], - render_graph: "mesh_draw_cbuffer_instanced" + render_graph: "mesh_draw_structured_buffer_instanced" } } #[export_update_fn] -pub fn setup_draw_cbuffer_instanced( +pub fn setup_draw_structured_buffer_instanced( mut device: bevy_ecs::change_detection::ResMut, mut commands: bevy_ecs::system::Commands) -> Result<(), hotline_rs::Error> { @@ -42,27 +42,29 @@ pub fn setup_draw_cbuffer_instanced( let mut rng = rand::thread_rng(); let size = 2.0; - let num = 32; // max number of bytes in cbuffer is 65536 + let num = 64; let instance_count = (num*num) as u32; let range = size * size * (num as f32); for mesh in meshes { + // Only one descriptor slot is needed: the per-batch instance buffer. instance_count is the + // element count inside that buffer, not the number of heap descriptors. let mut heap = device.create_heap(&gfx::HeapInfo { heap_type: gfx::HeapType::Shader, - num_descriptors: instance_count as usize, + num_descriptors: 1, debug_name: Some("instance_buffer_heap".to_string()) }); let parent = commands.spawn(InstanceBatch { mesh: MeshComponent(mesh.clone()), - pipeline: PipelineComponent("mesh_cbuffer_instanced".to_string()), - instance_buffer: InstanceBuffer { + pipeline: PipelineComponent("mesh_structured_buffer_instanced".to_string()), + instance_buffer: InstanceBuffer { buffer: device.create_buffer_with_heap(&gfx::BufferInfo{ - usage: gfx::BufferUsage::CONSTANT_BUFFER, + usage: gfx::BufferUsage::SHADER_RESOURCE, cpu_access: gfx::CpuAccessFlags::WRITE, format: gfx::Format::Unknown, stride: std::mem::size_of::(), num_elements: instance_count as usize, - initial_state: gfx::ResourceState::VertexConstantBuffer + initial_state: gfx::ResourceState::ShaderResource }, hotline_rs::data![], &mut heap).unwrap(), instance_count, heap: Some(heap) @@ -70,7 +72,7 @@ pub fn setup_draw_cbuffer_instanced( }).id(); for _ in 0..num { for _ in 0..num { - // spawn a bunch of entites with slightly randomised + // spawn a bunch of entites with slightly randomised let pos = vec3f(rng.gen(), rng.gen(), rng.gen()) * splat3f(range) * 2.0 - vec3f(range, 0.0, range); let rot = vec3f(rng.gen(), rng.gen(), rng.gen()) * f32::pi() * 2.0; commands.spawn(Instance { @@ -87,30 +89,35 @@ pub fn setup_draw_cbuffer_instanced( Ok(()) } -/// Renders all scene instance batches with cbuffer instance buffer +/// Renders all scene instance batches, sourcing per-instance world matrices from a StructuredBuffer SRV #[export_render_fn] -pub fn draw_meshes_cbuffer_instanced( +pub fn draw_meshes_structured_buffer_instanced( pmfx: &Res, view: &pmfx::View, cmd_buf: &mut ::CmdBuf, instance_draw_query: Query<(&InstanceBuffer, &MeshComponent, &PipelineComponent)> ) -> Result<(), hotline_rs::Error> { - + let pmfx = &pmfx; let fmt = view.pass.get_format_hash(); let camera = pmfx.get_camera_constants(&view.camera)?; - + for (instance_batch, mesh, pipeline) in &instance_draw_query { // set pipeline per batch let pipeline = pmfx.get_render_pipeline_for_format(&pipeline.0, fmt)?; cmd_buf.set_render_pipeline(pipeline); cmd_buf.push_render_constants(pipeline, 0, 0, 16, 0, gfx::as_u8_slice(&camera.view_projection_matrix)); - // bind the constant buffer (cbv) on the slot for b1, space0 specified in the shader + // set_heap binds the heap's argument buffer and (on metal) calls use_resource_at over + // its buffer_slots so device-allocated buffers like our SRV are made resident on the encoder + let heap = instance_batch.heap.as_ref().unwrap(); + cmd_buf.set_heap(pipeline, heap); + + // bind the structured buffer as an SRV at t0, space0 specified in the shader cmd_buf.set_binding( - pipeline, 1, 0, gfx::DescriptorType::ConstantBuffer, - instance_batch.heap.as_ref().unwrap(), - instance_batch.buffer.get_cbv_index().unwrap() + pipeline, 0, 0, gfx::DescriptorType::ShaderResource, + heap, + instance_batch.buffer.get_srv_index().unwrap() ); // bind vb, ib and draw instanced @@ -120,4 +127,4 @@ pub fn draw_meshes_cbuffer_instanced( } Ok(()) -} \ No newline at end of file +} diff --git a/plugins/ecs_examples/src/draw_vertex_buffer_instanced.rs b/plugins/ecs_examples/src/draw_vertex_buffer_instanced.rs index e29aae58..47e137ee 100644 --- a/plugins/ecs_examples/src/draw_vertex_buffer_instanced.rs +++ b/plugins/ecs_examples/src/draw_vertex_buffer_instanced.rs @@ -1,6 +1,6 @@ /// /// Draw Vertex Buffer Instanced -/// +/// use crate::prelude::*; @@ -46,7 +46,7 @@ pub fn setup_draw_vertex_buffer_instanced( let parent = commands.spawn(InstanceBatch { mesh: MeshComponent(mesh.clone()), pipeline: PipelineComponent("mesh_vertex_buffer_instanced".to_string()), - instance_buffer: InstanceBuffer { + instance_buffer: InstanceBuffer { buffer: device.create_buffer(&gfx::BufferInfo{ usage: gfx::BufferUsage::VERTEX, cpu_access: gfx::CpuAccessFlags::WRITE, @@ -61,7 +61,7 @@ pub fn setup_draw_vertex_buffer_instanced( }).id(); for _ in 0..num { for _ in 0..num { - // spawn a bunch of entites with slightly randomised + // spawn a bunch of entites with slightly randomised let pos = vec3f(rng.gen(), rng.gen(), rng.gen()) * splat3f(range) * 2.0 - vec3f(range, 0.0, range); let rot = vec3f(rng.gen(), rng.gen(), rng.gen()) * f32::pi() * 2.0; commands.spawn(Instance { @@ -86,7 +86,7 @@ pub fn draw_meshes_vertex_buffer_instanced( cmd_buf: &mut ::CmdBuf, instance_draw_query: Query<(&InstanceBuffer, &MeshComponent, &PipelineComponent)> ) -> Result<(), hotline_rs::Error> { - + let pmfx = &pmfx; let fmt = view.pass.get_format_hash(); let camera = pmfx.get_camera_constants(&view.camera)?; diff --git a/plugins/ecs_examples/src/lib.rs b/plugins/ecs_examples/src/lib.rs index 36bd93ee..a468b95a 100644 --- a/plugins/ecs_examples/src/lib.rs +++ b/plugins/ecs_examples/src/lib.rs @@ -1,6 +1,3 @@ -// currently windows only because here we need a concrete gfx and os implementation -#![cfg(target_os = "windows")] - /// Contains basic examples and unit tests of rendering and ecs functionality mod error_tests; mod draw; @@ -9,7 +6,7 @@ mod draw_push_constants; mod draw_indirect; mod geometry_primitives; mod draw_vertex_buffer_instanced; -mod draw_cbuffer_instanced; +mod draw_structured_buffer_instanced; mod bindless_texture; mod tangent_space_normal_maps; mod bindless_material; @@ -57,8 +54,8 @@ pub fn load_material( if !map_path.is_empty() { textures.push( image::load_texture_from_file( - device, - &format!("{}/{}", dir, map_path[0]), + device, + &format!("{}/{}", dir, map_path[0]), Some(&mut pmfx.shader_heap) ).unwrap() ); @@ -68,7 +65,7 @@ pub fn load_material( if textures.len() != 3 { return Err(hotline_rs::Error { msg: format!( - "hotline_rs::ecs:: error: material '{}' does not contain enough maps ({}/3)", + "hotline_rs::ecs:: error: material '{}' does not contain enough maps ({}/3)", dir, textures.len() ) @@ -207,8 +204,8 @@ const fn unit_aabb_corners() -> [Vec3f; 8] { pub fn batch_bindless_draw_data( mut pmfx: ResMut, draw_query: Query<(&WorldMatrix, &Extents)>) -> Result<(), hotline_rs::Error> { - - let world_buffers = pmfx.get_world_buffers_mut(); + + let world_buffers = pmfx.get_world_buffers_mut(); world_buffers.draw.clear(); world_buffers.extent.clear(); @@ -222,7 +219,7 @@ pub fn batch_bindless_draw_data( let emin = extents.aabb_min; let emax = extents.aabb_max; - + let transform_min = corners.iter().fold( Vec3f::max_value(), |acc, x| min(acc, world_matrix.0 * (emin + (emax - emin) * *x))); let transform_max = corners.iter().fold(-Vec3f::max_value(), |acc, x| max(acc, world_matrix.0 * (emin + (emax - emin) * *x))); @@ -249,7 +246,7 @@ pub fn render_meshes_bindless( Query<(&MeshComponent, &WorldMatrix), Without> ) ) -> Result<(), hotline_rs::Error> { - + let (instance_draw_query, single_draw_query) = queries; let pmfx = &pmfx; @@ -314,7 +311,7 @@ pub fn render_meshes( Query<(&WorldMatrix, &MeshComponent), (With, Without)>, Query<(&WorldMatrix, &MeshComponent), With>, )) -> Result<(), hotline_rs::Error> { - + let fmt = view.pass.get_format_hash(); let pipeline = pmfx.get_render_pipeline_for_format(&view.view_pipeline, fmt)?; let camera = pmfx.get_camera_constants(&view.camera)?; @@ -379,7 +376,7 @@ pub fn render_debug( if session_info.debug_draw_flags.is_empty() { return Ok(()); } - + let fmt = view.pass.get_format_hash(); let pipeline = pmfx.get_render_pipeline_for_format("imdraw_3d", fmt)?; let camera = pmfx.get_camera_constants(&view.camera)?; @@ -401,7 +398,7 @@ pub fn render_debug( if i % 20 == 0 { tint *= 0.125; } - + imdraw.add_line_3d(Vec3f::new(offset, 0.0, -scale), Vec3f::new(offset, 0.0, scale), Vec4f::from(tint)); imdraw.add_line_3d(Vec3f::new(-scale, 0.0, offset), Vec3f::new(scale, 0.0, offset), Vec4f::from(tint)); } @@ -436,13 +433,14 @@ pub fn render_debug( imdraw.add_frustum(constants.view_projection_matrix, Vec4f::white()); } } - + // submit the buffers imdraw.submit(&mut device.0, bb as usize).unwrap(); // draw cmd_buf.set_render_pipeline(&pipeline); cmd_buf.push_render_constants(pipeline, 0, 0, 16, 0, &camera.view_projection_matrix); + imdraw.draw_3d(cmd_buf, bb as usize); Ok(()) @@ -455,7 +453,7 @@ pub fn blit( view: &pmfx::View, cmd_buf: &mut ::CmdBuf ) -> Result<(), hotline_rs::Error> { - + let pmfx = &pmfx; let fmt = view.pass.get_format_hash(); let pipeline = pmfx.get_render_pipeline_for_format("imdraw_blit", fmt)?; @@ -487,7 +485,7 @@ pub fn cubemap_clear( view: &pmfx::View, cmd_buf: &mut ::CmdBuf ) -> Result<(), hotline_rs::Error> { - + let pmfx = &pmfx; let fmt = view.pass.get_format_hash(); let pipeline = pmfx.get_render_pipeline_for_format("cubemap_clear", fmt)?; @@ -538,8 +536,12 @@ pub fn dispatch_compute( ); } + // stash animation time (user_data.x) in the last resource slot (`resources.input7.index`). + // current compute passes use at most 5 of the 8 input slots, so input7 is always free. + cmd_buf.push_compute_constants(pipeline, 0, 1, 1, 7 * 4, gfx::as_u8_slice(&[pmfx.push_constant_user_data[0]])); + cmd_buf.set_heap(pipeline, &pmfx.shader_heap); - + cmd_buf.dispatch( pass.group_count, pass.numthreads @@ -621,7 +623,7 @@ pub fn update_tlas( } ); } - + if let Some(tlas) = t.tlas.as_ref() { let instance_buffer = device.create_raytracing_instance_buffer(&instances)?; cmd_buf.update_raytracing_tlas(tlas, &instance_buffer, instances.len(), gfx::AccelerationStructureRebuildMode::Refit); @@ -642,7 +644,7 @@ pub fn get_demos_ecs_examples() -> Vec { "draw_indirect", "geometry_primitives", "draw_vertex_buffer_instanced", - "draw_cbuffer_instanced", + "draw_structured_buffer_instanced", "bindless_texture", "tangent_space_normal_maps", "bindless_material", diff --git a/plugins/ecs_examples/src/point_lights.rs b/plugins/ecs_examples/src/point_lights.rs index 3e159d0c..9a17ae7e 100644 --- a/plugins/ecs_examples/src/point_lights.rs +++ b/plugins/ecs_examples/src/point_lights.rs @@ -1,11 +1,8 @@ -// currently windows only because here we need a concrete gfx and os implementation -#![cfg(target_os = "windows")] - use crate::prelude::*; /// /// Point Lights -/// +/// /// Init function for primitives demo #[no_mangle] @@ -65,7 +62,7 @@ pub fn setup_point_lights( let irc = rc as i32; let size = 10.0; - let half_size = size * 0.5; + let half_size = size * 0.5; let step = size * half_size; let half_extent = (rc-1.0) * step * 0.5; let start_pos = vec3f(-half_extent, size, -half_extent); @@ -97,9 +94,9 @@ pub fn setup_point_lights( #[export_update_fn] pub fn animate_point_lights( - time: Res, + time: Res, mut light_query: Query<&mut Position, With>) -> Result<(), hotline_rs::Error> { - + let t = time.accumulated; let r = sin(t); @@ -107,7 +104,7 @@ pub fn animate_point_lights( let rot1 = sin(-t); let rot2 = sin(t * 0.5); let rot3 = sin(-t * 0.5); - + let step = 1.0 / 16.0; let mut f = 0.0; let mut i = 0; diff --git a/plugins/ecs_examples/src/read_write_texture.rs b/plugins/ecs_examples/src/read_write_texture.rs index fc515e08..bdb144f9 100644 --- a/plugins/ecs_examples/src/read_write_texture.rs +++ b/plugins/ecs_examples/src/read_write_texture.rs @@ -12,11 +12,25 @@ pub fn read_write_texture(client: &mut Client, + mut pmfx: ResMut) -> Result<(), hotline_rs::Error> { + + // pack accumulated time into user_data so cs_write_texture3d can animate the volume + pmfx.push_constant_user_data[0] = time.accumulated.to_bits(); + + Ok(()) +} + #[export_update_fn] pub fn setup_read_write_texture( mut device: ResMut, diff --git a/plugins/ecs_examples/src/spot_lights.rs b/plugins/ecs_examples/src/spot_lights.rs index 61401bfa..7c372dad 100644 --- a/plugins/ecs_examples/src/spot_lights.rs +++ b/plugins/ecs_examples/src/spot_lights.rs @@ -114,7 +114,7 @@ pub fn setup_spot_lights( let size = 10.0; let height = 50.0; - let half_size = size * 0.5; + let half_size = size * 0.5; let step = size * half_size; let half_extent = (rc-1.0) * step * 0.5; let start_pos = vec3f(-half_extent, size, -half_extent); @@ -146,12 +146,12 @@ pub fn setup_spot_lights( #[export_update_fn] pub fn animate_spot_lights( - time: Res, + time: Res, mut light_query: Query<&mut Position, With>) -> Result<(), hotline_rs::Error> { - + let t = time.accumulated; let rot0 = t; - + let mut i = 0; for mut position in &mut light_query { if i < 16 { @@ -161,7 +161,7 @@ pub fn animate_spot_lights( let ss = 300.0 * ts; position.x = sin(fi * f32::two_pi()) * f32::tau() * ss; position.z = cos(fi * f32::two_pi()) * f32::tau() * ss; - + let pr = rotate_2d(position.xz(), rot0); position.set_xz(pr); } @@ -172,7 +172,7 @@ pub fn animate_spot_lights( let ss = 300.0 * ts; position.x = -sin(fi * f32::two_pi()) * f32::tau() * ss; position.z = cos(fi * f32::two_pi()) * f32::tau() * ss; - + let pr = rotate_2d(position.xz(), -rot0); position.set_xz(pr); } @@ -183,7 +183,7 @@ pub fn animate_spot_lights( let ss = 300.0 * ts; position.x = sin(fi * f32::two_pi()) * f32::tau() * ss; position.z = -cos(fi * f32::two_pi()) * f32::tau() * ss; - + let pr = rotate_2d(position.xz(), -rot0); position.set_xz(pr); } @@ -194,7 +194,7 @@ pub fn animate_spot_lights( let ss = 300.0 * ts; position.x = -sin(fi * f32::two_pi()) * f32::tau() * ss; position.z = -cos(fi * f32::two_pi()) * f32::tau() * ss; - + let pr = rotate_2d(position.xz(), rot0); position.set_xz(pr); } diff --git a/plugins/ecs_examples/src/tangent_space_normal_maps.rs b/plugins/ecs_examples/src/tangent_space_normal_maps.rs index a55b7954..40b7e2dd 100644 --- a/plugins/ecs_examples/src/tangent_space_normal_maps.rs +++ b/plugins/ecs_examples/src/tangent_space_normal_maps.rs @@ -1,13 +1,10 @@ -// currently windows only because here we need a concrete gfx and os implementation -#![cfg(target_os = "windows")] - -/// +/// /// Tangent Space Normal Maps -/// +/// -use crate::prelude::*; +use crate::prelude::*; -/// Init function for tangent space normal maps to debug tangents +/// Init function for tangent space normal maps to debug tangents #[no_mangle] pub fn tangent_space_normal_maps(client: &mut Client) -> ScheduleInfo { client.pmfx.load(hotline_rs::get_data_path("shaders/ecs_examples").as_str()).unwrap(); @@ -28,10 +25,10 @@ pub fn setup_tangent_space_normal_maps( mut device: ResMut, mut pmfx: ResMut, mut commands: Commands) -> Result<(), hotline_rs::Error> { - + let textures = [ - TextureComponent(image::load_texture_from_file(&mut device, - &hotline_rs::get_data_path("textures/pbr/antique-grate1/antique-grate1_normal.dds"), + TextureComponent(image::load_texture_from_file(&mut device, + &hotline_rs::get_data_path("textures/pbr/antique-grate1/antique-grate1_normal.dds"), Some(&mut pmfx.shader_heap) ).unwrap()) ]; @@ -55,7 +52,7 @@ pub fn render_meshes_debug_tangent_space( Query<&TextureComponent>, Query<(&WorldMatrix, &MeshComponent)> )) -> Result<(), hotline_rs::Error> { - + let pmfx = &pmfx; let fmt = view.pass.get_format_hash(); let pipeline = pmfx.get_render_pipeline_for_format(&view.view_pipeline, fmt)?; diff --git a/readme.md b/readme.md index fc5b2370..c848ad98 100644 --- a/readme.md +++ b/readme.md @@ -7,7 +7,7 @@ Hotline is a graphics engine and live coding tool that allows you to edit code, shaders, and render state without restating the application. - + Checkout a live [demo](https://www.youtube.com/watch?v=8a_qcqmpZlg)!. @@ -34,17 +34,17 @@ Windows with Direct3D12 is the first fully supported platform, macOS with metal For the time being it is recommended to use the repository from GitHub if you want to use the example `plugins` or standalone `examples`. If you just want to use the library then `crates.io` is suitable. There are some difficulties with publishing data and plugins which I hope to iron out in time. -### Building / Fetching Data +### Building / Fetching Data The [hotline-data](https://github.com/polymonster/hotline-data) repository is required to build and serve data for the examples and the example plugins, it is included as a submodule of this repository, you can clone with submodules as so: -``` +```text git clone https://github.com/polymonster/hotline.git --recursive ``` You can add the submodule after cloning or update the submodule to keep it in-sync with the main repository as follows: -``` +```text git submodule update --init --recursive ``` @@ -70,7 +70,7 @@ Any code changes made to the plugin libs will cause a rebuild and reload to happ ### Building from Visual Studio Code -There are included `tasks` and `launch` files for vscode including configurations for the client and the examples. Launching the `client` from vscode in debug or release will build the core hotline `lib`, `client`, `data` and `plugins`. +There are included `tasks` and `launch` files for vscode including configurations for the client and the examples. Launching the `client` from vscode in debug or release will build the core hotline `lib`, `client`, `data` and `plugins`. ## Adding Plugins @@ -100,7 +100,7 @@ impl Plugin for EmptyPlugin { } } - fn setup(&mut self, client: Client) + fn setup(&mut self, client: Client) -> Client { println!("plugin setup"); client @@ -159,7 +159,7 @@ You can then provide an initialisation function named after the demo this return pub fn primitives(client: &mut Client) -> ScheduleInfo { // load resources we may need client.pmfx.load(&hotline_rs::get_data_path("shaders/debug").as_str()).unwrap(); - + // fill out info ScheduleInfo { setup: systems![ @@ -209,10 +209,10 @@ You can also supply your own `update` systems to animate and move your entities. ```rust #[export_update_fn] fn update_cameras( - app: Res, - main_window: Res, + app: Res, + main_window: Res, mut query: Query<(&mut Position, &mut Rotation, &mut ViewProjectionMatrix), With> -) -> Result<(), { +) -> Result<(), { let app = &app.0; for (mut position, mut rotation, mut view_proj) in &mut query { // .. @@ -231,7 +231,7 @@ pub fn render_meshes( view: &pmfx::View, cmd_buf: &mut ::CmdBuf, mesh_draw_query: Query<(&WorldMatrix, &MeshComponent)>) -> Result<(), hotline_rs::Error> { - + let fmt = view.pass.get_format_hash(); let mesh_debug = pmfx.get_render_pipeline_for_format(&view.view_pipeline, fmt)?; let camera = pmfx.get_camera_constants(&view.camera)?; @@ -282,7 +282,7 @@ pub fn dispatch_compute( } cmd_buf.set_heap(pipeline, &pmfx.shader_heap); - + cmd_buf.dispatch( pass.group_count, pass.numthreads @@ -360,7 +360,7 @@ A quick example of a basic application setup: // include prelude for convenience use hotline_rs::prelude::*; -pub fn main() -> Result<(), hotline_rs::Error> { +pub fn main() -> Result<(), hotline_rs::Error> { // Create an Application let mut app = os_platform::App::create(os::AppInfo { name: String::from("triangle"), @@ -391,7 +391,7 @@ pub fn main() -> Result<(), hotline_rs::Error> { ..Default::default() }; let mut swap_chain = device.create_swap_chain::(&swap_chain_info, &window)?; - + /// Create a command buffer let mut cmd = device.create_cmd_buf(num_buffers); @@ -670,11 +670,11 @@ while app.run() { if player.is_ended() { // .. handle case where video is ended } - + // get texture if let Some(video_tex) = &player.get_texture() { let size = player.get_size(); - + // .. render } } @@ -844,7 +844,7 @@ This example provides instanced draws by updating entity world matrices on the C -Bindless texturing example - uses push constants to push a per draw call texture id for each entity. The texture id (shader resource view index) is used to lookup the texture inside an unbounded descriptor array in the fragment shader. +Bindless texturing example - uses push constants to push a per draw call texture id for each entity. The texture id (shader resource view index) is used to lookup the texture inside an unbounded descriptor array in the fragment shader. ### Tangent Space Normal Maps @@ -968,7 +968,7 @@ Inline raytracing shadows with mixed raster and raytracing workload. Raster pipe ## Tests -There are standalone tests and client/plugin tests to test graphics API features. This requires a test runner which has a GPU and is not headless, so I am using my home machine as a self-hosted actions runner. You can run the tests yourself but because of the requirement of a GPU device and plugin loading the tests need to be ran single threaded. +There are standalone tests and client/plugin tests to test graphics API features. This requires a test runner which has a GPU and is not headless, so I am using my home machine as a self-hosted actions runner. You can run the tests yourself but because of the requirement of a GPU device and plugin loading the tests need to be ran single threaded. ```text cargo test -- --test-threads=1 @@ -982,4 +982,4 @@ pmbuild test ## Contributing -Contributions of all kinds are welcome, you can make a fork and send a PR if you want to submit small fixes or improvements. Anyone interested in being more involved in development I am happy to take on people to help with the project of all experience levels, especially people with more experience in Rust. You can contact me if interested via [Twitter](twitter.com/polymonster) or [Discord](https://discord.com/invite/3yjXwJ8wJC). +Contributions of all kinds are welcome, you can make a fork and send a PR if you want to submit small fixes or improvements. Anyone interested in being more involved in development I am happy to take on people to help with the project of all experience levels, especially people with more experience in Rust. You can contact me if interested via [Twitter](twitter.com/polymonster) or [Discord](https://discord.com/invite/3yjXwJ8wJC). diff --git a/shaders/bindful.hlsl b/shaders/bindful.hlsl index 6dda7c24..14a01990 100644 --- a/shaders/bindful.hlsl +++ b/shaders/bindful.hlsl @@ -1,7 +1,7 @@ //#pragma argument(developmentfeatures) struct vs_input { - float3 position : POSITION; + float3 position : POSITION; float4 uv : TEXCOORD0; }; diff --git a/shaders/draw.hlsl b/shaders/draw.hlsl index 19feca05..6c3773c3 100644 --- a/shaders/draw.hlsl +++ b/shaders/draw.hlsl @@ -19,7 +19,7 @@ vs_output vs_mesh_identity(vs_input_mesh input) { output.texcoord = float4(input.texcoord, 0.0, 0.0); output.colour = float4(1.0, 1.0, 1.0, 1.0); output.normal = input.normal.xyz; - + return output; } @@ -39,13 +39,13 @@ vs_output vs_mesh(vs_input_mesh input) { output.texcoord = float4(input.texcoord, 0.0, 0.0); output.colour = material_colour; output.normal = normalize(mul(rot, input.normal.xyz)); - + return output; } // // textureles checkboard shader -// +// float4 ps_checkerboard(vs_output input) : SV_Target { float4 output = float4(input.normal.xyz * 0.5 + 0.5, 1.0); @@ -73,7 +73,7 @@ float4 ps_checkerboard(vs_output input) : SV_Target { // debug switches // u gradient //colour.rgb = uv_gradient(u % 1.0); - + // v gradient //colour.rgb = uv_gradient(v % 1.0); diff --git a/shaders/draw_instanced.hlsl b/shaders/draw_instanced.hlsl index d1913f73..ef179a36 100644 --- a/shaders/draw_instanced.hlsl +++ b/shaders/draw_instanced.hlsl @@ -6,44 +6,39 @@ struct vs_input_instance { float4 row0: TEXCOORD4; float4 row1: TEXCOORD5; float4 row2: TEXCOORD6; - float4 row3: TEXCOORD7; }; vs_output vs_mesh_vertex_buffer_instanced(vs_input_mesh input, vs_input_instance instance_input) { vs_output output; - float3x4 instance_matrix; - instance_matrix[0] = instance_input.row0; - instance_matrix[1] = instance_input.row1; - instance_matrix[2] = instance_input.row2; - float4 pos = float4(input.position.xyz, 1.0); - pos.xyz = mul(instance_matrix, pos); - output.position = mul(view_projection_matrix, pos); + float3 transformed; + transformed.x = dot(instance_input.row0, pos); + transformed.y = dot(instance_input.row1, pos); + transformed.z = dot(instance_input.row2, pos); + pos.xyz = transformed; + + output.position = mul(view_projection_matrix, float4(pos.xyz, 1.0)); output.world_pos = pos; output.texcoord = float4(input.texcoord, 0.0, 0.0); output.colour = float4(input.normal.xyz * 0.5 + 0.5, 1.0); output.normal = input.normal.xyz; - + return output; } // -// example using a cbuffer to lookup instance info from SV_InstanceID +// example using a structured buffer to lookup instance info from SV_InstanceID // -struct cbuffer_instance_data { - float3x4 cbuffer_world_matrix[1024]; -}; - -ConstantBuffer cbuffer_instance : register(b1); +StructuredBuffer instance_world_matrices : register(t0); -vs_output vs_mesh_cbuffer_instanced(vs_input_mesh input, uint iid: SV_InstanceID) { +vs_output vs_mesh_structured_buffer_instanced(vs_input_mesh input, uint iid: SV_InstanceID) { vs_output output; float4 pos = float4(input.position.xyz, 1.0); - pos.xyz = mul(cbuffer_instance.cbuffer_world_matrix[iid], pos); + pos.xyz = mul(instance_world_matrices[iid], pos); output.position = mul(view_projection_matrix, pos); output.world_pos = pos; diff --git a/shaders/draw_instanced.jsn b/shaders/draw_instanced.jsn index e583a002..82d36ae9 100644 --- a/shaders/draw_instanced.jsn +++ b/shaders/draw_instanced.jsn @@ -20,8 +20,8 @@ raster_state: cull_back topology: "TriangleList" } - mesh_cbuffer_instanced: { - vs: vs_mesh_cbuffer_instanced + mesh_structured_buffer_instanced: { + vs: vs_mesh_structured_buffer_instanced ps: ps_checkerboard push_constants: [ "view_push_constants" @@ -46,7 +46,7 @@ depends_on: ["debug"] } } - mesh_draw_cbuffer_instanced: { + mesh_draw_structured_buffer_instanced: { debug: { view: "main_view" pipelines: ["imdraw_3d"] @@ -54,8 +54,8 @@ } meshes: { view: "main_view_no_clear" - pipelines: ["mesh_cbuffer_instanced"] - function: "draw_meshes_cbuffer_instanced" + pipelines: ["mesh_structured_buffer_instanced"] + function: "draw_meshes_structured_buffer_instanced" depends_on: ["debug"] } } diff --git a/shaders/ecs.hlsl b/shaders/ecs.hlsl index 15e16a21..b11a9ec9 100644 --- a/shaders/ecs.hlsl +++ b/shaders/ecs.hlsl @@ -2,7 +2,7 @@ // contains core descriptor layout to be used among different / shared ecs systems // -// generic (fat) + non-skinned mesh vertex layout +// generic (fat) + non-skinned mesh vertex layout struct vs_input_mesh { float3 position: POSITION; float2 texcoord: TEXCOORD0; @@ -18,13 +18,13 @@ struct ps_output { // per view constants with basic camera transforms cbuffer view_push_constants : register(b0) { - float4x4 view_projection_matrix; + row_major float4x4 view_projection_matrix; float4 view_position; } // per entity draw constants used in CPU draw calls cbuffer draw_push_constants : register(b1) { - float3x4 world_matrix; + row_major float3x4 world_matrix; float4 material_colour; uint4 draw_indices; } @@ -61,7 +61,7 @@ ConstantBuffer resources: register(b0, space1); // bindless draw data for entites to look up by ID struct draw_data { - float3x4 world_matrix; + row_major float3x4 world_matrix; } // bindless material ID's which can be looked up into textures array @@ -118,7 +118,7 @@ struct directional_light_data { // camera data struct camera_data { - float4x4 view_projection_matrix; + row_major float4x4 view_projection_matrix; float4 view_position; float4 planes[6]; } @@ -136,21 +136,21 @@ StructuredBuffer materials[] : register(t0, space2); StructuredBuffer point_lights[] : register(t0, space3); StructuredBuffer spot_lights[] : register(t0, space4); StructuredBuffer directional_lights[] : register(t0, space5); -StructuredBuffer shadow_matrices[] : register(t0, space6); +StructuredBuffer shadow_matrices[] : register(t0, space6); -// textures -Texture2D textures[] : register(t0, space7); -Texture2DMS msaa8x_textures[] : register(t0, space8); -TextureCube cubemaps[] : register(t0, space9); -Texture2DArray texture_arrays[] : register(t0, space10); -Texture3D volume_textures[] : register(t0, space11); +// textures +Texture2D textures[] : register(t1, space7); +Texture2DMS msaa8x_textures[] : register(t1, space8); +TextureCube cubemaps[] : register(t1, space9); +Texture2DArray texture_arrays[] : register(t1, space10); +Texture3D volume_textures[] : register(t1, space11); // tlas RaytracingAccelerationStructure scene_tlas[] : register(t0, space12); // uav textures -RWTexture2D rw_textures[] : register(u0, space0); -RWTexture3D rw_volume_textures[] : register(u0, space1); +RWTexture2D rw_textures[] : register(u1, space0); +RWTexture3D rw_volume_textures[] : register(u1, space1); // main constants to obtain the indices of the buffer types ConstantBuffer world_buffer_info : register(b2); @@ -185,6 +185,6 @@ camera_data get_camera_data() { } // utility to return a shadow matrix by index -float4x4 get_shadow_matrix(uint shadow_index) { +row_major float4x4 get_shadow_matrix(uint shadow_index) { return shadow_matrices[world_buffer_info.shadow_matrix.x][shadow_index]; } \ No newline at end of file diff --git a/shaders/imdraw.hlsl b/shaders/imdraw.hlsl index bdba2069..135e3c16 100644 --- a/shaders/imdraw.hlsl +++ b/shaders/imdraw.hlsl @@ -28,7 +28,7 @@ struct ps_output { }; cbuffer view_push_constants : register(b0) { - float4x4 projection_matrix; + row_major float4x4 projection_matrix; }; struct vs_input_mesh { @@ -41,16 +41,16 @@ struct vs_input_mesh { vs_output vs_2d( vs_input_2d input ) { vs_output output; - + output.position = mul(projection_matrix, float4(input.position.xy, 0.0, 1.0)); output.colour = input.colour; - + return output; } ps_output ps_main( vs_output input ) { ps_output output; - + output.colour = input.colour; return output; @@ -63,7 +63,7 @@ vs_output vs_3d( vs_input_3d input ) float4 pos = float4(input.position.xyz, 1.0); output.position = mul(projection_matrix, pos); output.colour = input.colour; - + return output; } diff --git a/shaders/imgui.hlsl b/shaders/imgui.hlsl index 64601d0c..814f314c 100644 --- a/shaders/imgui.hlsl +++ b/shaders/imgui.hlsl @@ -1,6 +1,6 @@ cbuffer proj_matrix : register(b0, space0) { - float4x4 ProjectionMatrix; + row_major float4x4 ProjectionMatrix; }; struct VS_INPUT diff --git a/shaders/imgui.pmfx b/shaders/imgui.pmfx index 009cc1f1..ba5a3b7d 100644 --- a/shaders/imgui.pmfx +++ b/shaders/imgui.pmfx @@ -10,6 +10,18 @@ address_w: "Wrap", } } + render_target_blend_states: { + alpha: { + blend_enabled: true + src_blend: SrcAlpha + dst_blend: InvSrcAlpha + } + } + blend_states: { + alpha: { + render_target: ["alpha"] + } + } pipelines: { default: { vs: vs_main, @@ -19,7 +31,8 @@ ], static_samplers: { sampler0: "wrap_linear" - } + }, + blend_state: "alpha" } } } diff --git a/shaders/julia.hlsl b/shaders/julia.hlsl new file mode 100644 index 00000000..be93cf3f --- /dev/null +++ b/shaders/julia.hlsl @@ -0,0 +1,80 @@ +// +// animated julia set computed into a read-write texture, then blitted to the back buffer +// + +// bindless read-write texture array - the compute kernel writes the fractal here +RWTexture2D rw_texture[] : register(u0, space0); + +cbuffer julia_constants : register(b0) { + uint output_index; // uav index of the rw_texture to write into + uint output_width; + uint output_height; + float cr; // animated complex constant (real) + float ci; // animated complex constant (imaginary) +}; + +// simple hue ramp for colouring iteration counts +float3 palette(float t) { + return 0.5 + 0.5 * cos(6.28318 * (float3(1.0, 1.0, 1.0) * t + float3(0.0, 0.33, 0.67))); +} + +[numthreads(8, 8, 1)] +void cs_julia(uint2 did : SV_DispatchThreadID) { + if(did.x >= output_width || did.y >= output_height) { + return; + } + + // map pixel to complex plane [-1.5, 1.5] x [-1.0, 1.0] + float aspect = float(output_width) / float(output_height); + float2 uv = float2(did.xy) / float2(output_width, output_height); + float2 z; + z.x = (uv.x * 2.0 - 1.0) * 1.5 * aspect; + z.y = (uv.y * 2.0 - 1.0) * 1.5; + + const int max_iter = 256; + int i = 0; + for(; i < max_iter; ++i) { + float x = z.x * z.x - z.y * z.y + cr; + float y = 2.0 * z.x * z.y + ci; + z = float2(x, y); + if(dot(z, z) > 4.0) { + break; + } + } + + float t = float(i) / float(max_iter); + float3 colour = (i == max_iter) ? float3(0.0, 0.0, 0.0) : palette(t); + rw_texture[output_index][did.xy] = float4(colour, 1.0); +} + +// +// fullscreen blit of the compute output to the back buffer (bindless texture sample) +// + +struct vs_input { + float2 position : POSITION; + float2 texcoord : TEXCOORD; +}; + +struct ps_input { + float4 position : SV_POSITION; + float2 texcoord : TEXCOORD; +}; + +cbuffer blit_constants : register(b0) { + int4 blit_srv_index; // srv index of the compute output texture +}; + +Texture2D blit_textures[] : register(t0); +SamplerState blit_sampler : register(s0); + +ps_input vs_blit(vs_input input) { + ps_input output; + output.position = float4(input.position, 0.0, 1.0); + output.texcoord = input.texcoord; + return output; +} + +float4 ps_blit(ps_input input) : SV_Target { + return blit_textures[blit_srv_index[0]].Sample(blit_sampler, input.texcoord); +} diff --git a/shaders/julia.pmfx b/shaders/julia.pmfx new file mode 100644 index 00000000..52ba7b50 --- /dev/null +++ b/shaders/julia.pmfx @@ -0,0 +1,32 @@ +{ + include: [ + "julia.hlsl" + ] + sampler_states: { + linear_clamp: { + filter: "Linear", + address_u: "Clamp", + address_v: "Clamp", + address_w: "Clamp", + } + } + pipelines: { + julia: { + cs: cs_julia, + push_constants: [ + "julia_constants" + ] + } + blit: { + vs: vs_blit, + ps: ps_blit, + push_constants: [ + "blit_constants" + ] + static_samplers: { + blit_sampler: "linear_clamp" + } + topology: "TriangleList" + } + } +} diff --git a/shaders/material.hlsl b/shaders/material.hlsl index 36f0c548..fcb60757 100644 --- a/shaders/material.hlsl +++ b/shaders/material.hlsl @@ -18,22 +18,23 @@ vs_output_material vs_mesh_material(vs_input_mesh input, vs_input_entity_ids ent // get draw call info and transform world matrix draw_data draw = get_draw_data(entity_input.ids[0]); - float4 pos = float4(input.position.xyz, 1.0); + float4 pos = float4(input.position.xyz, 1.0); + pos.xyz = mul(draw.world_matrix, pos); output.position = mul(view_projection_matrix, pos); output.world_pos = pos; output.texcoord = float4(input.texcoord, 0.0, 0.0); - + float3x3 rot = (float3x3)draw.world_matrix; output.normal = normalize(mul(rot, input.normal)); output.tangent = normalize(mul(rot, input.tangent)); output.bitangent = normalize(mul(rot, input.bitangent)); - + // mat material_data mat = get_material_data(entity_input.ids[1]); output.ids = uint4(mat.albedo_id, mat.normal_id, mat.roughness_id, mat.padding); - + return output; } @@ -42,7 +43,7 @@ vs_output_material vs_mesh_material_indirect(vs_input_mesh input) { // get draw call info and transform world matrix draw_data draw = get_draw_data(indirect_ids.x); - float4 pos = float4(input.position.xyz, 1.0); + float4 pos = float4(input.position.xyz, 1.0); pos.xyz = mul(draw.world_matrix, pos); // get camera data and transform projection matrix @@ -60,14 +61,14 @@ vs_output_material vs_mesh_material_indirect(vs_input_mesh input) { material_data mat = get_material_data(indirect_ids.y); output.ids = uint4(mat.albedo_id, mat.normal_id, mat.roughness_id, mat.padding); - + return output; } vs_output_material vs_mesh_lit(vs_input_mesh input) { vs_output_material output; - float3x4 wm = world_matrix; + row_major float3x4 wm = world_matrix; float4 pos = float4(input.position.xyz, 1.0); pos.xyz = mul(wm, pos); @@ -86,7 +87,7 @@ vs_output_material vs_mesh_lit(vs_input_mesh input) { ps_output ps_mesh_debug_tangent_space(vs_output_material input) { ps_output output; - output.colour = float4(0.0, 0.0, 0.0, 0.0); + output.colour = float4(0.0, 0.0, 0.0, 1.0); float3 ts_normal = textures[draw_indices.x].Sample(sampler_wrap_linear, input.texcoord.xy).xyz; ts_normal = ts_normal * 2.0 - 1.0; @@ -109,7 +110,7 @@ ps_output ps_mesh_material(vs_output_material input) { output.colour = float4(0.0, 0.0, 0.0, 0.0); float2 tc = input.texcoord.xy; - + // sample maps // albedo @@ -150,16 +151,18 @@ ps_output ps_mesh_material(vs_output_material input) { light.radius, input.world_pos.xyz ); - + output.colour += atteniuation * light.colour * diffuse * albedo; output.colour += atteniuation * light.colour * specular; } } + output.colour.a = 1.0; return output; } float4 ps_mesh_material_instanced_ibl(vs_output_material input) : SV_TARGET { + float2 tc = input.texcoord.xy; // albedo @@ -240,7 +243,7 @@ ps_output ps_mesh_lit(vs_output input) { light.radius, input.world_pos.xyz ); - + output.colour += atteniuation * light.colour * diffuse; output.colour += atteniuation * light.colour * specular; } @@ -262,7 +265,7 @@ ps_output ps_mesh_lit(vs_output input) { light.cutoff, light.falloff ); - + output.colour += atteniuation * light.colour * diffuse; output.colour += atteniuation * light.colour * specular; } @@ -291,7 +294,7 @@ float4 ps_mesh_pbr_ibl(vs_output input) : SV_TARGET { float3 v = normalize(input.world_pos.xyz - view_position.xyz); float3 n = input.normal; - + float3 albedo = float3(1.0, 0.5, 0.0); float3 f0 = lerp(float3(0.04, 0.04, 0.04), albedo, metalness); @@ -299,7 +302,7 @@ float4 ps_mesh_pbr_ibl(vs_output input) : SV_TARGET { float3 rd = normalize(input.world_pos.xyz - view_position.xyz) * float3(1.0, 1.0, -1.0); float3 nd = normalize(input.normal.xyz * float3(1.0, 1.0, -1.0)); - float3 r = reflect(rd, nd); + float3 r = reflect(rd, nd); r.z *= -1.0; // irradiance / diffuse @@ -383,10 +386,10 @@ void scene_raygen_shader() // unproject ray float4 near = float4(ndc.x, ndc.y, 0.0, 1.0); float4 far = float4(ndc.x, ndc.y, 1.0, 1.0); - + float4 wnear = mul(inverse_wvp, near); wnear /= wnear.w; - + float4 wfar = mul(inverse_wvp, far); wfar /= wfar.w; @@ -399,13 +402,13 @@ void scene_raygen_shader() RayPayload payload = default_payload(); TraceRay( - scene_tlas[resource_indices.y], - RAY_FLAG_NONE, - 0xff, + scene_tlas[resource_indices.y], + RAY_FLAG_NONE, + 0xff, 0, 2, - 0, - ray, + 0, + ray, payload ); @@ -419,13 +422,13 @@ void scene_raygen_shader() payload = default_payload(); payload.bounce_count = bounce_count; TraceRay( - scene_tlas[resource_indices.y], - RAY_FLAG_NONE, - 0xff, + scene_tlas[resource_indices.y], + RAY_FLAG_NONE, + 0xff, 0, 2, - 0, - bounce_ray, + 0, + bounce_ray, payload ); @@ -543,7 +546,7 @@ void scene_closest_hit_shader(inout RayPayload payload, in BuiltInTriangleInters float3 ray_dir = refract(rd, geo_normal, refidx); float3 ray_start = ip + rd * 0.001; - + if(length(ray_dir) == 0.0) { ray_dir = reflect(rd, geo_normal); @@ -569,7 +572,7 @@ void scene_closest_hit_shader(inout RayPayload payload, in BuiltInTriangleInters return; } - + float2 tx = v0.texcoord * u + v1.texcoord * w + v2.texcoord * v; // checkerboard uv @@ -591,7 +594,7 @@ void scene_closest_hit_shader(inout RayPayload payload, in BuiltInTriangleInters float rxy = rx + ry > 1.0 ? 0.0 : rx + ry; float3 checkerboard = rxy < 0.001 ? 0.66 : 1.0; - + payload.col = float4(geo_normal, 1.0); payload.col.xyz = payload.col.xyz * 0.5 + 0.5 * checkerboard; @@ -667,7 +670,7 @@ ps_output ps_mesh_lit_rt_shadow(vs_output input) { light_colour += atteniuation * light.colour * specular; bool occluded = is_occluded(input.world_pos.xyz + input.normal * 0.1, -l, 0.1, rl + 0.1); - + if(!occluded) { output.colour += light_colour; } diff --git a/shaders/render_targets.hlsl b/shaders/render_targets.hlsl index dad76f20..075bcb0c 100644 --- a/shaders/render_targets.hlsl +++ b/shaders/render_targets.hlsl @@ -12,7 +12,7 @@ vs_output vs_heightmap(vs_input_mesh input) { float height = 200.0; float3 p1 = pos.xyz; - + float h = fbm(p1.xz + fbm(p1.xz + fbm(p1.xz, 6), 6), 6) * height; p1.y += h; @@ -60,7 +60,7 @@ ps_output_mrt ps_heightmap_example_mrt(vs_output input) { void cs_heightmap_mrt_resolve(uint2 did: SV_DispatchThreadID, uint2 group_id: SV_GroupID) { // grab the output dimension from input0 (which we write to) uint2 half_dim = resources.input0.dimension.xy / 2; - + // render into 4 quadrants float4 final = float4(0.0, 0.0, 0.0, 0.0); if(did.x < half_dim.x && did.y < half_dim.y) { diff --git a/shaders/shadows.hlsl b/shaders/shadows.hlsl index cd3afdd8..465c785a 100644 --- a/shaders/shadows.hlsl +++ b/shaders/shadows.hlsl @@ -44,7 +44,7 @@ float sample_shadow_pcf_9(float3 sp, uint sm_index, float2 sm_size) { samples[6] = float2(1.0, -1.0) * inv_sm_size; samples[7] = float2(1.0, 0.0) * inv_sm_size; samples[8] = float2(1.0, 1.0) * inv_sm_size; - + float shadow = 0.0; [unroll] @@ -73,7 +73,7 @@ float sample_shadow_cube_pcf_9(float3 cv, float d, uint sm_index, float sm_size) samples[6] = (b2 * 1.0 + t * -1.0) * inv_sm_size; samples[7] = (b2 * 1.0 + t * 0.0) * inv_sm_size; samples[8] = (b2 * 1.0 + t * -1.0) * inv_sm_size; - + float shadow = 0.0; [unroll] @@ -82,7 +82,6 @@ float sample_shadow_cube_pcf_9(float3 cv, float d, uint sm_index, float sm_size) } shadow /= 9.0; - // shadow = cubemaps[shadow_map_index].SampleCmp(sampler_shadow_compare, cv, d).r; return shadow; } @@ -104,7 +103,7 @@ float4 ps_single_directional_shadow(vs_output input) : SV_Target { int shadow_map_index = light.shadow_map.srv_index; float4x4 shadow_matrix = get_shadow_matrix(light.shadow_map.matrix_index); - + // project shadow coord float4 offset_pos = float4(input.world_pos.xyz, 1.0); @@ -113,10 +112,12 @@ float4 ps_single_directional_shadow(vs_output input) : SV_Target { sp.y *= -1.0; sp.xy = sp.xy * 0.5 + 0.5; - float shadow_sample = textures[shadow_map_index].Sample(sampler_clamp_point, sp.xy).r; - float shadow = sp.z >= shadow_sample ? 0.0 : 1.0; + // sample (non cmp) + //float shadow_sample = textures[shadow_map_index].Sample(sampler_clamp_point, sp.xy).r; + //float shadow = sp.z >= shadow_sample ? 0.0 : 1.0; - shadow = sample_shadow_pcf_9(sp, shadow_map_index, float2(4096.0, 4096.0)); + float2 sm_size = float2(4096.0, 4096.0); + float shadow = sample_shadow_pcf_9(sp, shadow_map_index, sm_size); float3 l = light.dir.xyz; float diffuse = lambert(l, n); @@ -128,12 +129,13 @@ float4 ps_single_directional_shadow(vs_output input) : SV_Target { float4 lit_colour = light.colour * diffuse + light.colour * specular; output = lit_colour * shadow + light.colour * 0.2; + output.a = 1.0; return output; } float4 ps_single_omni_shadow(vs_output input) : SV_Target { - + int i = 0; float ks = 2.0; float roughness = 0.9; @@ -158,7 +160,7 @@ float4 ps_single_omni_shadow(vs_output input) : SV_Target { light.radius, input.world_pos.xyz ); - + output += atteniuation * light.colour * diffuse; output += atteniuation * light.colour * specular; @@ -196,7 +198,7 @@ float4 ps_single_omni_shadow(vs_output input) : SV_Target { samples[6] = (b2 * 1.0 + t * -1.0) * inv_sm_size; samples[7] = (b2 * 1.0 + t * 0.0) * inv_sm_size; samples[8] = (b2 * 1.0 + t * -1.0) * inv_sm_size; - + float shadow = 0.0; [unroll] @@ -208,6 +210,9 @@ float4 ps_single_omni_shadow(vs_output input) : SV_Target { if(dot(n, l) >= 0.0) { shadow = 0.0; } - - return output * shadow; + + output.rgb *= shadow; + output.a = 1.0; + + return output; } \ No newline at end of file diff --git a/shaders/texture.hlsl b/shaders/texture.hlsl index d0a8bcfd..40414ac6 100644 --- a/shaders/texture.hlsl +++ b/shaders/texture.hlsl @@ -12,7 +12,7 @@ float4 ps_texture2d(vs_output input) : SV_Target { // // cubemap with bindless lookup -// +// float4 ps_cubemap(vs_output input) : SV_Target { float4 col = cubemaps[draw_indices.x] @@ -46,8 +46,8 @@ float4 ps_texture2d_array(vs_output input) : SV_Target { vs_output vs_texture3d(vs_input_mesh input) { vs_output output; - float3x4 wm = world_matrix; - + row_major float3x4 wm = world_matrix; + float4 pos = float4(input.position.xyz, 1.0); pos.xyz = mul(wm, pos); @@ -56,7 +56,7 @@ vs_output vs_texture3d(vs_input_mesh input) { output.texcoord = float4(input.position, 0.0); output.colour = float4(input.normal.xyz, 1.0); output.normal = input.normal.xyz; - + return output; } @@ -67,7 +67,7 @@ ps_output ps_volume_texture_ray_march_sdf(vs_output input) { float3 v = input.texcoord.xyz; float3 chebyshev_norm = chebyshev_normalize(v); float3 uvw = chebyshev_norm * 0.5 + 0.5; - + float max_samples = 64.0; float3x3 inv_rot; @@ -77,45 +77,45 @@ ps_output ps_volume_texture_ray_march_sdf(vs_output input) { inv_rot = transpose(inv_rot); float3 ray_dir = normalize(input.world_pos.xyz - view_position.xyz); - + ray_dir = mul(inv_rot, ray_dir); ray_dir = normalize(ray_dir); - + float3 vddx = ddx( uvw ); float3 vddy = ddy( uvw ); - + float3 scale = float3( - length(world_matrix[0].xyz), - length(world_matrix[1].xyz), + length(world_matrix[0].xyz), + length(world_matrix[1].xyz), length(world_matrix[2].xyz) ) * 2.0; - + float d = volume_textures[draw_indices.x].SampleGrad(sampler_wrap_linear, uvw, vddx, vddy).r; - + float3 col = float3( 0.0, 0.0, 0.0 ); float3 ray_pos = input.world_pos.xyz; float taken = 0.0; - float3 min_step = (scale / max_samples); - + float3 min_step = (scale / max_samples); + for( int s = 0; s < int(max_samples); ++s ) - { + { taken += 1.0 / max_samples; - + d = volume_textures[draw_indices.x].SampleGrad(sampler_wrap_linear, uvw, vddx, vddy).r; - + float3 step = ray_dir.xyz * float3(d / scale) * 0.5; - + uvw += step; - + if(uvw.x >= 1.0 || uvw.x <= 0.0) discard; - + if(uvw.y >= 1.0 || uvw.y <= 0.0) discard; - + if(uvw.z >= 1.0 || uvw.z <= 0.0) discard; - + if( d <= 0.3 ) break; } @@ -129,56 +129,56 @@ ps_output ps_volume_texture_ray_march_sdf(vs_output input) { ps_output ps_volume_texture_ray_march(vs_output input) { ps_output output; - + float depth = 1.0; float max_samples = 256.0; - + float3 v = input.texcoord.xyz; float3 chebyshev_norm = chebyshev_normalize(v); float3 uvw = chebyshev_norm * 0.5 + 0.5; - + float3x3 inv_rot; inv_rot[0] = world_matrix[0].xyz; inv_rot[1] = world_matrix[1].xyz; inv_rot[2] = world_matrix[2].xyz; inv_rot = transpose(inv_rot); - - float3 ray_dir = normalize(input.world_pos.xyz - view_position.xyz); + + float3 ray_dir = normalize(input.world_pos.xyz - view_position.xyz); ray_dir = mul( inv_rot, ray_dir ); - + float3 ray_step = chebyshev_normalize(ray_dir.xyz) / max_samples; - + float depth_step = 1.0 / max_samples; - + float3 vddx = ddx( uvw ); float3 vddy = ddy( uvw ); - + for(int s = 0; s < int(max_samples); ++s ) { - output.colour = + output.colour = volume_textures[draw_indices.x].SampleGrad(sampler_wrap_linear, uvw, vddx, vddy); - + if(output.colour.a != 0.0) break; - + depth -= depth_step; uvw += ray_step; - + if(uvw.x > 1.0 || uvw.x < 0.0) discard; - + if(uvw.y > 1.0 || uvw.y < 0.0) discard; - + if(uvw.z > 1.0 || uvw.z < 0.0) discard; - + if(s == int(max_samples)-1) discard; } - + output.colour.rgb *= lerp( 0.5, 1.0, depth ); - + return output; } @@ -188,24 +188,32 @@ ps_output ps_volume_texture_ray_march(vs_output input) { [numthreads(8, 8, 8)] void cs_write_texture3d(uint3 did : SV_DispatchThreadID) { - float3 dim = float3(64.0, 64.0, 64.0); - float3 grid_pos = did.xyz * 2.0 - float3(64.0, 64.0, 64.0); - - float4 sphere; - float d = 1.0; - - float nxz = voronoise(did.xz / 8.0, 1.0, 0.0); - float nxy = voronoise(did.xy / 8.0, 1.0, 0.0); - float nyz = voronoise(did.yz / 8.0, 1.0, 0.0); + // animated time, packed into the spare resource slot by `dispatch_compute` + float t = asfloat(resources.input7.index); + float3 grid_pos = did.xyz * 2.0 - float3(64.0, 64.0, 64.0); float3 n = normalize(grid_pos); - float nn = - abs(dot(n, float3(0.0, 1.0, 0.0))) * nxz + // scroll the noise fields over time so the volume churns + float2 flow = float2(t * 4.0, sin(t) * 4.0); + float nxz = voronoise(did.xz / 8.0 + flow, 1.0, 0.0); + float nxy = voronoise(did.xy / 8.0 - flow, 1.0, 0.0); + float nyz = voronoise(did.yz / 8.0 + flow.yx, 1.0, 0.0); + + float nn = + abs(dot(n, float3(0.0, 1.0, 0.0))) * nxz + abs(dot(n, float3(0.0, 0.0, 1.0))) * nxy + abs(dot(n, float3(1.0, 0.0, 0.0))) * nyz; - rw_volume_textures[resources.input0.index][did.xyz] = float4(nn, 0.0, 0.0, nn < 0.9 ? 0.0 : 1.0); + // a travelling spherical shell pulses in and out, carving the surface threshold + float radius = length(grid_pos) / 64.0; + float pulse = 0.6 + 0.35 * sin(t * 1.5 - radius * 4.0); + + // colour shifts through the spectrum over time for a livelier look + float3 colour = 0.5 + 0.5 * cos(t + nn * 6.0 + float3(0.0, 2.094, 4.188)); + + rw_volume_textures[resources.input0.index][did.xyz] = + float4(colour * nn, nn < pulse ? 0.0 : 1.0); } // diff --git a/shaders/util.hlsl b/shaders/util.hlsl index bf5a7ecf..6c30fde5 100644 --- a/shaders/util.hlsl +++ b/shaders/util.hlsl @@ -23,14 +23,18 @@ void cs_mip_chain_texture2d(uint2 did: SV_DispatchThreadID) { offsets[7] = uint2( 1, -1); offsets[8] = uint2( 0, -1); +#ifndef __spirv__ + // DXC's SPIR-V backend can't cast a groupshared (address-space 3) lvalue to void, + // which is what pmfx_touch expands to. Skip it on the spirv/metal path for now. pmfx_touch(group_accumulated[0]); +#endif float4 level_up = float4(0.0, 0.0, 0.0, 0.0); - + for(int i = 0; i < 9; ++i) { level_up += rw_texture[read][did.xy * 2]; } - + rw_texture[write][did.xy] = level_up / 9.0; } @@ -56,21 +60,21 @@ vs_output_ndc vs_ndc(vs_input_2d_texcoord input) { } cbuffer cubemap_clear_constants : register(b0) { - float4x4 inverse_wvp; + row_major float4x4 inverse_wvp; }; TextureCube cubemap : register(t0); -SamplerState sampler_wrap_linear : register(s0); +SamplerState sampler_wrap_linear : register(s0); float4 ps_cubemap_clear(vs_output_ndc input) : SV_Target { // unproject ray float2 ndc = input.ndc; float4 near = float4(ndc.x, ndc.y, 0.0, 1.0); float4 far = float4(ndc.x, ndc.y, 1.0, 1.0); - + float4 wnear = mul(inverse_wvp, near); wnear /= wnear.w; - + float4 wfar = mul(inverse_wvp, far); wfar /= wfar.w; diff --git a/src/client.rs b/src/client.rs index 5f70b44f..b66197c2 100644 --- a/src/client.rs +++ b/src/client.rs @@ -547,8 +547,15 @@ impl Client where D: gfx::Device, A: os::App, D::RenderPipeline: gfx ], }; - if !std::path::Path::new(&lib_path).join(name.to_string() + ".dll").exists() { - println!("hotline_rs::client:: plugin not found: {}/{}", lib_path, name); + #[cfg(target_os = "windows")] + let lib_file = format!("{}.dll", name); + #[cfg(target_os = "macos")] + let lib_file = format!("lib{}.dylib", name); + #[cfg(target_os = "linux")] + let lib_file = format!("lib{}.so", name); + + if !std::path::Path::new(&lib_path).join(&lib_file).exists() { + println!("hotline_rs::client:: plugin not found: {}/{}", lib_path, lib_file); return; } diff --git a/src/gfx.rs b/src/gfx.rs index e56dcec9..3c588dfb 100644 --- a/src/gfx.rs +++ b/src/gfx.rs @@ -5,6 +5,10 @@ pub mod null; #[cfg(target_os = "windows")] pub mod d3d12; +/// Implemets this interface with a metal backend. +#[cfg(target_os = "macos")] +pub mod mtl; + use crate::os; use std::any::Any; use serde::{Deserialize, Serialize}; @@ -12,8 +16,6 @@ use std::hash::Hash; use std::sync::{Arc, Mutex}; use maths_rs::max; -use null::*; - type Error = super::Error; // @@ -118,7 +120,7 @@ impl DropList { let to_drop: Vec> = { let mut drop_list = self.list.lock().unwrap(); let initial_len = drop_list.len(); - + // Partition: extract items ready to drop, keep the rest let mut ready_to_drop = Vec::new(); drop_list.retain_mut(|drop_res| { @@ -144,12 +146,12 @@ impl DropList { } } }); - + if initial_len > 0 || ready_to_drop.len() > 0 { - println!("DropList cleanup: {} items, dropping {}, {} remaining", + println!("DropList cleanup: {} items, dropping {}, {} remaining", initial_len, ready_to_drop.len(), drop_list.len()); } - + ready_to_drop }; // Resources drop here, OUTSIDE the mutex lock @@ -387,7 +389,7 @@ pub struct PipelineStatistics { /// Information to pass to `Device::create_swap_chain`. pub struct SwapChainInfo { - /// Number of internal buffers to keep behind the scenes, which are swapped between each frame + /// Number of internal buffers to keep behind the scenes, which are swapped between each frame /// to allow overlapped CPU/GPU command buffer producer / consumer pub num_buffers: u32, /// Must be BGRA8n, RGBA8n or RGBA16f. @@ -434,7 +436,7 @@ pub struct ShaderCompileInfo { } /// The stage to which a shader will bind itself. -#[derive(Copy, Clone)] +#[derive(Copy, Clone, PartialEq, Eq, Hash)] pub enum ShaderType { Vertex, Fragment, @@ -569,7 +571,7 @@ pub struct PipelineLayout { pub bindings: Option>, /// Small amounts of data that can be pushed into a command buffer and available as data in shaders pub push_constants: Option>, - /// Static samplers that come along with the pipeline, + /// Static samplers that come along with the pipeline, pub static_samplers: Option>, } @@ -585,16 +587,19 @@ pub enum ResourceType { /// Cubemap texture (TextureCube) TextureCube, /// Multi-sampled 2D texture (Texture2DMS) - #[serde(rename = "Texture2DMS")] Texture2DMS, /// Read-write 2D texture (RWTexture2D) RWTexture2D, - /// Read-write 3D texture (RWTexture3D) + /// Read-write 3D texture (RWTexture3D) RWTexture3D, /// Structured buffer (StructuredBuffer) StructuredBuffer, /// Read-write structured buffer (RWStructuredBuffer) RWStructuredBuffer, + /// Append structured buffer (AppendStructuredBuffer) + AppendStructuredBuffer, + /// Consume structured buffer (ConsumeStructuredBuffer) + ConsumeStructuredBuffer, /// Constant buffer (cbuffer/ConstantBuffer) #[serde(alias = "cbuffer")] ConstantBuffer, @@ -627,7 +632,7 @@ pub struct DescriptorBinding { } /// Describes the type of descriptor binding to create. -#[derive(Clone, Copy, Serialize, Deserialize, Hash)] +#[derive(Clone, Copy, Serialize, Deserialize, Hash, PartialEq, Eq)] pub enum DescriptorType { /// Used for textures or structured buffers. ShaderResource, @@ -667,7 +672,7 @@ pub struct ResourceViewInfo { pub num_elements: usize } -/// Describes space in the shader to send data to via `CmdBuf::push_constants`. +/// Describes space in the shader to send data to via `CmdBuf::push_constants`. #[derive(Clone, Serialize, Deserialize)] pub struct PushConstantInfo { /// The shader stage the constants will be accessible to. @@ -686,7 +691,7 @@ pub struct PipelineSlotInfo { /// The slot in the pipeline layout to bind to pub index: u32, /// The number of descriptors or the number of 32-bit push constant values, if `None` the table is unbounded - pub count: Option + pub count: Option, } /// Input layout describes the layout of vertex buffers bound to the input assembler. @@ -870,7 +875,7 @@ pub struct DepthStencilInfo { } /// Write to the depth buffer, or omit writes and just perform depth testing -#[derive(Clone, Copy, Serialize, Deserialize)] +#[derive(Clone, Copy, PartialEq, Serialize, Deserialize)] pub enum DepthWriteMask { Zero, All, @@ -1194,7 +1199,7 @@ pub enum ResourceState { CopyDst, /// Used as destination to read back data from buffers / queries GenericRead, - /// Used for argument buffer in `execute_indirect` calls + /// Used for argument buffer in `execute_indirect` calls IndirectArgument, /// Used for destination acceleration structure buffers AccelerationStructure @@ -1333,7 +1338,7 @@ pub struct IndirectArgument { pub arguments: Option } -/// Structure of arguments which can be used to execute `draw_instanced` calls indirectly +/// Structure of arguments which can be used to execute `draw_instanced` calls indirectly #[repr(C)] #[derive(Clone, Copy)] pub struct DrawArguments { @@ -1343,7 +1348,7 @@ pub struct DrawArguments { pub start_instance_location: u32 } -/// Structure of arguments which can be used to execute `draw_indexed_instanced` calls indirectly +/// Structure of arguments which can be used to execute `draw_indexed_instanced` calls indirectly #[repr(C)] #[derive(Clone, Copy)] pub struct DrawIndexedArguments { @@ -1429,7 +1434,7 @@ pub trait Device: 'static + Send + Sync + Sized + Any + Clone { window: &A::Window, ) -> Result; /// Create a new `CmdBuf` with `num_buffers` internal buffers, the buffers can be swapped and syncronised - /// with a new `SwapChain` to allow in-flight gpu/cpu overlapped prodicer consumers + /// with a new `SwapChain` to allow in-flight gpu/cpu overlapped prodicer consumers fn create_cmd_buf(&self, num_buffers: u32) -> Self::CmdBuf; /// Create a new `Shader` from `ShaderInfo` fn create_shader(&self, info: &ShaderInfo, src: &[T]) -> Result; @@ -1522,7 +1527,7 @@ pub trait Device: 'static + Send + Sync + Sized + Any + Clone { ) -> Result; /// Create a command signature for `execute_indirect` commands associated on the `RenderPipeline` fn create_indirect_render_command( - &mut self, + &mut self, arguments: Vec, pipeline: Option<&Self::RenderPipeline> ) -> Result; @@ -1545,7 +1550,7 @@ pub trait Device: 'static + Send + Sync + Sized + Any + Clone { /// Read back u64 timestamp values as values in seconds, the vector will be empty if the buffer is yet to be written /// on the GPU fn read_timestamps(&self, swap_chain: &Self::SwapChain, buffer: &Self::Buffer, size_bytes: usize, frame_written_fence: u64) -> Vec; - /// Read back a single pipeline statistics query, assuming `buffer` was created with `create_read_back_buffer` + /// Read back a single pipeline statistics query, assuming `buffer` was created with `create_read_back_buffer` /// and is of size `get_pipeline_statistics_size_bytes()`. None is returned if the buffer is not ready fn read_pipeline_statistics(&self, swap_chain: &Self::SwapChain, buffer: &Self::Buffer, frame_written_fence: u64) -> Option; /// Reorts internal graphics api backend resources @@ -1591,7 +1596,7 @@ pub trait SwapChain: 'static + Sized + Any + Send + Sync + Clone { /// Call swap at the end of the frame to swap the back buffer, we rotate through n-buffers fn swap(&mut self, device: &mut D); } - + /// Responsible for buffering graphics commands. Internally it will contain a platform specific /// command list for each buffer in the associated swap chain. /// At the start of each frame `reset` must be called with an associated swap chain to internally switch @@ -1601,10 +1606,10 @@ pub trait CmdBuf: Send + Sync + Clone { /// Reset the `CmdBuf` for use on a new frame, it will be syncronised with the `SwapChain` so that /// in-flight command buffers are not overwritten fn reset(&mut self, swap_chain: &D::SwapChain); - /// Call close to the command buffer after all commands have been added and before passing to `Device::execute` + /// Call close to the command buffer after all commands have been added and before passing to `Device::execute` fn close(&mut self) -> Result<(), Error>; /// Internally the `CmdBuf` contains a set of buffers which it rotates through to allow inflight operations - /// to complete, this value indicates the buffer number you should `write` to during the current frame + /// to complete, this value indicates the buffer number you should `write` to during the current frame fn get_backbuffer_index(&self) -> u32; /// Begins a render pass, end must be called fn begin_render_pass(&mut self, render_pass: &D::RenderPass); @@ -1621,7 +1626,7 @@ pub trait CmdBuf: Send + Sync + Clone { fn timestamp_query(&mut self, heap: &mut D::QueryHeap, resolve_buffer: &mut D::Buffer); /// Begin a new query in the heap, it will allocate an index which is returned as `usize` fn begin_query(&mut self, heap: &mut D::QueryHeap, query_type: QueryType) -> usize; - /// End a query that was made on the heap results will be pushed into the `resolve_buffer` + /// End a query that was made on the heap results will be pushed into the `resolve_buffer` /// the data can be read by `Device::read_buffer` or specialisations such as `read_pipeline_statistics` fn end_query(&mut self, heap: &mut D::QueryHeap, query_type: QueryType, index: usize, resolve_buffer: &mut D::Buffer); /// Add a transition barrier for resources to change states based on info supplied in `TransitionBarrier` @@ -1679,9 +1684,9 @@ pub trait CmdBuf: Send + Sync + Clone { /// Issue indirect commands with signature created from `create_indirect_render_command` fn execute_indirect( &mut self, - command: &D::CommandSignature, - max_command_count: u32, - argument_buffer: &D::Buffer, + command: &D::CommandSignature, + max_command_count: u32, + argument_buffer: &D::Buffer, argument_buffer_offset: usize, counter_buffer: Option<&D::Buffer>, counter_buffer_offset: usize @@ -1692,21 +1697,21 @@ pub trait CmdBuf: Send + Sync + Clone { fn update_raytracing_tlas(&mut self, tlas: &D::RaytracingTLAS, instance_buffer: &D::Buffer, instance_count: usize, mode: AccelerationStructureRebuildMode); /// Resolves the `subresource` (mip index, 3d texture slice or array slice) fn resolve_texture_subresource(&mut self, texture: &D::Texture, subresource: u32) -> Result<(), Error>; - /// Generates a full mip chain for the specified `texture` where `heap` is the shader heap the texture was created on + /// Generates a full mip chain for the specified `texture` where `heap` is the shader heap the texture was created on fn generate_mip_maps(&mut self, texture: &D::Texture, device: &D, heap: &D::Heap) -> Result<(), Error>; /// Read back the swapchains contents to CPU fn read_back_backbuffer(&mut self, swap_chain: &D::SwapChain) -> Result; /// Copy from one buffer to another with offsets fn copy_buffer_region( - &mut self, - dst_buffer: &D::Buffer, - dst_offset: usize, - src_buffer: &D::Buffer, + &mut self, + dst_buffer: &D::Buffer, + dst_offset: usize, + src_buffer: &D::Buffer, src_offset: usize, num_bytes: usize ); /// Copy from one texture to another with offsets, if `None` is specified for `src_region` - /// it will copy the full size of src + /// it will copy the full size of src fn copy_texture_region( &mut self, dst_texture: &D::Texture, @@ -1726,7 +1731,7 @@ pub trait Buffer: Send + Sync { /// this function internally will map and unmap fn update(&mut self, offset: usize, data: &[T]) -> Result<(), Error>; // TODO: should be mut surely? // write data directly to the buffer, the buffer is required to be persistently mapped - fn write(&mut self, offset: usize, data: &[T]) -> Result<(), Error>; + fn write(&mut self, offset: usize, data: &[T]) -> Result<(), Error>; /// maps the entire buffer for reading or writing... see MapInfo fn map(&mut self, info: &MapInfo) -> *mut u8; /// unmap buffer... see UnmapInfo diff --git a/src/gfx/d3d12.rs b/src/gfx/d3d12.rs index 275d57a6..abe78e2e 100644 --- a/src/gfx/d3d12.rs +++ b/src/gfx/d3d12.rs @@ -1290,7 +1290,7 @@ impl Device { let h = get_binding_descriptor_hash(constants.shader_register, constants.register_space, super::DescriptorType::PushConstants); lookup.insert(h, PipelineSlotInfo { index: slot_iter, - count: Some(constants.num_values) + count: Some(constants.num_values), }); slot_iter += 1; } @@ -1357,7 +1357,7 @@ impl Device { let h = get_binding_descriptor_hash(binding.shader_register, binding.register_space, binding.binding_type); lookup.entry(h).or_insert(PipelineSlotInfo { index: slot_iter, - count: binding.num_descriptors + count: binding.num_descriptors, }); } slot_iter += 1; diff --git a/src/gfx/mtl.rs b/src/gfx/mtl.rs index 76c16407..6180fb9d 100644 --- a/src/gfx/mtl.rs +++ b/src/gfx/mtl.rs @@ -1,4 +1,7 @@ #![cfg(target_os = "macos")] +// objc 0.2.x's sel_impl macro uses cfg(cargo-clippy) which is not a valid identifier +// so it cannot be declared via check-cfg; suppress the lint for this file. +#![allow(unexpected_cfgs)] use crate::os_platform; use crate::os::Window; @@ -24,6 +27,8 @@ use std::collections::HashMap; use std::result; use cocoa::{appkit::NSView, base::id as cocoa_id}; +#[allow(unused_imports)] +use objc::{msg_send, sel, sel_impl}; use core_graphics_types::geometry::CGSize; use std::path::Path; @@ -49,7 +54,7 @@ const fn to_mtl_vertex_format(format: super::Format) -> MTLVertexFormat { super::Format::RGB32u => MTLVertexFormat::UInt3, super::Format::RGB32i => MTLVertexFormat::Int3, super::Format::RGB32f => MTLVertexFormat::Float3, - super::Format::RGBA8n => MTLVertexFormat::Char4Normalized, + super::Format::RGBA8n => MTLVertexFormat::UChar4Normalized, super::Format::RGBA8u => MTLVertexFormat::UChar4, super::Format::RGBA8i => MTLVertexFormat::Char4, super::Format::RGBA16u => MTLVertexFormat::UShort4, @@ -62,6 +67,78 @@ const fn to_mtl_vertex_format(format: super::Format) -> MTLVertexFormat { } } +fn to_mtl_primitive_type(topology: Topology) -> metal::MTLPrimitiveType { + match topology { + Topology::PointList => metal::MTLPrimitiveType::Point, + Topology::LineList => metal::MTLPrimitiveType::Line, + Topology::LineStrip => metal::MTLPrimitiveType::LineStrip, + Topology::TriangleList => metal::MTLPrimitiveType::Triangle, + Topology::TriangleStrip => metal::MTLPrimitiveType::TriangleStrip, + _ => metal::MTLPrimitiveType::Triangle, + } +} + +fn to_mtl_blend_factor(factor: &super::BlendFactor) -> metal::MTLBlendFactor { + match factor { + super::BlendFactor::Zero => metal::MTLBlendFactor::Zero, + super::BlendFactor::One => metal::MTLBlendFactor::One, + super::BlendFactor::SrcColour => metal::MTLBlendFactor::SourceColor, + super::BlendFactor::InvSrcColour => metal::MTLBlendFactor::OneMinusSourceColor, + super::BlendFactor::SrcAlpha => metal::MTLBlendFactor::SourceAlpha, + super::BlendFactor::InvSrcAlpha => metal::MTLBlendFactor::OneMinusSourceAlpha, + super::BlendFactor::DstAlpha => metal::MTLBlendFactor::DestinationAlpha, + super::BlendFactor::InvDstAlpha => metal::MTLBlendFactor::OneMinusDestinationAlpha, + super::BlendFactor::DstColour => metal::MTLBlendFactor::DestinationColor, + super::BlendFactor::InvDstColour => metal::MTLBlendFactor::OneMinusDestinationColor, + super::BlendFactor::SrcAlphaSat => metal::MTLBlendFactor::SourceAlphaSaturated, + super::BlendFactor::BlendFactor => metal::MTLBlendFactor::BlendColor, + super::BlendFactor::InvBlendFactor => metal::MTLBlendFactor::OneMinusBlendColor, + super::BlendFactor::Src1Colour => metal::MTLBlendFactor::Source1Color, + super::BlendFactor::InvSrc1Colour => metal::MTLBlendFactor::OneMinusSource1Color, + super::BlendFactor::Src1Alpha => metal::MTLBlendFactor::Source1Alpha, + super::BlendFactor::InvSrc1Alpha => metal::MTLBlendFactor::OneMinusSource1Alpha, + } +} + +fn to_mtl_blend_op(op: &super::BlendOp) -> metal::MTLBlendOperation { + match op { + super::BlendOp::Add => metal::MTLBlendOperation::Add, + super::BlendOp::Subtract => metal::MTLBlendOperation::Subtract, + super::BlendOp::RevSubtract => metal::MTLBlendOperation::ReverseSubtract, + super::BlendOp::Min => metal::MTLBlendOperation::Min, + super::BlendOp::Max => metal::MTLBlendOperation::Max, + } +} + +fn to_mtl_write_mask(mask: &super::WriteMask) -> metal::MTLColorWriteMask { + let mut mtl_mask = metal::MTLColorWriteMask::empty(); + if mask.contains(super::WriteMask::RED) { + mtl_mask |= metal::MTLColorWriteMask::Red; + } + if mask.contains(super::WriteMask::GREEN) { + mtl_mask |= metal::MTLColorWriteMask::Green; + } + if mask.contains(super::WriteMask::BLUE) { + mtl_mask |= metal::MTLColorWriteMask::Blue; + } + if mask.contains(super::WriteMask::ALPHA) { + mtl_mask |= metal::MTLColorWriteMask::Alpha; + } + mtl_mask +} + +fn to_mtl_texture_type(tex_type: super::TextureType) -> metal::MTLTextureType { + match tex_type { + super::TextureType::Texture1D => metal::MTLTextureType::D1, + super::TextureType::Texture1DArray => metal::MTLTextureType::D1Array, + super::TextureType::Texture2D => metal::MTLTextureType::D2, + super::TextureType::Texture2DArray => metal::MTLTextureType::D2Array, + super::TextureType::Texture3D => metal::MTLTextureType::D3, + super::TextureType::TextureCube => metal::MTLTextureType::Cube, + super::TextureType::TextureCubeArray => metal::MTLTextureType::CubeArray, + } +} + fn to_mtl_texture_usage(usage: TextureUsage) -> MTLTextureUsage { let mut mtl_usage : MTLTextureUsage = MTLTextureUsage::Unknown; if usage.contains(super::TextureUsage::SHADER_RESOURCE) { @@ -80,14 +157,192 @@ fn to_mtl_texture_usage(usage: TextureUsage) -> MTLTextureUsage { mtl_usage } +fn to_mtl_compare_func(func: super::ComparisonFunc) -> metal::MTLCompareFunction { + match func { + super::ComparisonFunc::Never => metal::MTLCompareFunction::Never, + super::ComparisonFunc::Less => metal::MTLCompareFunction::Less, + super::ComparisonFunc::Equal => metal::MTLCompareFunction::Equal, + super::ComparisonFunc::LessEqual => metal::MTLCompareFunction::LessEqual, + super::ComparisonFunc::Greater => metal::MTLCompareFunction::Greater, + super::ComparisonFunc::NotEqual => metal::MTLCompareFunction::NotEqual, + super::ComparisonFunc::GreaterEqual => metal::MTLCompareFunction::GreaterEqual, + super::ComparisonFunc::Always => metal::MTLCompareFunction::Always, + } +} + +fn to_mtl_sampler_address_mode(mode: super::SamplerAddressMode) -> metal::MTLSamplerAddressMode { + match mode { + super::SamplerAddressMode::Wrap => metal::MTLSamplerAddressMode::Repeat, + super::SamplerAddressMode::Mirror => metal::MTLSamplerAddressMode::MirrorRepeat, + super::SamplerAddressMode::Clamp => metal::MTLSamplerAddressMode::ClampToEdge, + super::SamplerAddressMode::Border => metal::MTLSamplerAddressMode::ClampToBorderColor, + super::SamplerAddressMode::MirrorOnce => metal::MTLSamplerAddressMode::MirrorClampToEdge, + } +} + +fn to_mtl_sampler_min_mag_filter(filter: super::SamplerFilter) -> metal::MTLSamplerMinMagFilter { + match filter { + super::SamplerFilter::Point => metal::MTLSamplerMinMagFilter::Nearest, + super::SamplerFilter::Linear | super::SamplerFilter::Anisotropic => metal::MTLSamplerMinMagFilter::Linear, + } +} + +fn to_mtl_sampler_mip_filter(filter: super::SamplerFilter) -> metal::MTLSamplerMipFilter { + match filter { + super::SamplerFilter::Point => metal::MTLSamplerMipFilter::Nearest, + super::SamplerFilter::Linear | super::SamplerFilter::Anisotropic => metal::MTLSamplerMipFilter::Linear, + } +} + +fn to_mtl_stencil_op(op: super::StencilOp) -> metal::MTLStencilOperation { + match op { + super::StencilOp::Keep => metal::MTLStencilOperation::Keep, + super::StencilOp::Zero => metal::MTLStencilOperation::Zero, + super::StencilOp::Replace => metal::MTLStencilOperation::Replace, + super::StencilOp::IncrSat => metal::MTLStencilOperation::IncrementClamp, + super::StencilOp::DecrSat => metal::MTLStencilOperation::DecrementClamp, + super::StencilOp::Invert => metal::MTLStencilOperation::Invert, + super::StencilOp::Incr => metal::MTLStencilOperation::IncrementWrap, + super::StencilOp::Decr => metal::MTLStencilOperation::DecrementWrap, + } +} + +fn has_stencil_component(format: metal::MTLPixelFormat) -> bool { + matches!(format, + metal::MTLPixelFormat::Depth32Float_Stencil8 + ) +} + +fn is_depth_format(format: metal::MTLPixelFormat) -> bool { + matches!(format, + metal::MTLPixelFormat::Depth32Float_Stencil8 + | metal::MTLPixelFormat::Depth32Float + | metal::MTLPixelFormat::Depth16Unorm + ) +} + +fn to_mtl_cull_mode(cull_mode: super::CullMode) -> metal::MTLCullMode { + match cull_mode { + super::CullMode::None => metal::MTLCullMode::None, + super::CullMode::Front => metal::MTLCullMode::Front, + super::CullMode::Back => metal::MTLCullMode::Back, + } +} + +fn to_mtl_winding(front_ccw: bool) -> metal::MTLWinding { + if front_ccw { + metal::MTLWinding::CounterClockwise + } else { + metal::MTLWinding::Clockwise + } +} + +fn to_mtl_triangle_fill_mode(fill_mode: super::FillMode) -> metal::MTLTriangleFillMode { + match fill_mode { + super::FillMode::Solid => metal::MTLTriangleFillMode::Fill, + super::FillMode::Wireframe => metal::MTLTriangleFillMode::Lines, + } +} + +fn to_mtl_index_type(stride: usize) -> metal::MTLIndexType { + match stride { + 2 => metal::MTLIndexType::UInt16, + 4 => metal::MTLIndexType::UInt32, + _ => panic!("Invalid index stride: {}, expected 2 or 4", stride), + } +} + +fn to_mtl_pixel_format(format: super::Format) -> metal::MTLPixelFormat { + match format { + super::Format::Unknown => metal::MTLPixelFormat::Invalid, + super::Format::R16n => metal::MTLPixelFormat::R16Unorm, + super::Format::R16u => metal::MTLPixelFormat::R16Uint, + super::Format::R16i => metal::MTLPixelFormat::R16Sint, + super::Format::R16f => metal::MTLPixelFormat::R16Float, + super::Format::R32u => metal::MTLPixelFormat::R32Uint, + super::Format::R32i => metal::MTLPixelFormat::R32Sint, + super::Format::R32f => metal::MTLPixelFormat::R32Float, + super::Format::RG16f => metal::MTLPixelFormat::RG16Float, + super::Format::RG16u => metal::MTLPixelFormat::RG16Uint, + super::Format::RG16i => metal::MTLPixelFormat::RG16Sint, + super::Format::RG32u => metal::MTLPixelFormat::RG32Uint, + super::Format::RG32i => metal::MTLPixelFormat::RG32Sint, + super::Format::RG32f => metal::MTLPixelFormat::RG32Float, + super::Format::RGB32u | + super::Format::RGB32i | + super::Format::RGB32f => panic!("hotline_rs::gfx::mtl RGB32 formats not supported in Metal"), + super::Format::RGBA8nSRGB => metal::MTLPixelFormat::RGBA8Unorm_sRGB, + super::Format::RGBA8n => metal::MTLPixelFormat::RGBA8Unorm, + super::Format::RGBA8u => metal::MTLPixelFormat::RGBA8Uint, + super::Format::RGBA8i => metal::MTLPixelFormat::RGBA8Sint, + super::Format::BGRA8n => metal::MTLPixelFormat::BGRA8Unorm, + super::Format::BGRX8n => metal::MTLPixelFormat::BGRA8Unorm, + super::Format::BGRA8nSRGB => metal::MTLPixelFormat::BGRA8Unorm_sRGB, + super::Format::BGRX8nSRGB => metal::MTLPixelFormat::BGRA8Unorm_sRGB, + super::Format::RGBA16u => metal::MTLPixelFormat::RGBA16Uint, + super::Format::RGBA16i => metal::MTLPixelFormat::RGBA16Sint, + super::Format::RGBA16f => metal::MTLPixelFormat::RGBA16Float, + super::Format::RGBA32u => metal::MTLPixelFormat::RGBA32Uint, + super::Format::RGBA32i => metal::MTLPixelFormat::RGBA32Sint, + super::Format::RGBA32f => metal::MTLPixelFormat::RGBA32Float, + super::Format::D32fS8X24u => metal::MTLPixelFormat::Depth32Float_Stencil8, + super::Format::D32f => metal::MTLPixelFormat::Depth32Float, + super::Format::D24nS8u => metal::MTLPixelFormat::Depth32Float_Stencil8, // D24S8 not supported on Apple Silicon + super::Format::D16n => metal::MTLPixelFormat::Depth16Unorm, + super::Format::BC1n => metal::MTLPixelFormat::BC1_RGBA, + super::Format::BC1nSRGB => metal::MTLPixelFormat::BC1_RGBA_sRGB, + super::Format::BC2n => metal::MTLPixelFormat::BC2_RGBA, + super::Format::BC2nSRGB => metal::MTLPixelFormat::BC2_RGBA_sRGB, + super::Format::BC3n => metal::MTLPixelFormat::BC3_RGBA, + super::Format::BC3nSRGB => metal::MTLPixelFormat::BC3_RGBA_sRGB, + super::Format::BC4n => metal::MTLPixelFormat::BC4_RUnorm, + super::Format::BC5n => metal::MTLPixelFormat::BC5_RGUnorm, + } +} + +fn to_mtl_data_type(resource_type: super::ResourceType) -> metal::MTLDataType { + match resource_type { + super::ResourceType::StructuredBuffer | + super::ResourceType::RWStructuredBuffer | + super::ResourceType::AppendStructuredBuffer | + super::ResourceType::ConsumeStructuredBuffer | + super::ResourceType::ConstantBuffer | + super::ResourceType::ByteAddressBuffer | + super::ResourceType::RWByteAddressBuffer | + super::ResourceType::Buffer => metal::MTLDataType::Pointer, + _ => metal::MTLDataType::Texture, // Texture2D, RWTexture2D, etc. + } +} + +// HLSL register kind ('t', 'u', 'b', 's') for a DescriptorType. Used to group bindings into MSL +// descriptor sets keyed by (kind, register_number) so e.g. t0 and u0 never share a [[buffer(N)]]. +fn descriptor_register_kind(ty: super::DescriptorType) -> char { + match ty { + super::DescriptorType::ShaderResource => 't', + super::DescriptorType::UnorderedAccess => 'u', + super::DescriptorType::ConstantBuffer | super::DescriptorType::PushConstants => 'b', + super::DescriptorType::Sampler => 's', + } +} + #[derive(Clone)] pub struct Device { metal_device: metal::Device, command_queue: metal::CommandQueue, shader_heap: Heap, - adapter_info: AdapterInfo + adapter_info: AdapterInfo, + heap_alloc_id: u16, + /// True when the GPU can sample timestamp counters at encoder stage boundaries, which lets us + /// take real per-encoder GPU timestamps. When false we fall back to MTLCommandBuffer's whole-CB + /// GPUStartTime / GPUEndTime (see timestamp_query / read_timestamps). + supports_stage_boundary_timestamps: bool, } +/// MTLCounterSamplingPoint::atStageBoundary — sampling at the boundary between encoder stages. +const MTL_COUNTER_SAMPLING_POINT_AT_STAGE_BOUNDARY: NSUInteger = 0; +/// MTLCounterDontSample sentinel: a stage index that should not record a timestamp. +const MTL_COUNTER_DONT_SAMPLE: NSUInteger = NSUInteger::MAX; + #[derive(Clone)] pub struct SwapChain { layer: metal::MetalLayer, @@ -97,6 +352,13 @@ pub struct SwapChain { backbuffer_texture: Texture, backbuffer_pass: RenderPass, backbuffer_pass_no_clear: RenderPass, + num_buffers: u32, + // GPU-side fence: present CB signals, each new CB waits — serialises GPU frames without blocking CPU + frame_event: metal::Event, + frame_value: u64, + // CPU-side ring: blocks the CPU only when num_buffers frames are already in flight, + // preventing shared-memory DynamicBuffer slots from being overwritten before the GPU is done + in_flight: std::sync::Arc>>, } impl super::SwapChain for SwapChain { @@ -104,19 +366,27 @@ impl super::SwapChain for SwapChain { } fn wait_for_last_frame(&self) { + let mut in_flight = self.in_flight.lock().unwrap(); + if in_flight.len() >= self.num_buffers as usize { + if let Some(oldest) = in_flight.pop_front() { + drop(in_flight); + oldest.wait_until_completed(); + } + } } fn get_num_buffers(&self) -> u32 { - 3 + self.num_buffers } fn get_frame_fence_value(&self) -> u64 { - 0 + self.frame_value } fn update(&mut self, device: &mut Device, window: &A::Window, cmd: &mut CmdBuf) -> bool { objc::rc::autoreleasepool(|| { let draw_size = window.get_size(); + self.layer.set_contents_scale(window.get_dpi_scale() as f64); self.layer.set_drawable_size(CGSize::new(draw_size.x as f64, draw_size.y as f64)); let drawable = self.layer.next_drawable() @@ -126,7 +396,11 @@ impl super::SwapChain for SwapChain { self.backbuffer_texture = Texture { metal_texture: drawable.texture().to_owned(), + resolved_texture: None, srv_index: None, + msaa_srv_index: None, + uav_index: None, + resolvable: false, heap_id: None }; @@ -164,32 +438,304 @@ impl super::SwapChain for SwapChain { fn swap(&mut self, device: &mut Device) { objc::rc::autoreleasepool(|| { - let cmd = device.command_queue.new_command_buffer(); + self.frame_value += 1; + let in_flight_count = self.in_flight.lock().unwrap().len(); + let cmd = device.command_queue.new_command_buffer().to_owned(); cmd.present_drawable(&self.drawable); + cmd.encode_signal_event(&self.frame_event, self.frame_value); cmd.commit(); + self.in_flight.lock().unwrap().push_back(cmd); }); } } -#[derive(Clone)] pub struct CmdBuf { cmd_queue: metal::CommandQueue, cmd: Option, render_encoder: Option, compute_encoder: Option, bound_index_buffer: Option, - bound_index_stride: usize + bound_index_stride: usize, + bound_render_pipeline: Option<*const RenderPipeline>, + bound_compute_pipeline: Option<*const ComputePipeline>, + metal_device: metal::Device, + transient_buffers: Vec, + vertex_binder: HashMap, + fragment_binder: HashMap, + compute_binder: HashMap, + deferred_ops: Vec, + pending_timestamp: Option<(metal::CounterSampleBuffer, NSUInteger)>, +} + +/// A unit of render-graph barrier work to replay each frame on Metal. Transition barriers are +/// absent because Metal tracks hazards automatically for resources in a `Tracked` heap. +#[derive(Clone)] +enum DeferredBarrierOp { + /// Resolve an MSAA texture into its single-sample resolve backing via a load/no-clear pass. + Resolve { + msaa: metal::Texture, + resolve: metal::Texture, + }, + /// Regenerate the mip chain of a sampled texture from mip 0 with a blit encoder. + GenerateMips { + texture: metal::Texture, + }, +} + +impl Clone for CmdBuf { + fn clone(&self) -> Self { + CmdBuf { + cmd_queue: self.cmd_queue.clone(), + cmd: self.cmd.clone(), + render_encoder: self.render_encoder.clone(), + compute_encoder: self.compute_encoder.clone(), + bound_index_buffer: self.bound_index_buffer.clone(), + bound_index_stride: self.bound_index_stride, + bound_render_pipeline: self.bound_render_pipeline, + bound_compute_pipeline: self.bound_compute_pipeline, + metal_device: self.metal_device.clone(), + transient_buffers: self.transient_buffers.clone(), + vertex_binder: self.vertex_binder.clone(), + fragment_binder: self.fragment_binder.clone(), + compute_binder: self.compute_binder.clone(), + deferred_ops: self.deferred_ops.clone(), + pending_timestamp: self.pending_timestamp.clone(), + } + } +} + +impl CmdBuf { + fn allocate_stage_bindings( + &mut self, + binder: &HashMap, + stage: super::ShaderType, + ) { + let encoder = match self.render_encoder.as_ref() { + Some(e) => e, + None => return, + }; + + // Bind push constants using setVertexBytes/setFragmentBytes (zero allocations) + // Skip binders that have not changed since the last draw + for b in binder.values() { + if let PipelineStageBinder::PushConstants(pc) = b { + if !pc.dirty { + continue; + } + let data_size = (pc.num_32_bit_constants * 4) as u64; + let data_ptr = pc.data.as_ptr() as *const std::ffi::c_void; + + match stage { + super::ShaderType::Vertex => { + encoder.set_vertex_bytes(pc.buffer_index as u64, data_size, data_ptr); + } + super::ShaderType::Fragment => { + encoder.set_fragment_bytes(pc.buffer_index as u64, data_size, data_ptr); + } + _ => unimplemented!(), + } + } + } + + // Group resource bindings by buffer_index, skipping groups where nothing is dirty + let mut groups: HashMap> = HashMap::new(); + for b in binder.values() { + if let PipelineStageBinder::Resource(rb) = b { + if rb.bound_resource.is_some() && rb.dirty { + groups.entry(rb.buffer_index).or_default().push(rb); + } + } + } + + // TODO: move to to_mtl_stage function + implement the others + let render_stage = match stage { + super::ShaderType::Vertex => metal::MTLRenderStages::Vertex, + super::ShaderType::Fragment => metal::MTLRenderStages::Fragment, + _ => unimplemented!(), + }; + + // Allocate resource bindings (grouped by buffer_index) + for (buffer_index, mut binders) in groups { + // Sort by binding_index to ensure deterministic order + binders.sort_by_key(|rb| rb.binding_index); + + let arg_descs: Vec = binders.iter().map(|rb| { + let arg_desc = metal::ArgumentDescriptor::new(); + arg_desc.set_index(rb.binding_index as u64); + arg_desc.set_data_type(rb.data_type); + arg_desc.set_array_length(rb.array_length); + arg_desc.set_access(metal::MTLArgumentAccess::ReadOnly); + arg_desc.to_owned() + }).collect(); + + let arg_encoder = self.metal_device.new_argument_encoder( + metal::Array::from_owned_slice(&arg_descs) + ); + let arg_buffer = self.metal_device.new_buffer( + arg_encoder.encoded_length(), + metal::MTLResourceOptions::StorageModeShared + ); + arg_encoder.set_argument_buffer(&arg_buffer, 0); + + for rb in &binders { + if let Some(ref binding) = rb.bound_resource { + let heap = unsafe { &*binding.heap_ptr }; + encoder.use_heap_at(&heap.mtl_heap, render_stage); + + match rb.data_type { + metal::MTLDataType::Texture => { + if let Some(texture) = heap.texture_slots.get(binding.offset).and_then(|t| t.as_ref()) { + arg_encoder.set_texture(rb.binding_index as u64, texture); + } + } + metal::MTLDataType::Pointer => { + if let Some(buffer) = heap.buffer_slots.get(binding.offset).and_then(|b| b.as_ref()) { + arg_encoder.set_buffer(rb.binding_index as u64, buffer, 0); + } + } + _ => {} + } + } + } + + match stage { + super::ShaderType::Vertex => encoder.set_vertex_buffer(buffer_index as u64, Some(&arg_buffer), 0), + super::ShaderType::Fragment => encoder.set_fragment_buffer(buffer_index as u64, Some(&arg_buffer), 0), + _ => unimplemented!(), + } + self.transient_buffers.push(arg_buffer); + } + } + + fn allocate_stage_resources(&mut self) { + let vertex_binder = self.vertex_binder.clone(); + let fragment_binder = self.fragment_binder.clone(); + + self.allocate_stage_bindings(&vertex_binder, super::ShaderType::Vertex); + self.allocate_stage_bindings(&fragment_binder, super::ShaderType::Fragment); + + // Clear dirty flags on originals now that encoding is done + for b in self.vertex_binder.values_mut() { + match b { + PipelineStageBinder::PushConstants(pc) => pc.dirty = false, + PipelineStageBinder::Resource(rb) => rb.dirty = false, + } + } + for b in self.fragment_binder.values_mut() { + match b { + PipelineStageBinder::PushConstants(pc) => pc.dirty = false, + PipelineStageBinder::Resource(rb) => rb.dirty = false, + } + } + } + + /// Flush the compute binder state onto the active compute encoder before a dispatch. Push + /// constants go via setBytes; explicitly bound (non-bindless) resources are encoded into a + /// transient argument buffer - bindless heap argument buffers are already bound by `set_heap`. + fn allocate_compute_resources(&mut self) { + let encoder = match self.compute_encoder.as_ref() { + Some(e) => e, + None => return, + }; + let binder = self.compute_binder.clone(); + + // push constants + for b in binder.values() { + if let PipelineStageBinder::PushConstants(pc) = b { + if !pc.dirty { + continue; + } + let data_size = (pc.num_32_bit_constants * 4) as u64; + let data_ptr = pc.data.as_ptr() as *const std::ffi::c_void; + encoder.set_bytes(pc.buffer_index as u64, data_size, data_ptr); + } + } + + // explicitly bound resources, grouped by buffer_index + let mut groups: HashMap> = HashMap::new(); + for b in binder.values() { + if let PipelineStageBinder::Resource(rb) = b { + if rb.bound_resource.is_some() && rb.dirty { + groups.entry(rb.buffer_index).or_default().push(rb); + } + } + } + + for (buffer_index, mut binders) in groups { + binders.sort_by_key(|rb| rb.binding_index); + + let arg_descs: Vec = binders.iter().map(|rb| { + let arg_desc = metal::ArgumentDescriptor::new(); + arg_desc.set_index(rb.binding_index as u64); + arg_desc.set_data_type(rb.data_type); + arg_desc.set_array_length(rb.array_length); + arg_desc.set_access(metal::MTLArgumentAccess::ReadWrite); + arg_desc.to_owned() + }).collect(); + + let arg_encoder = self.metal_device.new_argument_encoder( + metal::Array::from_owned_slice(&arg_descs) + ); + let arg_buffer = self.metal_device.new_buffer( + arg_encoder.encoded_length(), + metal::MTLResourceOptions::StorageModeShared + ); + arg_encoder.set_argument_buffer(&arg_buffer, 0); + + for rb in &binders { + if let Some(ref binding) = rb.bound_resource { + let heap = unsafe { &*binding.heap_ptr }; + encoder.use_heap(&heap.mtl_heap); + + match rb.data_type { + metal::MTLDataType::Texture => { + if let Some(texture) = heap.texture_slots.get(binding.offset).and_then(|t| t.as_ref()) { + arg_encoder.set_texture(rb.binding_index as u64, texture); + } + } + metal::MTLDataType::Pointer => { + if let Some(buffer) = heap.buffer_slots.get(binding.offset).and_then(|b| b.as_ref()) { + arg_encoder.set_buffer(rb.binding_index as u64, buffer, 0); + } + } + _ => {} + } + } + } + + encoder.set_buffer(buffer_index as u64, Some(&arg_buffer), 0); + self.transient_buffers.push(arg_buffer); + } + + // clear dirty flags now encoding is done + for b in self.compute_binder.values_mut() { + match b { + PipelineStageBinder::PushConstants(pc) => pc.dirty = false, + PipelineStageBinder::Resource(rb) => rb.dirty = false, + } + } + } } impl super::CmdBuf for CmdBuf { fn reset(&mut self, swap_chain: &SwapChain) { objc::rc::autoreleasepool(|| { - self.cmd = Some(self.cmd_queue.new_command_buffer().to_owned()); + let cmd = self.cmd_queue.new_command_buffer().to_owned(); + // GPU waits for the previous frame's present signal before executing any work + if swap_chain.frame_value > 0 { + cmd.encode_wait_for_event(&swap_chain.frame_event, swap_chain.frame_value); + } + self.cmd = Some(cmd); + self.transient_buffers.clear(); }); } fn close(&mut self) -> result::Result<(), super::Error> { objc::rc::autoreleasepool(|| { + // close any open compute encoder before committing + if let Some(enc) = self.compute_encoder.take() { + enc.end_encoding(); + } self.cmd.as_ref().expect("hotline_rs::gfx::mtl expected call to CmdBuf::reset before close").commit(); self.cmd = None; Ok(()) @@ -206,6 +752,24 @@ impl super::CmdBuf for CmdBuf { assert!(self.render_encoder.is_none(), "hotline_rs::gfx::mtl begin_render_pass called without matching CmdBuf::end_render_pass"); + // close any open compute encoder - Metal forbids two live encoders on one cmd buffer + if let Some(enc) = self.compute_encoder.take() { + enc.end_encoding(); + } + + // if a timestamp pair is armed, sample the GPU clock at this encoder's stage boundaries: + // start_of_vertex = pass start, end_of_fragment = pass end (the inner boundaries are left + // as MTLCounterDontSample). Set on the descriptor before the encoder is created. + if let Some((sample_buffer, start)) = self.pending_timestamp.take() { + if let Some(att) = render_pass.desc.sample_buffer_attachments().object_at(0) { + att.set_sample_buffer(&sample_buffer); + att.set_start_of_vertex_sample_index(start); + att.set_end_of_vertex_sample_index(MTL_COUNTER_DONT_SAMPLE); + att.set_start_of_fragment_sample_index(MTL_COUNTER_DONT_SAMPLE); + att.set_end_of_fragment_sample_index(start + 1); + } + } + // catch mismatched close/reset let render_encoder = self.cmd.as_ref() .expect("hotline_rs::gfx::mtl expected call to CmdBuf::reset after close") @@ -226,12 +790,44 @@ impl super::CmdBuf for CmdBuf { } fn begin_event(&mut self, colour: u32, name: &str) { + if let Some(enc) = self.render_encoder.as_ref() { + enc.push_debug_group(name); + } else if let Some(enc) = self.compute_encoder.as_ref() { + enc.push_debug_group(name); + } else if let Some(cmd) = self.cmd.as_ref() { + cmd.push_debug_group(name); + } } fn end_event(&mut self) { + if let Some(enc) = self.render_encoder.as_ref() { + enc.pop_debug_group(); + } else if let Some(enc) = self.compute_encoder.as_ref() { + enc.pop_debug_group(); + } else if let Some(cmd) = self.cmd.as_ref() { + cmd.pop_debug_group(); + } } fn timestamp_query(&mut self, heap: &mut QueryHeap, resolve_buffer: &mut Buffer) { + let idx = heap.alloc_index; + heap.alloc_index += 1; + resolve_buffer.counter_sample_index = idx; + resolve_buffer.counter_cmd = self.cmd.clone(); + + if let Some(sample_buffer) = heap.sample_buffer.as_ref() { + // counter-sampling path: tag the buffer so read_timestamps resolves the counter, and on + // the start sample of a pair arm the next encoder to record both stage-boundary samples + // ([idx, idx + 1]). Multiple start/end pairs in one CB each arm their own encoder. + resolve_buffer.counter_sample_buffer = Some(sample_buffer.to_owned()); + if idx % 2 == 0 { + self.pending_timestamp = Some((sample_buffer.to_owned(), idx as NSUInteger)); + } + } + else { + // fallback path: no counter buffer, read GPUStartTime / GPUEndTime of the pass CB. + resolve_buffer.counter_sample_buffer = None; + } } fn begin_query(&mut self, heap: &mut QueryHeap, query_type: QueryType) -> usize { @@ -251,7 +847,7 @@ impl super::CmdBuf for CmdBuf { unimplemented!() } - fn set_viewport(&self, viewport: &super::Viewport) { + fn set_viewport(&mut self, viewport: &super::Viewport) { objc::rc::autoreleasepool(|| { self.render_encoder .as_ref() @@ -267,7 +863,7 @@ impl super::CmdBuf for CmdBuf { }); } - fn set_scissor_rect(&self, scissor_rect: &super::ScissorRect) { + fn set_scissor_rect(&mut self, scissor_rect: &super::ScissorRect) { objc::rc::autoreleasepool(|| { self.render_encoder .as_ref() @@ -281,7 +877,7 @@ impl super::CmdBuf for CmdBuf { }); } - fn set_vertex_buffer(&self, buffer: &Buffer, slot: u32) { + fn set_vertex_buffer(&mut self, buffer: &Buffer, slot: u32) { objc::rc::autoreleasepool(|| { self.render_encoder .as_ref() @@ -295,126 +891,274 @@ impl super::CmdBuf for CmdBuf { self.bound_index_stride = buffer.element_stride; } - fn set_render_pipeline(&self, pipeline: &RenderPipeline) { + fn set_render_pipeline(&mut self, pipeline: &RenderPipeline) { objc::rc::autoreleasepool(|| { - self.render_encoder + let encoder = self.render_encoder .as_ref() - .expect("hotline_rs::gfx::metal expected a call to begin render pass before using render commands") - .set_render_pipeline_state(&pipeline.pipeline_state); + .expect("hotline_rs::gfx::metal expected a call to begin render pass before using render commands"); - // bind static samplers - for sampler in &pipeline.static_samplers { - self.render_encoder.as_ref().unwrap().set_fragment_sampler_state( - sampler.slot as u64, Some(&sampler.sampler)) + encoder.set_render_pipeline_state(&pipeline.pipeline_state); + + // Set depth stencil state + encoder.set_depth_stencil_state(&pipeline.depth_stencil_state); + + // Set rasterizer state + let raster = &pipeline.raster_info; + encoder.set_cull_mode(to_mtl_cull_mode(raster.cull_mode)); + encoder.set_front_facing_winding(to_mtl_winding(raster.front_ccw)); + encoder.set_triangle_fill_mode(to_mtl_triangle_fill_mode(raster.fill_mode)); + encoder.set_depth_bias(raster.depth_bias as f32, raster.slope_scaled_depth_bias, raster.depth_bias_clamp); + + // Bind sampler argument buffer at buffer(0) in fragment shader + if let Some(ref sampler_arg_buffer) = pipeline.sampler_argument_buffer { + encoder.set_fragment_buffer( + 0, + Some(sampler_arg_buffer), + 0 + ); } + + // store pipeline pointer for push_render_constants + self.bound_render_pipeline = Some(pipeline as *const RenderPipeline); + + // Clone binder templates from pipeline to command buffer + self.vertex_binder = pipeline.vertex_binder.clone(); + self.fragment_binder = pipeline.fragment_binder.clone(); }); } - fn set_compute_pipeline(&self, pipeline: &ComputePipeline) { + fn set_compute_pipeline(&mut self, pipeline: &ComputePipeline) { + objc::rc::autoreleasepool(|| { + // open a compute encoder lazily; reused across dispatches until a render pass or close + if self.compute_encoder.is_none() { + let cmd = self.cmd.as_ref() + .expect("hotline_rs::gfx::mtl expected a call to CmdBuf::reset before set_compute_pipeline"); + // if a timestamp pair is armed, sample the GPU clock at this encoder's boundaries + let encoder = if let Some((sample_buffer, start)) = self.pending_timestamp.take() { + let desc = metal::ComputePassDescriptor::new(); + if let Some(att) = desc.sample_buffer_attachments().object_at(0) { + att.set_sample_buffer(&sample_buffer); + att.set_start_of_encoder_sample_index(start); + att.set_end_of_encoder_sample_index(start + 1); + } + cmd.compute_command_encoder_with_descriptor(desc).to_owned() + } + else { + cmd.new_compute_command_encoder().to_owned() + }; + self.compute_encoder = Some(encoder); + } + + self.compute_encoder.as_ref().unwrap() + .set_compute_pipeline_state(&pipeline.pipeline_state); + // store pipeline pointer and clone binder template into command buffer state + self.bound_compute_pipeline = Some(pipeline as *const ComputePipeline); + self.compute_binder = pipeline.compute_binder.clone(); + }); } - fn set_raytracing_pipeline(&self, pipeline: &RaytracingPipeline) { + fn set_raytracing_pipeline(&mut self, pipeline: &RaytracingPipeline) { unimplemented!() } - fn set_heap(&self, pipeline: &T, heap: &Heap) { - - } + fn set_heap(&mut self, pipeline: &T, heap: &Heap) { + // compute pipelines bind the heap argument buffers on the compute encoder (single stage) + if matches!(T::get_pipeline_type(), super::PipelineType::Compute) { + let encoder = self.compute_encoder + .as_ref() + .expect("hotline_rs::gfx::metal expected a call to set_compute_pipeline before set_heap"); + let cp: &ComputePipeline = unsafe { std::mem::transmute(pipeline) }; + + encoder.use_heap(&heap.mtl_heap); + // Structured buffers are device-allocated (not part of mtl_heap), so use_heap does not + // make them resident - they are reached indirectly through the bindless buffer argument + // buffer, so without this the GPU can read unmapped memory. Textures live in mtl_heap + // and are covered by use_heap above. + for buffer in heap.buffer_slots.iter().flatten() { + encoder.use_resource(buffer, metal::MTLResourceUsage::Read | metal::MTLResourceUsage::Write); + } + for (_key, slot) in &cp.compute_binder { + if let PipelineStageBinder::Resource(res) = slot { + let arg_buffer = match res.data_type { + metal::MTLDataType::Texture => heap.get_texture_argument_buffer(), + metal::MTLDataType::Pointer => heap.get_buffer_argument_buffer(), + _ => continue, + }; + encoder.set_buffer(res.buffer_index as u64, Some(arg_buffer), 0); + } + } + return; + } - fn set_heap_render(&self, pipeline: &RenderPipeline, heap: &Heap) { - self.render_encoder + let encoder = self.render_encoder .as_ref() - .expect("hotline_rs::gfx::metal expected a call to begin render pass before using render commands") - .use_heap_at(&heap.mtl_heap, metal::MTLRenderStages::Fragment); - - - pipeline.fragment_descriptor_slots.iter().enumerate().for_each(|(slot_index, slot)| { - if let Some(slot) = slot { - slot.argument_encoder.set_argument_buffer(&slot.argument_buffer, 0); + .expect("hotline_rs::gfx::metal expected a call to begin render pass before using render commands"); + + // Cast pipeline to RenderPipeline to access slot_lookup + let rp: &RenderPipeline = unsafe { std::mem::transmute(pipeline) }; + + // Structured buffers are device-allocated (not part of mtl_heap), so use_heap_at does not + // make them resident - they are reached indirectly through the bindless buffer argument + // buffer, so without this the GPU can read unmapped memory. Textures live in mtl_heap and + // are covered by use_heap_at below. + for buffer in heap.buffer_slots.iter().flatten() { + encoder.use_resource_at( + buffer, + metal::MTLResourceUsage::Read | metal::MTLResourceUsage::Write, + metal::MTLRenderStages::Vertex | metal::MTLRenderStages::Fragment, + ); + } - // TODO: need to know data types (Texture, Buffer) - // assign textures to slots - heap.texture_slots.iter().enumerate().for_each(|(index, texture)| { - if let Some(texture) = texture { - slot.argument_encoder.set_texture(index as u64, texture); - } - }); + // vertex bindings + encoder.use_heap_at(&heap.mtl_heap, metal::MTLRenderStages::Vertex); + for (key, slot) in &rp.vertex_binder { + match slot { + PipelineStageBinder::Resource(res) => { + let arg_buffer = match res.data_type { + metal::MTLDataType::Texture => heap.get_texture_argument_buffer(), + metal::MTLDataType::Pointer => heap.get_buffer_argument_buffer(), + _ => continue, + }; + encoder.set_vertex_buffer(res.buffer_index as u64, Some(arg_buffer), 0); + } + _ => {} + } + } - // TODO: need to know what stages to bind on - self.render_encoder.as_ref().unwrap().set_fragment_buffer(slot_index as u64, Some(&slot.argument_buffer), 0); + // fragment bindings + encoder.use_heap_at(&heap.mtl_heap, metal::MTLRenderStages::Fragment); + for (key, slot) in &rp.fragment_binder { + match slot { + PipelineStageBinder::Resource(res) => { + let arg_buffer = match res.data_type { + metal::MTLDataType::Texture => heap.get_texture_argument_buffer(), + metal::MTLDataType::Pointer => heap.get_buffer_argument_buffer(), + _ => continue, + }; + encoder.set_fragment_buffer(res.buffer_index as u64, Some(arg_buffer), 0); + } + _ => {} } - }); + } + } - /* - pipeline.vertex_descriptor_slots.iter().enumerate().for_each(|(slot_index, slot)| { - if let Some(slot) = slot { - slot.argument_encoder.set_argument_buffer(&slot.argument_buffer, 0); + fn set_binding(&mut self, _pipeline: &T, register: u32, space: u32, descriptor_type: super::DescriptorType, heap: &Heap, offset: usize) -> Option<()> { + let key: SlotKey = (register, space, descriptor_type); + let heap_ptr = heap as *const Heap; - // TODO: need to know data types (Texture, Buffer) - // assign textures to slots - heap.texture_slots.iter().enumerate().for_each(|(index, texture)| { - slot.argument_encoder.set_texture(index as u64, texture); - }); + // Write to vertex binder if present + if let Some(binder) = self.vertex_binder.get_mut(&key) { + if let PipelineStageBinder::Resource(ref mut rb) = binder { + rb.bound_resource = Some(ResourceBinding { heap_ptr, offset }); + rb.dirty = true; + } + } - // TODO: need to know what stages to bind on - self.render_encoder.as_ref().unwrap().set_vertex_buffer(slot_index as u64, Some(&slot.argument_buffer), 0); + // Write to fragment binder if present + if let Some(binder) = self.fragment_binder.get_mut(&key) { + if let PipelineStageBinder::Resource(ref mut rb) = binder { + rb.bound_resource = Some(ResourceBinding { heap_ptr, offset }); + rb.dirty = true; } - }); - */ - } - - // TODO: needs stage - fn set_binding(&self, pipeline: &T, register: u32, space: u32, descriptor_type: super::DescriptorType, heap: &Heap, offset: usize) -> Option<()> { - let slot = pipeline.get_pipeline_slot(register, space, descriptor_type)?; - let rp : &RenderPipeline = unsafe { std::mem::transmute(pipeline) }; - if rp.fragment_descriptor_slots.len() > 0 { - if let Some(d) = rp.fragment_descriptor_slots[0].as_ref() { - d.argument_encoder.set_argument_buffer(&d.argument_buffer, 0); - if let Some(texture) = heap.texture_slots[offset].as_ref() { - d.argument_encoder.set_texture(slot.index as u64, &texture); - } + } + + // Write to compute binder if present + if let Some(binder) = self.compute_binder.get_mut(&key) { + if let PipelineStageBinder::Resource(ref mut rb) = binder { + rb.bound_resource = Some(ResourceBinding { heap_ptr, offset }); + rb.dirty = true; } } - Some(()) - } - fn set_texture(&mut self, texture: &Texture, slot: u32) { - self.render_encoder.as_ref().unwrap().set_fragment_texture(slot as u64, Some(&texture.metal_texture)); + Some(()) } fn set_marker(&mut self, colour: u32, name: &str) { } - fn push_render_constants(&mut self, pipeline: &P, register: u32, space: u32, num_values: u32, dest_offset: u32, data: &[T]) -> Option<()> { - let _slot = pipeline.get_pipeline_slot(register, space, super::DescriptorType::PushConstants)?; - // TODO: need to know the stages and the buffer offset - self.render_encoder - .as_ref() - .expect("hotline_rs::gfx::metal expected a call to begin render pass before using render commands") - // .set_fragment_bytes(slot.index as u64 + 2, num_values as u64 * 4, data.as_ptr() as _); - // .set_vertex_bytes(slot.index as u64 + 2, num_values as u64 * 4, data.as_ptr() as _); - .set_vertex_bytes(1, num_values as u64 * 4, data.as_ptr() as _); - Some(()) + fn push_render_constants(&mut self, _pipeline: &P, register: u32, space: u32, num_values: u32, dest_offset: u32, data: &[T]) -> Option<()> { + let key = (register, space, super::DescriptorType::PushConstants); + + let data_size_dwords = num_values as usize; + let data_u32 = unsafe { + std::slice::from_raw_parts( + data.as_ptr() as *const u32, + data_size_dwords + ) + }; + + let mut result = None; + + // Write to vertex binder if matching key found + if let Some(PipelineStageBinder::PushConstants(ref mut pc)) = self.vertex_binder.get_mut(&key) { + let dest_start = dest_offset as usize; + let dest_end = dest_start + data_size_dwords; + if dest_end <= pc.data.len() { + pc.data[dest_start..dest_end].copy_from_slice(data_u32); + pc.dirty = true; + } + result = Some(()); + } + + // Write to fragment binder if matching key found + if let Some(PipelineStageBinder::PushConstants(ref mut pc)) = self.fragment_binder.get_mut(&key) { + let dest_start = dest_offset as usize; + let dest_end = dest_start + data_size_dwords; + if dest_end <= pc.data.len() { + pc.data[dest_start..dest_end].copy_from_slice(data_u32); + pc.dirty = true; + } + result = Some(()); + } + + result } - fn push_compute_constants(&mut self, _pipeline: &P, _register: u32, _space: u32, _num_values: u32, _dest_offset: u32, _data: &[T]) -> Option<()> { + fn push_compute_constants(&mut self, _pipeline: &P, register: u32, space: u32, num_values: u32, dest_offset: u32, data: &[T]) -> Option<()> { + let key = (register, space, super::DescriptorType::PushConstants); + + let data_size_dwords = num_values as usize; + let data_u32 = unsafe { + std::slice::from_raw_parts( + data.as_ptr() as *const u32, + data_size_dwords + ) + }; + + if let Some(PipelineStageBinder::PushConstants(ref mut pc)) = self.compute_binder.get_mut(&key) { + let dest_start = dest_offset as usize; + let dest_end = dest_start + data_size_dwords; + if dest_end <= pc.data.len() { + pc.data[dest_start..dest_end].copy_from_slice(data_u32); + pc.dirty = true; + } + return Some(()); + } + None } fn draw_instanced( - &self, + &mut self, vertex_count: u32, instance_count: u32, start_vertex: u32, start_instance: u32, ) { objc::rc::autoreleasepool(|| { + self.allocate_stage_resources(); + + let primitive_type = self.bound_render_pipeline + .map(|p| unsafe { (*p).topology }) + .map(to_mtl_primitive_type) + .unwrap_or(metal::MTLPrimitiveType::Triangle); + self.render_encoder .as_ref() .expect("hotline_rs::gfx::metal expected a call to begin render pass before using render commands") .draw_primitives_instanced_base_instance( - metal::MTLPrimitiveType::TriangleStrip, + primitive_type, start_vertex as u64, vertex_count as u64, instance_count as u64, @@ -424,7 +1168,7 @@ impl super::CmdBuf for CmdBuf { } fn draw_indexed_instanced( - &self, + &mut self, index_count: u32, instance_count: u32, start_index: u32, @@ -432,13 +1176,20 @@ impl super::CmdBuf for CmdBuf { start_instance: u32, ) { objc::rc::autoreleasepool(|| { + self.allocate_stage_resources(); + + let primitive_type = self.bound_render_pipeline + .map(|p| unsafe { (*p).topology }) + .map(to_mtl_primitive_type) + .unwrap_or(metal::MTLPrimitiveType::Triangle); + self.render_encoder .as_ref() .expect("hotline_rs::gfx::metal expected a call to begin render pass before using render commands") .draw_indexed_primitives_instanced_base_instance( - metal::MTLPrimitiveType::TriangleStrip, + primitive_type, index_count as u64, - metal::MTLIndexType::UInt16, + to_mtl_index_type(self.bound_index_stride), &self.bound_index_buffer.as_ref().unwrap(), start_index as u64 * self.bound_index_stride as u64, instance_count as u64, @@ -448,11 +1199,24 @@ impl super::CmdBuf for CmdBuf { }) } - fn dispatch(&self, group_count: Size3, _numthreads: Size3) { + fn dispatch(&mut self, group_count: Size3, numthreads: Size3) { + objc::rc::autoreleasepool(|| { + self.allocate_compute_resources(); + + let threadgroups = metal::MTLSize::new( + group_count.x as u64, group_count.y as u64, group_count.z as u64); + let threads_per_group = metal::MTLSize::new( + numthreads.x as u64, numthreads.y as u64, numthreads.z as u64); + + self.compute_encoder + .as_ref() + .expect("hotline_rs::gfx::metal expected a call to set_compute_pipeline before dispatch") + .dispatch_thread_groups(threadgroups, threads_per_group); + }); } fn execute_indirect( - &self, + &mut self, command: &CommandSignature, max_command_count: u32, argument_buffer: &Buffer, @@ -468,11 +1232,28 @@ impl super::CmdBuf for CmdBuf { }) } - fn resolve_texture_subresource(&self, texture: &Texture, subresource: u32) -> result::Result<(), super::Error> { + fn resolve_texture_subresource(&mut self, texture: &Texture, _subresource: u32) -> result::Result<(), super::Error> { + // Record the resolve as deferred barrier work; Device::execute replays it into a fresh + // command buffer each frame (a committed Metal command buffer can't be re-submitted). + if let Some(resolve) = texture.resolved_texture.as_ref() { + self.deferred_ops.push(DeferredBarrierOp::Resolve { + msaa: texture.metal_texture.to_owned(), + resolve: resolve.to_owned(), + }); + } Ok(()) } - fn generate_mip_maps(&mut self, texture: &Texture, device: &Device, heap: &Heap) -> result::Result<(), super::Error> { + fn generate_mip_maps(&mut self, texture: &Texture, _device: &Device, _heap: &Heap) -> result::Result<(), super::Error> { + // Record mip generation as deferred barrier work (replayed per-frame by Device::execute). + // Generate on the texture shaders actually sample: the resolve backing for an MSAA target + // (its mip 0 is filled by the preceding resolve op), otherwise the texture itself. + let target = texture.resolved_texture.as_ref().unwrap_or(&texture.metal_texture); + if target.mipmap_level_count() > 1 { + self.deferred_ops.push(DeferredBarrierOp::GenerateMips { + texture: target.to_owned(), + }); + } Ok(()) } @@ -498,7 +1279,7 @@ impl super::CmdBuf for CmdBuf { ) { } - fn dispatch_rays(&self, sbt: &RaytracingShaderBindingTable, numthreads: Size3) { + fn dispatch_rays(&mut self, sbt: &RaytracingShaderBindingTable, numthreads: Size3) { unimplemented!() } @@ -507,34 +1288,44 @@ impl super::CmdBuf for CmdBuf { } } +#[derive(Clone)] pub struct Buffer { metal_buffer: metal::Buffer, - element_stride: usize + element_stride: usize, + srv_index: Option, + uav_index: Option, + cbv_index: Option, + counter_sample_buffer: Option, + counter_sample_index: usize, + // Metal substitute for a D3D12 GPU fence: wait_until_completed before resolving counter data + counter_cmd: Option, } impl super::Buffer for Buffer { fn update(&mut self, offset: usize, data: &[T]) -> result::Result<(), super::Error> { unsafe { - let data_ptr = self.metal_buffer.contents(); - std::ptr::copy_nonoverlapping(data.as_ptr() as *mut u8, data_ptr as *mut u8, data.len()); + let data_ptr = self.metal_buffer.contents() as *mut u8; + let dest_ptr = data_ptr.add(offset); + let byte_len = data.len() * std::mem::size_of::(); + std::ptr::copy_nonoverlapping(data.as_ptr() as *const u8, dest_ptr, byte_len); } Ok(()) } fn write(&mut self, offset: usize, data: &[T]) -> result::Result<(), super::Error> { - Ok(()) + self.update(offset, data) } fn get_cbv_index(&self) -> Option { - None + self.cbv_index } fn get_srv_index(&self) -> Option { - None + self.srv_index } fn get_uav_index(&self) -> Option { - None + self.uav_index } fn get_vbv(&self) -> Option { @@ -570,62 +1361,66 @@ struct MetalSamplerBinding { sampler: metal::SamplerState } +/// Push constants binder - uses setVertexBytes/setFragmentBytes for zero-allocation binding #[derive(Clone)] -pub struct DescriptorMember { - offset: u32, - num: u32, - info: PipelineSlotInfo +struct PushConstantsBinder { + pub data: Vec, + pub num_32_bit_constants: u32, + pub buffer_index: u32, + pub dirty: bool, +} + +#[derive(Clone, Copy)] +struct ResourceBinding { + pub heap_ptr: *const Heap, + pub offset: usize, } -type DescriptorMemberArray = Vec>; #[derive(Clone)] -pub struct DescriptorSlot { - argument_buffer: metal::Buffer, - argument_encoder: metal::ArgumentEncoder, - members: Vec>, +struct ResourceBinder { + pub buffer_index: u32, + pub binding_index: u32, + pub data_type: metal::MTLDataType, + pub array_length: u64, + pub bound_resource: Option, + pub dirty: bool, } -type DescriptorSlotArray = Vec>; -pub struct PushConstantSlot { - buffer: metal::Buffer, - slot: u32, - visibility: ShaderVisibility +#[derive(Clone)] +enum PipelineStageBinder { + PushConstants(PushConstantsBinder), + Resource(ResourceBinder), } + +/// Key for slot lookup: (register, space, descriptor_type) +type SlotKey = (u32, u32, DescriptorType); + pub struct RenderPipeline { pipeline_state: metal::RenderPipelineState, - static_samplers: Vec, slots: Vec, - vertex_descriptor_slots: DescriptorSlotArray, - fragment_descriptor_slots: DescriptorSlotArray, - vertex_push_constant_slots: Vec, - fragment_push_constant_slots: Vec + /// Primitive topology for draw calls + topology: Topology, + /// Unified slot lookup by (register, space, descriptor_type) + slot_lookup: HashMap, + /// Static samplers + static_samplers: Vec, + /// Sampler argument buffer + sampler_argument_buffer: Option, + /// Vertex stage binders for push constants, keyed by (register, space, descriptor_type) + vertex_binder: HashMap, + /// Fragment stage binders for push constants, keyed by (register, space, descriptor_type) + fragment_binder: HashMap, + /// Depth stencil state + depth_stencil_state: metal::DepthStencilState, + /// Rasterizer state (applied dynamically on encoder in Metal) + raster_info: super::RasterInfo, } impl super::RenderPipeline for RenderPipeline {} impl super::Pipeline for RenderPipeline { fn get_pipeline_slot(&self, register: u32, space: u32, descriptor_type: DescriptorType) -> Option<&super::PipelineSlotInfo> { - if (space as usize) < self.fragment_descriptor_slots.len() { - if let Some(set) = self.fragment_descriptor_slots[space as usize].as_ref() { - if (register as usize) < set.members.len() { - if let Some(member) = set.members[(register as usize)].as_ref() { - Some(&member.info) - } - else { - None - } - } - else { - None - } - } - else { - None - } - } - else { - None - } + self.slot_lookup.get(&(register, space, descriptor_type)) } fn get_pipeline_slots(&self) -> &Vec { @@ -640,7 +1435,15 @@ impl super::Pipeline for RenderPipeline { #[derive(Clone)] pub struct Texture { metal_texture: metal::Texture, + /// Single-sample resolve backing for an MSAA texture (samples > 1); also the texture sampled + /// when reading a resolvable target normally + resolved_texture: Option, + /// Bindless index of the resolved / non-MSAA view (returned by `get_srv_index`) srv_index: Option, + /// Bindless index of the MSAA view, for `Texture2DMS` reads (returned by `get_msaa_srv_index`) + msaa_srv_index: Option, + uav_index: Option, + resolvable: bool, heap_id: Option } @@ -654,23 +1457,27 @@ impl super::Texture for Texture { } fn get_msaa_srv_index(&self) -> Option { - None + self.msaa_srv_index } fn get_uav_index(&self) -> Option { - None + self.uav_index } fn clone_inner(&self) -> Texture { Texture { metal_texture: self.metal_texture.clone(), + resolved_texture: self.resolved_texture.clone(), srv_index: self.srv_index, + msaa_srv_index: self.msaa_srv_index, + uav_index: self.uav_index, + resolvable: self.resolvable, heap_id: self.heap_id } } fn is_resolvable(&self) -> bool { - false + self.resolvable } fn get_shader_heap_id(&self) -> Option { @@ -706,7 +1513,12 @@ impl super::ReadBackRequest for ReadBackRequest { #[derive(Clone)] pub struct RenderPass { - desc: metal::RenderPassDescriptor + desc: metal::RenderPassDescriptor, + /// Colour attachment formats, one per MRT target (index 0 = SV_Target0) + pixel_formats: Vec, + depth_format: Option, + /// MSAA sample count shared by all attachments in the pass (1 = no MSAA) + sample_count: u32, } impl super::RenderPass for RenderPass { @@ -716,12 +1528,17 @@ impl super::RenderPass for RenderPass { } pub struct ComputePipeline { - slots: Vec + pipeline_state: metal::ComputePipelineState, + slots: Vec, + /// Unified slot lookup by (register, space, descriptor_type) + slot_lookup: HashMap, + /// Single-stage binders for push constants and resource bindings + compute_binder: HashMap, } impl super::Pipeline for ComputePipeline { fn get_pipeline_slot(&self, register: u32, space: u32, descriptor_type: DescriptorType) -> Option<&super::PipelineSlotInfo> { - None + self.slot_lookup.get(&(register, space, descriptor_type)) } fn get_pipeline_slots(&self) -> &Vec { @@ -747,7 +1564,15 @@ pub struct Heap { buffer_slots: Vec>, resource_type: Vec, offset: usize, - id: u16 + id: u16, + /// Argument encoder for bindless texture access (pre-encodes all textures) + texture_argument_encoder: metal::ArgumentEncoder, + /// Pre-encoded argument buffer containing all texture references + texture_argument_buffer: metal::Buffer, + /// Argument encoder for bindless buffer access + buffer_argument_encoder: metal::ArgumentEncoder, + /// Pre-encoded argument buffer containing all buffer references + buffer_argument_buffer: metal::Buffer, } impl Heap { @@ -761,9 +1586,31 @@ impl Heap { self.resource_type.resize(self.offset, HeapResourceType::None); srv } -} -impl super::Heap for Heap { + /// Encode a texture into the heap's argument buffer at the given index (for bindless) + fn encode_texture(&self, index: usize, texture: &metal::Texture) { + self.texture_argument_encoder.set_argument_buffer(&self.texture_argument_buffer, 0); + self.texture_argument_encoder.set_texture(index as u64, texture); + } + + /// Encode a buffer into the heap's argument buffer at the given index (for bindless) + fn encode_buffer(&self, index: usize, buffer: &metal::Buffer) { + self.buffer_argument_encoder.set_argument_buffer(&self.buffer_argument_buffer, 0); + self.buffer_argument_encoder.set_buffer(index as u64, buffer, 0); + } + + /// Get the pre-encoded texture argument buffer for binding + pub fn get_texture_argument_buffer(&self) -> &metal::Buffer { + &self.texture_argument_buffer + } + + /// Get the pre-encoded buffer argument buffer for binding + pub fn get_buffer_argument_buffer(&self) -> &metal::Buffer { + &self.buffer_argument_buffer + } +} + +impl super::Heap for Heap { fn deallocate(&mut self, index: usize) { } @@ -772,16 +1619,20 @@ impl super::Heap for Heap { } fn get_heap_id(&self) -> u16 { - 0 + self.id } } pub struct QueryHeap { - + heap_type: super::QueryType, + sample_buffer: Option, + alloc_index: usize, + capacity: usize, } impl super::QueryHeap for QueryHeap { fn reset(&mut self) { + self.alloc_index = 0; } } @@ -806,6 +1657,16 @@ pub struct RaytracingTLAS { } impl Device { + /// Largest texture sample count <= `requested` that this device supports (always >= 1). + /// Apple GPUs commonly cap at 4x, so an 8x request is clamped down rather than asserting. + fn supported_sample_count(&self, requested: u32) -> u32 { + let mut count = requested.max(1); + while count > 1 && !self.metal_device.supports_texture_sample_count(count as NSUInteger) { + count /= 2; + } + count + } + fn create_render_pass_for_swap_chain( &self, texture: &Texture, @@ -832,137 +1693,394 @@ impl Device { texture_descriptor.set_depth(1); texture_descriptor.set_texture_type(metal::MTLTextureType::D2); texture_descriptor.set_pixel_format(metal::MTLPixelFormat::RGBA8Unorm); - texture_descriptor.set_storage_mode(metal::MTLStorageMode::Shared); + // Private storage: required for MSAA textures (which can't be Shared) and faster for + // GPU sampling on Apple Silicon. Texture data is uploaded via a staging buffer + blit. + texture_descriptor.set_storage_mode(metal::MTLStorageMode::Private); // Determine the size required for the heap for the given descriptor let size_and_align = mtl_device.heap_texture_size_and_align(&texture_descriptor); let texture_size = align_pow2(size_and_align.size, size_and_align.align); - let heap_size = texture_size * info.num_descriptors.max(1) as u64; + // The 512x512 RGBA8 reference (~1MB) underestimates real descriptors: 2k material + // textures, IBL cubemaps and MSAA render targets are far larger. Oversize the heap so + // the bindless descriptor pool doesn't run out of memory when many/large textures load. + const HEAP_OVERSIZE_FACTOR: u64 = 2; + let heap_size = texture_size * info.num_descriptors.max(1) as u64 * HEAP_OVERSIZE_FACTOR; let heap_descriptor = metal::HeapDescriptor::new(); - heap_descriptor.set_storage_mode(metal::MTLStorageMode::Shared); + heap_descriptor.set_storage_mode(metal::MTLStorageMode::Private); heap_descriptor.set_size(heap_size); + // Enable hazard tracking so Metal automatically synchronizes heap-allocated + // textures across command buffers (by default heaps are MTLHazardTrackingModeUntracked) + unsafe { let _: () = msg_send![&*heap_descriptor, setHazardTrackingMode: metal::MTLHazardTrackingMode::Tracked]; }; + + /* + // newHeapWithDescriptor: returns nil on allocation failure (e.g. requested size + // exceeds what the GPU can back). metal-rs wraps the result without checking, so + // every later deref of a nil heap would trip foreign-types' from_ptr assert with + // no useful context. Probe first via raw msg_send so we can report what was + // actually requested before falling back to metal-rs's wrapper. + let probe_ptr: *mut objc::runtime::Object = unsafe { + msg_send![&*mtl_device, newHeapWithDescriptor: &*heap_descriptor] + }; + assert!( + !probe_ptr.is_null(), + "hotline_rs::gfx::mtl: failed to allocate MTLHeap ({:.1} MB, {} descriptors, storage Private). \ + Requested size = texture_size({} B) * num_descriptors({}) * HEAP_OVERSIZE_FACTOR({}). \ + Reduce num_descriptors in HeapInfo, or use a smaller per-batch heap for buffer-only allocations.", + heap_size as f64 / (1024.0 * 1024.0), + info.num_descriptors, + texture_size, + info.num_descriptors.max(1), + HEAP_OVERSIZE_FACTOR + ); + // probe_ptr is a +1 retained heap; release it and let metal-rs allocate again so + // we keep using its wrapper type without juggling raw pointer ownership. + unsafe { let _: () = msg_send![probe_ptr, release]; } + */ + let heap = mtl_device.new_heap(&heap_descriptor); + // Create texture argument encoder for bindless access + let max_resources = info.num_descriptors.max(1) as u64; + let tex_arg_desc = metal::ArgumentDescriptor::new(); + tex_arg_desc.set_index(0); + tex_arg_desc.set_data_type(metal::MTLDataType::Texture); + tex_arg_desc.set_array_length(max_resources); + tex_arg_desc.set_access(metal::MTLArgumentAccess::ReadOnly); + + let texture_argument_encoder = mtl_device.new_argument_encoder( + metal::Array::from_owned_slice(&[tex_arg_desc.to_owned()]) + ); + let texture_argument_buffer = mtl_device.new_buffer( + texture_argument_encoder.encoded_length(), + metal::MTLResourceOptions::StorageModeShared + ); + + // Create buffer argument encoder for bindless access + let buf_arg_desc = metal::ArgumentDescriptor::new(); + buf_arg_desc.set_index(0); + buf_arg_desc.set_data_type(metal::MTLDataType::Pointer); + buf_arg_desc.set_array_length(max_resources); + buf_arg_desc.set_access(metal::MTLArgumentAccess::ReadOnly); + + let buffer_argument_encoder = mtl_device.new_argument_encoder( + metal::Array::from_owned_slice(&[buf_arg_desc.to_owned()]) + ); + let buffer_argument_buffer = mtl_device.new_buffer( + buffer_argument_encoder.encoded_length(), + metal::MTLResourceOptions::StorageModeShared + ); + Heap { mtl_heap: heap, texture_slots: Vec::new(), buffer_slots: Vec::new(), resource_type: Vec::new(), offset: 0, - id + id, + texture_argument_encoder, + texture_argument_buffer, + buffer_argument_encoder, + buffer_argument_buffer, } } - fn to_mtl_descriptor_slot(&self, visibility: super::ShaderVisibility, pipeline_bindings: &Option>) -> DescriptorSlotArray { - // argument buffer to descriptor slot style - let mut descriptor_slots : DescriptorSlotArray = Vec::new(); + /// Build unified slot lookup + fn build_slot_lookup( + &self, + pipeline_bindings: &Option>, + pipeline_push_constants: &Option>, + ) -> HashMap { + let mut slot_lookup: HashMap = HashMap::new(); + + // hardcoded sampler offsets + let vertex_samplers_offset: u32 = 2; + let fragment_samplers_offset: u32 = 0; + let mut vertex_binding_offset: u32 = vertex_samplers_offset + 1; + let mut fragment_binding_offset: u32 = fragment_samplers_offset + 1; + + // Add push constant slots first (they come before regular bindings in htwv) + if let Some(push_constants) = pipeline_push_constants.as_ref() { + for push_constant in push_constants { + // Determine stage indices based on visibility, using per-stage offsets + let (vertex_idx, fragment_idx, canonical_index) = match push_constant.visibility { + ShaderVisibility::Vertex => { + let idx = vertex_binding_offset; + vertex_binding_offset += 1; + (Some(idx), None, idx) + }, + ShaderVisibility::Fragment => { + let idx = fragment_binding_offset; + fragment_binding_offset += 1; + (None, Some(idx), idx) + }, + ShaderVisibility::All => { + let v_idx = vertex_binding_offset; + let f_idx = fragment_binding_offset; + vertex_binding_offset += 1; + fragment_binding_offset += 1; + // Use vertex index as canonical for lookup + (Some(v_idx), Some(f_idx), v_idx) + }, + _ => (None, None, 0), + }; - // register spaces, and shader registers may not be ordered and may not be sequential or have gaps - if let Some(bindings) = pipeline_bindings.as_ref() { - // make space for enough shader register spaces - let mut space_count = 0; - for binding in bindings.iter().filter(|b| b.visibility == visibility || b.visibility == ShaderVisibility::All) { - space_count = binding.register_space.max(space_count); + slot_lookup.insert( + (push_constant.shader_register, push_constant.register_space, DescriptorType::PushConstants), + PipelineSlotInfo { + index: canonical_index, + count: Some(push_constant.num_values), + }, + ); } - descriptor_slots.resize((space_count + 1) as usize, None); - - // iterate over descriptor slots and find members - descriptor_slots.iter_mut().enumerate().for_each(|(space, descriptor_slot)| { - let mut members : DescriptorMemberArray = Vec::new(); - for binding in bindings.iter().filter(|b| b.visibility == visibility || b.visibility == ShaderVisibility::All) { - if binding.register_space == space as u32 { - if members.len() < (binding.shader_register + 1) as usize { - members.resize((binding.shader_register + 1) as usize, None); - } + } - // get num - let num = if let Some(num) = binding.num_descriptors { - num + // Add regular binding slots, grouped by (register_kind, shader_register, register_space) to + // mirror the descriptor-set layout produced by htwv's MSL codegen. Each (kind, register, + // space) becomes its own MSL [[buffer(N)]] slot so the heap's texture and buffer argument + // buffers never share a slot, and bindless arrays sharing a register but differing by space + // (e.g. textures t1/space7, cubemaps t1/space9) each get their own set at id(0). + if let Some(bindings) = pipeline_bindings.as_ref() { + if !bindings.is_empty() { + // (register_kind, shader_register, register_space) -> buffer_index + let mut v_groups: HashMap<(char, u32, u32), u32> = HashMap::new(); + let mut f_groups: HashMap<(char, u32, u32), u32> = HashMap::new(); + + for binding in bindings { + let key = (descriptor_register_kind(binding.binding_type), binding.shader_register, binding.register_space); + + let v_slot = if matches!(binding.visibility, ShaderVisibility::Vertex | ShaderVisibility::All) { + Some(*v_groups.entry(key).or_insert_with(|| { + let idx = vertex_binding_offset; + vertex_binding_offset += 1; + idx + })) + } else { None }; + + let f_slot = if matches!(binding.visibility, ShaderVisibility::Fragment | ShaderVisibility::All) { + Some(*f_groups.entry(key).or_insert_with(|| { + let idx = fragment_binding_offset; + fragment_binding_offset += 1; + idx + })) + } else { None }; + + let canonical_index = v_slot.or(f_slot).unwrap_or(0); + slot_lookup.insert( + (binding.shader_register, binding.register_space, binding.binding_type), + PipelineSlotInfo { + index: canonical_index, + count: binding.num_descriptors, } - else { - 1 - }; - - // assign member info - members[binding.shader_register as usize] = Some( - DescriptorMember { - offset: 0, - num: num, - info: PipelineSlotInfo { - index: binding.shader_register, - count: binding.num_descriptors - } - } - ); - } + ); } + } + } - // now work out the offsets of the members within the space - let mut offset = 0; - for member in &mut members { - if let Some(member) = member { - member.offset = offset; - offset += member.num; - } - } + slot_lookup + } - // finally if we have members and not an empty space - // create an argument buffer - if members.len() > 0 { - let mut member_descriptors = Vec::new(); + fn build_stage_binders( + &self, + pipeline_bindings: &Option>, + pipeline_push_constants: &Option>, + ) -> (HashMap, HashMap) { + const MAX_BINDLESS_TEXTURES: u64 = 1024; - let mut total_num = 0; - for member in &members { - if let Some(member) = member { - let descriptor = metal::ArgumentDescriptor::new(); - descriptor.set_index(member.offset as u64); - descriptor.set_array_length(member.num as u64); + let mut vertex_binder: HashMap = HashMap::new(); + let mut fragment_binder: HashMap = HashMap::new(); - // TODO: types / access - descriptor.set_data_type(metal::MTLDataType::Texture); - descriptor.set_access(metal::MTLArgumentAccess::ReadOnly); + let vertex_samplers_offset: u32 = 2; + let fragment_samplers_offset: u32 = 0; - // push metal argument descriptor - member_descriptors.push(descriptor.to_owned()); + let mut vertex_binding_offset: u32 = vertex_samplers_offset + 1; + let mut fragment_binding_offset: u32 = fragment_samplers_offset + 1; - total_num += member.num; - } - } + // Add push constant binders (no ArgumentEncoder needed - uses setVertexBytes/setFragmentBytes) + if let Some(push_constants) = pipeline_push_constants.as_ref() { + for push_constant in push_constants { + let key: SlotKey = ( + push_constant.shader_register, + push_constant.register_space, + DescriptorType::PushConstants + ); - // create encoder and argument buffer - let argument_encoder = self.metal_device.new_argument_encoder(metal::Array::from_owned_slice(member_descriptors.as_slice())); - let argument_buffer_size = argument_encoder.encoded_length() * total_num as u64; - let argument_buffer = self.metal_device.new_buffer(argument_buffer_size, metal::MTLResourceOptions::empty()); + match push_constant.visibility { + ShaderVisibility::Vertex => { + let buffer_index = vertex_binding_offset; + vertex_binding_offset += 1; + + vertex_binder.insert(key, PipelineStageBinder::PushConstants(PushConstantsBinder { + data: vec![0u32; push_constant.num_values as usize], + num_32_bit_constants: push_constant.num_values, + buffer_index, + dirty: true, + })); + }, + ShaderVisibility::Fragment => { + let buffer_index = fragment_binding_offset; + fragment_binding_offset += 1; + + fragment_binder.insert(key, PipelineStageBinder::PushConstants(PushConstantsBinder { + data: vec![0u32; push_constant.num_values as usize], + num_32_bit_constants: push_constant.num_values, + buffer_index, + dirty: true, + })); + }, + ShaderVisibility::All => { + let v_buffer_index = vertex_binding_offset; + let f_buffer_index = fragment_binding_offset; + vertex_binding_offset += 1; + fragment_binding_offset += 1; + + vertex_binder.insert(key, PipelineStageBinder::PushConstants(PushConstantsBinder { + data: vec![0u32; push_constant.num_values as usize], + num_32_bit_constants: push_constant.num_values, + buffer_index: v_buffer_index, + dirty: true, + })); + + fragment_binder.insert(key, PipelineStageBinder::PushConstants(PushConstantsBinder { + data: vec![0u32; push_constant.num_values as usize], + num_32_bit_constants: push_constant.num_values, + buffer_index: f_buffer_index, + dirty: true, + })); + }, + _ => {}, + } + } + } - *descriptor_slot = Some( - DescriptorSlot { - argument_encoder, - argument_buffer, - members - } - ) + // Add resource binders, grouped by (register_kind, shader_register, register_space). Each + // (kind, register, space) gets its own [[buffer(N)]] slot per stage so the heap's texture + // and buffer argument buffers are bound to distinct slots, and bindless arrays sharing a + // register but differing by space each get their own set. Each group holds exactly one + // binding, so it always sits at id(0) - binding_index is always 0. + if let Some(bindings) = pipeline_bindings.as_ref() { + if !bindings.is_empty() { + // (register_kind, shader_register, register_space) -> buffer_index + let mut v_groups: HashMap<(char, u32, u32), u32> = HashMap::new(); + let mut f_groups: HashMap<(char, u32, u32), u32> = HashMap::new(); + + for binding in bindings { + let key: SlotKey = (binding.shader_register, binding.register_space, binding.binding_type); + let group_key = (descriptor_register_kind(binding.binding_type), binding.shader_register, binding.register_space); + let data_type = to_mtl_data_type( + binding.resource_type.expect("hotline_rs::gfx::mtl: requires resource type for binding") + ); + let array_length = binding.num_descriptors.map(|n| n as u64).unwrap_or(MAX_BINDLESS_TEXTURES); + + if matches!(binding.visibility, ShaderVisibility::Vertex | ShaderVisibility::All) { + let buffer_index = *v_groups.entry(group_key).or_insert_with(|| { + let idx = vertex_binding_offset; + vertex_binding_offset += 1; + idx + }); + vertex_binder.insert(key, PipelineStageBinder::Resource(ResourceBinder { + buffer_index, + binding_index: 0, + data_type, + array_length, + bound_resource: None, + dirty: true, + })); + } + + if matches!(binding.visibility, ShaderVisibility::Fragment | ShaderVisibility::All) { + let buffer_index = *f_groups.entry(group_key).or_insert_with(|| { + let idx = fragment_binding_offset; + fragment_binding_offset += 1; + idx + }); + fragment_binder.insert(key, PipelineStageBinder::Resource(ResourceBinder { + buffer_index, + binding_index: 0, + data_type, + array_length, + bound_resource: None, + dirty: true, + })); + } } - }); + } } - descriptor_slots + (vertex_binder, fragment_binder) } - fn to_mtl_push_constant_slot(&self, visibility: super::ShaderVisibility, pipeline_push_constants: &Option>, binding_offset: u32) -> Vec { - let mut push_constant_slots : Vec = Vec::new(); + /// Build a single-stage binder for a compute pipeline. Mirrors `build_stage_binders` but emits + /// one map: push constants and resource bindings share a single MSL [[buffer(N)]] namespace + /// (no vertex/fragment split). Buffer indices begin at `COMPUTE_BINDING_BASE` which must match + /// the [[buffer(N)]] slots htwv emits for the compute kernel. + fn build_compute_binder( + &self, + pipeline_bindings: &Option>, + pipeline_push_constants: &Option>, + ) -> HashMap { + const MAX_BINDLESS_TEXTURES: u64 = 1024; + // Compute follows the same MSL [[buffer(N)]] layout htwv emits for the fragment stage: + // buffer(0) is reserved for the sampler descriptor set, push constants take buffer(1), and + // space0 resource descriptor sets follow at buffer(2)+. So start binding indices at 1. + const COMPUTE_BINDING_BASE: u32 = 1; + + let mut binder: HashMap = HashMap::new(); + let mut binding_offset: u32 = COMPUTE_BINDING_BASE; + + // Push constants (use setBytes - no ArgumentEncoder) if let Some(push_constants) = pipeline_push_constants.as_ref() { - push_constants.iter().filter(|b| b.visibility == visibility || b.visibility == ShaderVisibility::All).enumerate().for_each(|(index, push_constant)| { - push_constant_slots.push(PushConstantSlot{ - buffer: self.metal_device.new_buffer(push_constant.num_values as u64 * 4, metal::MTLResourceOptions::StorageModeShared), - slot: binding_offset + index as u32, - visibility: ShaderVisibility::All - }) - }); + for push_constant in push_constants { + let key: SlotKey = ( + push_constant.shader_register, + push_constant.register_space, + DescriptorType::PushConstants, + ); + let buffer_index = binding_offset; + binding_offset += 1; + binder.insert(key, PipelineStageBinder::PushConstants(PushConstantsBinder { + data: vec![0u32; push_constant.num_values as usize], + num_32_bit_constants: push_constant.num_values, + buffer_index, + dirty: true, + })); + } + } + + // Resource bindings grouped by (register_kind, shader_register, register_space) - one + // [[buffer(N)]] per group, so arrays sharing a register but differing by space each get + // their own set. Each group holds exactly one binding, so it always sits at id(0). + if let Some(bindings) = pipeline_bindings.as_ref() { + if !bindings.is_empty() { + let mut groups: HashMap<(char, u32, u32), u32> = HashMap::new(); + for binding in bindings { + let key: SlotKey = (binding.shader_register, binding.register_space, binding.binding_type); + let group_key = (descriptor_register_kind(binding.binding_type), binding.shader_register, binding.register_space); + let data_type = to_mtl_data_type( + binding.resource_type.expect("hotline_rs::gfx::mtl: requires resource type for binding") + ); + let array_length = binding.num_descriptors.map(|n| n as u64).unwrap_or(MAX_BINDLESS_TEXTURES); + + let buffer_index = *groups.entry(group_key).or_insert_with(|| { + let idx = binding_offset; + binding_offset += 1; + idx + }); + binder.insert(key, PipelineStageBinder::Resource(ResourceBinder { + buffer_index, + binding_index: 0, + data_type, + array_length, + bound_resource: None, + dirty: true, + })); + } + } } - push_constant_slots + + binder } } @@ -1002,7 +2120,13 @@ impl super::Device for Device { // feature info let tier = device.argument_buffers_support(); - assert_eq!(metal::MTLArgumentBuffersTier::Tier2, tier); //TODO: message + assert_eq!(metal::MTLArgumentBuffersTier::Tier2, tier); + + // Can the GPU sample timestamp counters at encoder stage boundaries? (Apple Silicon + // typically can; older/other GPUs may not, in which case we fall back to whole-CB times.) + let supports_stage_boundary_timestamps: bool = unsafe { + msg_send![&*device, supportsCounterSampling: MTL_COUNTER_SAMPLING_POINT_AT_STAGE_BOUNDARY] + }; Device { command_queue: command_queue, @@ -1012,7 +2136,9 @@ impl super::Device for Device { debug_name: Some("mtl device: shader heap".to_string()) }, 1), adapter_info: adapter_info, - metal_device: device + metal_device: device, + heap_alloc_id: 2, + supports_stage_boundary_timestamps, } }) } @@ -1022,12 +2148,31 @@ impl super::Device for Device { } fn create_heap(&mut self, info: &HeapInfo) -> Heap { - Self::create_heap_mtl(&self.metal_device, &info, 2) + let id = self.heap_alloc_id; + self.heap_alloc_id += 1; + Self::create_heap_mtl(&self.metal_device, &info, id) } fn create_query_heap(&self, info: &QueryHeapInfo) -> QueryHeap { + let sample_buffer = if info.heap_type == super::QueryType::Timestamp + && self.supports_stage_boundary_timestamps { + let counter_sets = self.metal_device.counter_sets(); + let ts_set = counter_sets.iter().find(|cs| cs.name().eq_ignore_ascii_case("timestamp")); + ts_set.and_then(|cs| { + let desc = metal::CounterSampleBufferDescriptor::new(); + desc.set_counter_set(cs); + desc.set_sample_count(info.num_queries as _); + desc.set_storage_mode(metal::MTLStorageMode::Shared); + self.metal_device.new_counter_sample_buffer_with_descriptor(&desc).ok() + }) + } else { + None + }; QueryHeap { - + heap_type: info.heap_type, + sample_buffer, + alloc_index: 0, + capacity: info.num_queries, } } @@ -1051,6 +2196,7 @@ impl super::Device for Device { view.setLayer(std::mem::transmute(layer.as_ref())); let draw_size = win.get_size(); + layer.set_contents_scale(win.get_dpi_scale() as f64); layer.set_drawable_size(CGSize::new(draw_size.x as f64, draw_size.y as f64)); let drawable = layer.next_drawable() @@ -1058,7 +2204,11 @@ impl super::Device for Device { let backbuffer_texture = Texture { metal_texture: drawable.texture().to_owned(), + resolved_texture: None, srv_index: None, + msaa_srv_index: None, + uav_index: None, + resolvable: false, heap_id: None }; let render_pass = self.create_render_pass_for_swap_chain(&backbuffer_texture, info.clear_colour); @@ -1073,6 +2223,10 @@ impl super::Device for Device { backbuffer_texture: backbuffer_texture, backbuffer_pass: render_pass, backbuffer_pass_no_clear: render_pass_no_clear, + num_buffers: info.num_buffers, + frame_event: self.metal_device.new_event(), + frame_value: 0, + in_flight: std::sync::Arc::new(std::sync::Mutex::new(std::collections::VecDeque::new())), }) }) } @@ -1080,13 +2234,25 @@ impl super::Device for Device { fn create_cmd_buf(&self, num_buffers: u32) -> CmdBuf { objc::rc::autoreleasepool(|| { + let cmd_queue = self.command_queue.clone(); + let cmd = cmd_queue.new_command_buffer().to_owned(); + CmdBuf { - cmd_queue: self.command_queue.clone(), - cmd: None, + cmd_queue, + cmd: Some(cmd), render_encoder: None, compute_encoder: None, bound_index_buffer: None, - bound_index_stride: 0 + bound_index_stride: 0, + bound_render_pipeline: None, + bound_compute_pipeline: None, + metal_device: self.metal_device.clone(), + transient_buffers: Vec::new(), + vertex_binder: HashMap::new(), + fragment_binder: HashMap::new(), + compute_binder: HashMap::new(), + deferred_ops: Vec::new(), + pending_timestamp: None, } }) } @@ -1101,14 +2267,16 @@ impl super::Device for Device { if let Some(vs) = info.vs { unsafe { let lib = self.metal_device.new_library_with_data(std::slice::from_raw_parts(vs.data, vs.data_size))?; - let vvs = lib.get_function("vs_main", None).unwrap(); + let name = &lib.function_names()[0]; + let vvs = lib.get_function(name, None).unwrap(); pipeline_state_descriptor.set_vertex_function(Some(&vvs)); } }; if let Some(fs) = info.fs { unsafe { let lib = self.metal_device.new_library_with_data(std::slice::from_raw_parts(fs.data, fs.data_size))?; - let pps = lib.get_function("ps_main", None).unwrap(); + let name = &lib.function_names()[0]; + let pps = lib.get_function(name, None).unwrap(); pipeline_state_descriptor.set_fragment_function(Some(&pps)); } }; @@ -1117,15 +2285,21 @@ impl super::Device for Device { let vertex_desc = metal::VertexDescriptor::new(); let mut attrib_index = 0; - // make spaces for slots to calculate the stride from offsets + size - let mut slot_strides = Vec::new(); + // track stride, step function, and step rate per slot + struct SlotLayout { + stride: u32, + input_slot_class: super::InputSlotClass, + step_rate: u32, + } + let mut slot_layouts: Vec> = Vec::new(); for element in &info.input_layout { - if slot_strides.len() < (element.input_slot + 1) as usize { - slot_strides.resize((element.input_slot + 1) as usize, 0); + let slot = element.input_slot as usize; + if slot_layouts.len() <= slot { + slot_layouts.resize_with(slot + 1, || None); } } - // make the idividual attributes and track the stride of each slot + // make the individual attributes and track the stride/stepping of each slot for element in &info.input_layout { let attribute = metal::VertexAttributeDescriptor::new(); attribute.set_format(to_mtl_vertex_format(element.format)); @@ -1135,46 +2309,136 @@ impl super::Device for Device { attrib_index += 1; let stride = element.aligned_byte_offset + block_size_for_format(element.format); - slot_strides[element.input_slot as usize] = max(slot_strides[element.input_slot as usize], stride); + let slot = element.input_slot as usize; + if let Some(ref mut layout) = slot_layouts[slot] { + layout.stride = max(layout.stride, stride); + } else { + slot_layouts[slot] = Some(SlotLayout { + stride, + input_slot_class: element.input_slot_class, + step_rate: element.step_rate, + }); + } } - // vertex layouts; TODO: work out MTLVertexStepFunction - let layout_desc = metal::VertexBufferLayoutDescriptor::new(); - layout_desc.set_step_function(metal::MTLVertexStepFunction::PerVertex); - layout_desc.set_stride(slot_strides[0] as NSUInteger); - vertex_desc.layouts().set_object_at(0, Some(&layout_desc)); + // create vertex buffer layouts for each slot + for (slot, layout_opt) in slot_layouts.iter().enumerate() { + if let Some(layout) = layout_opt { + let layout_desc = metal::VertexBufferLayoutDescriptor::new(); + layout_desc.set_stride(layout.stride as NSUInteger); + match layout.input_slot_class { + super::InputSlotClass::PerVertex => { + layout_desc.set_step_function(metal::MTLVertexStepFunction::PerVertex); + layout_desc.set_step_rate(1); + } + super::InputSlotClass::PerInstance => { + layout_desc.set_step_function(metal::MTLVertexStepFunction::PerInstance); + layout_desc.set_step_rate(layout.step_rate as NSUInteger); + } + } + vertex_desc.layouts().set_object_at(slot as NSUInteger, Some(&layout_desc)); + } + } pipeline_state_descriptor.set_vertex_descriptor(Some(&vertex_desc)); - // TODO: attachments - let attachment = pipeline_state_descriptor - .color_attachments() - .object_at(0) - .unwrap(); - attachment.set_pixel_format(metal::MTLPixelFormat::BGRA8Unorm); - attachment.set_blending_enabled(false); - attachment.set_rgb_blend_operation(metal::MTLBlendOperation::Add); - attachment.set_alpha_blend_operation(metal::MTLBlendOperation::Add); - attachment.set_source_rgb_blend_factor(metal::MTLBlendFactor::SourceAlpha); - attachment.set_source_alpha_blend_factor(metal::MTLBlendFactor::SourceAlpha); - attachment.set_destination_rgb_blend_factor(metal::MTLBlendFactor::OneMinusSourceAlpha); - attachment.set_destination_alpha_blend_factor(metal::MTLBlendFactor::OneMinusSourceAlpha); - - // TODO: depth stencil - - // TODO: raster - - // TODO: samplers? + // colour attachments - one per MRT target from the pass (SV_Target0..N). With no pass + // (eg. depth-only / default) fall back to a single BGRA8 attachment. + let pixel_formats: Vec = info.pass + .map(|p| p.pixel_formats.clone()) + .filter(|f| !f.is_empty()) + .unwrap_or_else(|| vec![metal::MTLPixelFormat::BGRA8Unorm]); + + for (i, &pixel_format) in pixel_formats.iter().enumerate() { + let attachment = pipeline_state_descriptor + .color_attachments() + .object_at(i as u64) + .unwrap(); + attachment.set_pixel_format(pixel_format); + + if pixel_format == metal::MTLPixelFormat::Invalid { + continue; + } + + // per-target blend state (falls back to the first / disabled) + let blend = info.blend_info.render_target.get(i) + .or_else(|| info.blend_info.render_target.first()); + if let Some(b) = blend { + attachment.set_blending_enabled(b.blend_enabled); + attachment.set_rgb_blend_operation(to_mtl_blend_op(&b.blend_op)); + attachment.set_alpha_blend_operation(to_mtl_blend_op(&b.blend_op_alpha)); + attachment.set_source_rgb_blend_factor(to_mtl_blend_factor(&b.src_blend)); + attachment.set_source_alpha_blend_factor(to_mtl_blend_factor(&b.src_blend_alpha)); + attachment.set_destination_rgb_blend_factor(to_mtl_blend_factor(&b.dst_blend)); + attachment.set_destination_alpha_blend_factor(to_mtl_blend_factor(&b.dst_blend_alpha)); + attachment.set_write_mask(to_mtl_write_mask(&b.write_mask)); + } else { + attachment.set_blending_enabled(false); + attachment.set_write_mask(metal::MTLColorWriteMask::all()); + } + } + + // Set depth format + MSAA sample count on pipeline descriptor to match the pass + if let Some(pass) = &info.pass { + if let Some(depth_format) = pass.depth_format { + pipeline_state_descriptor.set_depth_attachment_pixel_format(depth_format); + if has_stencil_component(depth_format) { + pipeline_state_descriptor.set_stencil_attachment_pixel_format(depth_format); + } + } + pipeline_state_descriptor.set_sample_count(pass.sample_count as NSUInteger); + } + + // Create depth stencil state + let depth_stencil_state = { + let ds_info = &info.depth_stencil_info; + let ds_desc = metal::DepthStencilDescriptor::new(); + + ds_desc.set_depth_compare_function(to_mtl_compare_func(ds_info.depth_func)); + ds_desc.set_depth_write_enabled(ds_info.depth_write_mask == super::DepthWriteMask::All); + + if ds_info.stencil_enabled { + // Front face + let front = metal::StencilDescriptor::new(); + front.set_stencil_compare_function(to_mtl_compare_func(ds_info.front_face.func)); + front.set_stencil_failure_operation(to_mtl_stencil_op(ds_info.front_face.fail)); + front.set_depth_failure_operation(to_mtl_stencil_op(ds_info.front_face.depth_fail)); + front.set_depth_stencil_pass_operation(to_mtl_stencil_op(ds_info.front_face.pass)); + front.set_read_mask(ds_info.stencil_read_mask as u32); + front.set_write_mask(ds_info.stencil_write_mask as u32); + ds_desc.set_front_face_stencil(Some(&front)); + + // Back face + let back = metal::StencilDescriptor::new(); + back.set_stencil_compare_function(to_mtl_compare_func(ds_info.back_face.func)); + back.set_stencil_failure_operation(to_mtl_stencil_op(ds_info.back_face.fail)); + back.set_depth_failure_operation(to_mtl_stencil_op(ds_info.back_face.depth_fail)); + back.set_depth_stencil_pass_operation(to_mtl_stencil_op(ds_info.back_face.pass)); + back.set_read_mask(ds_info.stencil_read_mask as u32); + back.set_write_mask(ds_info.stencil_write_mask as u32); + ds_desc.set_back_face_stencil(Some(&back)); + } + + self.metal_device.new_depth_stencil_state(&ds_desc) + }; + + // Create static samplers and argument buffer (bound at fragment buffer(0)) let mut pipeline_static_samplers = Vec::new(); + let mut sampler_argument_buffer = None; + if let Some(static_samplers) = &info.pipeline_layout.static_samplers { for sampler in static_samplers { + let si = &sampler.sampler_info; let desc = metal::SamplerDescriptor::new(); - desc.set_address_mode_r(metal::MTLSamplerAddressMode::Repeat); - desc.set_address_mode_s(metal::MTLSamplerAddressMode::Repeat); - desc.set_address_mode_t(metal::MTLSamplerAddressMode::Repeat); - desc.set_min_filter(metal::MTLSamplerMinMagFilter::Linear); - desc.set_mag_filter(metal::MTLSamplerMinMagFilter::Linear); - desc.set_mip_filter(metal::MTLSamplerMipFilter::Linear); + desc.set_address_mode_r(to_mtl_sampler_address_mode(si.address_w)); + desc.set_address_mode_s(to_mtl_sampler_address_mode(si.address_u)); + desc.set_address_mode_t(to_mtl_sampler_address_mode(si.address_v)); + desc.set_min_filter(to_mtl_sampler_min_mag_filter(si.filter)); + desc.set_mag_filter(to_mtl_sampler_min_mag_filter(si.filter)); + desc.set_mip_filter(to_mtl_sampler_mip_filter(si.filter)); + if let Some(func) = si.comparison { + desc.set_compare_function(to_mtl_compare_func(func)); + } desc.set_support_argument_buffers(true); pipeline_static_samplers.push(MetalSamplerBinding { @@ -1182,12 +2446,48 @@ impl super::Device for Device { sampler: self.metal_device.new_sampler(&desc) }) } + + // Create argument buffer for samplers. SPIRV-Cross repacks the samplers actually + // used by a shader into spvDescriptorSetBuffer0 with sequential ids starting at 0 + // (it does NOT preserve the HLSL register, eg. sampler_wrap_linear at s1 becomes + // [[id(0)]]). So encode each sampler at its position in the static_samplers list, + // which matches that packing order. + if !pipeline_static_samplers.is_empty() { + let arg_desc = metal::ArgumentDescriptor::new(); + arg_desc.set_index(0); + arg_desc.set_data_type(metal::MTLDataType::Sampler); + arg_desc.set_array_length(pipeline_static_samplers.len() as u64); + arg_desc.set_access(metal::MTLArgumentAccess::ReadOnly); + + let argument_encoder = self.metal_device.new_argument_encoder( + metal::Array::from_owned_slice(&[arg_desc.to_owned()]) + ); + let arg_buffer = self.metal_device.new_buffer( + argument_encoder.encoded_length(), + metal::MTLResourceOptions::StorageModeShared + ); + + // Encode each sampler at its packed id (list position) + argument_encoder.set_argument_buffer(&arg_buffer, 0); + for (id, s) in pipeline_static_samplers.iter().enumerate() { + argument_encoder.set_sampler_state(id as u64, &s.sampler); + } + + sampler_argument_buffer = Some(arg_buffer); + } } - let vertex_descriptor_slots = self.to_mtl_descriptor_slot(ShaderVisibility::Vertex, &info.pipeline_layout.bindings); - let fragment_descriptor_slots = self.to_mtl_descriptor_slot(ShaderVisibility::Fragment, &info.pipeline_layout.bindings); - let vertex_push_constant_slots = self.to_mtl_push_constant_slot(ShaderVisibility::Vertex, &info.pipeline_layout.push_constants, vertex_descriptor_slots.len() as u32); - let fragment_push_constant_slots = self.to_mtl_push_constant_slot(ShaderVisibility::Fragment, &info.pipeline_layout.push_constants, fragment_descriptor_slots.len() as u32); + // Build unified slot lookup + let slot_lookup = self.build_slot_lookup( + &info.pipeline_layout.bindings, + &info.pipeline_layout.push_constants, + ); + + // Build stage binders for push constants and resource bindings + let (vertex_binder, fragment_binder) = self.build_stage_binders( + &info.pipeline_layout.bindings, + &info.pipeline_layout.push_constants, + ); let pipeline_state = self.metal_device.new_render_pipeline_state(&pipeline_state_descriptor)?; @@ -1195,10 +2495,13 @@ impl super::Device for Device { pipeline_state, slots: Vec::new(), static_samplers: pipeline_static_samplers, - fragment_descriptor_slots, - vertex_descriptor_slots, - vertex_push_constant_slots, - fragment_push_constant_slots + slot_lookup, + vertex_binder, + fragment_binder, + sampler_argument_buffer, + topology: info.topology, + depth_stencil_state, + raster_info: info.raster_info, }) }) } @@ -1266,8 +2569,12 @@ impl super::Device for Device { heap: &mut Heap ) -> result::Result { objc::rc::autoreleasepool(|| { + // StorageModeShared: CPU and GPU share the same physical memory — no didModifyRange + // needed and no stale-copy hazard. StorageModeManaged has a separate GPU copy that + // requires an explicit sync notification after every CPU write; without it the GPU + // reads stale data, causing tearing let opt = metal::MTLResourceOptions::CPUCacheModeDefaultCache | - metal::MTLResourceOptions::StorageModeManaged; + metal::MTLResourceOptions::StorageModeShared; let byte_len = (info.stride * info.num_elements) as NSUInteger; @@ -1279,39 +2586,58 @@ impl super::Device for Device { self.metal_device.new_buffer(byte_len, opt) }; - Ok(Buffer{ - metal_buffer: buf, - element_stride: info.stride - }) - }) - } + // allocate on the heap + let alloc_index = heap.allocate(); + heap.buffer_slots[alloc_index] = Some(buf.to_owned()); + heap.encode_buffer(alloc_index, &buf); - fn create_buffer( - &mut self, - info: &super::BufferInfo, - data: Option<&[T]>, - ) -> result::Result { - objc::rc::autoreleasepool(|| { - let opt = metal::MTLResourceOptions::CPUCacheModeDefaultCache | - metal::MTLResourceOptions::StorageModeManaged; + // assign srv or uav + let srv_index = if info.usage.contains(BufferUsage::SHADER_RESOURCE) { + Some(alloc_index) + } + else { + None + }; - let byte_len = (info.stride * info.num_elements) as NSUInteger; + let uav_index = if info.usage.contains(BufferUsage::UNORDERED_ACCESS) { + Some(alloc_index) + } + else { + None + }; - let buf = if let Some(data) = data { - let bytes = data.as_ptr() as *const std::ffi::c_void; - self.metal_device.new_buffer_with_data(bytes, byte_len, opt) + let cbv_index = if info.usage.contains(BufferUsage::CONSTANT_BUFFER) { + Some(alloc_index) } else { - self.metal_device.new_buffer(byte_len, opt) + None }; Ok(Buffer{ metal_buffer: buf, - element_stride: info.stride + element_stride: info.stride, + srv_index, + uav_index, + cbv_index, + counter_sample_buffer: None, + counter_sample_index: 0, + counter_cmd: None, }) }) } + fn create_buffer( + &mut self, + info: &super::BufferInfo, + data: Option<&[T]>, + ) -> result::Result { + self.create_buffer_with_heap( + info, + data, + &mut self.shader_heap.clone() + ) + } + fn create_read_back_buffer( &mut self, size: usize, @@ -1320,12 +2646,19 @@ impl super::Device for Device { let opt = metal::MTLResourceOptions::CPUCacheModeDefaultCache | metal::MTLResourceOptions::StorageModeManaged; - let byte_len = size as NSUInteger; + // Metal doesn't allow zero-size buffers + let byte_len = size.max(1) as NSUInteger; let buf = self.metal_device.new_buffer(byte_len, opt); Ok(Buffer{ metal_buffer: buf, - element_stride: size + element_stride: size, + srv_index: None, + uav_index: None, + cbv_index: None, + counter_sample_buffer: None, + counter_sample_index: 0, + counter_cmd: None, }) }) } @@ -1335,74 +2668,194 @@ impl super::Device for Device { info: &super::TextureInfo, data: Option<&[T]>, ) -> result::Result { + self.create_texture_with_heaps( + info, + TextureHeapInfo::default(), + data, + ) + } + + fn create_texture_with_heaps( + &mut self, + info: &TextureInfo, + heaps: TextureHeapInfo, + data: Option<&[T]>, + ) -> result::Result { objc::rc::autoreleasepool(|| { let desc = TextureDescriptor::new(); - // TODO: - // tex_type - // format - // initial_state + // clamp requested MSAA to what the device supports (eg. 8x -> 4x on most Apple GPUs) + let sample_count = self.supported_sample_count(info.samples); + let msaa = sample_count > 1; // desc - desc.set_pixel_format(metal::MTLPixelFormat::RGBA8Unorm); // TODO: format - + desc.set_pixel_format(to_mtl_pixel_format(info.format)); desc.set_width(info.width as NSUInteger); desc.set_height(info.height as NSUInteger); desc.set_depth(info.depth as NSUInteger); - desc.set_array_length(info.array_layers as NSUInteger); - desc.set_mipmap_level_count(info.mip_levels as NSUInteger); - desc.set_sample_count(info.samples as NSUInteger); + // MSAA textures cannot have a mip chain + desc.set_mipmap_level_count(if msaa { 1 } else { info.mip_levels as NSUInteger }); desc.set_usage(to_mtl_texture_usage(info.usage)); - desc.set_storage_mode(metal::MTLStorageMode::Shared); - desc.set_texture_type(metal::MTLTextureType::D2); + // Must match the (Private) heap the texture is allocated from + desc.set_storage_mode(metal::MTLStorageMode::Private); + // MSAA Texture2D uses the D2Multisample type + desc.set_texture_type(if msaa && matches!(info.tex_type, super::TextureType::Texture2D) { + metal::MTLTextureType::D2Multisample + } else { + to_mtl_texture_type(info.tex_type) + }); + + // For cubemaps, arrayLength must be 1 (6 faces are implicit) + // For cube arrays, arrayLength is the number of cubemaps (not faces) + let array_length = match info.tex_type { + super::TextureType::TextureCube => 1, + super::TextureType::TextureCubeArray => info.array_layers / 6, + _ => info.array_layers, + }; + desc.set_array_length(array_length as NSUInteger); + + desc.set_sample_count(sample_count as NSUInteger); + + // use supplied heap or fallback to the device default + let shader_heap = if let Some(shader_heap) = heaps.shader { + shader_heap + } + else { + &mut self.shader_heap + }; // heap bindless - let tex = self.shader_heap.mtl_heap.new_texture(&desc) + let tex = shader_heap.mtl_heap.new_texture(&desc) .expect("hotline_rs::gfx::mtl failed to allocate texture in heap!"); - // data + // upload texture data with support for mips, cubemaps, and array slices. + // The heap is Private (not CPU-writable), so stage the bytes in a Shared buffer and + // blit each subresource into the texture on a one-shot command buffer. if let Some(data) = data { - tex.replace_region( - metal::MTLRegion { - origin: metal::MTLOrigin { x: 0, y: 0, z: 0 }, - size: metal::MTLSize { - width: info.width, - height: info.height, - depth: info.depth as u64, - }, - }, - 0, - data.as_ptr() as _, - info.width * 4, // TODO size from format + let block_size = super::block_size_for_format(info.format) as u64; + let tpb = super::texels_per_block_for_format(info.format); + + let bytes = unsafe { + std::slice::from_raw_parts( + data.as_ptr() as *const u8, + std::mem::size_of_val(data) + ) + }; + let staging = self.metal_device.new_buffer_with_data( + bytes.as_ptr() as *const std::ffi::c_void, + bytes.len() as NSUInteger, + metal::MTLResourceOptions::StorageModeShared ); + + let cmd = self.command_queue.new_command_buffer(); + let blit = cmd.new_blit_command_encoder(); + + let mut data_offset: u64 = 0; + for a in 0..info.array_layers { + let mut mip_w = info.width; + let mut mip_h = info.height; + let mut mip_d = info.depth as u64; + + for mip in 0..info.mip_levels { + let pitch = block_size * (mip_w / tpb).max(1); + let depth_pitch = pitch * (mip_h / tpb).max(1); + + blit.copy_from_buffer_to_texture( + &staging, + data_offset as NSUInteger, + pitch as NSUInteger, + depth_pitch as NSUInteger, + metal::MTLSize { width: mip_w, height: mip_h, depth: mip_d }, + &tex, + a as NSUInteger, + mip as NSUInteger, + metal::MTLOrigin { x: 0, y: 0, z: 0 }, + metal::MTLBlitOption::empty(), + ); + + data_offset += depth_pitch * mip_d.max(1); + + // halve dimensions for next mip (non-pot safe) + mip_w = (mip_w / 2).max(1); + mip_h = (mip_h / 2).max(1); + mip_d = (mip_d / 2).max(1); + } + } + + blit.end_encoding(); + cmd.commit(); + cmd.wait_until_completed(); } - // srv - let srv_index = self.shader_heap.allocate(); - self.shader_heap.texture_slots[srv_index] = Some(tex.to_owned()); + // allocate on the heap + let alloc_index = shader_heap.allocate(); + shader_heap.texture_slots[alloc_index] = Some(tex.to_owned()); - Ok(Texture{ - metal_texture: tex, - srv_index: Some(srv_index), - heap_id: Some(self.shader_heap.id) - }) - }) - } + // Encode texture into heap's argument buffer for bindless access + shader_heap.encode_texture(alloc_index, &tex); - fn create_texture_with_heaps( - &mut self, - info: &TextureInfo, - heaps: TextureHeapInfo, - data: Option<&[T]>, - ) -> result::Result { - objc::rc::autoreleasepool(|| { - let desc = TextureDescriptor::new(); - let tex = self.metal_device.new_texture(&desc); - Ok(Texture{ - metal_texture: tex, - srv_index: None, - heap_id: Some(self.shader_heap.id) - }) + let shader_resource = info.usage.contains(TextureUsage::SHADER_RESOURCE); + + // UAV only applies to the (non-MSAA) texture + let uav_index = if info.usage.contains(TextureUsage::UNORDERED_ACCESS) { + Some(alloc_index) + } + else { + None + }; + + if msaa { + // The primary texture is the MSAA view (read as Texture2DMS via get_msaa_srv_index). + let msaa_srv_index = if shader_resource { Some(alloc_index) } else { None }; + + // Create a single-sample resolve backing so the texture can be read normally and + // resolved via resolve_texture_subresource (matches the D3D12 resolve concept). + let mut resolved_texture = None; + let mut srv_index = None; + if shader_resource { + let rdesc = TextureDescriptor::new(); + rdesc.set_pixel_format(to_mtl_pixel_format(info.format)); + rdesc.set_width(info.width as NSUInteger); + rdesc.set_height(info.height as NSUInteger); + rdesc.set_depth(info.depth as NSUInteger); + rdesc.set_mipmap_level_count(info.mip_levels as NSUInteger); + rdesc.set_usage(to_mtl_texture_usage(info.usage)); + rdesc.set_storage_mode(metal::MTLStorageMode::Private); + rdesc.set_texture_type(to_mtl_texture_type(info.tex_type)); + rdesc.set_array_length(array_length as NSUInteger); + rdesc.set_sample_count(1); + + let resolve_tex = shader_heap.mtl_heap.new_texture(&rdesc) + .expect("hotline_rs::gfx::mtl failed to allocate resolve texture in heap!"); + let resolve_index = shader_heap.allocate(); + shader_heap.texture_slots[resolve_index] = Some(resolve_tex.to_owned()); + shader_heap.encode_texture(resolve_index, &resolve_tex); + srv_index = Some(resolve_index); + resolved_texture = Some(resolve_tex); + } + + Ok(Texture{ + metal_texture: tex, + resolved_texture, + srv_index, + msaa_srv_index, + uav_index, + resolvable: shader_resource, + heap_id: Some(shader_heap.id) + }) + } + else { + let srv_index = if shader_resource { Some(alloc_index) } else { None }; + Ok(Texture{ + metal_texture: tex, + resolved_texture: None, + srv_index, + msaa_srv_index: None, + uav_index, + resolvable: false, + heap_id: Some(shader_heap.id) + }) + } }) } @@ -1414,24 +2867,85 @@ impl super::Device for Device { // new desc let descriptor = metal::RenderPassDescriptor::new(); - // colour attachments - for rt in &info.render_targets { - let color_attachment = descriptor.color_attachments().object_at(0).unwrap(); + // colour attachments - one per MRT target (SV_Target0..N) + let mut pixel_formats = Vec::new(); + for (i, rt) in info.render_targets.iter().enumerate() { + let color_attachment = descriptor.color_attachments().object_at(i as u64).unwrap(); color_attachment.set_texture(Some(&rt.metal_texture)); + color_attachment.set_slice(info.array_slice as u64); if let Some(cc) = info.rt_clear { color_attachment.set_load_action(metal::MTLLoadAction::Clear); color_attachment.set_clear_color(metal::MTLClearColor::new(cc.r as f64, cc.g as f64, cc.b as f64, 1.0)); - color_attachment.set_store_action(metal::MTLStoreAction::Store); } else { color_attachment.set_load_action(metal::MTLLoadAction::Load); - color_attachment.set_store_action(metal::MTLStoreAction::Store); } + + // Keep the rendered (MSAA) samples. The MSAA resolve and any mip downsample are + // driven by the render graph barriers (see resolve_texture_subresource / + // generate_mip_maps), not baked into every pass, so the barrier can decide when they + // happen (eg. only after the last of several passes that target the same resource). + color_attachment.set_store_action(metal::MTLStoreAction::Store); + + pixel_formats.push(rt.metal_texture.pixel_format()); } + // sample count shared by all attachments (read from the first colour/depth target) + let sample_count = info.render_targets.first() + .map(|rt| rt.metal_texture.sample_count() as u32) + .or_else(|| info.depth_stencil.map(|ds| ds.metal_texture.sample_count() as u32)) + .unwrap_or(1); + + // Handle depth stencil attachment + let depth_format = if let Some(ds_texture) = &info.depth_stencil { + let depth_attachment = descriptor.depth_attachment().unwrap(); + depth_attachment.set_texture(Some(&ds_texture.metal_texture)); + depth_attachment.set_slice(info.array_slice as u64); + + if let Some(ds_clear) = &info.ds_clear { + if let Some(depth_val) = ds_clear.depth { + depth_attachment.set_load_action(metal::MTLLoadAction::Clear); + depth_attachment.set_clear_depth(depth_val as f64); + } else { + depth_attachment.set_load_action(metal::MTLLoadAction::Load); + } + } else { + depth_attachment.set_load_action(metal::MTLLoadAction::Load); + } + depth_attachment.set_store_action(metal::MTLStoreAction::Store); + + let format = ds_texture.metal_texture.pixel_format(); + + // Handle stencil if format has stencil component + if has_stencil_component(format) { + let stencil_attachment = descriptor.stencil_attachment().unwrap(); + stencil_attachment.set_texture(Some(&ds_texture.metal_texture)); + stencil_attachment.set_slice(info.array_slice as u64); + + if let Some(ds_clear) = &info.ds_clear { + if let Some(stencil_val) = ds_clear.stencil { + stencil_attachment.set_load_action(metal::MTLLoadAction::Clear); + stencil_attachment.set_clear_stencil(stencil_val as u32); + } else { + stencil_attachment.set_load_action(metal::MTLLoadAction::Load); + } + } else { + stencil_attachment.set_load_action(metal::MTLLoadAction::Load); + } + stencil_attachment.set_store_action(metal::MTLStoreAction::Store); + } + + Some(format) + } else { + None + }; + Ok(RenderPass{ - desc: descriptor.to_owned() + desc: descriptor.to_owned(), + pixel_formats, + depth_format, + sample_count, }) }) } @@ -1457,19 +2971,40 @@ impl super::Device for Device { unimplemented!() } - fn create_raytracing_tlas( - &mut self, - info: &RaytracingTLASInfo - ) -> result::Result { - unimplemented!() - } - fn create_compute_pipeline( &self, info: &super::ComputePipelineInfo, ) -> result::Result { - Ok(ComputePipeline{ - slots: Vec::new() + objc::rc::autoreleasepool(|| { + // load the compute kernel function from the shader library + let function = unsafe { + let cs = info.cs; + let lib = self.metal_device.new_library_with_data( + std::slice::from_raw_parts(cs.data, cs.data_size) + )?; + let name = &lib.function_names()[0]; + lib.get_function(name, None).unwrap() + }; + + let pipeline_state = self.metal_device.new_compute_pipeline_state_with_function(&function)?; + + // unified slot lookup + single-stage binder, both keyed by (register, space, type) + let slot_lookup = self.build_slot_lookup( + &info.pipeline_layout.bindings, + &info.pipeline_layout.push_constants, + ); + + let compute_binder = self.build_compute_binder( + &info.pipeline_layout.bindings, + &info.pipeline_layout.push_constants, + ); + + Ok(ComputePipeline { + pipeline_state, + slots: Vec::new(), + slot_lookup, + compute_binder, + }) }) } @@ -1482,7 +3017,58 @@ impl super::Device for Device { } fn execute(&mut self, cmd: &CmdBuf) { + // Pass command buffers commit themselves in CmdBuf::close, so there is nothing to submit + // here for them. Barrier command buffers instead carry deferred ops (transition / resolve / + // generate mips) which we replay into a fresh command buffer every frame, mirroring how + // D3D12 re-executes a pre-recorded barrier command list. + if cmd.deferred_ops.is_empty() { + return; + } + objc::rc::autoreleasepool(|| { + let metal_cmd = self.command_queue.new_command_buffer(); + for op in &cmd.deferred_ops { + match op { + DeferredBarrierOp::Resolve { msaa, resolve } => { + // a load/no-clear pass with a MultisampleResolve store action resolves the + // MSAA samples into the single-sample backing without drawing anything. + let descriptor = metal::RenderPassDescriptor::new(); + // depth/stencil targets must resolve through the depth (and stencil) + // attachments, not a color attachment - a depth format on color + // attachment 0 is "not color renderable" and trips Metal validation, + // blocking GPU captures. + if is_depth_format(msaa.pixel_format()) { + let depth = descriptor.depth_attachment().unwrap(); + depth.set_texture(Some(msaa)); + depth.set_resolve_texture(Some(resolve)); + depth.set_load_action(metal::MTLLoadAction::Load); + depth.set_store_action(metal::MTLStoreAction::MultisampleResolve); + if has_stencil_component(msaa.pixel_format()) { + let stencil = descriptor.stencil_attachment().unwrap(); + stencil.set_texture(Some(msaa)); + stencil.set_resolve_texture(Some(resolve)); + stencil.set_load_action(metal::MTLLoadAction::Load); + stencil.set_store_action(metal::MTLStoreAction::MultisampleResolve); + } + } else { + let attachment = descriptor.color_attachments().object_at(0).unwrap(); + attachment.set_texture(Some(msaa)); + attachment.set_resolve_texture(Some(resolve)); + attachment.set_load_action(metal::MTLLoadAction::Load); + attachment.set_store_action(metal::MTLStoreAction::StoreAndMultisampleResolve); + } + let encoder = metal_cmd.new_render_command_encoder(&descriptor); + encoder.end_encoding(); + } + DeferredBarrierOp::GenerateMips { texture } => { + let blit = metal_cmd.new_blit_command_encoder(); + blit.generate_mipmaps(texture); + blit.end_encoding(); + } + } + } + metal_cmd.commit(); + }); } fn report_live_objects(&self) -> result::Result<(), super::Error> { @@ -1513,7 +3099,49 @@ impl super::Device for Device { None } - fn read_timestamps(&self, swap_chain: &SwapChain, buffer: &Self::Buffer, size_bytes: usize, frame_written_fence: u64) -> Vec { + fn read_timestamps(&self, _swap_chain: &SwapChain, buffer: &Self::Buffer, _size_bytes: usize, _frame_written_fence: u64) -> Vec { + // Metal has no GPU-signalled fence; wait for the pass command buffer to finish before reading + // its timestamps (equivalent to D3D12's GPU fence check). + if let Some(cmd) = &buffer.counter_cmd { + cmd.wait_until_completed(); + } + + if let Some(sample_buffer) = &buffer.counter_sample_buffer { + // counter-sampling path: resolve the one timestamp this buffer points at. The GPU + // timestamp is in nanoseconds on Apple Silicon; gather_stats wants seconds. + unsafe { + let range = metal::NSRange { + location: buffer.counter_sample_index as _, + length: 1, + }; + let ns_data: *mut objc::runtime::Object = + msg_send![sample_buffer.as_ref(), resolveCounterRange: range]; + if !ns_data.is_null() { + let bytes: *const u8 = msg_send![ns_data, bytes]; + let len: usize = msg_send![ns_data, length]; + if len >= std::mem::size_of::() { + let nanos = (bytes as *const u64).read_unaligned(); + // MTLCounterErrorValue marks a sample the GPU could not record - treat as none + if nanos != u64::MAX { + return vec![nanos as f64 / 1_000_000_000.0]; + } + } + } + } + return vec![]; + } + + // fallback path: whole-CB timing. index 0 = start of pass, index 1 = end of pass. + if let Some(cmd) = &buffer.counter_cmd { + let seconds: f64 = unsafe { + if buffer.counter_sample_index == 0 { + msg_send![cmd.as_ref(), GPUStartTime] + } else { + msg_send![cmd.as_ref(), GPUEndTime] + } + }; + return vec![seconds]; + } vec![] } @@ -1522,7 +3150,7 @@ impl super::Device for Device { } fn get_timestamp_size_bytes() -> usize { - 0 + 8 // u64; matches D3D12 — Metal uses CounterSampleBuffer, not this backing store } fn get_pipeline_statistics_size_bytes() -> usize { @@ -1536,6 +3164,45 @@ impl super::Device for Device { fn get_counter_alignment() -> usize { 0 } + + fn create_upload_buffer( + &mut self, + data: &[T] + ) -> Result { + unimplemented!() + } + + fn create_raytracing_instance_buffer( + &mut self, + instances: &Vec> + ) -> Result { + unimplemented!() + } + + fn create_raytracing_tlas( + &mut self, + info: &RaytracingTLASInfo + ) -> Result { + unimplemented!() + } + + fn create_resource_view( + &mut self, + info: &ResourceViewInfo, + resource: Resource, + heap: &mut Heap + ) -> Result { + unimplemented!() + } + + fn create_raytracing_tlas_with_heap( + &mut self, + info: &RaytracingTLASInfo, + heap: &mut Heap + ) -> Result { + unimplemented!() + } + } unsafe impl Send for Device {} diff --git a/src/image.rs b/src/image.rs index a8316add..653ab206 100644 --- a/src/image.rs +++ b/src/image.rs @@ -29,7 +29,7 @@ pub fn load_from_file(filename: &str) -> Result { let mut f = fs::File::open(path).expect("hotline_rs::image:: File not found"); // dds file if filename.ends_with(".dds") { - let dds = DDS::read(f)?; + let dds = DDS::read(f)?; Ok(ImageData { info: TextureInfo { tex_type: to_gfx_texture_type(&dds), @@ -72,7 +72,7 @@ pub fn load_from_file(filename: &str) -> Result { let data_size_bytes = x * y * 4; data_out.resize(data_size_bytes as usize, 0); std::ptr::copy_nonoverlapping(img, data_out.as_mut_ptr(), data_size_bytes as usize); - + // cleanup stb_image_rust::c_runtime::free(img); @@ -96,14 +96,13 @@ pub fn load_from_file(filename: &str) -> Result { } /// Loads an image from file and creates a shader resource on the specified heap, or on the device heap if `heap.is_none()` -#[cfg(target_os = "windows")] pub fn load_texture_from_file( device: &mut crate::gfx_platform::Device, file: &str, heap: Option<&mut crate::gfx_platform::Heap>) -> Result { let image = load_from_file(file)?; device.create_texture_with_heaps( - &image.info, + &image.info, gfx::TextureHeapInfo { shader: heap, ..Default::default() @@ -391,7 +390,7 @@ fn to_gfx_format(dds: &DDS) -> gfx::Format { DxgiFormat::Y216 => panic!(), DxgiFormat::NV11 => panic!(), DxgiFormat::AI44 => panic!(), - DxgiFormat::IA44 => panic!(), + DxgiFormat::IA44 => panic!(), DxgiFormat::P8 => panic!(), DxgiFormat::A8P8 => panic!(), DxgiFormat::B4G4R4A4_UNorm => panic!(), diff --git a/src/imgui.rs b/src/imgui.rs index a38551df..922a15b0 100644 --- a/src/imgui.rs +++ b/src/imgui.rs @@ -13,6 +13,7 @@ use crate::gfx::CmdBuf; use crate::gfx::Device; use crate::gfx::SwapChain; use crate::gfx::Texture; +use crate::gfx::Pipeline; use maths_rs::Vec4f; @@ -454,6 +455,7 @@ fn render_draw_data( cmd.set_vertex_buffer(&buffers.vb, 0); cmd.set_index_buffer(&buffers.ib); cmd.set_render_pipeline(pipeline); + cmd.push_render_constants(pipeline, 0, 0, 16, 0, &mvp); let clip_off = draw_data.DisplayPos; @@ -492,7 +494,7 @@ fn render_draw_data( cmd.set_binding(pipeline, 0, 0, gfx::DescriptorType::ShaderResource, device.get_shader_heap(), srv); } else { - // bund srv in another heap + // bound srv in another heap for heap in image_heaps { if heap.get_heap_id() == heap_id { cmd.set_binding(pipeline, 0, 0, gfx::DescriptorType::ShaderResource, heap, srv); @@ -591,7 +593,7 @@ impl ImGui where D: Device, A: App, D::RenderPipeline: gfx::Pipeline let io = &mut *igGetIO(); io.ConfigFlags |= ImGuiConfigFlags_DockingEnable as i32; - io.ConfigFlags |= ImGuiConfigFlags_ViewportsEnable as i32; + // io.ConfigFlags |= ImGuiConfigFlags_ViewportsEnable as i32; // construct path for ini to be along side the exe let exe_path = std::env::current_exe().ok().unwrap(); @@ -1568,10 +1570,23 @@ impl ImGui where D: Device, A: App, D::RenderPipeline: gfx::Pipeline impl Drop for ImGui where D: Device, A: App { fn drop(&mut self) { unsafe { - igDestroyPlatformWindows(); + // Clean up main viewport's user data (set during create) + let main_vp = &mut *igGetMainViewport(); + if !main_vp.PlatformUserData.is_null() { + std::ptr::drop_in_place(main_vp.PlatformUserData as *mut ViewportData); + main_vp.PlatformUserData = std::ptr::null_mut(); + } + if !main_vp.PlatformHandle.is_null() { + std::ptr::drop_in_place(main_vp.PlatformHandle as *mut A::NativeHandle); + main_vp.PlatformHandle = std::ptr::null_mut(); + } + let platform_io = &mut *igGetPlatformIO(); std::ptr::drop_in_place(platform_io.Monitors.Data as *mut ImGuiPlatformMonitor); platform_io.Monitors.Data = std::ptr::null_mut(); + + // Destroy non-main viewport windows + igDestroyPlatformWindows(); } } } diff --git a/src/lib.rs b/src/lib.rs index 97061d60..8eb23049 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -242,16 +242,33 @@ pub mod prelude { // modules gfx, os, + client, + plugin, pmfx, imgui, + image, + + // platform specific + gfx_platform, + os_platform, // traits + ecs_base::*, gfx::{Device, SwapChain, CmdBuf, Texture, RenderPass, Pipeline, Buffer}, + pmfx::{DrawData, MaterialData, PointLightData, SpotLightData, DirectionalLightData, WorldBufferReserveInfo, WorldBufferInfo}, os::{App, Window}, pmfx::Pmfx, imgui::ImGui, imdraw::ImDraw, + client::{Client, HotlineInfo, PluginInfo}, + plugin::{Plugin}, av::{VideoPlayer}, + + // macros + hotline_plugin, + system_func, + demos, + systems }; } @@ -261,7 +278,7 @@ pub use os::macos as os_platform; /// This is a hardcoded compile time selection of os backend for macos as null #[cfg(target_os = "macos")] -pub use gfx::null as gfx_platform; +pub use gfx::mtl as gfx_platform; /// This is a hardcoded compile time selection of os backend for macos as null #[cfg(target_os = "macos")] diff --git a/src/os/macos.rs b/src/os/macos.rs index 7f9c2a0d..6fb64f7e 100644 --- a/src/os/macos.rs +++ b/src/os/macos.rs @@ -3,21 +3,91 @@ extern crate objc; use core::panic; +use std::collections::HashMap; use std::time::Duration; use std::sync::Arc; use std::sync::RwLock; use winit::{ - dpi::{LogicalPosition, LogicalSize}, event::{Event, WindowEvent}, event_loop::{self, ControlFlow}, platform::pump_events::{EventLoopExtPumpEvents, PumpStatus}, raw_window_handle::{HasWindowHandle, RawWindowHandle} + dpi::{PhysicalPosition, PhysicalSize}, + event::{WindowEvent, ElementState}, + event_loop::ActiveEventLoop, + keyboard::{Key, PhysicalKey, KeyCode}, + raw_window_handle::{HasWindowHandle, RawWindowHandle}, }; -use cocoa::{appkit::{NSEvent, NSLeftArrowFunctionKey, NSNextFunctionKey, NSView}, base::id as cocoa_id }; +use cocoa::base::id as cocoa_id; use crate::os::Rect; +/// Input state tracking (similar to ProcData in win32.rs) +#[derive(Clone)] +struct InputState { + // Mouse state - screen coordinates + mouse_pos: super::Point, + mouse_pos_prev: super::Point, + // Mouse state - window-local coordinates (from CursorMoved) + mouse_client_pos: super::Point, + mouse_down: [bool; super::MouseButton::Count as usize], + mouse_wheel: f32, + mouse_hwheel: f32, + hovered_window_id: Option, + + // Keyboard state + key_down: [bool; 256], + key_press: [bool; 256], + key_debounce: [bool; 256], + + // System keys (Ctrl, Shift, Alt) + sys_key_down: [bool; super::SysKey::Count as usize], + sys_key_press: [bool; super::SysKey::Count as usize], + sys_key_debounce: [bool; super::SysKey::Count as usize], + + // Text input + utf16_inputs: Vec, + + // Input enabled flags + keyboard_enabled: bool, + mouse_enabled: bool, +} + +impl InputState { + fn new() -> Self { + InputState { + mouse_pos: super::Point { x: 0, y: 0 }, + mouse_pos_prev: super::Point { x: 0, y: 0 }, + mouse_client_pos: super::Point { x: 0, y: 0 }, + mouse_down: [false; super::MouseButton::Count as usize], + mouse_wheel: 0.0, + mouse_hwheel: 0.0, + hovered_window_id: None, + key_down: [false; 256], + key_press: [false; 256], + key_debounce: [false; 256], + sys_key_down: [false; super::SysKey::Count as usize], + sys_key_press: [false; super::SysKey::Count as usize], + sys_key_debounce: [false; super::SysKey::Count as usize], + utf16_inputs: Vec::new(), + keyboard_enabled: true, + mouse_enabled: true, + } + } +} + #[derive(Clone)] pub struct App { - event_loop: Arc>> + event_loop: Arc>>, + input_state: Arc>, + windows: Arc>>>, + monitors: Arc>>, + // Per-window cells, mirrored on `Window`. We cache state that's normally read from winit + // (scale, size, position) because the equivalent winit calls dispatch objc messages + // (`backingScaleFactor`, `frame`) to the wrong receiver on macOS and panic. + window_scales: Arc>>>>, + window_sizes: Arc>>>>>, + window_positions: Arc>>>>>, + /// When false, windows render at 1x (non-retina) for lower GPU cost (eg. with MSAA) + dpi_aware: bool, } unsafe impl Send for App {} @@ -25,7 +95,15 @@ unsafe impl Sync for App {} #[derive(Clone)] pub struct Window { - winit_window: Arc + winit_window: Arc, + window_id: winit::window::WindowId, + input_state: Arc>, + events: Arc>, + cached_scale: Arc>, + cached_size: Arc>>, + cached_position: Arc>>, + /// Inherited from `AppInfo.dpi_aware`; false = render at 1x (logical pixels) + dpi_aware: bool, } unsafe impl Send for Window {} @@ -56,6 +134,209 @@ pub fn isize_id_from_window(window: &Window) -> isize { } } +impl InputState { + /// Debounce logic for keys - updates press and debounce based on current down state + fn debounce_keys(&mut self) { + for i in 0..256 { + if self.key_down[i] && !self.key_press[i] && !self.key_debounce[i] { + // First frame key down: trigger press + self.key_press[i] = true; + self.key_debounce[i] = true; + } else if self.key_press[i] { + // Subsequent frames: clear press + self.key_press[i] = false; + } else if !self.key_down[i] { + // Key released: clear debounce + self.key_debounce[i] = false; + } + } + } + + /// Debounce logic for system keys + fn debounce_sys_keys(&mut self) { + for i in 0..super::SysKey::Count as usize { + if self.sys_key_down[i] && !self.sys_key_press[i] && !self.sys_key_debounce[i] { + self.sys_key_press[i] = true; + self.sys_key_debounce[i] = true; + } else if self.sys_key_press[i] { + self.sys_key_press[i] = false; + } else if !self.sys_key_down[i] { + self.sys_key_debounce[i] = false; + } + } + } +} + +impl App { + /// Update input state at the start of each frame + fn update_input_state(&self) { + let mut state = self.input_state.write().unwrap(); + + // Reset per-frame values + state.mouse_wheel = 0.0; + state.mouse_hwheel = 0.0; + state.utf16_inputs.clear(); + + // Store previous mouse position for delta calculation + // Current mouse_pos is computed on-demand in get_mouse_pos() + state.mouse_pos_prev = state.mouse_pos; + + // Compute current screen position from window + client pos. + // Cached position is physical; mouse_client_pos is already in render-coord units + // (logical when !dpi_aware), so scale position to match. + if let Some(window_id) = state.hovered_window_id { + if let Some(pos_cell) = self.window_positions.read().unwrap().get(&window_id) { + let pos = *pos_cell.read().unwrap(); + let scale = if self.dpi_aware { + 1.0 + } else { + self.window_scales.read().unwrap() + .get(&window_id) + .map(|s| (*s.read().unwrap()).max(1.0)) + .unwrap_or(1.0) + }; + state.mouse_pos = super::Point { + x: (pos.x as f64 / scale).round() as i32 + state.mouse_client_pos.x, + y: (pos.y as f64 / scale).round() as i32 + state.mouse_client_pos.y, + }; + } + } + + // Debounce keys + state.debounce_keys(); + state.debounce_sys_keys(); + } +} + +struct FrameHandler<'a> { + resume: &'a mut bool, + input_state: Arc>, + monitors: Arc>>, + window_scales: Arc>>>>, + window_sizes: Arc>>>>>, + window_positions: Arc>>>>>, + dpi_aware: bool, +} + +impl winit::application::ApplicationHandler for FrameHandler<'_> { + fn resumed(&mut self, elwt: &ActiveEventLoop) { + let primary = elwt.primary_monitor(); + *self.monitors.write().unwrap() = elwt.available_monitors().map(|m| { + let winit::dpi::PhysicalSize { width, height } = m.size(); + let winit::dpi::PhysicalPosition { x, y } = m.position(); + super::MonitorInfo { + rect: Rect { x, y, width: width as i32, height: height as i32 }, + client_rect: Rect { x, y, width: width as i32, height: height as i32 }, + dpi_scale: m.scale_factor() as f32, + primary: primary.as_ref() == Some(&m), + } + }).collect(); + } + + fn window_event(&mut self, _elwt: &ActiveEventLoop, window_id: winit::window::WindowId, event: WindowEvent) { + if let WindowEvent::ScaleFactorChanged { scale_factor, .. } = event { + if let Some(scale) = self.window_scales.read().unwrap().get(&window_id) { + *scale.write().unwrap() = scale_factor; + } + return; + } + if let WindowEvent::Resized(new_size) = event { + if let Some(size) = self.window_sizes.read().unwrap().get(&window_id) { + *size.write().unwrap() = new_size; + } + return; + } + if let WindowEvent::Moved(new_pos) = event { + if let Some(pos) = self.window_positions.read().unwrap().get(&window_id) { + *pos.write().unwrap() = new_pos; + } + return; + } + + let mut state = self.input_state.write().unwrap(); + match event { + WindowEvent::CloseRequested => { + *self.resume = false; + } + WindowEvent::RedrawRequested => {} + WindowEvent::CursorMoved { position, .. } => { + // winit reports physical pixels; convert to logical (1x) when !dpi_aware + // so coords match the render size returned to clients. + let scale = if self.dpi_aware { + 1.0 + } else { + self.window_scales.read().unwrap() + .get(&window_id) + .map(|s| (*s.read().unwrap()).max(1.0)) + .unwrap_or(1.0) + }; + state.mouse_client_pos = super::Point { + x: (position.x / scale).round() as i32, + y: (position.y / scale).round() as i32, + }; + state.hovered_window_id = Some(window_id); + } + WindowEvent::CursorEntered { .. } => { + state.hovered_window_id = Some(window_id); + } + WindowEvent::CursorLeft { .. } => { + if state.hovered_window_id == Some(window_id) { + state.hovered_window_id = None; + } + } + WindowEvent::MouseInput { state: element_state, button, .. } => { + let pressed = element_state == ElementState::Pressed; + let index = match button { + winit::event::MouseButton::Left => Some(super::MouseButton::Left as usize), + winit::event::MouseButton::Middle => Some(super::MouseButton::Middle as usize), + winit::event::MouseButton::Right => Some(super::MouseButton::Right as usize), + winit::event::MouseButton::Back => Some(super::MouseButton::X1 as usize), + winit::event::MouseButton::Forward => Some(super::MouseButton::X2 as usize), + winit::event::MouseButton::Other(_) => None, + }; + if let Some(idx) = index { + state.mouse_down[idx] = pressed; + } + } + WindowEvent::MouseWheel { delta, .. } => { + match delta { + winit::event::MouseScrollDelta::LineDelta(h, v) => { + state.mouse_wheel += v; + state.mouse_hwheel += h; + } + winit::event::MouseScrollDelta::PixelDelta(pos) => { + state.mouse_wheel += (pos.y / 20.0) as f32; + state.mouse_hwheel += (pos.x / 20.0) as f32; + } + } + } + WindowEvent::KeyboardInput { event, .. } => { + let pressed = event.state == ElementState::Pressed; + if let PhysicalKey::Code(key_code) = event.physical_key { + let code = key_code as usize; + if code < 256 { + state.key_down[code] = pressed; + } + } + if pressed { + if let Key::Character(ref c) = event.logical_key { + for ch in c.encode_utf16() { + state.utf16_inputs.push(ch); + } + } + } + } + WindowEvent::ModifiersChanged(modifiers) => { + let mods = modifiers.state(); + state.sys_key_down[super::SysKey::Ctrl as usize] = mods.control_key(); + state.sys_key_down[super::SysKey::Shift as usize] = mods.shift_key(); + state.sys_key_down[super::SysKey::Alt as usize] = mods.alt_key(); + } + _ => {} + } + } +} + impl super::App for App { type Window = Window; type NativeHandle = NativeHandle; @@ -63,20 +344,68 @@ impl super::App for App { /// Create an application instance fn create(info: super::AppInfo) -> Self { App { - event_loop: Arc::new(RwLock::new(winit::event_loop::EventLoop::new().unwrap())) + event_loop: Arc::new(RwLock::new(winit::event_loop::EventLoop::new().unwrap())), + input_state: Arc::new(RwLock::new(InputState::new())), + windows: Arc::new(RwLock::new(HashMap::new())), + monitors: Arc::new(RwLock::new(Vec::new())), + window_scales: Arc::new(RwLock::new(HashMap::new())), + window_sizes: Arc::new(RwLock::new(HashMap::new())), + window_positions: Arc::new(RwLock::new(HashMap::new())), + dpi_aware: info.dpi_aware, } } /// Create a new operating system window fn create_window(&mut self, info: super::WindowInfo) -> Self::Window { - let window = winit::window::WindowBuilder::new() - .with_inner_size(winit::dpi::LogicalSize::new(info.rect.width, info.rect.height)) - .with_position(winit::dpi::LogicalPosition::new(info.rect.x, info.rect.y)) - .with_title(info.title) - .build(&*self.event_loop.read().unwrap()) + #[allow(deprecated)] + let window = self.event_loop.read().unwrap() + .create_window( + winit::window::Window::default_attributes() + .with_inner_size(PhysicalSize::new(info.rect.width, info.rect.height)) + .with_position(PhysicalPosition::new(info.rect.x, info.rect.y)) + .with_title(info.title) + ) .unwrap(); + let window_id = window.id(); + let winit_window = Arc::new(window); + + // Register window for position lookups + self.windows.write().unwrap().insert(window_id, winit_window.clone()); + + // Seed scale from the primary monitor; updated by ScaleFactorChanged events. + // Avoid calling winit_window.scale_factor() here — on macOS it can dispatch + // backingScaleFactor to the application delegate and panic. + let initial_scale = self.monitors.read().unwrap() + .iter() + .find(|m| m.primary) + .map(|m| m.dpi_scale as f64) + .unwrap_or(1.0); + let cached_scale = Arc::new(RwLock::new(initial_scale)); + self.window_scales.write().unwrap().insert(window_id, cached_scale.clone()); + + // Seed size/position from creation info; updated by Resized / Moved events. + // Avoid winit_window.inner_size()/outer_position() — both dispatch `frame` + // through paths that can panic on macOS for the same objc-receiver reason. + let cached_size = Arc::new(RwLock::new(PhysicalSize::new( + info.rect.width.max(0) as u32, + info.rect.height.max(0) as u32, + ))); + self.window_sizes.write().unwrap().insert(window_id, cached_size.clone()); + + let cached_position = Arc::new(RwLock::new(PhysicalPosition::new( + info.rect.x, info.rect.y, + ))); + self.window_positions.write().unwrap().insert(window_id, cached_position.clone()); + Window { - winit_window: Arc::new(window) + winit_window, + window_id, + input_state: self.input_state.clone(), + events: Arc::new(RwLock::new(super::WindowEventFlags::NONE)), + cached_scale, + cached_size, + cached_position, + dpi_aware: self.dpi_aware, } } @@ -88,26 +417,22 @@ impl super::App for App { /// Call to update windows and os state each frame, when false is returned the app has been requested to close fn run(&mut self) -> bool { objc::rc::autoreleasepool(|| { - let mut resume = true; - let _ = self.event_loop.write().and_then(|mut event_loop| { - let status = event_loop.pump_events(Some(Duration::ZERO), |event, elwt| { - match event { - Event::WindowEvent { event, .. } => match event { - WindowEvent::CloseRequested => { - resume = false; - } - WindowEvent::RedrawRequested => { - } - _ => { - - } - } - _ => { + self.update_input_state(); - } - } - }); + let mut resume = true; + let mut handler = FrameHandler { + resume: &mut resume, + input_state: self.input_state.clone(), + monitors: self.monitors.clone(), + window_scales: self.window_scales.clone(), + window_sizes: self.window_sizes.clone(), + window_positions: self.window_positions.clone(), + dpi_aware: self.dpi_aware, + }; + let _ = self.event_loop.write().and_then(|mut event_loop| { + use winit::platform::pump_events::EventLoopExtPumpEvents; + event_loop.pump_app_events(Some(Duration::ZERO), &mut handler); Ok(()) }); resume @@ -121,70 +446,57 @@ impl super::App for App { /// Retuns the mouse in screen coordinates fn get_mouse_pos(&self) -> super::Point { - // TODO: - super::Point { - x: 0, - y: 0 - } + self.input_state.read().unwrap().mouse_pos } /// Retuns the mouse vertical wheel position fn get_mouse_wheel(&self) -> f32 { - //panic!(); - // TODO: - 0.0 + self.input_state.read().unwrap().mouse_wheel } /// Retuns the mouse horizontal wheel positions fn get_mouse_hwheel(&self) -> f32 { - // panic!(); - 0.0 + self.input_state.read().unwrap().mouse_hwheel } /// Retuns the mouse button states, up or down fn get_mouse_buttons(&self) -> [bool; super::MouseButton::Count as usize] { - // panic!(); - [false; 5] + self.input_state.read().unwrap().mouse_down } /// Returns the distance the mouse has moved since the last frame fn get_mouse_pos_delta(&self) -> super::Size { - // panic!(); + let state = self.input_state.read().unwrap(); super::Size { - x: 0, - y: 0 + x: state.mouse_pos.x - state.mouse_pos_prev.x, + y: state.mouse_pos.y - state.mouse_pos_prev.y, } } /// Returns a vector of utf-16 characters that have been input since the last frame fn get_utf16_input(&self) -> Vec { - // panic!(); - vec![] + self.input_state.read().unwrap().utf16_inputs.clone() } /// Returns an array of bools containing 0-256 keys down (true) or up (false) fn get_keys_down(&self) -> [bool; 256] { - // panic!(); - [false; 256] + self.input_state.read().unwrap().key_down } /// Returns an array of bools containing 0-256 of keys pressed, will trigger only once and then require debouce fn get_keys_pressed(&self) -> [bool; 256] { - // panic!(); - [false; 256] + self.input_state.read().unwrap().key_press } /// Returns true if the sys key is down and false if the key is up fn is_sys_key_down(&self, key: super::SysKey) -> bool { - // panic!(); - false + self.input_state.read().unwrap().sys_key_down[key as usize] } /// Returns true if the sys key is pressed this frame and /// requires debounce until it is pressed again fn is_sys_key_pressed(&self, key: super::SysKey) -> bool { - // panic!(); - false + self.input_state.read().unwrap().sys_key_press[key as usize] } /// Get os system virtual key code from Key @@ -211,36 +523,42 @@ impl super::App for App { /// Set's whethere input from keybpard or mouse is available or not fn set_input_enabled(&mut self, keyboard: bool, mouse: bool) { - panic!(); + let mut state = self.input_state.write().unwrap(); + state.keyboard_enabled = keyboard; + state.mouse_enabled = mouse; } /// Get value for whether (keyboard, mouse) input is enabled fn get_input_enabled(&self) -> (bool, bool) { - panic!(); - (false, false) + let state = self.input_state.read().unwrap(); + (state.keyboard_enabled, state.mouse_enabled) } fn enumerate_display_monitors(&self) -> Vec { - let event_loop = &*self.event_loop.read().unwrap(); - let primary_monitor = event_loop.primary_monitor(); - event_loop.available_monitors().map(|monitor| { - let winit::dpi::PhysicalSize { width, height } = monitor.size(); - let winit::dpi::PhysicalPosition { x, y } = monitor.position(); - super::MonitorInfo { - rect: Rect { - x, y, - width: width as i32, - height: height as i32 - }, - client_rect: Rect { - x, y, - width: width as i32, - height: height as i32 - }, - dpi_scale: monitor.scale_factor() as f32, - primary: primary_monitor.as_ref() == Some(&monitor) + { + let cached = self.monitors.read().unwrap(); + if !cached.is_empty() { + return cached.clone(); } - }).collect() + } + // Monitors aren't populated until the first pump_app_events triggers `resumed`. + // Pump once here so the cache gets filled before the caller needs it. + let mut dummy_resume = true; + let mut handler = FrameHandler { + resume: &mut dummy_resume, + input_state: self.input_state.clone(), + monitors: self.monitors.clone(), + window_scales: self.window_scales.clone(), + window_sizes: self.window_sizes.clone(), + window_positions: self.window_positions.clone(), + dpi_aware: self.dpi_aware, + }; + let _ = self.event_loop.write().and_then(|mut event_loop| { + use winit::platform::pump_events::EventLoopExtPumpEvents; + event_loop.pump_app_events(Some(Duration::ZERO), &mut handler); + Ok(()) + }); + self.monitors.read().unwrap().clone() } /// Sets the mouse cursor @@ -266,7 +584,25 @@ impl super::App for App { /// Sets the console window rect that belongs to this app fn set_console_window_rect(&self, rect: super::Rect) { - panic!(); + // stub + + } +} + +impl Window { + /// Render-target size for this window: physical pixels when dpi-aware, otherwise the logical + /// 1x size (physical / scale_factor) so retina + MSAA isn't allocated at full backing scale. + fn render_size(&self) -> super::Size { + let size = *self.cached_size.read().unwrap(); + if self.dpi_aware { + super::Size { x: size.width as i32, y: size.height as i32 } + } else { + let scale = (*self.cached_scale.read().unwrap()).max(1.0); + super::Size { + x: ((size.width as f64 / scale).round() as i32).max(1), + y: ((size.height as f64 / scale).round() as i32).max(1), + } + } } } @@ -322,8 +658,8 @@ impl super::Window for Window { /// Returns true if the mouse if hovering this window fn is_mouse_hovered(&self) -> bool { - // TODO: - false + let state = self.input_state.read().unwrap(); + state.hovered_window_id == Some(self.window_id) } /// Set the window display title that appears on the title bar @@ -333,7 +669,7 @@ impl super::Window for Window { /// Set window position in screen space fn set_pos(&self, pos: super::Point) { - self.winit_window.set_outer_position(LogicalPosition { + self.winit_window.set_outer_position(PhysicalPosition { x: pos.x, y: pos.y }); @@ -341,12 +677,12 @@ impl super::Window for Window { /// Set window size in screen coordinates fn set_size(&self, size: super::Size) { - self.winit_window.request_inner_size(LogicalSize::new(size.x, size.y)); + self.winit_window.request_inner_size(PhysicalSize::new(size.x, size.y)); } /// Returns the screen position for the top-left corner of the window fn get_pos(&self) -> super::Point { - let pos = self.winit_window.outer_position().unwrap(); + let pos = *self.cached_position.read().unwrap(); super::Point { x: pos.x, y: pos.y @@ -355,28 +691,24 @@ impl super::Window for Window { /// Returns a gfx friendly full window rect to use as `gfx::Viewport` or `gfx::Scissor` fn get_viewport_rect(&self) -> super::Rect { - let size = self.winit_window.inner_size(); + let size = self.render_size(); super::Rect { x: 0, y: 0, - width: size.width as i32, - height: size.height as i32 + width: size.x, + height: size.y } } - /// Returns the screen position for the top-left corner of the window + /// Returns the render size of the window (physical pixels, or logical 1x when !dpi_aware) fn get_size(&self) -> super::Size { - let size = self.winit_window.inner_size(); - super::Size { - x: size.width as i32, - y: size.height as i32 - } + self.render_size() } /// Returns the screen rect of the window screen pos x, y , size x, y. fn get_window_rect(&self) -> super::Rect { - let pos = self.winit_window.outer_position().unwrap(); - let size = self.winit_window.inner_size(); + let pos = *self.cached_position.read().unwrap(); + let size = *self.cached_size.read().unwrap(); super::Rect { x: pos.x, y: pos.y, @@ -386,17 +718,16 @@ impl super::Window for Window { } /// Return mouse position in relative coordinates from the top left corner of the window - fn get_mouse_client_pos(&self, mouse_pos: super::Point) -> super::Point { - // panic!(); - super::Point { - x: 0, - y: 0 - } + fn get_mouse_client_pos(&self, _mouse_pos: super::Point) -> super::Point { + // Use the tracked client position from CursorMoved events + // which is already in window-local coordinates + self.input_state.read().unwrap().mouse_client_pos } - /// Return the dpi scale for the current monitor the window is on + /// Return the dpi scale for the current monitor the window is on. When the app is not + /// dpi-aware the engine operates entirely in logical (1x) pixels, so the effective scale is 1. fn get_dpi_scale(&self) -> f32 { - self.winit_window.scale_factor() as f32 + if self.dpi_aware { *self.cached_scale.read().unwrap() as f32 } else { 1.0 } } /// Gets the internal native handle @@ -408,14 +739,12 @@ impl super::Window for Window { /// Gets window events tracked from os update, to handle events inside external systems fn get_events(&self) -> super::WindowEventFlags { - panic!(); - super::WindowEventFlags { - bits: 0 - } + *self.events.read().unwrap() } + /// Clears events after they have been responded to fn clear_events(&mut self) { - panic!(); + *self.events.write().unwrap() = super::WindowEventFlags::NONE; } /// Const pointer diff --git a/src/pmfx.rs b/src/pmfx.rs index 6bcec06e..99ae95fb 100644 --- a/src/pmfx.rs +++ b/src/pmfx.rs @@ -1,4 +1,4 @@ -#![allow(clippy::collapsible_if)] +#![allow(clippy::collapsible_if)] use crate::gfx::Buffer; use crate::gfx::PipelineStatistics; @@ -36,7 +36,7 @@ use std::hash::{Hash, Hasher}; pub struct ResourceUse { pub index: u32, pub dimension: Vec3u -} +} /// Everything you need to render a world view; command buffers will be automatically reset and submitted for you. pub struct View { @@ -121,7 +121,7 @@ struct TrackedTexture { _tex_type: gfx::TextureType, } -/// Information to track changes to +/// Information to track changes to struct PmfxTrackingInfo { /// Filepath to the data which the pmfx File was deserialised from filepath: std::path::PathBuf, @@ -132,7 +132,7 @@ struct PmfxTrackingInfo { // pipelines (name) > permutation (mask : u32) which is tuple (build_hash, pipeline) type FormatPipelineMap = HashMap>; -// hash of the view in .0, the view itself in .1 the source view name which was used to generate the instance is stored in .2, +// hash of the view in .0, the view itself in .1 the source view name which was used to generate the instance is stored in .2, type TrackedView = (PmfxHash, Arc>>, String); type TrackedComputePass = (PmfxHash, Arc>>); @@ -147,8 +147,8 @@ pub struct Pmfx { pmfx: File, /// Tracking info for check on data reloads, grouped by pmfx name pmfx_tracking: HashMap, - /// Folder paths for - pmfx_folders: HashMap, + /// Folder paths for + pmfx_folders: HashMap, /// Updated by calling 'update_window' this will cause any tracked textures to check for resizes and rebuild textures if necessary window_sizes: HashMap, /// Nested structure of: format (u64) > FormatPipelineMap @@ -203,7 +203,7 @@ pub struct TotalStats { pub gpu_start: f64, /// Time of the final submission in seconds pub gpu_end: f64, - /// Total pipeline statistics + /// Total pipeline statistics pub pipeline_stats: PipelineStatistics } @@ -237,7 +237,7 @@ impl PassStats where D: gfx::Device { pub fn new_query_buffer(device: &mut D, elem_size: usize, num_elems: usize) -> D::Buffer { device.create_read_back_buffer(elem_size * num_elems).unwrap() } - + pub fn new(device: &mut D, num_buffers: usize) -> Self { let mut timestamp_buffers = Vec::new(); let mut fences = Vec::new(); @@ -291,7 +291,7 @@ struct File { dependencies: Vec } -/// pmfx File serialisation, +/// pmfx File serialisation, impl File { /// creates a new empty pmfx fn new() -> Self { @@ -358,7 +358,7 @@ struct RaytracingShaderBindingTableInfo { callable_shaders: Vec } -/// Enum for possible pipeline types +/// Enum for possible pipeline types #[derive(Clone)] pub enum PipelineType { None, @@ -462,6 +462,8 @@ impl DynamicBuffer where D: gfx::Device, T: Sized { } } + pub fn get_bb(&self) -> usize { self.bb } + /// Swap buffers once a frame for safe CPU writes an GPU in flight reads pub fn swap(&mut self) { self.bb = (self.bb + 1) % self.num_buffers @@ -561,7 +563,7 @@ impl DynamicBuffer where D: gfx::Device, T: Sized { pub fn get_lookup(&self) -> GpuBufferLookup { GpuBufferLookup { index: self.get_index() as u32, - count: self.len as u32 + count: self.len as u32, } } } @@ -588,14 +590,14 @@ pub struct DynamicWorldBuffers { impl Default for DynamicWorldBuffers where D: gfx::Device { fn default() -> Self { Self { - draw: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), - extent: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), - material: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), - point_light: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), - spot_light: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), + draw: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), + extent: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), + material: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), + point_light: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), + spot_light: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), directional_light: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), - camera: DynamicBuffer::::new(gfx::BufferUsage::CONSTANT_BUFFER, 3), - shadow_matrix: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), + camera: DynamicBuffer::::new(gfx::BufferUsage::CONSTANT_BUFFER, 3), + shadow_matrix: DynamicBuffer::::new(gfx::BufferUsage::SHADER_RESOURCE, 3), } } } @@ -627,7 +629,7 @@ pub struct CameraConstants { #[derive(Clone)] pub struct DrawData { /// World matrix for transforming entity - pub world_matrix: Mat34f, + pub world_matrix: Mat34f, } /// GPU friendly struct containing single entity draw data @@ -637,10 +639,10 @@ pub struct ExtentData { /// Centre pos of aabb pub pos: Vec3f, /// Half extent of aabb - pub extent: Vec3f + pub extent: Vec3f } -/// GPU friendly structure containing lookup id's for bindless materials +/// GPU friendly structure containing lookup id's for bindless materials #[repr(C)] #[derive(Clone)] pub struct MaterialData { @@ -735,7 +737,7 @@ pub fn cubemap_camera_face(face: usize, pos: Vec3f, near: f32, far: f32) -> Came vec3f(1.0, 0.0, 0.0), //+x vec3f(-1.0, 0.0, 0.0), //-x vec3f(0.0, 1.0, 0.0), //+y - vec3f(0.0, -1.0, 0.0), //-y + vec3f(0.0, -1.0, 0.0), //-y vec3f(0.0, 0.0, 1.0), //+z vec3f(0.0, 0.0, -1.0) //-z ]; @@ -779,7 +781,7 @@ pub fn cubemap_camera_face(face: usize, pos: Vec3f, near: f32, far: f32) -> Came fn create_shader_from_file(device: &D, folder: &Path, file: Option) -> Result, super::Error> { if let Some(shader) = file { let shader_filepath = folder.join(shader); - let shader_data = fs::read(shader_filepath)?; + let shader_data = fs::read(shader_filepath)?; let shader_info = gfx::ShaderInfo { shader_type: gfx::ShaderType::Vertex, compile_info: None @@ -851,9 +853,9 @@ fn to_gfx_texture_info(pmfx_texture: &TextureInfo, ratio_size: (u64, u64)) -> gf let (width, height) = ratio_size; // infer texture type from dimensions - let tex_type = if pmfx_texture.cubemap { + let tex_type = if pmfx_texture.cubemap { gfx::TextureType::TextureCube - } + } else if pmfx_texture.depth > 1 { gfx::TextureType::Texture3D } @@ -1047,24 +1049,23 @@ impl Pmfx where D: gfx::Device { /// Retunrs a `WorldBufferInfo` that contains the serv index and count of the various world buffers used /// during rendering pub fn get_world_buffer_info(&self) -> WorldBufferInfo { - // construct on the fly WorldBufferInfo { - draw: self.world_buffers.draw.get_lookup(), - extent: self.world_buffers.extent.get_lookup(), - material: self.world_buffers.material.get_lookup(), - point_light: self.world_buffers.point_light.get_lookup(), - spot_light: self.world_buffers.spot_light.get_lookup(), + draw: self.world_buffers.draw.get_lookup(), + extent: self.world_buffers.extent.get_lookup(), + material: self.world_buffers.material.get_lookup(), + point_light: self.world_buffers.point_light.get_lookup(), + spot_light: self.world_buffers.spot_light.get_lookup(), directional_light: self.world_buffers.directional_light.get_lookup(), - camera: self.world_buffers.camera.get_lookup(), - shadow_matrix: self.world_buffers.shadow_matrix.get_lookup(), - user_data: self.push_constant_user_data + camera: self.world_buffers.camera.get_lookup(), + shadow_matrix: self.world_buffers.shadow_matrix.get_lookup(), + user_data: self.push_constant_user_data, } } /// Load a pmfx from a folder, where the folder contains a pmfx info.json and shader binaries in separate files within the directory /// You can load multiple pmfx files which will be merged together, shaders are grouped by pmfx_name/ps_main.psc - /// Render graphs and pipleines must have unique names, if multiple pmfx name a pipeline the same name - pub fn load(&mut self, filepath: &str) -> Result<(), super::Error> { + /// Render graphs and pipleines must have unique names, if multiple pmfx name a pipeline the same name + pub fn load(&mut self, filepath: &str) -> Result<(), super::Error> { // get the name for indexing by pmfx name/folder let folder = Path::new(filepath); let pmfx_name = if let Some(name) = folder.file_name() { @@ -1081,19 +1082,19 @@ impl Pmfx where D: gfx::Device { let info_filepath = folder.join(format!("{}.json", pmfx_name)); let pmfx_data = fs::read(&info_filepath)?; let file : File = serde_json::from_slice(&pmfx_data)?; - + // create tracking info to check if the pmfx has been rebuilt let file_metadata = fs::metadata(&info_filepath)?; e.insert(PmfxTrackingInfo { modified_time: file_metadata.modified()?, filepath: info_filepath }); - + // add files from pmfx for tracking for dep in &file.dependencies { self.reloader.add_file(dep); } - + // merge into pmfx self.merge_pmfx(file, filepath); } @@ -1314,7 +1315,7 @@ impl Pmfx where D: gfx::Device { /// Retruns a vector of resource use indices specified in `pmfx` pass and based on `ResourceUsage` /// creates resources that do not yet exist fn get_resource_use_indices(&mut self, device: &mut D, info: &GraphPassInfo) -> Result, super::Error> { - // create textures we may use + // create textures we may use let mut use_indices = Vec::new(); if let Some(uses) = &info.uses { for (resource, usage) in uses { @@ -1371,15 +1372,15 @@ impl Pmfx where D: gfx::Device { } fn create_view_pass_inner( - &mut self, device: - &mut D, view_name: &str, - graph_pass_name: &str, + &mut self, device: + &mut D, view_name: &str, + graph_pass_name: &str, info: &GraphPassInfo, pmfx_view: &ViewInfo, array_slice: usize, cubemap: bool ) -> Result<(), super::Error> { - + // make a custom name for multi pass let graph_pass_multi_name = if array_slice > 0 { format!("{}_{}", graph_pass_name, array_slice) @@ -1515,7 +1516,7 @@ impl Pmfx where D: gfx::Device { self.pass_stats.insert(graph_pass_multi_name.to_string(), PassStats::new(device, 2)); Ok(()) - } + } /// Create a view pass from information specified in pmfx file fn create_view_pass(&mut self, device: &mut D, view_name: &str, graph_pass_name: &str, info: &GraphPassInfo) -> Result<(), super::Error> { @@ -1526,7 +1527,7 @@ impl Pmfx where D: gfx::Device { // create pass from targets let pmfx_view = self.pmfx.views[view_name].clone(); - // create textures for view + // create textures for view let mut cubemap = false; for name in &pmfx_view.render_target { self.create_texture(device, name)?; @@ -1635,7 +1636,7 @@ impl Pmfx where D: gfx::Device { let mut hash = DefaultHasher::new(); graph_pass_name.hash(&mut hash); let colour_hash : u32 = hash.finish() as u32 | 0xff000000; - + let pass = ComputePass { phantom_data: std::marker::PhantomData, pass_pipline: pass_pipeline, @@ -1745,9 +1746,9 @@ impl Pmfx where D: gfx::Device { fn create_resolve_transition( &mut self, device: &mut D, - texture_barriers: &mut HashMap, - view_name: &str, - texture_name: &str, + texture_barriers: &mut HashMap, + view_name: &str, + texture_name: &str, target_state: ResourceState) -> Result<(), super::Error> { if texture_barriers.contains_key(texture_name) { let state = texture_barriers[texture_name]; @@ -1780,7 +1781,7 @@ impl Pmfx where D: gfx::Device { }, Subresource::ResolveResource ); - + // perform the resolve cmd_buf.resolve_texture_subresource(tex, 0)?; @@ -1814,21 +1815,21 @@ impl Pmfx where D: gfx::Device { fn create_texture_transition_barrier( &mut self, device: &mut D, - texture_barriers: &mut HashMap, - view_name: &str, - texture_name: &str, + texture_barriers: &mut HashMap, + view_name: &str, + texture_name: &str, target_state: ResourceState) -> Result<(), super::Error> { if texture_barriers.contains_key(texture_name) { let state = texture_barriers[texture_name]; if state != target_state { // add barrier placeholder in the command_queue let barrier_name = format!("barrier_{}-{} ({:?})", view_name, texture_name, target_state); - self.command_queue.push(barrier_name.to_string()); + self.command_queue.push(barrier_name.to_string()); // create a command buffer let mut cmd_buf = device.create_cmd_buf(1); cmd_buf.begin_event( - 0xfff1b023, + 0xfff1b023, &format!("transition_barrier: {} ({} -> {})", &texture_name, state, target_state) ); cmd_buf.transition_barrier(&gfx::TransitionBarrier { @@ -1840,7 +1841,7 @@ impl Pmfx where D: gfx::Device { cmd_buf.end_event(); cmd_buf.close()?; self.barriers.insert(barrier_name, cmd_buf); - + // update track state texture_barriers.remove(texture_name); texture_barriers.insert(texture_name.to_string(), target_state); @@ -1857,8 +1858,8 @@ impl Pmfx where D: gfx::Device { } /// Create a render graph wih automatic resource barrier generation from info specified insie .pmfx file - pub fn create_render_graph(&mut self, device: &mut D, graph_name: &str) -> Result<(), super::Error> { - // go through the graph sequentially, as the command lists are executed in order but generated + pub fn create_render_graph(&mut self, device: &mut D, graph_name: &str) -> Result<(), super::Error> { + // go through the graph sequentially, as the command lists are executed in order but generated if self.pmfx.render_graphs.contains_key(graph_name) { // create views for any nodes in the graph @@ -1869,16 +1870,16 @@ impl Pmfx where D: gfx::Device { self.command_queue.clear(); let mut barriers = self.pmfx.textures.iter().filter(|tex|{ - tex.1.usage.contains(&ResourceState::ShaderResource) || + tex.1.usage.contains(&ResourceState::ShaderResource) || tex.1.usage.contains(&ResourceState::RenderTarget) || tex.1.usage.contains(&ResourceState::DepthStencil) }).map(|tex|{ - (tex.0.to_string(), ResourceState::ShaderResource) + (tex.0.to_string(), ResourceState::ShaderResource) }).collect::>(); // loop over the graph multiple times adding views in depends on order, until we add all the views let mut to_add = self.pmfx.render_graphs[graph_name].len(); - + let mut added = 0; let mut dependencies = HashSet::new(); while added < to_add { @@ -1893,12 +1894,12 @@ impl Pmfx where D: gfx::Device { continue; } } - + // already added this pass if dependencies.contains(graph_pass_name) { continue; } - + // wait for dependencies if let Some(depends_on) = &instance.depends_on { let mut passes = false; @@ -1906,7 +1907,7 @@ impl Pmfx where D: gfx::Device { for d in depends_on { if !pmfx_graph.contains_key(d) { passes = true; - println!("hotline_rs::pmfx:: [warning] graph pass {} missing dependency {}. ignoring", + println!("hotline_rs::pmfx:: [warning] graph pass {} missing dependency {}. ignoring", graph_pass_name, d); } else if dependencies.contains(d) { @@ -1962,21 +1963,21 @@ impl Pmfx where D: gfx::Device { // resolve and generate mips if resolve { self.create_resolve_transition( - device, - &mut barriers, - &graph_pass_name, + device, + &mut barriers, + &graph_pass_name, &u.0, ResourceState::ShaderResource, )?; } - + // generate mips on non msaa resources if gen_mips { // generate_mip_maps mips expects us to be in ShaderResource state self.create_texture_transition_barrier( - device, - &mut barriers, - &graph_pass_name, + device, + &mut barriers, + &graph_pass_name, &u.0, ResourceState::ShaderResource)?; @@ -1988,14 +1989,14 @@ impl Pmfx where D: gfx::Device { // transition to target state self.create_texture_transition_barrier( - device, - &mut barriers, - &graph_pass_name, + device, + &mut barriers, + &graph_pass_name, &u.0, res_state)?; } } - + if let Some(view) = &instance.view { // create transitions by inspecting view info let pmfx_view = self.pmfx.views[view].clone(); @@ -2004,14 +2005,14 @@ impl Pmfx where D: gfx::Device { for rt_name in pmfx_view.render_target { self.create_texture_transition_barrier( device, &mut barriers, view, &rt_name, ResourceState::RenderTarget)?; - + } - + // same for depth stencils for ds_name in pmfx_view.depth_stencil { self.create_texture_transition_barrier( device, &mut barriers, view, &ds_name, ResourceState::DepthStencil)?; - + } // create pipelines requested for this view instance with the pass format @@ -2052,7 +2053,7 @@ impl Pmfx where D: gfx::Device { dependencies.insert(graph_pass_name.to_string()); } } - + // finally all targets which are in the 'barriers' array are transitioned to shader resources (for debug views) let srvs = barriers.keys().map(|k|{ k.to_string() @@ -2061,7 +2062,7 @@ impl Pmfx where D: gfx::Device { for name in srvs { let result = self.create_resolve_transition( device, &mut barriers, "eof", &name, ResourceState::ShaderResource); - + if result.is_err() { // TODO: tell user without spewing out errors } @@ -2109,7 +2110,7 @@ impl Pmfx where D: gfx::Device { } /// Create a ComputePipeline instance for the combination of pmfx_pipeline settings - pub fn create_compute_pipeline(&mut self, device: &D, pipeline_name: &str) -> Result<(), super::Error> { + pub fn create_compute_pipeline(&mut self, device: &D, pipeline_name: &str) -> Result<(), super::Error> { if self.pmfx.pipelines.contains_key(pipeline_name) { // first create shaders if necessary let folder = self.pmfx_folders.get(pipeline_name) @@ -2119,7 +2120,7 @@ impl Pmfx where D: gfx::Device { self.create_shader(device, Path::new(&folder), &pipeline.cs)?; } - for (_, pipeline) in self.pmfx.pipelines[pipeline_name].clone() { + for (_, pipeline) in self.pmfx.pipelines[pipeline_name].clone() { let cs = self.get_shader(&pipeline.cs); if let Some(cs) = cs { let pso = device.create_compute_pipeline(&gfx::ComputePipelineInfo { @@ -2131,7 +2132,7 @@ impl Pmfx where D: gfx::Device { // TODO: permutations //let mask = permutation.parse().unwrap(); //permutations.insert(mask, (pipeline.hash, pso)); - + self.compute_pipelines.insert(pipeline_name.to_string(), (pipeline.hash, pso)); } } @@ -2146,27 +2147,27 @@ impl Pmfx where D: gfx::Device { } /// Create a RenderPipeline instance for the combination of pmfx_pipeline settings and an associated RenderPass - pub fn create_render_pipeline(&mut self, device: &D, pipeline_name: &str, pass: &D::RenderPass) -> Result<(), super::Error> { + pub fn create_render_pipeline(&mut self, device: &D, pipeline_name: &str, pass: &D::RenderPass) -> Result<(), super::Error> { if self.pmfx.pipelines.contains_key(pipeline_name) { // first create shaders if necessary let folder = self.pmfx_folders.get(pipeline_name) .unwrap_or_else(|| panic!("hotline_rs::pmfx:: expected to find pipeline {} in pmfx_folders", pipeline_name)).to_string(); - + for (_, pipeline) in self.pmfx.pipelines[pipeline_name].clone() { self.create_shader(device, Path::new(&folder), &pipeline.vs)?; self.create_shader(device, Path::new(&folder), &pipeline.ps)?; } - + // create entry for this format if it does not exist let fmt = pass.get_format_hash(); let format_pipeline = self.render_pipelines.entry(fmt).or_insert(HashMap::new()); - + // create entry for this pipeline permutation set if it does not exist if !format_pipeline.contains_key(pipeline_name) { println!("hotline_rs::pmfx:: creating pipeline: {}", pipeline_name); format_pipeline.insert(pipeline_name.to_string(), HashMap::new()); // we create a pipeline per-permutation - for (permutation, pipeline) in self.pmfx.pipelines[pipeline_name].clone() { + for (permutation, pipeline) in self.pmfx.pipelines[pipeline_name].clone() { let vertex_layout = pipeline.vertex_layout.as_ref().unwrap(); let pso = device.create_render_pipeline(&gfx::RenderPipelineInfo { vs: self.get_shader(&pipeline.vs), @@ -2182,10 +2183,10 @@ impl Pmfx where D: gfx::Device { pass: Some(pass), ..Default::default() })?; - + println!("hotline_rs::pmfx:: compiled render pipeline: {}", pipeline_name); let format_pipeline = self.render_pipelines.get_mut(&fmt).unwrap(); - let permutations = format_pipeline.get_mut(pipeline_name).unwrap(); + let permutations = format_pipeline.get_mut(pipeline_name).unwrap(); let mask = permutation.parse().unwrap(); permutations.insert(mask, (pipeline.hash, pso)); @@ -2215,9 +2216,9 @@ impl Pmfx where D: gfx::Device { self.create_shader(device, Path::new(&folder), &Some(shader))?; } } - + // for each permutation create a pipeline - for (_, pipeline) in self.pmfx.pipelines[pipeline_name].clone() { + for (_, pipeline) in self.pmfx.pipelines[pipeline_name].clone() { let shaders = pipeline.lib.expect("hotline_rs::pmfx:: ray tracing pipeline expects a lib member with a set of raytacing shaders") .iter() .map(|x| gfx::RaytracingShader { @@ -2262,7 +2263,7 @@ impl Pmfx where D: gfx::Device { self.get_render_pipeline_permutation_for_format(pipeline_name, 0, format_hash) } - /// Returns a pmfx pipline for a random / unknown render target format... prefer to use `get_render_pipeline_for_format` + /// Returns a pmfx pipline for a random / unknown render target format... prefer to use `get_render_pipeline_for_format` /// if you know the format the target you are rendering in to. pub fn get_render_pipeline<'stack>(&'stack self, pipeline_name: &str) -> Result<&'stack D::RenderPipeline, super::Error> { for format in self.render_pipelines.values() { @@ -2305,7 +2306,7 @@ impl Pmfx where D: gfx::Device { }) } } - + /// Fetch a prebuilt RaytracingPipelineBinding which is contains a RaytracingPipeline and RaytracingShaderBindingTable pub fn get_raytracing_pipeline<'stack>(&'stack self, pipeline_name: &str) -> Result<&'stack RaytracingPipelineBinding, super::Error> { if self.raytracing_pipelines.contains_key(pipeline_name) { @@ -2336,7 +2337,7 @@ impl Pmfx where D: gfx::Device { if !timestamps.is_empty() { stats.start_timestamp = timestamps[0]; min_frame_timestamp = min(stats.start_timestamp, min_frame_timestamp); - + } // end timestamp let timestamps = device.read_timestamps( @@ -2387,7 +2388,7 @@ impl Pmfx where D: gfx::Device { } /// Reload all active resources based on hashes - pub fn reload(&mut self, device: &mut D) -> Result<(), super::Error> { + pub fn reload(&mut self, device: &mut D) -> Result<(), super::Error> { let reload_paths = self.pmfx_tracking.iter_mut().filter(|(_, tracking)| { fs::metadata(&tracking.filepath).unwrap().modified().unwrap() > tracking.modified_time }).map(|tracking| { @@ -2399,7 +2400,7 @@ impl Pmfx where D: gfx::Device { if !reload_filepath.is_empty() { println!("hotline_rs::pmfx:: reload from {}", reload_filepath); let pmfx_data = fs::read(&reload_filepath).expect("hotline_rs::pmfx:: failed to read file"); - + let file : File = serde_json::from_slice(&pmfx_data)?; self.merge_pmfx(file, PathBuf::from(&reload_filepath).parent().unwrap().to_str().unwrap()); @@ -2474,11 +2475,11 @@ impl Pmfx where D: gfx::Device { println!("hotline::pmfx:: reloading shader: {}", shader); self.shaders.remove(shader); } - + // reload pipelines tuple = (format_hash, pipeline_name, permutation_mask) for pipeline in &reload_pipelines { println!("hotline::pmfx:: reloading pipeline: {}", pipeline.1); - + // TODO: here we could only remove affected permutations let format_pipelines = self.render_pipelines.get_mut(&pipeline.0).unwrap(); format_pipelines.remove(&pipeline.1); @@ -2507,7 +2508,7 @@ impl Pmfx where D: gfx::Device { }); } - // + // if rebuild_graph { self.create_render_graph(device, &self.active_render_graph.to_string())?; } @@ -2637,7 +2638,7 @@ impl Pmfx where D: gfx::Device { // view pipeline stats pass_stats.pipeline_stats_heap.reset(); pass_stats.pipeline_query_index = cmd_buf.begin_query( - &mut pass_stats.pipeline_stats_heap, + &mut pass_stats.pipeline_stats_heap, gfx::QueryType::PipelineStatistics ); } @@ -2651,7 +2652,7 @@ impl Pmfx where D: gfx::Device { if pass_stats.pipeline_query_index != usize::max_value() { let buf = &mut pass_stats.pipeline_stats_buffers[pass_stats.write_index]; cmd_buf.end_query( - &mut pass_stats.pipeline_stats_heap, + &mut pass_stats.pipeline_stats_heap, gfx::QueryType::PipelineStatistics, pass_stats.pipeline_query_index, buf, @@ -2668,9 +2669,10 @@ impl Pmfx where D: gfx::Device { cmd_buf.reset(swap_chain); // inserts markers for timing and tracking pipeline stats - let mut stats = self.pass_stats.remove(name).unwrap(); - Self::stats_start(cmd_buf, &mut stats); - self.pass_stats.insert(name.to_string(), stats); + if let Some(mut stats) = self.pass_stats.remove(name) { + Self::stats_start(cmd_buf, &mut stats); + self.pass_stats.insert(name.to_string(), stats); + } } } } @@ -2728,9 +2730,10 @@ impl Pmfx where D: gfx::Device { } else if let Some(cmd_buf) = cmd_bufs.get_mut(node) { // inserts markers for timing and tracking pipeline stats - let mut stats = self.pass_stats.remove(node).unwrap(); - Self::stats_end(cmd_buf, &mut stats); - self.pass_stats.insert(node.to_string(), stats); + if let Some(mut stats) = self.pass_stats.remove(node) { + Self::stats_end(cmd_buf, &mut stats); + self.pass_stats.insert(node.to_string(), stats); + } cmd_buf.close().unwrap(); device.execute(cmd_buf); @@ -2758,7 +2761,7 @@ impl imgui::UserInterface for Pmfx where D: gfx::Device, A: os::A let mut imgui_open = open; if imgui.begin("textures", &mut imgui_open, imgui::WindowFlags::ALWAYS_HORIZONTAL_SCROLLBAR) { for texture in self.textures.values() { - + let thumb_size = 256.0; let aspect = texture.1.size.0 as f32 / texture.1.size.1 as f32; let w = thumb_size * aspect; @@ -2818,7 +2821,7 @@ impl imgui::UserInterface for Pmfx where D: gfx::Device, A: os::A imgui.end(); imgui_open - } + } else { false } @@ -2842,7 +2845,7 @@ impl PmfxReloadResponder { impl ReloadResponder for PmfxReloadResponder { fn add_file(&mut self, filepath: &str) { self.files.push(filepath.to_string()); - } + } fn get_files(&self) -> Vec { self.files.to_vec() @@ -2853,21 +2856,23 @@ impl ReloadResponder for PmfxReloadResponder { } fn build(&mut self) -> std::process::ExitStatus { + // Shader/data compilation is hooked into the crate's build.rs (under the `build_data` + // feature), so reuse that single path here. `cargo build --lib` re-runs the build script - + // which recompiles the data - while skipping the example/bin targets; when only files under + // `shaders/` changed cargo leaves the crate artifacts untouched, so this rebuilds data only. let hotline_path = super::get_data_path("../.."); - let pmbuild = super::get_data_path("../../hotline-data/pmbuild.cmd"); - let output = std::process::Command::new(pmbuild) + let output = std::process::Command::new("cargo") .current_dir(hotline_path) - .arg("win32-data") - .arg("-pmfx") + .args(["build", "--lib"]) .output() .expect("hotline::hot_lib:: hot pmfx failed to compile!"); if !output.stdout.is_empty() { - println!("{}", String::from_utf8(output.stdout).unwrap()); + println!("{}", String::from_utf8_lossy(&output.stdout)); } if !output.stderr.is_empty() { - println!("{}", String::from_utf8(output.stderr).unwrap()); + println!("{}", String::from_utf8_lossy(&output.stderr)); } if output.status.success() { diff --git a/tests/tests.rs b/tests/tests.rs index 305f1ea9..d3659870 100644 --- a/tests/tests.rs +++ b/tests/tests.rs @@ -415,8 +415,8 @@ fn pmfx() -> Result<(), hotline_rs::Error> { let texture_pipeline = ctx.pmfx.get_render_pipeline_for_format("texture2d_array", fmt)?; // texture array pipeline has 2 sets of push constants, so the resources bind onto 2 - // t0, space9 is Texture2DArray - let slots = texture_pipeline.get_pipeline_slot(0, 10, gfx::DescriptorType::ShaderResource); + // t1, space10 is Texture2DArray + let slots = texture_pipeline.get_pipeline_slot(1, 10, gfx::DescriptorType::ShaderResource); assert!(slots.is_some()); if let Some(slots) = slots { @@ -677,8 +677,8 @@ fn draw_vertex_buffer_instanced() -> Result<(), hotline_rs::Error> { } #[test] -fn draw_cbuffer_instanced() -> Result<(), hotline_rs::Error> { - boot_client_ecs_plugin_demo("draw_cbuffer_instanced") +fn draw_structured_buffer_instanced() -> Result<(), hotline_rs::Error> { + boot_client_ecs_plugin_demo("draw_structured_buffer_instanced") } #[test] diff --git a/todo.txt b/todo.txt index 1c39cece..dc01d74f 100644 --- a/todo.txt +++ b/todo.txt @@ -1,4 +1,33 @@ -// TODO: +macos +x perf issues push constants +x hot reload +x window pos restore +x shader errors don't properly display +x majority shader compilation +x dynamic cubemap +x tangent space normal map is black +x shadow map is black +x omni shadow is black +x compute pipeline +x RW texture + better demo +x MSAA +x MRT +x mip downsample +x gpu timestamp +x pmfx hotreload +x material ibl +x bindless material +x sample cmp shadow +x cbuffer instanced causes huge perf issues (should be structured tbh) + +- cleanup mem + +- pmbuild needs universal install into hotline-data +- single shader compile, htwv + +- draw indirect +- video player +- resource tests // issues // - swap between dynamic cube and PBR causes inconsitency in the cubemap texture @@ -14,36 +43,30 @@ // - area light // - disney brdf -// platforms -// - bring across macos shader compilation -// - integrate macos changes and fixup win32/d3d12 -// - try update windows-rs - // engine // - reverse depth // - visibility buffer // - mesh shader -// - hello triangle (ray tracing) // - HDR pipeline // - glft // - set name on resources // - lazy init print function -// - example triangle culling via execute indirect // gfx +// - cross platform append/consume buffer (buffer + atomic counter). metal has no +// AppendStructuredBuffer; currently mapped to a plain buffer so counter semantics are lost // - Alpha to coverage // - Stencil Ref / Buffer // - API for fence // ui / debug -// - per demo settings // - thread stats // - view menu + saving state // - imgui not tracked within draw call stats -// build - // DONE: +// x per demo settings +// x hello triangle (ray tracing) // x create tlas with heap? // x raytraced triangle // x test shader compilation with mac changes @@ -158,7 +181,7 @@ // x fix issues with crash when a return happens between begin/end render pass // x gfx::buffer should need to be mut for update // x propagate sample mask?? -// x Draw Instanced +// x Draw Instanced // x draw instanced // x pmfx v2 docs (view) // x batch vertex instance buffer