encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Sarah 2021-06-23 19:19:06 +00:00 committed by Dawn LUCI CQ
parent 1317636fc8
commit 2a57db73cb
26 changed files with 72 additions and 70 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/tests/DawnTest.cpp
View File

@ -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<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(

9
src/tests/end2end/BindGroupTests.cpp
View File

@ -127,7 +127,7 @@ TEST_P(BindGroupTests, ReusedBindGroupSingleSubmit) {
};
[[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 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) {
[[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 @@ 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";

8
src/tests/end2end/BufferZeroInitTests.cpp
View File

@ -427,7 +427,7 @@ class BufferZeroInitTest : public DawnTest {
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 @@ TEST_P(BufferZeroInitTest, BoundAsUniformBuffer) {
[[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 @@ TEST_P(BufferZeroInitTest, BoundAsReadonlyStorageBuffer) {
[[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 @@ TEST_P(BufferZeroInitTest, BoundAsStorageBuffer) {
[[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));

6
src/tests/end2end/ComputeCopyStorageBufferTests.cpp
View File

@ -99,7 +99,7 @@ TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfBasic) {
[[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 @@ TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfStruct) {
[[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 @@ TEST_P(ComputeCopyStorageBufferTests, UnsizedArrayOfBasic) {
[[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; }

14
src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
View File

@ -38,7 +38,7 @@ TEST_P(ComputeStorageBufferBarrierTests, AddIncrement) {
[[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 @@ TEST_P(ComputeStorageBufferBarrierTests, AddPingPong) {
[[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 @@ TEST_P(ComputeStorageBufferBarrierTests, StorageAndReadonlyStoragePingPongInOneP
[[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 @@ TEST_P(ComputeStorageBufferBarrierTests, UniformToStorageAddPingPong) {
[[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 @@ TEST_P(ComputeStorageBufferBarrierTests, UniformToStorageAddPingPongInOnePass) {
[[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 @@ TEST_P(ComputeStorageBufferBarrierTests, IndirectBufferCorrectBarrier) {
};
[[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 @@ TEST_P(ComputeStorageBufferBarrierTests, IndirectBufferCorrectBarrier) {
};
[[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;

10
src/tests/end2end/CreatePipelineAsyncTests.cpp
View File

@ -80,7 +80,7 @@ TEST_P(CreatePipelineAsyncTest, BasicUseOfCreateComputePipelineAsync) {
};
[[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 @@ TEST_P(CreatePipelineAsyncTest, CreateComputePipelineFailed) {
};
[[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, 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<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 @@ TEST_P(CreatePipelineAsyncTest, CreateSamePipelineTwiceAtSameTime) {
};
[[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";

4
src/tests/end2end/D3D12CachingTests.cpp
View File

@ -213,11 +213,11 @@ TEST_P(D3D12CachingTests, ReuseShaderWithMultipleEntryPoints) {
};
[[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;
}
)");

2
src/tests/end2end/DeprecatedAPITests.cpp
View File

@ -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";

5
src/tests/end2end/DepthStencilSamplingTests.cpp
View File

@ -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<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);
})");

4
src/tests/end2end/DeviceLostTests.cpp
View File

@ -110,7 +110,7 @@ TEST_P(DeviceLostTest, GetBindGroupLayoutFails) {
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 @@ 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;

2
src/tests/end2end/DynamicBufferOffsetTests.cpp
View File

@ -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<u32>(multipleNumber, multipleNumber) * (uBuffer.value.xy + sBufferNotDynamic.value.xy);
}

4
src/tests/end2end/EntryPointTests.cpp
View File

@ -66,12 +66,12 @@ TEST_P(EntryPointTests, TwoComputeInModule) {
};
[[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;
}

10
src/tests/end2end/GpuMemorySynchronizationTests.cpp
View File

@ -40,7 +40,7 @@ class GpuMemorySyncTests : public DawnTest {
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 @@ TEST_P(GpuMemorySyncTests, SampledAndROStorageTextureInComputePass) {
[[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 @@ class StorageToUniformSyncTests : public DawnTest {
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 @@ TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) {
[[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 @@ TEST_P(MultipleWriteThenMultipleReadTests, OneBuffer) {
[[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);

2
src/tests/end2end/MultisampledSamplingTests.cpp
View File

@ -102,7 +102,7 @@ class MultisampledSamplingTest : public DawnTest {
};
[[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;

8
src/tests/end2end/ObjectCachingTests.cpp
View File

@ -124,16 +124,16 @@ TEST_P(ObjectCachingTest, ShaderModuleDeduplication) {
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 @@ TEST_P(ObjectCachingTest, ComputePipelineDeduplicationOnLayout) {
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;
})");

2
src/tests/end2end/OpArrayLengthTests.cpp
View File

@ -130,7 +130,7 @@ TEST_P(OpArrayLengthTest, Compute) {
};
[[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);

4
src/tests/end2end/ShaderTests.cpp
View File

@ -37,7 +37,7 @@ TEST_P(ShaderTests, ComputeLog2) {
[[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 @@ 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;

14
src/tests/end2end/StorageTextureTests.cpp
View File

@ -358,7 +358,7 @@ fn IsEqualTo(pixel : vec4<f32>, expected : vec4<f32>) -> 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<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 @@ TEST_P(StorageTextureTests, ReadonlyStorageTextureInComputeShader) {
[[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 @@ TEST_P(StorageTextureTests, Readonly2DArrayStorageTexture) {
[[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 @@ TEST_P(StorageTextureTests, ReadonlyAndWriteonlyStorageTexturePingPong) {
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 @@ TEST_P(StorageTextureTests, SampledAndWriteonlyStorageTexturePingPong) {
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 @@ fn doTest() -> bool {
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 @@ TEST_P(StorageTextureZeroInitTests, ReadonlyStorageTextureClearsToZeroInComputeP
[[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 {

2
src/tests/end2end/TextureZeroInitTests.cpp
View File

@ -979,7 +979,7 @@ TEST_P(TextureZeroInitTest, ComputePassSampledTextureClear) {
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);
}
)";

4
src/tests/unittests/validation/BindGroupValidationTests.cpp
View File

@ -1981,7 +1981,7 @@ TEST_F(BindGroupLayoutCompatibilityTest, TextureViewDimension) {
})";
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 @@ TEST_F(BindGroupLayoutCompatibilityTest, TextureViewDimension) {
})";
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));
})";

6
src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
View File

@ -152,7 +152,7 @@ TEST_F(GetBindGroupLayoutTests, ComputePipeline) {
};
[[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 @@ TEST_F(GetBindGroupLayoutTests, FromCorrectEntryPoint) {
[[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;
}
)");

2
src/tests/unittests/validation/MultipleDeviceTests.cpp
View File

@ -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() {
}
)";

6
src/tests/unittests/validation/QueueSubmitValidationTests.cpp
View File

@ -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;

2
src/tests/unittests/validation/ResourceUsageTrackingTests.cpp
View File

@ -63,7 +63,7 @@ namespace {
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));

8
src/tests/unittests/validation/StorageTextureValidationTests.cpp
View File

@ -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<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 @@ TEST_F(StorageTextureValidationTests, ComputePipeline) {
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 @@ TEST_F(StorageTextureValidationTests, ReadWriteStorageTexture) {
{
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));
})"));
}

2
src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
View File

@ -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.