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 <kainino@chromium.org>
Commit-Queue: Kai Ninomiya <kainino@chromium.org>
diff --git a/src/tests/DawnTest.cpp b/src/tests/DawnTest.cpp
index adec6d0..3b9b5a9 100644
--- a/src/tests/DawnTest.cpp
+++ b/src/tests/DawnTest.cpp
@@ -1104,7 +1104,7 @@
[[group(0), binding(0)]] var tex : texture_depth_2d;
[[group(0), binding(1)]] var<storage, read_write> result : Result;
- [[stage(compute)]] fn main(
+ [[stage(compute), workgroup_size(1)]] fn main(
[[builtin(global_invocation_id)]] GlobalInvocationId : vec3<u32>
) {
result.values[GlobalInvocationId.y * width + GlobalInvocationId.x] = textureLoad(
diff --git a/src/tests/end2end/BindGroupTests.cpp b/src/tests/end2end/BindGroupTests.cpp
index cfd0c70..279c764 100644
--- a/src/tests/end2end/BindGroupTests.cpp
+++ b/src/tests/end2end/BindGroupTests.cpp
@@ -127,7 +127,7 @@
};
[[group(0), binding(0)]] var <uniform> contents: Contents;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
var f : f32 = contents.f;
})");
@@ -841,7 +841,7 @@
[[group(0), binding(0)]] var<storage, read> buffer0 : Buffer0;
[[group(0), binding(4)]] var<storage, read_write> outputBuffer : OutputBuffer;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
outputBuffer.value = vec3<u32>(buffer0.value, buffer2.value, buffer3.value);
})");
pipelineDescriptor.compute.entryPoint = "main";
@@ -1066,7 +1066,7 @@
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 @@
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 29e4933..39bd98a 100644
--- a/src/tests/end2end/BufferZeroInitTests.cpp
+++ b/src/tests/end2end/BufferZeroInitTests.cpp
@@ -427,7 +427,7 @@
const char* computeShader = R"(
[[group(0), binding(0)]] var outImage : texture_storage_2d<rgba8unorm, write>;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(1.0, 0.0, 0.0, 1.0));
})";
@@ -999,7 +999,7 @@
[[group(0), binding(0)]] var<uniform> ubo : UBO;
[[group(0), binding(1)]] var outImage : texture_storage_2d<rgba8unorm, write>;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
if (all(ubo.value == vec4<u32>(0u, 0u, 0u, 0u))) {
textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(0.0, 1.0, 0.0, 1.0));
} else {
@@ -1038,7 +1038,7 @@
[[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
[[group(0), binding(1)]] var outImage : texture_storage_2d<rgba8unorm, write>;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
if (all(ssbo.value == vec4<u32>(0u, 0u, 0u, 0u))) {
textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(0.0, 1.0, 0.0, 1.0));
} else {
@@ -1077,7 +1077,7 @@
[[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
[[group(0), binding(1)]] var outImage : texture_storage_2d<rgba8unorm, write>;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
if (all(ssbo.value[0] == vec4<u32>(0u, 0u, 0u, 0u)) &&
all(ssbo.value[1] == vec4<u32>(0u, 0u, 0u, 0u))) {
textureStore(outImage, vec2<i32>(0, 0), vec4<f32>(0.0, 1.0, 0.0, 1.0));
diff --git a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp
index 86e6aac..a118c29 100644
--- a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp
+++ b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp
@@ -99,7 +99,7 @@
[[set(0), binding(0)]] var<storage, read_write> src : Buf1;
[[set(0), binding(1)]] var<storage, read_write> dst : Buf2;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
let index : u32 = GlobalInvocationID.x;
if (index >= 4u) { return; }
@@ -126,7 +126,7 @@
[[set(0), binding(0)]] var<storage, read_write> src : Buf1;
[[set(0), binding(1)]] var<storage, read_write> dst : Buf2;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
let index : u32 = GlobalInvocationID.x;
if (index >= 4u) { return; }
@@ -148,7 +148,7 @@
[[set(0), binding(0)]] var<storage, read_write> src : Buf1;
[[set(0), binding(1)]] var<storage, read_write> dst : Buf2;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
let index : u32 = GlobalInvocationID.x;
if (index >= 4u) { return; }
diff --git a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
index 970fdec..2cd7aa4 100644
--- a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
+++ b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
@@ -38,7 +38,7 @@
[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
buf.data[GlobalInvocationID.x] = buf.data[GlobalInvocationID.x] + 0x1234u;
}
@@ -94,7 +94,7 @@
[[group(0), binding(0)]] var<storage, read_write> src : Src;
[[group(0), binding(1)]] var<storage, read_write> dst : Dst;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] + 0x1234u;
}
@@ -165,7 +165,7 @@
[[group(0), binding(0)]] var<storage, read> src : Src;
[[group(0), binding(1)]] var<storage, read_write> dst : Dst;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] + 0x1234u;
}
@@ -233,7 +233,7 @@
[[group(0), binding(0)]] var<uniform> src : Buf;
[[group(0), binding(1)]] var<storage, read_write> dst : Buf;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] +
vec4<u32>(0x1234u, 0x1234u, 0x1234u, 0x1234u);
@@ -301,7 +301,7 @@
[[group(0), binding(0)]] var<uniform> src : Buf;
[[group(0), binding(1)]] var<storage, read_write> dst : Buf;
- [[stage(compute)]]
+ [[stage(compute), workgroup_size(1)]]
fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
dst.data[GlobalInvocationID.x] = src.data[GlobalInvocationID.x] +
vec4<u32>(0x1234u, 0x1234u, 0x1234u, 0x1234u);
@@ -361,7 +361,7 @@
};
[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
buf.data = array<u32, 3>(1u, 1u, 1u);
}
)");
@@ -380,7 +380,7 @@
};
[[group(0), binding(1)]] var<storage, read_write> 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 ea4b3e1..2023884 100644
--- a/src/tests/end2end/CreatePipelineAsyncTests.cpp
+++ b/src/tests/end2end/CreatePipelineAsyncTests.cpp
@@ -80,7 +80,7 @@
};
[[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
ssbo.value = 1u;
})");
csDesc.compute.entryPoint = "main";
@@ -115,7 +115,7 @@
};
[[group(0), binding(0)]] var<storage, read_write> 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, 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 @@
};
[[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
ssbo.value = 1u;
})");
csDesc.compute.entryPoint = "main";
@@ -355,7 +355,7 @@
};
[[group(0), binding(0)]] var<storage, read_write> 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 d3c72eb..00eab66 100644
--- a/src/tests/end2end/D3D12CachingTests.cpp
+++ b/src/tests/end2end/D3D12CachingTests.cpp
@@ -213,11 +213,11 @@
};
[[binding(0), group(0)]] var<storage, read_write> 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 83346a7..645214b 100644
--- a/src/tests/end2end/DeprecatedAPITests.cpp
+++ b/src/tests/end2end/DeprecatedAPITests.cpp
@@ -89,7 +89,7 @@
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 c6feabe..ad727c30 100644
--- a/src/tests/end2end/DepthStencilSamplingTests.cpp
+++ b/src/tests/end2end/DepthStencilSamplingTests.cpp
@@ -163,7 +163,8 @@
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 @@
};
[[group(0), binding(3)]] var<storage, read_write> samplerResult : SamplerResult;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
samplerResult.value = textureSampleCompare(tex, samp, vec2<f32>(0.5, 0.5), uniforms.compareRef);
})");
diff --git a/src/tests/end2end/DeviceLostTests.cpp b/src/tests/end2end/DeviceLostTests.cpp
index 8513c16..3dd136f 100644
--- a/src/tests/end2end/DeviceLostTests.cpp
+++ b/src/tests/end2end/DeviceLostTests.cpp
@@ -110,7 +110,7 @@
pos : vec4<f32>;
};
[[group(0), binding(0)]] var<uniform> ubo : UniformBuffer;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
})");
wgpu::ComputePipelineDescriptor descriptor;
@@ -448,7 +448,7 @@
// 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 d8ae960..f735a55 100644
--- a/src/tests/end2end/DynamicBufferOffsetTests.cpp
+++ b/src/tests/end2end/DynamicBufferOffsetTests.cpp
@@ -209,7 +209,7 @@
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<u32>(multipleNumber, multipleNumber) * (uBuffer.value.xy + sBufferNotDynamic.value.xy);
}
diff --git a/src/tests/end2end/EntryPointTests.cpp b/src/tests/end2end/EntryPointTests.cpp
index c97244d..b951341 100644
--- a/src/tests/end2end/EntryPointTests.cpp
+++ b/src/tests/end2end/EntryPointTests.cpp
@@ -66,12 +66,12 @@
};
[[binding(0), group(0)]] var<storage, read_write> 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 035b647..4989ce1 100644
--- a/src/tests/end2end/GpuMemorySynchronizationTests.cpp
+++ b/src/tests/end2end/GpuMemorySynchronizationTests.cpp
@@ -40,7 +40,7 @@
a : i32;
};
[[group(0), binding(0)]] var<storage, read_write> data : Data;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
data.a = data.a + 1;
})");
@@ -259,7 +259,7 @@
[[group(0), binding(1)]] var sampledTex : texture_2d<u32>;
[[group(0), binding(2)]] var storageTex : texture_storage_2d<r32uint, read>;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
output.sampledOut = textureLoad(sampledTex, vec2<i32>(0, 0), 0).x;
output.storageOut = textureLoad(storageTex, vec2<i32>(0, 0)).x;
}
@@ -316,7 +316,7 @@
a : f32;
};
[[group(0), binding(0)]] var<storage, read_write> data : Data;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
data.a = 1.0;
})");
@@ -529,7 +529,7 @@
[[group(0), binding(2)]] var<storage, read_write> uniformContents : ColorContents1;
[[group(0), binding(3)]] var<storage, read_write> storageContents : ColorContents2;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
vbContents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0);
vbContents.pos[1] = vec4<f32>(1.0, 1.0, 0.0, 1.0);
vbContents.pos[2] = vec4<f32>(1.0, -1.0, 0.0, 1.0);
@@ -644,7 +644,7 @@
[[group(0), binding(0)]] var<storage, read_write> contents : Contents;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
contents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0);
contents.pos[1] = vec4<f32>(1.0, 1.0, 0.0, 1.0);
contents.pos[2] = vec4<f32>(1.0, -1.0, 0.0, 1.0);
diff --git a/src/tests/end2end/MultisampledSamplingTests.cpp b/src/tests/end2end/MultisampledSamplingTests.cpp
index 4d913c1..fe4ad4e 100644
--- a/src/tests/end2end/MultisampledSamplingTests.cpp
+++ b/src/tests/end2end/MultisampledSamplingTests.cpp
@@ -102,7 +102,7 @@
};
[[group(0), binding(2)]] var<storage, read_write> 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<i32>(0, 0), i).x;
results.depthSamples[i] = textureLoad(texture1, vec2<i32>(0, 0), i).x;
diff --git a/src/tests/end2end/ObjectCachingTests.cpp b/src/tests/end2end/ObjectCachingTests.cpp
index fb6c2dd..a42e7f0 100644
--- a/src/tests/end2end/ObjectCachingTests.cpp
+++ b/src/tests/end2end/ObjectCachingTests.cpp
@@ -124,16 +124,16 @@
TEST_P(ObjectCachingTest, ComputePipelineDeduplicationOnShaderModule) {
wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
var<workgroup> i : u32;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
i = 0u;
})");
wgpu::ShaderModule sameModule = utils::CreateShaderModule(device, R"(
var<workgroup> 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 @@
desc.compute.entryPoint = "main";
desc.compute.module = utils::CreateShaderModule(device, R"(
var<workgroup> 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 3a219b4..2f2ad52 100644
--- a/src/tests/end2end/OpArrayLengthTests.cpp
+++ b/src/tests/end2end/OpArrayLengthTests.cpp
@@ -130,7 +130,7 @@
};
[[group(1), binding(0)]] var<storage, read_write> 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 590e485..4fb95b3 100644
--- a/src/tests/end2end/ShaderTests.cpp
+++ b/src/tests/end2end/ShaderTests.cpp
@@ -37,7 +37,7 @@
[[group(0), binding(0)]] var<storage, read_write> 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 @@
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 4eed7e3..d6e966a 100644
--- a/src/tests/end2end/StorageTextureTests.cpp
+++ b/src/tests/end2end/StorageTextureTests.cpp
@@ -358,7 +358,7 @@
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<i32> = 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 @@
[[group(0), binding(1)]] var<storage, read_write> 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 @@
[[group(0), binding(1)]] var<storage, read_write> 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 @@
wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
[[group(0), binding(0)]] var Src : texture_storage_2d<r32uint, read>;
[[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
-[[stage(compute)]] fn main() {
+[[stage(compute), workgroup_size(1)]] fn main() {
var srcValue : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0));
srcValue.x = srcValue.x + 1u;
textureStore(Dst, vec2<i32>(0, 0), srcValue);
@@ -1056,7 +1056,7 @@
wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
[[group(0), binding(0)]] var Src : texture_2d<u32>;
[[group(0), binding(1)]] var Dst : texture_storage_2d<r32uint, write>;
-[[stage(compute)]] fn main() {
+[[stage(compute), workgroup_size(1)]] fn main() {
var srcValue : vec4<u32> = textureLoad(Src, vec2<i32>(0, 0), 0);
srcValue.x = srcValue.x + 1u;
textureStore(Dst, vec2<i32>(0, 0), srcValue);
@@ -1161,7 +1161,7 @@
const char* kCommonWriteOnlyZeroInitTestCodeCompute = R"(
[[group(0), binding(0)]] var dstImage : texture_storage_2d<r32uint, write>;
-[[stage(compute)]] fn main() {
+[[stage(compute), workgroup_size(1)]] fn main() {
textureStore(dstImage, vec2<i32>(0, 0), vec4<u32>(1u, 0u, 0u, 1u));
})";
};
@@ -1204,7 +1204,7 @@
[[group(0), binding(0)]] var srcImage : texture_storage_2d<r32uint, read>;
[[group(0), binding(1)]] var<storage, read_write> 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 00d61f7..8c657ea 100644
--- a/src/tests/end2end/TextureZeroInitTests.cpp
+++ b/src/tests/end2end/TextureZeroInitTests.cpp
@@ -979,7 +979,7 @@
value : vec4<f32>;
};
[[group(0), binding(1)]] var<storage, read_write> result : Result;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
result.value = textureLoad(tex, vec2<i32>(0,0), 0);
}
)";
diff --git a/src/tests/unittests/validation/BindGroupValidationTests.cpp b/src/tests/unittests/validation/BindGroupValidationTests.cpp
index 759b265..1057567 100644
--- a/src/tests/unittests/validation/BindGroupValidationTests.cpp
+++ b/src/tests/unittests/validation/BindGroupValidationTests.cpp
@@ -1981,7 +1981,7 @@
})";
constexpr char kTexture2DShaderCS[] = R"(
[[group(0), binding(0)]] var myTexture : texture_2d<f32>;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
ignore(textureDimensions(myTexture));
})";
@@ -2020,7 +2020,7 @@
})";
constexpr char kTexture2DArrayShaderCS[] = R"(
[[group(0), binding(0)]] var myTexture : texture_2d_array<f32>;
- [[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 987599e..13d4ddd 100644
--- a/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
+++ b/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
@@ -152,7 +152,7 @@
};
[[group(0), binding(0)]] var<uniform> uniforms : S;
- [[stage(compute)]] fn main() {
+ [[stage(compute), workgroup_size(1)]] fn main() {
var pos : vec4<f32> = uniforms.pos;
})");
@@ -915,11 +915,11 @@
[[group(0), binding(0)]] var<storage, read_write> data0 : Data;
[[group(0), binding(1)]] var<storage, read_write> 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 8e9637d..c94e94a 100644
--- a/src/tests/unittests/validation/MultipleDeviceTests.cpp
+++ b/src/tests/unittests/validation/MultipleDeviceTests.cpp
@@ -33,7 +33,7 @@
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 0b6d591..3f5ad4b 100644
--- a/src/tests/unittests/validation/QueueSubmitValidationTests.cpp
+++ b/src/tests/unittests/validation/QueueSubmitValidationTests.cpp
@@ -211,7 +211,7 @@
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 @@
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 @@
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 eed4393..ffc1fe4 100644
--- a/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp
+++ b/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp
@@ -63,7 +63,7 @@
wgpu::ComputePipeline CreateNoOpComputePipeline(std::vector<wgpu::BindGroupLayout> 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 b532492..919ad60 100644
--- a/src/tests/unittests/validation/StorageTextureValidationTests.cpp
+++ b/src/tests/unittests/validation/StorageTextureValidationTests.cpp
@@ -85,7 +85,7 @@
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 @@
};
[[group(0), binding(1)]] var<storage, read_write> buf : Buf;
- [[stage(compute)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) {
+ [[stage(compute), workgroup_size(1)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) {
buf.data = textureLoad(image0, vec2<i32>(LocalInvocationID.xy)).x;
})");
@@ -211,7 +211,7 @@
wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
[[group(0), binding(0)]] var image0 : texture_storage_2d<rgba8unorm, write>;
- [[stage(compute)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) {
+ [[stage(compute), workgroup_size(1)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) {
textureStore(image0, vec2<i32>(LocalInvocationID.xy), vec4<f32>(0.0, 0.0, 0.0, 0.0));
})");
@@ -248,7 +248,7 @@
{
ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"(
[[group(0), binding(0)]] var image0 : texture_storage_2d<rgba8unorm, read_write>;
- [[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 12bd37d..61b3271 100644
--- a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
+++ b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
@@ -136,7 +136,7 @@
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.
