From 2a57db73cbfc35bc9446bdfee691359934b99340 Mon Sep 17 00:00:00 2001 From: Sarah Date: Wed, 23 Jun 2021 19:19:06 +0000 Subject: [PATCH] update compute shaders to include workgroup_size in their attribute Change-Id: Ia9027f8ec9c4a293ae1bef735c8b5eeea84e81e2 Bug: tint:884 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/55685 Reviewed-by: Kai Ninomiya Commit-Queue: Kai Ninomiya --- src/tests/DawnTest.cpp | 2 +- src/tests/end2end/BindGroupTests.cpp | 9 +++++---- src/tests/end2end/BufferZeroInitTests.cpp | 8 ++++---- .../end2end/ComputeCopyStorageBufferTests.cpp | 6 +++--- .../end2end/ComputeStorageBufferBarrierTests.cpp | 14 +++++++------- src/tests/end2end/CreatePipelineAsyncTests.cpp | 10 +++++----- src/tests/end2end/D3D12CachingTests.cpp | 4 ++-- src/tests/end2end/DeprecatedAPITests.cpp | 2 +- src/tests/end2end/DepthStencilSamplingTests.cpp | 5 +++-- src/tests/end2end/DeviceLostTests.cpp | 4 ++-- src/tests/end2end/DynamicBufferOffsetTests.cpp | 2 +- src/tests/end2end/EntryPointTests.cpp | 4 ++-- .../end2end/GpuMemorySynchronizationTests.cpp | 10 +++++----- src/tests/end2end/MultisampledSamplingTests.cpp | 2 +- src/tests/end2end/ObjectCachingTests.cpp | 8 ++++---- src/tests/end2end/OpArrayLengthTests.cpp | 2 +- src/tests/end2end/ShaderTests.cpp | 4 ++-- src/tests/end2end/StorageTextureTests.cpp | 14 +++++++------- src/tests/end2end/TextureZeroInitTests.cpp | 2 +- .../validation/BindGroupValidationTests.cpp | 4 ++-- .../GetBindGroupLayoutValidationTests.cpp | 6 +++--- .../unittests/validation/MultipleDeviceTests.cpp | 2 +- .../validation/QueueSubmitValidationTests.cpp | 6 +++--- .../validation/ResourceUsageTrackingTests.cpp | 2 +- .../validation/StorageTextureValidationTests.cpp | 8 ++++---- .../validation/UnsafeAPIValidationTests.cpp | 2 +- 26 files changed, 72 insertions(+), 70 deletions(-) diff --git a/src/tests/DawnTest.cpp b/src/tests/DawnTest.cpp index adec6d0a92..3b9b5a92ad 100644 --- a/src/tests/DawnTest.cpp +++ b/src/tests/DawnTest.cpp @@ -1104,7 +1104,7 @@ std::ostringstream& DawnTestBase::ExpectSampledDepthData(wgpu::Texture texture, [[group(0), binding(0)]] var tex : texture_depth_2d; [[group(0), binding(1)]] var result : Result; - [[stage(compute)]] fn main( + [[stage(compute), workgroup_size(1)]] fn main( [[builtin(global_invocation_id)]] GlobalInvocationId : vec3 ) { result.values[GlobalInvocationId.y * width + GlobalInvocationId.x] = textureLoad( diff --git a/src/tests/end2end/BindGroupTests.cpp b/src/tests/end2end/BindGroupTests.cpp index cfd0c70c4d..279c764108 100644 --- a/src/tests/end2end/BindGroupTests.cpp +++ b/src/tests/end2end/BindGroupTests.cpp @@ -127,7 +127,7 @@ TEST_P(BindGroupTests, ReusedBindGroupSingleSubmit) { }; [[group(0), binding(0)]] var contents: Contents; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { var f : f32 = contents.f; })"); @@ -841,7 +841,7 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) { [[group(0), binding(0)]] var buffer0 : Buffer0; [[group(0), binding(4)]] var outputBuffer : OutputBuffer; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { outputBuffer.value = vec3(buffer0.value, buffer2.value, buffer3.value); })"); pipelineDescriptor.compute.entryPoint = "main"; @@ -1066,7 +1066,7 @@ TEST_P(BindGroupTests, EmptyLayout) { pipelineDesc.layout = utils::MakeBasicPipelineLayout(device, &bgl); pipelineDesc.compute.entryPoint = "main"; pipelineDesc.compute.module = utils::CreateShaderModule(device, R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDesc); @@ -1260,7 +1260,8 @@ TEST_P(BindGroupTests, ReallyLargeBindGroup) { body << "result.value = 1u;\n"; - std::string shader = interface.str() + "[[stage(compute)]] fn main() {\n" + body.str() + "}\n"; + std::string shader = interface.str() + "[[stage(compute), workgroup_size(1)]] fn main() {\n" + + body.str() + "}\n"; wgpu::ComputePipelineDescriptor cpDesc; cpDesc.compute.module = utils::CreateShaderModule(device, shader.c_str()); cpDesc.compute.entryPoint = "main"; diff --git a/src/tests/end2end/BufferZeroInitTests.cpp b/src/tests/end2end/BufferZeroInitTests.cpp index 29e4933271..39bd98a314 100644 --- a/src/tests/end2end/BufferZeroInitTests.cpp +++ b/src/tests/end2end/BufferZeroInitTests.cpp @@ -427,7 +427,7 @@ class BufferZeroInitTest : public DawnTest { const char* computeShader = R"( [[group(0), binding(0)]] var outImage : texture_storage_2d; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { textureStore(outImage, vec2(0, 0), vec4(1.0, 0.0, 0.0, 1.0)); })"; @@ -999,7 +999,7 @@ TEST_P(BufferZeroInitTest, BoundAsUniformBuffer) { [[group(0), binding(0)]] var ubo : UBO; [[group(0), binding(1)]] var outImage : texture_storage_2d; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { if (all(ubo.value == vec4(0u, 0u, 0u, 0u))) { textureStore(outImage, vec2(0, 0), vec4(0.0, 1.0, 0.0, 1.0)); } else { @@ -1038,7 +1038,7 @@ TEST_P(BufferZeroInitTest, BoundAsReadonlyStorageBuffer) { [[group(0), binding(0)]] var ssbo : SSBO; [[group(0), binding(1)]] var outImage : texture_storage_2d; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { if (all(ssbo.value == vec4(0u, 0u, 0u, 0u))) { textureStore(outImage, vec2(0, 0), vec4(0.0, 1.0, 0.0, 1.0)); } else { @@ -1077,7 +1077,7 @@ TEST_P(BufferZeroInitTest, BoundAsStorageBuffer) { [[group(0), binding(0)]] var ssbo : SSBO; [[group(0), binding(1)]] var outImage : texture_storage_2d; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { if (all(ssbo.value[0] == vec4(0u, 0u, 0u, 0u)) && all(ssbo.value[1] == vec4(0u, 0u, 0u, 0u))) { textureStore(outImage, vec2(0, 0), vec4(0.0, 1.0, 0.0, 1.0)); diff --git a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp index 86e6aac4aa..a118c29506 100644 --- a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp +++ b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp @@ -99,7 +99,7 @@ TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfBasic) { [[set(0), binding(0)]] var src : Buf1; [[set(0), binding(1)]] var dst : Buf2; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { let index : u32 = GlobalInvocationID.x; if (index >= 4u) { return; } @@ -126,7 +126,7 @@ TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfStruct) { [[set(0), binding(0)]] var src : Buf1; [[set(0), binding(1)]] var dst : Buf2; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { let index : u32 = GlobalInvocationID.x; if (index >= 4u) { return; } @@ -148,7 +148,7 @@ TEST_P(ComputeCopyStorageBufferTests, UnsizedArrayOfBasic) { [[set(0), binding(0)]] var src : Buf1; [[set(0), binding(1)]] var dst : Buf2; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { let index : u32 = GlobalInvocationID.x; if (index >= 4u) { return; } diff --git a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp index 970fdec3eb..2cd7aa4c38 100644 --- a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp +++ b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp @@ -38,7 +38,7 @@ TEST_P(ComputeStorageBufferBarrierTests, AddIncrement) { [[group(0), binding(0)]] var buf : Buf; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { buf.data[GlobalInvocationID.x] = buf.data[GlobalInvocationID.x] + 0x1234u; } @@ -94,7 +94,7 @@ TEST_P(ComputeStorageBufferBarrierTests, AddPingPong) { [[group(0), binding(0)]] var src : Src; [[group(0), binding(1)]] var dst : Dst; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] + 0x1234u; } @@ -165,7 +165,7 @@ TEST_P(ComputeStorageBufferBarrierTests, StorageAndReadonlyStoragePingPongInOneP [[group(0), binding(0)]] var src : Src; [[group(0), binding(1)]] var dst : Dst; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] + 0x1234u; } @@ -233,7 +233,7 @@ TEST_P(ComputeStorageBufferBarrierTests, UniformToStorageAddPingPong) { [[group(0), binding(0)]] var src : Buf; [[group(0), binding(1)]] var dst : Buf; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] + vec4(0x1234u, 0x1234u, 0x1234u, 0x1234u); @@ -301,7 +301,7 @@ TEST_P(ComputeStorageBufferBarrierTests, UniformToStorageAddPingPongInOnePass) { [[group(0), binding(0)]] var src : Buf; [[group(0), binding(1)]] var dst : Buf; - [[stage(compute)]] + [[stage(compute), workgroup_size(1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3) { dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] + vec4(0x1234u, 0x1234u, 0x1234u, 0x1234u); @@ -361,7 +361,7 @@ TEST_P(ComputeStorageBufferBarrierTests, IndirectBufferCorrectBarrier) { }; [[group(0), binding(0)]] var buf : Buf; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { buf.data = array(1u, 1u, 1u); } )"); @@ -380,7 +380,7 @@ TEST_P(ComputeStorageBufferBarrierTests, IndirectBufferCorrectBarrier) { }; [[group(0), binding(1)]] var result : Result; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { result.data = 2u; if (buf.data[0] == 1u && buf.data[1] == 1u && buf.data[2] == 1u) { result.data = 1u; diff --git a/src/tests/end2end/CreatePipelineAsyncTests.cpp b/src/tests/end2end/CreatePipelineAsyncTests.cpp index ea4b3e1017..2023884b56 100644 --- a/src/tests/end2end/CreatePipelineAsyncTests.cpp +++ b/src/tests/end2end/CreatePipelineAsyncTests.cpp @@ -80,7 +80,7 @@ TEST_P(CreatePipelineAsyncTest, BasicUseOfCreateComputePipelineAsync) { }; [[group(0), binding(0)]] var ssbo : SSBO; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { ssbo.value = 1u; })"); csDesc.compute.entryPoint = "main"; @@ -115,7 +115,7 @@ TEST_P(CreatePipelineAsyncTest, CreateComputePipelineFailed) { }; [[group(0), binding(0)]] var ssbo : SSBO; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { ssbo.value = 1u; })"); csDesc.compute.entryPoint = "main0"; @@ -253,7 +253,7 @@ TEST_P(CreatePipelineAsyncTest, CreateRenderPipelineFailed) { TEST_P(CreatePipelineAsyncTest, ReleaseDeviceBeforeCallbackOfCreateComputePipelineAsync) { wgpu::ComputePipelineDescriptor csDesc; csDesc.compute.module = utils::CreateShaderModule(device, R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); csDesc.compute.entryPoint = "main"; @@ -314,7 +314,7 @@ TEST_P(CreatePipelineAsyncTest, CreateSameComputePipelineTwice) { }; [[group(0), binding(0)]] var ssbo : SSBO; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { ssbo.value = 1u; })"); csDesc.compute.entryPoint = "main"; @@ -355,7 +355,7 @@ TEST_P(CreatePipelineAsyncTest, CreateSamePipelineTwiceAtSameTime) { }; [[group(0), binding(0)]] var ssbo : SSBO; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { ssbo.value = 1u; })"); csDesc.compute.entryPoint = "main"; diff --git a/src/tests/end2end/D3D12CachingTests.cpp b/src/tests/end2end/D3D12CachingTests.cpp index d3c72ebb92..00eab66cb5 100644 --- a/src/tests/end2end/D3D12CachingTests.cpp +++ b/src/tests/end2end/D3D12CachingTests.cpp @@ -213,11 +213,11 @@ TEST_P(D3D12CachingTests, ReuseShaderWithMultipleEntryPoints) { }; [[binding(0), group(0)]] var data : Data; - [[stage(compute)]] fn write1() { + [[stage(compute), workgroup_size(1)]] fn write1() { data.data = 1u; } - [[stage(compute)]] fn write42() { + [[stage(compute), workgroup_size(1)]] fn write42() { data.data = 42u; } )"); diff --git a/src/tests/end2end/DeprecatedAPITests.cpp b/src/tests/end2end/DeprecatedAPITests.cpp index 83346a7494..645214bfaa 100644 --- a/src/tests/end2end/DeprecatedAPITests.cpp +++ b/src/tests/end2end/DeprecatedAPITests.cpp @@ -89,7 +89,7 @@ TEST_P(DeprecationTests, SetAttachmentDescriptorAttachment) { TEST_P(DeprecationTests, ComputeStage) { wgpu::ComputePipelineDescriptor csDesc; csDesc.computeStage.module = utils::CreateShaderModule(device, R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); csDesc.computeStage.entryPoint = "main"; diff --git a/src/tests/end2end/DepthStencilSamplingTests.cpp b/src/tests/end2end/DepthStencilSamplingTests.cpp index c6feabee89..ad727c30a8 100644 --- a/src/tests/end2end/DepthStencilSamplingTests.cpp +++ b/src/tests/end2end/DepthStencilSamplingTests.cpp @@ -163,7 +163,8 @@ class DepthStencilSamplingTest : public DawnTest { index++; } - shaderSource << "[[stage(compute)]] fn main() { " << shaderBody.str() << "\n}"; + shaderSource << "[[stage(compute), workgroup_size(1)]] fn main() { " << shaderBody.str() + << "\n}"; wgpu::ShaderModule csModule = utils::CreateShaderModule(device, shaderSource.str().c_str()); @@ -223,7 +224,7 @@ class DepthStencilSamplingTest : public DawnTest { }; [[group(0), binding(3)]] var samplerResult : SamplerResult; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { samplerResult.value = textureSampleCompare(tex, samp, vec2(0.5, 0.5), uniforms.compareRef); })"); diff --git a/src/tests/end2end/DeviceLostTests.cpp b/src/tests/end2end/DeviceLostTests.cpp index 8513c16578..3dd136fc95 100644 --- a/src/tests/end2end/DeviceLostTests.cpp +++ b/src/tests/end2end/DeviceLostTests.cpp @@ -110,7 +110,7 @@ TEST_P(DeviceLostTest, GetBindGroupLayoutFails) { pos : vec4; }; [[group(0), binding(0)]] var ubo : UniformBuffer; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); wgpu::ComputePipelineDescriptor descriptor; @@ -448,7 +448,7 @@ TEST_P(DeviceLostTest, DeviceLostDoesntCallUncapturedError) { // before the callback of Create*PipelineAsync() is called. TEST_P(DeviceLostTest, DeviceLostBeforeCreatePipelineAsyncCallback) { wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); wgpu::ComputePipelineDescriptor descriptor; diff --git a/src/tests/end2end/DynamicBufferOffsetTests.cpp b/src/tests/end2end/DynamicBufferOffsetTests.cpp index d8ae96086f..f735a55b42 100644 --- a/src/tests/end2end/DynamicBufferOffsetTests.cpp +++ b/src/tests/end2end/DynamicBufferOffsetTests.cpp @@ -209,7 +209,7 @@ class DynamicBufferOffsetTests : public DawnTest { cs << "let multipleNumber : u32 = " << multipleNumber << "u;\n"; cs << R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { sBufferNotDynamic.value = uBufferNotDynamic.value.xy; sBuffer.value = vec2(multipleNumber, multipleNumber) * (uBuffer.value.xy + sBufferNotDynamic.value.xy); } diff --git a/src/tests/end2end/EntryPointTests.cpp b/src/tests/end2end/EntryPointTests.cpp index c97244db6d..b951341ce8 100644 --- a/src/tests/end2end/EntryPointTests.cpp +++ b/src/tests/end2end/EntryPointTests.cpp @@ -66,12 +66,12 @@ TEST_P(EntryPointTests, TwoComputeInModule) { }; [[binding(0), group(0)]] var data : Data; - [[stage(compute)]] fn write1() { + [[stage(compute), workgroup_size(1)]] fn write1() { data.data = 1u; return; } - [[stage(compute)]] fn write42() { + [[stage(compute), workgroup_size(1)]] fn write42() { data.data = 42u; return; } diff --git a/src/tests/end2end/GpuMemorySynchronizationTests.cpp b/src/tests/end2end/GpuMemorySynchronizationTests.cpp index 035b647ec1..4989ce1baf 100644 --- a/src/tests/end2end/GpuMemorySynchronizationTests.cpp +++ b/src/tests/end2end/GpuMemorySynchronizationTests.cpp @@ -40,7 +40,7 @@ class GpuMemorySyncTests : public DawnTest { a : i32; }; [[group(0), binding(0)]] var data : Data; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { data.a = data.a + 1; })"); @@ -259,7 +259,7 @@ TEST_P(GpuMemorySyncTests, SampledAndROStorageTextureInComputePass) { [[group(0), binding(1)]] var sampledTex : texture_2d; [[group(0), binding(2)]] var storageTex : texture_storage_2d; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { output.sampledOut = textureLoad(sampledTex, vec2(0, 0), 0).x; output.storageOut = textureLoad(storageTex, vec2(0, 0)).x; } @@ -316,7 +316,7 @@ class StorageToUniformSyncTests : public DawnTest { a : f32; }; [[group(0), binding(0)]] var data : Data; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { data.a = 1.0; })"); @@ -529,7 +529,7 @@ TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) { [[group(0), binding(2)]] var uniformContents : ColorContents1; [[group(0), binding(3)]] var storageContents : ColorContents2; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { vbContents.pos[0] = vec4(-1.0, 1.0, 0.0, 1.0); vbContents.pos[1] = vec4(1.0, 1.0, 0.0, 1.0); vbContents.pos[2] = vec4(1.0, -1.0, 0.0, 1.0); @@ -644,7 +644,7 @@ TEST_P(MultipleWriteThenMultipleReadTests, OneBuffer) { [[group(0), binding(0)]] var contents : Contents; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { contents.pos[0] = vec4(-1.0, 1.0, 0.0, 1.0); contents.pos[1] = vec4(1.0, 1.0, 0.0, 1.0); contents.pos[2] = vec4(1.0, -1.0, 0.0, 1.0); diff --git a/src/tests/end2end/MultisampledSamplingTests.cpp b/src/tests/end2end/MultisampledSamplingTests.cpp index 4d913c18c3..fe4ad4eddf 100644 --- a/src/tests/end2end/MultisampledSamplingTests.cpp +++ b/src/tests/end2end/MultisampledSamplingTests.cpp @@ -102,7 +102,7 @@ class MultisampledSamplingTest : public DawnTest { }; [[group(0), binding(2)]] var results : Results; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { for (var i : i32 = 0; i < 4; i = i + 1) { results.colorSamples[i] = textureLoad(texture0, vec2(0, 0), i).x; results.depthSamples[i] = textureLoad(texture1, vec2(0, 0), i).x; diff --git a/src/tests/end2end/ObjectCachingTests.cpp b/src/tests/end2end/ObjectCachingTests.cpp index fb6c2ddc2b..a42e7f0ce3 100644 --- a/src/tests/end2end/ObjectCachingTests.cpp +++ b/src/tests/end2end/ObjectCachingTests.cpp @@ -124,16 +124,16 @@ TEST_P(ObjectCachingTest, ShaderModuleDeduplication) { TEST_P(ObjectCachingTest, ComputePipelineDeduplicationOnShaderModule) { wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( var i : u32; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { i = 0u; })"); wgpu::ShaderModule sameModule = utils::CreateShaderModule(device, R"( var i : u32; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { i = 0u; })"); wgpu::ShaderModule otherModule = utils::CreateShaderModule(device, R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); EXPECT_NE(module.Get(), otherModule.Get()); @@ -176,7 +176,7 @@ TEST_P(ObjectCachingTest, ComputePipelineDeduplicationOnLayout) { desc.compute.entryPoint = "main"; desc.compute.module = utils::CreateShaderModule(device, R"( var i : u32; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { i = 0u; })"); diff --git a/src/tests/end2end/OpArrayLengthTests.cpp b/src/tests/end2end/OpArrayLengthTests.cpp index 3a219b4ce4..2f2ad52af7 100644 --- a/src/tests/end2end/OpArrayLengthTests.cpp +++ b/src/tests/end2end/OpArrayLengthTests.cpp @@ -130,7 +130,7 @@ TEST_P(OpArrayLengthTest, Compute) { }; [[group(1), binding(0)]] var result : ResultBuffer; )" + mShaderInterface + R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { result.data[0] = arrayLength(&buffer1.data); result.data[1] = arrayLength(&buffer2.data); result.data[2] = arrayLength(&buffer3.data); diff --git a/src/tests/end2end/ShaderTests.cpp b/src/tests/end2end/ShaderTests.cpp index 590e4859c7..4fb95b3981 100644 --- a/src/tests/end2end/ShaderTests.cpp +++ b/src/tests/end2end/ShaderTests.cpp @@ -37,7 +37,7 @@ TEST_P(ShaderTests, ComputeLog2) { [[group(0), binding(0)]] var buf : Buf; -[[stage(compute)]] fn main() { +[[stage(compute), workgroup_size(1)]] fn main() { let factor : f32 = 1.0001; buf.data[0] = u32(log2(1.0 * factor)); @@ -313,7 +313,7 @@ TEST_P(ShaderTests, PipelineOverridableUsed) { std::string shader = R"( [[override]] let foo : f32; -[[stage(compute)]] +[[stage(compute), workgroup_size(1)]] fn ep_func() { var local_foo : f32; local_foo = foo; diff --git a/src/tests/end2end/StorageTextureTests.cpp b/src/tests/end2end/StorageTextureTests.cpp index 4eed7e33d3..d6e966a2cf 100644 --- a/src/tests/end2end/StorageTextureTests.cpp +++ b/src/tests/end2end/StorageTextureTests.cpp @@ -358,7 +358,7 @@ fn IsEqualTo(pixel : vec4, expected : vec4) -> bool { std::ostringstream ostream; ostream << GetImageDeclaration(format, "write", is2DArray, 0) << "\n"; ostream << GetImageDeclaration(format, "read", is2DArray, 1) << "\n"; - ostream << "[[stage(compute)]] fn main() {\n"; + ostream << "[[stage(compute), workgroup_size(1)]] fn main() {\n"; ostream << " let size : vec2 = textureDimensions(storageImage0);\n"; ostream << " let layerCount : i32 = " << layerCount << ";\n"; ostream << " for (var layer : i32 = 0; layer < layerCount; layer = layer + 1) {\n"; @@ -713,7 +713,7 @@ TEST_P(StorageTextureTests, ReadonlyStorageTextureInComputeShader) { [[group(0), binding(1)]] var dstBuffer : DstBuffer; )" << CommonReadOnlyTestCode(format) << R"( -[[stage(compute)]] fn main() { +[[stage(compute), workgroup_size(1)]] fn main() { if (doTest()) { dstBuffer.result = 1u; } else { @@ -937,7 +937,7 @@ TEST_P(StorageTextureTests, Readonly2DArrayStorageTexture) { [[group(0), binding(1)]] var dstBuffer : DstBuffer; )" << CommonReadOnlyTestCode(kTextureFormat, true) << R"( -[[stage(compute)]] fn main() { +[[stage(compute), workgroup_size(1)]] fn main() { if (doTest()) { dstBuffer.result = 1u; } else { @@ -982,7 +982,7 @@ TEST_P(StorageTextureTests, ReadonlyAndWriteonlyStorageTexturePingPong) { wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( [[group(0), binding(0)]] var Src : texture_storage_2d; [[group(0), binding(1)]] var Dst : texture_storage_2d; -[[stage(compute)]] fn main() { +[[stage(compute), workgroup_size(1)]] fn main() { var srcValue : vec4 = textureLoad(Src, vec2(0, 0)); srcValue.x = srcValue.x + 1u; textureStore(Dst, vec2(0, 0), srcValue); @@ -1056,7 +1056,7 @@ TEST_P(StorageTextureTests, SampledAndWriteonlyStorageTexturePingPong) { wgpu::ShaderModule module = utils::CreateShaderModule(device, R"( [[group(0), binding(0)]] var Src : texture_2d; [[group(0), binding(1)]] var Dst : texture_storage_2d; -[[stage(compute)]] fn main() { +[[stage(compute), workgroup_size(1)]] fn main() { var srcValue : vec4 = textureLoad(Src, vec2(0, 0), 0); srcValue.x = srcValue.x + 1u; textureStore(Dst, vec2(0, 0), srcValue); @@ -1161,7 +1161,7 @@ fn doTest() -> bool { const char* kCommonWriteOnlyZeroInitTestCodeCompute = R"( [[group(0), binding(0)]] var dstImage : texture_storage_2d; -[[stage(compute)]] fn main() { +[[stage(compute), workgroup_size(1)]] fn main() { textureStore(dstImage, vec2(0, 0), vec4(1u, 0u, 0u, 1u)); })"; }; @@ -1204,7 +1204,7 @@ TEST_P(StorageTextureZeroInitTests, ReadonlyStorageTextureClearsToZeroInComputeP [[group(0), binding(0)]] var srcImage : texture_storage_2d; [[group(0), binding(1)]] var dstBuffer : DstBuffer; )") + kCommonReadOnlyZeroInitTestCode + R"( -[[stage(compute)]] fn main() { +[[stage(compute), workgroup_size(1)]] fn main() { if (doTest()) { dstBuffer.result = 1u; } else { diff --git a/src/tests/end2end/TextureZeroInitTests.cpp b/src/tests/end2end/TextureZeroInitTests.cpp index 00d61f7ab8..8c657eaff4 100644 --- a/src/tests/end2end/TextureZeroInitTests.cpp +++ b/src/tests/end2end/TextureZeroInitTests.cpp @@ -979,7 +979,7 @@ TEST_P(TextureZeroInitTest, ComputePassSampledTextureClear) { value : vec4; }; [[group(0), binding(1)]] var result : Result; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { result.value = textureLoad(tex, vec2(0,0), 0); } )"; diff --git a/src/tests/unittests/validation/BindGroupValidationTests.cpp b/src/tests/unittests/validation/BindGroupValidationTests.cpp index 759b2659c2..1057567807 100644 --- a/src/tests/unittests/validation/BindGroupValidationTests.cpp +++ b/src/tests/unittests/validation/BindGroupValidationTests.cpp @@ -1981,7 +1981,7 @@ TEST_F(BindGroupLayoutCompatibilityTest, TextureViewDimension) { })"; constexpr char kTexture2DShaderCS[] = R"( [[group(0), binding(0)]] var myTexture : texture_2d; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { ignore(textureDimensions(myTexture)); })"; @@ -2020,7 +2020,7 @@ TEST_F(BindGroupLayoutCompatibilityTest, TextureViewDimension) { })"; constexpr char kTexture2DArrayShaderCS[] = R"( [[group(0), binding(0)]] var myTexture : texture_2d_array; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { ignore(textureDimensions(myTexture)); })"; diff --git a/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp b/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp index 987599ee01..13d4ddd839 100644 --- a/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp +++ b/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp @@ -152,7 +152,7 @@ TEST_F(GetBindGroupLayoutTests, ComputePipeline) { }; [[group(0), binding(0)]] var uniforms : S; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { var pos : vec4 = uniforms.pos; })"); @@ -915,11 +915,11 @@ TEST_F(GetBindGroupLayoutTests, FromCorrectEntryPoint) { [[group(0), binding(0)]] var data0 : Data; [[group(0), binding(1)]] var data1 : Data; - [[stage(compute)]] fn compute0() { + [[stage(compute), workgroup_size(1)]] fn compute0() { data0.data = 0.0; } - [[stage(compute)]] fn compute1() { + [[stage(compute), workgroup_size(1)]] fn compute1() { data1.data = 0.0; } )"); diff --git a/src/tests/unittests/validation/MultipleDeviceTests.cpp b/src/tests/unittests/validation/MultipleDeviceTests.cpp index 8e9637dd8c..c94e94aa88 100644 --- a/src/tests/unittests/validation/MultipleDeviceTests.cpp +++ b/src/tests/unittests/validation/MultipleDeviceTests.cpp @@ -33,7 +33,7 @@ TEST_F(MultipleDeviceTest, ValidatesSameDevice) { TEST_F(MultipleDeviceTest, ValidatesSameDeviceCreatePipelineAsync) { wgpu::ShaderModuleWGSLDescriptor wgslDesc = {}; wgslDesc.source = R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1, 1, 1)]] fn main() { } )"; diff --git a/src/tests/unittests/validation/QueueSubmitValidationTests.cpp b/src/tests/unittests/validation/QueueSubmitValidationTests.cpp index 0b6d591694..3f5ad4bfb1 100644 --- a/src/tests/unittests/validation/QueueSubmitValidationTests.cpp +++ b/src/tests/unittests/validation/QueueSubmitValidationTests.cpp @@ -211,7 +211,7 @@ namespace { wgpu::ComputePipelineDescriptor descriptor; descriptor.compute.module = utils::CreateShaderModule(device, R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); descriptor.compute.entryPoint = "main"; device.CreateComputePipelineAsync(&descriptor, callback, &callbackData); @@ -236,7 +236,7 @@ namespace { cpDesc.layout = utils::MakePipelineLayout(device, {emptyBGL, testBGL}); cpDesc.compute.entryPoint = "main"; cpDesc.compute.module = - utils::CreateShaderModule(device, "[[stage(compute)]] fn main() {}"); + utils::CreateShaderModule(device, "[[stage(compute), workgroup_size(1)]] fn main() {}"); wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cpDesc); wgpu::BufferDescriptor bufDesc; @@ -304,7 +304,7 @@ namespace { cpDesc.layout = utils::MakePipelineLayout(device, {emptyBGL, emptyBGL, testBGL}); cpDesc.compute.entryPoint = "main"; cpDesc.compute.module = - utils::CreateShaderModule(device, "[[stage(compute)]] fn main() {}"); + utils::CreateShaderModule(device, "[[stage(compute), workgroup_size(1)]] fn main() {}"); wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cpDesc); wgpu::TextureDescriptor texDesc; diff --git a/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp b/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp index eed4393e80..ffc1fe46cd 100644 --- a/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp +++ b/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp @@ -63,7 +63,7 @@ namespace { wgpu::ComputePipeline CreateNoOpComputePipeline(std::vector bgls) { wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"( - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { })"); wgpu::ComputePipelineDescriptor pipelineDescriptor; pipelineDescriptor.layout = utils::MakePipelineLayout(device, std::move(bgls)); diff --git a/src/tests/unittests/validation/StorageTextureValidationTests.cpp b/src/tests/unittests/validation/StorageTextureValidationTests.cpp index b532492320..919ad601a3 100644 --- a/src/tests/unittests/validation/StorageTextureValidationTests.cpp +++ b/src/tests/unittests/validation/StorageTextureValidationTests.cpp @@ -85,7 +85,7 @@ class StorageTextureValidationTests : public ValidationTest { ostream << "[[group(0), binding(0)]] var image0 : " << imageTypeDeclaration << "<" << imageFormatQualifier << ", " << access << ">;\n" - "[[stage(compute)]] fn main() {\n" + "[[stage(compute), workgroup_size(1)]] fn main() {\n" " ignore(textureDimensions(image0));\n" "}\n"; @@ -194,7 +194,7 @@ TEST_F(StorageTextureValidationTests, ComputePipeline) { }; [[group(0), binding(1)]] var buf : Buf; - [[stage(compute)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3) { + [[stage(compute), workgroup_size(1)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3) { buf.data = textureLoad(image0, vec2(LocalInvocationID.xy)).x; })"); @@ -211,7 +211,7 @@ TEST_F(StorageTextureValidationTests, ComputePipeline) { wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"( [[group(0), binding(0)]] var image0 : texture_storage_2d; - [[stage(compute)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3) { + [[stage(compute), workgroup_size(1)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3) { textureStore(image0, vec2(LocalInvocationID.xy), vec4(0.0, 0.0, 0.0, 0.0)); })"); @@ -248,7 +248,7 @@ TEST_F(StorageTextureValidationTests, ReadWriteStorageTexture) { { ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"( [[group(0), binding(0)]] var image0 : texture_storage_2d; - [[stage(compute)]] fn main() { + [[stage(compute), workgroup_size(1)]] fn main() { ignore(textureDimensions(image0)); })")); } diff --git a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp index 12bd37d979..61b327154a 100644 --- a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp +++ b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp @@ -136,7 +136,7 @@ TEST_F(UnsafeAPIValidationTest, DispatchIndirectDisallowed) { wgpu::ComputePipelineDescriptor pipelineDesc; pipelineDesc.compute.entryPoint = "main"; pipelineDesc.compute.module = - utils::CreateShaderModule(device, "[[stage(compute)]] fn main() {}"); + utils::CreateShaderModule(device, "[[stage(compute), workgroup_size(1)]] fn main() {}"); wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDesc); // Control case: dispatch is allowed.