Skip to content

release: v0.34.6 - #368

Open
kolkov wants to merge 9 commits into
mainfrom
release/v0.34.6
Open

kolkov wants to merge 9 commits into
mainfrom
release/v0.34.6

Conversation

@kolkov

@kolkov kolkov commented Oct 18, 2026 •

Copy link
Copy Markdown
Contributor

Patch release: pipeline-overridable constants now work on every backend, a GLES crash fix, and the naga and goffi bumps. The v0.34.6 CHANGELOG section also covers what merged since v0.34.5: #367 (usage scopes per pass), #366, #363, #362 and #359.

Fixed

  • Pipeline-overridable constants were ignored; @workgroup_size(<override>) compiled as 1. No backend processed WGSL override declarations, so ComputePipelineDescriptor.Constants had no effect and an override-sized workgroup ran with size 1: a @workgroup_size(n, 2) shader with n = 4 ran 2 invocations instead of 8 on Vulkan, DX12 and GLES. Each backend now specializes the module per pipeline stage with naga.ProcessOverrides before the shader backend, as Rust wgpu-hal does (vulkan/device.rs, dx12/device.rs, gles/device.rs, metal/device.rs); modules without overrides keep the compile-once path. On DX12 the constants are part of the DXIL cache key.

  • GLES fence polling without the GL context. Queue.PollCompleted, Device.Wait and Device.GetFenceStatus called glGetSynciv / glClientWaitSync without taking the adapter context, so the call ran on whatever OS thread the caller was on, with no GL context current. With Intel's Windows GL driver Device.Poll (reached from Buffer.Map) crashed with an access violation. They now take the context first, as Rust wgpu-hal does (gles/device.rs get_fence_value, wait).

  • GLES texture readback came back mirrored, with R and B swapped. CopyTextureToBuffer read with a hard-coded GL_BGRA, flipped the rows (shaders already use naga's ADJUST_COORDINATE_SPACE, so texture row 0 is the image top) and ignored bytesPerRow: examples/triangle-headless on GLES returned a blue, upside-down triangle, and 256-byte-pitch readbacks were sheared. A copy that did not fill the whole destination buffer read back zeros, because MapBuffer uses the CPU shadow only when it covers the buffer. It now reads in the texture's own transfer format and type, keeps memory row order, writes rows at bytesPerRow and keeps the rest of the buffer; CopyBufferToTexture sets GL_UNPACK_ROW_LENGTH from bytesPerRow. Both as in Rust wgpu-hal gles/queue.rs. The GLES output of triangle-headless is now byte-identical to Vulkan's.

  • GLES (Linux) CreateTexture failed on an earlier call's GL error. On Mesa's d3d12 driver in WSL the first texture after device creation failed with GL_INVALID_ENUM left by a previous call; stale errors are now drained and logged before the allocation, so triangle-headless runs there.

  • Software CopyBufferToTexture ignored bytesPerRow, rowsPerImage and the origin, copying one tight block to the start of the texture. Rows are now copied into the region at the texture origin.

Changed

  • Unknown pipeline-constant keys are an error. A Constants key that names no override now fails pipeline creation with naga's *PipelineConstantError instead of being ignored. Per the WebGPU spec the key of @id(N) override x is "N".
  • Dependencies: naga v0.19.0 → v0.20.3 (naga.ProcessOverrides), goffi v0.6.3 → v0.6.7. goffi v0.6.6+ refuses calls whose argument count differs from the call interface; a static check of all 770 ffi.CallFunction sites found no mismatch. The Android preview check pins goffi v0.6.7, since it asserts the exact version.
  • Dependency versions in README, AGENTS.md, ROADMAP, ARCHITECTURE and the Android preview docs.

Testing

  • New override tests through the public API and on every adapter: workgroup size from defaults and from constants, overrides in arithmetic, by @id and by name, an array sized by an anonymous override, unknown keys, render stages. 116 pass on Windows (Vulkan on Intel and NVIDIA, DX12 FXC and DXIL on both, GLES, Software). The workgroup cases fail on main on Vulkan, DX12 FXC, DX12 DXIL and GLES.

  • TestFencePollingTakesAdapterContext holds the adapter context from another goroutine and checks that each of the three calls waits for it; it fails on main for all three.

  • Builds for windows/amd64, linux/amd64, linux/arm64, darwin/arm64, darwin/amd64, js/wasm and -tags rust; golangci-lint clean on Windows, Linux and macOS; tests on Windows and Linux.

  • Adapters without compute shaders (Mesa's d3d12 GL driver in WSL) skip the compute cases, as Rust's test harness does with downlevel_flags.

  • TestTextureCopyRoundTrip uploads four colors with WriteTexture and with a 256-byte-pitch CopyBufferToTexture and reads them back with a 256-byte pitch on every backend; on main it fails for GLES (both uploads) and for the software CopyBufferToTexture.

  • examples/triangle-headless on Vulkan (validation layer clean), DX12, GLES and Software on Windows, and GLES on Mesa d3d12 in WSL: identical output.

Not verified

  • macOS/Metal: the override change and goffi v0.6.7 are not yet run on a Mac.

…rocessOverrides

Every HAL backend now lowers WGSL to naga IR once per shader module and,
for a module with overrides or a stage with pipeline constants, specializes
it per pipeline stage with naga.ProcessOverrides before the shader backend,
as Rust wgpu-hal does (vulkan/device.rs:804, dx12/device.rs:349,
gles/device.rs:385, metal/device.rs:174). Modules without overrides keep
the compile-once path.

- internal/shaderir: Lower, LowerValidated, PerStage, Specialize,
  BackendError (unresolved overrides from a backend = HAL bug), SPIRV.
- Vulkan: temporary VkShaderModule per stage for modules with overrides.
- DX12: per-stage HLSL/DXIL; pipeline constants are part of the DXIL
  shader-cache key.
- GLES: WGSL lowered at module creation; GLSL per stage.
- Metal: MSL per stage, workgroup size from the specialized entry point,
  resource slots numbered from the whole module (mslmap.StageOptions);
  mslmap uses ArraySize.IsRuntimeSized.
- Software: SPIR-V specialized and parsed per pipeline.
- Unknown pipeline-constant keys fail pipeline creation with naga's
  *PipelineConstantError (also for modules without overrides).

Fixes BUG-OVERRIDE-WORKGROUP-001: @workgroup_size(<override>) compiled as 1.
Resolves TODO(compute-constants) in all backends.

Includes the naga v0.19.0 -> v0.20.3 bump that provides naga.ProcessOverrides
and ir.ArraySize.IsRuntimeSized.
- hal/allbackends: public API tests (workgroup size from an override,
  pipeline constants by @id and by name, an array sized by an anonymous
  override, unknown keys, overrides in render stages) on the adapter that
  RequestAdapter picks per backend, and a HAL test that runs the same
  compute cases on every adapter of every backend. A wrong result is
  checked against a control shader with const instead of override: when
  the control fails too, the failure is not an override failure and the
  case is skipped with both readbacks.
- gles: GLSL written for an override workgroup size and constants.
- mslmap: specialized modules keep the slots of the whole module;
  override-sized arrays get no _buffer_sizes field.
- shaderir: unit tests.
v0.6.6 refuses calls whose argument count differs from the call interface;
a static check of all 770 ffi.CallFunction sites found no mismatch (768
matched, 2 dynamic objc_msgSend calls sized by construction). v0.6.6 also
rewrote the arm64 argument layout, and v0.6.7 fixes struct results and
arguments for hand-built call interfaces; wgpu builds its struct-returning
interfaces with PrepareCallInterface, so the new Size/Alignment rule does
not apply.

Verified on Windows (Vulkan with validation layer, DX12, GLES compute) and
Linux (Vulkan llvmpipe). macOS/Metal not yet run on hardware.
…aders

Mesa's d3d12 GL driver in WSL exposes no compute shaders; wgpu reports
it (DownlevelFlagsComputeShaders unset, CreateComputePipeline refuses), and
the override tests failed there instead of skipping. Both the public-API
and the per-adapter HAL tests now skip such adapters, as Rust wgpu's test
harness does with TestParameters::downlevel_flags.
Queue.PollCompleted, Device.Wait and Device.GetFenceStatus called
glGetSynciv / glClientWaitSync without taking the adapter context, so the
call ran on whatever OS thread the caller's goroutine was on, with no GL
context current. On Intel's Windows ICD this crashed with an access
violation in Device.Poll (seen from Buffer.Map); it also read the pending
fence list without ordering against Submit. Rust wgpu-hal locks the
context for both calls (gles/device.rs get_fence_value, wait).

TestFencePollingTakesAdapterContext holds the context from another
goroutine and checks that each of the three calls blocks until it is
released; it fails on the previous code for all three.
CHANGELOG v0.34.6 section; dependency versions in README, AGENTS.md,
ROADMAP, ARCHITECTURE and the Android preview docs; the Android preview
check script pins goffi v0.6.7 (it asserts the exact selected version).
@kolkov kolkov mentioned this pull request Oct 18, 2026
1 of 2 tasks
@codecov

codecov Bot commented Oct 18, 2026

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 63.49206% with 23 lines in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
internal/shaderir/shaderir.go 63.49% 18 Missing and 5 partials ⚠️

📢 Thoughts on this report? Let us know!

…ferToTexture

CopyBufferToTexture copied width*height*layers bytes as one tight block to
the start of the texture, so a source with padded rows (bytesPerRow is a
multiple of 256 for multi-row copies) was sheared and the destination
origin was ignored. Rows are now copied one by one at the source
bytesPerRow / rowsPerImage into the region at the texture origin, per
layer, mirroring CopyTextureToBuffer.
CopyTextureToBuffer read with a hard-coded GL_BGRA (R and B swapped for
RGBA8), flipped the rows (shaders already use ADJUST_COORDINATE_SPACE, so
texture row 0 is the image top) and ignored bytesPerRow, so readback came
back mirrored, blue for red, and sheared with 256-byte rows. It also
returned zeros when the copy did not fill the whole destination buffer:
MapBuffer uses the CPU shadow only when it covers the buffer. It now reads
in the texture's transfer format and type, keeps memory row order, writes
rows at bytesPerRow, and starts a full-size shadow from the GL buffer.
CopyBufferToTexture now sets GL_UNPACK_ROW_LENGTH from bytesPerRow. Both
follow Rust wgpu-hal gles/queue.rs.

On Linux, CreateTexture reported a GL error left by an earlier call as its
own (Mesa's d3d12 driver in WSL: the first texture after device creation
failed with GL_INVALID_ENUM, so examples/triangle-headless could not run);
it now drains and logs stale errors before allocating.

TestTextureCopyRoundTrip uploads four colors with WriteTexture and with a
256-byte-pitch CopyBufferToTexture and reads them back with a 256-byte
pitch on every backend. On main it fails for GLES (both uploads) and for
the software CopyBufferToTexture.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant