From ecb2480b822cefd70e2177c72e8cc98c39553efb Mon Sep 17 00:00:00 2001 From: Mark Date: Tue, 14 Jul 2026 15:23:14 +0300 Subject: [PATCH 1/3] fix(dx12): match indirect descriptor ABI size --- hal/dx12/d3d12/types.go | 6 +++++- hal/dx12/d3d12/types_test.go | 17 +++++++++++++++++ 2 files changed, 22 insertions(+), 1 deletion(-) create mode 100644 hal/dx12/d3d12/types_test.go diff --git a/hal/dx12/d3d12/types.go b/hal/dx12/d3d12/types.go index 62055fb..504b56b 100644 --- a/hal/dx12/d3d12/types.go +++ b/hal/dx12/d3d12/types.go @@ -597,7 +597,11 @@ type D3D12_COMMAND_SIGNATURE_DESC struct { type D3D12_INDIRECT_ARGUMENT_DESC struct { Type D3D12_INDIRECT_ARGUMENT_TYPE // Union for different argument types - Union [8]byte + // The largest member (D3D12_INDIRECT_ARGUMENT_DESC_CONSTANT) is three + // uint32 values, so the C union occupies 12 bytes and the full descriptor + // is 16 bytes including Type. Keeping the exact size matters because + // CreateCommandSignature reads the native descriptor layout directly. + Union [12]byte } // D3D12_DISCARD_REGION describes a discard region. diff --git a/hal/dx12/d3d12/types_test.go b/hal/dx12/d3d12/types_test.go new file mode 100644 index 0000000..16d5052 --- /dev/null +++ b/hal/dx12/d3d12/types_test.go @@ -0,0 +1,17 @@ +//go:build windows && !(js && wasm) + +package d3d12 + +import ( + "testing" + "unsafe" +) + +func TestIndirectArgumentDescABI(t *testing.T) { + if got := unsafe.Sizeof(D3D12_INDIRECT_ARGUMENT_DESC{}); got != 16 { + t.Fatalf("D3D12_INDIRECT_ARGUMENT_DESC size = %d, want 16", got) + } + if got := unsafe.Offsetof(D3D12_INDIRECT_ARGUMENT_DESC{}.Union); got != 4 { + t.Fatalf("D3D12_INDIRECT_ARGUMENT_DESC.Union offset = %d, want 4", got) + } +} From 19534523036ff4a84247fb19c95dfa7d6807eb13 Mon Sep 17 00:00:00 2001 From: Mark Date: Tue, 14 Jul 2026 15:54:14 +0300 Subject: [PATCH 2/3] feat: add fixed-count indirect draws --- CHANGELOG.md | 14 ++ core/command.go | 20 +- core/device_hal_test.go | 4 +- core/indirect_count_test.go | 122 +++++++++++ hal/command.go | 8 +- hal/dx12/adapter.go | 4 + hal/dx12/command.go | 39 +++- hal/gles/command.go | 16 +- hal/gles/gl/context_linux.go | 1 - hal/gles/indirect_test.go | 26 +++ hal/metal/encoder.go | 58 ++++- hal/metal/indirect_count_test.go | 14 ++ hal/noop/command.go | 4 +- hal/noop/noop_test.go | 4 +- hal/software/command.go | 4 +- hal/software/software_test.go | 4 +- hal/vulkan/adapter.go | 2 + hal/vulkan/command.go | 105 ++++++++- hal/vulkan/device.go | 20 +- hal/vulkan/indirect_count_test.go | 107 +++++++++ indirect.go | 53 +++++ indirect_count_integration_test.go | 218 +++++++++++++++++++ internal/browser/renderpass.go | 2 +- internal/browser/renderpass_indirect_test.go | 56 +++++ renderpass_browser.go | 35 ++- renderpass_indirect_browser_test.go | 54 +++++ renderpass_indirect_rust_test.go | 46 ++++ renderpass_indirect_test.go | 77 +++++++ renderpass_indirect_wrappers_test.go | 45 ++++ renderpass_native.go | 38 +++- renderpass_rust.go | 47 +++- wgpu_test.go | 141 ++++++++++++ 32 files changed, 1315 insertions(+), 73 deletions(-) create mode 100644 core/indirect_count_test.go create mode 100644 hal/gles/indirect_test.go create mode 100644 hal/metal/indirect_count_test.go create mode 100644 hal/vulkan/indirect_count_test.go create mode 100644 indirect.go create mode 100644 indirect_count_integration_test.go create mode 100644 internal/browser/renderpass_indirect_test.go create mode 100644 renderpass_indirect_browser_test.go create mode 100644 renderpass_indirect_rust_test.go create mode 100644 renderpass_indirect_test.go create mode 100644 renderpass_indirect_wrappers_test.go diff --git a/CHANGELOG.md b/CHANGELOG.md index 7cb3465..09ea7f4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -5,6 +5,20 @@ All notable changes to this project will be documented in this file. The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.1.0/), and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html). +## [Unreleased] + +### Changed + +- **Counted indirect draws** — added `RenderPassEncoder.MultiDrawIndirect` and + `MultiDrawIndexedIndirect` for consecutive 16-byte and 20-byte argument + records. Existing two-argument `DrawIndirect` and `DrawIndexedIndirect` + remain single-draw APIs. The Vulkan `FeatureMultiDrawIndirect` capability is + used only as a performance hint; Vulkan multi-draw calls fall back to exact + loops when unavailable or over the device limit. GLES indirect drawing + remains unsupported as before. `FeatureMultiDrawIndirectCount` remains + reserved for future GPU-driven count buffers. External HAL adapters must add + the `drawCount` parameter to both indirect draw methods. + ## [0.30.20] - 2026-07-14 ### Fixed diff --git a/core/command.go b/core/command.go index adf704b..a6cedac 100644 --- a/core/command.go +++ b/core/command.go @@ -799,7 +799,11 @@ func (p *CoreRenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstInde // DrawIndirect draws primitives with GPU-generated parameters. func (p *CoreRenderPassEncoder) DrawIndirect(buffer *Buffer, offset uint64) { - if p.ended { + p.MultiDrawIndirect(buffer, offset, 1) +} + +func (p *CoreRenderPassEncoder) MultiDrawIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if p.ended || drawCount == 0 { return } if p.raw != nil && buffer != nil { @@ -807,14 +811,20 @@ func (p *CoreRenderPassEncoder) DrawIndirect(buffer *Buffer, offset uint64) { defer guard.Release() halBuffer := buffer.Raw(guard) if halBuffer != nil { - p.raw.DrawIndirect(halBuffer, offset) + p.raw.DrawIndirect(halBuffer, offset, drawCount) } } } -// DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. +// DrawIndexedIndirect draws one indexed primitive with GPU-generated parameters. func (p *CoreRenderPassEncoder) DrawIndexedIndirect(buffer *Buffer, offset uint64) { - if p.ended { + p.MultiDrawIndexedIndirect(buffer, offset, 1) +} + +// MultiDrawIndexedIndirect draws consecutive indexed primitives with +// GPU-generated parameters. +func (p *CoreRenderPassEncoder) MultiDrawIndexedIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if p.ended || drawCount == 0 { return } if p.raw != nil && buffer != nil { @@ -822,7 +832,7 @@ func (p *CoreRenderPassEncoder) DrawIndexedIndirect(buffer *Buffer, offset uint6 defer guard.Release() halBuffer := buffer.Raw(guard) if halBuffer != nil { - p.raw.DrawIndexedIndirect(halBuffer, offset) + p.raw.DrawIndexedIndirect(halBuffer, offset, drawCount) } } } diff --git a/core/device_hal_test.go b/core/device_hal_test.go index a091384..0abae96 100644 --- a/core/device_hal_test.go +++ b/core/device_hal_test.go @@ -93,8 +93,8 @@ func (mockRenderPassEncoder) SetBlendConstant(_ *gputypes.Color) func (mockRenderPassEncoder) SetStencilReference(_ uint32) {} func (mockRenderPassEncoder) Draw(_, _, _, _ uint32) {} func (mockRenderPassEncoder) DrawIndexed(_, _, _ uint32, _ int32, _ uint32) {} -func (mockRenderPassEncoder) DrawIndirect(_ hal.Buffer, _ uint64) {} -func (mockRenderPassEncoder) DrawIndexedIndirect(_ hal.Buffer, _ uint64) {} +func (mockRenderPassEncoder) DrawIndirect(_ hal.Buffer, _ uint64, _ uint32) {} +func (mockRenderPassEncoder) DrawIndexedIndirect(_ hal.Buffer, _ uint64, _ uint32) {} func (mockRenderPassEncoder) ExecuteBundle(_ hal.RenderBundle) {} // mockComputePassEncoder implements hal.ComputePassEncoder (minimal) diff --git a/core/indirect_count_test.go b/core/indirect_count_test.go new file mode 100644 index 0000000..c3f0391 --- /dev/null +++ b/core/indirect_count_test.go @@ -0,0 +1,122 @@ +//go:build !(js && wasm) + +package core + +import ( + "fmt" + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +type countedRenderPassEncoder struct { + mockRenderPassEncoder + drawCalls int + drawOffset uint64 + drawCount uint32 + calls int + offset uint64 + count uint32 +} + +func (p *countedRenderPassEncoder) DrawIndirect(_ hal.Buffer, offset uint64, count uint32) { + p.drawCalls++ + p.drawOffset = offset + p.drawCount = count +} + +func TestCoreRenderPassEncoderDrawIndirectForwardsCountOnce(t *testing.T) { + device := NewDevice(&mockHALDevice{}, &Adapter{}, gputypes.Features(0), gputypes.DefaultLimits(), "TestDevice") + buffer := NewBuffer(mockBuffer{}, device, gputypes.BufferUsageIndirect, 64, "indirect") + recorder := &countedRenderPassEncoder{} + pass := &CoreRenderPassEncoder{raw: recorder, device: device} + + pass.DrawIndirect(buffer, 4) + if recorder.drawCalls != 1 || recorder.drawOffset != 4 || recorder.drawCount != 1 { + t.Fatalf("single forwarding = calls %d offset %d count %d; want 1, 4, 1", recorder.drawCalls, recorder.drawOffset, recorder.drawCount) + } + pass.MultiDrawIndirect(buffer, 4, 3) + if recorder.drawCalls != 2 || recorder.drawCount != 3 { + t.Fatalf("multi forwarding = calls %d count %d; want 2, 3", recorder.drawCalls, recorder.drawCount) + } +} + +func (p *countedRenderPassEncoder) DrawIndexedIndirect(_ hal.Buffer, offset uint64, count uint32) { + p.calls++ + p.offset = offset + p.count = count +} + +func TestCoreRenderPassEncoderDrawIndexedIndirectForwardsCountOnce(t *testing.T) { + for _, drawCount := range []uint32{1, 3} { + t.Run(fmt.Sprintf("count-%d", drawCount), func(t *testing.T) { + device := NewDevice(&mockHALDevice{}, &Adapter{}, gputypes.Features(0), gputypes.DefaultLimits(), "TestDevice") + buffer := NewBuffer(mockBuffer{}, device, gputypes.BufferUsageIndirect, 64, "indirect") + recorder := &countedRenderPassEncoder{} + pass := &CoreRenderPassEncoder{raw: recorder, device: device} + + pass.MultiDrawIndexedIndirect(buffer, 4, drawCount) + if recorder.calls != 1 { + t.Fatalf("HAL DrawIndexedIndirect calls = %d, want 1", recorder.calls) + } + if recorder.offset != 4 || recorder.count != drawCount { + t.Fatalf("HAL arguments = offset %d, count %d; want 4, %d", recorder.offset, recorder.count, drawCount) + } + }) + } +} + +func TestCoreRenderPassEncoderDrawIndexedIndirectZeroCountIsNoOp(t *testing.T) { + device := NewDevice(&mockHALDevice{}, &Adapter{}, gputypes.Features(0), gputypes.DefaultLimits(), "TestDevice") + recorder := &countedRenderPassEncoder{} + pass := &CoreRenderPassEncoder{raw: recorder, device: device} + + pass.MultiDrawIndexedIndirect(nil, ^uint64(0), 0) + if recorder.calls != 0 { + t.Fatalf("HAL DrawIndexedIndirect calls = %d, want 0", recorder.calls) + } +} + +func TestCoreRenderPassEncoderDrawIndexedIndirectAllocations(t *testing.T) { + device := NewDevice(&mockHALDevice{}, &Adapter{}, gputypes.Features(0), gputypes.DefaultLimits(), "TestDevice") + buffer := NewBuffer(mockBuffer{}, device, gputypes.BufferUsageIndirect, 64, "indirect") + recorder := &countedRenderPassEncoder{} + pass := &CoreRenderPassEncoder{raw: recorder, device: device} + + for _, drawCount := range []uint32{1, 3} { + t.Run(fmt.Sprintf("count-%d", drawCount), func(t *testing.T) { + if allocs := testing.AllocsPerRun(1000, func() { + pass.MultiDrawIndexedIndirect(buffer, 4, drawCount) + }); allocs != 0 { + t.Fatalf("CoreRenderPassEncoder.DrawIndexedIndirect allocations = %v, want 0", allocs) + } + }) + } +} + +func BenchmarkCoreRenderPassEncoderDrawIndexedIndirectCount1(b *testing.B) { + benchmarkCoreRenderPassEncoderDrawIndexedIndirect(b, 1) +} + +func BenchmarkCoreRenderPassEncoderDrawIndexedIndirectCountN(b *testing.B) { + benchmarkCoreRenderPassEncoderDrawIndexedIndirect(b, 1024) +} + +func benchmarkCoreRenderPassEncoderDrawIndexedIndirect(b *testing.B, drawCount uint32) { + device := NewDevice(&mockHALDevice{}, &Adapter{}, gputypes.Features(0), gputypes.DefaultLimits(), "TestDevice") + bufferSize := uint64(drawCount)*20 + 4 + buffer := NewBuffer(mockBuffer{}, device, gputypes.BufferUsageIndirect, bufferSize, "indirect") + recorder := &countedRenderPassEncoder{} + pass := &CoreRenderPassEncoder{raw: recorder, device: device} + + b.ReportAllocs() + b.ResetTimer() + for i := 0; i < b.N; i++ { + pass.MultiDrawIndexedIndirect(buffer, 4, drawCount) + } + b.StopTimer() + if recorder.calls != b.N { + b.Fatalf("HAL DrawIndexedIndirect calls = %d, want %d", recorder.calls, b.N) + } +} diff --git a/hal/command.go b/hal/command.go index 347366a..2598909 100644 --- a/hal/command.go +++ b/hal/command.go @@ -124,12 +124,12 @@ type RenderPassEncoder interface { DrawIndexed(indexCount, instanceCount, firstIndex uint32, baseVertex int32, firstInstance uint32) // DrawIndirect draws primitives with GPU-generated parameters. - // buffer contains DrawIndirectArgs at the given offset. - DrawIndirect(buffer Buffer, offset uint64) + // buffer contains drawCount consecutive 16-byte DrawIndirectArgs records. + DrawIndirect(buffer Buffer, offset uint64, drawCount uint32) // DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. - // buffer contains DrawIndexedIndirectArgs at the given offset. - DrawIndexedIndirect(buffer Buffer, offset uint64) + // buffer contains drawCount consecutive 20-byte DrawIndexedIndirectArgs records. + DrawIndexedIndirect(buffer Buffer, offset uint64, drawCount uint32) // ExecuteBundle executes a pre-recorded render bundle. // Bundles are an optimization for repeated draw calls. diff --git a/hal/dx12/adapter.go b/hal/dx12/adapter.go index de4b6fe..fb9a679 100644 --- a/hal/dx12/adapter.go +++ b/hal/dx12/adapter.go @@ -239,6 +239,9 @@ func (a *Adapter) Info() gputypes.AdapterInfo { // Features returns supported WebGPU features. func (a *Adapter) Features() gputypes.Features { var features gputypes.Features + // ExecuteIndirect supports counted draws natively. This feature is a + // performance hint; the public MultiDraw APIs do not gate on it. + features |= gputypes.Features(gputypes.FeatureMultiDrawIndirect) // Map D3D12 capabilities to WebGPU features // Feature level 11.0+ guarantees basic compute and texture compression @@ -588,6 +591,7 @@ func (a *AdapterLegacy) toExposedAdapter() hal.ExposedAdapter { // Features returns supported WebGPU features for legacy adapter. func (a *AdapterLegacy) Features() gputypes.Features { var features gputypes.Features + features |= gputypes.Features(gputypes.FeatureMultiDrawIndirect) if a.capabilities.FeatureLevel >= d3d12.D3D_FEATURE_LEVEL_11_0 { features |= gputypes.Features(gputypes.FeatureTextureCompressionBC) } diff --git a/hal/dx12/command.go b/hal/dx12/command.go index 543affe..88e45f8 100644 --- a/hal/dx12/command.go +++ b/hal/dx12/command.go @@ -895,33 +895,60 @@ func (e *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex ui // DrawIndirect draws primitives with GPU-generated parameters. // The buffer must contain a D3D12_DRAW_ARGUMENTS struct at the given offset // (vertexCountPerInstance, instanceCount, startVertexLocation, startInstanceLocation — 16 bytes). -func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64) { +func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { buf, ok := buffer.(*Buffer) - if !ok || !e.encoder.isRecording { + if !ok || !e.encoder.isRecording || drawCount == 0 { + return + } + if !indirectRangeFits(buf.size, offset, 16, drawCount) { + return + } + if _, ok := indirectRecordOffset(offset, 16, drawCount-1); !ok { return } e.encoder.cmdList.ExecuteIndirect( e.encoder.device.cmdSignatures.draw, - 1, buf.raw, offset, nil, 0, + drawCount, buf.raw, offset, nil, 0, ) } // DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. // The buffer must contain a D3D12_DRAW_INDEXED_ARGUMENTS struct at the given offset // (indexCountPerInstance, instanceCount, startIndexLocation, baseVertexLocation, startInstanceLocation — 20 bytes). -func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64) { +func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { buf, ok := buffer.(*Buffer) - if !ok || !e.encoder.isRecording { + if !ok || !e.encoder.isRecording || drawCount == 0 { + return + } + if !indirectRangeFits(buf.size, offset, 20, drawCount) { + return + } + if _, ok := indirectRecordOffset(offset, 20, drawCount-1); !ok { return } e.encoder.cmdList.ExecuteIndirect( e.encoder.device.cmdSignatures.drawIndexed, - 1, buf.raw, offset, nil, 0, + drawCount, buf.raw, offset, nil, 0, ) } +func indirectRecordOffset(offset, stride uint64, index uint32) (uint64, bool) { + delta := uint64(index) * stride + if offset > ^uint64(0)-delta { + return 0, false + } + return offset + delta, true +} + +func indirectRangeFits(bufferSize, offset, stride uint64, drawCount uint32) bool { + if offset > bufferSize { + return false + } + return uint64(drawCount) <= (bufferSize-offset)/stride +} + // ExecuteBundle executes a pre-recorded render bundle. func (e *RenderPassEncoder) ExecuteBundle(bundle hal.RenderBundle) { // Note: DX12 bundles use ID3D12GraphicsCommandList created with D3D12_COMMAND_LIST_TYPE_BUNDLE. diff --git a/hal/gles/command.go b/hal/gles/command.go index 9c80a94..76fb568 100644 --- a/hal/gles/command.go +++ b/hal/gles/command.go @@ -650,20 +650,20 @@ func (e *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex ui }) } -// DrawIndirect draws primitives with GPU-generated parameters. -// Note: Requires GL_ARB_draw_indirect (GL 4.0+ / GLES 3.1+). -// Currently not implemented - use direct Draw calls instead. -func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64) { +// DrawIndirect is not implemented by the GLES backend. +func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { _ = buffer _ = offset + _ = drawCount } -// DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. -// Note: Requires GL_ARB_draw_indirect (GL 4.0+ / GLES 3.1+). -// Currently not implemented - use direct DrawIndexed calls instead. -func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64) { +// DrawIndexedIndirect is not implemented by the GLES backend. +// OpenGL cannot apply a WebGPU index-buffer base offset to an indirect record +// without translating that record or the bound index buffer. +func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { _ = buffer _ = offset + _ = drawCount } // ExecuteBundle executes a pre-recorded render bundle. diff --git a/hal/gles/gl/context_linux.go b/hal/gles/gl/context_linux.go index b6c3414..2e9e20f 100644 --- a/hal/gles/gl/context_linux.go +++ b/hal/gles/gl/context_linux.go @@ -172,7 +172,6 @@ func initCommonCallInterfaces() error { if err != nil { return err } - // int32 fn(uint32, void*) err = ffi.PrepareCallInterface(&cifInt322, types.DefaultCall, types.SInt32TypeDescriptor, diff --git a/hal/gles/indirect_test.go b/hal/gles/indirect_test.go new file mode 100644 index 0000000..29c49ea --- /dev/null +++ b/hal/gles/indirect_test.go @@ -0,0 +1,26 @@ +//go:build (windows || linux) && !(js && wasm) + +package gles + +import ( + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +func TestRenderPassEncoderCountedIndirectRemainsUnsupported(t *testing.T) { + enc := &CommandEncoder{} + if err := enc.BeginEncoding("indirect"); err != nil { + t.Fatal(err) + } + pass := enc.BeginRenderPass(&hal.RenderPassDescriptor{ColorAttachments: []hal.RenderPassColorAttachment{}}) + + pass.DrawIndirect(&Buffer{id: 7, size: 64}, 0, 2) + pass.SetIndexBuffer(&Buffer{id: 9, size: 64}, gputypes.IndexFormatUint32, 16) + pass.DrawIndexedIndirect(&Buffer{id: 8, size: 64}, 0, 2) + + if len(enc.commands) != 1 { + t.Fatalf("commands = %d, want only SetIndexBufferCommand", len(enc.commands)) + } +} diff --git a/hal/metal/encoder.go b/hal/metal/encoder.go index 73f1391..5398d21 100644 --- a/hal/metal/encoder.go +++ b/hal/metal/encoder.go @@ -592,24 +592,66 @@ func (e *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex ui } // DrawIndirect draws primitives with GPU-generated parameters. -func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64) { +func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { buf, ok := buffer.(*Buffer) - if !ok || buf == nil { + if !ok || buf == nil || drawCount == 0 { + return + } + if !indirectRangeFits(buf.size, offset, 16, drawCount) { + return + } + if _, ok := indirectRecordOffset(offset, 16, drawCount-1); !ok { return } - _ = MsgSend(e.raw, Sel("drawPrimitives:indirectBuffer:indirectBufferOffset:"), - uintptr(MTLPrimitiveTypeTriangle), uintptr(buf.raw), uintptr(offset)) + for i := uint32(0); i < drawCount; i++ { + recordOffset, _ := indirectRecordOffset(offset, 16, i) + _ = MsgSend(e.raw, Sel("drawPrimitives:indirectBuffer:indirectBufferOffset:"), + uintptr(MTLPrimitiveTypeTriangle), uintptr(buf.raw), uintptr(recordOffset)) + } } // DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. -func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64) { +// Metal exposes only the single-record indirect operation, so count is lowered +// to consecutive 20-byte calls. +func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { buf, ok := buffer.(*Buffer) - if !ok || buf == nil || e.indexBuffer == nil { + if !ok || buf == nil || e.indexBuffer == nil || drawCount == 0 { + return + } + if !indirectRangeFits(buf.size, offset, 20, drawCount) { + return + } + if _, ok := indirectRecordOffset(offset, 20, drawCount-1); !ok { return } indexType := indexFormatToMTL(e.indexFormat) - _ = MsgSend(e.raw, Sel("drawIndexedPrimitives:indexType:indexBuffer:indexBufferOffset:indirectBuffer:indirectBufferOffset:"), - uintptr(MTLPrimitiveTypeTriangle), uintptr(indexType), uintptr(e.indexBuffer.raw), uintptr(e.indexOffset), uintptr(buf.raw), uintptr(offset)) + for i := uint32(0); i < drawCount; i++ { + recordOffset, ok := indirectRecordOffset(offset, 20, i) + if !ok { + return + } + _ = MsgSend(e.raw, Sel("drawIndexedPrimitives:indexType:indexBuffer:indexBufferOffset:indirectBuffer:indirectBufferOffset:"), + uintptr(MTLPrimitiveTypeTriangle), uintptr(indexType), uintptr(e.indexBuffer.raw), uintptr(e.indexOffset), uintptr(buf.raw), uintptr(recordOffset)) + } +} + +func indirectRecordOffset(offset, stride uint64, index uint32) (uint64, bool) { + delta := uint64(index) * stride + if offset > ^uint64(0)-delta { + return 0, false + } + return offset + delta, true +} + +func indirectRangeFits(bufferSize, offset, stride uint64, drawCount uint32) bool { + if offset > bufferSize { + return false + } + return uint64(drawCount) <= (bufferSize-offset)/stride +} + +func indexedIndirectRecordOffset(offset uint64, index uint32) (uint64, bool) { + return indirectRecordOffset(offset, 20, index) } // ExecuteBundle executes a pre-recorded render bundle. diff --git a/hal/metal/indirect_count_test.go b/hal/metal/indirect_count_test.go new file mode 100644 index 0000000..88c7731 --- /dev/null +++ b/hal/metal/indirect_count_test.go @@ -0,0 +1,14 @@ +//go:build darwin && !(js && wasm) + +package metal + +import "testing" + +func TestIndexedIndirectRecordOffset(t *testing.T) { + if got, ok := indexedIndirectRecordOffset(4, 3); !ok || got != 64 { + t.Fatalf("indexedIndirectRecordOffset(4, 3) = %d, %t; want 64, true", got, ok) + } + if _, ok := indexedIndirectRecordOffset(^uint64(0)-3, 1); ok { + t.Fatal("indexedIndirectRecordOffset should reject uint64 overflow") + } +} diff --git a/hal/noop/command.go b/hal/noop/command.go index 5d1cf77..5dee2cc 100644 --- a/hal/noop/command.go +++ b/hal/noop/command.go @@ -102,10 +102,10 @@ func (r *RenderPassEncoder) Draw(_, _, _, _ uint32) {} func (r *RenderPassEncoder) DrawIndexed(_, _, _ uint32, _ int32, _ uint32) {} // DrawIndirect is a no-op. -func (r *RenderPassEncoder) DrawIndirect(_ hal.Buffer, _ uint64) {} +func (r *RenderPassEncoder) DrawIndirect(_ hal.Buffer, _ uint64, _ uint32) {} // DrawIndexedIndirect is a no-op. -func (r *RenderPassEncoder) DrawIndexedIndirect(_ hal.Buffer, _ uint64) {} +func (r *RenderPassEncoder) DrawIndexedIndirect(_ hal.Buffer, _ uint64, _ uint32) {} // ExecuteBundle is a no-op. func (r *RenderPassEncoder) ExecuteBundle(_ hal.RenderBundle) {} diff --git a/hal/noop/noop_test.go b/hal/noop/noop_test.go index c02bb37..e00ebaa 100644 --- a/hal/noop/noop_test.go +++ b/hal/noop/noop_test.go @@ -1562,8 +1562,8 @@ func TestNoopRenderPassEncoder(t *testing.T) { pass.SetStencilReference(0xFF) pass.Draw(6, 1, 0, 0) pass.DrawIndexed(6, 1, 0, 0, 0) - pass.DrawIndirect(buf, 0) - pass.DrawIndexedIndirect(buf, 0) + pass.DrawIndirect(buf, 0, 1) + pass.DrawIndexedIndirect(buf, 0, 1) pass.ExecuteBundle(nil) pass.End() } diff --git a/hal/software/command.go b/hal/software/command.go index 3213cff..71c9093 100644 --- a/hal/software/command.go +++ b/hal/software/command.go @@ -481,10 +481,10 @@ func (r *RenderPassEncoder) drawVertexIndex(firstVertex, pos uint32) uint32 { } // DrawIndirect is a no-op. -func (r *RenderPassEncoder) DrawIndirect(_ hal.Buffer, _ uint64) {} +func (r *RenderPassEncoder) DrawIndirect(_ hal.Buffer, _ uint64, _ uint32) {} // DrawIndexedIndirect is a no-op. -func (r *RenderPassEncoder) DrawIndexedIndirect(_ hal.Buffer, _ uint64) {} +func (r *RenderPassEncoder) DrawIndexedIndirect(_ hal.Buffer, _ uint64, _ uint32) {} // ExecuteBundle is a no-op. func (r *RenderPassEncoder) ExecuteBundle(_ hal.RenderBundle) {} diff --git a/hal/software/software_test.go b/hal/software/software_test.go index 59cf880..df54683 100644 --- a/hal/software/software_test.go +++ b/hal/software/software_test.go @@ -972,8 +972,8 @@ func TestRenderPassEncoderAllNoOps(t *testing.T) { pass.SetStencilReference(0xFF) pass.Draw(3, 1, 0, 0) pass.DrawIndexed(6, 1, 0, 0, 0) - pass.DrawIndirect(buf, 0) - pass.DrawIndexedIndirect(buf, 0) + pass.DrawIndirect(buf, 0, 1) + pass.DrawIndexedIndirect(buf, 0, 1) pass.ExecuteBundle(nil) pass.End() } diff --git a/hal/vulkan/adapter.go b/hal/vulkan/adapter.go index 0a972e3..24b5798 100644 --- a/hal/vulkan/adapter.go +++ b/hal/vulkan/adapter.go @@ -148,6 +148,8 @@ func (a *Adapter) Open(features gputypes.Features, limits gputypes.Limits) (hal. instance: a.instance, graphicsFamily: uint32(graphicsFamily), cmds: &deviceCmds, + supportsMultiDrawIndirect: a.features.MultiDrawIndirect != 0, + maxDrawIndirectCount: a.properties.Limits.MaxDrawIndirectCount, supportsIncrementalPresent: hasIncrementalPresent, } diff --git a/hal/vulkan/command.go b/hal/vulkan/command.go index 15864bf..a64bc05 100644 --- a/hal/vulkan/command.go +++ b/hal/vulkan/command.go @@ -866,6 +866,12 @@ type RenderPassEncoder struct { framebuffer vk.Framebuffer } +const ( + drawIndirectStride = uint32(16) + drawIndexedIndirectStride = uint32(20) + indexedIndirectStride = drawIndexedIndirectStride +) + // End finishes the render pass. // Returns the encoder to the pool for reuse (VK-PERF-006). func (e *RenderPassEncoder) End() { @@ -1030,23 +1036,108 @@ func (e *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex ui } // DrawIndirect draws primitives with GPU-generated parameters. -func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64) { +func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { buf, ok := buffer.(*Buffer) - if !ok || e.encoder.active == 0 { + if !ok || e.encoder.active == 0 || drawCount == 0 { return } - - vkCmdDrawIndirect(e.encoder.device.cmds, e.encoder.active, buf.handle, vk.DeviceSize(offset), 1, 0) + if !indirectRangeFits(buf.size, offset, uint64(drawIndirectStride), drawCount) { + return + } + call, batched, ok := indirectCallPlan(e.encoder.device.supportsMultiDrawIndirect, + e.encoder.device.maxDrawIndirectCount, offset, drawCount, uint32(drawIndirectStride)) + if !ok { + return + } + if batched { + vkCmdDrawIndirect(e.encoder.device.cmds, e.encoder.active, buf.handle, vk.DeviceSize(call.offset), call.count, call.stride) + return + } + for i := uint32(0); i < drawCount; i++ { + recordOffset, _ := indirectRecordOffset(offset, uint64(drawIndirectStride), i) + vkCmdDrawIndirect(e.encoder.device.cmds, e.encoder.active, buf.handle, vk.DeviceSize(recordOffset), 1, drawIndirectStride) + } } // DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. -func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64) { +// Vulkan can encode a span in one command only when the optional +// multiDrawIndirect feature is enabled and the count fits the device limit; +// otherwise emit the exact single-record loop. +func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64, drawCount uint32) { buf, ok := buffer.(*Buffer) - if !ok || e.encoder.active == 0 { + if !ok || e.encoder.active == 0 || drawCount == 0 { + return + } + if !indirectRangeFits(buf.size, offset, uint64(drawIndexedIndirectStride), drawCount) { + return + } + call, batched, ok := indirectCallPlan( + e.encoder.device.supportsMultiDrawIndirect, + e.encoder.device.maxDrawIndirectCount, + offset, + drawCount, + uint32(drawIndexedIndirectStride), + ) + if !ok { + return + } + + if batched { + vkCmdDrawIndexedIndirect(e.encoder.device.cmds, e.encoder.active, buf.handle, vk.DeviceSize(call.offset), call.count, call.stride) return } + for i := uint32(0); i < drawCount; i++ { + recordOffset, ok := indirectRecordOffset(offset, uint64(drawIndexedIndirectStride), i) + if !ok { + return + } + vkCmdDrawIndexedIndirect(e.encoder.device.cmds, e.encoder.active, buf.handle, vk.DeviceSize(recordOffset), call.count, call.stride) + } +} + +type indexedIndirectCall struct { + offset uint64 + count uint32 + stride uint32 +} + +// indexedIndirectCallPlan returns the first native call shape for an indexed +// indirect draw. The plan is pure so count/stride policy can be tested without +// constructing a Vulkan command buffer or invoking FFI. +func indirectCallPlan(supportsMultiDraw bool, maxDrawCount uint32, offset uint64, drawCount, stride uint32) (indexedIndirectCall, bool, bool) { + if drawCount == 0 { + return indexedIndirectCall{}, false, false + } + if _, ok := indirectRecordOffset(offset, uint64(stride), drawCount-1); !ok { + return indexedIndirectCall{}, false, false + } + if supportsMultiDraw && drawCount <= maxDrawCount { + return indexedIndirectCall{offset: offset, count: drawCount, stride: stride}, true, true + } + return indexedIndirectCall{offset: offset, count: 1, stride: stride}, false, true +} + +func indexedIndirectCallPlan(supportsMultiDraw bool, maxDrawCount uint32, offset uint64, drawCount uint32) (indexedIndirectCall, bool, bool) { + return indirectCallPlan(supportsMultiDraw, maxDrawCount, offset, drawCount, uint32(drawIndexedIndirectStride)) +} + +func indirectRecordOffset(offset, stride uint64, index uint32) (uint64, bool) { + delta := uint64(index) * stride + if offset > ^uint64(0)-delta { + return 0, false + } + return offset + delta, true +} + +func indirectRangeFits(bufferSize, offset, stride uint64, drawCount uint32) bool { + if offset > bufferSize { + return false + } + return uint64(drawCount) <= (bufferSize-offset)/stride +} - vkCmdDrawIndexedIndirect(e.encoder.device.cmds, e.encoder.active, buf.handle, vk.DeviceSize(offset), 1, 0) +func indexedIndirectRecordOffset(offset uint64, index uint32) (uint64, bool) { + return indirectRecordOffset(offset, uint64(drawIndexedIndirectStride), index) } // ExecuteBundle executes a pre-recorded render bundle. diff --git a/hal/vulkan/device.go b/hal/vulkan/device.go index 67a4ca5..a0f3f56 100644 --- a/hal/vulkan/device.go +++ b/hal/vulkan/device.go @@ -56,15 +56,17 @@ var renderPassPool = sync.Pool{ // Device implements hal.Device for Vulkan. type Device struct { - handle vk.Device - physicalDevice vk.PhysicalDevice - instance *Instance - graphicsFamily uint32 - allocator *memory.GpuAllocator - cmds *vk.Commands - descriptorAllocator *DescriptorAllocator // Descriptor pool management for bind groups - queue *Queue // Primary queue (for swapchain synchronization) - renderPassCache *RenderPassCache // Cache for VkRenderPass and VkFramebuffer objects + handle vk.Device + physicalDevice vk.PhysicalDevice + instance *Instance + graphicsFamily uint32 + allocator *memory.GpuAllocator + cmds *vk.Commands + supportsMultiDrawIndirect bool + maxDrawIndirectCount uint32 + descriptorAllocator *DescriptorAllocator // Descriptor pool management for bind groups + queue *Queue // Primary queue (for swapchain synchronization) + renderPassCache *RenderPassCache // Cache for VkRenderPass and VkFramebuffer objects // supportsIncrementalPresent is true when VK_KHR_incremental_present // is enabled on this device. When true, Present can chain diff --git a/hal/vulkan/indirect_count_test.go b/hal/vulkan/indirect_count_test.go new file mode 100644 index 0000000..10b2ad8 --- /dev/null +++ b/hal/vulkan/indirect_count_test.go @@ -0,0 +1,107 @@ +//go:build !(js && wasm) + +package vulkan + +import "testing" + +func TestIndexedIndirectCallPlanCapturesCountAndStride(t *testing.T) { + tests := []struct { + name string + supports bool + max uint32 + drawCount uint32 + wantOffset uint64 + wantCount uint32 + wantBatch bool + }{ + {name: "native multi draw", supports: true, max: 8, drawCount: 3, wantOffset: 4, wantCount: 3, wantBatch: true}, + {name: "single draw fallback", supports: false, max: 8, drawCount: 3, wantOffset: 4, wantCount: 1, wantBatch: false}, + {name: "count one", supports: false, max: 0, drawCount: 1, wantOffset: 4, wantCount: 1, wantBatch: false}, + } + + for _, test := range tests { + t.Run(test.name, func(t *testing.T) { + call, batched, ok := indexedIndirectCallPlan(test.supports, test.max, test.wantOffset, test.drawCount) + if !ok { + t.Fatal("indexedIndirectCallPlan rejected a valid range") + } + if batched != test.wantBatch { + t.Fatalf("batched = %t, want %t", batched, test.wantBatch) + } + if call.offset != test.wantOffset || call.count != test.wantCount || call.stride != indexedIndirectStride { + t.Fatalf("captured call = offset %d, count %d, stride %d; want offset %d, count %d, stride %d", call.offset, call.count, call.stride, test.wantOffset, test.wantCount, indexedIndirectStride) + } + }) + } +} + +func TestIndirectCallPlanUsesFixedRecordStrideForBothDrawKinds(t *testing.T) { + for _, test := range []struct { + name string + stride uint32 + }{ + {name: "draw", stride: drawIndirectStride}, + {name: "indexed", stride: drawIndexedIndirectStride}, + } { + t.Run(test.name, func(t *testing.T) { + call, batched, ok := indirectCallPlan(true, 8, 4, 3, test.stride) + if !ok || !batched { + t.Fatalf("indirectCallPlan = %#v, batched=%t, ok=%t", call, batched, ok) + } + if call.offset != 4 || call.count != 3 || call.stride != test.stride { + t.Fatalf("call = %#v, want offset=4 count=3 stride=%d", call, test.stride) + } + }) + } +} + +func TestIndexedIndirectRecordOffset(t *testing.T) { + if got, ok := indexedIndirectRecordOffset(4, 3); !ok || got != 64 { + t.Fatalf("indexedIndirectRecordOffset(4, 3) = %d, %t; want 64, true", got, ok) + } + if _, ok := indexedIndirectRecordOffset(^uint64(0)-3, 1); ok { + t.Fatal("indexedIndirectRecordOffset should reject uint64 overflow") + } +} + +func TestIndexedIndirectHelpersAllocateZero(t *testing.T) { + var ( + offset uint64 + ok bool + ) + if allocs := testing.AllocsPerRun(1000, func() { + offset, ok = indexedIndirectRecordOffset(4, 3) + }); allocs != 0 { + t.Fatalf("indexedIndirectRecordOffset allocations = %v, want 0", allocs) + } + if offset != 64 || !ok { + t.Fatalf("indexedIndirectRecordOffset result = %d, %t; want 64, true", offset, ok) + } + if allocs := testing.AllocsPerRun(1000, func() { + _, _, ok = indexedIndirectCallPlan(true, 8, 4, 3) + }); allocs != 0 { + t.Fatalf("indexedIndirectCallPlan allocations = %v, want 0", allocs) + } +} + +func BenchmarkIndexedIndirectCallPlanCount1(b *testing.B) { + benchmarkIndexedIndirectCallPlan(b, 1) +} + +func BenchmarkIndexedIndirectCallPlanCountN(b *testing.B) { + benchmarkIndexedIndirectCallPlan(b, 1024) +} + +func benchmarkIndexedIndirectCallPlan(b *testing.B, drawCount uint32) { + b.ReportAllocs() + var call indexedIndirectCall + var ok bool + b.ResetTimer() + for i := 0; i < b.N; i++ { + call, _, ok = indexedIndirectCallPlan(true, ^uint32(0), 4, drawCount) + } + b.StopTimer() + if !ok || call.count != drawCount || call.stride != indexedIndirectStride { + b.Fatalf("indexedIndirectCallPlan result = %#v, %t", call, ok) + } +} diff --git a/indirect.go b/indirect.go new file mode 100644 index 0000000..7987a3b --- /dev/null +++ b/indirect.go @@ -0,0 +1,53 @@ +package wgpu + +const ( + drawIndirectRecordSize = uint64(16) + drawIndexedIndirectRecordSize = uint64(20) + indexedIndirectRecordSize = drawIndexedIndirectRecordSize +) + +func indirectRangeFits(bufferSize, offset, recordSize uint64, drawCount uint32) bool { + if offset > bufferSize { + return false + } + return uint64(drawCount) <= (bufferSize-offset)/recordSize +} + +func indirectRecordOffset(offset, recordSize uint64, index uint32) (uint64, bool) { + delta := uint64(index) * recordSize + if offset > ^uint64(0)-delta { + return 0, false + } + return offset + delta, true +} + +func drawIndirectRangeFits(bufferSize, offset uint64, drawCount uint32) bool { + return indirectRangeFits(bufferSize, offset, drawIndirectRecordSize, drawCount) +} + +func drawIndirectRecordOffset(offset uint64, index uint32) (uint64, bool) { + return indirectRecordOffset(offset, drawIndirectRecordSize, index) +} + +func indirectDelegatedValidationOffset(bufferSize, offset, recordSize uint64, drawCount uint32) uint64 { + lastOffset, ok := indirectRecordOffset(offset, recordSize, drawCount-1) + if !ok { + return bufferSize + } + return lastOffset +} + +// indexedIndirectRangeFits reports whether drawCount consecutive indexed +// indirect argument records fit in a buffer without overflowing uint64 math. +func indexedIndirectRangeFits(bufferSize, offset uint64, drawCount uint32) bool { + if offset > bufferSize { + return false + } + return indirectRangeFits(bufferSize, offset, drawIndexedIndirectRecordSize, drawCount) +} + +// indexedIndirectRecordOffset returns the byte offset of one record in a +// counted indexed-indirect span without allowing uint64 wraparound. +func indexedIndirectRecordOffset(offset uint64, index uint32) (uint64, bool) { + return indirectRecordOffset(offset, drawIndexedIndirectRecordSize, index) +} diff --git a/indirect_count_integration_test.go b/indirect_count_integration_test.go new file mode 100644 index 0000000..d1a5a83 --- /dev/null +++ b/indirect_count_integration_test.go @@ -0,0 +1,218 @@ +//go:build integration && darwin && !rust && !(js && wasm) + +package wgpu_test + +import ( + "context" + "encoding/binary" + "math" + "testing" + "time" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu" + _ "github.com/gogpu/wgpu/hal/metal" +) + +const countedIndexedShader = ` +@vertex +fn vs_main(@location(0) pos: vec2) -> @builtin(position) vec4 { + return vec4(pos, 0.0, 1.0); +} + +@fragment +fn fs_main() -> @location(0) vec4 { + return vec4(0.0, 1.0, 0.0, 1.0); +} +` + +func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { + instance, err := wgpu.CreateInstance(&wgpu.InstanceDescriptor{Backends: wgpu.BackendsPrimary}) + if err != nil { + t.Skipf("CreateInstance: %v", err) + } + defer instance.Release() + + adapter, err := instance.RequestAdapter(nil) + if err != nil { + t.Skipf("RequestAdapter: %v", err) + } + defer adapter.Release() + if info := adapter.Info(); info.Backend != gputypes.BackendMetal { + t.Skipf("native Metal adapter unavailable: %+v", info) + } + + device, err := adapter.RequestDevice(nil) + if err != nil { + t.Skipf("RequestDevice: %v", err) + } + defer device.Release() + queue := device.Queue() + + vertices := []float32{-1, -1, 1, -1, 1, 1, -1, 1} + vertexData := make([]byte, len(vertices)*4) + for i, value := range vertices { + binary.LittleEndian.PutUint32(vertexData[i*4:], math.Float32bits(value)) + } + indexData := []byte{0, 0, 1, 0, 2, 0, 2, 0, 3, 0, 0, 0} + + // Four padding bytes make the starting offset observable. The two records + // select different triangles; both are required to fill the target. + indirectData := make([]byte, 4+2*20) + putCountedIndexedIndirectRecord(indirectData[4:], 3, 1, 0, 0, 0) + putCountedIndexedIndirectRecord(indirectData[24:], 3, 1, 3, 0, 0) + + vertexBuffer := createCountedIndirectTestBuffer(t, device, queue, vertexData, wgpu.BufferUsageVertex|wgpu.BufferUsageCopyDst) + defer vertexBuffer.Release() + indexBuffer := createCountedIndirectTestBuffer(t, device, queue, indexData, wgpu.BufferUsageIndex|wgpu.BufferUsageCopyDst) + defer indexBuffer.Release() + indirectBuffer := createCountedIndirectTestBuffer(t, device, queue, indirectData, wgpu.BufferUsageIndirect|wgpu.BufferUsageCopyDst) + defer indirectBuffer.Release() + + shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{WGSL: countedIndexedShader}) + if err != nil { + t.Fatalf("CreateShaderModule: %v", err) + } + defer shader.Release() + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{}) + if err != nil { + t.Fatalf("CreatePipelineLayout: %v", err) + } + defer layout.Release() + pipeline, err := device.CreateRenderPipeline(&wgpu.RenderPipelineDescriptor{ + Layout: layout, + Vertex: wgpu.VertexState{ + Module: shader, EntryPoint: "vs_main", + Buffers: []gputypes.VertexBufferLayout{{ + ArrayStride: 8, + StepMode: gputypes.VertexStepModeVertex, + Attributes: []gputypes.VertexAttribute{{ + Format: gputypes.VertexFormatFloat32x2, ShaderLocation: 0, + }}, + }}, + }, + Fragment: &wgpu.FragmentState{ + Module: shader, EntryPoint: "fs_main", + Targets: []gputypes.ColorTargetState{{ + Format: gputypes.TextureFormatRGBA8Unorm, WriteMask: gputypes.ColorWriteMaskAll, + }}, + }, + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList, CullMode: gputypes.CullModeNone}, + Multisample: gputypes.MultisampleState{Count: 1, Mask: 0xFFFFFFFF}, + }) + if err != nil { + t.Fatalf("CreateRenderPipeline: %v", err) + } + defer pipeline.Release() + + const width, height = uint32(64), uint32(16) + target, err := device.CreateTexture(&wgpu.TextureDescriptor{ + Size: wgpu.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: %v", err) + } + defer target.Release() + targetView, err := device.CreateTextureView(target, nil) + if err != nil { + t.Fatalf("CreateTextureView: %v", err) + } + defer targetView.Release() + + encoder, err := device.CreateCommandEncoder(&wgpu.CommandEncoderDescriptor{}) + if err != nil { + t.Fatalf("CreateCommandEncoder: %v", err) + } + pass, err := encoder.BeginRenderPass(&wgpu.RenderPassDescriptor{ + ColorAttachments: []wgpu.RenderPassColorAttachment{{ + View: targetView, LoadOp: gputypes.LoadOpClear, StoreOp: gputypes.StoreOpStore, + ClearValue: gputypes.Color{A: 1}, + }}, + }) + if err != nil { + t.Fatalf("BeginRenderPass: %v", err) + } + pass.SetPipeline(pipeline) + pass.SetVertexBuffer(0, vertexBuffer, 0) + pass.SetIndexBuffer(indexBuffer, gputypes.IndexFormatUint16, 0) + pass.MultiDrawIndexedIndirect(indirectBuffer, 4, 2) + if err := pass.End(); err != nil { + t.Fatalf("RenderPass.End: %v", err) + } + encoder.TransitionTextures([]wgpu.TextureBarrier{{ + Texture: target, + Usage: wgpu.TextureUsageTransition{ + OldUsage: gputypes.TextureUsageRenderAttachment, + NewUsage: gputypes.TextureUsageCopySrc, + }, + }}) + + readbackSize := uint64(width * height * 4) + readback, err := device.CreateBuffer(&wgpu.BufferDescriptor{ + Size: readbackSize, Usage: gputypes.BufferUsageMapRead | gputypes.BufferUsageCopyDst, + }) + if err != nil { + t.Fatalf("CreateBuffer(readback): %v", err) + } + defer readback.Release() + encoder.CopyTextureToBuffer(target, readback, []wgpu.BufferTextureCopy{{ + BufferLayout: wgpu.ImageDataLayout{BytesPerRow: width * 4, RowsPerImage: height}, + TextureBase: wgpu.ImageCopyTexture{Texture: target}, + Size: wgpu.Extent3D{Width: width, Height: height, DepthOrArrayLayers: 1}, + }}) + + commands, err := encoder.Finish() + if err != nil { + t.Fatalf("Finish: %v", err) + } + if _, err := queue.Submit(commands); err != nil { + t.Fatalf("Submit: %v", err) + } + mapContext, cancelMap := context.WithTimeout(context.Background(), 5*time.Second) + defer cancelMap() + if err := readback.Map(mapContext, wgpu.MapModeRead, 0, readbackSize); err != nil { + t.Fatalf("Map: %v", err) + } + mapped, err := readback.MappedRange(0, readbackSize) + if err != nil { + _ = readback.Unmap() + t.Fatalf("MappedRange: %v", err) + } + pixels := mapped.Bytes() + filled := 0 + for pixel := 0; pixel < int(width*height); pixel++ { + if pixels[pixel*4+1] > 128 { + filled++ + } + } + if err := readback.Unmap(); err != nil { + t.Fatalf("Unmap: %v", err) + } + if minimum := int(width*height) * 3 / 4; filled < minimum { + t.Fatalf("counted indexed indirect filled %d/%d pixels, want at least %d; both records did not render", filled, width*height, minimum) + } +} + +func createCountedIndirectTestBuffer(t *testing.T, device *wgpu.Device, queue *wgpu.Queue, data []byte, usage wgpu.BufferUsage) *wgpu.Buffer { + t.Helper() + buffer, err := device.CreateBuffer(&wgpu.BufferDescriptor{Size: uint64(len(data)), Usage: usage}) + if err != nil { + t.Fatalf("CreateBuffer: %v", err) + } + if err := queue.WriteBuffer(buffer, 0, data); err != nil { + buffer.Release() + t.Fatalf("WriteBuffer: %v", err) + } + return buffer +} + +func putCountedIndexedIndirectRecord(dst []byte, indexCount, instanceCount, firstIndex uint32, baseVertex int32, firstInstance uint32) { + binary.LittleEndian.PutUint32(dst[0:], indexCount) + binary.LittleEndian.PutUint32(dst[4:], instanceCount) + binary.LittleEndian.PutUint32(dst[8:], firstIndex) + binary.LittleEndian.PutUint32(dst[12:], uint32(baseVertex)) + binary.LittleEndian.PutUint32(dst[16:], firstInstance) +} diff --git a/internal/browser/renderpass.go b/internal/browser/renderpass.go index dddc4a5..18c2be0 100644 --- a/internal/browser/renderpass.go +++ b/internal/browser/renderpass.go @@ -118,7 +118,7 @@ func (p *RenderPassEncoder) DrawIndirect(buffer js.Value, offset uint64) { p.fnDrawIndirect.Invoke(buffer, float64(offset)) } -// DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. +// DrawIndexedIndirect draws one indexed primitive with GPU-generated parameters. func (p *RenderPassEncoder) DrawIndexedIndirect(buffer js.Value, offset uint64) { p.fnDrawIndexedIndirect.Invoke(buffer, float64(offset)) } diff --git a/internal/browser/renderpass_indirect_test.go b/internal/browser/renderpass_indirect_test.go new file mode 100644 index 0000000..2190f8c --- /dev/null +++ b/internal/browser/renderpass_indirect_test.go @@ -0,0 +1,56 @@ +//go:build js && wasm + +package browser + +import ( + "syscall/js" + "testing" +) + +func TestDrawIndexedIndirectOverflowDelegatesOneValidationCall(t *testing.T) { + var offsets []float64 + draw := js.FuncOf(func(_ js.Value, args []js.Value) any { + offsets = append(offsets, args[1].Float()) + return nil + }) + defer draw.Release() + + pass := &RenderPassEncoder{fnDrawIndexedIndirect: draw.Value} + pass.DrawIndexedIndirect(js.Undefined(), ^uint64(0)-3) + + if len(offsets) != 1 { + t.Fatalf("drawIndexedIndirect calls = %d, want one delegated validation call", len(offsets)) + } +} + +func TestDrawIndexedIndirectCountOnePreservesOffset(t *testing.T) { + var offsets []float64 + draw := js.FuncOf(func(_ js.Value, args []js.Value) any { + offsets = append(offsets, args[1].Float()) + return nil + }) + defer draw.Release() + + pass := &RenderPassEncoder{fnDrawIndexedIndirect: draw.Value} + pass.DrawIndexedIndirect(js.Undefined(), 8) + + if len(offsets) != 1 || offsets[0] != 8 { + t.Fatalf("drawIndexedIndirect offsets = %v, want [8]", offsets) + } +} + +func TestDrawIndexedIndirectSingleRecordCallsOnce(t *testing.T) { + var calls int + draw := js.FuncOf(func(_ js.Value, _ []js.Value) any { + calls++ + return nil + }) + defer draw.Release() + + pass := &RenderPassEncoder{fnDrawIndexedIndirect: draw.Value} + pass.DrawIndexedIndirect(js.Undefined(), 8) + + if calls != 1 { + t.Fatalf("drawIndexedIndirect calls = %d, want one", calls) + } +} diff --git a/renderpass_browser.go b/renderpass_browser.go index 8c7f4c3..c73c2e1 100644 --- a/renderpass_browser.go +++ b/renderpass_browser.go @@ -84,18 +84,49 @@ func (p *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex ui // DrawIndirect draws primitives with GPU-generated parameters. func (p *RenderPassEncoder) DrawIndirect(buffer *Buffer, offset uint64) { + p.MultiDrawIndirect(buffer, offset, 1) +} + +// MultiDrawIndirect draws consecutive primitives with GPU-generated parameters. +func (p *RenderPassEncoder) MultiDrawIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if drawCount == 0 { + return + } if buffer == nil || buffer.browser == nil { return } - p.browser.DrawIndirect(buffer.browser.Ref(), offset) + if !drawIndirectRangeFits(buffer.Size(), offset, drawCount) { + p.browser.DrawIndirect(buffer.browser.Ref(), indirectDelegatedValidationOffset(buffer.Size(), offset, drawIndirectRecordSize, drawCount)) + return + } + for i := uint32(0); i < drawCount; i++ { + recordOffset, _ := drawIndirectRecordOffset(offset, i) + p.browser.DrawIndirect(buffer.browser.Ref(), recordOffset) + } } // DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. func (p *RenderPassEncoder) DrawIndexedIndirect(buffer *Buffer, offset uint64) { + p.MultiDrawIndexedIndirect(buffer, offset, 1) +} + +// MultiDrawIndexedIndirect draws consecutive indexed primitives with +// GPU-generated parameters. +func (p *RenderPassEncoder) MultiDrawIndexedIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if drawCount == 0 { + return + } if buffer == nil || buffer.browser == nil { return } - p.browser.DrawIndexedIndirect(buffer.browser.Ref(), offset) + if !indexedIndirectRangeFits(buffer.Size(), offset, drawCount) { + p.browser.DrawIndexedIndirect(buffer.browser.Ref(), indirectDelegatedValidationOffset(buffer.Size(), offset, drawIndexedIndirectRecordSize, drawCount)) + return + } + for i := uint32(0); i < drawCount; i++ { + recordOffset, _ := indexedIndirectRecordOffset(offset, i) + p.browser.DrawIndexedIndirect(buffer.browser.Ref(), recordOffset) + } } // End ends the render pass. diff --git a/renderpass_indirect_browser_test.go b/renderpass_indirect_browser_test.go new file mode 100644 index 0000000..f144cd2 --- /dev/null +++ b/renderpass_indirect_browser_test.go @@ -0,0 +1,54 @@ +//go:build js && wasm + +package wgpu + +import ( + "slices" + "syscall/js" + "testing" + + "github.com/gogpu/wgpu/internal/browser" +) + +func TestBrowserDrawIndexedIndirectInvalidSpanDelegatesOneFailingCall(t *testing.T) { + tests := []struct { + name string + bufferSize uint64 + offset uint64 + drawCount uint32 + want []float64 + }{ + {name: "later record fails", bufferSize: 59, offset: 0, drawCount: 3, want: []float64{40}}, + {name: "count one preserves offset", bufferSize: 20, offset: 24, drawCount: 1, want: []float64{24}}, + {name: "arithmetic overflow", bufferSize: 64, offset: ^uint64(0) - 3, drawCount: 2, want: []float64{64}}, + } + + for _, test := range tests { + t.Run(test.name, func(t *testing.T) { + var offsets []float64 + draw := js.FuncOf(func(_ js.Value, args []js.Value) any { + offsets = append(offsets, args[1].Float()) + return nil + }) + defer draw.Release() + + rawPass := js.Global().Get("Object").New() + rawPass.Set("drawIndexedIndirect", draw) + pass := &RenderPassEncoder{browser: browser.NewRenderPassEncoder(rawPass)} + + rawBuffer := js.Global().Get("Object").New() + rawBuffer.Set("size", float64(test.bufferSize)) + rawBuffer.Set("usage", 0) + buffer := &Buffer{ + browser: browser.NewBuffer(rawBuffer), + size: test.bufferSize, + } + + pass.MultiDrawIndexedIndirect(buffer, test.offset, test.drawCount) + + if !slices.Equal(offsets, test.want) { + t.Fatalf("underlying offsets = %v, want %v", offsets, test.want) + } + }) + } +} diff --git a/renderpass_indirect_rust_test.go b/renderpass_indirect_rust_test.go new file mode 100644 index 0000000..53a3686 --- /dev/null +++ b/renderpass_indirect_rust_test.go @@ -0,0 +1,46 @@ +//go:build rust + +package wgpu + +import ( + "slices" + "testing" +) + +func TestLowerRustIndexedIndirectInvalidSpanDelegatesOneFailingCall(t *testing.T) { + tests := []struct { + name string + bufferSize uint64 + offset uint64 + drawCount uint32 + want []uint64 + }{ + {name: "later record fails", bufferSize: 59, offset: 0, drawCount: 3, want: []uint64{40}}, + {name: "count one preserves offset", bufferSize: 20, offset: 24, drawCount: 1, want: []uint64{24}}, + {name: "arithmetic overflow", bufferSize: 64, offset: ^uint64(0) - 3, drawCount: 2, want: []uint64{64}}, + } + + for _, test := range tests { + t.Run(test.name, func(t *testing.T) { + var offsets []uint64 + lowerRustIndexedIndirect(test.bufferSize, test.offset, test.drawCount, func(offset uint64) { + offsets = append(offsets, offset) + }) + + if !slices.Equal(offsets, test.want) { + t.Fatalf("underlying offsets = %v, want %v", offsets, test.want) + } + }) + } +} + +func TestLowerRustIndexedIndirectValidSpanEmitsEveryRecordInOrder(t *testing.T) { + var offsets []uint64 + lowerRustIndexedIndirect(64, 4, 3, func(offset uint64) { + offsets = append(offsets, offset) + }) + + if want := []uint64{4, 24, 44}; !slices.Equal(offsets, want) { + t.Fatalf("underlying offsets = %v, want %v", offsets, want) + } +} diff --git a/renderpass_indirect_test.go b/renderpass_indirect_test.go new file mode 100644 index 0000000..a87b542 --- /dev/null +++ b/renderpass_indirect_test.go @@ -0,0 +1,77 @@ +//go:build !rust && !(js && wasm) + +package wgpu + +import "testing" + +func TestIndexedIndirectRangeFits(t *testing.T) { + tests := []struct { + name string + bufferSize uint64 + offset uint64 + drawCount uint32 + want bool + }{ + {name: "one record", bufferSize: 20, offset: 0, drawCount: 1, want: true}, + {name: "three records", bufferSize: 60, offset: 0, drawCount: 3, want: true}, + {name: "tail records", bufferSize: 80, offset: 20, drawCount: 3, want: true}, + {name: "one byte short", bufferSize: 59, offset: 0, drawCount: 3, want: false}, + {name: "offset past end", bufferSize: 20, offset: 21, drawCount: 1, want: false}, + {name: "offset overflow", bufferSize: ^uint64(0), offset: ^uint64(0) - 3, drawCount: 1, want: false}, + {name: "count overflow", bufferSize: ^uint64(0), offset: 0, drawCount: ^uint32(0), want: true}, + } + + for _, test := range tests { + t.Run(test.name, func(t *testing.T) { + if got := indexedIndirectRangeFits(test.bufferSize, test.offset, test.drawCount); got != test.want { + t.Fatalf("indexedIndirectRangeFits(%d, %d, %d) = %t, want %t", test.bufferSize, test.offset, test.drawCount, got, test.want) + } + }) + } +} + +func TestIndexedIndirectHelpersAllocateZero(t *testing.T) { + var ( + fits bool + offset uint64 + ok bool + ) + if allocs := testing.AllocsPerRun(1000, func() { + fits = indexedIndirectRangeFits(64, 4, 3) + }); allocs != 0 { + t.Fatalf("indexedIndirectRangeFits allocations = %v, want 0", allocs) + } + if !fits { + t.Fatal("indexedIndirectRangeFits rejected a valid range") + } + if allocs := testing.AllocsPerRun(1000, func() { + offset, ok = indexedIndirectRecordOffset(4, 2) + }); allocs != 0 { + t.Fatalf("indexedIndirectRecordOffset allocations = %v, want 0", allocs) + } + if offset != 44 || !ok { + t.Fatalf("indexedIndirectRecordOffset result = %d, %t; want 44, true", offset, ok) + } +} + +func BenchmarkIndexedIndirectRangeFitsCount1(b *testing.B) { + benchmarkIndexedIndirectRangeFits(b, 1) +} + +func BenchmarkIndexedIndirectRangeFitsCountN(b *testing.B) { + benchmarkIndexedIndirectRangeFits(b, 1024) +} + +func benchmarkIndexedIndirectRangeFits(b *testing.B, drawCount uint32) { + b.ReportAllocs() + var fits bool + bufferSize := uint64(drawCount)*indexedIndirectRecordSize + 4 + b.ResetTimer() + for i := 0; i < b.N; i++ { + fits = indexedIndirectRangeFits(bufferSize, 4, drawCount) + } + b.StopTimer() + if !fits { + b.Fatal("indexedIndirectRangeFits rejected a valid range") + } +} diff --git a/renderpass_indirect_wrappers_test.go b/renderpass_indirect_wrappers_test.go new file mode 100644 index 0000000..eef1ef5 --- /dev/null +++ b/renderpass_indirect_wrappers_test.go @@ -0,0 +1,45 @@ +//go:build rust || (js && wasm) + +package wgpu + +import "testing" + +func TestIndexedIndirectDelegatedValidationOffset(t *testing.T) { + tests := []struct { + name string + bufferSize uint64 + offset uint64 + drawCount uint32 + want uint64 + }{ + { + name: "count one preserves caller offset", + bufferSize: 20, + offset: 24, + drawCount: 1, + want: 24, + }, + { + name: "later record delegates only failing offset", + bufferSize: 59, + offset: 0, + drawCount: 3, + want: 40, + }, + { + name: "arithmetic overflow uses end of buffer", + bufferSize: 64, + offset: ^uint64(0) - 3, + drawCount: 2, + want: 64, + }, + } + + for _, test := range tests { + t.Run(test.name, func(t *testing.T) { + if got := indirectDelegatedValidationOffset(test.bufferSize, test.offset, drawIndexedIndirectRecordSize, test.drawCount); got != test.want { + t.Fatalf("indexedIndirectDelegatedValidationOffset(%d, %d, %d) = %d, want %d", test.bufferSize, test.offset, test.drawCount, got, test.want) + } + }) + } +} diff --git a/renderpass_native.go b/renderpass_native.go index f3e50a4..48cf720 100644 --- a/renderpass_native.go +++ b/renderpass_native.go @@ -246,6 +246,15 @@ func (p *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex ui // DrawIndirect draws primitives with GPU-generated parameters. func (p *RenderPassEncoder) DrawIndirect(buffer *Buffer, offset uint64) { + p.MultiDrawIndirect(buffer, offset, 1) +} + +// MultiDrawIndirect draws consecutive primitives with GPU-generated parameters. +// Each argument record is 16 bytes. +func (p *RenderPassEncoder) MultiDrawIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if drawCount == 0 { + return + } if !p.validateDrawState("DrawIndirect") { return } @@ -271,19 +280,28 @@ func (p *RenderPassEncoder) DrawIndirect(buffer *Buffer, offset uint64) { } // VAL-B3: Validate indirect args fit within buffer. // DrawIndirect args: 4 × uint32 = 16 bytes. Matches Rust render.rs:2772-2779. - if offset+16 > buffer.Size() { + if !drawIndirectRangeFits(buffer.Size(), offset, drawCount) { p.encoder.setError(fmt.Errorf( - "wgpu: RenderPass.DrawIndirect: offset %d + 16 bytes exceeds buffer size %d: %w", - offset, buffer.Size(), ErrDrawIndirectBufferOverrun)) + "wgpu: RenderPass.MultiDrawIndirect: offset %d + %d draw(s) exceeds buffer size %d: %w", + offset, drawCount, buffer.Size(), ErrDrawIndirectBufferOverrun)) return } p.trackRef(buffer.core.Ref) p.encoder.trackBuffer(buffer) - p.core.DrawIndirect(buffer.coreBuffer(), offset) + p.core.MultiDrawIndirect(buffer.coreBuffer(), offset, drawCount) } // DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. func (p *RenderPassEncoder) DrawIndexedIndirect(buffer *Buffer, offset uint64) { + p.MultiDrawIndexedIndirect(buffer, offset, 1) +} + +// MultiDrawIndexedIndirect draws consecutive indexed primitives with +// GPU-generated parameters. Each argument record is 20 bytes. +func (p *RenderPassEncoder) MultiDrawIndexedIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if drawCount == 0 { + return + } if !p.validateDrawState("DrawIndexedIndirect") { return } @@ -320,17 +338,17 @@ func (p *RenderPassEncoder) DrawIndexedIndirect(buffer *Buffer, offset uint64) { offset, ErrDrawIndirectOffsetAlignment)) return } - // VAL-B3: Validate indirect args fit within buffer. - // DrawIndexedIndirect args: 5 × uint32 = 20 bytes. Matches Rust render.rs:2772-2779. - if offset+20 > buffer.Size() { + // VAL-B3: Validate all indirect args fit within the buffer without allowing + // offset+drawCount*20 to wrap around uint64. + if !indexedIndirectRangeFits(buffer.Size(), offset, drawCount) { p.encoder.setError(fmt.Errorf( - "wgpu: RenderPass.DrawIndexedIndirect: offset %d + 20 bytes exceeds buffer size %d: %w", - offset, buffer.Size(), ErrDrawIndirectBufferOverrun)) + "wgpu: RenderPass.MultiDrawIndexedIndirect: offset %d + %d draw(s) exceeds buffer size %d: %w", + offset, drawCount, buffer.Size(), ErrDrawIndirectBufferOverrun)) return } p.trackRef(buffer.core.Ref) p.encoder.trackBuffer(buffer) - p.core.DrawIndexedIndirect(buffer.coreBuffer(), offset) + p.core.MultiDrawIndexedIndirect(buffer.coreBuffer(), offset, drawCount) } // End ends the render pass. diff --git a/renderpass_rust.go b/renderpass_rust.go index f973847..80f0f3e 100644 --- a/renderpass_rust.go +++ b/renderpass_rust.go @@ -90,18 +90,61 @@ func (p *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex ui // DrawIndirect draws primitives with GPU-generated parameters. func (p *RenderPassEncoder) DrawIndirect(buffer *Buffer, offset uint64) { + p.MultiDrawIndirect(buffer, offset, 1) +} + +// MultiDrawIndirect draws consecutive primitives with GPU-generated parameters. +func (p *RenderPassEncoder) MultiDrawIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if drawCount == 0 { + return + } if buffer == nil || buffer.r == nil { return } - p.r.DrawIndirect(buffer.r, offset) + if !drawIndirectRangeFits(buffer.Size(), offset, drawCount) { + p.r.DrawIndirect(buffer.r, indirectDelegatedValidationOffset(buffer.Size(), offset, drawIndirectRecordSize, drawCount)) + return + } + for i := uint32(0); i < drawCount; i++ { + recordOffset, _ := drawIndirectRecordOffset(offset, i) + p.r.DrawIndirect(buffer.r, recordOffset) + } } // DrawIndexedIndirect draws indexed primitives with GPU-generated parameters. func (p *RenderPassEncoder) DrawIndexedIndirect(buffer *Buffer, offset uint64) { + p.MultiDrawIndexedIndirect(buffer, offset, 1) +} + +// MultiDrawIndexedIndirect draws consecutive indexed primitives with +// GPU-generated parameters. +func (p *RenderPassEncoder) MultiDrawIndexedIndirect(buffer *Buffer, offset uint64, drawCount uint32) { + if drawCount == 0 { + return + } if buffer == nil || buffer.r == nil { return } - p.r.DrawIndexedIndirect(buffer.r, offset) + lowerRustIndexedIndirect(buffer.Size(), offset, drawCount, func(recordOffset uint64) { + p.r.DrawIndexedIndirect(buffer.r, recordOffset) + }) +} + +// lowerRustIndexedIndirect lowers one counted span through the Rust adapter's +// single-record interface. Invalid positive spans delegate exactly one failing +// record before any valid record can be emitted. +func lowerRustIndexedIndirect(bufferSize, offset uint64, drawCount uint32, draw func(uint64)) { + if drawCount == 0 { + return + } + if !indexedIndirectRangeFits(bufferSize, offset, drawCount) { + draw(indirectDelegatedValidationOffset(bufferSize, offset, drawIndexedIndirectRecordSize, drawCount)) + return + } + for i := uint32(0); i < drawCount; i++ { + recordOffset, _ := indexedIndirectRecordOffset(offset, i) + draw(recordOffset) + } } // End ends the render pass. diff --git a/wgpu_test.go b/wgpu_test.go index 8ca789f..552236c 100644 --- a/wgpu_test.go +++ b/wgpu_test.go @@ -1663,6 +1663,147 @@ func TestRenderPassDrawIndexedIndirectNilDeferredError(t *testing.T) { } } +func TestRenderPassDrawIndexedIndirectZeroCountIsNoOp(t *testing.T) { + device, encoder, pass := newEncoderWithRenderPass(t) + defer device.Release() + + // Zero count is intentionally checked before state and buffer validation. + pass.MultiDrawIndexedIndirect(nil, ^uint64(0)-3, 0) + if err := pass.End(); err != nil { + t.Fatalf("End: %v", err) + } + if _, err := encoder.Finish(); err != nil { + t.Fatalf("zero-count DrawIndexedIndirect recorded an error: %v", err) + } +} + +func TestRenderPassMultiDrawIndirectZeroCountIsNoOp(t *testing.T) { + device, encoder, pass := newEncoderWithRenderPass(t) + defer device.Release() + + pass.MultiDrawIndirect(nil, ^uint64(0)-3, 0) + if err := pass.End(); err != nil { + t.Fatalf("End: %v", err) + } + if _, err := encoder.Finish(); err != nil { + t.Fatalf("zero-count MultiDrawIndirect recorded an error: %v", err) + } +} + +func TestRenderPassMultiDrawIndirectRangeValidation(t *testing.T) { + device, encoder, pass := newEncoderWithRenderPass(t) + defer device.Release() + + pipeline := &wgpu.RenderPipeline{} + pipeline.SetTestRequiredVertexBuffers(0) + pass.SetPipeline(pipeline) + buf, err := device.CreateBuffer(&wgpu.BufferDescriptor{ + Label: "count-indirect-buf", + Size: 31, + Usage: wgpu.BufferUsageIndirect, + }) + if err != nil { + t.Fatalf("CreateBuffer: %v", err) + } + defer buf.Release() + + pass.MultiDrawIndirect(buf, 0, 2) + _ = pass.End() + if _, err := encoder.Finish(); !errors.Is(err, wgpu.ErrDrawIndirectBufferOverrun) { + t.Fatalf("Finish error = %v, want ErrDrawIndirectBufferOverrun", err) + } +} + +func TestRenderPassDrawIndexedIndirectCountRangeValidation(t *testing.T) { + tests := []struct { + name string + bufferSize uint64 + offset uint64 + drawCount uint32 + }{ + {name: "count exceeds buffer", bufferSize: 59, offset: 0, drawCount: 3}, + {name: "offset arithmetic overflow", bufferSize: 64, offset: ^uint64(0) - 3, drawCount: 1}, + } + + for _, test := range tests { + t.Run(test.name, func(t *testing.T) { + device, encoder, pass := newEncoderWithRenderPass(t) + defer device.Release() + + pipeline := &wgpu.RenderPipeline{} + pipeline.SetTestRequiredVertexBuffers(0) + pass.SetPipeline(pipeline) + + idxBuf, err := device.CreateBuffer(&wgpu.BufferDescriptor{ + Label: "count-index-buffer", + Size: 64, + Usage: wgpu.BufferUsageIndex, + }) + if err != nil { + t.Fatalf("CreateBuffer(index): %v", err) + } + defer idxBuf.Release() + pass.SetIndexBuffer(idxBuf, gputypes.IndexFormatUint16, 0) + + indirectBuf, err := device.CreateBuffer(&wgpu.BufferDescriptor{ + Label: "count-indirect-buffer", + Size: test.bufferSize, + Usage: wgpu.BufferUsageIndirect, + }) + if err != nil { + t.Fatalf("CreateBuffer(indirect): %v", err) + } + defer indirectBuf.Release() + + pass.MultiDrawIndexedIndirect(indirectBuf, test.offset, test.drawCount) + _ = pass.End() + _, finishErr := encoder.Finish() + if finishErr == nil { + t.Fatal("Finish() should report an indirect buffer overrun") + } + if !errors.Is(finishErr, wgpu.ErrDrawIndirectBufferOverrun) { + t.Fatalf("error = %v, want ErrDrawIndirectBufferOverrun", finishErr) + } + }) + } +} + +func TestRenderPassDrawIndexedIndirectCountMany(t *testing.T) { + device, encoder, pass := newEncoderWithRenderPass(t) + defer device.Release() + + pipeline := &wgpu.RenderPipeline{} + pipeline.SetTestRequiredVertexBuffers(0) + pass.SetPipeline(pipeline) + + idxBuf, err := device.CreateBuffer(&wgpu.BufferDescriptor{ + Label: "count-index-buffer", + Size: 64, + Usage: wgpu.BufferUsageIndex, + }) + if err != nil { + t.Fatalf("CreateBuffer(index): %v", err) + } + defer idxBuf.Release() + pass.SetIndexBuffer(idxBuf, gputypes.IndexFormatUint16, 0) + + indirectBuf, err := device.CreateBuffer(&wgpu.BufferDescriptor{ + Label: "count-indirect-buffer", + Size: 60, + Usage: wgpu.BufferUsageIndirect, + }) + if err != nil { + t.Fatalf("CreateBuffer(indirect): %v", err) + } + defer indirectBuf.Release() + + pass.MultiDrawIndexedIndirect(indirectBuf, 0, 3) + _ = pass.End() + if _, err := encoder.Finish(); err != nil { + t.Fatalf("counted DrawIndexedIndirect failed: %v", err) + } +} + func TestComputePassSetPipelineNilDeferredError(t *testing.T) { device, encoder, pass := newEncoderWithComputePass(t) defer device.Release() From 3bbeb6ceab20cb7779e01179449adc47a2398123 Mon Sep 17 00:00:00 2001 From: Mark Date: Wed, 15 Jul 2026 14:39:21 +0300 Subject: [PATCH 3/3] feat(metal): lower large indexed multi-draws to ICBs --- hal/metal/device.go | 39 ++- hal/metal/encoder.go | 321 +++++++++++++++++++--- hal/metal/encoder_test.go | 88 ++++++ hal/metal/icb_indexed.go | 420 +++++++++++++++++++++++++++++ hal/metal/icb_indexed_test.go | 114 ++++++++ hal/metal/icb_pipeline.go | 60 +++++ hal/metal/icb_pipeline_test.go | 115 ++++++++ hal/metal/metal.go | 1 + hal/metal/resource.go | 11 +- hal/metal/types.go | 7 + indirect_count_integration_test.go | 252 +++++++++++++++-- metal_icb_benchmark_test.go | 190 +++++++++++++ 12 files changed, 1549 insertions(+), 69 deletions(-) create mode 100644 hal/metal/icb_indexed.go create mode 100644 hal/metal/icb_indexed_test.go create mode 100644 hal/metal/icb_pipeline.go create mode 100644 hal/metal/icb_pipeline_test.go create mode 100644 metal_icb_benchmark_test.go diff --git a/hal/metal/device.go b/hal/metal/device.go index 66cabfa..6bc6dce 100644 --- a/hal/metal/device.go +++ b/hal/metal/device.go @@ -9,6 +9,7 @@ import ( "fmt" "runtime" "strings" + "sync" "time" "unsafe" @@ -38,6 +39,8 @@ type Device struct { // and direct CPU writes without a staging blit, which eliminates the main source // of resize-induced memory growth. hasUnifiedMemory bool + icbTranslatorMu sync.Mutex + icbTranslators map[gputypes.IndexFormat]indexedICBTranslator } // newDevice creates a new Device from a Metal device. @@ -755,10 +758,23 @@ func (d *Device) CreateRenderPipeline(desc *hal.RenderPipelineDescriptor) (hal.R } _ = MsgSend(pipelineDesc, Sel("setSampleCount:"), uintptr(sampleCount)) - // Create pipeline state + // Create pipeline state. ICB support stays entirely private: eligible + // pipelines get one flagged attempt and transparently retry ordinary + // creation if Metal rejects the stricter descriptor. var errorPtr ID - pipelineState := MsgSend(d.raw, Sel("newRenderPipelineStateWithDescriptor:error:"), - uintptr(pipelineDesc), uintptr(unsafe.Pointer(&errorPtr))) + icbCandidate := d.canCreateICBPipeline(desc, pipelineDesc) + pipelineState, icbCompatible := createRenderPipelineState(icbCandidate, func(icb bool) ID { + errorPtr = 0 + if icbCandidate { + var enabled uintptr + if icb { + enabled = 1 + } + _ = MsgSend(pipelineDesc, Sel("setSupportIndirectCommandBuffers:"), enabled) + } + return MsgSend(d.raw, Sel("newRenderPipelineStateWithDescriptor:error:"), + uintptr(pipelineDesc), uintptr(unsafe.Pointer(&errorPtr))) + }) if pipelineState == 0 { errMsg := unknownError @@ -783,11 +799,12 @@ func (d *Device) CreateRenderPipeline(desc *hal.RenderPipelineDescriptor) (hal.R pipeLayout = pl } return &RenderPipeline{ - raw: pipelineState, - device: d, - layout: pipeLayout, - cullMode: cullModeToMTL(desc.Primitive.CullMode), - frontFace: frontFaceToMTL(desc.Primitive.FrontFace), + raw: pipelineState, + device: d, + layout: pipeLayout, + cullMode: cullModeToMTL(desc.Primitive.CullMode), + frontFace: frontFaceToMTL(desc.Primitive.FrontFace), + icbCompatible: icbCompatible, depthStencil: depthStencilState, depthBias: depthBias, @@ -1168,10 +1185,7 @@ func (d *Device) FreeCommandBuffer(cmdBuffer hal.CommandBuffer) { if !ok || cb == nil { return } - if cb.raw != 0 { - Release(cb.raw) - cb.raw = 0 - } + cb.Destroy() } // CreateRenderBundleEncoder is not supported in Metal backend. @@ -1254,6 +1268,7 @@ func (d *Device) WaitIdle() error { // Destroy releases the device and associated resources. func (d *Device) Destroy() { hal.Logger().Debug("metal: device destroyed") + d.releaseIndexedICBTranslators() if d.eventListener != 0 { Release(d.eventListener) d.eventListener = 0 diff --git a/hal/metal/encoder.go b/hal/metal/encoder.go index 5398d21..01cc112 100644 --- a/hal/metal/encoder.go +++ b/hal/metal/encoder.go @@ -7,6 +7,7 @@ package metal import ( "fmt" + "sync" "github.com/gogpu/gputypes" "github.com/gogpu/wgpu/hal" @@ -21,6 +22,10 @@ type CommandEncoder struct { device *Device cmdBuffer ID label string + icbOwners []*indexedICBOwnership + finished *CommandBuffer + passState renderPassPendingState + recordErr error } // IsRecording returns true if the encoder has an active command buffer. @@ -36,6 +41,7 @@ func (e *CommandEncoder) BeginEncoding(label string) error { return fmt.Errorf("metal: encoder is already recording") } e.label = label + e.recordErr = nil // Scoped autorelease pool — drain immediately after creating the command buffer. // The command buffer is Retained so it survives the pool drain. @@ -64,8 +70,18 @@ func (e *CommandEncoder) EndEncoding() (hal.CommandBuffer, error) { if e.cmdBuffer == 0 { return nil, fmt.Errorf("metal: command encoder is not recording") } - cb := &CommandBuffer{raw: e.cmdBuffer, device: e.device} + if e.recordErr != nil { + err := e.recordErr + e.recordErr = nil + Release(e.cmdBuffer) + e.cmdBuffer = 0 + e.releaseICBOwners() + return nil, err + } + cb := &CommandBuffer{raw: e.cmdBuffer, device: e.device, icbOwners: e.icbOwners} e.cmdBuffer = 0 // Recording state becomes false + e.icbOwners = nil + e.finished = cb hal.Logger().Debug("metal: encoding ended") return cb, nil } @@ -78,13 +94,58 @@ func (e *CommandEncoder) DiscardEncoding() { Release(e.cmdBuffer) e.cmdBuffer = 0 // Recording state becomes false } + e.releaseICBOwners() + e.recordErr = nil } // ResetAll resets command buffers for reuse. -func (e *CommandEncoder) ResetAll(_ []hal.CommandBuffer) {} +func (e *CommandEncoder) ResetAll(commandBuffers []hal.CommandBuffer) { + if len(commandBuffers) == 0 { + if e.finished != nil { + e.finished.Destroy() + e.finished = nil + } + return + } + for _, raw := range commandBuffers { + if cb, ok := raw.(*CommandBuffer); ok && cb != nil { + cb.Destroy() + if e.finished == cb { + e.finished = nil + } + } + } +} -// Destroy is a no-op for Metal (command buffers are managed by MTLCommandQueue). -func (e *CommandEncoder) Destroy() {} +// Destroy releases recording, finished, and private ICB state. +func (e *CommandEncoder) Destroy() { + if e == nil { + return + } + if e.cmdBuffer != 0 { + Release(e.cmdBuffer) + e.cmdBuffer = 0 + } + e.releaseICBOwners() + e.recordErr = nil + if e.finished != nil { + e.finished.Destroy() + e.finished = nil + } +} + +func (e *CommandEncoder) failRecording(err error) { + if e != nil && err != nil && e.recordErr == nil { + e.recordErr = err + } +} + +func (e *CommandEncoder) releaseICBOwners() { + for _, owner := range e.icbOwners { + owner.release() + } + e.icbOwners = nil +} // TransitionBuffers transitions buffer states for synchronization. func (e *CommandEncoder) TransitionBuffers(_ []hal.BufferBarrier) {} @@ -265,12 +326,10 @@ func (e *CommandEncoder) BeginRenderPass(desc *hal.RenderPassDescriptor) hal.Ren if e.cmdBuffer == 0 { return nil } - // Scoped pool: rpDesc and other autoreleased objects are only needed during - // encoder creation. The encoder itself is Retained to survive pool drain. - pool := NewAutoreleasePool() - rpDesc := MsgSend(ID(GetClass("MTLRenderPassDescriptor")), Sel("renderPassDescriptor")) + // Use an owned descriptor so deferring native encoder creation does not + // require keeping an autorelease pool open across caller work. + rpDesc := MsgSend(ID(GetClass("MTLRenderPassDescriptor")), Sel("new")) if rpDesc == 0 { - pool.Drain() return nil } colorAttachments := MsgSend(rpDesc, Sel("colorAttachments")) @@ -335,14 +394,12 @@ func (e *CommandEncoder) BeginRenderPass(desc *hal.RenderPassDescriptor) hal.Ren } _ = MsgSend(stencilAttachment, Sel("setStoreAction:"), uintptr(storeOpToMTL(dsa.StencilStoreOp))) } - encoder := MsgSend(e.cmdBuffer, Sel("renderCommandEncoderWithDescriptor:"), uintptr(rpDesc)) - if encoder == 0 { - pool.Drain() - return nil - } - Retain(encoder) - pool.Drain() // drain now — encoder is Retained, rpDesc no longer needed - return &RenderPassEncoder{raw: encoder, device: e.device} + // Keep the descriptor alive but delay creation of the native render encoder + // until the first draw. Metal requires the ICB translator to run on a compute + // encoder before the render encoder exists; retaining the descriptor gives + // that backend-private lowering a narrow seam without journaling draw calls. + e.passState = renderPassPendingState{} + return &RenderPassEncoder{descriptor: rpDesc, commandEncoder: e, device: e.device, pending: &e.passState} } // BeginComputePass begins a compute pass. @@ -370,17 +427,28 @@ func (e *CommandEncoder) BeginComputePass(desc *hal.ComputePassDescriptor) hal.C // CommandBuffer implements hal.CommandBuffer for Metal. type CommandBuffer struct { - raw ID - device *Device - drawable ID // Attached drawable for presentation + raw ID + device *Device + drawable ID // Attached drawable for presentation + icbOwners []*indexedICBOwnership + destroy sync.Once } // Destroy releases the command buffer. func (cb *CommandBuffer) Destroy() { - if cb.raw != 0 { - Release(cb.raw) - cb.raw = 0 + if cb == nil { + return } + cb.destroy.Do(func() { + if cb.raw != 0 { + Release(cb.raw) + cb.raw = 0 + } + for _, owner := range cb.icbOwners { + owner.release() + } + cb.icbOwners = nil + }) } // SetDrawable attaches a drawable for presentation. @@ -391,22 +459,121 @@ func (cb *CommandBuffer) SetDrawable(drawable ID) { // RenderPassEncoder implements hal.RenderPassEncoder for Metal. type RenderPassEncoder struct { - raw ID - device *Device - pipeline *RenderPipeline - currentLayout *PipelineLayout // set by SetPipeline for SetBindGroup slot offsets - indexBuffer *Buffer - indexFormat gputypes.IndexFormat - indexOffset uint64 + raw ID + descriptor ID + commandEncoder *CommandEncoder + device *Device + pipeline *RenderPipeline + currentLayout *PipelineLayout // set by SetPipeline for SetBindGroup slot offsets + indexBuffer *Buffer + indexFormat gputypes.IndexFormat + indexOffset uint64 + pending *renderPassPendingState +} + +const ( + maxRenderBindGroups = 4 + maxRenderDynamicOffsets = 16 +) + +type renderBindGroupState struct { + group *BindGroup + offsets [maxRenderDynamicOffsets]uint32 + offsetCount uint8 + set bool +} + +type renderVertexBufferState struct { + buffer *Buffer + offset uint64 + set bool +} + +type renderPassPendingState struct { + bindGroups [maxRenderBindGroups]renderBindGroupState + vertexBuffers [maxVertexBuffers]renderVertexBufferState + viewport MTLViewport + scissor MTLScissorRect + blend gputypes.Color + stencil uint32 + viewportSet bool + scissorSet bool + blendSet bool + stencilSet bool +} + +func (e *RenderPassEncoder) beginNative() bool { + if e == nil { + return false + } + if e.raw != 0 { + return true + } + if e.descriptor == 0 || e.commandEncoder == nil || e.commandEncoder.cmdBuffer == 0 { + return false + } + + pool := NewAutoreleasePool() + encoder := MsgSend(e.commandEncoder.cmdBuffer, Sel("renderCommandEncoderWithDescriptor:"), uintptr(e.descriptor)) + Release(e.descriptor) + e.descriptor = 0 + if encoder == 0 { + e.commandEncoder.failRecording(fmt.Errorf("metal: failed to create render command encoder")) + pool.Drain() + return false + } + Retain(encoder) + e.raw = encoder + pool.Drain() + e.replayPendingState() + return true +} + +func (e *RenderPassEncoder) replayPendingState() { + if e.pipeline != nil { + e.applyPipeline(e.pipeline) + } + if e.pending == nil { + return + } + for i := range e.pending.bindGroups { + state := &e.pending.bindGroups[i] + if state.set { + e.applyBindGroup(uint32(i), state.group, state.offsets[:state.offsetCount]) + } + } + for i := range e.pending.vertexBuffers { + state := &e.pending.vertexBuffers[i] + if state.set { + e.applyVertexBuffer(uint32(i), state.buffer, state.offset) + } + } + if e.pending.viewportSet { + msgSendVoid(e.raw, Sel("setViewport:"), argStruct(e.pending.viewport, mtlViewportType)) + } + if e.pending.scissorSet { + msgSendVoid(e.raw, Sel("setScissorRect:"), argStruct(e.pending.scissor, mtlScissorRectType)) + } + if e.pending.blendSet { + e.applyBlendConstant(e.pending.blend) + } + if e.pending.stencilSet { + _ = MsgSend(e.raw, Sel("setStencilReferenceValue:"), uintptr(e.pending.stencil)) + } } // End finishes the render pass. func (e *RenderPassEncoder) End() { + e.beginNative() if e.raw != 0 { _ = MsgSend(e.raw, Sel("endEncoding")) Release(e.raw) e.raw = 0 } + if e.descriptor != 0 { + Release(e.descriptor) + e.descriptor = 0 + } } // SetPipeline sets the render pipeline. @@ -417,6 +584,13 @@ func (e *RenderPassEncoder) SetPipeline(pipeline hal.RenderPipeline) { } e.pipeline = p e.currentLayout = p.layout // store for SetBindGroup slot offset lookup + if e.raw == 0 { + return + } + e.applyPipeline(p) +} + +func (e *RenderPassEncoder) applyPipeline(p *RenderPipeline) { _ = MsgSend(e.raw, Sel("setRenderPipelineState:"), uintptr(p.raw)) _ = MsgSend(e.raw, Sel("setCullMode:"), uintptr(p.cullMode)) _ = MsgSend(e.raw, Sel("setFrontFacingWinding:"), uintptr(p.frontFace)) @@ -470,6 +644,21 @@ func (e *RenderPassEncoder) SetBindGroup(index uint32, group hal.BindGroup, offs if !ok || bg == nil { return } + if e.pending != nil && index < maxRenderBindGroups && len(offsets) <= maxRenderDynamicOffsets { + state := &e.pending.bindGroups[index] + state.group = bg + state.offsetCount = uint8(len(offsets)) + state.set = true + copy(state.offsets[:], offsets) + clear(state.offsets[len(offsets):]) + } + if e.raw == 0 { + return + } + e.applyBindGroup(index, bg, offsets) +} + +func (e *RenderPassEncoder) applyBindGroup(index uint32, bg *BindGroup, offsets []uint32) { // Metal uses per-type sequential indices: [[buffer(N)]], [[texture(M)]], [[sampler(K)]]. // naga MSL generates these indices sequentially across ALL bind groups in the @@ -521,9 +710,22 @@ func (e *RenderPassEncoder) SetBindGroup(index uint32, group hal.BindGroup, offs // SetVertexBuffer sets a vertex buffer. func (e *RenderPassEncoder) SetVertexBuffer(slot uint32, buffer hal.Buffer, offset uint64) { buf, ok := buffer.(*Buffer) - if !ok || buf == nil { + if !ok || buf == nil || slot >= maxVertexBuffers { return } + if e.pending != nil { + state := &e.pending.vertexBuffers[slot] + state.buffer = buf + state.offset = offset + state.set = true + } + if e.raw == 0 { + return + } + e.applyVertexBuffer(slot, buf, offset) +} + +func (e *RenderPassEncoder) applyVertexBuffer(slot uint32, buf *Buffer, offset uint64) { bufIdx := maxVertexBuffers - 1 - slot _ = MsgSend(e.raw, Sel("setVertexBuffer:offset:atIndex:"), uintptr(buf.raw), uintptr(offset), uintptr(bufIdx)) } @@ -542,12 +744,26 @@ func (e *RenderPassEncoder) SetIndexBuffer(buffer hal.Buffer, format gputypes.In // SetViewport sets the viewport. func (e *RenderPassEncoder) SetViewport(x, y, width, height, minDepth, maxDepth float32) { viewport := MTLViewport{OriginX: float64(x), OriginY: float64(y), Width: float64(width), Height: float64(height), ZNear: float64(minDepth), ZFar: float64(maxDepth)} + if e.pending != nil { + e.pending.viewport = viewport + e.pending.viewportSet = true + } + if e.raw == 0 { + return + } msgSendVoid(e.raw, Sel("setViewport:"), argStruct(viewport, mtlViewportType)) } // SetScissorRect sets the scissor rectangle. func (e *RenderPassEncoder) SetScissorRect(x, y, width, height uint32) { scissor := MTLScissorRect{X: NSUInteger(x), Y: NSUInteger(y), Width: NSUInteger(width), Height: NSUInteger(height)} + if e.pending != nil { + e.pending.scissor = scissor + e.pending.scissorSet = true + } + if e.raw == 0 { + return + } msgSendVoid(e.raw, Sel("setScissorRect:"), argStruct(scissor, mtlScissorRectType)) } @@ -556,6 +772,17 @@ func (e *RenderPassEncoder) SetBlendConstant(color *gputypes.Color) { if color == nil { return } + if e.pending != nil { + e.pending.blend = *color + e.pending.blendSet = true + } + if e.raw == 0 { + return + } + e.applyBlendConstant(*color) +} + +func (e *RenderPassEncoder) applyBlendConstant(color gputypes.Color) { msgSendVoid(e.raw, Sel("setBlendColorRed:green:blue:alpha:"), argFloat32(float32(color.R)), argFloat32(float32(color.G)), @@ -566,18 +793,28 @@ func (e *RenderPassEncoder) SetBlendConstant(color *gputypes.Color) { // SetStencilReference sets the stencil reference value. func (e *RenderPassEncoder) SetStencilReference(ref uint32) { + if e.pending != nil { + e.pending.stencil = ref + e.pending.stencilSet = true + } + if e.raw == 0 { + return + } _ = MsgSend(e.raw, Sel("setStencilReferenceValue:"), uintptr(ref)) } // Draw draws primitives. func (e *RenderPassEncoder) Draw(vertexCount, instanceCount, firstVertex, firstInstance uint32) { + if !e.beginNative() { + return + } _ = MsgSend(e.raw, Sel("drawPrimitives:vertexStart:vertexCount:instanceCount:baseInstance:"), uintptr(MTLPrimitiveTypeTriangle), uintptr(firstVertex), uintptr(vertexCount), uintptr(instanceCount), uintptr(firstInstance)) } // DrawIndexed draws indexed primitives. func (e *RenderPassEncoder) DrawIndexed(indexCount, instanceCount, firstIndex uint32, baseVertex int32, firstInstance uint32) { - if e.indexBuffer == nil { + if e.indexBuffer == nil || !e.beginNative() { return } indexType := indexFormatToMTL(e.indexFormat) @@ -603,6 +840,9 @@ func (e *RenderPassEncoder) DrawIndirect(buffer hal.Buffer, offset uint64, drawC if _, ok := indirectRecordOffset(offset, 16, drawCount-1); !ok { return } + if !e.beginNative() { + return + } for i := uint32(0); i < drawCount; i++ { recordOffset, _ := indirectRecordOffset(offset, 16, i) _ = MsgSend(e.raw, Sel("drawPrimitives:indirectBuffer:indirectBufferOffset:"), @@ -624,6 +864,12 @@ func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64 if _, ok := indirectRecordOffset(offset, 20, drawCount-1); !ok { return } + if e.tryFirstIndexedICB(buf, offset, drawCount) { + return + } + if !e.beginNative() { + return + } indexType := indexFormatToMTL(e.indexFormat) for i := uint32(0); i < drawCount; i++ { recordOffset, ok := indirectRecordOffset(offset, 20, i) @@ -635,6 +881,17 @@ func (e *RenderPassEncoder) DrawIndexedIndirect(buffer hal.Buffer, offset uint64 } } +func (e *RenderPassEncoder) tryFirstIndexedICB(buffer *Buffer, offset uint64, count uint32) bool { + commands, ok := e.prepareIndexedICB(buffer, offset, count) + if !ok { + return false + } + if !e.beginNative() { + return true + } + return e.executeIndexedICB(commands, buffer) +} + func indirectRecordOffset(offset, stride uint64, index uint32) (uint64, bool) { delta := uint64(index) * stride if offset > ^uint64(0)-delta { diff --git a/hal/metal/encoder_test.go b/hal/metal/encoder_test.go index aaaa56b..9314b40 100644 --- a/hal/metal/encoder_test.go +++ b/hal/metal/encoder_test.go @@ -6,11 +6,22 @@ package metal import ( + "errors" "testing" "github.com/gogpu/gputypes" ) +func TestCommandEncoderRecordingErrorKeepsFirstFailure(t *testing.T) { + first := errors.New("first") + encoder := &CommandEncoder{} + encoder.failRecording(first) + encoder.failRecording(errors.New("second")) + if !errors.Is(encoder.recordErr, first) { + t.Fatalf("recording error = %v, want first failure", encoder.recordErr) + } +} + // TestCommandEncoder_RecordingState is a regression test for Issue #24. // The IsRecording() method returns cmdBuffer != 0. // @@ -102,6 +113,83 @@ func TestCommandEncoder_IsRecordingMethod(t *testing.T) { } } +func TestRenderPassDeferredStateKeepsLatestValues(t *testing.T) { + firstPipeline := &RenderPipeline{} + latestPipeline := &RenderPipeline{} + firstGroup := &BindGroup{} + latestGroup := &BindGroup{} + firstVertex := &Buffer{} + latestVertex := &Buffer{} + index := &Buffer{} + state := renderPassPendingState{} + pass := &RenderPassEncoder{pending: &state} + + pass.SetPipeline(firstPipeline) + pass.SetPipeline(latestPipeline) + pass.SetBindGroup(1, firstGroup, []uint32{4, 8}) + pass.SetBindGroup(1, latestGroup, []uint32{12}) + pass.SetVertexBuffer(2, firstVertex, 16) + pass.SetVertexBuffer(2, latestVertex, 24) + pass.SetIndexBuffer(index, gputypes.IndexFormatUint32, 32) + pass.SetViewport(1, 2, 3, 4, 0.25, 0.75) + pass.SetScissorRect(5, 6, 7, 8) + pass.SetBlendConstant(&gputypes.Color{R: 0.1, G: 0.2, B: 0.3, A: 0.4}) + pass.SetStencilReference(9) + + if pass.pipeline != latestPipeline { + t.Fatal("deferred pipeline did not keep latest value") + } + group := pass.pending.bindGroups[1] + if !group.set || group.group != latestGroup || group.offsetCount != 1 || group.offsets[0] != 12 { + t.Fatalf("deferred bind group = %#v, want latest group and offsets", group) + } + vertex := pass.pending.vertexBuffers[2] + if !vertex.set || vertex.buffer != latestVertex || vertex.offset != 24 { + t.Fatalf("deferred vertex buffer = %#v, want latest buffer and offset", vertex) + } + if pass.indexBuffer != index || pass.indexFormat != gputypes.IndexFormatUint32 || pass.indexOffset != 32 { + t.Fatal("deferred index buffer state was not retained") + } + if !pass.pending.viewportSet || pass.pending.viewport != (MTLViewport{OriginX: 1, OriginY: 2, Width: 3, Height: 4, ZNear: 0.25, ZFar: 0.75}) { + t.Fatalf("deferred viewport = %#v", pass.pending.viewport) + } + if !pass.pending.scissorSet || pass.pending.scissor != (MTLScissorRect{X: 5, Y: 6, Width: 7, Height: 8}) { + t.Fatalf("deferred scissor = %#v", pass.pending.scissor) + } + if !pass.pending.blendSet || pass.pending.blend != (gputypes.Color{R: 0.1, G: 0.2, B: 0.3, A: 0.4}) { + t.Fatalf("deferred blend constant = %#v", pass.pending.blend) + } + if !pass.pending.stencilSet || pass.pending.stencil != 9 { + t.Fatalf("deferred stencil reference = %d", pass.pending.stencil) + } +} + +func TestRenderPassDeferredStateUsesBoundedStorage(t *testing.T) { + group := &BindGroup{} + buffer := &Buffer{} + pipeline := &RenderPipeline{} + state := renderPassPendingState{} + var pass RenderPassEncoder + offsets := []uint32{1, 2, 3, 4} + + allocs := testing.AllocsPerRun(100, func() { + clear(state.bindGroups[:]) + clear(state.vertexBuffers[:]) + pass = RenderPassEncoder{pending: &state} + pass.SetPipeline(pipeline) + pass.SetBindGroup(0, group, offsets) + pass.SetVertexBuffer(0, buffer, 12) + pass.SetIndexBuffer(buffer, gputypes.IndexFormatUint16, 4) + pass.SetViewport(0, 0, 10, 10, 0, 1) + pass.SetScissorRect(0, 0, 10, 10) + pass.SetBlendConstant(&gputypes.Color{A: 1}) + pass.SetStencilReference(2) + }) + if allocs != 0 { + t.Fatalf("deferred state allocated %.2f objects per recording", allocs) + } +} + // TestComputeBindSlots_PerTypeSequentialIndexing is a regression test for the // SetBindGroup slot indexing bug. // diff --git a/hal/metal/icb_indexed.go b/hal/metal/icb_indexed.go new file mode 100644 index 0000000..654cb2a --- /dev/null +++ b/hal/metal/icb_indexed.go @@ -0,0 +1,420 @@ +// Copyright 2025 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build darwin && !(js && wasm) + +package metal + +import ( + "fmt" + "sync" + "unsafe" + + "github.com/gogpu/gputypes" +) + +const ( + indexedICBMinCommands uint32 = 1024 + indexedICBMaxCommands uint32 = 52428 // floor(1 MiB / 20-byte records) + indexedICBMaxRetainedIDs = 128 + indexedICBThreadsPerGroup uint32 = 64 +) + +type indexedICBTranslator struct { + library ID + pipeline ID +} + +type indexedICBParams struct { + commandBase uint32 + count uint32 +} + +type indexedICBArena struct { + icb ID + argumentBuffer ID + capacity uint32 +} + +type indexedICBCommands struct { + icb ID + count uint32 + owner *indexedICBOwnership +} + +// indexedICBOwnership is the single terminal-path owner for native objects and +// explicitly retained buffers referenced by one ICB execution. The fixed array +// makes residency bounded before any optimized work is encoded. +type indexedICBOwnership struct { + once sync.Once + objects [indexedICBMaxRetainedIDs]ID + objectCount int + releaseObject func(ID) +} + +// retainOwned adopts an existing +1 Objective-C reference. +func (o *indexedICBOwnership) retainOwned(id ID) bool { + if o == nil || id == 0 || o.objectCount >= len(o.objects) { + return false + } + o.objects[o.objectCount] = id + o.objectCount++ + return true +} + +func (o *indexedICBOwnership) retainReference(id ID) bool { + if o == nil || id == 0 || o.objectCount >= len(o.objects) { + return false + } + Retain(id) + return o.retainOwned(id) +} + +func (o *indexedICBOwnership) release() { + if o == nil { + return + } + o.once.Do(func() { + release := o.releaseObject + if release == nil { + release = Release + } + for i := o.objectCount - 1; i >= 0; i-- { + release(o.objects[i]) + o.objects[i] = 0 + } + o.objectCount = 0 + }) +} + +func indexedICBCountEligible(count uint32) bool { + return count >= indexedICBMinCommands && count <= indexedICBMaxCommands +} + +func indexedICBArenaCapacity(count, max uint32) uint32 { + if count == 0 || max == 0 || count > max { + return 0 + } + capacity := indexedICBMinCommands + if capacity > max { + capacity = max + } + for capacity < count { + next := capacity * 2 + if next < capacity || next > max { + capacity = max + break + } + capacity = next + } + if capacity < count { + return 0 + } + return capacity +} + +func indexedICBDispatchGroups(count uint32) uint32 { + if count == 0 { + return 0 + } + return (count + indexedICBThreadsPerGroup - 1) / indexedICBThreadsPerGroup +} + +func (d *Device) indexedICBTranslator(format gputypes.IndexFormat) (indexedICBTranslator, error) { + d.icbTranslatorMu.Lock() + defer d.icbTranslatorMu.Unlock() + if d.icbTranslators == nil { + d.icbTranslators = make(map[gputypes.IndexFormat]indexedICBTranslator, 2) + } + if translator, ok := d.icbTranslators[format]; ok { + return translator, nil + } + translator, err := compileIndexedICBTranslator(d, format) + if err != nil { + return indexedICBTranslator{}, err + } + d.icbTranslators[format] = translator + return translator, nil +} + +func (d *Device) releaseIndexedICBTranslators() { + d.icbTranslatorMu.Lock() + defer d.icbTranslatorMu.Unlock() + for format, translator := range d.icbTranslators { + if translator.pipeline != 0 { + Release(translator.pipeline) + } + if translator.library != 0 { + Release(translator.library) + } + delete(d.icbTranslators, format) + } +} + +func compileIndexedICBTranslator(d *Device, format gputypes.IndexFormat) (indexedICBTranslator, error) { + indexType := "uint" + switch format { + case gputypes.IndexFormatUint16: + indexType = "ushort" + case gputypes.IndexFormatUint32: + default: + return indexedICBTranslator{}, fmt.Errorf("metal: unsupported ICB index format %d", format) + } + source := fmt.Sprintf(` +#include +#include +using namespace metal; +struct DrawIndexedArgs { uint indexCount; uint instanceCount; uint firstIndex; int baseVertex; uint firstInstance; }; +struct ICBContainer { command_buffer commandBuffer [[id(0)]]; }; +struct IndexedICBParams { uint commandBase; uint count; }; +kernel void wgpu_translate_indexed_icb(device const DrawIndexedArgs* args [[buffer(0)]], + device ICBContainer* container [[buffer(1)]], device const %s* indices [[buffer(2)]], + constant IndexedICBParams& params [[buffer(3)]], uint id [[thread_position_in_grid]]) { + if (id >= params.count) return; + DrawIndexedArgs a = args[id]; + render_command command(container->commandBuffer, params.commandBase + id); + command.draw_indexed_primitives(primitive_type::triangle, a.indexCount, + indices + a.firstIndex, a.instanceCount, a.baseVertex, a.firstInstance); +} +`, indexType) + + str := NSString(source) + defer Release(str) + var errorPtr ID + library := MsgSend(d.raw, Sel("newLibraryWithSource:options:error:"), uintptr(str), 0, uintptr(unsafe.Pointer(&errorPtr))) + if library == 0 { + return indexedICBTranslator{}, fmt.Errorf("metal: indexed ICB translator compilation failed: %s", formatNSError(errorPtr)) + } + name := NSString("wgpu_translate_indexed_icb") + function := MsgSend(library, Sel("newFunctionWithName:"), uintptr(name)) + Release(name) + if function == 0 { + Release(library) + return indexedICBTranslator{}, fmt.Errorf("metal: indexed ICB translator function missing") + } + pipeline := MsgSend(d.raw, Sel("newComputePipelineStateWithFunction:error:"), uintptr(function), uintptr(unsafe.Pointer(&errorPtr))) + Release(function) + if pipeline == 0 { + Release(library) + return indexedICBTranslator{}, fmt.Errorf("metal: indexed ICB translator pipeline creation failed: %s", formatNSError(errorPtr)) + } + return indexedICBTranslator{library: library, pipeline: pipeline}, nil +} + +func (d *Device) newIndexedICBArena(translator indexedICBTranslator, count uint32) (*indexedICBArena, error) { + capacity := indexedICBArenaCapacity(count, indexedICBMaxCommands) + if capacity == 0 { + return nil, fmt.Errorf("metal: indexed ICB command count out of range") + } + descriptor := MsgSend(ID(GetClass("MTLIndirectCommandBufferDescriptor")), Sel("new")) + if descriptor == 0 { + return nil, fmt.Errorf("metal: failed to create ICB descriptor") + } + defer Release(descriptor) + _ = MsgSend(descriptor, Sel("setCommandTypes:"), uintptr(MTLIndirectCommandTypeDrawIndexed)) + _ = MsgSend(descriptor, Sel("setInheritPipelineState:"), uintptr(YES)) + _ = MsgSend(descriptor, Sel("setInheritBuffers:"), uintptr(YES)) + icb := MsgSend(d.raw, Sel("newIndirectCommandBufferWithDescriptor:maxCommandCount:options:"), + uintptr(descriptor), uintptr(capacity), 0) + if icb == 0 { + return nil, fmt.Errorf("metal: failed to create indexed ICB") + } + + name := NSString("wgpu_translate_indexed_icb") + function := MsgSend(translator.library, Sel("newFunctionWithName:"), uintptr(name)) + Release(name) + if function == 0 { + Release(icb) + return nil, fmt.Errorf("metal: indexed ICB translator function missing") + } + defer Release(function) + argumentEncoder := MsgSend(function, Sel("newArgumentEncoderWithBufferIndex:"), uintptr(1)) + if argumentEncoder == 0 { + Release(icb) + return nil, fmt.Errorf("metal: failed to create ICB argument encoder") + } + defer Release(argumentEncoder) + encodedLength := MsgSendUint(argumentEncoder, Sel("encodedLength")) + argumentBuffer := MsgSend(d.raw, Sel("newBufferWithLength:options:"), uintptr(encodedLength), uintptr(MTLResourceStorageModePrivate)) + if argumentBuffer == 0 { + Release(icb) + return nil, fmt.Errorf("metal: failed to create ICB argument buffer") + } + _ = MsgSend(argumentEncoder, Sel("setArgumentBuffer:offset:"), uintptr(argumentBuffer), 0) + _ = MsgSend(argumentEncoder, Sel("setIndirectCommandBuffer:atIndex:"), uintptr(icb), 0) + return &indexedICBArena{icb: icb, argumentBuffer: argumentBuffer, capacity: capacity}, nil +} + +func (d *Device) indexedICBSelectorsAvailable() bool { + if d == nil || d.raw == 0 || !DeviceSupportsFamily(d.raw, MTLGPUFamilyApple7) { + return false + } + for _, name := range []string{ + "newIndirectCommandBufferWithDescriptor:maxCommandCount:options:", + "newLibraryWithSource:options:error:", + "newComputePipelineStateWithFunction:error:", + "newBufferWithLength:options:", + } { + selector := Sel(name) + if selector == 0 || !MsgSendBool(d.raw, Sel("respondsToSelector:"), uintptr(selector)) { + return false + } + } + class := GetClass("MTLIndirectCommandBufferDescriptor") + if class == 0 { + return false + } + for _, name := range []string{"setCommandTypes:", "setInheritPipelineState:", "setInheritBuffers:"} { + selector := Sel(name) + if selector == 0 || !MsgSendBool(ID(class), Sel("instancesRespondToSelector:"), uintptr(selector)) { + return false + } + } + for _, name := range []string{ + "resetWithRange:", "executeCommandsInBuffer:withRange:", + "newArgumentEncoderWithBufferIndex:", "setArgumentBuffer:offset:", + "setIndirectCommandBuffer:atIndex:", + } { + if Sel(name) == 0 { + return false + } + } + return true +} + +func (e *RenderPassEncoder) prepareIndexedICB(arguments *Buffer, offset uint64, count uint32) (*indexedICBCommands, bool) { + if e == nil || e.commandEncoder == nil || e.commandEncoder.cmdBuffer == 0 || e.device == nil || + e.raw != 0 || e.pipeline == nil || !e.pipeline.icbCompatible || e.indexBuffer == nil || + arguments == nil || arguments.raw == 0 || e.indexBuffer.raw == 0 || !indexedICBCountEligible(count) || + !e.device.indexedICBSelectorsAvailable() { + return nil, false + } + if e.indexFormat != gputypes.IndexFormatUint16 && e.indexFormat != gputypes.IndexFormatUint32 { + return nil, false + } + if !indirectRangeFits(arguments.size, offset, 20, count) { + return nil, false + } + + translator, err := e.device.indexedICBTranslator(e.indexFormat) + if err != nil { + return nil, false + } + arena, err := e.device.newIndexedICBArena(translator, count) + if err != nil { + return nil, false + } + owner := &indexedICBOwnership{} + if !owner.retainOwned(arena.icb) || !owner.retainOwned(arena.argumentBuffer) { + owner.release() + return nil, false + } + if !e.retainIndexedICBResources(owner, arguments) { + owner.release() + return nil, false + } + + rng := NSRange{Length: NSUInteger(count)} + msgSendVoid(arena.icb, Sel("resetWithRange:"), argStruct(rng, nsRangeType)) + pool := NewAutoreleasePool() + compute := MsgSend(e.commandEncoder.cmdBuffer, Sel("computeCommandEncoder")) + if compute == 0 { + pool.Drain() + owner.release() + return nil, false + } + Retain(compute) + pool.Drain() + _ = MsgSend(compute, Sel("setComputePipelineState:"), uintptr(translator.pipeline)) + _ = MsgSend(compute, Sel("setBuffer:offset:atIndex:"), uintptr(arguments.raw), uintptr(offset), 0) + _ = MsgSend(compute, Sel("setBuffer:offset:atIndex:"), uintptr(arena.argumentBuffer), 0, 1) + _ = MsgSend(compute, Sel("setBuffer:offset:atIndex:"), uintptr(e.indexBuffer.raw), uintptr(e.indexOffset), 2) + params := indexedICBParams{count: count} + _ = MsgSend(compute, Sel("setBytes:length:atIndex:"), uintptr(unsafe.Pointer(¶ms)), uintptr(unsafe.Sizeof(params)), 3) + _ = MsgSend(compute, Sel("useResource:usage:"), uintptr(arguments.raw), 1) + _ = MsgSend(compute, Sel("useResource:usage:"), uintptr(arena.argumentBuffer), 1) + _ = MsgSend(compute, Sel("useResource:usage:"), uintptr(e.indexBuffer.raw), 1) + _ = MsgSend(compute, Sel("useResource:usage:"), uintptr(arena.icb), 2) + groups := MTLSize{Width: NSUInteger(indexedICBDispatchGroups(count)), Height: 1, Depth: 1} + threads := MTLSize{Width: NSUInteger(indexedICBThreadsPerGroup), Height: 1, Depth: 1} + msgSendVoid(compute, Sel("dispatchThreadgroups:threadsPerThreadgroup:"), + argStruct(groups, mtlSizeType), argStruct(threads, mtlSizeType)) + _ = MsgSend(compute, Sel("endEncoding")) + Release(compute) + + e.commandEncoder.icbOwners = append(e.commandEncoder.icbOwners, owner) + return &indexedICBCommands{icb: arena.icb, count: count, owner: owner}, true +} + +func (e *RenderPassEncoder) executeIndexedICB(commands *indexedICBCommands, arguments *Buffer) bool { + if commands == nil || commands.icb == 0 || commands.count == 0 || e.raw == 0 { + return false + } + e.declareIndexedICBResources(arguments) + rng := NSRange{Length: NSUInteger(commands.count)} + msgSendVoid(e.raw, Sel("executeCommandsInBuffer:withRange:"), + argPointer(uintptr(commands.icb)), argStruct(rng, nsRangeType)) + return true +} + +func (e *RenderPassEncoder) declareIndexedICBResources(arguments *Buffer) { + declare := func(id ID) { + if id != 0 { + _ = MsgSend(e.raw, Sel("useResource:usage:"), uintptr(id), 1) + } + } + if arguments != nil { + declare(arguments.raw) + } + if e.indexBuffer != nil { + declare(e.indexBuffer.raw) + } + if e.pending == nil { + return + } + for i := range e.pending.vertexBuffers { + state := &e.pending.vertexBuffers[i] + if state.set && state.buffer != nil { + declare(state.buffer.raw) + } + } + for i := range e.pending.bindGroups { + state := &e.pending.bindGroups[i] + if !state.set || state.group == nil { + continue + } + for _, entry := range state.group.entries { + if binding, ok := entry.Resource.(gputypes.BufferBinding); ok { + declare(ID(binding.Buffer)) + } + } + } +} + +func (e *RenderPassEncoder) retainIndexedICBResources(owner *indexedICBOwnership, arguments *Buffer) bool { + if !owner.retainReference(arguments.raw) || !owner.retainReference(e.indexBuffer.raw) { + return false + } + if e.pending == nil { + return true + } + for i := range e.pending.vertexBuffers { + state := &e.pending.vertexBuffers[i] + if state.set && state.buffer != nil && !owner.retainReference(state.buffer.raw) { + return false + } + } + for i := range e.pending.bindGroups { + state := &e.pending.bindGroups[i] + if !state.set || state.group == nil { + continue + } + for _, entry := range state.group.entries { + binding, ok := entry.Resource.(gputypes.BufferBinding) + if !ok || binding.Buffer == 0 || !owner.retainReference(ID(binding.Buffer)) { + return false + } + } + } + return true +} diff --git a/hal/metal/icb_indexed_test.go b/hal/metal/icb_indexed_test.go new file mode 100644 index 0000000..052f61f --- /dev/null +++ b/hal/metal/icb_indexed_test.go @@ -0,0 +1,114 @@ +// Copyright 2025 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build darwin && !(js && wasm) + +package metal + +import ( + "sync" + "testing" + + "github.com/gogpu/wgpu/hal" +) + +func TestIndexedICBArenaCapacityIsGeometricAndBounded(t *testing.T) { + const max = uint32(52428) + for _, test := range []struct { + count uint32 + want uint32 + }{ + {1024, 1024}, + {1025, 2048}, + {10000, 16384}, + {max, max}, + {max + 1, 0}, + } { + if got := indexedICBArenaCapacity(test.count, max); got != test.want { + t.Fatalf("capacity(%d) = %d, want %d", test.count, got, test.want) + } + } +} + +func TestIndexedICBCountGate(t *testing.T) { + for count, want := range map[uint32]bool{ + 0: false, 1: false, 20: false, 1023: false, 1024: true, + indexedICBMaxCommands: true, indexedICBMaxCommands + 1: false, + } { + if got := indexedICBCountEligible(count); got != want { + t.Fatalf("eligible(%d) = %v, want %v", count, got, want) + } + } +} + +func TestIndexedICBOwnershipReleasesExactlyOnceAcrossRaces(t *testing.T) { + var mu sync.Mutex + released := make(map[ID]int) + owner := &indexedICBOwnership{releaseObject: func(id ID) { + mu.Lock() + released[id]++ + mu.Unlock() + }} + for _, id := range []ID{1, 2, 3} { + if !owner.retainOwned(id) { + t.Fatalf("failed to retain object %d", id) + } + } + + var wg sync.WaitGroup + for range 8 { + wg.Add(1) + go func() { + defer wg.Done() + owner.release() + }() + } + wg.Wait() + for _, id := range []ID{1, 2, 3} { + if released[id] != 1 { + t.Fatalf("object %d released %d times, want once", id, released[id]) + } + } +} + +func TestCommandEncoderResetAllReleasesFinishedICBOwnership(t *testing.T) { + releases := 0 + owner := &indexedICBOwnership{releaseObject: func(ID) { releases++ }} + owner.retainOwned(1) + cb := &CommandBuffer{icbOwners: []*indexedICBOwnership{owner}} + encoder := &CommandEncoder{finished: cb} + + encoder.ResetAll(nil) + encoder.ResetAll([]hal.CommandBuffer{cb}) + if releases != 1 { + t.Fatalf("release count = %d, want one", releases) + } + if encoder.finished != nil || cb.icbOwners != nil { + t.Fatal("finished ICB ownership was retained after reset") + } +} + +func TestCommandBufferDestroyReleasesExactlyOnceAcrossRaces(t *testing.T) { + var mu sync.Mutex + releases := 0 + owner := &indexedICBOwnership{releaseObject: func(ID) { + mu.Lock() + releases++ + mu.Unlock() + }} + owner.retainOwned(1) + cb := &CommandBuffer{icbOwners: []*indexedICBOwnership{owner}} + + var wg sync.WaitGroup + for range 8 { + wg.Add(1) + go func() { + defer wg.Done() + cb.Destroy() + }() + } + wg.Wait() + if releases != 1 { + t.Fatalf("release count = %d, want one", releases) + } +} diff --git a/hal/metal/icb_pipeline.go b/hal/metal/icb_pipeline.go new file mode 100644 index 0000000..d995fa4 --- /dev/null +++ b/hal/metal/icb_pipeline.go @@ -0,0 +1,60 @@ +// Copyright 2025 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build darwin && !(js && wasm) + +package metal + +import ( + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +func renderPipelineICBCandidate(desc *hal.RenderPipelineDescriptor) bool { + if desc == nil || desc.Primitive.Topology != gputypes.PrimitiveTopologyTriangleList { + return false + } + layout, ok := desc.Layout.(*PipelineLayout) + if !ok || layout == nil { + return false + } + for _, rawLayout := range layout.layouts { + group, ok := rawLayout.(*BindGroupLayout) + if !ok || group == nil { + return false + } + for _, entry := range group.entries { + if entry.Buffer == nil || entry.Buffer.Type == gputypes.BufferBindingTypeStorage || + entry.Texture != nil || entry.StorageTexture != nil || entry.Sampler != nil { + return false + } + } + } + return true +} + +func (d *Device) canCreateICBPipeline(desc *hal.RenderPipelineDescriptor, pipelineDesc ID) bool { + if d == nil || d.raw == 0 || pipelineDesc == 0 || !renderPipelineICBCandidate(desc) { + return false + } + // Metal family support is cumulative. Apple8+ devices also report Apple7, + // so this admits the M1 proof family and later Apple GPU families while + // leaving Intel and AMD devices on the ordinary path. + if !DeviceSupportsFamily(d.raw, MTLGPUFamilyApple7) { + return false + } + setter := Sel("setSupportIndirectCommandBuffers:") + return setter != 0 && MsgSendBool(pipelineDesc, Sel("respondsToSelector:"), uintptr(setter)) +} + +// createRenderPipelineState centralizes the fail-closed policy: an ICB +// candidate gets one flagged attempt, followed by the ordinary attempt on any +// failure. Callers only observe failure when the ordinary pipeline also fails. +func createRenderPipelineState(candidate bool, create func(icb bool) ID) (ID, bool) { + if candidate { + if pipeline := create(true); pipeline != 0 { + return pipeline, true + } + } + return create(false), false +} diff --git a/hal/metal/icb_pipeline_test.go b/hal/metal/icb_pipeline_test.go new file mode 100644 index 0000000..efb4822 --- /dev/null +++ b/hal/metal/icb_pipeline_test.go @@ -0,0 +1,115 @@ +// Copyright 2025 The GoGPU Authors +// SPDX-License-Identifier: MIT + +//go:build darwin && !(js && wasm) + +package metal + +import ( + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu/hal" +) + +func TestRenderPipelineICBCandidateIsPrivateAndBufferOnly(t *testing.T) { + bufferLayout := &BindGroupLayout{entries: []gputypes.BindGroupLayoutEntry{{Buffer: &gputypes.BufferBindingLayout{}}}} + readOnlyStorageLayout := &BindGroupLayout{entries: []gputypes.BindGroupLayoutEntry{{Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeReadOnlyStorage}}}} + writableStorageLayout := &BindGroupLayout{entries: []gputypes.BindGroupLayoutEntry{{Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeStorage}}}} + textureLayout := &BindGroupLayout{entries: []gputypes.BindGroupLayoutEntry{{Texture: &gputypes.TextureBindingLayout{}}}} + + tests := []struct { + name string + desc hal.RenderPipelineDescriptor + want bool + }{ + { + name: "triangle buffer-only", + desc: hal.RenderPipelineDescriptor{ + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}, + Layout: &PipelineLayout{layouts: []hal.BindGroupLayout{bufferLayout}}, + }, + want: true, + }, + { + name: "triangle without bindings", + desc: hal.RenderPipelineDescriptor{ + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}, + Layout: &PipelineLayout{}, + }, + want: true, + }, + { + name: "read-only storage binding", + desc: hal.RenderPipelineDescriptor{ + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}, + Layout: &PipelineLayout{layouts: []hal.BindGroupLayout{readOnlyStorageLayout}}, + }, + want: true, + }, + { + name: "writable storage binding", + desc: hal.RenderPipelineDescriptor{ + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}, + Layout: &PipelineLayout{layouts: []hal.BindGroupLayout{writableStorageLayout}}, + }, + }, + { + name: "texture binding", + desc: hal.RenderPipelineDescriptor{ + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}, + Layout: &PipelineLayout{layouts: []hal.BindGroupLayout{textureLayout}}, + }, + }, + { + name: "non-triangle topology", + desc: hal.RenderPipelineDescriptor{ + Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyLineList}, + Layout: &PipelineLayout{layouts: []hal.BindGroupLayout{bufferLayout}}, + }, + }, + { + name: "foreign layout", + desc: hal.RenderPipelineDescriptor{Primitive: gputypes.PrimitiveState{Topology: gputypes.PrimitiveTopologyTriangleList}}, + }, + } + + for _, tt := range tests { + t.Run(tt.name, func(t *testing.T) { + if got := renderPipelineICBCandidate(&tt.desc); got != tt.want { + t.Fatalf("candidate = %v, want %v", got, tt.want) + } + }) + } +} + +func TestCreateRenderPipelineStateRetriesOrdinaryAfterICBFailure(t *testing.T) { + var attempts []bool + pipeline, compatible := createRenderPipelineState(true, func(icb bool) ID { + attempts = append(attempts, icb) + if icb { + return 0 + } + return 42 + }) + if pipeline != 42 || compatible { + t.Fatalf("result = (%d, %v), want (42, false)", pipeline, compatible) + } + if len(attempts) != 2 || !attempts[0] || attempts[1] { + t.Fatalf("attempts = %v, want [true false]", attempts) + } +} + +func TestCreateRenderPipelineStateSkipsFlagForIneligiblePipeline(t *testing.T) { + var attempts []bool + pipeline, compatible := createRenderPipelineState(false, func(icb bool) ID { + attempts = append(attempts, icb) + return 7 + }) + if pipeline != 7 || compatible { + t.Fatalf("result = (%d, %v), want (7, false)", pipeline, compatible) + } + if len(attempts) != 1 || attempts[0] { + t.Fatalf("attempts = %v, want [false]", attempts) + } +} diff --git a/hal/metal/metal.go b/hal/metal/metal.go index 3d7981d..80a0e8b 100644 --- a/hal/metal/metal.go +++ b/hal/metal/metal.go @@ -114,6 +114,7 @@ func preRegisterSelectors() { "newSamplerStateWithDescriptor:", "newCommandQueue", "newRenderPipelineStateWithDescriptor:error:", + "setSupportIndirectCommandBuffers:", "newLibraryWithSource:options:error:", "newFunctionWithName:", "supportsFamily:", diff --git a/hal/metal/resource.go b/hal/metal/resource.go index b5ad827..fdaf577 100644 --- a/hal/metal/resource.go +++ b/hal/metal/resource.go @@ -198,11 +198,12 @@ func (l *PipelineLayout) Destroy() { // RenderPipeline implements hal.RenderPipeline for Metal. type RenderPipeline struct { - raw ID // id - device *Device - layout *PipelineLayout // for SetBindGroup slot offset lookup - cullMode MTLCullMode - frontFace MTLWinding + raw ID // id + device *Device + layout *PipelineLayout // for SetBindGroup slot offset lookup + cullMode MTLCullMode + frontFace MTLWinding + icbCompatible bool depthStencil ID // id depthBias float32 diff --git a/hal/metal/types.go b/hal/metal/types.go index 8419fad..b5d3d00 100644 --- a/hal/metal/types.go +++ b/hal/metal/types.go @@ -82,6 +82,13 @@ type MTLScissorRect struct { Width, Height NSUInteger } +// MTLIndirectCommandType values accepted by MTLIndirectCommandBufferDescriptor. +type MTLIndirectCommandType NSUInteger + +const ( + MTLIndirectCommandTypeDrawIndexed MTLIndirectCommandType = 1 << 1 +) + // MTLPixelFormat represents a pixel format. type MTLPixelFormat NSUInteger diff --git a/indirect_count_integration_test.go b/indirect_count_integration_test.go index d1a5a83..59ca24a 100644 --- a/indirect_count_integration_test.go +++ b/indirect_count_integration_test.go @@ -3,6 +3,7 @@ package wgpu_test import ( + "bytes" "context" "encoding/binary" "math" @@ -15,9 +16,21 @@ import ( ) const countedIndexedShader = ` +struct Params { + expected_instance: u32, + _padding: vec3, +}; + +@group(0) @binding(0) +var params: Params; + +@group(0) @binding(1) +var shifts: array; + @vertex -fn vs_main(@location(0) pos: vec2) -> @builtin(position) vec4 { - return vec4(pos, 0.0, 1.0); +fn vs_main(@location(0) pos: vec2, @builtin(instance_index) instance: u32) -> @builtin(position) vec4 { + let x = select(pos.x + 4.0, pos.x + shifts[0], instance == params.expected_instance); + return vec4(x, pos.y, 0.0, 1.0); } @fragment @@ -26,7 +39,54 @@ fn fs_main() -> @location(0) vec4 { } ` +const countedIndirectWriterShader = ` +@group(0) @binding(0) +var args: array; + +@compute @workgroup_size(1) +fn main() { + args[1] = 1u; + args[6] = 1u; +} +` + func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { + uint16Indices := []byte{0, 0, 1, 0, 2, 0, 2, 0, 3, 0, 0, 0} + uint32Indices := make([]byte, 4+6*4) + for i, index := range []uint32{0, 1, 2, 2, 3, 0} { + binary.LittleEndian.PutUint32(uint32Indices[4+i*4:], index) + } + for _, test := range []struct { + name string + vertices []float32 + indices []byte + format gputypes.IndexFormat + indexOffset uint64 + baseVertex int32 + }{ + { + name: "uint16", + vertices: []float32{-1, -1, 1, -1, 1, 1, -1, 1}, + indices: uint16Indices, format: gputypes.IndexFormatUint16, + }, + { + name: "uint32_nonzero_index_offset_base_vertex", + vertices: []float32{9, 9, -1, -1, 1, -1, 1, 1, -1, 1}, + indices: uint32Indices, format: gputypes.IndexFormatUint32, indexOffset: 4, baseVertex: 1, + }, + } { + t.Run(test.name, func(t *testing.T) { + icbPixels := runCountedIndexedICBCase(t, test.vertices, test.indices, test.format, test.indexOffset, test.baseVertex, false) + loopPixels := runCountedIndexedICBCase(t, test.vertices, test.indices, test.format, test.indexOffset, test.baseVertex, true) + if !bytes.Equal(icbPixels, loopPixels) { + t.Fatal("first-draw ICB pixels differ from later-draw loop oracle") + } + }) + } +} + +func runCountedIndexedICBCase(t *testing.T, vertices []float32, indexData []byte, indexFormat gputypes.IndexFormat, indexOffset uint64, baseVertex int32, forceLoop bool) []byte { + t.Helper() instance, err := wgpu.CreateInstance(&wgpu.InstanceDescriptor{Backends: wgpu.BackendsPrimary}) if err != nil { t.Skipf("CreateInstance: %v", err) @@ -49,32 +109,85 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { defer device.Release() queue := device.Queue() - vertices := []float32{-1, -1, 1, -1, 1, 1, -1, 1} vertexData := make([]byte, len(vertices)*4) for i, value := range vertices { binary.LittleEndian.PutUint32(vertexData[i*4:], math.Float32bits(value)) } - indexData := []byte{0, 0, 1, 0, 2, 0, 2, 0, 3, 0, 0, 0} - - // Four padding bytes make the starting offset observable. The two records - // select different triangles; both are required to fill the target. - indirectData := make([]byte, 4+2*20) - putCountedIndexedIndirectRecord(indirectData[4:], 3, 1, 0, 0, 0) - putCountedIndexedIndirectRecord(indirectData[24:], 3, 1, 3, 0, 0) + // Aligned padding makes the starting offset observable. The first two + // records select different triangles; the remaining zero-instance records + // keep the 1,024-command batch visually inert while exercising Metal's ICB + // threshold. A compute pass changes the first two instance counts from zero + // to one before the render pass, proving the translator consumes GPU-written + // arguments. + const drawCount = uint32(1024) + const argumentsOffset = uint64(256) + indirectData := make([]byte, argumentsOffset+uint64(drawCount)*20) + putCountedIndexedIndirectRecord(indirectData[argumentsOffset:], 3, 0, 0, baseVertex, 3) + putCountedIndexedIndirectRecord(indirectData[argumentsOffset+20:], 3, 0, 3, baseVertex, 3) vertexBuffer := createCountedIndirectTestBuffer(t, device, queue, vertexData, wgpu.BufferUsageVertex|wgpu.BufferUsageCopyDst) defer vertexBuffer.Release() indexBuffer := createCountedIndirectTestBuffer(t, device, queue, indexData, wgpu.BufferUsageIndex|wgpu.BufferUsageCopyDst) defer indexBuffer.Release() - indirectBuffer := createCountedIndirectTestBuffer(t, device, queue, indirectData, wgpu.BufferUsageIndirect|wgpu.BufferUsageCopyDst) + indirectBuffer := createCountedIndirectTestBuffer(t, device, queue, indirectData, wgpu.BufferUsageIndirect|wgpu.BufferUsageStorage|wgpu.BufferUsageCopyDst) defer indirectBuffer.Release() + writerShader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{WGSL: countedIndirectWriterShader}) + if err != nil { + t.Fatalf("CreateShaderModule(writer): %v", err) + } + defer writerShader.Release() + writerBGL, err := device.CreateBindGroupLayout(&wgpu.BindGroupLayoutDescriptor{Entries: []wgpu.BindGroupLayoutEntry{{ + Binding: 0, Visibility: wgpu.ShaderStageCompute, + Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeStorage}, + }}}) + if err != nil { + t.Fatalf("CreateBindGroupLayout(writer): %v", err) + } + defer writerBGL.Release() + writerLayout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{BindGroupLayouts: []*wgpu.BindGroupLayout{writerBGL}}) + if err != nil { + t.Fatalf("CreatePipelineLayout(writer): %v", err) + } + defer writerLayout.Release() + writerPipeline, err := device.CreateComputePipeline(&wgpu.ComputePipelineDescriptor{ + Layout: writerLayout, Module: writerShader, EntryPoint: "main", + }) + if err != nil { + t.Fatalf("CreateComputePipeline(writer): %v", err) + } + defer writerPipeline.Release() + writerGroup, err := device.CreateBindGroup(&wgpu.BindGroupDescriptor{ + Layout: writerBGL, + Entries: []wgpu.BindGroupEntry{{ + Binding: 0, Buffer: indirectBuffer, Offset: argumentsOffset, Size: uint64(drawCount) * 20, + }}, + }) + if err != nil { + t.Fatalf("CreateBindGroup(writer): %v", err) + } + defer writerGroup.Release() + shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{WGSL: countedIndexedShader}) if err != nil { t.Fatalf("CreateShaderModule: %v", err) } defer shader.Release() - layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{}) + renderBGL, err := device.CreateBindGroupLayout(&wgpu.BindGroupLayoutDescriptor{Entries: []wgpu.BindGroupLayoutEntry{ + { + Binding: 0, Visibility: gputypes.ShaderStageVertex, + Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeUniform, MinBindingSize: 16}, + }, + { + Binding: 1, Visibility: gputypes.ShaderStageVertex, + Buffer: &gputypes.BufferBindingLayout{Type: gputypes.BufferBindingTypeReadOnlyStorage, MinBindingSize: 4}, + }, + }}) + if err != nil { + t.Fatalf("CreateBindGroupLayout(render): %v", err) + } + defer renderBGL.Release() + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{BindGroupLayouts: []*wgpu.BindGroupLayout{renderBGL}}) if err != nil { t.Fatalf("CreatePipelineLayout: %v", err) } @@ -105,6 +218,25 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { } defer pipeline.Release() + uniformData := make([]byte, 16) + binary.LittleEndian.PutUint32(uniformData, 3) + storageData := make([]byte, 4) + renderUniform := createCountedIndirectTestBuffer(t, device, queue, uniformData, wgpu.BufferUsageUniform|wgpu.BufferUsageCopyDst) + defer renderUniform.Release() + renderStorage := createCountedIndirectTestBuffer(t, device, queue, storageData, wgpu.BufferUsageStorage|wgpu.BufferUsageCopyDst) + defer renderStorage.Release() + renderGroup, err := device.CreateBindGroup(&wgpu.BindGroupDescriptor{ + Layout: renderBGL, + Entries: []wgpu.BindGroupEntry{ + {Binding: 0, Buffer: renderUniform, Size: 16}, + {Binding: 1, Buffer: renderStorage, Size: 4}, + }, + }) + if err != nil { + t.Fatalf("CreateBindGroup(render): %v", err) + } + defer renderGroup.Release() + const width, height = uint32(64), uint32(16) target, err := device.CreateTexture(&wgpu.TextureDescriptor{ Size: wgpu.Extent3D{Width: width, Height: height, DepthOrArrayLayers: 1}, @@ -121,11 +253,36 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { t.Fatalf("CreateTextureView: %v", err) } defer targetView.Release() + clearOnlyTarget, err := device.CreateTexture(&wgpu.TextureDescriptor{ + Size: wgpu.Extent3D{Width: width, Height: 1, DepthOrArrayLayers: 1}, + MipLevelCount: 1, SampleCount: 1, Dimension: gputypes.TextureDimension2D, + Format: gputypes.TextureFormatRGBA8Unorm, + Usage: gputypes.TextureUsageRenderAttachment | gputypes.TextureUsageCopySrc, + }) + if err != nil { + t.Fatalf("CreateTexture(clear-only): %v", err) + } + defer clearOnlyTarget.Release() + clearOnlyView, err := device.CreateTextureView(clearOnlyTarget, nil) + if err != nil { + t.Fatalf("CreateTextureView(clear-only): %v", err) + } + defer clearOnlyView.Release() encoder, err := device.CreateCommandEncoder(&wgpu.CommandEncoderDescriptor{}) if err != nil { t.Fatalf("CreateCommandEncoder: %v", err) } + compute, err := encoder.BeginComputePass(&wgpu.ComputePassDescriptor{}) + if err != nil { + t.Fatalf("BeginComputePass: %v", err) + } + compute.SetPipeline(writerPipeline) + compute.SetBindGroup(0, writerGroup, nil) + compute.Dispatch(1, 1, 1) + if err := compute.End(); err != nil { + t.Fatalf("ComputePass.End: %v", err) + } pass, err := encoder.BeginRenderPass(&wgpu.RenderPassDescriptor{ ColorAttachments: []wgpu.RenderPassColorAttachment{{ View: targetView, LoadOp: gputypes.LoadOpClear, StoreOp: gputypes.StoreOpStore, @@ -136,19 +293,44 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { t.Fatalf("BeginRenderPass: %v", err) } pass.SetPipeline(pipeline) + pass.SetBindGroup(0, renderGroup, nil) pass.SetVertexBuffer(0, vertexBuffer, 0) - pass.SetIndexBuffer(indexBuffer, gputypes.IndexFormatUint16, 0) - pass.MultiDrawIndexedIndirect(indirectBuffer, 4, 2) + pass.SetIndexBuffer(indexBuffer, indexFormat, indexOffset) + if forceLoop { + pass.Draw(0, 0, 0, 0) + } + pass.MultiDrawIndexedIndirect(indirectBuffer, argumentsOffset, drawCount) if err := pass.End(); err != nil { t.Fatalf("RenderPass.End: %v", err) } - encoder.TransitionTextures([]wgpu.TextureBarrier{{ - Texture: target, - Usage: wgpu.TextureUsageTransition{ - OldUsage: gputypes.TextureUsageRenderAttachment, - NewUsage: gputypes.TextureUsageCopySrc, + clearOnlyPass, err := encoder.BeginRenderPass(&wgpu.RenderPassDescriptor{ + ColorAttachments: []wgpu.RenderPassColorAttachment{{ + View: clearOnlyView, LoadOp: gputypes.LoadOpClear, StoreOp: gputypes.StoreOpStore, + ClearValue: gputypes.Color{R: 1, G: 0.25, B: 0.5, A: 1}, + }}, + }) + if err != nil { + t.Fatalf("BeginRenderPass(clear-only): %v", err) + } + if err := clearOnlyPass.End(); err != nil { + t.Fatalf("RenderPass.End(clear-only): %v", err) + } + encoder.TransitionTextures([]wgpu.TextureBarrier{ + { + Texture: target, + Usage: wgpu.TextureUsageTransition{ + OldUsage: gputypes.TextureUsageRenderAttachment, + NewUsage: gputypes.TextureUsageCopySrc, + }, }, - }}) + { + Texture: clearOnlyTarget, + Usage: wgpu.TextureUsageTransition{ + OldUsage: gputypes.TextureUsageRenderAttachment, + NewUsage: gputypes.TextureUsageCopySrc, + }, + }, + }) readbackSize := uint64(width * height * 4) readback, err := device.CreateBuffer(&wgpu.BufferDescriptor{ @@ -163,6 +345,18 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { TextureBase: wgpu.ImageCopyTexture{Texture: target}, Size: wgpu.Extent3D{Width: width, Height: height, DepthOrArrayLayers: 1}, }}) + clearOnlyReadback, err := device.CreateBuffer(&wgpu.BufferDescriptor{ + Size: uint64(width * 4), Usage: gputypes.BufferUsageMapRead | gputypes.BufferUsageCopyDst, + }) + if err != nil { + t.Fatalf("CreateBuffer(clear-only readback): %v", err) + } + defer clearOnlyReadback.Release() + encoder.CopyTextureToBuffer(clearOnlyTarget, clearOnlyReadback, []wgpu.BufferTextureCopy{{ + BufferLayout: wgpu.ImageDataLayout{BytesPerRow: width * 4, RowsPerImage: 1}, + TextureBase: wgpu.ImageCopyTexture{Texture: clearOnlyTarget}, + Size: wgpu.Extent3D{Width: width, Height: 1, DepthOrArrayLayers: 1}, + }}) commands, err := encoder.Finish() if err != nil { @@ -171,6 +365,22 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { if _, err := queue.Submit(commands); err != nil { t.Fatalf("Submit: %v", err) } + clearMapContext, cancelClearMap := context.WithTimeout(context.Background(), 5*time.Second) + defer cancelClearMap() + if err := clearOnlyReadback.Map(clearMapContext, wgpu.MapModeRead, 0, uint64(width*4)); err != nil { + t.Fatalf("Map(clear-only readback): %v", err) + } + clearMapped, err := clearOnlyReadback.MappedRange(0, uint64(width*4)) + if err != nil { + t.Fatalf("MappedRange(clear-only readback): %v", err) + } + clearPixel := clearMapped.Bytes()[:4] + if clearPixel[0] < 250 || clearPixel[1] < 60 || clearPixel[1] > 68 || clearPixel[2] < 124 || clearPixel[2] > 132 || clearPixel[3] < 250 { + t.Fatalf("clear-only pass pixel = %v, want approximately [255 64 128 255]", clearPixel) + } + if err := clearOnlyReadback.Unmap(); err != nil { + t.Fatalf("Unmap(clear-only readback): %v", err) + } mapContext, cancelMap := context.WithTimeout(context.Background(), 5*time.Second) defer cancelMap() if err := readback.Map(mapContext, wgpu.MapModeRead, 0, readbackSize); err != nil { @@ -182,6 +392,7 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { t.Fatalf("MappedRange: %v", err) } pixels := mapped.Bytes() + result := append([]byte(nil), pixels...) filled := 0 for pixel := 0; pixel < int(width*height); pixel++ { if pixels[pixel*4+1] > 128 { @@ -194,6 +405,7 @@ func TestMultiDrawIndexedIndirectRendersDistinctRecords(t *testing.T) { if minimum := int(width*height) * 3 / 4; filled < minimum { t.Fatalf("counted indexed indirect filled %d/%d pixels, want at least %d; both records did not render", filled, width*height, minimum) } + return result } func createCountedIndirectTestBuffer(t *testing.T, device *wgpu.Device, queue *wgpu.Queue, data []byte, usage wgpu.BufferUsage) *wgpu.Buffer { diff --git a/metal_icb_benchmark_test.go b/metal_icb_benchmark_test.go new file mode 100644 index 0000000..71933cb --- /dev/null +++ b/metal_icb_benchmark_test.go @@ -0,0 +1,190 @@ +//go:build integration && darwin && !rust && !(js && wasm) + +package wgpu_test + +import ( + "encoding/binary" + "runtime" + "testing" + + "github.com/gogpu/gputypes" + "github.com/gogpu/wgpu" + _ "github.com/gogpu/wgpu/hal/metal" +) + +const metalICBBenchmarkShader = ` +@vertex +fn vs_main(@location(0) pos: vec2) -> @builtin(position) vec4 { + return vec4(pos, 0.0, 1.0); +} +@fragment +fn fs_main() -> @location(0) vec4 { + return vec4(0.0, 1.0, 0.0, 1.0); +} +` + +func BenchmarkMetalIndexedFrame(b *testing.B) { + for _, test := range []struct { + name string + count uint32 + forceLoop bool + }{ + {name: "direct"}, + {name: "count_1", count: 1}, + {name: "count_20", count: 20}, + {name: "count_1024", count: 1024}, + {name: "count_1024_loop", count: 1024, forceLoop: true}, + {name: "count_10000", count: 10000}, + {name: "count_10000_loop", count: 10000, forceLoop: true}, + } { + b.Run(test.name, func(b *testing.B) { + benchmarkMetalIndexedFrame(b, test.count, test.forceLoop) + }) + } +} + +func benchmarkMetalIndexedFrame(b *testing.B, count uint32, forceLoop bool) { + runtime.LockOSThread() + defer runtime.UnlockOSThread() + + instance, err := wgpu.CreateInstance(&wgpu.InstanceDescriptor{Backends: wgpu.BackendsPrimary}) + if err != nil { + b.Skipf("CreateInstance: %v", err) + } + defer instance.Release() + adapter, err := instance.RequestAdapter(nil) + if err != nil { + b.Skipf("RequestAdapter: %v", err) + } + defer adapter.Release() + if info := adapter.Info(); info.Backend != gputypes.BackendMetal { + b.Skipf("native Metal adapter unavailable: %+v", info) + } + device, err := adapter.RequestDevice(nil) + if err != nil { + b.Fatal(err) + } + defer device.Release() + queue := device.Queue() + + shader, err := device.CreateShaderModule(&wgpu.ShaderModuleDescriptor{WGSL: metalICBBenchmarkShader}) + if err != nil { + b.Fatal(err) + } + defer shader.Release() + layout, err := device.CreatePipelineLayout(&wgpu.PipelineLayoutDescriptor{}) + if err != nil { + b.Fatal(err) + } + defer layout.Release() + pipeline, err := device.CreateRenderPipeline(&wgpu.RenderPipelineDescriptor{ + Layout: layout, + Vertex: wgpu.VertexState{ + Module: shader, EntryPoint: "vs_main", + Buffers: []gputypes.VertexBufferLayout{{ + ArrayStride: 8, + Attributes: []gputypes.VertexAttribute{{Format: gputypes.VertexFormatFloat32x2}}, + }}, + }, + Fragment: &wgpu.FragmentState{ + Module: shader, EntryPoint: "fs_main", + 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 { + b.Fatal(err) + } + defer pipeline.Release() + + vertex := benchmarkMetalBuffer(b, device, queue, make([]byte, 3*8), wgpu.BufferUsageVertex|wgpu.BufferUsageCopyDst) + defer vertex.Release() + index := benchmarkMetalBuffer(b, device, queue, make([]byte, 8), wgpu.BufferUsageIndex|wgpu.BufferUsageCopyDst) + defer index.Release() + argumentCount := max(count, 1) + arguments := make([]byte, uint64(argumentCount)*20) + for i := uint32(0); i < count; i++ { + binary.LittleEndian.PutUint32(arguments[uint64(i)*20:], 3) + } + indirect := benchmarkMetalBuffer(b, device, queue, arguments, wgpu.BufferUsageIndirect|wgpu.BufferUsageCopyDst) + defer indirect.Release() + target, err := device.CreateTexture(&wgpu.TextureDescriptor{ + Size: wgpu.Extent3D{Width: 1, Height: 1, DepthOrArrayLayers: 1}, + MipLevelCount: 1, SampleCount: 1, Dimension: gputypes.TextureDimension2D, + Format: gputypes.TextureFormatRGBA8Unorm, + Usage: gputypes.TextureUsageRenderAttachment, + }) + if err != nil { + b.Fatal(err) + } + defer target.Release() + view, err := device.CreateTextureView(target, nil) + if err != nil { + b.Fatal(err) + } + defer view.Release() + + frame := func() { + encoder, err := device.CreateCommandEncoder(&wgpu.CommandEncoderDescriptor{}) + if err != nil { + b.Fatal(err) + } + pass, err := encoder.BeginRenderPass(&wgpu.RenderPassDescriptor{ + ColorAttachments: []wgpu.RenderPassColorAttachment{{ + View: view, LoadOp: gputypes.LoadOpClear, StoreOp: gputypes.StoreOpStore, + }}, + }) + if err != nil { + b.Fatal(err) + } + pass.SetPipeline(pipeline) + pass.SetVertexBuffer(0, vertex, 0) + pass.SetIndexBuffer(index, gputypes.IndexFormatUint16, 0) + if count == 0 { + pass.DrawIndexed(3, 1, 0, 0, 0) + } else { + if forceLoop { + pass.Draw(0, 0, 0, 0) + } + pass.MultiDrawIndexedIndirect(indirect, 0, count) + } + if err := pass.End(); err != nil { + b.Fatal(err) + } + commands, err := encoder.Finish() + if err != nil { + b.Fatal(err) + } + if _, err := queue.Submit(commands); err != nil { + b.Fatal(err) + } + if err := device.WaitIdle(); err != nil { + b.Fatal(err) + } + } + + for range 5 { + frame() + } + b.ReportAllocs() + b.ResetTimer() + for range b.N { + frame() + } +} + +func benchmarkMetalBuffer(b *testing.B, device *wgpu.Device, queue *wgpu.Queue, data []byte, usage wgpu.BufferUsage) *wgpu.Buffer { + b.Helper() + buffer, err := device.CreateBuffer(&wgpu.BufferDescriptor{Size: uint64(len(data)), Usage: usage}) + if err != nil { + b.Fatal(err) + } + if err := queue.WriteBuffer(buffer, 0, data); err != nil { + buffer.Release() + b.Fatal(err) + } + return buffer +}