From 87a10619b4f723d5300e81328f727119a48294d0 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 16:21:18 +0300 Subject: [PATCH 01/16] fix: a query set outlives its Release until the work that recorded it completes Release is Rust's drop, not WebGPU's destroy (ADR-056). A query set released after timestamp writes or a resolve were recorded was destroyed at the previous submission index, so the HAL query pool went away under the submitted work: VUID-vkDestroyQueryPool-queryPool-00793 on Vulkan and zero timestamps on DX12. CreateQuerySet gives the set a core.ResourceRef and Release drops it. The command encoder clones the reference when a render or compute pass records timestamp writes or ResolveQuerySet records a resolve (Rust wgpu-core's command buffer trackers hold Arc), and Submit rejects only a set whose HAL query set is gone, as query_set.try_raw does (device/queue.rs:2167-2168). Test first: hal/allbackends/query_set_release_test.go failed on Vulkan (VVL) and DX12 before this change. P1-CORE-021. --- encoder_native.go | 43 ++++++- error.go | 7 ++ hal/allbackends/query_set_release_test.go | 141 ++++++++++++++++++++++ query_native.go | 30 ++++- queue_native.go | 11 ++ 5 files changed, 228 insertions(+), 4 deletions(-) create mode 100644 hal/allbackends/query_set_release_test.go diff --git a/encoder_native.go b/encoder_native.go index 12a683b9..16de67c1 100644 --- a/encoder_native.go +++ b/encoder_native.go @@ -48,6 +48,12 @@ type CommandEncoder struct { // (device/queue.rs:1815-1817). usedBindGroups map[*BindGroup]struct{} + // usedQuerySets tracks query sets referenced during encoding for + // submit-time validation. At Submit, a set whose HAL query set is gone + // causes an error, as Rust wgpu-core's query_set.try_raw does + // (device/queue.rs:2167-2168). + usedQuerySets map[*QuerySet]struct{} + // internal holds the resources the frontend created for the encoder's own // commands (the stencil upload blit's staging texture, views, uniform // buffer and bind groups). They have no application owner, so the command @@ -143,6 +149,25 @@ func (e *CommandEncoder) trackBindGroup(bg *BindGroup) { e.usedBindGroups[bg] = struct{}{} } +// trackQuerySet records a query set reference for submit-time validation and, +// the first time, clones the set's reference so it stays alive until the +// command buffer's submission completes, as Rust wgpu-core's command buffer +// trackers hold an Arc. The map is lazily initialized to avoid +// allocation when no query sets are used (P1-CORE-021). +func (e *CommandEncoder) trackQuerySet(qs *QuerySet) { + if qs == nil || qs.core == nil { + return + } + if e.usedQuerySets == nil { + e.usedQuerySets = make(map[*QuerySet]struct{}) + } + if _, seen := e.usedQuerySets[qs]; seen { + return + } + e.usedQuerySets[qs] = struct{}{} + e.trackRef(qs.core.Ref) +} + // copyTextureUsage describes one texture endpoint of a copy command. type copyTextureUsage struct { texture *Texture @@ -343,6 +368,9 @@ func (e *CommandEncoder) BeginRenderPass(desc *RenderPassDescriptor) (*RenderPas ds := desc.DepthStencilAttachment pass.depthReadOnly, pass.stencilReadOnly = ds.DepthReadOnly, ds.StencilReadOnly } + if desc != nil && desc.TimestampWrites != nil { + e.trackQuerySet(desc.TimestampWrites.QuerySet) + } return pass, nil } @@ -371,6 +399,9 @@ func (e *CommandEncoder) BeginComputePass(desc *ComputePassDescriptor) (*Compute return nil, err } + if desc != nil && desc.TimestampWrites != nil { + e.trackQuerySet(desc.TimestampWrites.QuerySet) + } return &ComputePassEncoder{core: corePass, encoder: e}, nil } @@ -814,7 +845,7 @@ func (e *CommandEncoder) ResolveQuerySet(querySet *QuerySet, firstQuery, queryCo } e.trackRef(dst.core.Ref) e.trackBuffer(dst) - e.trackRef(querySet.core.Ref) + e.trackQuerySet(querySet) raw.ResolveQuerySet(halQuerySet, firstQuery, queryCount, halDestination, dstOffset) } @@ -907,6 +938,7 @@ func (e *CommandEncoder) Finish() (*CommandBuffer, error) { usedBuffers: e.usedBuffers, usedTextures: e.usedTextures, usedBindGroups: e.usedBindGroups, + usedQuerySets: e.usedQuerySets, internal: e.internal, } e.trackedRefs = nil @@ -915,6 +947,7 @@ func (e *CommandEncoder) Finish() (*CommandBuffer, error) { e.usedBuffers = nil // ownership transferred e.usedTextures = nil // ownership transferred e.usedBindGroups = nil // ownership transferred + e.usedQuerySets = nil // ownership transferred if e.device != nil { e.device.recordings.addBuffer(cb) } @@ -1056,6 +1089,11 @@ type CommandBuffer struct { // (device/queue.rs:1815-1817). usedBindGroups map[*BindGroup]struct{} + // usedQuerySets tracks all query sets referenced during encoding. + // Validated at Submit time: a set whose HAL query set is gone causes an + // error, as Rust wgpu-core's query_set.try_raw does (device/queue.rs:2167-2168). + usedQuerySets map[*QuerySet]struct{} + // internal is CommandEncoder.internal, released once the command buffer // is spent. internal []internalResource @@ -1123,7 +1161,7 @@ func releaseSubmittedInternal(commandBuffers []*CommandBuffer) { } // dropUsedSets releases the encode-time validation sets. usedBuffers, -// usedTextures and usedBindGroups exist only for +// usedTextures, usedBindGroups and usedQuerySets exist only for // validateCommandBufferForSubmit; once a command buffer is spent — submitted // or released — they are hard references pinning every resource the frame // touched for as long as the command buffer stays reachable. @@ -1131,6 +1169,7 @@ func (cb *CommandBuffer) dropUsedSets() { cb.usedBuffers = nil cb.usedTextures = nil cb.usedBindGroups = nil + cb.usedQuerySets = nil } // halBufferList returns all HAL command buffers in submission order. diff --git a/error.go b/error.go index 8cab12dc..b5025476 100644 --- a/error.go +++ b/error.go @@ -173,6 +173,13 @@ var ( // matches it still compiles. ErrSubmitBindGroupDestroyed = errors.New("wgpu: Submit: command buffer references destroyed bind group") + // ErrSubmitQuerySetDestroyed reports a submitted command buffer that + // references a query set whose HAL query set is gone. Release() after + // recording is not that: the command buffer keeps the set alive until its + // work completes. Matches Rust QueueSubmitError::DestroyedResource for + // query sets (query_set.try_raw, device/queue.rs:2167-2168). + ErrSubmitQuerySetDestroyed = errors.New("wgpu: Submit: command buffer references destroyed query set") + // ErrSubmitCommandBufferInvalid is returned when a command buffer has // already been submitted or was never properly finished. // Matches Rust QueueSubmitError::CommandEncoder for invalid command buffers. diff --git a/hal/allbackends/query_set_release_test.go b/hal/allbackends/query_set_release_test.go new file mode 100644 index 00000000..0379b1bc --- /dev/null +++ b/hal/allbackends/query_set_release_test.go @@ -0,0 +1,141 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build !android && !(js && wasm) && !rust + +package allbackends_test + +import ( + "context" + "encoding/binary" + "testing" + "time" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu" +) + +// Query indices of the compute and the render pass of the release test. +const ( + releaseComputeBegin = iota + releaseComputeEnd + releaseRenderBegin + releaseRenderEnd + releaseQueryCount +) + +// TestReleaseQuerySetBeforeSubmit: a timestamp query set written at the +// boundaries of a compute pass and of a render pass, then resolved, is +// released before Submit. Release is a drop, not a destroy (ADR-056): the +// command buffer holds the query set's reference, so the HAL query pool must +// stay alive until the submitted work completes (P1-CORE-021; +// VUID-vkDestroyQueryPool-queryPool-00793). The readback also proves the +// timestamps were still written. +func TestReleaseQuerySetBeforeSubmit(t *testing.T) { + forEachOverrideBackendFeatures(t, 0, gputypes.Features(gputypes.FeatureTimestampQuery), + func(t *testing.T, device *wgpu.Device) { + if !device.Features().Contains(gputypes.FeatureTimestampQuery) { + t.Skip("adapter does not offer TimestampQuery") + } + querySet, err := device.CreateQuerySet(&wgpu.QuerySetDescriptor{ + Label: "release-timestamps", + Type: wgpu.QueryTypeTimestamp, + Count: releaseQueryCount, + }) + if err != nil { + t.Fatalf("CreateQuerySet: %v", err) + } + resolve := newPublicBuffer(t, device, releaseQueryCount*8, + wgpu.BufferUsageQueryResolve|wgpu.BufferUsageCopySrc) + staging := newPublicBuffer(t, device, releaseQueryCount*8, + wgpu.BufferUsageCopyDst|wgpu.BufferUsageMapRead) + + enc := newEncoder(t, device) + computeBegin, computeEnd := uint32(releaseComputeBegin), uint32(releaseComputeEnd) + cpass, err := enc.BeginComputePass(&wgpu.ComputePassDescriptor{ + Label: "release-timestamps-compute", + TimestampWrites: &wgpu.ComputePassTimestampWrites{ + QuerySet: querySet, + BeginningOfPassWriteIndex: &computeBegin, + EndOfPassWriteIndex: &computeEnd, + }, + }) + if err != nil { + t.Fatalf("BeginComputePass: %v", err) + } + if err := cpass.End(); err != nil { + t.Fatalf("compute End: %v", err) + } + + target := newReleaseTexture(t, device, wgpu.TextureUsageRenderAttachment) + t.Cleanup(target.Release) + view, err := device.CreateTextureView(target, nil) + if err != nil { + t.Fatalf("CreateTextureView: %v", err) + } + t.Cleanup(view.Release) + renderBegin, renderEnd := uint32(releaseRenderBegin), uint32(releaseRenderEnd) + rpass, err := enc.BeginRenderPass(&wgpu.RenderPassDescriptor{ + Label: "release-timestamps-render", + ColorAttachments: []wgpu.RenderPassColorAttachment{{ + View: view, LoadOp: gputypes.LoadOpClear, StoreOp: gputypes.StoreOpStore, + ClearValue: gputypes.Color{R: 0, G: 0, B: 1, A: 1}, + }}, + TimestampWrites: &wgpu.RenderPassTimestampWrites{ + QuerySet: querySet, + BeginningOfPassWriteIndex: &renderBegin, + EndOfPassWriteIndex: &renderEnd, + }, + }) + if err != nil { + t.Fatalf("BeginRenderPass: %v", err) + } + if err := rpass.End(); err != nil { + t.Fatalf("render End: %v", err) + } + + enc.ResolveQuerySet(querySet, 0, releaseQueryCount, resolve, 0) + enc.CopyBufferToBuffer(resolve, 0, staging, 0, releaseQueryCount*8) + + // Released after the work that uses it was recorded, before Submit. + querySet.Release() + submit(t, device, enc) + + ts := readQueryTicks(t, staging, releaseQueryCount) + for _, pair := range [][2]int{ + {releaseComputeBegin, releaseComputeEnd}, + {releaseRenderBegin, releaseRenderEnd}, + } { + if ts[pair[1]] == 0 { + t.Errorf("timestamp %d is 0: the end-of-pass write did not happen", pair[1]) + } + if ts[pair[1]] < ts[pair[0]] { + t.Errorf("timestamp %d (%d) is before timestamp %d (%d)", + pair[1], ts[pair[1]], pair[0], ts[pair[0]]) + } + } + }) +} + +// readQueryTicks maps the staging buffer and returns its timestamps in ticks. +func readQueryTicks(t *testing.T, staging *wgpu.Buffer, count int) []uint64 { + t.Helper() + size := uint64(count * 8) + ctx, cancel := context.WithTimeout(context.Background(), 10*time.Second) + defer cancel() + if err := staging.Map(ctx, wgpu.MapModeRead, 0, size); err != nil { + t.Fatalf("Map: %v", err) + } + rng, err := staging.MappedRange(0, size) + if err != nil { + t.Fatalf("MappedRange: %v", err) + } + ts := make([]uint64, count) + for i := range ts { + ts[i] = binary.LittleEndian.Uint64(rng.Bytes()[i*8:]) + } + if err := staging.Unmap(); err != nil { + t.Fatalf("Unmap: %v", err) + } + return ts +} diff --git a/query_native.go b/query_native.go index 89369812..47f7a693 100644 --- a/query_native.go +++ b/query_native.go @@ -55,8 +55,27 @@ func (qs *QuerySet) liveCore() *core.QuerySet { return qs.core } -// Release destroys the query set. Destruction is deferred until the GPU -// completes any submission that may reference this query set. +// newQuerySetRef returns the reference count of a HAL query set: when its last +// owner drops it, the set is destroyed after the submissions in flight, as a +// sampler or a bind group is. The owners are the application until Release and +// every command buffer that recorded the set's timestamp writes or a resolve +// of it (P1-CORE-021, Rust wgpu-core's QuerySet is held by an Arc). +func newQuerySetRef(d *Device, coreQS *core.QuerySet, label string) *core.ResourceRef { + return core.NewResourceRef("QuerySet:"+label, func() { + dq := d.destroyQueue() + if dq == nil { + coreQS.Destroy() + return + } + dq.Defer(d.lastSubmissionIndex(), "QuerySet", coreQS.Destroy) + }) +} + +// Release drops the application's reference to the query set. The HAL query +// set is destroyed once no command buffer that recorded it references it any +// more and the GPU has completed those submissions: a query set released after +// its timestamp writes or a resolve were recorded stays valid for that work +// (P1-CORE-021). func (qs *QuerySet) Release() { if qs == nil || qs.released { return @@ -65,6 +84,12 @@ func (qs *QuerySet) Release() { if qs.core == nil { return } + if qs.core.Ref != nil { + qs.core.Ref.Drop() + return + } + + // No reference count: destroy after the submissions in flight. halDevice := qs.device.halDevice() if halDevice == nil { qs.core.Destroy() @@ -109,5 +134,6 @@ func (d *Device) CreateQuerySet(desc *QuerySetDescriptor) (*QuerySet, error) { } coreQS := core.NewQuerySet(halQS, d.core, halDesc.Type, halDesc.Count, desc.Label) + coreQS.Ref = newQuerySetRef(d, coreQS, desc.Label) return &QuerySet{core: coreQS, device: d}, nil } diff --git a/queue_native.go b/queue_native.go index 157a2734..08563af9 100644 --- a/queue_native.go +++ b/queue_native.go @@ -576,6 +576,17 @@ func validateCommandBufferForSubmit(cb *CommandBuffer, index int) error { } } + // 5. Check the query sets the command buffer writes or resolves, as Rust's + // query_set.try_raw does (queue.rs:2167-2168). Release() is a drop: the + // command buffer holds the set's reference until its submission completes + // (P1-CORE-021). Only a query set whose HAL query set is gone fails. + for qs := range cb.usedQuerySets { + if qs.core == nil || qs.core.IsDestroyed() { + return fmt.Errorf("wgpu: Submit: command buffer at index %d references destroyed query set: %w", + index, ErrSubmitQuerySetDestroyed) + } + } + return nil } From 111fbbd0e768ec29d3b3891a8e83736ac75b60de Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 16:42:49 +0300 Subject: [PATCH 02/16] fix: pipelines keep their pipeline layout alive RenderPipeline and ComputePipeline kept no reference to the PipelineLayout they were created from. WebGPU has no destroy() for a layout, so the only way an application lets one go is Release, which is Rust's drop: it deferred vkDestroyPipelineLayout, and every later SetBindGroup of the pipeline passed the destroyed handle to vkCmdBindDescriptorSets (VUID-vkCmdBindDescriptorSets- layout-parameter; on the test machine the Vulkan driver crashed with an access violation). Rust wgpu-core holds an Arc in both pipelines (pipeline.rs:496, :1155). PipelineLayout now has a core.ResourceRef whose onZero defers the HAL destroy through the destroy queue at the last submission index, as a sampler's does, and both pipeline kinds clone it. The pipeline drops it in the same onZero that destroys its own HAL pipeline: pipeline first, then layout, so a released layout stays valid for the pipelines created with it. A pipeline created from an already released layout is rejected with ErrReleased instead of naming a dying HAL handle and reviving its reference count. Bind group layouts keep no reference: VkPipelineLayoutCreateInfo says the implementation must not access the pSetLayouts objects outside the duration of vkCreatePipelineLayout, so a VkDescriptorSetLayout may be destroyed after a dependent pipeline layout exists. Rust and Dawn hold their own bind group layout objects for validation, not to keep that Vulkan handle alive. Test-first in hal/allbackends: both pipeline kinds are created from an explicit layout, the layout is released, WaitIdle runs the deferred destruction, and only then does a new pass bind and draw or dispatch. Unfixed the Vulkan driver crashes in vkCmdBindDescriptorSets; fixed, every backend reads back the expected values with the validation layer quiet. P2-CORE-022 --- bind_native.go | 35 ++- device_native.go | 65 +++- .../pipeline_layout_release_test.go | 279 ++++++++++++++++++ 3 files changed, 365 insertions(+), 14 deletions(-) create mode 100644 hal/allbackends/pipeline_layout_release_test.go diff --git a/bind_native.go b/bind_native.go index 40746090..f0176770 100644 --- a/bind_native.go +++ b/bind_native.go @@ -129,16 +129,47 @@ type PipelineLayout struct { // bindGroupLayouts stores the layouts used to create this pipeline layout. // Used by the binder for draw-time compatibility validation. bindGroupLayouts []*BindGroupLayout + // ref counts the layout's owners: the application until Release, and + // every pipeline created with it until that pipeline is destroyed, as + // Rust wgpu-core's RenderPipeline and ComputePipeline hold an + // Arc (pipeline.rs:496, :1155). The HAL layout is + // destroyed when the last one drops (ADR-056, P2-CORE-022). + ref *core.ResourceRef } -// Release destroys the pipeline layout. Destruction is deferred until the -// GPU completes any submission that may reference this layout. +// newPipelineLayoutRef returns the reference count of a HAL pipeline layout: +// when its last owner drops it, the layout is destroyed after the submissions +// in flight, as a sampler or a bind group is. The owners are the application +// until Release and every pipeline created with the layout, which passes it to +// vkCmdBindDescriptorSets on every bind (P2-CORE-022). +func newPipelineLayoutRef(d *Device, halDevice hal.Device, halLayout hal.PipelineLayout, label string) *core.ResourceRef { + return core.NewResourceRef("PipelineLayout:"+label, func() { + dq := d.destroyQueue() + if dq == nil { + halDevice.DestroyPipelineLayout(halLayout) + return + } + dq.Defer(d.lastSubmissionIndex(), "PipelineLayout", func() { + halDevice.DestroyPipelineLayout(halLayout) + }) + }) +} + +// Release drops the application's reference to the pipeline layout. The HAL +// layout is destroyed once no pipeline created with it exists any more and the +// GPU has completed every submission that may use it: a layout released after +// its pipeline was created stays valid for that pipeline's binds (P2-CORE-022). func (l *PipelineLayout) Release() { if l.released { return } l.released = true + if l.ref != nil { + l.ref.Drop() + return + } + // No reference count: destroy after the submissions in flight. halDevice := l.device.halDevice() if halDevice == nil { return diff --git a/device_native.go b/device_native.go index 68b2357d..e55de2ef 100644 --- a/device_native.go +++ b/device_native.go @@ -417,6 +417,7 @@ func (d *Device) CreatePipelineLayout(desc *PipelineLayoutDescriptor) (*Pipeline device: d, bindGroupCount: uint32(len(desc.BindGroupLayouts)), //nolint:gosec // layout count fits uint32 bindGroupLayouts: bgLayouts, + ref: newPipelineLayoutRef(d, halDevice, halLayout, desc.Label), }, nil } @@ -654,6 +655,12 @@ func (d *Device) CreateRenderPipeline(desc *RenderPipelineDescriptor) (*RenderPi if desc == nil { return nil, fmt.Errorf("wgpu: render pipeline descriptor is nil") } + // A released layout may already be queued for destruction; the pipeline + // would be created from a dying HAL layout and would later revive its + // reference count (P2-CORE-022). + if desc.Layout != nil && desc.Layout.released { + return nil, ErrReleased + } halDevice := d.halDevice() if halDevice == nil { @@ -678,7 +685,16 @@ func (d *Device) CreateRenderPipeline(desc *RenderPipelineDescriptor) (*RenderPi var bgCount uint32 var bgLayouts []*BindGroupLayout + var layoutRef *core.ResourceRef if desc.Layout != nil { + // The pipeline keeps its layout alive until its own HAL pipeline is + // destroyed: every later SetBindGroup passes the layout to + // vkCmdBindDescriptorSets. Rust wgpu-core holds an Arc + // in the pipeline (pipeline.rs:496). + layoutRef = desc.Layout.ref + if layoutRef != nil { + layoutRef.Clone() + } bgCount = desc.Layout.bindGroupCount bgLayouts = desc.Layout.bindGroupLayouts } @@ -712,15 +728,20 @@ func (d *Device) CreateRenderPipeline(desc *RenderPipelineDescriptor) (*RenderPi // DestroyQueue. ADR-056: unified resource lifecycle. halRP := halPipeline rpOnZero := func() { + // Pipeline first, then the layout it was created from: the layout's + // own destruction waits for the submissions in flight. + destroy := func() { + halDevice.DestroyRenderPipeline(halRP) + if layoutRef != nil { + layoutRef.Drop() + } + } dq := d.destroyQueue() if dq != nil { - subIdx := d.lastSubmissionIndex() - dq.Defer(subIdx, "RenderPipeline", func() { - halDevice.DestroyRenderPipeline(halRP) - }) - } else { - halDevice.DestroyRenderPipeline(halRP) + dq.Defer(d.lastSubmissionIndex(), "RenderPipeline", destroy) + return } + destroy() } passContext, writesDepth, writesStencil := renderPipelineTargets(desc) @@ -805,6 +826,12 @@ func (d *Device) CreateComputePipeline(desc *ComputePipelineDescriptor) (*Comput if desc == nil { return nil, fmt.Errorf("wgpu: compute pipeline descriptor is nil") } + // A released layout may already be queued for destruction; the pipeline + // would be created from a dying HAL layout and would later revive its + // reference count (P2-CORE-022). + if desc.Layout != nil && desc.Layout.released { + return nil, ErrReleased + } // Validate that the device supports compute shaders (downlevel check). // Matches Rust wgpu-core resource.rs:4367 — first validation in create_compute_pipeline_or_error. @@ -834,7 +861,16 @@ func (d *Device) CreateComputePipeline(desc *ComputePipelineDescriptor) (*Comput var bgCount uint32 var bgLayouts []*BindGroupLayout + var layoutRef *core.ResourceRef if desc.Layout != nil { + // The pipeline keeps its layout alive until its own HAL pipeline is + // destroyed: every later SetBindGroup passes the layout to + // vkCmdBindDescriptorSets. Rust wgpu-core holds an Arc + // in the pipeline (pipeline.rs:1155). + layoutRef = desc.Layout.ref + if layoutRef != nil { + layoutRef.Clone() + } bgCount = desc.Layout.bindGroupCount bgLayouts = desc.Layout.bindGroupLayouts } @@ -848,15 +884,20 @@ func (d *Device) CreateComputePipeline(desc *ComputePipelineDescriptor) (*Comput // ADR-056: unified resource lifecycle. halCP := halPipeline cpOnZero := func() { + // Pipeline first, then the layout it was created from: the layout's + // own destruction waits for the submissions in flight. + destroy := func() { + halDevice.DestroyComputePipeline(halCP) + if layoutRef != nil { + layoutRef.Drop() + } + } dq := d.destroyQueue() if dq != nil { - subIdx := d.lastSubmissionIndex() - dq.Defer(subIdx, "ComputePipeline", func() { - halDevice.DestroyComputePipeline(halCP) - }) - } else { - halDevice.DestroyComputePipeline(halCP) + dq.Defer(d.lastSubmissionIndex(), "ComputePipeline", destroy) + return } + destroy() } return &ComputePipeline{ diff --git a/hal/allbackends/pipeline_layout_release_test.go b/hal/allbackends/pipeline_layout_release_test.go new file mode 100644 index 00000000..ad557119 --- /dev/null +++ b/hal/allbackends/pipeline_layout_release_test.go @@ -0,0 +1,279 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build !android && !(js && wasm) && !rust + +package allbackends_test + +import ( + "errors" + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu" +) + +// A pipeline keeps its pipeline layout alive. WebGPU has no destroy() for a +// GPUPipelineLayout, so the only way an application lets one go is Release, +// which is Rust's drop (ADR-056). A pipeline created with that layout still +// needs it: every later SetBindGroup passes the layout to +// vkCmdBindDescriptorSets, and the pipeline was created from it. Rust +// wgpu-core holds an Arc in the pipeline +// (pipeline.rs:496, :1155), so the HAL layout is destroyed only with the last +// pipeline. Without the same reference the released layout is destroyed at +// the next GPU completion and the following bind names a destroyed +// VkPipelineLayout (VUID-vkCmdBindDescriptorSets-layout-parameter), which the +// validation layer reports and this harness turns into a test failure +// (P2-CORE-022). + +const pipelineLayoutReleaseComputeWGSL = ` +@group(0) @binding(0) var out: array; + +@compute @workgroup_size(4) +fn main(@builtin(global_invocation_id) id: vec3) { + out[id.x] = id.x + 1u; +} +` + +// TestReleasedPipelineLayoutComputePipeline creates a compute pipeline from a +// pipeline layout, releases the layout, lets the deferred destruction run, and +// only then dispatches through the pipeline. +func TestReleasedPipelineLayoutComputePipeline(t *testing.T) { + forEachOverrideBackend(t, gputypes.DownlevelFlagsComputeShaders, func(t *testing.T, device *wgpu.Device) { + bgl := newBindGroupLayout(t, device, gputypes.BindGroupLayoutEntry{ + Binding: 0, Visibility: wgpu.ShaderStageCompute, + Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeStorage}, + }) + out := newPublicBuffer(t, device, releaseSize*4, wgpu.BufferUsageStorage|wgpu.BufferUsageCopySrc) + staging := newPublicBuffer(t, device, releaseSize*4, wgpu.BufferUsageCopyDst|wgpu.BufferUsageMapRead) + group := newBindGroup(t, device, bgl, wgpu.BindGroupEntry{Binding: 0, Buffer: out, Size: releaseSize * 4}) + + shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{ + Label: "pipeline-layout-release", WGSL: pipelineLayoutReleaseComputeWGSL, + }) + if err != nil { + t.Fatalf("CreateShaderModule: %v", err) + } + t.Cleanup(shader.Release) + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{ + Label: "released", + BindGroupLayouts: []*wgpu.BindGroupLayout{bgl}, + }) + if err != nil { + t.Fatalf("CreatePipelineLayout: %v", err) + } + pipeline, err := device.CreateComputePipeline(&wgpu.ComputePipelineDescriptor{ + Label: "pipeline-of-a-released-layout", Layout: layout, Module: shader, EntryPoint: "main", + }) + if err != nil { + t.Fatalf("CreateComputePipeline: %v", err) + } + t.Cleanup(pipeline.Release) + + releasePipelineLayout(t, device, layout) + + enc := newEncoder(t, device) + pass, err := enc.BeginComputePass(nil) + if err != nil { + t.Fatalf("BeginComputePass: %v", err) + } + pass.SetPipeline(pipeline) + pass.SetBindGroup(0, group, nil) + pass.Dispatch(1, 1, 1) + if err := pass.End(); err != nil { + t.Fatalf("ComputePass.End: %v", err) + } + enc.CopyBufferToBuffer(out, 0, staging, 0, releaseSize*4) + submit(t, device, enc) + + checkWords(t, "dispatch through a pipeline whose layout was released", + readPublicBuffer(t, staging, releaseSize), []uint32{1, 2, 3, 4}) + }) +} + +const pipelineLayoutReleaseRenderWGSL = ` +@group(0) @binding(0) var color: vec4; + +@vertex +fn vs(@builtin(vertex_index) i: u32) -> @builtin(position) vec4 { + var p = array, 3>(vec2(-1.0, -1.0), vec2(3.0, -1.0), vec2(-1.0, 3.0)); + return vec4(p[i], 0.0, 1.0); +} + +@fragment +fn fs() -> @location(0) vec4 { + return color; +} +` + +// TestReleasedPipelineLayoutRenderPipeline is the render pipeline's half of +// TestReleasedPipelineLayoutComputePipeline: the layout is released, the +// deferred destruction runs, and only then does a pass bind the pipeline and +// its bind group and draw. +func TestReleasedPipelineLayoutRenderPipeline(t *testing.T) { + forEachOverrideBackend(t, 0, func(t *testing.T, device *wgpu.Device) { + bgl := newBindGroupLayout(t, device, gputypes.BindGroupLayoutEntry{ + Binding: 0, Visibility: wgpu.ShaderStageFragment, + Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeUniform}, + }) + uniform := newPublicBuffer(t, device, 16, wgpu.BufferUsageUniform|wgpu.BufferUsageCopyDst) + writeWords(t, device, uniform, mathFloatBits(0, 1, 0, 1)...) // green + group := newBindGroup(t, device, bgl, wgpu.BindGroupEntry{Binding: 0, Buffer: uniform, Size: 16}) + + shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{ + Label: "pipeline-layout-release", WGSL: pipelineLayoutReleaseRenderWGSL, + }) + if err != nil { + t.Fatalf("CreateShaderModule: %v", err) + } + t.Cleanup(shader.Release) + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{ + Label: "released", + BindGroupLayouts: []*wgpu.BindGroupLayout{bgl}, + }) + if err != nil { + t.Fatalf("CreatePipelineLayout: %v", err) + } + pipeline, err := device.CreateRenderPipeline(&wgpu.RenderPipelineDescriptor{ + Label: "pipeline-of-a-released-layout", + Layout: layout, + Vertex: wgpu.VertexState{Module: shader, EntryPoint: "vs"}, + Fragment: &wgpu.FragmentState{ + Module: shader, + EntryPoint: "fs", + Targets: []gputypes.ColorTargetState{ + {Format: gputypes.TextureFormatRGBA8Unorm, WriteMask: gputypes.ColorWriteMaskAll}, + }, + }, + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}, + Multisample: gputypes.MultisampleState{Count: 1, Mask: 0xFFFFFFFF}, + }) + if err != nil { + t.Fatalf("CreateRenderPipeline: %v", err) + } + t.Cleanup(pipeline.Release) + + releasePipelineLayout(t, device, layout) + + target := newPublicBufferTarget(t, device) + staging := newPublicBuffer(t, device, pipelineLayoutTargetRowPitch*overrideTargetHeight, + wgpu.BufferUsageCopyDst|wgpu.BufferUsageMapRead) + + enc := newEncoder(t, device) + pass, err := enc.BeginRenderPass(&wgpu.RenderPassDescriptor{ + ColorAttachments: []wgpu.RenderPassColorAttachment{{ + View: target.view, + LoadOp: gputypes.LoadOpClear, + StoreOp: gputypes.StoreOpStore, + ClearValue: gputypes.Color{R: 1, G: 0, B: 0, A: 1}, + }}, + }) + if err != nil { + t.Fatalf("BeginRenderPass: %v", err) + } + pass.SetPipeline(pipeline) + pass.SetBindGroup(0, group, nil) + pass.Draw(gputypes.DrawArgs{VertexCount: 3, InstanceCount: 1}) + if err := pass.End(); err != nil { + t.Fatalf("RenderPass.End: %v", err) + } + enc.CopyTextureToBuffer(target.texture, staging, []wgpu.BufferTextureCopy{{ + BufferLayout: wgpu.ImageDataLayout{BytesPerRow: pipelineLayoutTargetRowPitch, RowsPerImage: overrideTargetHeight}, + TextureBase: wgpu.ImageCopyTexture{Texture: target.texture}, + Size: wgpu.Extent3D{Width: overrideTargetWidth, Height: overrideTargetHeight, DepthOrArrayLayers: 1}, + }}) + submit(t, device, enc) + + // Texel (0, 0) of row 0 is the first word of the staging copy. + pixel := readPublicBuffer(t, staging, overrideTargetWidth)[0] + if want := uint32(0xFF00FF00); pixel != want { + t.Errorf("texel (0,0) = %#08x, want %#08x (green)", pixel, want) + } + }) +} + +// TestReleasedPipelineLayoutCannotCreatePipeline: a released layout may +// already be queued for destruction, so a pipeline must not be created from +// it — the pipeline would name a dying HAL layout and revive its reference +// count, destroying it twice. A released sampler is rejected the same way by +// CreateBindGroup. +func TestReleasedPipelineLayoutCannotCreatePipeline(t *testing.T) { + forEachOverrideBackend(t, 0, func(t *testing.T, device *wgpu.Device) { + shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{ + Label: "pipeline-layout-released", WGSL: pipelineLayoutReleaseRenderWGSL, + }) + if err != nil { + t.Fatalf("CreateShaderModule: %v", err) + } + t.Cleanup(shader.Release) + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{Label: "released"}) + if err != nil { + t.Fatalf("CreatePipelineLayout: %v", err) + } + layout.Release() + + pipeline, err := device.CreateRenderPipeline(&wgpu.RenderPipelineDescriptor{ + Label: "of-a-released-layout", + Layout: layout, + Vertex: wgpu.VertexState{Module: shader, EntryPoint: "vs"}, + Fragment: &wgpu.FragmentState{ + Module: shader, + EntryPoint: "fs", + Targets: []gputypes.ColorTargetState{ + {Format: gputypes.TextureFormatRGBA8Unorm, WriteMask: gputypes.ColorWriteMaskAll}, + }, + }, + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}, + Multisample: gputypes.MultisampleState{Count: 1, Mask: 0xFFFFFFFF}, + }) + if !errors.Is(err, wgpu.ErrReleased) { + if pipeline != nil { + pipeline.Release() + } + t.Fatalf("CreateRenderPipeline with a released layout = %v, want ErrReleased", err) + } + }) +} + +const pipelineLayoutTargetRowPitch = 256 // overrideTargetWidth RGBA8 texels + +type publicBufferTarget struct { + texture *wgpu.Texture + view *wgpu.TextureView +} + +// newPublicBufferTarget creates the 64x4 RGBA8 render target these tests draw +// into, with a copy source and a view of its only mip level. +func newPublicBufferTarget(t *testing.T, device *wgpu.Device) publicBufferTarget { + t.Helper() + texture, err := device.CreateTexture(&wgpu.TextureDescriptor{ + Label: "pipeline-layout-release", + Size: wgpu.Extent3D{Width: overrideTargetWidth, Height: overrideTargetHeight, DepthOrArrayLayers: 1}, + MipLevelCount: 1, + SampleCount: 1, + Dimension: gputypes.TextureDimension2D, + Format: gputypes.TextureFormatRGBA8Unorm, + Usage: wgpu.TextureUsageRenderAttachment | wgpu.TextureUsageCopySrc, + }) + if err != nil { + t.Fatalf("CreateTexture: %v", err) + } + t.Cleanup(texture.Release) + view, err := device.CreateTextureView(texture, nil) + if err != nil { + t.Fatalf("CreateTextureView: %v", err) + } + t.Cleanup(view.Release) + return publicBufferTarget{texture: texture, view: view} +} + +// releasePipelineLayout releases a pipeline layout and lets the deferred +// destruction run, so a layout that no longer has an owner is really destroyed +// before the work recorded afterwards. +func releasePipelineLayout(t *testing.T, device *wgpu.Device, layout *wgpu.PipelineLayout) { + t.Helper() + layout.Release() + if err := device.WaitIdle(); err != nil { + t.Fatalf("WaitIdle: %v", err) + } +} From 12358bedbf7a713a09d5c48d030cdd62ecd299e0 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 16:48:37 +0300 Subject: [PATCH 03/16] =?UTF-8?q?fix:=20build=20on=20GOARCH=3D386=20?= =?UTF-8?q?=E2=80=94=20int=20constants=20no=20longer=20overflow?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Four comparisons against TrackerIndex bounds were written as int constants: int(^TrackerIndex(0)-1) is 4294967294 and does not fit a 32-bit int, so core/track (buffer, texture, encoder tracker) and internal/raytracing (math.MaxUint32 in safeIntToUint32) failed to compile for GOARCH=386. Compare in uint64 instead: the tracker guard is now uint64(i) >= uint64(InvalidTrackerIndex), which rejects the same values on 64-bit and is vacuous on 32-bit, where an int cannot reach the invalid index; the raytracing clamp converts v to uint64 after its negative check. Semantics on amd64, arm64 and 386 are unchanged; 386 was simply never buildable. Reported by gogpu (BUG-ELSEWHERE, v0.34.10). --- core/track/buffer.go | 2 +- core/track/encoder_tracker.go | 2 +- core/track/texture.go | 2 +- internal/raytracing/build.go | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) diff --git a/core/track/buffer.go b/core/track/buffer.go index a3c92b41..409fe9d3 100644 --- a/core/track/buffer.go +++ b/core/track/buffer.go @@ -176,7 +176,7 @@ func (t *BufferTracker) Merge(scope *BufferUsageScope) []PendingTransition { var transitions []PendingTransition for i := range scope.states { - if i < 0 || i > int(^TrackerIndex(0)-1) { + if i < 0 || uint64(i) >= uint64(InvalidTrackerIndex) { continue // Skip if index would overflow TrackerIndex } index := TrackerIndex(i) diff --git a/core/track/encoder_tracker.go b/core/track/encoder_tracker.go index d4094f94..b3a24c96 100644 --- a/core/track/encoder_tracker.go +++ b/core/track/encoder_tracker.go @@ -265,7 +265,7 @@ func (t *EncoderBufferTracker) ensureSize(size int) { // Positions outside the TrackerIndex range (never produced by the allocator) // are reported invalid instead of being truncated. func trackerIndexAt(i int) (TrackerIndex, bool) { - if i < 0 || i > int(^TrackerIndex(0)-1) { + if i < 0 || uint64(i) >= uint64(InvalidTrackerIndex) { return 0, false } return TrackerIndex(i), true diff --git a/core/track/texture.go b/core/track/texture.go index cc959565..4824b55d 100644 --- a/core/track/texture.go +++ b/core/track/texture.go @@ -247,7 +247,7 @@ func (t *TextureTracker) Merge(scope *TextureUsageScope) []TexturePendingTransit var transitions []TexturePendingTransition for i := range scope.states { - if i < 0 || i > int(^TrackerIndex(0)-1) { + if i < 0 || uint64(i) >= uint64(InvalidTrackerIndex) { continue } index := TrackerIndex(i) diff --git a/internal/raytracing/build.go b/internal/raytracing/build.go index 3719ba7f..adc8ac70 100644 --- a/internal/raytracing/build.go +++ b/internal/raytracing/build.go @@ -195,7 +195,7 @@ func safeIntToUint32(v int) uint32 { return 0 } - if v > math.MaxUint32 { + if uint64(v) > math.MaxUint32 { return math.MaxUint32 } From 28e7af19f5d21b92688612a9a92668f112f25ece Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 16:55:31 +0300 Subject: [PATCH 04/16] fix(rust): an unmapped error filter is an error, not a silent Validation scope errorFilterToRust mapped any value outside the three WebGPU filters to Validation, so PushErrorScope(ErrorFilter(42)) silently pushed a scope that could never capture what the caller asked for. Never swallow an error. PushErrorScope now returns an error in every build: the native path validates in core (ValidateErrorFilter, the spec's pushErrorScope TypeError), the rust path rejects the value before it reaches wgpu-native, which panics on an unknown filter and the panic aborts the process across the FFI boundary. The browser stub keeps its panic and only gains the return type. Callers that call PushErrorScope as a statement are unaffected; none in the ecosystem used the return value before because there was none. Tests: a table over every filter constant plus unmapped values in core and under -tags rust. P3-API-025. --- core/error_scope.go | 23 ++++++++++++++++++++++- core/error_scope_test.go | 18 ++++++++++++++++++ device_browser.go | 2 +- device_native.go | 7 +++++-- device_rust.go | 28 +++++++++++++++++++--------- enum_values_rust_test.go | 15 ++++++++++++++- 6 files changed, 79 insertions(+), 14 deletions(-) diff --git a/core/error_scope.go b/core/error_scope.go index f7f101b5..84efd2bf 100644 --- a/core/error_scope.go +++ b/core/error_scope.go @@ -177,6 +177,23 @@ func (m *ErrorScopeManager) ScopeDepth() int { // PushErrorScope pushes a new error scope onto this device's error scope stack. // +// ValidateErrorFilter reports whether f is one of WebGPU's three error +// filters. pushErrorScope throws a TypeError for any other value +// (WebGPU spec, pushErrorScope step 1); in Go the caller gets this error. A +// filter outside the three would never match a captured error, so a scope +// pushed with it would silently capture nothing. +func ValidateErrorFilter(f ErrorFilter) error { + switch f { + case ErrorFilterValidation, ErrorFilterOutOfMemory, ErrorFilterInternal: + return nil + } + return fmt.Errorf("wgpu: PushErrorScope: unknown error filter %d", int(f)) +} + +// PushErrorScope pushes a new error scope onto the device's error scope stack. +// It returns an error for a filter that is not one of the three WebGPU error +// filters, and pushes nothing then. +// // The scope will capture the first error matching the specified filter type. // Error scopes are LIFO (stack-based) -- the last pushed scope is checked first. // @@ -190,8 +207,12 @@ func (m *ErrorScopeManager) ScopeDepth() int { // if gpuErr != nil { // log.Printf("Validation error: %s", gpuErr.Message) // } -func (d *Device) PushErrorScope(filter ErrorFilter) { +func (d *Device) PushErrorScope(filter ErrorFilter) error { + if err := ValidateErrorFilter(filter); err != nil { + return err + } d.errorScopes().PushErrorScope(filter) + return nil } // PopErrorScope pops the most recently pushed error scope and returns the diff --git a/core/error_scope_test.go b/core/error_scope_test.go index ba847a92..a92b71c0 100644 --- a/core/error_scope_test.go +++ b/core/error_scope_test.go @@ -463,3 +463,21 @@ func TestDeviceErrorScopeLazyInit(t *testing.T) { // Clean up _, _ = device.errorScopes().PopErrorScope() } + +// TestValidateErrorFilter checks that every WebGPU error filter passes and +// that a value outside the three constants is an error, never a silent +// Validation fallback (P3-API-025): a scope pushed with an unknown filter +// would capture nothing. +func TestValidateErrorFilter(t *testing.T) { + valid := []ErrorFilter{ErrorFilterValidation, ErrorFilterOutOfMemory, ErrorFilterInternal} + for _, f := range valid { + if err := ValidateErrorFilter(f); err != nil { + t.Errorf("ValidateErrorFilter(%s) = %v, want nil", f, err) + } + } + for _, f := range []ErrorFilter{-1, 3, 42} { + if err := ValidateErrorFilter(f); err == nil { + t.Errorf("ValidateErrorFilter(%d) = nil, want an error", int(f)) + } + } +} diff --git a/device_browser.go b/device_browser.go index 580a54ee..2b2b141a 100644 --- a/device_browser.go +++ b/device_browser.go @@ -332,7 +332,7 @@ func (d *Device) WaitForFence(_ *Fence, _ uint64, _ time.Duration) (bool, error) // PushErrorScope pushes a new error scope onto the device's error scope stack. // Phase 2 — not yet implemented for browser. -func (d *Device) PushErrorScope(filter ErrorFilter) { +func (d *Device) PushErrorScope(filter ErrorFilter) error { panic("wgpu: browser PushErrorScope not yet implemented (Phase 2)") } diff --git a/device_native.go b/device_native.go index e55de2ef..00aad822 100644 --- a/device_native.go +++ b/device_native.go @@ -1159,8 +1159,11 @@ func (d *Device) FreeCommandBuffer(cb *CommandBuffer) { } // PushErrorScope pushes a new error scope onto the device's error scope stack. -func (d *Device) PushErrorScope(filter ErrorFilter) { - d.core.PushErrorScope(filter) +// It returns an error for a filter that is not one of the three WebGPU error +// filters, and pushes nothing then: the spec throws a TypeError for that value +// (pushErrorScope step 1), and a Go API reports it as an error. +func (d *Device) PushErrorScope(filter ErrorFilter) error { + return d.core.PushErrorScope(filter) } // PopErrorScope pops the most recently pushed error scope. diff --git a/device_rust.go b/device_rust.go index 8bb3cc49..5aa3d568 100644 --- a/device_rust.go +++ b/device_rust.go @@ -444,26 +444,36 @@ func (d *Device) WaitForFence(f *Fence, _ uint64, _ time.Duration) (bool, error) } // PushErrorScope pushes a new error scope onto the device's error scope stack. -func (d *Device) PushErrorScope(filter ErrorFilter) { +// It returns an error for a filter that is not one of the three WebGPU error +// filters, and pushes nothing then. The return value is what makes the reason +// reach the caller: wgpu-native panics on an unknown filter and a panic across +// the FFI boundary aborts the process, so the value is rejected here instead. +func (d *Device) PushErrorScope(filter ErrorFilter) error { + rFilter, err := errorFilterToRust(filter) + if err != nil { + return err + } if d.r != nil { - d.r.PushErrorScope(errorFilterToRust(filter)) + d.r.PushErrorScope(rFilter) } + return nil } // errorFilterToRust maps an ErrorFilter (core's numbering, Validation = 0, as // on the native build) to the webgpu.h member of the same name, which counts // from 1. wgpu-native panics on a filter it does not know, and a panic across -// the FFI boundary aborts the process, so a value outside the three constants -// is sent as Validation rather than as an invalid member. -func errorFilterToRust(filter ErrorFilter) rwgpu.ErrorFilter { +// the FFI boundary aborts the process, so an unmapped value is an error here +// and never reaches the binding (P3-API-025). +func errorFilterToRust(filter ErrorFilter) (rwgpu.ErrorFilter, error) { switch filter { + case ErrorFilterValidation: + return rwgpu.ErrorFilterValidation, nil case ErrorFilterOutOfMemory: - return rwgpu.ErrorFilterOutOfMemory + return rwgpu.ErrorFilterOutOfMemory, nil case ErrorFilterInternal: - return rwgpu.ErrorFilterInternal - default: - return rwgpu.ErrorFilterValidation + return rwgpu.ErrorFilterInternal, nil } + return 0, fmt.Errorf("wgpu: PushErrorScope: unknown error filter %d", int(filter)) } // PopErrorScope pops the most recently pushed error scope. diff --git a/enum_values_rust_test.go b/enum_values_rust_test.go index 45a1e860..e4ced146 100644 --- a/enum_values_rust_test.go +++ b/enum_values_rust_test.go @@ -46,10 +46,23 @@ func TestRustErrorFilterToNative(t *testing.T) { {ErrorFilterOutOfMemory, rwgpu.ErrorFilterOutOfMemory}, {ErrorFilterInternal, rwgpu.ErrorFilterInternal}, } { - if got := errorFilterToRust(c.filter); got != c.want { + got, err := errorFilterToRust(c.filter) + if err != nil { + t.Errorf("%v: errorFilterToRust = %v, want wgpu-native filter %d", c.filter, err, c.want) + continue + } + if got != c.want { t.Errorf("%v: wgpu-native filter %d, want %d", c.filter, got, c.want) } } + // An unmapped filter is an error, never a silent Validation fallback: + // wgpu-native panics on an unknown filter and the panic aborts the process + // across the FFI boundary (P3-API-025). + for _, f := range []ErrorFilter{-1, 3, 42} { + if _, err := errorFilterToRust(f); err == nil { + t.Errorf("errorFilterToRust(%d) = nil error, want an error", int(f)) + } + } } // TestRustMapModeValues pins the -tags rust MapMode constants, which have no From 026ddf0c8ac3cda100685d805667106d500d3f7b Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 17:08:05 +0300 Subject: [PATCH 05/16] fix: a released bind group layout cannot back a new bind group or pipeline layout MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Release is Rust's drop (ADR-056), so a layout the application released may have its HAL destruction queued. CreateBindGroup passed desc.Layout.hal to vkAllocateDescriptorSets anyway: on the unfixed code the test crashed with an access violation inside the Vulkan driver (the layout was already destroyed), VUID-vkAllocateDescriptorSets-pAllocateInfo's class. CreatePipelineLayout read the same handles when building the pipeline layout. Both now refuse a released layout with ErrReleased, as CreateBindGroup already refuses a released sampler. No reference count is added: the HAL bind group does not keep the layout — Rust wgpu-hal's Vulkan BindGroup holds only the descriptor set — and the layout is read at creation only, so a group created while the layout was alive is unaffected by a later Release. P2-CORE-023. Test first: hal/allbackends/released_layout_test.go crashed on Vulkan before the fix and passes on vulkan, dx12 (FXC and DXIL), gles and software after it. --- device_native.go | 13 ++++ hal/allbackends/released_layout_test.go | 91 +++++++++++++++++++++++++ 2 files changed, 104 insertions(+) create mode 100644 hal/allbackends/released_layout_test.go diff --git a/device_native.go b/device_native.go index 00aad822..bccbcb3c 100644 --- a/device_native.go +++ b/device_native.go @@ -391,6 +391,11 @@ func (d *Device) CreatePipelineLayout(desc *PipelineLayoutDescriptor) (*Pipeline if layout == nil { return nil, fmt.Errorf("wgpu: bind group layout at index %d is nil", i) } + // A released layout may already be queued for destruction; its HAL + // handle must not back a new pipeline layout (P2-CORE-023). + if layout.released { + return nil, ErrReleased + } halLayouts[i] = layout.hal } @@ -441,6 +446,14 @@ func (d *Device) CreateBindGroup(desc *BindGroupDescriptor) (*BindGroup, error) Label: desc.Label, } } + // A released layout may already be queued for destruction; its HAL handle + // must not back a new descriptor set (P2-CORE-023). The HAL bind group does + // not keep the layout — Rust wgpu-hal's Vulkan BindGroup holds only the + // descriptor set — so a group created while the layout was alive is + // unaffected by a later Release. + if desc.Layout.released { + return nil, ErrReleased + } halEntries := make([]gputypes.BindGroupEntry, len(desc.Entries)) for i, entry := range desc.Entries { diff --git a/hal/allbackends/released_layout_test.go b/hal/allbackends/released_layout_test.go new file mode 100644 index 00000000..b8e4c6d3 --- /dev/null +++ b/hal/allbackends/released_layout_test.go @@ -0,0 +1,91 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build !android && !(js && wasm) && !rust + +package allbackends_test + +import ( + "errors" + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu" +) + +// A bind group layout has no destroy() in WebGPU; Release is Rust's drop +// (ADR-056). An object created from it after that must be refused, as Rust +// refuses anything that needs a value it no longer has: the HAL bind group does +// not keep the layout (Rust wgpu-hal's Vulkan BindGroup holds only the +// descriptor set), so a new bind group or pipeline layout would name a HAL +// layout whose destruction is already queued. P2-CORE-023. + +// TestCreateBindGroupWithReleasedLayout: CreateBindGroup must refuse a layout +// the application released, instead of allocating a descriptor set from a +// destroyed VkDescriptorSetLayout (VUID-vkAllocateDescriptorSets-pAllocateInfo). +func TestCreateBindGroupWithReleasedLayout(t *testing.T) { + forEachOverrideBackend(t, 0, func(t *testing.T, device *wgpu.Device) { + bgl := newBindGroupLayout(t, device, gputypes.BindGroupLayoutEntry{ + Binding: 0, Visibility: wgpu.ShaderStageCompute, + Sampler: &gputypes.SamplerBindingLayout{Type: gputypes.SamplerBindingTypeNonFiltering}, + }) + sampler, err := device.CreateSampler(&wgpu.SamplerDescriptor{Label: "released-layout-sampler"}) + if err != nil { + t.Fatalf("CreateSampler: %v", err) + } + t.Cleanup(sampler.Release) + + bgl.Release() + flushDeferredDestruction(t, device) + + group, err := device.CreateBindGroup(&wgpu.BindGroupDescriptor{ + Label: "released-layout-group", + Layout: bgl, + Entries: []wgpu.BindGroupEntry{{Binding: 0, Sampler: sampler}}, + }) + if !errors.Is(err, wgpu.ErrReleased) { + if group != nil { + group.Release() + } + t.Fatalf("CreateBindGroup with a released layout = %v, want ErrReleased", err) + } + }) +} + +// TestCreatePipelineLayoutWithReleasedLayout: CreatePipelineLayout must refuse +// a bind group layout the application released, as its HAL call reads the +// descriptor set layouts the pipeline layout is built from. +func TestCreatePipelineLayoutWithReleasedLayout(t *testing.T) { + forEachOverrideBackend(t, 0, func(t *testing.T, device *wgpu.Device) { + bgl := newBindGroupLayout(t, device, gputypes.BindGroupLayoutEntry{ + Binding: 0, Visibility: wgpu.ShaderStageCompute, + Sampler: &gputypes.SamplerBindingLayout{Type: gputypes.SamplerBindingTypeNonFiltering}, + }) + bgl.Release() + flushDeferredDestruction(t, device) + + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{ + Label: "released-layout-pipeline", + BindGroupLayouts: []*wgpu.BindGroupLayout{bgl}, + }) + if !errors.Is(err, wgpu.ErrReleased) { + if layout != nil { + layout.Release() + } + t.Fatalf("CreatePipelineLayout with a released layout = %v, want ErrReleased", err) + } + }) +} + +// flushDeferredDestruction submits an empty command buffer and waits for the +// device, so the destroy queue's Triage runs the destruction a Release +// deferred. Without it the HAL layout may still be alive and the check under +// test would not be exercised. +func flushDeferredDestruction(t *testing.T, device *wgpu.Device) { + t.Helper() + enc := newEncoder(t, device) + submit(t, device, enc) + if err := device.WaitIdle(); err != nil { + t.Fatalf("WaitIdle: %v", err) + } +} From 2c2c73f79f02909b99e39a9b1e4ac3a049280018 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 17:13:39 +0300 Subject: [PATCH 06/16] fix(core): validate SetViewport and vertex attribute formats SetViewport recorded whatever it was given: a negative or non-finite size, an origin outside the spec's viewport range, or a depth range outside [0, 1] went straight to the HAL (invalid VkViewport values on Vulkan, undefined behaviour on DX12) and, since the software backend honours MinDepth/MaxDepth, silently changed the rendered depth. Validate the viewport as GPURenderPassEncoder.setViewport requires (spec/index.bs:13695-13708) and as Rust wgpu-core does (command/render.rs:2986): the first failed rule puts the encoder into its deferred error state, which End and Finish report, and nothing is recorded into the HAL pass. Rust lets a NaN origin or size through because its f32 API has no IDL boundary in front of it; the spec rejects it, so the check is explicit here. Vertex attribute formats get the same treatment at pipeline creation: an undefined or unknown gputypes.VertexFormat has size 0, which every existing vertex check treats as a no-op, and the backends then map it to a wrong format (GLES and Vulkan Float32x4, Metal MTLVertexFormatInvalid, DX12 DXGI_FORMAT_UNKNOWN, software zeros). Rust's VertexFormat enum cannot express such a value and Dawn validates it (RenderPipeline.cpp:118); core reports it as an invalid vertex format instead. P2-CORE-019. --- core/command.go | 16 ++ core/error.go | 86 ++++++- core/render_pass_viewport_test.go | 390 ++++++++++++++++++++++++++++++ core/validate.go | 44 ++-- core/validate_test.go | 59 +++++ core/validate_viewport.go | 101 ++++++++ 6 files changed, 680 insertions(+), 16 deletions(-) create mode 100644 core/render_pass_viewport_test.go create mode 100644 core/validate_viewport.go diff --git a/core/command.go b/core/command.go index 01ee6fd8..fb136f1c 100644 --- a/core/command.go +++ b/core/command.go @@ -588,6 +588,7 @@ func (e *CoreCommandEncoder) BeginRenderPass(desc *RenderPassDescriptor) (*CoreR raw: halPass, encoder: e, device: e.device, + label: desc.Label, passContext: passCtx, scope: newPassUsageScopeWithTextures(attachments.scope), } @@ -1403,6 +1404,10 @@ type CoreRenderPassEncoder struct { // device is the parent device. device *Device + // label is the render pass's debug label, from the descriptor. Validation + // errors of pass commands carry it (Rust wgpu-core RenderPass::label). + label string + // pipeline is the currently bound render pipeline. pipeline *RenderPipeline @@ -1540,10 +1545,21 @@ func (p *CoreRenderPassEncoder) SetIndexBuffer(buffer *Buffer, format gputypes.I } // SetViewport sets the viewport. +// +// The viewport is validated as the WebGPU spec requires +// (GPURenderPassEncoder.setViewport valid usage, spec/index.bs:13695-13708). +// An invalid viewport is not recorded: it puts the encoder into its deferred +// error state, which End and Finish report. +// +// Reference: wgpu-core command/render.rs:2986-3026 (set_viewport). func (p *CoreRenderPassEncoder) SetViewport(vp gputypes.Viewport) { if p.ended { return } + if err := validateViewport(vp, p.device.Limits, p.label); err != nil { + p.encoder.SetError(fmt.Errorf("SetViewport: %w", err)) + return + } if p.raw != nil { p.raw.SetViewport(vp) } diff --git a/core/error.go b/core/error.go index 028d8783..1c922eab 100644 --- a/core/error.go +++ b/core/error.go @@ -701,6 +701,15 @@ const ( // constants of the stage named by Stage are invalid, or do not apply to // the overrides of that stage's entry point. Cause holds the reason. CreateRenderPipelineErrorPipelineConstants + // CreateRenderPipelineErrorInvalidVertexFormat indicates a vertex + // attribute format that gputypes does not define (VertexFormatUndefined + // included). + // WebGPU's IDL rejects an unknown GPUVertexFormat before the algorithm + // runs; Dawn validates it as a creation error + // (dawn/src/dawn/native/RenderPipeline.cpp:118, ValidateVertexFormat). + // Rust wgpu-types codes VertexFormat as an enum, so wgpu-core needs no + // equivalent check. + CreateRenderPipelineErrorInvalidVertexFormat ) // CreateRenderPipelineError represents an error during render pipeline creation. @@ -729,6 +738,10 @@ type CreateRenderPipelineError struct { // AttributeFormat is the attribute vertex format for // CreateRenderPipelineErrorVertexAttributeOutOfStride. AttributeFormat string + // AttributeFormatValue is the attribute's vertex format, beside its name. + // An unknown value has no printable name, so + // CreateRenderPipelineErrorInvalidVertexFormat reports the number. + AttributeFormatValue gputypes.VertexFormat // ShaderLocation is the attribute's shaderLocation for vertex attribute // offset, location and clash errors. ShaderLocation uint32 @@ -830,6 +843,9 @@ func (e *CreateRenderPipelineError) vertexAttributeMessage(label string) string label, e.Format) case CreateRenderPipelineErrorPipelineConstants: return fmt.Sprintf("render pipeline %q: %s stage: %v", label, e.Stage, e.Cause) + case CreateRenderPipelineErrorInvalidVertexFormat: + return fmt.Sprintf("render pipeline %q: vertex buffer [%d] attribute at shaderLocation %d: vertex format %d is not a valid GPUVertexFormat", + label, e.BufferIndex, e.ShaderLocation, e.AttributeFormatValue) default: return fmt.Sprintf("render pipeline %q: unknown error", label) } @@ -1321,9 +1337,34 @@ const ( // WebGPU spec: GPURenderPassColorAttachment Valid Usage. // Rust wgpu-core: RenderPassErrorInner::UnsupportedResolveTargetFormat. RenderPassErrorUnsupportedResolveFormat + + // RenderPassErrorInvalidViewportRectSize indicates a SetViewport width or + // height that is negative, not finite, or above + // Limits.MaxTextureDimension2D. + // WebGPU spec: GPURenderPassEncoder.setViewport valid usage. + // Rust wgpu-core: RenderCommandError::InvalidViewportRectSize + // (command/draw.rs:126). + RenderPassErrorInvalidViewportRectSize + + // RenderPassErrorInvalidViewportRectPosition indicates a SetViewport + // origin that is not finite, or an extent that leaves + // [-2 * maxTextureDimension2D, 2 * maxTextureDimension2D - 1]. + // WebGPU spec: GPURenderPassEncoder.setViewport valid usage. + // Rust wgpu-core: RenderCommandError::InvalidViewportRectPosition + // (command/draw.rs:128). + RenderPassErrorInvalidViewportRectPosition + + // RenderPassErrorInvalidViewportDepth indicates a SetViewport whose + // minDepth or maxDepth is outside [0, 1], not finite, or below the other. + // WebGPU spec: GPURenderPassEncoder.setViewport valid usage. + // Rust wgpu-core: RenderCommandError::InvalidViewportDepth + // (command/draw.rs:130). + RenderPassErrorInvalidViewportDepth ) -// RenderPassValidationError represents a validation error during render pass creation. +// RenderPassValidationError represents a validation error of a render pass: +// during its creation (attachments) or while commands are recorded on it +// (viewport). type RenderPassValidationError struct { Kind RenderPassValidationErrorKind Label string @@ -1368,6 +1409,15 @@ type RenderPassValidationError struct { // For InvalidResolveSampleCounts: ActualSamples is the resolve source's // sample count, ResolveSamples the resolve target's. ResolveSamples uint32 + + // For the viewport kinds: the rejected viewport, and the limit it was + // checked against. Limit is Limits.MaxTextureDimension2D for + // InvalidViewportRectSize; ViewportMin and ViewportMax are the allowed + // range of x and y for InvalidViewportRectPosition, zero when the device + // limit was unset and only finiteness was checked. + Viewport gputypes.Viewport + ViewportMin float32 + ViewportMax float32 } // attachmentLocation names the attachment of e the way Rust wgpu-core's @@ -1413,6 +1463,37 @@ func (e *RenderPassValidationError) notRenderableReason() string { } } +// viewportError words the viewport kinds of e the way Rust wgpu-core's +// RenderCommandError does (command/draw.rs:126-130), naming the ranges the +// WebGPU spec requires. +func (e *RenderPassValidationError) viewportError(label string) string { + const spec = "(WebGPU spec: GPURenderPassEncoder.setViewport valid usage)" + vp := e.Viewport + switch e.Kind { + case RenderPassErrorInvalidViewportRectSize: + // The limit is zero when the device's was never populated; then only + // the sign and finiteness of the size were checked. + if e.Limit == 0 { + return fmt.Sprintf("render pass %q: viewport size { w: %v, h: %v } is negative or not finite %s", + label, vp.Width, vp.Height, spec) + } + return fmt.Sprintf("render pass %q: viewport size { w: %v, h: %v } is negative, not finite, or greater than maxTextureDimension2D %d %s", + label, vp.Width, vp.Height, e.Limit, spec) + case RenderPassErrorInvalidViewportRectPosition: + // With an unset device limit only the finiteness of the origin is + // checked, and ViewportMin and ViewportMax stay zero. + if e.ViewportMin == 0 && e.ViewportMax == 0 { + return fmt.Sprintf("render pass %q: viewport rect { x: %v, y: %v, w: %v, h: %v } has an origin that is not finite %s", + label, vp.X, vp.Y, vp.Width, vp.Height, spec) + } + return fmt.Sprintf("render pass %q: viewport rect { x: %v, y: %v, w: %v, h: %v } must keep x and y at or above %v, and x + width and y + height at or below %v %s", + label, vp.X, vp.Y, vp.Width, vp.Height, e.ViewportMin, e.ViewportMax, spec) + default: + return fmt.Sprintf("render pass %q: viewport minDepth %v and/or maxDepth %v are not in [0, 1], or minDepth exceeds maxDepth %s", + label, vp.MinDepth, vp.MaxDepth, spec) + } +} + // formatError returns the message of the format and resolve kinds of e. func (e *RenderPassValidationError) formatError(label string) string { const spec = "(WebGPU spec: GPURenderPassColorAttachment Valid Usage)" @@ -1488,6 +1569,9 @@ func (e *RenderPassValidationError) Error() string { RenderPassErrorInvalidResolveSampleCounts, RenderPassErrorMismatchedResolveFormat, RenderPassErrorUnsupportedResolveFormat: return e.formatError(label) + case RenderPassErrorInvalidViewportRectSize, RenderPassErrorInvalidViewportRectPosition, + RenderPassErrorInvalidViewportDepth: + return e.viewportError(label) default: return fmt.Sprintf("render pass %q: unknown validation error", label) } diff --git a/core/render_pass_viewport_test.go b/core/render_pass_viewport_test.go new file mode 100644 index 00000000..d99761a9 --- /dev/null +++ b/core/render_pass_viewport_test.go @@ -0,0 +1,390 @@ +//go:build !(js && wasm) + +package core + +import ( + "errors" + "math" + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +// The recording HAL types below wrap the package's mocks so a test can see +// what reached the backend: every viewport the render pass was given. + +// recordingViewportPass records the viewports it is given. +type recordingViewportPass struct { + mockRenderPassEncoder + viewports []gputypes.Viewport +} + +var _ hal.RenderPassEncoder = (*recordingViewportPass)(nil) + +func (p *recordingViewportPass) SetViewport(vp gputypes.Viewport) { + p.viewports = append(p.viewports, vp) +} + +// recordingViewportEncoder hands out the shared recording pass. +type recordingViewportEncoder struct { + mockCommandEncoder + pass *recordingViewportPass +} + +var _ hal.CommandEncoder = (*recordingViewportEncoder)(nil) + +func (e *recordingViewportEncoder) BeginRenderPass(_ *hal.RenderPassDescriptor) hal.RenderPassEncoder { + return e.pass +} + +// recordingViewportDevice hands out recording viewport encoders. +type recordingViewportDevice struct { + mockHALDevice + pass *recordingViewportPass +} + +var _ hal.Device = (*recordingViewportDevice)(nil) + +func (d *recordingViewportDevice) CreateCommandEncoder(_ *hal.CommandEncoderDescriptor) (hal.CommandEncoder, error) { + return &recordingViewportEncoder{pass: d.pass}, nil +} + +// newViewportPass returns a command encoder and an open render pass on a +// device with the given limits, plus the HAL pass that records viewports. +func newViewportPass(t *testing.T, limits gputypes.Limits) (*CoreCommandEncoder, *CoreRenderPassEncoder, *recordingViewportPass) { + t.Helper() + rec := &recordingViewportPass{} + device := NewDevice(&recordingViewportDevice{pass: rec}, &Adapter{}, gputypes.Features(0), limits, "TestDevice") + encoder, err := device.CreateCommandEncoder("TestEncoder") + if err != nil { + t.Fatalf("CreateCommandEncoder: %v", err) + } + pass, err := encoder.BeginRenderPass(&RenderPassDescriptor{Label: "TestPass"}) + if err != nil { + t.Fatalf("BeginRenderPass: %v", err) + } + return encoder, pass, rec +} + +// viewportCases builds the valid viewport table. Each case is checked against +// the WebGPU spec's setViewport valid-usage ranges (spec/index.bs:13695-13708) +// and Rust wgpu-core's set_viewport (command/render.rs:2986-3026). +func validViewportCases() []struct { + name string + vp gputypes.Viewport + limits gputypes.Limits +} { + return []struct { + name string + vp gputypes.Viewport + limits gputypes.Limits + }{ + { + name: "common viewport", + vp: gputypes.Viewport{X: 0, Y: 0, Width: 800, Height: 600, MinDepth: 0, MaxDepth: 1}, + limits: gputypes.DefaultLimits(), + }, + { + name: "zero size", + vp: gputypes.Viewport{X: 4, Y: 5, Width: 0, Height: 0, MinDepth: 0, MaxDepth: 1}, + limits: gputypes.DefaultLimits(), + }, + { + name: "negative origin", + vp: gputypes.Viewport{X: -100, Y: -50, Width: 20, Height: 10, MinDepth: 0, MaxDepth: 1}, + limits: gputypes.DefaultLimits(), + }, + { + // maxTextureDimension2D itself is a legal width and height. + name: "size at maxTextureDimension2D", + vp: gputypes.Viewport{X: 0, Y: 0, Width: 8192, Height: 8192, MinDepth: 0, MaxDepth: 1}, + limits: gputypes.DefaultLimits(), + }, + { + // Origin -2 * maxTextureDimension2D is the lowest legal x and y. + name: "origin at the lower range edge", + vp: gputypes.Viewport{X: -16384, Y: -16384, Width: 0, Height: 0, MinDepth: 0, MaxDepth: 1}, + limits: gputypes.DefaultLimits(), + }, + { + // x + width and y + height may reach 2 * maxTextureDimension2D - 1. + name: "extent at the upper range edge", + vp: gputypes.Viewport{X: 8192, Y: 8192, Width: 8191, Height: 8191, MinDepth: 0, MaxDepth: 1}, + limits: gputypes.DefaultLimits(), + }, + { + name: "equal depths", + vp: gputypes.Viewport{Width: 10, Height: 10, MinDepth: 0.5, MaxDepth: 0.5}, + limits: gputypes.DefaultLimits(), + }, + { + // Limits that were never populated skip the range checks (the same + // convention as the vertex stride checks in validateVertexBuffers); + // only the spec's own ranges apply. + name: "unset limits skip the range checks", + vp: gputypes.Viewport{X: 1e9, Y: -1e9, Width: 1e9, Height: 1e9, MinDepth: 0, MaxDepth: 1}, + limits: gputypes.Limits{}, + }, + } +} + +func TestCoreRenderPassEncoder_SetViewport_Valid(t *testing.T) { + for _, tt := range validViewportCases() { + t.Run(tt.name, func(t *testing.T) { + encoder, pass, rec := newViewportPass(t, tt.limits) + + pass.SetViewport(tt.vp) + + if err := encoder.Error(); err != nil { + t.Fatalf("valid viewport recorded a deferred error: %v", err) + } + if got := encoder.Status(); got != CommandEncoderStatusLocked { + t.Fatalf("status = %v, want %v", got, CommandEncoderStatusLocked) + } + if len(rec.viewports) != 1 { + t.Fatalf("HAL viewports = %v, want the one valid viewport", rec.viewports) + } + if rec.viewports[0] != tt.vp { + t.Errorf("HAL viewport = %+v, want %+v", rec.viewports[0], tt.vp) + } + if err := pass.End(); err != nil { + t.Fatalf("End: %v", err) + } + if _, err := encoder.Finish(); err != nil { + t.Fatalf("Finish: %v", err) + } + }) + } +} + +func TestCoreRenderPassEncoder_SetViewport_Validation(t *testing.T) { + nan := float32(math.NaN()) + inf := float32(math.Inf(1)) + limits := gputypes.DefaultLimits() + valid := gputypes.Viewport{X: 0, Y: 0, Width: 100, Height: 100, MinDepth: 0, MaxDepth: 1} + with := func(change func(vp *gputypes.Viewport)) gputypes.Viewport { + vp := valid + change(&vp) + return vp + } + + tests := []struct { + name string + vp gputypes.Viewport + wantKind RenderPassValidationErrorKind + }{ + { + name: "negative width", + vp: with(func(vp *gputypes.Viewport) { vp.Width = -1 }), + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "negative height", + vp: with(func(vp *gputypes.Viewport) { vp.Height = -0.5 }), + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "NaN width", + vp: with(func(vp *gputypes.Viewport) { vp.Width = nan }), + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "infinite height", + vp: with(func(vp *gputypes.Viewport) { vp.Height = inf }), + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "width above maxTextureDimension2D", + vp: with(func(vp *gputypes.Viewport) { vp.Width = float32(limits.MaxTextureDimension2D) + 1 }), + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "height above maxTextureDimension2D", + vp: with(func(vp *gputypes.Viewport) { vp.Height = float32(limits.MaxTextureDimension2D) + 1 }), + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "NaN x", + vp: with(func(vp *gputypes.Viewport) { vp.X = nan }), + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "infinite y", + vp: with(func(vp *gputypes.Viewport) { vp.Y = inf }), + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "x below -2 * maxTextureDimension2D", + vp: with(func(vp *gputypes.Viewport) { vp.X = -2*float32(limits.MaxTextureDimension2D) - 1 }), + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "y below -2 * maxTextureDimension2D", + vp: with(func(vp *gputypes.Viewport) { vp.Y = -2*float32(limits.MaxTextureDimension2D) - 1 }), + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "x + width past 2 * maxTextureDimension2D - 1", + vp: with(func(vp *gputypes.Viewport) { vp.X = 16000; vp.Width = 1000 }), + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "y + height past 2 * maxTextureDimension2D - 1", + vp: with(func(vp *gputypes.Viewport) { vp.Y = 16000; vp.Height = 1000 }), + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "origin past the upper range edge", + vp: with(func(vp *gputypes.Viewport) { vp.X = 2 * float32(limits.MaxTextureDimension2D) }), + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "minDepth below 0", + vp: with(func(vp *gputypes.Viewport) { vp.MinDepth = -0.1 }), + wantKind: RenderPassErrorInvalidViewportDepth, + }, + { + name: "maxDepth above 1", + vp: with(func(vp *gputypes.Viewport) { vp.MaxDepth = 1.5 }), + wantKind: RenderPassErrorInvalidViewportDepth, + }, + { + name: "NaN minDepth", + vp: with(func(vp *gputypes.Viewport) { vp.MinDepth = nan }), + wantKind: RenderPassErrorInvalidViewportDepth, + }, + { + name: "minDepth above maxDepth", + vp: with(func(vp *gputypes.Viewport) { vp.MinDepth = 0.8; vp.MaxDepth = 0.2 }), + wantKind: RenderPassErrorInvalidViewportDepth, + }, + } + + for _, tt := range tests { + t.Run(tt.name, func(t *testing.T) { + encoder, pass, rec := newViewportPass(t, limits) + + pass.SetViewport(tt.vp) + + if len(rec.viewports) != 0 { + t.Errorf("invalid viewport reached the HAL pass: %+v", rec.viewports) + } + if got := encoder.Status(); got != CommandEncoderStatusError { + t.Fatalf("status = %v, want %v", got, CommandEncoderStatusError) + } + var rpve *RenderPassValidationError + if !errors.As(encoder.Error(), &rpve) { + t.Fatalf("deferred error = %T (%v), want *RenderPassValidationError", encoder.Error(), encoder.Error()) + } + if rpve.Kind != tt.wantKind { + t.Errorf("kind = %v, want %v (%v)", rpve.Kind, tt.wantKind, rpve) + } + // The invalid viewport must not produce a command buffer: End and + // Finish report the deferred error. + if err := pass.End(); !errors.As(err, &rpve) { + t.Errorf("End error = %v, want the deferred viewport error", err) + } + if _, err := encoder.Finish(); err == nil { + t.Error("Finish must fail after a deferred viewport error") + } + }) + } +} + +// TestCoreRenderPassEncoder_SetViewport_Validation_UnknownLimit checks the +// checks that do not depend on maxTextureDimension2D still run when the limit +// was never populated. +func TestCoreRenderPassEncoder_SetViewport_Validation_UnknownLimit(t *testing.T) { + tests := []struct { + name string + vp gputypes.Viewport + wantKind RenderPassValidationErrorKind + }{ + { + name: "negative width", + vp: gputypes.Viewport{Width: -10, Height: 10, MaxDepth: 1}, + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "NaN height", + vp: gputypes.Viewport{Width: 10, Height: float32(math.NaN()), MaxDepth: 1}, + wantKind: RenderPassErrorInvalidViewportRectSize, + }, + { + name: "NaN x", + vp: gputypes.Viewport{X: float32(math.NaN()), Width: 10, Height: 10, MaxDepth: 1}, + wantKind: RenderPassErrorInvalidViewportRectPosition, + }, + { + name: "minDepth above maxDepth", + vp: gputypes.Viewport{Width: 10, Height: 10, MinDepth: 0.9, MaxDepth: 0.1}, + wantKind: RenderPassErrorInvalidViewportDepth, + }, + } + + for _, tt := range tests { + t.Run(tt.name, func(t *testing.T) { + encoder, pass, rec := newViewportPass(t, gputypes.Limits{}) + pass.SetViewport(tt.vp) + + if len(rec.viewports) != 0 { + t.Errorf("invalid viewport reached the HAL pass: %+v", rec.viewports) + } + var rpve *RenderPassValidationError + if !errors.As(encoder.Error(), &rpve) { + t.Fatalf("deferred error = %v, want *RenderPassValidationError", encoder.Error()) + } + if rpve.Kind != tt.wantKind { + t.Errorf("kind = %v, want %v (%v)", rpve.Kind, tt.wantKind, rpve) + } + }) + } +} + +// TestRenderPassValidationError_ViewportMessages checks the wording of the +// three viewport kinds and, for the position kind, both forms: with a known +// limit and with the finite-only form of an unpopulated one. +func TestRenderPassValidationError_ViewportMessages(t *testing.T) { + vp := gputypes.Viewport{X: -1, Y: 2, Width: 30, Height: 40, MinDepth: 0.1, MaxDepth: 0.2} + tests := []struct { + name string + err *RenderPassValidationError + contains string + }{ + { + name: "size", + err: &RenderPassValidationError{Kind: RenderPassErrorInvalidViewportRectSize, Label: "p", Viewport: vp, Limit: 8192}, + contains: "maxTextureDimension2D 8192", + }, + { + name: "size without a limit", + err: &RenderPassValidationError{Kind: RenderPassErrorInvalidViewportRectSize, Label: "p", Viewport: vp}, + contains: "negative or not finite", + }, + { + name: "position with a known limit", + err: &RenderPassValidationError{Kind: RenderPassErrorInvalidViewportRectPosition, Label: "p", Viewport: vp, ViewportMin: -16384, ViewportMax: 16383}, + contains: "16383", + }, + { + name: "position without a limit", + err: &RenderPassValidationError{Kind: RenderPassErrorInvalidViewportRectPosition, Label: "p", Viewport: vp}, + contains: "not finite", + }, + { + name: "depth", + err: &RenderPassValidationError{Kind: RenderPassErrorInvalidViewportDepth, Label: "p", Viewport: vp}, + contains: "minDepth", + }, + } + + for _, tt := range tests { + t.Run(tt.name, func(t *testing.T) { + msg := tt.err.Error() + if !containsString(msg, tt.contains) { + t.Errorf("message %q does not contain %q", msg, tt.contains) + } + }) + } +} diff --git a/core/validate.go b/core/validate.go index 5e45098f..8d7bcc97 100644 --- a/core/validate.go +++ b/core/validate.go @@ -639,10 +639,12 @@ const maxVertexAttributeOffset = 0x10000000 // validateVertexAttributes checks the attributes of one vertex buffer in the // order of Rust wgpu-core create_render_pipeline -// (wgpu-core/src/device/resource.rs): per attribute the offset alignment -// (RP10f) and the location limit (RP10g), then per attribute the offset bound -// (RP10f) and location clashes with every attribute seen so far (RP10h). -// locations collects the shader locations across the pipeline's buffers. +// (wgpu-core/src/device/resource.rs): first whether the format is one gputypes +// defines (a check Rust needs neither, since its VertexFormat is an enum), +// then per attribute the offset alignment (RP10f) and the location limit +// (RP10g), then per attribute the offset bound (RP10f) and location clashes +// with every attribute seen so far (RP10h). locations collects the shader +// locations across the pipeline's buffers. // // Rust's later total-count check against maxVertexAttributes cannot fail // after these: unique locations below maxVertexAttributes bound the total. @@ -657,20 +659,32 @@ func validateVertexAttributes( ) error { attrError := func(kind CreateRenderPipelineErrorKind, j uint32, attr gputypes.VertexAttribute) error { return &CreateRenderPipelineError{ - Kind: kind, - Label: label, - BufferIndex: bufferIndex, - ArrayStride: vb.ArrayStride, - AttributeIndex: j, - AttributeOffset: attr.Offset, - AttributeFormat: attr.Format.String(), - ShaderLocation: attr.ShaderLocation, - Limit: limits.MaxVertexAttributes, + Kind: kind, + Label: label, + BufferIndex: bufferIndex, + ArrayStride: vb.ArrayStride, + AttributeIndex: j, + AttributeOffset: attr.Offset, + AttributeFormat: attr.Format.String(), + AttributeFormatValue: attr.Format, + ShaderLocation: attr.ShaderLocation, + Limit: limits.MaxVertexAttributes, } } for j, attr := range vb.Attributes { - // An undefined format has size 0; it then needs no alignment here. - if alignment := min(4, attr.Format.Size()); alignment > 0 && attr.Offset%alignment != 0 { + // The format must be one gputypes defines. This is not a spec + // validation rule: WebGPU's IDL rejects an unknown GPUVertexFormat + // before the algorithm runs. A Go caller can build one, and an + // unknown value has size 0, so every other check here treats it as a + // no-op and a backend receives a format it cannot map + // (hal/gles/command.go vertexFormatToGL falls back to Float32x4, + // reading 16 bytes per vertex). + if attr.Format.Size() == 0 { + return attrError(CreateRenderPipelineErrorInvalidVertexFormat, uint32(j), attr) + } + // Every defined format is at least 2 bytes, so the alignment is + // non-zero here. + if alignment := min(4, attr.Format.Size()); attr.Offset%alignment != 0 { return attrError(CreateRenderPipelineErrorVertexAttributeInvalidOffset, uint32(j), attr) } if attr.ShaderLocation >= limits.MaxVertexAttributes { diff --git a/core/validate_test.go b/core/validate_test.go index 1a34167b..bc467803 100644 --- a/core/validate_test.go +++ b/core/validate_test.go @@ -3370,6 +3370,65 @@ func TestValidateRenderPipelineDescriptor_VertexBuffers(t *testing.T) { } } +// TestValidateRenderPipelineDescriptor_VertexFormat checks that a vertex +// attribute format must be one gputypes defines. WebGPU's IDL rejects an +// unknown GPUVertexFormat before the algorithm runs (a TypeError, not a +// validation error); a Go caller can build one, and every other vertex check +// treats it as a no-op because its size is 0, so the pipeline would reach the +// backend with a format no backend maps (hal/gles/command.go vertexFormatToGL +// falls back to Float32x4). Dawn validates it as a creation error +// (dawn/src/dawn/native/RenderPipeline.cpp:118 ValidateVertexFormat). +func TestValidateRenderPipelineDescriptor_VertexFormat(t *testing.T) { + layout := func(format gputypes.VertexFormat) gputypes.VertexBufferLayout { + return gputypes.VertexBufferLayout{ + ArrayStride: 16, + StepMode: gputypes.VertexStepModeVertex, + Attributes: []gputypes.VertexAttribute{{ + Format: format, + ShaderLocation: 1, + }}, + } + } + + tests := []struct { + name string + format gputypes.VertexFormat + wantErr bool + }{ + {name: "first defined format", format: gputypes.VertexFormatUint8x2}, + {name: "last defined format", format: gputypes.VertexFormatUnorm1010102}, + {name: "undefined", format: gputypes.VertexFormatUndefined, wantErr: true}, + {name: "first value past the defined set", format: gputypes.VertexFormat(0x20), wantErr: true}, + {name: "arbitrary unknown value", format: gputypes.VertexFormat(0xDEADBEEF), wantErr: true}, + } + + for _, tt := range tests { + t.Run(tt.name, func(t *testing.T) { + desc := vertexStrideDesc(layout(tt.format)) + err := ValidateRenderPipelineDescriptor(desc, gputypes.DefaultLimits(), gputypes.Features(0)) + if !tt.wantErr { + if err != nil { + t.Fatalf("expected nil error, got: %v", err) + } + return + } + var crpe *CreateRenderPipelineError + if !errors.As(err, &crpe) { + t.Fatalf("expected CreateRenderPipelineError, got %T (%v)", err, err) + } + if crpe.Kind != CreateRenderPipelineErrorInvalidVertexFormat { + t.Fatalf("expected kind InvalidVertexFormat, got %v (%v)", crpe.Kind, err) + } + if crpe.AttributeIndex != 0 || crpe.ShaderLocation != 1 { + t.Errorf("AttributeIndex = %d, ShaderLocation = %d, want 0 and 1", crpe.AttributeIndex, crpe.ShaderLocation) + } + if !containsString(crpe.Error(), "not a valid GPUVertexFormat") { + t.Errorf("message %q does not say the format is invalid", crpe.Error()) + } + }) + } +} + func TestCreateRenderPipelineError_VertexBufferMessages(t *testing.T) { tests := []struct { name string diff --git a/core/validate_viewport.go b/core/validate_viewport.go new file mode 100644 index 00000000..34cba907 --- /dev/null +++ b/core/validate_viewport.go @@ -0,0 +1,101 @@ +//go:build !(js && wasm) + +package core + +import ( + "math" + + "github.com/gogpu/gputypes" +) + +// validateViewport validates a render pass viewport, returning a +// *RenderPassValidationError for the first failed rule. +// +// WebGPU spec: GPURenderPassEncoder.setViewport valid usage +// (spec/index.bs:13695-13708). The IDL takes x, y, width and height as +// restricted floats, so a NaN or infinite argument is a TypeError before the +// algorithm runs; a Go float32 can hold such a value, so the checks below +// reject it explicitly. +// +// Reference: wgpu-core command/render.rs:2986-3026 (set_viewport). Rust checks +// the same ranges, but its f32 API has no IDL boundary in front of it and the +// NaN comparisons all fail open, so a NaN origin or size passes Rust. Here the +// spec is normative, as for every other check in this package. +// +// The rules that depend on maxTextureDimension2D are skipped when the limit was +// never populated (0), the convention the vertex stride checks use too. +func validateViewport(vp gputypes.Viewport, limits gputypes.Limits, label string) error { + switch { + case !viewportSizeValid(vp, limits): + return &RenderPassValidationError{ + Kind: RenderPassErrorInvalidViewportRectSize, + Label: label, + Viewport: vp, + Limit: limits.MaxTextureDimension2D, + } + case !viewportPositionValid(vp, limits): + err := &RenderPassValidationError{ + Kind: RenderPassErrorInvalidViewportRectPosition, + Label: label, + Viewport: vp, + } + // Zero marks the "limit unknown" form of the message. + if maxDim := float32(limits.MaxTextureDimension2D); maxDim > 0 { + err.ViewportMin = -2 * maxDim + err.ViewportMax = 2*maxDim - 1 + } + return err + case !viewportDepthValid(vp): + return &RenderPassValidationError{ + Kind: RenderPassErrorInvalidViewportDepth, + Label: label, + Viewport: vp, + } + } + return nil +} + +// viewportSizeValid reports whether width and height are finite and within +// [0, maxTextureDimension2D]. +func viewportSizeValid(vp gputypes.Viewport, limits gputypes.Limits) bool { + if !isFinite32(vp.Width) || !isFinite32(vp.Height) || vp.Width < 0 || vp.Height < 0 { + return false + } + maxDim := float32(limits.MaxTextureDimension2D) + return maxDim <= 0 || (vp.Width <= maxDim && vp.Height <= maxDim) +} + +// viewportPositionValid reports whether x and y are finite and, when the device +// limit is known, whether the origin is at or above -2 * +// maxTextureDimension2D and the extent ends at or below +// 2 * maxTextureDimension2D - 1. +func viewportPositionValid(vp gputypes.Viewport, limits gputypes.Limits) bool { + if !isFinite32(vp.X) || !isFinite32(vp.Y) { + return false + } + maxDim := float32(limits.MaxTextureDimension2D) + if maxDim <= 0 { + return true + } + // The spec's terms are named as it words them: the origin may not go + // below -maxViewportRange, the extent may not end past + // maxViewportRange - 1. Both are compared as floats, so the extent bound + // keeps the spec's <= and is not folded into a < against the range. + lowestOrigin := -2 * maxDim + maxExtent := 2*maxDim - 1 + return vp.X >= lowestOrigin && vp.Y >= lowestOrigin && + vp.X+vp.Width <= maxExtent && vp.Y+vp.Height <= maxExtent +} + +// viewportDepthValid reports whether 0 <= minDepth <= maxDepth <= 1. A NaN +// fails every comparison, so it needs no separate finite test. +func viewportDepthValid(vp gputypes.Viewport) bool { + return vp.MinDepth >= 0 && vp.MinDepth <= 1 && + vp.MaxDepth >= 0 && vp.MaxDepth <= 1 && + vp.MinDepth <= vp.MaxDepth +} + +// isFinite32 reports whether f is neither NaN nor an infinity. +func isFinite32(f float32) bool { + return !math.IsNaN(float64(f)) && !math.IsInf(float64(f), 0) +} From 6d1a79e298e053d5c12165b99863cf0f8ae8c3ee Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 16:39:28 +0300 Subject: [PATCH 07/16] ci: lint windows and darwin targets, pin every action (P2-INFRA-007) CI linted one linux target, so code under darwin and windows build tags (hal/metal, hal/dx12, hal/gles/wgl) was never linted in CI: a PR could carry darwin or windows findings into main between releases (gogpu#519 hit this). - lint-targets: golangci-lint for windows/amd64 and darwin/arm64 on a linux runner, GOOS/GOARCH in the step env, fail-fast false. The module has no arch-specific files and no GOOS+GOARCH build tags, so one GOARCH per GOOS covers every file; js/wasm stays a build check in cross-compile. - golangci-lint 2.15.0 comes from gogpu/.github/actions/setup-golangci-lint (sha256-checked, cached, pinned by SHA) and the action runs with install-mode none: no more 'latest' drift, no per-run release download. - every action is pinned by commit SHA with a version comment; dependabot tracks github-actions and the /tools module, the root module stays on the ecosystem cascade schedule. - go-mod-outdated is pinned in tools/go.mod, replacing @latest and the retry loop. - test: the plain run is CGO_ENABLED=0 (goffi fakecgo, the path users ship); the race run stays CGO_ENABLED=1 with the toolchain detector. - lint-targets joins the Test gate. Job names stay Lint and Test: the ruleset matches that exact name and the switch is a separate decision. --- .github/dependabot.yml | 23 +++++++ .github/workflows/ci.yml | 140 +++++++++++++++++++++++++-------------- tools/go.mod | 12 ++++ tools/go.sum | 10 +++ 4 files changed, 137 insertions(+), 48 deletions(-) create mode 100644 .github/dependabot.yml create mode 100644 tools/go.mod create mode 100644 tools/go.sum diff --git a/.github/dependabot.yml b/.github/dependabot.yml new file mode 100644 index 00000000..c78d8e00 --- /dev/null +++ b/.github/dependabot.yml @@ -0,0 +1,23 @@ +version: 2 +updates: + # Every action in our workflows is pinned by commit SHA with a version + # comment; Dependabot moves the SHA and the comment together. + - package-ecosystem: github-actions + directory: / + schedule: + interval: weekly + groups: + actions: + patterns: ["*"] + + # CI tools (go-mod-outdated), pinned in tools/go.mod. The root module is left + # out on purpose: ecosystem dependencies move only to released versions, in + # the order the cascade tiers set, and the deps job already reports what is + # behind. + - package-ecosystem: gomod + directory: /tools + schedule: + interval: weekly + groups: + tools: + patterns: ["*"] diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 25f185ea..a21c7022 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -21,7 +21,6 @@ on: permissions: contents: read - id-token: write env: # goffi requires CGO_ENABLED=0 for pure Go FFI on Unix @@ -39,10 +38,10 @@ jobs: steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Checkout exact merged WebGPU Android surface source - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 with: repository: go-webgpu/webgpu ref: a801aed7399042e5564ef76fc9f075da5cb70081 @@ -54,7 +53,7 @@ jobs: run: test "$(git rev-parse HEAD)" = a801aed7399042e5564ef76fc9f075da5cb70081 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: ${{ matrix.go-version }} cache: true @@ -85,10 +84,10 @@ jobs: steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: ${{ matrix.go-version }} cache: true @@ -115,10 +114,10 @@ jobs: - { goos: linux, goarch: amd64, tags: 'rust', name: 'rust' } steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: '1.25' cache: true @@ -134,6 +133,10 @@ jobs: test: name: Test - ${{ matrix.os }} runs-on: ${{ matrix.os }} + # Codecov uploads with OIDC; no other job needs an id token. + permissions: + contents: read + id-token: write strategy: # Each OS runs different backends (DX12, Metal, Vulkan/GLES), so one # OS failing must not cancel the others' results. @@ -144,10 +147,10 @@ jobs: steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: ${{ matrix.go-version }} cache: true @@ -161,14 +164,20 @@ jobs: run: | go vet $(go list ./... | grep -v '/hal/vulkan/vk$' | grep -v '/hal/gles') + # The build users ship: CGO_ENABLED=0, so goffi runs its fakecgo path. + - name: Run tests + shell: bash + run: go test -v -count=1 -coverprofile=coverage.txt -covermode=atomic ./... + + # Race detection uses the toolchain's -race (ThreadSanitizer, needs cgo). + # With CGO_ENABLED=1 goffi links the real runtime/cgo instead of its + # fakecgo, so this step tests that path, not the one above. A non-std cgo + # package in the test binary would fail to link here on Linux (goffi#72). - name: Run tests with the race detector shell: bash - # The standard race detector needs cgo at test time only; the library - # itself stays cgo-free (goffi supports CGO_ENABLED=1 too). A data race - # fails the job. env: CGO_ENABLED: 1 - run: go test -race -v -coverprofile=coverage.txt -covermode=atomic ./... + run: go test -race -count=1 ./... - name: Verify coverage file if: matrix.os == 'ubuntu-latest' @@ -183,7 +192,7 @@ jobs: - name: Upload coverage to Codecov if: matrix.os == 'ubuntu-latest' - uses: codecov/codecov-action@v5 + uses: codecov/codecov-action@0fb7174895f61a3b6b78fc075e0cd60383518dac # v5.5.5 with: use_oidc: true files: ./coverage.txt @@ -197,10 +206,10 @@ jobs: steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: '1.25' cache: true @@ -274,18 +283,70 @@ jobs: steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: '1.25' cache: true + # The golangci-lint version and its sha256 are pinned in the composite + # action, which restores the verified binary from the Actions cache, so + # lint does not depend on GitHub Releases being reachable on every run + # (on 2026-11-08 the release download URL returned 504 for over an hour). + # Bump golangci-lint in gogpu/.github, then move this pin; keep + # lint-targets on the same commit. + - name: Set up golangci-lint + uses: gogpu/.github/actions/setup-golangci-lint@cdb7bd4a5bd3852541a05f0fa1fa92b09d1ee732 # main 2026-11-08, golangci-lint v2.15.0 + - name: Run golangci-lint - uses: golangci/golangci-lint-action@v8 + uses: golangci/golangci-lint-action@ba0d7d2ec06a0ea1cb5fa41b2e4a3ab91d21278a # v9.3.0 with: - version: latest + install-mode: none + args: --timeout=5m + + # Lint the other desktop targets. The job above type-checks only the linux + # files, so code under darwin or windows build tags (hal/metal, hal/dx12, + # hal/gles/wgl) is never linted otherwise. The code is pure Go + # (CGO_ENABLED=0), so GOOS/GOARCH on a linux runner load exactly that + # target's files; a target that fails to type-check exits non-zero, never + # "0 issues". One GOARCH per GOOS is enough: the module has no arch-specific + # files and no darwin/windows + GOARCH build tags (checked 2026-11-08), so + # only darwin/arm64 and windows/amd64 are added. GOOS-named tags in the + # .golangci.yml build-tags would not switch the target and are not used. + lint-targets: + name: Lint - ${{ matrix.goos }}/${{ matrix.goarch }} + runs-on: ubuntu-latest + strategy: + fail-fast: false + matrix: + include: + - { goos: windows, goarch: amd64 } + - { goos: darwin, goarch: arm64 } + + steps: + - name: Checkout code + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 + + - name: Set up Go + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 + with: + go-version: '1.25' + cache: true + + - name: Set up golangci-lint + uses: gogpu/.github/actions/setup-golangci-lint@cdb7bd4a5bd3852541a05f0fa1fa92b09d1ee732 # main 2026-11-08, golangci-lint v2.15.0 + + # GOOS/GOARCH go on the lint step only: the binary above stays linux/amd64. + - name: Run golangci-lint + uses: golangci/golangci-lint-action@ba0d7d2ec06a0ea1cb5fa41b2e4a3ab91d21278a # v9.3.0 + env: + GOOS: ${{ matrix.goos }} + GOARCH: ${{ matrix.goarch }} + CGO_ENABLED: "0" + with: + install-mode: none args: --timeout=5m # Code formatting check @@ -295,10 +356,10 @@ jobs: steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: '1.25' cache: true @@ -323,38 +384,20 @@ jobs: steps: - name: Checkout code - uses: actions/checkout@v4 + uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 - name: Set up Go - uses: actions/setup-go@v5 + uses: actions/setup-go@40f1582b2485089dde7abd97c1529aa768e1baff # v5.6.0 with: go-version: '1.25' cache: true - # Pin the module version so each run does not re-resolve @latest. - # Retry: proxy.golang.org occasionally returns INTERNAL_ERROR / stream - # resets (seen as red Dependencies checks on otherwise-green PRs). - # continue-on-error: the direct-deps step is advisory only; ecosystem - # checks below do not need this binary. - - name: Install go-mod-outdated - id: install-outdated - continue-on-error: true - run: | - for attempt in 1 2 3; do - if go install github.com/psampaz/go-mod-outdated@v0.9.0; then - exit 0 - fi - echo "install attempt ${attempt} failed; retrying..." - sleep $((attempt * 5)) - done - echo "::warning::Failed to install go-mod-outdated after retries; skipping direct-deps check" - exit 1 - + # go-mod-outdated is pinned in tools/go.mod (checksums in tools/go.sum, + # updated by Dependabot), never installed @latest. - name: Check direct dependencies for updates - if: steps.install-outdated.outcome == 'success' run: | echo "## Direct Dependencies with Updates" - OUTDATED=$(go list -u -m -json all 2>/dev/null | go-mod-outdated -update -direct || true) + OUTDATED=$(go list -u -m -json all 2>/dev/null | go tool -modfile=tools/go.mod go-mod-outdated -update -direct || true) if [ -n "$OUTDATED" ]; then echo "$OUTDATED" echo "::warning::Some direct dependencies have updates available" @@ -391,14 +434,15 @@ jobs: [ $WARNINGS -gt 0 ] && echo "::warning::$WARNINGS ecosystem dep(s) outdated" exit 0 - # The one check the main-protection ruleset requires. Jobs and the matrix can + # The one check the main-protection ruleset requires, by this exact job name: + # renaming the job silently drops the requirement. Jobs and the matrix can # change without touching the ruleset; a new required job goes into `needs`. # `if: always()` and the explicit result check matter: a job skipped because a # job it needs failed counts as passed for a required check. gate: name: Test if: always() - needs: [android-arm64-preview, build, cross-compile, test, integration-gles, lint, formatting, deps] + needs: [android-arm64-preview, build, cross-compile, test, integration-gles, lint, lint-targets, formatting, deps] runs-on: ubuntu-latest steps: - name: Every required job succeeded diff --git a/tools/go.mod b/tools/go.mod new file mode 100644 index 00000000..cd573374 --- /dev/null +++ b/tools/go.mod @@ -0,0 +1,12 @@ +module github.com/gogpu/wgpu/tools + +go 1.25.0 + +tool github.com/psampaz/go-mod-outdated + +require ( + github.com/mattn/go-runewidth v0.0.14 // indirect + github.com/olekukonko/tablewriter v0.0.5 // indirect + github.com/psampaz/go-mod-outdated v0.9.0 // indirect + github.com/rivo/uniseg v0.4.3 // indirect +) diff --git a/tools/go.sum b/tools/go.sum new file mode 100644 index 00000000..4265e619 --- /dev/null +++ b/tools/go.sum @@ -0,0 +1,10 @@ +github.com/mattn/go-runewidth v0.0.9/go.mod h1:H031xJmbD/WCDINGzjvQ9THkh0rPKHF+m2gUSrubnMI= +github.com/mattn/go-runewidth v0.0.14 h1:+xnbZSEeDbOIg5/mE6JF0w6n9duR1l3/WmbinWVwUuU= +github.com/mattn/go-runewidth v0.0.14/go.mod h1:Jdepj2loyihRzMpdS35Xk/zdY8IAYHsh153qUoGf23w= +github.com/olekukonko/tablewriter v0.0.5 h1:P2Ga83D34wi1o9J6Wh1mRuqd4mF/x/lgBS7N7AbDhec= +github.com/olekukonko/tablewriter v0.0.5/go.mod h1:hPp6KlRPjbx+hW8ykQs1w3UBbZlj6HuIJcUGPhkA7kY= +github.com/psampaz/go-mod-outdated v0.9.0 h1:P3f6z6NrAgG1kq1W4xcsa/lL8SM2SEZxAlRvG1AMFBs= +github.com/psampaz/go-mod-outdated v0.9.0/go.mod h1:FcfE/igcl0GuLxemNXSL7r+rconnPmFP8kmO/Th3/To= +github.com/rivo/uniseg v0.2.0/go.mod h1:J6wj4VEh+S6ZtnVlnTBMWIodfgj8LQOQFoIToxlJtxc= +github.com/rivo/uniseg v0.4.3 h1:utMvzDsuh3suAEnhH0RdHmoPbU648o6CvXxTx4SBMOw= +github.com/rivo/uniseg v0.4.3/go.mod h1:FN3SvrM+Zdj16jyLfmOkMNblXMcoc8DfTHruCPUcx88= From 6b28f8b87fe30c1495ca75ac412b5d4e73c5c0e4 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 16:57:40 +0300 Subject: [PATCH 08/16] ci: build windows/386 in the cross-compile job MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 32-bit int is a build class CI did not cover: four constants that fit a 64-bit int broke the 386 build silently until gogpu reported it (P1-INFRA-008). The leg builds only — no GPU and no tests there. windows and not linux: goffi's internal/syscall has no linux/386 files, so the linux/386 combination cannot build in the ecosystem today; the 32-bit int checks apply either way. needs the P1-INFRA-008 fix (1183df5) to pass. --- .github/workflows/ci.yml | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index a21c7022..64f474dd 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -112,6 +112,12 @@ jobs: include: - { goos: js, goarch: wasm, tags: '', name: 'browser' } - { goos: linux, goarch: amd64, tags: 'rust', name: 'rust' } + # 32-bit int: a constant that fits int on 64-bit targets does not + # there (P1-INFRA-008: TrackerIndex bounds and math.MaxUint32). Build + # only — there is no GPU on 386 and no test runs. windows rather than + # linux: goffi's internal/syscall has no linux/386 files, so the + # linux/386 combination cannot build in the ecosystem. + - { goos: windows, goarch: "386", tags: '', name: '386' } steps: - name: Checkout code uses: actions/checkout@11d5960a326750d5838078e36cf38b85af677262 # v4.4.0 From bc2e22be7e163d2a86f9ed74200c2258f93e6a58 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 17:58:29 +0300 Subject: [PATCH 09/16] test(vulkan): type-check the vulkan tests for 32-bit targets Two test files used values that do not fit 32-bit: a t.Errorf argument of 4294967296 (4 << 30, defaults to int) in memory/types_test.go, and a 64-bit ANativeWindow handle built with a constant shift in android_surface_policy_test.go. GOARCH=386 go vet cannot compile either, so the 386 leg added for P1-INFRA-008 would fail on the tests rather than on the code. The heap size is printed as uint64; the handle's shift count is a variable, so the constant no longer overflows when the package is type-checked for a 32-bit target, where the test skips as it already did for the arm64 ABI case. --- CHANGELOG.md | 26 +++++++++++++++++++++++ hal/vulkan/android_surface_policy_test.go | 9 +++++++- hal/vulkan/memory/types_test.go | 2 +- 3 files changed, 35 insertions(+), 2 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 8677d6d1..d3b13b38 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -7,6 +7,32 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0 ## [Unreleased] +## [0.34.11] - 2026-11-08 + +A lifetime and validation patch: resources released after recording stay alive through the work that uses them, +pass input is validated as the spec requires, and the module builds on 32-bit. + +### Fixed + +- **A query set released after recording stays alive until that work completes** (P1-CORE-021, closes the v0.34.10 known issue). `QuerySet.Release` destroyed the HAL query pool at the previous submission index, so a set released after timestamp writes or a `ResolveQuerySet` were recorded went away under the submitted command buffer: `VUID-vkDestroyQueryPool-queryPool-00793` on Vulkan, zero timestamps on DX12. Query sets are reference counted as samplers are; the command encoder holds the set's reference until the submission completes, and `Submit` refuses only a set whose HAL resource is gone (`ErrSubmitQuerySetDestroyed`). +- **A render or compute pipeline keeps its pipeline layout alive** (P2-CORE-022, closes the second v0.34.10 known issue). A `PipelineLayout` released after the pipeline was created was destroyed while later `SetBindGroup` calls still named it — an invalid `VkPipelineLayout` at `vkCmdBindDescriptorSets`. Pipelines now hold their layout's reference and drop it when their own resource is destroyed. +- **A released bind group layout cannot back a new bind group or pipeline layout** (P2-CORE-023). `CreateBindGroup` and `CreatePipelineLayout` now return `ErrReleased`, as `CreateBindGroup` already did for a released sampler. Before, the released layout's HAL handle reached `vkAllocateDescriptorSets` and crashed inside the driver. +- **`SetViewport` is validated as the spec requires** (P2-CORE-019). Non-finite or negative sizes, an origin outside the viewport range, or a depth range outside `[0, 1]` is a `RenderPassValidationError` on the encoder instead of invalid values reaching the driver. Where the spec and Rust wgpu differ, the spec decides. +- **A vertex format gputypes does not define is a pipeline-creation error** (P2-CORE-019, `CreateRenderPipelineErrorInvalidVertexFormat`). Each backend used to substitute a format of its own (float32x4 on Vulkan and GLES, an invalid Metal format, an unknown DXGI format). +- **The module builds on GOARCH=386** (P1-INFRA-008). Four comparisons against `TrackerIndex` and `math.MaxUint32` were written as `int` constants and overflowed a 32-bit `int`, so every 386 build of wgpu failed to compile. Comparison is done in `uint64` now; behaviour on 64-bit targets is unchanged. Build only — there is no GPU on 386, and ecosystem CI needs goffi's own 386 support before the other repositories can run the leg. + +### Changed + +- **`Device.PushErrorScope` returns an error for a filter outside the three WebGPU filters** (P3-API-025). It used to map an unknown value to `Validation` silently, pushing a scope that could never capture what the caller asked for. The native build validates in core, `-tags rust` rejects the value before it reaches wgpu-native, which panics on an unknown filter and whose panic would abort the process across the FFI boundary. Calls written as statements are unaffected. + +### Infrastructure + +- CI lints windows/amd64 and darwin/arm64 in addition to the host, pins every action by SHA, installs the pinned golangci-lint from the ecosystem action, adds a dependabot configuration, and builds windows/386 (P2-INFRA-007, P1-INFRA-008). The ruleset still requires the `Test` gate; switching it to `Required checks` is a separate, coordinated step. + +### Known issues + +- Scissor rects are not validated yet; they go to the HAL unvalidated, as viewports did before this release (P2-CORE-024, next patch). + ## [0.34.10] - 2026-11-08 A patch for frame-batching consumers and a security update of naga. diff --git a/hal/vulkan/android_surface_policy_test.go b/hal/vulkan/android_surface_policy_test.go index fdee5376..cb476d8d 100644 --- a/hal/vulkan/android_surface_policy_test.go +++ b/hal/vulkan/android_surface_policy_test.go @@ -32,7 +32,14 @@ func TestAndroidSurfaceRequestIsWindowOnlyAndStateless(t *testing.T) { } func TestAndroidSurfaceCreateInfoPreservesIndependentRawWindows(t *testing.T) { - firstWindow := uintptr(0x87654321) | uintptr(0x10)<<32 + if unsafe.Sizeof(uintptr(0)) != 8 { + t.Skip("a 64-bit ANativeWindow handle does not fit a 32-bit uintptr") + } + // The shift count is a variable so the constant does not overflow uintptr + // when the package is type-checked for a 32-bit target, where the test + // skips above. + highBits := uint(32) + firstWindow := uintptr(0x87654321) | uintptr(0x10)< Date: Sun, 8 Nov 2026 17:48:44 +0300 Subject: [PATCH 10/16] fix(gles): MSAA, buffer mapping and storage textures on OpenGL ES P1-GLES-025. Three defects found by running the backend on an OpenGL ES 3.2 context (Mesa llvmpipe, the wgpu_gles_force_es seam): - A multisampled render target was created with glTexImage2DMultisample, which OpenGL ES does not have, so every MSAA texture failed with GL_INVALID_OPERATION 0x502 and a pass that resolved MSAA could not be created at all. It is a renderbuffer now, attached to the framebuffer with glFramebufferRenderbuffer, as Rust wgpu-hal creates render targets (gles/device.rs create_texture). - Device.MapBuffer mapped with glMapBuffer, which OpenGL ES does not have either (Mesa exposes no glMapBufferOES on ES 3.x): the mapping failed silently and the CPU copy stayed zero-filled. Every buffer readback of GPU-written data on ES therefore read zeros, which is the whole class of "compute and storage readbacks are all zeros" failures. Mapping goes through glMapBufferRange, as in Rust wgpu-hal (gles/device.rs map_buffer). - Textures were allocated with mutable TexImage storage. OpenGL ES rejects an image binding (glBindImageTexture) of a texture with mutable storage with GL_INVALID_OPERATION, so imageStore and imageLoad wrote nowhere. 2D, cube and layered textures use immutable storage (glTexStorage2D/3D) when the context has it, which is Rust wgpu-hal's TEXTURE_STORAGE path (gles/adapter.rs, device.rs create_texture). Tests, all forced-ES (integration && linux && wgpu_gles_force_es): an MSAA clear-and-resolve readback, three compute readbacks over a constant, a storage buffer and a uniform block, and a storage-texture write. Each fails on the old code (GL errors and zero readbacks) and passes now; the public API tests TestMSAAStoreWithResolveIsLoadable, TestSampledRenderTarget, TestBindGroupRebindCompute, TestWriteBufferThenDispatchReadsIt, TestStorageTextureWriteThenSample and TestStorageWriteArrayLayer go green on that context as well. --- hal/gles/adapter.go | 2 + hal/gles/adapter_linux.go | 2 + hal/gles/capabilities.go | 21 +++ hal/gles/command.go | 6 +- hal/gles/compute_readback_es_linux_test.go | 177 +++++++++++++++++++++ hal/gles/device.go | 79 +++++++-- hal/gles/device_linux.go | 109 +++++++++++-- hal/gles/gl/const.go | 11 ++ hal/gles/gl/context.go | 60 +++++-- hal/gles/gl/context_linux.go | 95 +++++++++-- hal/gles/layered_texture.go | 8 +- hal/gles/msaa_es_linux_test.go | 164 +++++++++++++++++++ hal/gles/resource.go | 40 ++++- hal/gles/storage_texture_es_linux_test.go | 127 +++++++++++++++ 14 files changed, 840 insertions(+), 61 deletions(-) create mode 100644 hal/gles/compute_readback_es_linux_test.go create mode 100644 hal/gles/msaa_es_linux_test.go create mode 100644 hal/gles/storage_texture_es_linux_test.go diff --git a/hal/gles/adapter.go b/hal/gles/adapter.go index 9a5cfe85..0aba05df 100644 --- a/hal/gles/adapter.go +++ b/hal/gles/adapter.go @@ -74,6 +74,8 @@ func (a *Adapter) Open(_ gputypes.Features, _ gputypes.Limits) (hal.OpenDevice, depthStencilTextureMode: glVersionAtLeast(a.caps.GLMajor, a.caps.GLMinor, a.caps.IsES, [2]int{3, 1}, [2]int{4, 3}), maxColorAttachments: a.caps.Limits.MaxColorAttachments, + isES: a.caps.IsES, + immutableStorage: a.caps.TextureStorage, } queue := &Queue{ diff --git a/hal/gles/adapter_linux.go b/hal/gles/adapter_linux.go index f905f605..56816d74 100644 --- a/hal/gles/adapter_linux.go +++ b/hal/gles/adapter_linux.go @@ -100,6 +100,8 @@ func (a *Adapter) open() (hal.OpenDevice, error) { depthStencilTextureMode: glVersionAtLeast(a.caps.GLMajor, a.caps.GLMinor, a.caps.IsES, [2]int{3, 1}, [2]int{4, 3}), maxColorAttachments: a.caps.Limits.MaxColorAttachments, + isES: a.caps.IsES, + immutableStorage: a.caps.TextureStorage, } queue := &Queue{ diff --git a/hal/gles/capabilities.go b/hal/gles/capabilities.go index b675a46a..ec220de2 100644 --- a/hal/gles/capabilities.go +++ b/hal/gles/capabilities.go @@ -70,6 +70,14 @@ type AdapterCapabilities struct { // emulated: the naga_vs_first_instance uniform and instance-stepped // vertex buffer offsets. FullyFeaturedInstancing bool + + // TextureStorage is Rust wgpu-hal's TEXTURE_STORAGE: the context allocates + // immutable texture storage with glTexStorage2D (GL 4.2+, ES 3.0+, or + // GL_ARB_texture_storage). Textures are allocated with it when present + // (CreateTexture), as Rust does: an OpenGL ES image binding + // (glBindImageTexture, storage textures in compute) is rejected for a + // texture with mutable storage, so a storage texture needs it. + TextureStorage bool } // fullyFeaturedInstancing reports Rust wgpu-hal's FULLY_FEATURED_INSTANCING @@ -80,6 +88,18 @@ func fullyFeaturedInstancing(exts map[string]bool, glMajor, glMinor int, isES bo hasExtension(exts, "GL_ARB_shader_draw_parameters") } +// hasTextureStorage reports Rust wgpu-hal's TEXTURE_STORAGE private capability +// (src/gles/adapter.rs): immutable texture storage, core in desktop GL 4.2 and +// ES 3.0, and available through GL_ARB_texture_storage before that. The entry +// point is the second check: an ES 3.0 context (or an ARB_texture_storage +// context) has glTexStorage2D, and without a texture storage allocation the +// driver may still refuse an image binding. +func hasTextureStorage(glCtx *gl.Context, exts map[string]bool, glMajor, glMinor int, isES bool) bool { + storage := glVersionAtLeast(glMajor, glMinor, isES, [2]int{3, 0}, [2]int{4, 2}) || + hasExtension(exts, "GL_ARB_texture_storage") + return storage && glCtx.HasImmutableStorage() +} + // queryAdapterCapabilities queries all adapter capabilities from the GL context. // This is the single entry point for adapter detection; callers pass the result // to buildExposedAdapter to construct the ExposedAdapter. @@ -116,6 +136,7 @@ func queryAdapterCapabilities(glCtx *gl.Context) (AdapterCapabilities, error) { // --- 4. Detect features --- caps.FullyFeaturedInstancing = fullyFeaturedInstancing(caps.Extensions, caps.GLMajor, caps.GLMinor, caps.IsES) && glCtx.HasBaseInstance() + caps.TextureStorage = hasTextureStorage(glCtx, caps.Extensions, caps.GLMajor, caps.GLMinor, caps.IsES) caps.Features = queryFeatures(caps.Extensions, caps.GLMajor, caps.GLMinor, caps.IsES, glCtx) // --- 5. Query limits --- diff --git a/hal/gles/command.go b/hal/gles/command.go index 8f0fb07d..fd0afccb 100644 --- a/hal/gles/command.go +++ b/hal/gles/command.go @@ -2313,13 +2313,17 @@ func (c *CopyTextureToBufferCommand) bindSourceFBO(ctx *gl.Context, attachment, // readGLBuffer returns a copy of buf's full GL contents (zeros if the buffer // cannot be mapped). Caller holds the context. +// +// The mapping goes through glMapBufferRange: it is the mapping entry point of +// OpenGL ES 3.0, which has no glMapBuffer, and Rust wgpu-hal maps with it on +// every GL context (gles/device.rs map_buffer). func readGLBuffer(ctx *gl.Context, buf *Buffer) []byte { data := make([]byte, buf.size) if buf.id == 0 || buf.size == 0 { return data } ctx.BindBuffer(gl.COPY_READ_BUFFER, buf.id) - if ptr := ctx.MapBuffer(gl.COPY_READ_BUFFER, gl.READ_ONLY); ptr != 0 { + if ptr := ctx.MapBufferRange(gl.COPY_READ_BUFFER, 0, int64(buf.size), gl.MAP_READ_BIT); ptr != 0 { src := unsafe.Slice(*(**byte)(unsafe.Pointer(&ptr)), buf.size) copy(data, src) ctx.UnmapBuffer(gl.COPY_READ_BUFFER) diff --git a/hal/gles/compute_readback_es_linux_test.go b/hal/gles/compute_readback_es_linux_test.go new file mode 100644 index 00000000..13a659f8 --- /dev/null +++ b/hal/gles/compute_readback_es_linux_test.go @@ -0,0 +1,177 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build integration && linux && !(js && wasm) && wgpu_gles_force_es + +package gles + +import ( + "encoding/binary" + "testing" + "unsafe" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +// esComputeVariants are three small compute shaders, each reading one kind of +// binding and writing a storage buffer: a constant, a storage buffer read, and +// a uniform block read. All three are read back through CopyBufferToBuffer and +// MapBuffer. +var esComputeVariants = []struct { + name string + wgsl string + want uint32 // the value of the first output word +}{ + { + name: "constant", + wgsl: ` +@group(0) @binding(0) var output: array; +@compute @workgroup_size(4) +fn main(@builtin(global_invocation_id) id: vec3) { + output[id.x] = 7u; +}`, + want: 7, + }, + { + name: "storage_buffer", + wgsl: ` +@group(0) @binding(0) var output: array; +@group(0) @binding(1) var input: array; +@compute @workgroup_size(4) +fn main(@builtin(global_invocation_id) id: vec3) { + output[id.x] = input[id.x] + 1u; +}`, + want: 101, // input[0] is 100 + }, + { + name: "uniform", + wgsl: ` +@group(0) @binding(0) var output: array; +struct Params { count: u32 } +@group(0) @binding(1) var params: Params; +@compute @workgroup_size(4) +fn main(@builtin(global_invocation_id) id: vec3) { + output[id.x] = params.count; +}`, + want: 42, + }, +} + +// TestESComputeReadsItsBindings dispatches compute shaders on an OpenGL ES +// context and reads their output buffers back through MapBuffer. +// +// The readback used to return zeros for every buffer: Device.MapBuffer mapped +// with glMapBuffer, which OpenGL ES does not have (ES maps with +// glMapBufferRange), so the mapping failed silently and the CPU copy stayed +// zero-filled. Rendering and texture readbacks were unaffected because those +// go through the CPU shadow, which is why only compute and buffer readbacks +// looked broken on ES (P1-GLES-025 item 3). Every compute workload whose +// shader read or wrote only buffers read as zeros. +func TestESComputeReadsItsBindings(t *testing.T) { + _, dev, queue := openForcedESDevice(t) + + const words = 8 + for _, tt := range esComputeVariants { + t.Run(tt.name, func(t *testing.T) { + mk := func(label string, size uint64, usage gputypes.BufferUsage) hal.Buffer { + b, err := dev.CreateBuffer(&hal.BufferDescriptor{Label: label, Size: size, Usage: usage}) + if err != nil { + t.Fatalf("CreateBuffer(%s): %v", label, err) + } + t.Cleanup(func() { dev.DestroyBuffer(b) }) + return b + } + in := mk("input", words*4, gputypes.BufferUsageStorage|gputypes.BufferUsageCopyDst) + out := mk("output", words*4, gputypes.BufferUsageStorage|gputypes.BufferUsageCopySrc) + staging := mk("staging", words*4, gputypes.BufferUsageCopyDst|gputypes.BufferUsageMapRead) + params := mk("params", 4, gputypes.BufferUsageUniform|gputypes.BufferUsageCopyDst) + + inData := make([]byte, words*4) + for i := 0; i < words; i++ { + binary.LittleEndian.PutUint32(inData[i*4:], uint32(100+i)) + } + if err := queue.WriteBuffer(in, 0, inData); err != nil { + t.Fatalf("WriteBuffer(input): %v", err) + } + if err := queue.WriteBuffer(params, 0, binary.LittleEndian.AppendUint32(nil, 42)); err != nil { + t.Fatalf("WriteBuffer(params): %v", err) + } + + module, err := dev.CreateShaderModule(&hal.ShaderModuleDescriptor{Label: tt.name, Source: hal.ShaderSource{WGSL: tt.wgsl}}) + if err != nil { + t.Fatalf("CreateShaderModule: %v", err) + } + t.Cleanup(func() { dev.DestroyShaderModule(module) }) + + entries := []gputypes.BindGroupLayoutEntry{ + {Binding: 0, Visibility: gputypes.ShaderStageCompute, Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeStorage}}, + {Binding: 1, Visibility: gputypes.ShaderStageCompute, Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeReadOnlyStorage}}, + } + bound := []gputypes.BindGroupEntry{ + {Binding: 0, Resource: gputypes.BufferBinding{Buffer: out.NativeHandle(), Size: words * 4}}, + {Binding: 1, Resource: gputypes.BufferBinding{Buffer: in.NativeHandle(), Size: words * 4}}, + } + if tt.name == "uniform" { + entries[1] = gputypes.BindGroupLayoutEntry{Binding: 1, Visibility: gputypes.ShaderStageCompute, Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeUniform}} + bound[1] = gputypes.BindGroupEntry{Binding: 1, Resource: gputypes.BufferBinding{Buffer: params.NativeHandle(), Size: 4}} + } + bgl, err := dev.CreateBindGroupLayout(&hal.BindGroupLayoutDescriptor{Label: tt.name, Entries: entries}) + if err != nil { + t.Fatalf("CreateBindGroupLayout: %v", err) + } + bg, err := dev.CreateBindGroup(&hal.BindGroupDescriptor{Label: tt.name, Layout: bgl, Entries: bound}) + if err != nil { + t.Fatalf("CreateBindGroup: %v", err) + } + pl, err := dev.CreatePipelineLayout(&hal.PipelineLayoutDescriptor{Label: tt.name, BindGroupLayouts: []hal.BindGroupLayout{bgl}}) + if err != nil { + t.Fatalf("CreatePipelineLayout: %v", err) + } + pipe, err := dev.CreateComputePipeline(&hal.ComputePipelineDescriptor{ + Label: tt.name, Layout: pl, Compute: hal.ComputeState{Module: module, EntryPoint: "main"}, + }) + if err != nil { + t.Fatalf("CreateComputePipeline: %v", err) + } + + enc, err := dev.CreateCommandEncoder(&hal.CommandEncoderDescriptor{Label: tt.name}) + if err != nil { + t.Fatalf("CreateCommandEncoder: %v", err) + } + if err := enc.BeginEncoding(tt.name); err != nil { + t.Fatalf("BeginEncoding: %v", err) + } + pass := enc.BeginComputePass(&hal.ComputePassDescriptor{Label: tt.name}) + pass.SetPipeline(pipe) + pass.SetBindGroup(0, bg, nil) + pass.Dispatch(2, 1, 1) + pass.End() + enc.CopyBufferToBuffer(out, staging, []hal.BufferCopy{{Size: words * 4}}) + cmd, err := enc.EndEncoding() + if err != nil { + t.Fatalf("EndEncoding: %v", err) + } + if _, err := queue.Submit([]hal.CommandBuffer{cmd}); err != nil { + t.Fatalf("Submit: %v", err) + } + if err := dev.WaitIdle(); err != nil { + t.Fatalf("WaitIdle: %v", err) + } + + m, err := dev.MapBuffer(staging, 0, words*4) + if err != nil { + t.Fatalf("MapBuffer: %v", err) + } + got := make([]byte, words*4) + copy(got, unsafe.Slice((*byte)(m.Ptr), len(got))) + if err := dev.UnmapBuffer(staging); err != nil { + t.Fatalf("UnmapBuffer: %v", err) + } + const index = 0 + if v := binary.LittleEndian.Uint32(got[index*4:]); v != tt.want { + t.Errorf("output[%d] = %d, want %d (words %v)", index, v, tt.want, got) + } + }) + } +} diff --git a/hal/gles/device.go b/hal/gles/device.go index 93cfebf1..6eb9a4a7 100644 --- a/hal/gles/device.go +++ b/hal/gles/device.go @@ -58,6 +58,16 @@ type Device struct { // own (resetFramebufferAttachments). maxColorAttachments uint32 + // isES is true for an OpenGL ES context: a multisampled texture is then a + // renderbuffer, which ES has and multisample texture images it lacks + // (createMultisampleRenderbuffer). + isES bool + + // immutableStorage is AdapterCapabilities.TextureStorage: 2D and cube + // textures are allocated with glTexStorage2D, which an ES image binding + // (glBindImageTexture) requires. + immutableStorage bool + // views maps TextureView.NativeHandle values back to the views, for // CreateBindGroup (newBindGroup). views handletable.Table[TextureView] @@ -144,15 +154,9 @@ func (d *Device) MapBuffer(buffer hal.Buffer, offset, size uint64) (hal.BufferMa case len(buf.data) == int(buf.size): copy(buf.mapped, buf.data) case buf.id != 0: - glCtx.BindBuffer(gl.COPY_READ_BUFFER, buf.id) - glPtr := glCtx.MapBuffer(gl.COPY_READ_BUFFER, gl.READ_ONLY) - if glPtr != 0 { - basePtr := *(**byte)(unsafe.Pointer(&glPtr)) - src := unsafe.Slice(basePtr, buf.size) - copy(buf.mapped, src) - glCtx.UnmapBuffer(gl.COPY_READ_BUFFER) - } - glCtx.BindBuffer(gl.COPY_READ_BUFFER, 0) + // readGLBuffer maps with glMapBufferRange, the entry point + // OpenGL ES has in place of glMapBuffer. + copy(buf.mapped, readGLBuffer(glCtx, buf)) } d.ctx.Unlock() } @@ -203,8 +207,6 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { glCtx := d.ctx.Lock() defer d.ctx.Unlock() - id := glCtx.GenTextures(1) - sampleCount := desc.SampleCount if sampleCount == 0 { sampleCount = 1 @@ -212,10 +214,35 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { target := textureTarget(desc.Dimension, desc.Size, sampleCount, d.cubeArrays) - glCtx.BindTexture(target, id) - internalFormat, format, dataType := glFormat.internal, glFormat.external, glFormat.dataType + // A multisampled render target on an ES context is a renderbuffer: ES has + // no glTexImage2DMultisample (desktop only) and no multisample texture + // images before ES 3.1. See Texture.renderBuffer and + // createMultisampleRenderbuffer. No WGL context of this backend is an ES + // context today; the branch keeps both platform files one definition of + // what a multisampled texture is. + if d.isES && sampleCount > 1 { + rb, err := createMultisampleRenderbuffer(glCtx, sampleCount, internalFormat, desc.Size) + if err != nil { + return nil, err + } + return &Texture{ + id: rb, + target: target, + format: desc.Format, + dimension: desc.Dimension, + size: desc.Size, + mipLevels: desc.MipLevelCount, + sampleCount: sampleCount, + glCtx: glCtx, + renderBuffer: true, + }, nil + } + + id := glCtx.GenTextures(1) + glCtx.BindTexture(target, id) + // Allocate texture storage switch { case desc.Format.IsCompressed(): @@ -230,6 +257,14 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { int32(desc.Size.Width), int32(desc.Size.Height), true) case target == gl.TEXTURE_2D: + if d.immutableStorage { + // Immutable storage, as Rust wgpu-hal allocates a texture with + // TEXTURE_STORAGE (gles/device.rs create_texture); every later + // write goes through TexSubImage2D. + glCtx.TexStorage2D(target, int32(max(desc.MipLevelCount, 1)), internalFormat, + int32(desc.Size.Width), int32(desc.Size.Height)) + break + } for level := uint32(0); level < desc.MipLevelCount; level++ { width := maxInt32(1, int32(desc.Size.Width>>level)) height := maxInt32(1, int32(desc.Size.Height>>level)) @@ -238,6 +273,13 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { } case target == gl.TEXTURE_CUBE_MAP: + if d.immutableStorage { + // One glTexStorage2D allocates every face and level (Rust + // wgpu-hal gles/device.rs create_texture). + glCtx.TexStorage2D(target, int32(max(desc.MipLevelCount, 1)), internalFormat, + int32(desc.Size.Width), int32(desc.Size.Height)) + break + } for face := uint32(0); face < 6; face++ { faceTarget := gl.TEXTURE_CUBE_MAP_POSITIVE_X + face for level := uint32(0); level < desc.MipLevelCount; level++ { @@ -249,14 +291,19 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { } case isLayeredTarget(target): - // Every level, as Rust wgpu-hal does with tex_image_3d on a context - // without immutable storage (gles/device.rs create_texture); the 2D - // path above uses mutable TexImage2D storage as well. + // One glTexStorage3D allocates every level when the context has + // immutable storage; without it every level is allocated with + // TexImage3D, as Rust wgpu-hal does (gles/device.rs create_texture). if !glCtx.HasLayeredTextures() { glCtx.BindTexture(target, 0) glCtx.DeleteTextures(id) return nil, errNoLayeredTextures } + if d.immutableStorage { + w, h, depth := layeredLevelSize(target, desc.Size, 0) + glCtx.TexStorage3D(target, int32(max(desc.MipLevelCount, 1)), internalFormat, w, h, depth) + break + } for level := uint32(0); level < desc.MipLevelCount; level++ { w, h, depth := layeredLevelSize(target, desc.Size, level) glCtx.TexImage3D(target, int32(level), int32(internalFormat), w, h, depth, 0, format, dataType, nil) diff --git a/hal/gles/device_linux.go b/hal/gles/device_linux.go index dbfc20a5..0572c349 100644 --- a/hal/gles/device_linux.go +++ b/hal/gles/device_linux.go @@ -63,6 +63,16 @@ type Device struct { // own (resetFramebufferAttachments). maxColorAttachments uint32 + // isES is true for an OpenGL ES context: a multisampled texture is then a + // renderbuffer, which ES has and multisample texture images it lacks + // (createMultisampleRenderbuffer). + isES bool + + // immutableStorage is AdapterCapabilities.TextureStorage: 2D and cube + // textures are allocated with glTexStorage2D, which an ES image binding + // (glBindImageTexture) requires. + immutableStorage bool + // views maps TextureView.NativeHandle values back to the views, for // CreateBindGroup (newBindGroup). views handletable.Table[TextureView] @@ -162,15 +172,9 @@ func (d *Device) MapBuffer(buffer hal.Buffer, offset, size uint64) (hal.BufferMa case len(buf.data) == int(buf.size): copy(buf.mapped, buf.data) case buf.id != 0: - glCtx.BindBuffer(gl.COPY_READ_BUFFER, buf.id) - glPtr := glCtx.MapBuffer(gl.COPY_READ_BUFFER, gl.READ_ONLY) - if glPtr != 0 { - basePtr := *(**byte)(unsafe.Pointer(&glPtr)) - src := unsafe.Slice(basePtr, buf.size) - copy(buf.mapped, src) - glCtx.UnmapBuffer(gl.COPY_READ_BUFFER) - } - glCtx.BindBuffer(gl.COPY_READ_BUFFER, 0) + // glMapBuffer does not exist on OpenGL ES; readGLBuffer maps + // with glMapBufferRange, which every GL context here has. + copy(buf.mapped, readGLBuffer(glCtx, buf)) } } } @@ -226,7 +230,6 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { if err != nil { return nil, err } - id := glCtx.GenTextures(1) sampleCount := desc.SampleCount if sampleCount == 0 { @@ -243,8 +246,6 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { target := textureTarget(desc.Dimension, desc.Size, sampleCount, d.cubeArrays) - glCtx.BindTexture(target, id) - internalFormat, format, dataType := glFormat.internal, glFormat.external, glFormat.dataType // glGetError is sticky: an error left by an earlier call was reported as @@ -261,6 +262,40 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { "error", fmt.Sprintf("0x%x", stale)) } + // A multisampled render target on an ES context is a renderbuffer: ES has no + // glTexImage2DMultisample (desktop only), and no multisample texture images + // before ES 3.1. It needs no texture object, so no GenTextures/BindTexture + // either — on ES 3.0 the multisample target itself is an invalid enum. + // Rust wgpu-hal uses renderbuffers for render targets (gles/device.rs + // create_texture). + if d.isES && sampleCount > 1 { + rb, err := createMultisampleRenderbuffer(glCtx, sampleCount, internalFormat, desc.Size) + if err != nil { + return nil, err + } + hal.Logger().Debug("gles: multisampled renderbuffer created", + "label", desc.Label, + "format", desc.Format, + "width", desc.Size.Width, + "height", desc.Size.Height, + "samples", sampleCount, + ) + return &Texture{ + id: rb, + target: target, + format: desc.Format, + dimension: desc.Dimension, + size: desc.Size, + mipLevels: desc.MipLevelCount, + sampleCount: sampleCount, + glCtx: glCtx, + renderBuffer: true, + }, nil + } + + id := glCtx.GenTextures(1) + glCtx.BindTexture(target, id) + // Allocate texture storage switch { case desc.Format.IsCompressed(): @@ -285,6 +320,22 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { } case target == gl.TEXTURE_2D: + if d.immutableStorage { + // Immutable storage, as Rust wgpu-hal allocates a texture with + // TEXTURE_STORAGE (gles/device.rs create_texture): ES rejects an + // image binding of a texture with mutable storage, so a storage + // texture would fail to bind in compute. Every later write goes + // through TexSubImage2D, which immutable storage accepts. + glCtx.TexStorage2D(target, int32(max(desc.MipLevelCount, 1)), internalFormat, + int32(desc.Size.Width), int32(desc.Size.Height)) + if glErr := glCtx.GetError(); glErr != 0 { + glCtx.BindTexture(target, 0) + glCtx.DeleteTextures(id) + return nil, fmt.Errorf("gles: TexStorage2D failed: GL error 0x%x (format=0x%x, %dx%d, %d levels)", + glErr, internalFormat, desc.Size.Width, desc.Size.Height, desc.MipLevelCount) + } + break + } for level := uint32(0); level < desc.MipLevelCount; level++ { width := maxInt32(1, int32(desc.Size.Width>>level)) height := maxInt32(1, int32(desc.Size.Height>>level)) @@ -298,6 +349,19 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { } case target == gl.TEXTURE_CUBE_MAP: + if d.immutableStorage { + // One glTexStorage2D allocates every face and level (Rust + // wgpu-hal gles/device.rs create_texture). + glCtx.TexStorage2D(target, int32(max(desc.MipLevelCount, 1)), internalFormat, + int32(desc.Size.Width), int32(desc.Size.Height)) + if glErr := glCtx.GetError(); glErr != 0 { + glCtx.BindTexture(target, 0) + glCtx.DeleteTextures(id) + return nil, fmt.Errorf("gles: TexStorage2D(cube) failed: GL error 0x%x (format=0x%x, %dx%d, %d levels)", + glErr, internalFormat, desc.Size.Width, desc.Size.Height, desc.MipLevelCount) + } + break + } for face := uint32(0); face < 6; face++ { faceTarget := gl.TEXTURE_CUBE_MAP_POSITIVE_X + face for level := uint32(0); level < desc.MipLevelCount; level++ { @@ -309,7 +373,7 @@ func (d *Device) CreateTexture(desc *TextureDescriptor) (hal.Texture, error) { } case isLayeredTarget(target): - if err := allocateLayeredStorage(glCtx, target, desc, internalFormat, format, dataType); err != nil { + if err := allocateLayeredStorage(glCtx, target, desc, internalFormat, format, dataType, d.immutableStorage); err != nil { glCtx.BindTexture(target, 0) glCtx.DeleteTextures(id) return nil, err @@ -962,13 +1026,24 @@ func maxInt32(a, b int32) int32 { } // allocateLayeredStorage allocates every mip level of the 3D or 2D array -// texture bound to target, as Rust wgpu-hal does with tex_image_3d on a -// context without immutable storage (gles/device.rs create_texture); the 2D -// path uses mutable TexImage2D storage as well. -func allocateLayeredStorage(glCtx *gl.Context, target uint32, desc *TextureDescriptor, internalFormat, format, dataType uint32) error { +// texture bound to target. With immutable storage one glTexStorage3D allocates +// all levels, as Rust wgpu-hal does when TEXTURE_STORAGE is present +// (gles/device.rs create_texture); without it every level is allocated with +// TexImage3D. An image binding of a mutable layered texture is rejected on +// OpenGL ES, so a storage texture of a 2D array needs the immutable path. +func allocateLayeredStorage(glCtx *gl.Context, target uint32, desc *TextureDescriptor, internalFormat, format, dataType uint32, immutable bool) error { if !glCtx.HasLayeredTextures() { return errNoLayeredTextures } + if immutable { + w, h, depth := layeredLevelSize(target, desc.Size, 0) + glCtx.TexStorage3D(target, int32(max(desc.MipLevelCount, 1)), internalFormat, w, h, depth) + if glErr := glCtx.GetError(); glErr != 0 { + return fmt.Errorf("gles: TexStorage3D failed: GL error 0x%x (format=0x%x, %dx%dx%d, %d levels)", + glErr, internalFormat, w, h, depth, desc.MipLevelCount) + } + return nil + } for level := uint32(0); level < desc.MipLevelCount; level++ { w, h, depth := layeredLevelSize(target, desc.Size, level) glCtx.TexImage3D(target, int32(level), int32(internalFormat), w, h, depth, 0, format, dataType, nil) diff --git a/hal/gles/gl/const.go b/hal/gles/gl/const.go index 82de4232..e3549982 100644 --- a/hal/gles/gl/const.go +++ b/hal/gles/gl/const.go @@ -81,6 +81,17 @@ const ( DYNAMIC_READ = 0x88E9 DYNAMIC_COPY = 0x88EA + // Buffer mapping access flags (GL 3.0+ / ES 3.0+), taken by + // glMapBufferRange. The old glMapBuffer takes the READ_ONLY/WRITE_ONLY + // access enums instead, which OpenGL ES does not have: + // glMapBufferRange is the only mapping entry point there. + MAP_READ_BIT = 0x0001 + MAP_WRITE_BIT = 0x0002 + MAP_INVALIDATE_RANGE_BIT = 0x0004 + MAP_INVALIDATE_BUFFER_BIT = 0x0008 + MAP_FLUSH_EXPLICIT_BIT = 0x0010 + MAP_UNSYNCHRONIZED_BIT = 0x0020 + // Texture units TEXTURE0 = 0x84C0 diff --git a/hal/gles/gl/context.go b/hal/gles/gl/context.go index d812abb6..4bdc3727 100644 --- a/hal/gles/gl/context.go +++ b/hal/gles/gl/context.go @@ -63,13 +63,14 @@ type Context struct { glUniformMatrix4fv uintptr // Buffers (GL 1.5+) - glGenBuffers uintptr - glDeleteBuffers uintptr - glBindBuffer uintptr - glBufferData uintptr - glBufferSubData uintptr - glMapBuffer uintptr - glUnmapBuffer uintptr + glGenBuffers uintptr + glDeleteBuffers uintptr + glBindBuffer uintptr + glBufferData uintptr + glBufferSubData uintptr + glMapBuffer uintptr + glMapBufferRange uintptr + glUnmapBuffer uintptr // VAO (GL 3.0+) glGenVertexArrays uintptr @@ -109,11 +110,12 @@ type Context struct { glPixelStorei uintptr // Renderbuffers (GL 3.0+) - glGenRenderbuffers uintptr - glDeleteRenderbuffers uintptr - glBindRenderbuffer uintptr - glRenderbufferStorage uintptr - glFramebufferRenderbuffer uintptr + glGenRenderbuffers uintptr + glDeleteRenderbuffers uintptr + glBindRenderbuffer uintptr + glRenderbufferStorage uintptr + glRenderbufferStorageMultisample uintptr + glFramebufferRenderbuffer uintptr // Blending (GL 1.4+) glBlendFunc uintptr @@ -267,6 +269,7 @@ func (c *Context) Load(getProcAddr ProcAddressFunc) error { c.glBufferData = getProcAddr("glBufferData") c.glBufferSubData = getProcAddr("glBufferSubData") c.glMapBuffer = getProcAddr("glMapBuffer") + c.glMapBufferRange = getProcAddr("glMapBufferRange") c.glUnmapBuffer = getProcAddr("glUnmapBuffer") // VAO @@ -314,6 +317,7 @@ func (c *Context) Load(getProcAddr ProcAddressFunc) error { c.glDeleteRenderbuffers = getProcAddr("glDeleteRenderbuffers") c.glBindRenderbuffer = getProcAddr("glBindRenderbuffer") c.glRenderbufferStorage = getProcAddr("glRenderbufferStorage") + c.glRenderbufferStorageMultisample = getProcAddr("glRenderbufferStorageMultisample") c.glFramebufferRenderbuffer = getProcAddr("glFramebufferRenderbuffer") // Blending @@ -679,6 +683,20 @@ func (c *Context) MapBuffer(target, access uint32) uintptr { return r } +// MapBufferRange maps a range of the buffer bound to target. GL 3.0+ / ES 3.0+ +// core, and the only mapping entry point OpenGL ES has: it has no glMapBuffer. +// GLintptr and GLsizeiptr are pointer-sized, so offset and length are passed as +// uintptr. Returns 0 when the entry point is missing or the driver refuses the +// mapping. +func (c *Context) MapBufferRange(target uint32, offset, length int64, access uint32) uintptr { + if c.glMapBufferRange == 0 { + return 0 + } + r, _, _ := syscall.SyscallN(c.glMapBufferRange, uintptr(target), + uintptr(offset), uintptr(length), uintptr(access)) + return r +} + // UnmapBuffer releases the mapping of a buffer object's data store. // Returns true if the buffer was successfully unmapped, false if the buffer contents // became corrupt during the mapping (GL_FALSE from driver) or glUnmapBuffer is unavailable. @@ -850,6 +868,14 @@ func (c *Context) HasCompressedTextures() bool { c.glCompressedTexSubImage2D != 0 && c.glCompressedTexSubImage3D != 0 } +// HasImmutableStorage reports whether the context has glTexStorage2D, the +// immutable texture storage of GL 4.2 / ES 3.0 (GL_ARB_texture_storage). The +// adapter's version or extension check decides whether it is offered +// (hasTextureStorage); this only confirms the entry point resolved. +func (c *Context) HasImmutableStorage() bool { + return c.glTexStorage2D != 0 +} + // TexStorage2D allocates immutable storage for every level of a 2D texture // or a cube map. No-op if not supported (HasCompressedTextures). func (c *Context) TexStorage2D(target uint32, levels int32, internalformat uint32, width, height int32) { @@ -1071,6 +1097,16 @@ func (c *Context) RenderbufferStorage(target, internalFormat uint32, width, heig uintptr(width), uintptr(height)) } +// RenderbufferStorageMultisample allocates multisampled storage for the +// currently bound renderbuffer (GL 3.0+, ES 3.0+). No-op if not supported. +func (c *Context) RenderbufferStorageMultisample(target uint32, samples int32, internalFormat uint32, width, height int32) { + if c.glRenderbufferStorageMultisample == 0 { + return + } + syscall.SyscallN(c.glRenderbufferStorageMultisample, uintptr(target), uintptr(samples), + uintptr(internalFormat), uintptr(width), uintptr(height)) +} + // FramebufferRenderbuffer attaches a renderbuffer to the currently bound // framebuffer at the specified attachment point. func (c *Context) FramebufferRenderbuffer(target, attachment, renderbufferTarget, renderbuffer uint32) { diff --git a/hal/gles/gl/context_linux.go b/hal/gles/gl/context_linux.go index 217b8ed7..a05edfee 100644 --- a/hal/gles/gl/context_linux.go +++ b/hal/gles/gl/context_linux.go @@ -53,6 +53,7 @@ var ( cifVoid5MultiElem types.CallInterface // void fn(uint32, uint32, void*, int32, int32) - MultiDrawElementsIndirect cifPtr1 types.CallInterface // void* fn(uint32) cifPtr2UU types.CallInterface // void* fn(uint32, uint32) + cifPtr4MapRange types.CallInterface // void* fn(uint32, int64, int64, uint32) - MapBufferRange cifVoid7ReadPx types.CallInterface // void fn(int32, int32, int32, int32, uint32, uint32, void*) cifVoid6TexMS types.CallInterface // void fn(uint32, int32, uint32, int32, int32, uint8) - TexImage2DMultisample cifVoid7BindImage types.CallInterface // void fn(uint32, uint32, int32, uint8, int32, uint32, uint32) - BindImageTexture @@ -68,6 +69,7 @@ var ( cifVoid5TexStorage2D types.CallInterface // void fn(uint32, int32, uint32, int32, int32) - TexStorage2D cifVoid6TexStorage3D types.CallInterface // void fn(uint32, int32, uint32, int32, int32, int32) - TexStorage3D + cifVoid5RBStorageMS types.CallInterface // void fn(uint32, int32, uint32, int32, int32) - RenderbufferStorageMultisample cifVoid9CompSubImg2D types.CallInterface // void fn(uint32, int32*5, uint32, int32, void*) - CompressedTexSubImage2D cifVoid11CompSubImg3D types.CallInterface // void fn(uint32, int32*7, uint32, int32, void*) - CompressedTexSubImage3D @@ -554,6 +556,27 @@ func initCommonCallInterfaces() error { if err := initSyncCopyCallInterfaces(); err != nil { return err } + // void fn(uint32, int32, uint32, int32, int32) - RenderbufferStorageMultisample + if err := ffi.PrepareCallInterface(&cifVoid5RBStorageMS, types.DefaultCall, types.VoidTypeDescriptor, + []*types.TypeDescriptor{ + types.UInt32TypeDescriptor, // target + types.SInt32TypeDescriptor, // samples + types.UInt32TypeDescriptor, // internalformat + types.SInt32TypeDescriptor, // width + types.SInt32TypeDescriptor, // height + }); err != nil { + return err + } + // void* fn(uint32, int64, int64, uint32) - MapBufferRange + if err := ffi.PrepareCallInterface(&cifPtr4MapRange, types.DefaultCall, types.PointerTypeDescriptor, + []*types.TypeDescriptor{ + types.UInt32TypeDescriptor, // target + types.SInt64TypeDescriptor, // offset (GLintptr) + types.SInt64TypeDescriptor, // length (GLsizeiptr) + types.UInt32TypeDescriptor, // access + }); err != nil { + return err + } cifInitialized = true return nil @@ -674,13 +697,14 @@ type Context struct { glUniformMatrix4fv unsafe.Pointer // Buffers (GL 1.5+) - glGenBuffers unsafe.Pointer - glDeleteBuffers unsafe.Pointer - glBindBuffer unsafe.Pointer - glBufferData unsafe.Pointer - glBufferSubData unsafe.Pointer - glMapBuffer unsafe.Pointer - glUnmapBuffer unsafe.Pointer + glGenBuffers unsafe.Pointer + glDeleteBuffers unsafe.Pointer + glBindBuffer unsafe.Pointer + glBufferData unsafe.Pointer + glBufferSubData unsafe.Pointer + glMapBuffer unsafe.Pointer + glMapBufferRange unsafe.Pointer + glUnmapBuffer unsafe.Pointer // VAO (GL 3.0+) glGenVertexArrays unsafe.Pointer @@ -720,11 +744,12 @@ type Context struct { glPixelStorei unsafe.Pointer // Renderbuffers (GL 3.0+) - glGenRenderbuffers unsafe.Pointer - glDeleteRenderbuffers unsafe.Pointer - glBindRenderbuffer unsafe.Pointer - glRenderbufferStorage unsafe.Pointer - glFramebufferRenderbuffer unsafe.Pointer + glGenRenderbuffers unsafe.Pointer + glDeleteRenderbuffers unsafe.Pointer + glBindRenderbuffer unsafe.Pointer + glRenderbufferStorage unsafe.Pointer + glRenderbufferStorageMultisample unsafe.Pointer + glFramebufferRenderbuffer unsafe.Pointer // Blending (GL 1.4+) glBlendFunc unsafe.Pointer @@ -903,6 +928,7 @@ func (c *Context) Load(getProcAddr ProcAddressFunc) error { if gles && c.glMapBuffer == nil { c.glMapBuffer = getProcAddr("glMapBufferOES") } + c.glMapBufferRange = getProcAddr("glMapBufferRange") c.glUnmapBuffer = getProcAddr("glUnmapBuffer") // VAO — GLES may need OES suffix on some Mesa dispatch paths. @@ -961,6 +987,7 @@ func (c *Context) Load(getProcAddr ProcAddressFunc) error { c.glDeleteRenderbuffers = getProcAddr("glDeleteRenderbuffers") c.glBindRenderbuffer = getProcAddr("glBindRenderbuffer") c.glRenderbufferStorage = getProcAddr("glRenderbufferStorage") + c.glRenderbufferStorageMultisample = getProcAddr("glRenderbufferStorageMultisample") c.glFramebufferRenderbuffer = getProcAddr("glFramebufferRenderbuffer") // Blending @@ -1483,6 +1510,26 @@ func (c *Context) MapBuffer(target, access uint32) uintptr { return ptr } +// MapBufferRange maps a range of the buffer bound to target. GL 3.0+ / ES 3.0+ +// core, and the only mapping entry point OpenGL ES has: it has no glMapBuffer +// (Mesa exposes no ES glMapBufferOES either, so the fallback in Load leaves the +// pointer nil). GLintptr and GLsizeiptr are 64-bit on both platforms. Returns 0 +// when the entry point is missing or the driver refuses the mapping. +func (c *Context) MapBufferRange(target uint32, offset, length int64, access uint32) uintptr { + if c.glMapBufferRange == nil { + return 0 + } + var ptr uintptr + args := [4]unsafe.Pointer{ + unsafe.Pointer(&target), + unsafe.Pointer(&offset), + unsafe.Pointer(&length), + unsafe.Pointer(&access), + } + _, _ = ffi.CallFunction(&cifPtr4MapRange, c.glMapBufferRange, unsafe.Pointer(&ptr), args[:]) + return ptr +} + // UnmapBuffer releases the mapping of a buffer object's data store. // Returns true if the buffer was successfully unmapped, false if the buffer contents // became corrupt during the mapping (GL_FALSE from driver) or glUnmapBuffer is unavailable. @@ -1773,6 +1820,14 @@ func (c *Context) HasCompressedTextures() bool { c.glCompressedTexSubImage2D != nil && c.glCompressedTexSubImage3D != nil } +// HasImmutableStorage reports whether the context has glTexStorage2D, the +// immutable texture storage of GL 4.2 / ES 3.0 (GL_ARB_texture_storage). The +// adapter's version or extension check decides whether it is offered +// (hasTextureStorage); this only confirms the entry point resolved. +func (c *Context) HasImmutableStorage() bool { + return c.glTexStorage2D != nil +} + // TexStorage2D allocates immutable storage for every level of a 2D texture // or a cube map. No-op if not supported (HasCompressedTextures). func (c *Context) TexStorage2D(target uint32, levels int32, internalformat uint32, width, height int32) { @@ -2152,6 +2207,22 @@ func (c *Context) RenderbufferStorage(target, internalFormat uint32, width, heig _, _ = ffi.CallFunction(&cifVoid4, c.glRenderbufferStorage, nil, args[:]) } +// RenderbufferStorageMultisample allocates multisampled storage for the +// currently bound renderbuffer (GL 3.0+, ES 3.0+). No-op if not supported. +func (c *Context) RenderbufferStorageMultisample(target uint32, samples int32, internalFormat uint32, width, height int32) { + if c.glRenderbufferStorageMultisample == nil { + return + } + args := [5]unsafe.Pointer{ + unsafe.Pointer(&target), + unsafe.Pointer(&samples), + unsafe.Pointer(&internalFormat), + unsafe.Pointer(&width), + unsafe.Pointer(&height), + } + _, _ = ffi.CallFunction(&cifVoid5RBStorageMS, c.glRenderbufferStorageMultisample, nil, args[:]) +} + // FramebufferRenderbuffer attaches a renderbuffer to the currently bound // framebuffer at the specified attachment point. func (c *Context) FramebufferRenderbuffer(target, attachment, renderbufferTarget, renderbuffer uint32) { diff --git a/hal/gles/layered_texture.go b/hal/gles/layered_texture.go index c6e90763..8824ce9e 100644 --- a/hal/gles/layered_texture.go +++ b/hal/gles/layered_texture.go @@ -251,8 +251,14 @@ func colorAttachmentLayer(view *TextureView, depthSlice *uint32) uint32 { // attachColorLayer attaches color attachment index of the bound framebuffer // to one mip level of tex and, for a layered texture or a cube map, to one // layer of it: the depth slice of a 3D texture, an array layer or a cube -// face. Rust wgpu-hal gles/queue.rs set_attachment does the same. +// face. A renderbuffer (a multisampled target on ES) attaches at the same +// point with glFramebufferRenderbuffer; it has one level and layer. +// Rust wgpu-hal gles/queue.rs set_attachment does the same. func attachColorLayer(ctx *gl.Context, attachment uint32, tex *Texture, mip, layer uint32) { + if tex.renderBuffer { + ctx.FramebufferRenderbuffer(gl.FRAMEBUFFER, attachment, gl.RENDERBUFFER, tex.id) + return + } if isLayeredTarget(tex.target) { ctx.FramebufferTextureLayer(gl.FRAMEBUFFER, attachment, tex.id, int32(mip), int32(layer)) return diff --git a/hal/gles/msaa_es_linux_test.go b/hal/gles/msaa_es_linux_test.go new file mode 100644 index 00000000..22917570 --- /dev/null +++ b/hal/gles/msaa_es_linux_test.go @@ -0,0 +1,164 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build integration && linux && !(js && wasm) && wgpu_gles_force_es + +package gles + +import ( + "testing" + "unsafe" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +// openForcedESDevice opens a device on the forced-ES instance context. +func openForcedESDevice(t *testing.T) (*Adapter, hal.Device, hal.Queue) { + t.Helper() + inst, err := Backend{}.CreateInstance(nil) + if err != nil { + skipUnlessGLESRequired(t, "no EGL: "+err.Error()) + } + t.Cleanup(inst.Destroy) + adapters := inst.EnumerateAdapters(nil) + if len(adapters) == 0 { + skipUnlessGLESRequired(t, "no adapter from the instance context") + } + adapter, ok := adapters[0].Adapter.(*Adapter) + if !ok || adapter.ctx == nil { + skipUnlessGLESRequired(t, "no instance context (Wayland, or EGL without a surfaceless or pbuffer path)") + } + dev, err := adapter.Open(gputypes.Features(0), gputypes.Limits{}) + if err != nil { + t.Fatalf("Open: %v", err) + } + return adapter, dev.Device, dev.Queue +} + +// TestMSAAIsRenderbufferOnES checks that a multisampled render target on an +// OpenGL ES context is a renderbuffer and renders and resolves. ES has no +// glTexImage2DMultisample (desktop only) and ES 3.0 no multisample texture +// images, so CreateTexture used to fail with GL_INVALID_OPERATION 0x502 and +// every pass that resolved MSAA could not be created. Rust wgpu-hal creates +// renderbuffers for render targets (gles/device.rs create_texture). +func TestMSAAIsRenderbufferOnES(t *testing.T) { + adapter, dev, queue := openForcedESDevice(t) + if !adapter.caps.IsES { + t.Fatalf("context version %q is not OpenGL ES under wgpu_gles_force_es", adapter.caps.Version) + } + + const width, height = 8, 8 + msaa, err := dev.CreateTexture(&hal.TextureDescriptor{ + Label: "msaa", + Size: hal.Extent3D{Width: width, Height: height, DepthOrArrayLayers: 1}, + MipLevelCount: 1, + SampleCount: 4, + Dimension: gputypes.TextureDimension2D, + Format: gputypes.TextureFormatRGBA8Unorm, + Usage: gputypes.TextureUsageRenderAttachment, + }) + if err != nil { + t.Fatalf("CreateTexture(msaa 4x): %v", err) + } + t.Cleanup(func() { dev.DestroyTexture(msaa) }) + msaaTex, ok := msaa.(*Texture) + if !ok { + t.Fatalf("CreateTexture returned %T, want *Texture", msaa) + } + if !msaaTex.renderBuffer { + t.Errorf("MSAA texture on ES is not a renderbuffer: an ES context has no multisample texture images") + } + if msaaTex.sampleCount != 4 { + t.Errorf("sampleCount = %d, want 4", msaaTex.sampleCount) + } + + resolve, err := dev.CreateTexture(&hal.TextureDescriptor{ + Label: "resolve", + Size: hal.Extent3D{Width: width, Height: height, DepthOrArrayLayers: 1}, + MipLevelCount: 1, + SampleCount: 1, + Dimension: gputypes.TextureDimension2D, + Format: gputypes.TextureFormatRGBA8Unorm, + Usage: gputypes.TextureUsageRenderAttachment | gputypes.TextureUsageCopySrc, + }) + if err != nil { + t.Fatalf("CreateTexture(resolve): %v", err) + } + t.Cleanup(func() { dev.DestroyTexture(resolve) }) + + msaaView, err := dev.CreateTextureView(msaa, nil) + if err != nil { + t.Fatalf("CreateTextureView(msaa): %v", err) + } + resolveView, err := dev.CreateTextureView(resolve, nil) + if err != nil { + t.Fatalf("CreateTextureView(resolve): %v", err) + } + + const bytesPerRow = 256 + staging, err := dev.CreateBuffer(&hal.BufferDescriptor{ + Label: "staging", + Size: bytesPerRow * height, + Usage: gputypes.BufferUsageCopyDst | gputypes.BufferUsageMapRead, + }) + if err != nil { + t.Fatalf("CreateBuffer(staging): %v", err) + } + t.Cleanup(func() { dev.DestroyBuffer(staging) }) + + enc, err := dev.CreateCommandEncoder(&hal.CommandEncoderDescriptor{Label: "msaa"}) + if err != nil { + t.Fatalf("CreateCommandEncoder: %v", err) + } + if err := enc.BeginEncoding("msaa"); err != nil { + t.Fatalf("BeginEncoding: %v", err) + } + pass := enc.BeginRenderPass(&hal.RenderPassDescriptor{ + Label: "msaa", + ColorAttachments: []hal.RenderPassColorAttachment{{ + View: msaaView, + ResolveTarget: resolveView, + LoadOp: gputypes.LoadOpClear, + StoreOp: gputypes.StoreOpStore, + ClearValue: gputypes.Color{R: 1, G: 0.5, B: 0, A: 1}, + }}, + }) + pass.End() + enc.CopyTextureToBuffer(resolve, staging, []hal.BufferTextureCopy{{ + TextureBase: hal.ImageCopyTexture{Texture: resolve, MipLevel: 0, Aspect: gputypes.TextureAspectAll}, + BufferLayout: hal.ImageDataLayout{Offset: 0, BytesPerRow: bytesPerRow, RowsPerImage: height}, + Size: hal.Extent3D{Width: 4, Height: 4, DepthOrArrayLayers: 1}, + }}) + cmd, err := enc.EndEncoding() + if err != nil { + t.Fatalf("EndEncoding: %v", err) + } + if _, err := queue.Submit([]hal.CommandBuffer{cmd}); err != nil { + t.Fatalf("Submit: %v", err) + } + if err := dev.WaitIdle(); err != nil { + t.Fatalf("WaitIdle: %v", err) + } + + m, err := dev.MapBuffer(staging, 0, bytesPerRow*height) + if err != nil { + t.Fatalf("MapBuffer: %v", err) + } + out := make([]byte, bytesPerRow*height) + copy(out, unsafe.Slice((*byte)(m.Ptr), len(out))) + if err := dev.UnmapBuffer(staging); err != nil { + t.Fatalf("UnmapBuffer: %v", err) + } + // The clear value 1, 0.5, 0, 1 in RGBA8. + want := [4]byte{255, 128, 0, 255} + for y := 0; y < 4; y++ { + for x := 0; x < 4; x++ { + i := y*bytesPerRow + x*4 + got := [4]byte{out[i], out[i+1], out[i+2], out[i+3]} + if got != want { + t.Fatalf("resolved texel (%d,%d) = %v, want %v: the MSAA renderbuffer pass did not resolve", x, y, got, want) + } + } + } +} diff --git a/hal/gles/resource.go b/hal/gles/resource.go index 59f6e672..a2d011a2 100644 --- a/hal/gles/resource.go +++ b/hal/gles/resource.go @@ -44,7 +44,7 @@ func (b *Buffer) NativeHandle() uintptr { return uintptr(b.id) } // Texture implements hal.Texture for OpenGL. type Texture struct { - id uint32 // GL texture object ID + id uint32 // GL texture object ID, or renderbuffer name when renderBuffer target uint32 // GL_TEXTURE_2D, GL_TEXTURE_2D_MULTISAMPLE, etc. format gputypes.TextureFormat dimension gputypes.TextureDimension @@ -53,6 +53,20 @@ type Texture struct { sampleCount uint32 // 1 for regular textures, >1 for MSAA fbo uint32 // GL framebuffer object ID (0 = no FBO created) glCtx *gl.Context + + // renderBuffer is true when the texture is a GL renderbuffer rather than a + // texture object: a multisampled texture on an OpenGL ES context. ES has no + // glTexImage2DMultisample (desktop only) and ES 3.0 no multisample texture + // images at all, so a multisampled render target is a renderbuffer, which + // the framebuffer attaches with glFramebufferRenderbuffer. Rendering and + // the resolve blit read the attachment, not the texture object, so nothing + // else changes. WebGPU allows a multisampled texture only as a render + // attachment, so it is never sampled or copied. + // + // Rust wgpu-hal creates a renderbuffer for every render-usage texture on + // every GL version (gles/device.rs create_texture); we keep textures for + // single-sample targets and use renderbuffers where ES forces it. + renderBuffer bool } // CurrentUsage returns 0 — GLES has no explicit resource state tracking. @@ -68,12 +82,34 @@ func (t *Texture) Destroy() { t.fbo = 0 } if t.id != 0 { - t.glCtx.DeleteTextures(t.id) + if t.renderBuffer { + t.glCtx.DeleteRenderbuffers(t.id) + } else { + t.glCtx.DeleteTextures(t.id) + } t.id = 0 } } } +// createMultisampleRenderbuffer allocates a multisampled renderbuffer of the +// given internal format: the storage of a multisampled render target on a +// context without multisample texture images (OpenGL ES). The caller owns the +// name and releases it with DeleteRenderbuffers. +func createMultisampleRenderbuffer(glCtx *gl.Context, samples uint32, internalFormat uint32, size hal.Extent3D) (uint32, error) { + rb := glCtx.GenRenderbuffers(1) + glCtx.BindRenderbuffer(gl.RENDERBUFFER, rb) + glCtx.RenderbufferStorageMultisample(gl.RENDERBUFFER, int32(samples), internalFormat, + int32(size.Width), int32(size.Height)) + glCtx.BindRenderbuffer(gl.RENDERBUFFER, 0) + if glErr := glCtx.GetError(); glErr != 0 { + glCtx.DeleteRenderbuffers(rb) + return 0, fmt.Errorf("gles: multisampled renderbuffer allocation failed: GL error 0x%x (format=0x%x, samples=%d, %dx%d)", + glErr, internalFormat, samples, size.Width, size.Height) + } + return rb, nil +} + // NativeHandle returns the GL texture object ID. func (t *Texture) NativeHandle() uintptr { return uintptr(t.id) } diff --git a/hal/gles/storage_texture_es_linux_test.go b/hal/gles/storage_texture_es_linux_test.go new file mode 100644 index 00000000..ac22f465 --- /dev/null +++ b/hal/gles/storage_texture_es_linux_test.go @@ -0,0 +1,127 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build integration && linux && !(js && wasm) && wgpu_gles_force_es + +package gles + +import ( + "testing" + "unsafe" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +// esStorageWriteWGSL writes green to a storage texture from compute. +const esStorageWriteWGSL = ` +@group(0) @binding(0) var img: texture_storage_2d; + +@compute @workgroup_size(4) +fn main(@builtin(global_invocation_id) id: vec3) { + textureStore(img, vec2(id.xy), vec4(0.0, 1.0, 0.0, 1.0)); +}` + +// TestESStorageTextureIsWritten stores to a storage texture from compute and +// copies it back. +// +// It failed on OpenGL ES with GL_INVALID_OPERATION from glBindImageTexture: ES +// rejects an image binding of a texture with mutable (glTexImage2D) storage, so +// the image was never bound and the shader stored nowhere. Textures are now +// allocated with glTexStorage2D when the context has immutable storage, as Rust +// wgpu-hal does (gles/device.rs create_texture, the TEXTURE_STORAGE capability), +// and image stores and loads work. +func TestESStorageTextureIsWritten(t *testing.T) { + _, dev, queue := openForcedESDevice(t) + + const size = 8 + tex, err := dev.CreateTexture(&hal.TextureDescriptor{ + Label: "storage", Size: hal.Extent3D{Width: size, Height: size, DepthOrArrayLayers: 1}, + MipLevelCount: 1, SampleCount: 1, Dimension: gputypes.TextureDimension2D, + Format: gputypes.TextureFormatRGBA8Unorm, + Usage: gputypes.TextureUsageStorageBinding | gputypes.TextureUsageCopySrc, + }) + if err != nil { + t.Fatalf("CreateTexture: %v", err) + } + t.Cleanup(func() { dev.DestroyTexture(tex) }) + view, err := dev.CreateTextureView(tex, nil) + if err != nil { + t.Fatalf("CreateTextureView: %v", err) + } + + module, err := dev.CreateShaderModule(&hal.ShaderModuleDescriptor{Label: "store", Source: hal.ShaderSource{WGSL: esStorageWriteWGSL}}) + if err != nil { + t.Fatalf("CreateShaderModule: %v", err) + } + t.Cleanup(func() { dev.DestroyShaderModule(module) }) + bgl, err := dev.CreateBindGroupLayout(&hal.BindGroupLayoutDescriptor{Label: "store", Entries: []gputypes.BindGroupLayoutEntry{{ + Binding: 0, Visibility: gputypes.ShaderStageCompute, + StorageTexture: &gputypes.StorageTextureBindingLayout{ + Access: gputypes.StorageTextureAccessWriteOnly, Format: gputypes.TextureFormatRGBA8Unorm, + ViewDimension: gputypes.TextureViewDimension2D, + }, + }}}) + if err != nil { + t.Fatalf("CreateBindGroupLayout: %v", err) + } + bg, err := dev.CreateBindGroup(&hal.BindGroupDescriptor{Label: "store", Layout: bgl, Entries: []gputypes.BindGroupEntry{{ + Binding: 0, Resource: gputypes.TextureViewBinding{TextureView: view.NativeHandle()}, + }}}) + if err != nil { + t.Fatalf("CreateBindGroup: %v", err) + } + pl, err := dev.CreatePipelineLayout(&hal.PipelineLayoutDescriptor{Label: "store", BindGroupLayouts: []hal.BindGroupLayout{bgl}}) + if err != nil { + t.Fatalf("CreatePipelineLayout: %v", err) + } + pipe, err := dev.CreateComputePipeline(&hal.ComputePipelineDescriptor{Label: "store", Layout: pl, Compute: hal.ComputeState{Module: module, EntryPoint: "main"}}) + if err != nil { + t.Fatalf("CreateComputePipeline: %v", err) + } + staging, err := dev.CreateBuffer(&hal.BufferDescriptor{Label: "staging", Size: 256 * 4, Usage: gputypes.BufferUsageCopyDst | gputypes.BufferUsageMapRead}) + if err != nil { + t.Fatalf("CreateBuffer: %v", err) + } + t.Cleanup(func() { dev.DestroyBuffer(staging) }) + + enc, err := dev.CreateCommandEncoder(&hal.CommandEncoderDescriptor{Label: "store"}) + if err != nil { + t.Fatalf("CreateCommandEncoder: %v", err) + } + if err := enc.BeginEncoding("store"); err != nil { + t.Fatalf("BeginEncoding: %v", err) + } + pass := enc.BeginComputePass(&hal.ComputePassDescriptor{Label: "store"}) + pass.SetPipeline(pipe) + pass.SetBindGroup(0, bg, nil) + pass.Dispatch(2, 2, 1) + pass.End() + enc.CopyTextureToBuffer(tex, staging, []hal.BufferTextureCopy{{ + TextureBase: hal.ImageCopyTexture{Texture: tex, Aspect: gputypes.TextureAspectAll}, + BufferLayout: hal.ImageDataLayout{BytesPerRow: 256, RowsPerImage: 4}, + Size: hal.Extent3D{Width: 4, Height: 4, DepthOrArrayLayers: 1}, + }}) + cmd, err := enc.EndEncoding() + if err != nil { + t.Fatalf("EndEncoding: %v", err) + } + if _, err := queue.Submit([]hal.CommandBuffer{cmd}); err != nil { + t.Fatalf("Submit: %v", err) + } + if err := dev.WaitIdle(); err != nil { + t.Fatalf("WaitIdle: %v", err) + } + m, err := dev.MapBuffer(staging, 0, 256*4) + if err != nil { + t.Fatalf("MapBuffer: %v", err) + } + out := make([]byte, 32) + copy(out, unsafe.Slice((*byte)(m.Ptr), len(out))) + if err := dev.UnmapBuffer(staging); err != nil { + t.Fatalf("UnmapBuffer: %v", err) + } + if out[0] != 0 || out[1] != 255 || out[2] != 0 || out[3] != 255 { + t.Fatalf("texel (0,0) = % x, want 00 ff 00 ff", out[:4]) + } +} From d1e27ec2ea0b96b83f1fb3fc68ea6899bffd7605 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 17:49:53 +0300 Subject: [PATCH 11/16] fix(gles): detect BC textures from core GL versions and ARB names P2-GLES-024. FeatureTextureCompressionBC was offered only when the EXT extension names were listed. Intel's Windows GL 4.6 driver (Iris Xe) offers RGTC as core GL and lists GL_ARB_texture_compression_rgtc, so no BC feature was advertised there even though every BC format decodes. RGTC (BC4-5) is core in desktop GL 3.0 and BPTC (BC6-7) in GL 4.2, and the ARB names count as well; S3TC (BC1-3) stays an extension on both APIs, with the sRGB S3TC extension required on ES. This is better than Rust wgpu-hal, which reads the extension strings only and misses the same driver; the deviation is documented at the detection. Test: a nine-case table of the detection (desktop 4.6 with ARB/EXT names, core 4.2, 3.3 and 4.1 without BPTC, a missing S3TC, and ES with and without the sRGB and BPTC extensions). On Windows the cross-backend TestCompressedTextureDecode/gles-vs-software now compares 24 of 24 formats on the Iris Xe GL adapter, BC included. --- hal/gles/capabilities.go | 41 ++++++++++++------ hal/gles/capabilities_test.go | 81 +++++++++++++++++++++++++++++++++++ 2 files changed, 108 insertions(+), 14 deletions(-) diff --git a/hal/gles/capabilities.go b/hal/gles/capabilities.go index ec220de2..0b34c4b9 100644 --- a/hal/gles/capabilities.go +++ b/hal/gles/capabilities.go @@ -399,7 +399,7 @@ func queryFeatures(exts map[string]bool, glMajor, glMinor int, isES bool, glCtx // Compressed textures are allocated with glTexStorage and written with // glCompressedTexSubImage; without them no compression feature is offered. if hasCompressedTexturePath(exts, glMajor, glMinor, isES, glCtx) { - insertCompressionFeatures(&features, exts, isES) + insertCompressionFeatures(&features, exts, glMajor, glMinor, isES) } // Float32 filterable (linear filtering for float textures) @@ -458,20 +458,33 @@ func hasCompressedTexturePath(exts map[string]bool, glMajor, glMinor int, isES b // insertCompressionFeatures adds the texture compression features whose // formats the context has, as Rust wgpu-hal gles/adapter.rs detects them. -func insertCompressionFeatures(features *gputypes.Features, exts map[string]bool, isES bool) { - // BC: S3TC + RGTC + BPTC. On GLES the sRGB S3TC formats come from their - // own extension; Rust asks only for that one, we also ask for the linear - // S3TC extension that defines the unorm formats. +// +// BC (BC1-7) needs three format families, from extensions on ES and from core +// GL or an extension on desktop GL: S3TC (BC1-3), RGTC (BC4-5, core in desktop +// GL 3.0) and BPTC (BC6-7, core in desktop GL 4.2). Rust wgpu-hal reads the +// extension strings only, so it misses BC on a driver that offers RGTC as core +// GL and lists GL_ARB_texture_compression_rgtc instead of the EXT name (Intel's +// Windows GL 4.6 driver does). Offering the feature there is a deliberate +// deviation from Rust: the formats are there, and the feature is what a +// consumer checks before creating a BC texture. +func insertCompressionFeatures(features *gputypes.Features, exts map[string]bool, glMajor, glMinor int, isES bool) { + // S3TC supplies BC1-3, RGTC BC4-5 and BPTC BC6-7. On ES each family is its + // own extension, and the sRGB S3TC formats come from their own extension + // too (Rust asks only for that one; we also ask for the linear S3TC + // extension that defines the unorm formats). + s3tc := hasExtension(exts, "GL_EXT_texture_compression_s3tc") + var rgtc, bptc bool if isES { - if hasExtension(exts, "GL_EXT_texture_compression_s3tc") && - hasExtension(exts, "GL_EXT_texture_compression_s3tc_srgb") && - hasExtension(exts, "GL_EXT_texture_compression_rgtc") && - hasExtension(exts, "GL_EXT_texture_compression_bptc") { - features.Insert(gputypes.FeatureTextureCompressionBC) - } - } else if hasExtension(exts, "GL_EXT_texture_compression_s3tc") && - hasExtension(exts, "GL_EXT_texture_compression_rgtc") && - hasExtension(exts, "GL_ARB_texture_compression_bptc") { + s3tc = s3tc && hasExtension(exts, "GL_EXT_texture_compression_s3tc_srgb") + rgtc = hasExtension(exts, "GL_EXT_texture_compression_rgtc") + bptc = hasExtension(exts, "GL_EXT_texture_compression_bptc") + } else { + rgtc = glVersionAtLeast(glMajor, glMinor, false, [2]int{0, 0}, [2]int{3, 0}) || + hasExtension(exts, "GL_EXT_texture_compression_rgtc", "GL_ARB_texture_compression_rgtc") + bptc = glVersionAtLeast(glMajor, glMinor, false, [2]int{0, 0}, [2]int{4, 2}) || + hasExtension(exts, "GL_ARB_texture_compression_bptc", "GL_EXT_texture_compression_bptc") + } + if s3tc && rgtc && bptc { features.Insert(gputypes.FeatureTextureCompressionBC) } diff --git a/hal/gles/capabilities_test.go b/hal/gles/capabilities_test.go index f993bae7..5b6d4dd7 100644 --- a/hal/gles/capabilities_test.go +++ b/hal/gles/capabilities_test.go @@ -254,3 +254,84 @@ func TestGlVersionAtLeast(t *testing.T) { } } } + +// TestInsertCompressionFeaturesBC checks the BC detection table over the +// three format families: S3TC (BC1-3) from an extension on both APIs, RGTC +// (BC4-5) from core desktop GL 3.0 or an ARB/EXT name, BPTC (BC6-7) from core +// desktop GL 4.2 or an ARB/EXT name. A driver that offers RGTC as core GL and +// lists only GL_ARB_texture_compression_rgtc is the Intel Windows case the +// Rust detection misses. +func TestInsertCompressionFeaturesBC(t *testing.T) { + ext := func(names ...string) map[string]bool { + m := make(map[string]bool, len(names)) + for _, n := range names { + m[n] = true + } + return m + } + tests := []struct { + name string + exts map[string]bool + major, minor int + isES bool + wantBC bool + }{ + { + name: "desktop 4.6, S3TC and ARB RGTC and ARB BPTC", + exts: ext("GL_EXT_texture_compression_s3tc", "GL_ARB_texture_compression_rgtc", "GL_ARB_texture_compression_bptc"), + major: 4, minor: 6, wantBC: true, + }, + { + name: "desktop 4.6, S3TC and EXT RGTC and EXT BPTC", + exts: ext("GL_EXT_texture_compression_s3tc", "GL_EXT_texture_compression_rgtc", "GL_EXT_texture_compression_bptc"), + major: 4, minor: 6, wantBC: true, + }, + { + name: "desktop 4.2, S3TC only: RGTC and BPTC are core", + exts: ext("GL_EXT_texture_compression_s3tc"), + major: 4, minor: 2, wantBC: true, + }, + { + name: "desktop 3.3, S3TC and ARB RGTC: BPTC not core yet", + exts: ext("GL_EXT_texture_compression_s3tc", "GL_ARB_texture_compression_rgtc"), + major: 3, minor: 3, wantBC: false, + }, + { + name: "desktop 4.1, S3TC and core RGTC: BPTC needs 4.2", + exts: ext("GL_EXT_texture_compression_s3tc"), + major: 4, minor: 1, wantBC: false, + }, + { + name: "desktop 4.6, no S3TC", + exts: ext("GL_ARB_texture_compression_rgtc", "GL_ARB_texture_compression_bptc"), + major: 4, minor: 6, wantBC: false, + }, + { + name: "ES 3.2, all four EXT names", + exts: ext("GL_EXT_texture_compression_s3tc", "GL_EXT_texture_compression_s3tc_srgb", + "GL_EXT_texture_compression_rgtc", "GL_EXT_texture_compression_bptc"), + major: 3, minor: 2, isES: true, wantBC: true, + }, + { + name: "ES 3.2, no sRGB S3TC", + exts: ext("GL_EXT_texture_compression_s3tc", + "GL_EXT_texture_compression_rgtc", "GL_EXT_texture_compression_bptc"), + major: 3, minor: 2, isES: true, wantBC: false, + }, + { + name: "ES 3.2, no BPTC", + exts: ext("GL_EXT_texture_compression_s3tc", "GL_EXT_texture_compression_s3tc_srgb", + "GL_EXT_texture_compression_rgtc"), + major: 3, minor: 2, isES: true, wantBC: false, + }, + } + for _, tt := range tests { + t.Run(tt.name, func(t *testing.T) { + var features gputypes.Features + insertCompressionFeatures(&features, tt.exts, tt.major, tt.minor, tt.isES) + if got := features.Contains(gputypes.FeatureTextureCompressionBC); got != tt.wantBC { + t.Errorf("BC = %v, want %v", got, tt.wantBC) + } + }) + } +} From d5b21db93dac6efb851c682cabb06b4a7597c474 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 18:15:16 +0300 Subject: [PATCH 12/16] test(hal): pin stencil increment-wrap, decrement-wrap and the write mask The stencil tests covered increment-clamp and decrement-wrap, but nothing pinned the wrap operations' result under a write mask, which gg's EvenOdd stroke construct builds on: with stencilWriteMask 0x01 an increment of 1 lands in bit 1 and is dropped, and a decrement of 0 wraps to 255, whose bit 0 survives. TestStencilWrap draws a front- or back-facing triangle into a Depth24PlusStencil8 attachment cleared to a known value, with compare "always", and reads the stencil aspect back through CopyTextureToBuffer. Each case checks every covered texel against the operation and every uncovered one against the clear value, so a backend that runs no stencil operation, or that paints the whole target, fails. Expectations follow the WebGPU stencil operation definitions, VkStencilOp (INCREMENT_AND_WRAP, DECREMENT_AND_WRAP) with the write-mask formula sa = (sa & ~sw) | (sg & sw) of the Vulkan stencil test, and D3D12's D3D12_STENCIL_OP_INCR and DECR. --- hal/allbackends/stencil_wrap_test.go | 287 +++++++++++++++++++++++++++ 1 file changed, 287 insertions(+) create mode 100644 hal/allbackends/stencil_wrap_test.go diff --git a/hal/allbackends/stencil_wrap_test.go b/hal/allbackends/stencil_wrap_test.go new file mode 100644 index 00000000..aa9171c0 --- /dev/null +++ b/hal/allbackends/stencil_wrap_test.go @@ -0,0 +1,287 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build !android && !(js && wasm) && !rust + +package allbackends_test + +import ( + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu" +) + +// The stencil wrap tests pin the increment-wrap and decrement-wrap operations +// and the way the stencil write mask applies to their result on the combined +// Depth24PlusStencil8 format: the parity behavior gg's EvenOdd stroke +// construct builds on. +// +// Expected values, from the specs: +// +// - index.bs "GPUStencilOperation": "increment-wrap" "Increments the current +// stencil value, wrapping to zero if the value exceeds the maximum +// representable value of the ... stencil aspect"; "decrement-wrap" +// "Decrement the current stencil value, wrapping to the maximum +// representable value of the ... stencil aspect if the value goes below +// 0". The stencil aspect is 8 bits, so the maximum is 255. +// - index.bs "GPUDepthStencilState": "stencilWriteMask: Bitmask controlling +// which ... stencil value bits are written to when performing stencil +// operations", and "stencilFront"/"stencilBack" apply to front- and +// back-facing primitives. +// - Vulkan, VkStencilOp: VK_STENCIL_OP_INCREMENT_AND_WRAP "increments the +// current value and wraps to 0 when the maximum value would have been +// exceeded"; VK_STENCIL_OP_DECREMENT_AND_WRAP "decrements the current +// value and wraps to the maximum possible value when the value would go +// below 0"; "For purposes of increment and decrement, the stencil bits are +// considered as an unsigned integer." +// - Vulkan, "Depth/Stencil Operations" (Stencil Test): if the stencil test +// passes and the depth test passes, passOp is used; the result is stored as +// sa = (sa & ~sw) | (sg & sw), where sw is the face's writeMask. +// - D3D12, D3D12_STENCIL_OP: D3D12_STENCIL_OP_INCR "Increment the stencil +// value by 1, and wrap the result if necessary"; D3D12_STENCIL_OP_DECR +// "Decrement the stencil value by 1, and wrap the result if necessary"; +// D3D12_DEPTH_STENCILOP_DESC::StencilWriteMask masks the result. +// - index.bs "GPUStencilOperation"/"GPUFrontFace"/"Polygon Rasterization": +// with frontFace "ccw" a polygon whose area in framebuffer coordinates is +// positive is front-facing; those coordinates start at the top-left corner +// with y growing down (index.bs "Coordinates"). + +const ( + stencilWrapSize = 16 + stencilWrapBytesPerRow = 256 // the copy alignment; the 16 stencil texels of a row fit +) + +var stencilWrapExtent = wgpu.Extent3D{Width: stencilWrapSize, Height: stencilWrapSize, DepthOrArrayLayers: 1} + +// stencilWrapShader draws one triangle twice: vertices 0-2 are front-facing, +// 3-5 the same triangle in the reverse order, so back-facing with frontFace +// "ccw". In the 16x16 target it covers framebuffer x < 8 on every row: the +// edge at x = 8 falls between two columns of pixel centers, and the two other +// edges leave the target, so the covered set is exactly the left half. +const stencilWrapShader = ` +@vertex +fn vs(@builtin(vertex_index) i: u32) -> @builtin(position) vec4f { + var p = array( + vec2f(-3.0, 0.0), vec2f(0.0, -3.0), vec2f(0.0, 3.0), + vec2f(0.0, 3.0), vec2f(0.0, -3.0), vec2f(-3.0, 0.0), + ); + return vec4f(p[i], 0.0, 1.0); +} +@fragment fn fs() -> @location(0) vec4f { return vec4f(0.0, 1.0, 0.0, 1.0); } +` + +// TestStencilWrap draws a triangle into a stencil cleared to a known value and +// reads the stencil aspect back. Every covered texel must hold the value the +// operation produces, and every uncovered one must keep the clear value: a +// backend that runs no stencil operation, or that writes the whole target, +// fails here. +func TestStencilWrap(t *testing.T) { + forEachOverrideBackend(t, gputypes.DownlevelFlagsDepthTextureAndBufferCopies, func(t *testing.T, device *wgpu.Device) { + t.Run("increment-wrap-full-mask", func(t *testing.T) { + // 255 is the maximum of the 8-bit stencil aspect: a front-facing + // increment wraps it to 0. + got := drawStencilWrap(t, device, stencilWrapCase{ + clear: 255, passOp: wgpu.StencilOperationIncrementWrap, writeMask: 0xFF, + }) + checkStencilWrap(t, got, 0, 255) + }) + t.Run("decrement-wrap-full-mask", func(t *testing.T) { + // 0 is the minimum: a back-facing decrement wraps it to 255 and + // stencilBack, not stencilFront, supplies the operation. + got := drawStencilWrap(t, device, stencilWrapCase{ + clear: 0, passOp: wgpu.StencilOperationDecrementWrap, writeMask: 0xFF, back: true, + }) + checkStencilWrap(t, got, 255, 0) + }) + t.Run("increment-wrap-one-bit-mask", func(t *testing.T) { + t.Run("clear-0", func(t *testing.T) { + // 0 increments to 1, and bit 0 is the bit the mask keeps. + got := drawStencilWrap(t, device, stencilWrapCase{ + clear: 0, passOp: wgpu.StencilOperationIncrementWrap, writeMask: 0x01, + }) + checkStencilWrap(t, got, 1, 0) + }) + t.Run("clear-1", func(t *testing.T) { + // 1 increments to 2: the carry lands in bit 1, which the write + // mask drops, so the stored value is (1 & ^0x01) | (2 & 0x01) = 0. + got := drawStencilWrap(t, device, stencilWrapCase{ + clear: 1, passOp: wgpu.StencilOperationIncrementWrap, writeMask: 0x01, + }) + checkStencilWrap(t, got, 0, 1) + }) + }) + t.Run("decrement-wrap-one-bit-mask", func(t *testing.T) { + // 0 decrements to 255, whose bit 0 is 1. + got := drawStencilWrap(t, device, stencilWrapCase{ + clear: 0, passOp: wgpu.StencilOperationDecrementWrap, writeMask: 0x01, back: true, + }) + checkStencilWrap(t, got, 1, 0) + }) + }) +} + +// stencilWrapCovers reports whether stencilWrapShader's triangle covers the +// texel in column x. It covers every row of the left half. +func stencilWrapCovers(x int) bool { return x < stencilWrapSize/2 } + +// checkStencilWrap asserts that every covered texel reads op and every +// uncovered texel still holds the value the pass cleared the stencil to. +func checkStencilWrap(t *testing.T, got []byte, op, clearValue byte) { + t.Helper() + for y := 0; y < stencilWrapSize; y++ { + row := got[y*stencilWrapBytesPerRow:] + for x := 0; x < stencilWrapSize; x++ { + want := clearValue + if stencilWrapCovers(x) { + want = op + } + if v := row[x]; v != want { + t.Fatalf("texel (%d,%d) = %d, want %d (covered = %v, clear = %d, op = %d)", + x, y, v, want, stencilWrapCovers(x), clearValue, op) + } + } + } +} + +// stencilWrapCase is one draw into a stencil target cleared to clear. +type stencilWrapCase struct { + clear uint32 + passOp gputypes.StencilOperation + writeMask uint32 + back bool // draw the back-facing copy of the triangle +} + +// drawStencilWrap clears the stencil aspect of a new 16x16 +// Depth24PlusStencil8 texture to c.clear, draws the triangle with compare +// "always" so that every covered fragment reaches c.passOp, and returns the +// stencil aspect read back through CopyTextureToBuffer. The operation is put +// on stencilFront, or on stencilBack for the back-facing copy, whose other +// face keeps the stencil. +func drawStencilWrap(t *testing.T, device *wgpu.Device, c stencilWrapCase) []byte { + t.Helper() + ds, dsView := newStencilWrapTexture(t, device, "stencil-wrap", gputypes.TextureFormatDepth24PlusStencil8, + wgpu.TextureUsageRenderAttachment|wgpu.TextureUsageCopySrc) + _, colorView := newStencilWrapTexture(t, device, "stencil-wrap-color", gputypes.TextureFormatRGBA8Unorm, + wgpu.TextureUsageRenderAttachment) + + keep := wgpu.StencilFaceState{ + Compare: gputypes.CompareFunctionAlways, + FailOp: wgpu.StencilOperationKeep, DepthFailOp: wgpu.StencilOperationKeep, PassOp: wgpu.StencilOperationKeep, + } + op := keep + op.PassOp = c.passOp + front, back := keep, keep + if c.back { + back = op + } else { + front = op + } + pipeline := newStencilWrapPipeline(t, device, front, back, c.writeMask) + + submitCopy(t, device, func(enc *wgpu.CommandEncoder) { + pass, err := enc.BeginRenderPass(&wgpu.RenderPassDescriptor{ + ColorAttachments: []wgpu.RenderPassColorAttachment{{ + View: colorView, LoadOp: gputypes.LoadOpClear, StoreOp: gputypes.StoreOpStore, + ClearValue: gputypes.Color{A: 1}, + }}, + DepthStencilAttachment: &wgpu.RenderPassDepthStencilAttachment{ + View: dsView, + DepthLoadOp: gputypes.LoadOpClear, DepthStoreOp: gputypes.StoreOpStore, DepthClearValue: 1, + StencilLoadOp: gputypes.LoadOpClear, StencilStoreOp: gputypes.StoreOpStore, + StencilClearValue: c.clear, + }, + }) + if err != nil { + t.Fatalf("BeginRenderPass: %v", err) + } + pass.SetPipeline(pipeline) + args := gputypes.DrawArgs{VertexCount: 3, InstanceCount: 1} + if c.back { + args.FirstVertex = 3 + } + pass.Draw(args) + if err := pass.End(); err != nil { + t.Fatalf("End: %v", err) + } + }) + return readStencilWrap(t, device, ds) +} + +// newStencilWrapPipeline draws stencilWrapShader's triangle with compare +// "always" on both faces, the given stencil face states and write mask. +func newStencilWrapPipeline(t *testing.T, device *wgpu.Device, front, back wgpu.StencilFaceState, + writeMask uint32) *wgpu.RenderPipeline { + t.Helper() + shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{Label: "stencil-wrap", WGSL: stencilWrapShader}) + if err != nil { + t.Fatalf("CreateShaderModule: %v", err) + } + t.Cleanup(shader.Release) + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{Label: "stencil-wrap"}) + if err != nil { + t.Fatalf("CreatePipelineLayout: %v", err) + } + t.Cleanup(layout.Release) + pipeline, err := device.CreateRenderPipeline(&wgpu.RenderPipelineDescriptor{ + Label: "stencil-wrap", + Layout: layout, + Vertex: wgpu.VertexState{Module: shader, EntryPoint: "vs"}, + Fragment: &wgpu.FragmentState{ + Module: shader, EntryPoint: "fs", + Targets: []gputypes.ColorTargetState{{Format: gputypes.TextureFormatRGBA8Unorm, WriteMask: gputypes.ColorWriteMaskAll}}, + }, + Primitive: gputypes.PrimitiveState{ + Topology: gputypes.PrimitiveTopologyTriangleList, FrontFace: gputypes.FrontFaceCCW, CullMode: gputypes.CullModeNone, + }, + DepthStencil: &wgpu.DepthStencilState{ + Format: gputypes.TextureFormatDepth24PlusStencil8, + DepthCompare: gputypes.CompareFunctionAlways, + StencilFront: front, + StencilBack: back, + StencilReadMask: 0xFF, StencilWriteMask: writeMask, + }, + Multisample: gputypes.MultisampleState{Count: 1, Mask: 0xFFFFFFFF}, + }) + if err != nil { + t.Fatalf("CreateRenderPipeline: %v", err) + } + t.Cleanup(pipeline.Release) + return pipeline +} + +// newStencilWrapTexture creates a 16x16 texture of format with a view. +func newStencilWrapTexture(t *testing.T, device *wgpu.Device, label string, format gputypes.TextureFormat, + usage gputypes.TextureUsage) (*wgpu.Texture, *wgpu.TextureView) { + t.Helper() + tex, err := device.CreateTexture(&wgpu.TextureDescriptor{ + Label: label, Size: stencilWrapExtent, MipLevelCount: 1, SampleCount: 1, + Dimension: gputypes.TextureDimension2D, Format: format, Usage: usage, + }) + if err != nil { + t.Fatalf("CreateTexture %s: %v", label, err) + } + t.Cleanup(tex.Release) + view, err := device.CreateTextureView(tex, nil) + if err != nil { + t.Fatalf("CreateTextureView %s: %v", label, err) + } + t.Cleanup(view.Release) + return tex, view +} + +// readStencilWrap copies the stencil aspect of tex into a buffer at +// stencilWrapBytesPerRow and returns the buffer's bytes. +func readStencilWrap(t *testing.T, device *wgpu.Device, tex *wgpu.Texture) []byte { + t.Helper() + bufSize := uint64(stencilWrapBytesPerRow * stencilWrapSize) + staging := newPublicBuffer(t, device, bufSize, wgpu.BufferUsageCopyDst|wgpu.BufferUsageMapRead) + submitCopy(t, device, func(enc *wgpu.CommandEncoder) { + enc.CopyTextureToBuffer(tex, staging, []wgpu.BufferTextureCopy{{ + BufferLayout: wgpu.ImageDataLayout{BytesPerRow: stencilWrapBytesPerRow, RowsPerImage: stencilWrapSize}, + TextureBase: wgpu.ImageCopyTexture{Texture: tex, Aspect: gputypes.TextureAspectStencilOnly}, + Size: stencilWrapExtent, + }}) + }) + return mapStaging(t, staging, bufSize) +} From 7def6d9b5f4f8d87095061f44cfe1f444ce3829d Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 19:12:34 +0300 Subject: [PATCH 13/16] chore(deps): naga v0.21.6, and the DX12 position.w readback asserts 1/w naga v0.21.6 fixes the fragment position builtin's w on HLSL/DXIL (P1-HLSL-037): it passed SV_Position through, whose w is the clip-space w, while WGSL defines the fragment position's w as 1/w. wgpu excluded both dx12 legs from that check since the readback was written; the exclusion is gone and the assertion holds on dx12-fxc and dx12-dxil, verified against the pinned candidate 0b74bc3 and again against the released tag. Also in v0.21.6: the streaming WGSL lexer (behaviour-neutral, byte-identical IR over a 308-file differential), DXIL built-in index into a local array, DXIL refuses workgroup storage above the shader model's limit, HLSL storage stores nested past 20, and an opt-in SPIR-V demote option (off by default; wgpu does not set it yet). IR unchanged, one additive public field. Downstream gate: wgpu GO on the pinned SHA, gg GO on the same candidate. --- CHANGELOG.md | 5 ++++- go.mod | 2 +- go.sum | 4 ++-- hal/allbackends/perspective_readback_test.go | 16 +--------------- 4 files changed, 8 insertions(+), 19 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index d3b13b38..4d1ac735 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -19,10 +19,13 @@ pass input is validated as the spec requires, and the module builds on 32-bit. - **A released bind group layout cannot back a new bind group or pipeline layout** (P2-CORE-023). `CreateBindGroup` and `CreatePipelineLayout` now return `ErrReleased`, as `CreateBindGroup` already did for a released sampler. Before, the released layout's HAL handle reached `vkAllocateDescriptorSets` and crashed inside the driver. - **`SetViewport` is validated as the spec requires** (P2-CORE-019). Non-finite or negative sizes, an origin outside the viewport range, or a depth range outside `[0, 1]` is a `RenderPassValidationError` on the encoder instead of invalid values reaching the driver. Where the spec and Rust wgpu differ, the spec decides. - **A vertex format gputypes does not define is a pipeline-creation error** (P2-CORE-019, `CreateRenderPipelineErrorInvalidVertexFormat`). Each backend used to substitute a format of its own (float32x4 on Vulkan and GLES, an invalid Metal format, an unknown DXGI format). -- **The module builds on GOARCH=386** (P1-INFRA-008). Four comparisons against `TrackerIndex` and `math.MaxUint32` were written as `int` constants and overflowed a 32-bit `int`, so every 386 build of wgpu failed to compile. Comparison is done in `uint64` now; behaviour on 64-bit targets is unchanged. Build only — there is no GPU on 386, and ecosystem CI needs goffi's own 386 support before the other repositories can run the leg. +- **The module builds on GOARCH=386** (P1-INFRA-008). Four comparisons against `TrackerIndex` and `math.MaxUint32` were written as `int` constants and overflowed a 32-bit `int`, so every 386 build of wgpu failed to compile. Comparison is done in `uint64` now; behaviour on 64-bit targets is unchanged. Build only — there is no GPU on 386, and ecosystem CI needs goffi's own 386 support before the other repositories can run the leg. Two test files that did not type-check for 32-bit are fixed with it. +- **GLES: MSAA, buffer mapping and storage textures work on an OpenGL ES context** (P1-GLES-025, items 1 and 3). Multisampled render targets were created with `glTexImage2DMultisample`, which ES does not have, so every MSAA texture failed with `GL_INVALID_OPERATION`; they are renderbuffers now (`glRenderbufferStorageMultisample`), as in Rust wgpu-hal. `MapBuffer` used `glMapBuffer`, absent from ES, so every buffer readback of GPU-written data returned zeros while rendering was correct; it maps with `glMapBufferRange`. And `glBindImageTexture` rejects a texture with mutable storage on ES, so compute storage textures wrote nothing; textures are allocated with immutable storage (`glTexStorage2D`/`glTexStorage3D`) where the context offers it. +- **GLES: BC textures are detected from core GL versions and ARB names** (P2-GLES-024). Detection required only the `EXT_texture_compression_*` names, so a driver that lists `GL_ARB_texture_compression_rgtc` (Intel on Windows) advertised no BC formats at all. RGTC from desktop GL 3.0, BPTC from 4.2 and the `ARB` names are accepted now; the compressed-format readback compares all 24 formats against the software backend on that adapter. ### Changed +- **naga v0.21.6.** The DX12 fragment `@builtin(position).w` is now WGSL's `1/w` (naga P1-HLSL-037): the perspective readback asserts it on both DX12 legs instead of excluding them. Also in it: the streaming WGSL lexer (behaviour-neutral), a DXIL built-in used as a runtime array index, a DXIL refusal of workgroup storage above the model's limit instead of a bad module, HLSL storage stores nested past 20, and an opt-in SPIR-V demote-to-helper (`spirv.Options.DemoteToHelperInvocation`, off by default). - **`Device.PushErrorScope` returns an error for a filter outside the three WebGPU filters** (P3-API-025). It used to map an unknown value to `Validation` silently, pushing a scope that could never capture what the caller asked for. The native build validates in core, `-tags rust` rejects the value before it reaches wgpu-native, which panics on an unknown filter and whose panic would abort the process across the FFI boundary. Calls written as statements are unaffected. ### Infrastructure diff --git a/go.mod b/go.mod index bed8bc39..d7397a4e 100644 --- a/go.mod +++ b/go.mod @@ -7,6 +7,6 @@ require ( github.com/go-webgpu/webgpu v0.5.5 github.com/gogpu/gpucontext v0.31.3 github.com/gogpu/gputypes v0.8.2 - github.com/gogpu/naga v0.21.5 + github.com/gogpu/naga v0.21.6 golang.org/x/sys v0.47.0 ) diff --git a/go.sum b/go.sum index 4e444a09..4c8fcca8 100644 --- a/go.sum +++ b/go.sum @@ -6,7 +6,7 @@ github.com/gogpu/gpucontext v0.31.3 h1:SMnYOQgr+s5issi1R/Oezbo8lV5Lb6HWUZ4HujFgr github.com/gogpu/gpucontext v0.31.3/go.mod h1:pnyWQA9lpvESPAKoqzKlsg3pDMFSkxdjonu4gWvnNtA= github.com/gogpu/gputypes v0.8.2 h1:XGn8Kib6fXOREwKLzYyfEqaXTBCi9Zs2oAWj1+yW+ic= github.com/gogpu/gputypes v0.8.2/go.mod h1:cnXrDMwTpWTvJLW1Vreop3PcT6a2YP/i3s91rPaOavw= -github.com/gogpu/naga v0.21.5 h1:x001iV3GQdDvNJM+J/aQB5DFYbGSA/PWUq7s27C+5XQ= -github.com/gogpu/naga v0.21.5/go.mod h1:15sQaHKkbqXcwTN+hHYGLsA0WBBnkmYzne/eF5p5WEg= +github.com/gogpu/naga v0.21.6 h1:8dq5tHfta78aMdTb5cJlIfNRtxZ/IMMAl6h5ZXxf+UU= +github.com/gogpu/naga v0.21.6/go.mod h1:15sQaHKkbqXcwTN+hHYGLsA0WBBnkmYzne/eF5p5WEg= golang.org/x/sys v0.47.0 h1:o7XGOvZQCADBQQ4Y7VNq2dRWQR7JmOUW8Kxx4ZsNgWs= golang.org/x/sys v0.47.0/go.mod h1:4GL1E5IUh+htKOUEOaiffhrAeqysfVGipDYzABqnCmw= diff --git a/hal/allbackends/perspective_readback_test.go b/hal/allbackends/perspective_readback_test.go index 6e994957..11beee5f 100644 --- a/hal/allbackends/perspective_readback_test.go +++ b/hal/allbackends/perspective_readback_test.go @@ -7,7 +7,6 @@ package allbackends_test import ( "math" - "strings" "testing" "github.com/gogpu/gputypes" @@ -123,9 +122,7 @@ func TestPerspectiveInterpolation(t *testing.T) { perspectiveVaryingTolerance) checkDepth(t, x, y, img.depthAt(x, y), want.depth, viewportTolerance) checkTexel(t, "position.z", x, y, [4]float64{pos[0]}, [4]float64{want.depth}, viewportTolerance) - if !positionWDeviates(t) { - checkTexel(t, "position.w", x, y, [4]float64{pos[1]}, [4]float64{want.invW}, viewportTolerance) - } + checkTexel(t, "position.w", x, y, [4]float64{pos[1]}, [4]float64{want.invW}, viewportTolerance) checkTexel(t, "texture", x, y, [4]float64{pos[2], pos[3]}, [4]float64{want.texel[0], want.texel[1]}, perspectiveTextureTolerance) }) @@ -378,17 +375,6 @@ func perspectiveScene(t *testing.T, device *wgpu.Device, depthTest gputypes.Comp } } -// positionWDeviates reports whether the backend under test is known to give -// the fragment position builtin a w other than WGSL's 1/w: naga's HLSL and -// DXIL backends pass SV_Position through, whose w is the clip-space w (Rust -// naga back/hlsl/conv.rs has the same gap; Tint inverts it in -// src/tint/lang/hlsl/writer/raise/shader_io.cc). Tracked by naga card -// P1-HLSL-037; remove this exclusion once wgpu takes the naga release that -// fixes it. -func positionWDeviates(t *testing.T) bool { - return strings.Contains(t.Name(), "/dx12-") -} - // forEachPerspectivePixel calls fn for every pixel of the target whose // coverage by the rectangle is certain, with its center in NDC and the // triangle covering it (-1 for none). Pixels whose center lies on an edge, From ad143fc67e000baf77f1fc4ac36f4903c22f222d Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 20:35:42 +0300 Subject: [PATCH 14/16] test(hal): naga's flow/short-circuit-or-rt readback case The case naga added for P2-DXIL-107/092, taken verbatim: a short-circuit || chain over runtime values decides an if/else, and the chain's bool local meets at two if merges. As phi i1 WARP ran neither arm's stores; as alloca i1 D3D12's GPU-based validation dropped them. Four invocations, sentinel fill, so a wrong arm shows as a value and a missing store shows as the sentinel. It runs on every backend in the readback corpus. Without the D3D12 debug layer the case passes on naga v0.21.6 as well (that shape is correct there); the discriminating run is WARP + GPU-based validation, kept in the gate harness (P2-INFRA-009) rather than here. --- hal/allbackends/flow_short_circuit_test.go | 102 +++++++++++++++++++++ 1 file changed, 102 insertions(+) create mode 100644 hal/allbackends/flow_short_circuit_test.go diff --git a/hal/allbackends/flow_short_circuit_test.go b/hal/allbackends/flow_short_circuit_test.go new file mode 100644 index 00000000..012eec90 --- /dev/null +++ b/hal/allbackends/flow_short_circuit_test.go @@ -0,0 +1,102 @@ +// Copyright 2026 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build !android && !(js && wasm) && !rust + +package allbackends_test + +import ( + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu" +) + +// flowShortCircuitOrRTWGSL is naga's readback case flow/short-circuit-or-rt +// (naga testdata/readback/flow/short-circuit-or-rt.wgsl, naga fix for +// P2-DXIL-107/092, released in v0.21.7). A short-circuit `||` chain over +// runtime values decides an if/else, and the chain's bool local meets at two +// if merges. As `phi i1` WARP ran neither arm's stores; as `alloca i1` D3D12's +// GPU-based validation dropped them. The case is about values: naga's own +// dxil tests assert the emitted shape. +const flowShortCircuitOrRTWGSL = ` +@group(0) @binding(0) var src: array; +@group(0) @binding(1) var dst: array; + +@compute @workgroup_size(4) +fn main(@builtin(local_invocation_index) li: u32) { + let x = src[li * 3u]; + let y = src[li * 3u + 1u]; + let z = src[li * 3u + 2u]; + + let limit = 65535u; + if (x > limit || y > limit || z > limit) { + dst[li] = 0u; + } else { + dst[li] = x + y + z; + } +} +` + +const ( + flowSrcWords = 12 + flowDstWords = 5 + // flowFill pre-fills dst so "wrote zero" and "wrote nothing" cannot be + // confused: the sentinel surviving in word 4 also proves nothing wrote + // past the four invocations' range. + flowFill = 0xDEADBEEF +) + +// TestFlowShortCircuitOrRT runs naga's flow/short-circuit-or-rt case through +// the public API. src holds three words per invocation: (10,20,30) sums to 60, +// and each of the other three trips the chain on its first, middle or last +// term and must store 0. Word 1 is the one WARP got wrong before naga v0.21.7. +func TestFlowShortCircuitOrRT(t *testing.T) { + forEachOverrideBackend(t, gputypes.DownlevelFlagsComputeShaders, func(t *testing.T, device *wgpu.Device) { + src := newPublicBuffer(t, device, flowSrcWords*4, wgpu.BufferUsageStorage|wgpu.BufferUsageCopyDst) + writeWords(t, device, src, 10, 20, 30, 70000, 20, 30, 10, 70000, 30, 10, 20, 70000) + + dst := newPublicBuffer(t, device, flowDstWords*4, + wgpu.BufferUsageStorage|wgpu.BufferUsageCopySrc|wgpu.BufferUsageCopyDst) + fill := make([]uint32, flowDstWords) + for i := range fill { + fill[i] = flowFill + } + writeWords(t, device, dst, fill...) + + staging := newPublicBuffer(t, device, flowDstWords*4, wgpu.BufferUsageCopyDst|wgpu.BufferUsageMapRead) + + bgl := newBindGroupLayout(t, device, + gputypes.BindGroupLayoutEntry{Binding: 0, Visibility: wgpu.ShaderStageCompute, + Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeReadOnlyStorage}}, + gputypes.BindGroupLayoutEntry{Binding: 1, Visibility: wgpu.ShaderStageCompute, + Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeStorage}}, + ) + group, err := device.CreateBindGroup(&wgpu.BindGroupDescriptor{Layout: bgl, Entries: []wgpu.BindGroupEntry{ + {Binding: 0, Buffer: src, Size: flowSrcWords * 4}, + {Binding: 1, Buffer: dst, Size: flowDstWords * 4}, + }}) + if err != nil { + t.Fatalf("CreateBindGroup: %v", err) + } + t.Cleanup(group.Release) + pipeline := newComputePipeline(t, device, flowShortCircuitOrRTWGSL, bgl) + + enc := newEncoder(t, device) + pass, err := enc.BeginComputePass(nil) + if err != nil { + t.Fatalf("BeginComputePass: %v", err) + } + pass.SetPipeline(pipeline) + pass.SetBindGroup(0, group, nil) + pass.Dispatch(1, 1, 1) + if err := pass.End(); err != nil { + t.Fatalf("ComputePass.End: %v", err) + } + enc.CopyBufferToBuffer(dst, 0, staging, 0, flowDstWords*4) + submit(t, device, enc) + + checkWords(t, "flow/short-circuit-or-rt", readPublicBuffer(t, staging, flowDstWords), + []uint32{60, 0, 0, 0, flowFill}) + }) +} From 95fd262896822b0bb4864f1a6ca742fffe8d25b6 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 20:44:25 +0300 Subject: [PATCH 15/16] chore(deps): naga v0.21.7, the DXIL WARP pair, and its readback case v0.21.7 is v0.21.6 plus the DXIL mem2reg fix (naga P2-DXIL-107/092): bool locals merge as i32 instead of phi i1 or alloca i1, shapes WARP ran no store of and D3D12's GPU-based validation dropped. Verified against the tag the same way as the candidate: the gated SHA 8149de3 and the released module differ only in CHANGELOG.md and ROADMAP.md. The flow/short-circuit-or-rt readback joined the corpus in the previous commit; this one takes the release in. No IR or API change. --- CHANGELOG.md | 2 +- go.mod | 2 +- go.sum | 4 ++-- 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 4d1ac735..333ba63b 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -25,7 +25,7 @@ pass input is validated as the spec requires, and the module builds on 32-bit. ### Changed -- **naga v0.21.6.** The DX12 fragment `@builtin(position).w` is now WGSL's `1/w` (naga P1-HLSL-037): the perspective readback asserts it on both DX12 legs instead of excluding them. Also in it: the streaming WGSL lexer (behaviour-neutral), a DXIL built-in used as a runtime array index, a DXIL refusal of workgroup storage above the model's limit instead of a bad module, HLSL storage stores nested past 20, and an opt-in SPIR-V demote-to-helper (`spirv.Options.DemoteToHelperInvocation`, off by default). +- **naga v0.21.7.** The DX12 fragment `@builtin(position).w` is now WGSL's `1/w` (naga P1-HLSL-037): the perspective readback asserts it on both DX12 legs instead of excluding them. DXIL `mem2reg` promotes bool locals as `i32` (naga P2-DXIL-107/092): a short-circuit `||` chain's local used to be a `phi i1` or an `alloca i1`, and WARP, the D3D12 software adapter, ran neither arm's stores on the first shape while D3D12's GPU-based validation dropped them on the second. The new readback case `flow/short-circuit-or-rt` pins the values on every backend; the discriminating WARP + GPU-based-validation run is kept in the gate harness. Also in the release: the streaming WGSL lexer (behaviour-neutral), a DXIL built-in used as a runtime array index, a DXIL refusal of workgroup storage above the model's limit instead of a bad module, HLSL storage stores nested past 20, and an opt-in SPIR-V demote-to-helper (`spirv.Options.DemoteToHelperInvocation`, off by default). No IR change. - **`Device.PushErrorScope` returns an error for a filter outside the three WebGPU filters** (P3-API-025). It used to map an unknown value to `Validation` silently, pushing a scope that could never capture what the caller asked for. The native build validates in core, `-tags rust` rejects the value before it reaches wgpu-native, which panics on an unknown filter and whose panic would abort the process across the FFI boundary. Calls written as statements are unaffected. ### Infrastructure diff --git a/go.mod b/go.mod index d7397a4e..7790ea42 100644 --- a/go.mod +++ b/go.mod @@ -7,6 +7,6 @@ require ( github.com/go-webgpu/webgpu v0.5.5 github.com/gogpu/gpucontext v0.31.3 github.com/gogpu/gputypes v0.8.2 - github.com/gogpu/naga v0.21.6 + github.com/gogpu/naga v0.21.7 golang.org/x/sys v0.47.0 ) diff --git a/go.sum b/go.sum index 4c8fcca8..218f9840 100644 --- a/go.sum +++ b/go.sum @@ -6,7 +6,7 @@ github.com/gogpu/gpucontext v0.31.3 h1:SMnYOQgr+s5issi1R/Oezbo8lV5Lb6HWUZ4HujFgr github.com/gogpu/gpucontext v0.31.3/go.mod h1:pnyWQA9lpvESPAKoqzKlsg3pDMFSkxdjonu4gWvnNtA= github.com/gogpu/gputypes v0.8.2 h1:XGn8Kib6fXOREwKLzYyfEqaXTBCi9Zs2oAWj1+yW+ic= github.com/gogpu/gputypes v0.8.2/go.mod h1:cnXrDMwTpWTvJLW1Vreop3PcT6a2YP/i3s91rPaOavw= -github.com/gogpu/naga v0.21.6 h1:8dq5tHfta78aMdTb5cJlIfNRtxZ/IMMAl6h5ZXxf+UU= -github.com/gogpu/naga v0.21.6/go.mod h1:15sQaHKkbqXcwTN+hHYGLsA0WBBnkmYzne/eF5p5WEg= +github.com/gogpu/naga v0.21.7 h1:IUkPOC6jGHMuilF/WbkjsV54v0jXCI/ngfxoXw0TbSk= +github.com/gogpu/naga v0.21.7/go.mod h1:15sQaHKkbqXcwTN+hHYGLsA0WBBnkmYzne/eF5p5WEg= golang.org/x/sys v0.47.0 h1:o7XGOvZQCADBQQ4Y7VNq2dRWQR7JmOUW8Kxx4ZsNgWs= golang.org/x/sys v0.47.0/go.mod h1:4GL1E5IUh+htKOUEOaiffhrAeqysfVGipDYzABqnCmw= From cba00f7d164a0d0a103090d083aea6c44373d613 Mon Sep 17 00:00:00 2001 From: Andrey Kolkov Date: Sun, 8 Nov 2026 21:20:13 +0300 Subject: [PATCH 16/16] docs: v0.34.11 state, and the CI and GLES facts the release changed MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ROADMAP: current state and the released v0.34.11 list, with what moves to v0.34.12 (P2-CORE-026/027, P2-INFRA-009, GLES-023, Metal, ASTC). AGENTS.md version line: v0.34.11, naga v0.21.7. ARCHITECTURE: the GLES patterns the release introduced — multisampled renderbuffers, immutable storage, and glMapBufferRange, BC from core versions and ARB names. CONTRIBUTING: the CI jobs the lint-targets and 386 legs added. --- AGENTS.md | 2 +- CONTRIBUTING.md | 2 +- ROADMAP.md | 21 ++++++++++++++++----- docs/ARCHITECTURE.md | 6 +++++- 4 files changed, 23 insertions(+), 8 deletions(-) diff --git a/AGENTS.md b/AGENTS.md index 42e3a11e..b9b64e1b 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -76,7 +76,7 @@ Example: `examples/raytracing-headless/` builds a BLAS through the software HAL ## Current Version -v0.34.10 | Go 1.25+ | Dependencies: naga v0.21.5, gpucontext v0.31.3, gputypes v0.8.2, goffi v0.6.8, webgpu v0.5.5 +v0.34.11 | Go 1.25+ | Dependencies: naga v0.21.7, gpucontext v0.31.3, gputypes v0.8.2, goffi v0.6.8, webgpu v0.5.5 ## Build & Test diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 1ca3e8ad..502b2c18 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -127,7 +127,7 @@ Components: `core`, `hal`, `vulkan`, `metal`, `dx12`, `gles`, `software`, `noop` - Reference the Rust wgpu equivalent for non-trivial HAL/core changes - Add tests for new functionality - Update documentation if needed (CHANGELOG, docs/, README) -- Ensure all CI checks pass (build, cross-compile, tests with the race detector, GLES integration, lint, formatting, deps, Android preview) +- Ensure all CI checks pass (build across three OSes and windows/386, cross-compile for browser and the Rust binding, tests with the race detector, GLES integration, lint on windows/amd64 and darwin/arm64 as well as linux, formatting, deps, Android preview) - Reference related issues ## Testing diff --git a/ROADMAP.md b/ROADMAP.md index 9b331dd0..5116ef0d 100644 --- a/ROADMAP.md +++ b/ROADMAP.md @@ -19,8 +19,11 @@ --- -## Current State: v0.34.10 +## Current State: v0.34.11 +✅ **Resources released after recording stay alive through the work that uses them (v0.34.11)** — a `QuerySet` and a `PipelineLayout` are reference counted as samplers are; a released layout is refused at creation instead of reaching `vkAllocateDescriptorSets`; a pipeline keeps its layout until its own destruction +✅ **Pass input is validated as the spec requires (v0.34.11)** — `SetViewport` per spec section and Rust, vertex formats gputypes does not define refused at pipeline creation +✅ **GLES works on real OpenGL ES contexts (v0.34.11)** — MSAA as renderbuffers, `glMapBufferRange` for mapping, immutable storage textures for compute, BC from core GL versions and ARB names ✅ **Resources released after recording stay alive for the recorded work (v0.34.10)** — `Release()` is a drop, as in Rust wgpu: frame-batching consumers can free textures, views, samplers, bind groups and buffers before the frame's submit ✅ **Software backend renders like the GPU backends (v0.34.9)** — commands at `Queue.Submit`, clip-space clipping, viewport, perspective-correct interpolation, 2x2 quads for derivatives and implicit LOD, two-sided stencil, exact integer values, BC/ETC2 sampling; checked by readback against Vulkan, DX12 and GLES ✅ **Triple-backend architecture (ADR-038)** — Native Go, Rust FFI, Browser WASM via build tags @@ -135,12 +138,20 @@ Before v1.0, a MINOR release is reserved for a large new capability. Bug fixes, spec never allowed, small features, and features that were already claimed but not implemented ship as PATCH releases. -**Next patch release, v0.34.11:** -- [ ] `QuerySet` reference counted; pipelines keep their `PipelineLayout` alive (P1-CORE-021, P2-CORE-022) -- [ ] GLES on OpenGL ES contexts: remaining gaps found by forced-ES runs (P1-GLES-025) -- [ ] GLES: non-renderable copy sources, BC detection on core and ARB contexts (P2-GLES-023, P2-GLES-024) +**Next patch release, v0.34.12:** +- [ ] GLES: compute blit for non-renderable copy sources; snorm readback without EXT_render_snorm; vertex-storage flag tightened against Mesa d3d12 (P2-GLES-023, P1-GLES-025 items 2 and 4) - [ ] Metal: compressed small mips and HDR formats (P2-MTL-001) - [ ] Software: ASTC sampling (P1-SW-029 step 2) +- [ ] Release only: `BindGroupLayout` reference counted as Rust and Dawn do; the Release-vs-create race contract (P2-CORE-027, P2-CORE-026) +- [ ] A tracked WARP + GPU-based-validation test package (P2-INFRA-009, P1-CORE-028 for the selection API) + +**Patch release v0.34.11** (2026-11-08, naga v0.21.7): +- [x] `QuerySet` reference counted; pipelines keep their `PipelineLayout` alive; released layouts refused at creation (P1-CORE-021, P2-CORE-022, P2-CORE-023) +- [x] GLES on OpenGL ES contexts: MSAA, buffer mapping, storage textures (P1-GLES-025 items 1 and 3) +- [x] GLES: BC detection on core and ARB contexts (P2-GLES-024) +- [x] `SetViewport` and vertex formats validated (P2-CORE-019); `PushErrorScope` reports an unknown filter (P3-API-025) +- [x] The module builds on GOARCH=386 (P1-INFRA-008); CI lints windows and darwin targets and builds 386 (P2-INFRA-007) +- [x] naga v0.21.7: DX12 `position.w` is `1/w` (the dx12 exclusion in the readback is gone), DXIL bool locals merge as i32 with the `flow/short-circuit-or-rt` case; DXIL worker storage above the model's limit is refused **Patch release v0.34.10** (2026-11-08, naga v0.21.5): - [x] `Queue.Submit` accepts textures, views, bind groups and buffers released after recording (gogpu#461) diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md index 03d43c06..778f7851 100644 --- a/docs/ARCHITECTURE.md +++ b/docs/ARCHITECTURE.md @@ -166,7 +166,11 @@ Pure Go OpenGL 3.3+ / OpenGL ES 3.0+ implementation; adapters below that minimum - Texture updates via `glTexSubImage2D` (not `glTexImage2D`) — matches Rust wgpu-hal pattern - `GL_DYNAMIC_DRAW` for all writable buffers (Rust wgpu-hal parity — some vendors freeze STATIC_DRAW buffers) - Scissor Y-flip: WebGPU top-left → OpenGL bottom-left origin conversion -- MSAA resolve via `glBlitFramebuffer` +- MSAA resolve via `glBlitFramebuffer`; a multisampled render target is a renderbuffer (`glRenderbufferStorageMultisample` + `glFramebufferRenderbuffer`), as Rust wgpu-hal does — ES has no `glTexImage2DMultisample` +- Textures allocate immutable storage (`glTexStorage2D`/`glTexStorage3D`) wherever the context has it: ES rejects `glBindImageTexture` on a texture with mutable storage +- `glMapBufferRange` for buffer mapping (ES has no `glMapBuffer`) +- Buffer-texture copy sources that GL cannot render to: uploaded through an internal compute pass when the context has compute +- BC advertised from core desktop GL versions (RGTC 3.0, BPTC 4.2) and the `ARB` names as well as the `EXT` ones; S3TC stays an extension and needs sRGB support on ES - Texture unit validation: warns when binding exceeds GL_MAX_TEXTURE_IMAGE_UNITS - Texture formats from Rust wgpu-hal's `describe_texture_format` table; a format without a row is an error, never a silent RGBA8 - Compressed textures (BC, ETC2, ASTC): immutable storage (`glTexStorage`) and `glCompressedTexSubImage` uploads in block rows