WGSL: Migrate access control to var<> Spec change: https://github.com/gpuweb/gpuweb/pull/1735 Bug: tint:846 Change-Id: Id2eddc4e8f3bdb86027db8d61bb96b9b8ef9778f Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/53386 Reviewed-by: Corentin Wallez <cwallez@chromium.org> Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tests/DawnTest.cpp b/src/tests/DawnTest.cpp index 0d972fc..3de7da0 100644 --- a/src/tests/DawnTest.cpp +++ b/src/tests/DawnTest.cpp
@@ -1098,7 +1098,7 @@ }; [[group(0), binding(0)]] var tex : texture_depth_2d; - [[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result; + [[group(0), binding(1)]] var<storage, read_write> result : Result; [[stage(compute)]] fn main( [[builtin(global_invocation_id)]] GlobalInvocationId : vec3<u32>
diff --git a/src/tests/end2end/BindGroupTests.cpp b/src/tests/end2end/BindGroupTests.cpp index fbe6e52..4ab61c4 100644 --- a/src/tests/end2end/BindGroupTests.cpp +++ b/src/tests/end2end/BindGroupTests.cpp
@@ -72,8 +72,8 @@ << " : Buffer" << i << ";"; break; case wgpu::BufferBindingType::Storage: - fs << "\n[[group(" << i << "), binding(0)]] var<storage> buffer" << i - << " : [[access(read)]] Buffer" << i << ";"; + fs << "\n[[group(" << i << "), binding(0)]] var<storage, read> buffer" << i + << " : Buffer" << i << ";"; break; default: UNREACHABLE(); @@ -837,9 +837,9 @@ }; [[group(0), binding(2)]] var<uniform> buffer2 : Buffer2; - [[group(0), binding(3)]] var<storage> buffer3 : [[access(read)]] Buffer3; - [[group(0), binding(0)]] var<storage> buffer0 : [[access(read)]] Buffer0; - [[group(0), binding(4)]] var<storage> outputBuffer : [[access(read_write)]] OutputBuffer; + [[group(0), binding(3)]] var<storage, read> buffer3 : Buffer3; + [[group(0), binding(0)]] var<storage, read> buffer0 : Buffer0; + [[group(0), binding(4)]] var<storage, read_write> outputBuffer : OutputBuffer; [[stage(compute)]] fn main() { outputBuffer.value = vec3<u32>(buffer0.value, buffer2.value, buffer3.value); @@ -1103,7 +1103,7 @@ [[block]] struct Buffer0 { color : vec4<f32>; }; - [[group(0), binding(0)]] var<storage> buffer0 : [[access(read)]] Buffer0; + [[group(0), binding(0)]] var<storage, read> buffer0 : Buffer0; [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { return buffer0.color; @@ -1240,8 +1240,7 @@ }; )"; interface << "[[group(0), binding(" << binding++ << ")]] " - << "var<storage> sbuf" << i << " : [[access(read)]] ReadOnlyStorageBuffer" << i - << ";\n"; + << "var<storage, read> sbuf" << i << " : ReadOnlyStorageBuffer" << i << ";\n"; body << "if (sbuf" << i << ".value != " << expectedValue++ << "u) {\n"; body << " return;\n"; @@ -1257,7 +1256,7 @@ }; )"; interface << "[[group(0), binding(" << binding++ << ")]] " - << "var<storage> result : [[access(read_write)]] ReadWriteStorageBuffer;\n"; + << "var<storage, read_write> result : ReadWriteStorageBuffer;\n"; body << "result.value = 1u;\n";
diff --git a/src/tests/end2end/BufferZeroInitTests.cpp b/src/tests/end2end/BufferZeroInitTests.cpp index 55bdf90..78e1478 100644 --- a/src/tests/end2end/BufferZeroInitTests.cpp +++ b/src/tests/end2end/BufferZeroInitTests.cpp
@@ -1035,7 +1035,7 @@ [[block]] struct SSBO { value : vec4<u32>; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO; + [[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() { @@ -1074,7 +1074,7 @@ [[block]] struct SSBO { value : array<vec4<u32>, 2>; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO; + [[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() {
diff --git a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp index f41da7b..721ff21 100644 --- a/src/tests/end2end/ComputeCopyStorageBufferTests.cpp +++ b/src/tests/end2end/ComputeCopyStorageBufferTests.cpp
@@ -96,8 +96,8 @@ }; // TODO(crbug.com/tint/386): Use the same struct type - [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1; - [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2; + [[set(0), binding(0)]] var<storage, read_write> src : Buf1; + [[set(0), binding(1)]] var<storage, read_write> dst : Buf2; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) { @@ -123,8 +123,8 @@ }; // TODO(crbug.com/tint/386): Use the same struct type - [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1; - [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2; + [[set(0), binding(0)]] var<storage, read_write> src : Buf1; + [[set(0), binding(1)]] var<storage, read_write> dst : Buf2; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) { @@ -145,8 +145,8 @@ }; // TODO(crbug.com/tint/386): Use the same struct type - [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1; - [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2; + [[set(0), binding(0)]] var<storage, read_write> src : Buf1; + [[set(0), binding(1)]] var<storage, read_write> dst : Buf2; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
diff --git a/src/tests/end2end/ComputeDispatchTests.cpp b/src/tests/end2end/ComputeDispatchTests.cpp index 93d844f..860aaee 100644 --- a/src/tests/end2end/ComputeDispatchTests.cpp +++ b/src/tests/end2end/ComputeDispatchTests.cpp
@@ -37,7 +37,7 @@ }; [[group(0), binding(0)]] var<uniform> input : InputBuf; - [[group(0), binding(1)]] var<storage> output : [[access(read_write)]] OutputBuf; + [[group(0), binding(1)]] var<storage, read_write> output : OutputBuf; [[stage(compute), workgroup_size(1, 1, 1)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
diff --git a/src/tests/end2end/ComputeSharedMemoryTests.cpp b/src/tests/end2end/ComputeSharedMemoryTests.cpp index 8ac3212..3e7e6de 100644 --- a/src/tests/end2end/ComputeSharedMemoryTests.cpp +++ b/src/tests/end2end/ComputeSharedMemoryTests.cpp
@@ -78,7 +78,7 @@ x : u32; }; - [[group(0), binding(0)]] var<storage> dst : [[access(write)]] Dst; + [[group(0), binding(0)]] var<storage, write> dst : Dst; var<workgroup> tmp : u32; [[stage(compute), workgroup_size(4,4,1)]]
diff --git a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp index 3cfa6dc..7775974 100644 --- a/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp +++ b/src/tests/end2end/ComputeStorageBufferBarrierTests.cpp
@@ -36,7 +36,7 @@ data : array<u32, 100>; }; - [[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf; + [[group(0), binding(0)]] var<storage, read_write> buf : Buf; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) { @@ -91,8 +91,8 @@ data : array<u32, 100>; }; - [[group(0), binding(0)]] var<storage> src : [[access(read_write)]] Src; - [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Dst; + [[group(0), binding(0)]] var<storage, read_write> src : Src; + [[group(0), binding(1)]] var<storage, read_write> dst : Dst; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) { @@ -162,8 +162,8 @@ data : array<u32, 100>; }; - [[group(0), binding(0)]] var<storage> src : [[access(read)]] Src; - [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Dst; + [[group(0), binding(0)]] var<storage, read> src : Src; + [[group(0), binding(1)]] var<storage, read_write> dst : Dst; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) { @@ -231,7 +231,7 @@ }; [[group(0), binding(0)]] var<uniform> src : Buf; - [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf; + [[group(0), binding(1)]] var<storage, read_write> dst : Buf; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) { @@ -299,7 +299,7 @@ }; [[group(0), binding(0)]] var<uniform> src : Buf; - [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf; + [[group(0), binding(1)]] var<storage, read_write> dst : Buf; [[stage(compute)]] fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) { @@ -359,7 +359,7 @@ [[block]] struct Buf { data : array<u32, 3>; }; - [[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf; + [[group(0), binding(0)]] var<storage, read_write> buf : Buf; [[stage(compute)]] fn main() { buf.data = array<u32, 3>(1u, 1u, 1u); @@ -373,12 +373,12 @@ [[block]] struct Buf { data : array<u32, 3>; }; - [[group(0), binding(0)]] var<storage> buf : [[access(read)]] Buf; + [[group(0), binding(0)]] var<storage, read> buf : Buf; [[block]] struct Result { data : u32; }; - [[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result; + [[group(0), binding(1)]] var<storage, read_write> result : Result; [[stage(compute)]] fn main() { result.data = 2u;
diff --git a/src/tests/end2end/CopyTextureForBrowserTests.cpp b/src/tests/end2end/CopyTextureForBrowserTests.cpp index ffcf2fc..1cce29a 100644 --- a/src/tests/end2end/CopyTextureForBrowserTests.cpp +++ b/src/tests/end2end/CopyTextureForBrowserTests.cpp
@@ -146,7 +146,7 @@ }; [[group(0), binding(0)]] var src : texture_2d<f32>; [[group(0), binding(1)]] var dst : texture_2d<f32>; - [[group(0), binding(2)]] var<storage> output : [[access(read_write)]] OutputBuf; + [[group(0), binding(2)]] var<storage, read_write> output : OutputBuf; [[group(0), binding(3)]] var<uniform> uniforms : Uniforms; fn aboutEqual(value : f32, expect : f32) -> bool { // The value diff should be smaller than the hard coded tolerance. @@ -389,7 +389,6 @@ // Verify CopyTextureForBrowserTests works with internal pipeline. // The case do copy without any transform. TEST_P(CopyTextureForBrowserTests, PassthroughCopy) { - constexpr uint32_t kWidth = 10; constexpr uint32_t kHeight = 1; @@ -420,7 +419,6 @@ } TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) { - constexpr uint32_t kWidth = 899; constexpr uint32_t kHeight = 999;
diff --git a/src/tests/end2end/CreatePipelineAsyncTests.cpp b/src/tests/end2end/CreatePipelineAsyncTests.cpp index 44990e0..c5e76f8 100644 --- a/src/tests/end2end/CreatePipelineAsyncTests.cpp +++ b/src/tests/end2end/CreatePipelineAsyncTests.cpp
@@ -78,7 +78,7 @@ [[block]] struct SSBO { value : u32; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO; + [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO; [[stage(compute)]] fn main() { ssbo.value = 1u; @@ -113,7 +113,7 @@ [[block]] struct SSBO { value : u32; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO; + [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO; [[stage(compute)]] fn main() { ssbo.value = 1u; @@ -312,7 +312,7 @@ [[block]] struct SSBO { value : u32; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO; + [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO; [[stage(compute)]] fn main() { ssbo.value = 1u; @@ -353,7 +353,7 @@ [[block]] struct SSBO { value : u32; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO; + [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO; [[stage(compute)]] fn main() { ssbo.value = 1u;
diff --git a/src/tests/end2end/D3D12CachingTests.cpp b/src/tests/end2end/D3D12CachingTests.cpp index fde7d6c..ecf7c8b 100644 --- a/src/tests/end2end/D3D12CachingTests.cpp +++ b/src/tests/end2end/D3D12CachingTests.cpp
@@ -211,7 +211,7 @@ [[block]] struct Data { data : u32; }; - [[binding(0), group(0)]] var<storage> data : [[access(read_write)]] Data; + [[binding(0), group(0)]] var<storage, read_write> data : Data; [[stage(compute)]] fn write1() { data.data = 1u;
diff --git a/src/tests/end2end/DepthStencilSamplingTests.cpp b/src/tests/end2end/DepthStencilSamplingTests.cpp index 1b7e584..d0681a0 100644 --- a/src/tests/end2end/DepthStencilSamplingTests.cpp +++ b/src/tests/end2end/DepthStencilSamplingTests.cpp
@@ -141,8 +141,8 @@ << " : texture_2d<f32>;\n"; shaderSource << "[[group(0), binding(" << 2 * index + 1 - << ")]] var<storage> result" << index - << " : [[access(read_write)]] DepthResult;\n"; + << ")]] var<storage, read_write> result" << index + << " : DepthResult;\n"; shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index << ", vec2<i32>(0, 0), 0)[" << componentIndex << "];"; @@ -152,8 +152,8 @@ << " : texture_2d<u32>;\n"; shaderSource << "[[group(0), binding(" << 2 * index + 1 - << ")]] var<storage> result" << index - << " : [[access(read_write)]] StencilResult;\n"; + << ")]] var<storage, read_write> result" << index + << " : StencilResult;\n"; shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index << ", vec2<i32>(0, 0), 0)[" << componentIndex << "];"; @@ -221,7 +221,7 @@ [[block]] struct SamplerResult { value : f32; }; - [[group(0), binding(3)]] var<storage> samplerResult : [[access(read_write)]] SamplerResult; + [[group(0), binding(3)]] var<storage, read_write> samplerResult : SamplerResult; [[stage(compute)]] fn main() { samplerResult.value = textureSampleCompare(tex, samp, vec2<f32>(0.5, 0.5), uniforms.compareRef);
diff --git a/src/tests/end2end/DynamicBufferOffsetTests.cpp b/src/tests/end2end/DynamicBufferOffsetTests.cpp index 33de9ef..32821fd 100644 --- a/src/tests/end2end/DynamicBufferOffsetTests.cpp +++ b/src/tests/end2end/DynamicBufferOffsetTests.cpp
@@ -125,9 +125,9 @@ }; [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1; - [[group(0), binding(1)]] var<storage> sBufferNotDynamic : [[access(read_write)]] Buffer2; + [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buffer2; [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3; - [[group(0), binding(4)]] var<storage> sBuffer : [[access(read_write)]] Buffer4; + [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buffer4; )"; if (isInheritedPipeline) { @@ -192,9 +192,9 @@ }; [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1; - [[group(0), binding(1)]] var<storage> sBufferNotDynamic : [[access(read_write)]] Buffer2; + [[group(0), binding(1)]] var<storage, read_write> sBufferNotDynamic : Buffer2; [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3; - [[group(0), binding(4)]] var<storage> sBuffer : [[access(read_write)]] Buffer4; + [[group(0), binding(4)]] var<storage, read_write> sBuffer : Buffer4; )"; if (isInheritedPipeline) {
diff --git a/src/tests/end2end/EntryPointTests.cpp b/src/tests/end2end/EntryPointTests.cpp index 037d6ed..da5d5db 100644 --- a/src/tests/end2end/EntryPointTests.cpp +++ b/src/tests/end2end/EntryPointTests.cpp
@@ -64,7 +64,7 @@ [[block]] struct Data { data : u32; }; - [[binding(0), group(0)]] var<storage> data : [[access(read_write)]] Data; + [[binding(0), group(0)]] var<storage, read_write> data : Data; [[stage(compute)]] fn write1() { data.data = 1u;
diff --git a/src/tests/end2end/FirstIndexOffsetTests.cpp b/src/tests/end2end/FirstIndexOffsetTests.cpp index b6e56d7..506c0d3 100644 --- a/src/tests/end2end/FirstIndexOffsetTests.cpp +++ b/src/tests/end2end/FirstIndexOffsetTests.cpp
@@ -132,7 +132,7 @@ vertex_index : u32; instance_index : u32; }; -[[group(0), binding(0)]] var<storage> idx_vals : [[access(read_write)]] IndexVals; +[[group(0), binding(0)]] var<storage, read_write> idx_vals : IndexVals; struct FragInputs { )" + fragmentInputs.str() + R"(
diff --git a/src/tests/end2end/GpuMemorySynchronizationTests.cpp b/src/tests/end2end/GpuMemorySynchronizationTests.cpp index 1db5400..7113a13 100644 --- a/src/tests/end2end/GpuMemorySynchronizationTests.cpp +++ b/src/tests/end2end/GpuMemorySynchronizationTests.cpp
@@ -39,7 +39,7 @@ [[block]] struct Data { a : i32; }; - [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data; + [[group(0), binding(0)]] var<storage, read_write> data : Data; [[stage(compute)]] fn main() { data.a = data.a + 1; })"); @@ -66,7 +66,7 @@ [[block]] struct Data { i : i32; }; - [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data; + [[group(0), binding(0)]] var<storage, read_write> data : Data; [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { data.i = data.i + 1; return vec4<f32>(f32(data.i) / 255.0, 0.0, 0.0, 1.0); @@ -255,7 +255,7 @@ sampledOut: u32; storageOut: u32; }; - [[group(0), binding(0)]] var<storage> output : [[access(write)]] Output; + [[group(0), binding(0)]] var<storage, write> output : Output; [[group(0), binding(1)]] var sampledTex : texture_2d<u32>; [[group(0), binding(2)]] var storageTex : texture_storage_2d<r32uint, read>; @@ -315,7 +315,7 @@ [[block]] struct Data { a : f32; }; - [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data; + [[group(0), binding(0)]] var<storage, read_write> data : Data; [[stage(compute)]] fn main() { data.a = 1.0; })"); @@ -512,12 +512,12 @@ [[block]] struct VBContents { pos : array<vec4<f32>, 4>; }; - [[group(0), binding(0)]] var<storage> vbContents : [[access(read_write)]] VBContents; + [[group(0), binding(0)]] var<storage, read_write> vbContents : VBContents; [[block]] struct IBContents { indices : array<vec4<i32>, 2>; }; - [[group(0), binding(1)]] var<storage> ibContents : [[access(read_write)]] IBContents; + [[group(0), binding(1)]] var<storage, read_write> ibContents : IBContents; // TODO(crbug.com/tint/386): Use the same struct. [[block]] struct ColorContents1 { @@ -526,8 +526,8 @@ [[block]] struct ColorContents2 { color : f32; }; - [[group(0), binding(2)]] var<storage> uniformContents : [[access(read_write)]] ColorContents1; - [[group(0), binding(3)]] var<storage> storageContents : [[access(read_write)]] ColorContents2; + [[group(0), binding(2)]] var<storage, read_write> uniformContents : ColorContents1; + [[group(0), binding(3)]] var<storage, read_write> storageContents : ColorContents2; [[stage(compute)]] fn main() { vbContents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0); @@ -581,7 +581,7 @@ }; [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf; - [[group(0), binding(1)]] var<storage> storageBuffer : [[access(read)]] Buf; + [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf; [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0); @@ -642,7 +642,7 @@ [[align(256)]] color1 : f32; }; - [[group(0), binding(0)]] var<storage> contents : [[access(read_write)]] Contents; + [[group(0), binding(0)]] var<storage, read_write> contents : Contents; [[stage(compute)]] fn main() { contents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0); @@ -696,7 +696,7 @@ color : f32; }; [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf; - [[group(0), binding(1)]] var<storage> storageBuffer : [[access(read)]] Buf; + [[group(0), binding(1)]] var<storage, read> storageBuffer : Buf; [[stage(fragment)]] fn main() -> [[location(0)]] vec4<f32> { return vec4<f32>(uniformBuffer.color, storageBuffer.color, 0.0, 1.0);
diff --git a/src/tests/end2end/MultisampledSamplingTests.cpp b/src/tests/end2end/MultisampledSamplingTests.cpp index 055f2c1..39e6103 100644 --- a/src/tests/end2end/MultisampledSamplingTests.cpp +++ b/src/tests/end2end/MultisampledSamplingTests.cpp
@@ -100,7 +100,7 @@ colorSamples : array<f32, 4>; depthSamples : array<f32, 4>; }; - [[group(0), binding(2)]] var<storage> results : [[access(read_write)]] Results; + [[group(0), binding(2)]] var<storage, read_write> results : Results; [[stage(compute)]] fn main() { for (var i : i32 = 0; i < 4; i = i + 1) {
diff --git a/src/tests/end2end/OpArrayLengthTests.cpp b/src/tests/end2end/OpArrayLengthTests.cpp index c9d42f7..b7b55bf 100644 --- a/src/tests/end2end/OpArrayLengthTests.cpp +++ b/src/tests/end2end/OpArrayLengthTests.cpp
@@ -63,10 +63,10 @@ }; // The length should be 1 because the buffer is 4-byte long. - [[group(0), binding(0)]] var<storage> buffer1 : [[access(read)]] DataBuffer1; + [[group(0), binding(0)]] var<storage, read> buffer1 : DataBuffer1; // The length should be 64 because the buffer is 256 bytes long. - [[group(0), binding(1)]] var<storage> buffer2 : [[access(read)]] DataBuffer2; + [[group(0), binding(1)]] var<storage, read> buffer2 : DataBuffer2; // The length should be (512 - 16*4) / 8 = 56 because the buffer is 512 bytes long // and the structure is 8 bytes big. @@ -79,7 +79,7 @@ [[size(64)]] garbage : mat4x4<f32>; data : [[stride(8)]] array<Buffer3Data>; }; - [[group(0), binding(2)]] var<storage> buffer3 : [[access(read)]] Buffer3; + [[group(0), binding(2)]] var<storage, read> buffer3 : Buffer3; )"; // See comments in the shader for an explanation of these values @@ -128,7 +128,7 @@ [[block]] struct ResultBuffer { data : [[stride(4)]] array<u32, 3>; }; - [[group(1), binding(0)]] var<storage> result : [[access(read_write)]] ResultBuffer; + [[group(1), binding(0)]] var<storage, read_write> result : ResultBuffer; )" + mShaderInterface + R"( [[stage(compute)]] fn main() { result.data[0] = arrayLength(buffer1.data);
diff --git a/src/tests/end2end/ShaderTests.cpp b/src/tests/end2end/ShaderTests.cpp index 7c9475c..ace5c8f 100644 --- a/src/tests/end2end/ShaderTests.cpp +++ b/src/tests/end2end/ShaderTests.cpp
@@ -35,7 +35,7 @@ data : array<u32, 19>; }; -[[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf; +[[group(0), binding(0)]] var<storage, read_write> buf : Buf; [[stage(compute)]] fn main() { let factor : f32 = 1.0001;
diff --git a/src/tests/end2end/StorageTextureTests.cpp b/src/tests/end2end/StorageTextureTests.cpp index 0ab4a0a..a25aa53 100644 --- a/src/tests/end2end/StorageTextureTests.cpp +++ b/src/tests/end2end/StorageTextureTests.cpp
@@ -167,12 +167,12 @@ std::ostringstream ostream; ostream << "[[group(0), binding(" << binding << ")]] " << "var storageImage" << binding << " : " - << "[[access(" << accessQualifier << ")]] " << "texture_storage_2d"; if (is2DArray) { ostream << "_array"; } - ostream << "<" << utils::GetWGSLImageFormatQualifier(format) << ">;"; + ostream << "<" << utils::GetWGSLImageFormatQualifier(format) << ", "; + ostream << accessQualifier << ">;"; return ostream.str(); } @@ -710,7 +710,7 @@ result : u32; }; -[[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer; +[[group(0), binding(1)]] var<storage, read_write> dstBuffer : DstBuffer; )" << CommonReadOnlyTestCode(format) << R"( [[stage(compute)]] fn main() { @@ -934,7 +934,7 @@ result : u32; }; -[[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer; +[[group(0), binding(1)]] var<storage, read_write> dstBuffer : DstBuffer; )" << CommonReadOnlyTestCode(kTextureFormat, true) << R"( [[stage(compute)]] fn main() { @@ -1202,7 +1202,7 @@ }; [[group(0), binding(0)]] var srcImage : texture_storage_2d<r32uint, read>; -[[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer; +[[group(0), binding(1)]] var<storage, read_write> dstBuffer : DstBuffer; )") + kCommonReadOnlyZeroInitTestCode + R"( [[stage(compute)]] fn main() { if (doTest()) {
diff --git a/src/tests/end2end/TextureZeroInitTests.cpp b/src/tests/end2end/TextureZeroInitTests.cpp index 3f2e8ee..f8a03d5 100644 --- a/src/tests/end2end/TextureZeroInitTests.cpp +++ b/src/tests/end2end/TextureZeroInitTests.cpp
@@ -979,7 +979,7 @@ [[block]] struct Result { value : vec4<f32>; }; - [[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result; + [[group(0), binding(1)]] var<storage, read_write> result : Result; [[stage(compute)]] fn main() { result.value = textureLoad(tex, vec2<i32>(0,0), 0); }
diff --git a/src/tests/perf_tests/ShaderRobustnessPerf.cpp b/src/tests/perf_tests/ShaderRobustnessPerf.cpp index 0add757..f63af7b 100644 --- a/src/tests/perf_tests/ShaderRobustnessPerf.cpp +++ b/src/tests/perf_tests/ShaderRobustnessPerf.cpp
@@ -29,9 +29,9 @@ numbers: array<f32>; }; - [[group(0), binding(0)]] var<storage> firstMatrix : [[access(read)]] Matrix; - [[group(0), binding(1)]] var<storage> secondMatrix : [[access(read)]] Matrix; - [[group(0), binding(2)]] var<storage> resultMatrix : [[access(write)]] Matrix; + [[group(0), binding(0)]] var<storage, read> firstMatrix : Matrix; + [[group(0), binding(1)]] var<storage, read> secondMatrix : Matrix; + [[group(0), binding(2)]] var<storage, write> resultMatrix : Matrix; [[group(0), binding(3)]] var<uniform> uniforms : Uniforms; fn mm_readA(row : u32, col : u32) -> f32 { @@ -196,9 +196,9 @@ numbers: array<vec4<f32>>; }; - [[group(0), binding(0)]] var<storage> firstMatrix : [[access(read)]] Matrix; - [[group(0), binding(1)]] var<storage> secondMatrix : [[access(read)]] Matrix; - [[group(0), binding(2)]] var<storage> resultMatrix : [[access(write)]] Matrix; + [[group(0), binding(0)]] var<storage, read> firstMatrix : Matrix; + [[group(0), binding(1)]] var<storage, read> secondMatrix : Matrix; + [[group(0), binding(2)]] var<storage, write> resultMatrix : Matrix; [[group(0), binding(3)]] var<uniform> uniforms : Uniforms; fn mm_readA(row : u32, col : u32) -> vec4<f32> {
diff --git a/src/tests/unittests/validation/BindGroupValidationTests.cpp b/src/tests/unittests/validation/BindGroupValidationTests.cpp index 33c656f..b41b8c4 100644 --- a/src/tests/unittests/validation/BindGroupValidationTests.cpp +++ b/src/tests/unittests/validation/BindGroupValidationTests.cpp
@@ -1278,8 +1278,8 @@ [[group(0), binding(0)]] var<uniform> uBufferDynamic : S; [[group(0), binding(1)]] var<uniform> uBuffer : S; - [[group(0), binding(2)]] var<storage> sBufferDynamic : [[access(read_write)]] S; - [[group(0), binding(3)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S; + [[group(0), binding(2)]] var<storage, read_write> sBufferDynamic : S; + [[group(0), binding(3)]] var<storage, read> sReadonlyBufferDynamic : S; [[stage(fragment)]] fn main() { })"); @@ -1301,8 +1301,8 @@ [[group(0), binding(0)]] var<uniform> uBufferDynamic : S; [[group(0), binding(1)]] var<uniform> uBuffer : S; - [[group(0), binding(2)]] var<storage> sBufferDynamic : [[access(read_write)]] S; - [[group(0), binding(3)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S; + [[group(0), binding(2)]] var<storage, read_write> sBufferDynamic : S; + [[group(0), binding(3)]] var<storage, read> sReadonlyBufferDynamic : S; [[stage(compute), workgroup_size(4, 4, 1)]] fn main() { })"); @@ -1728,8 +1728,7 @@ ss << "[[group(" << l << "), binding(" << b << ")]] "; switch (binding) { case wgpu::BufferBindingType::Storage: - ss << "var<storage> set" << l << "_binding" << b - << " : [[access(read_write)]] S;"; + ss << "var<storage, read_write> set" << l << "_binding" << b << " : S;"; break; case wgpu::BufferBindingType::Uniform: ss << "var<uniform> set" << l << "_binding" << b << " : S;"; @@ -1897,8 +1896,8 @@ value : vec2<f32>; }; - [[group(0), binding(0)]] var<storage> sBufferDynamic : [[access(read_write)]] S; - [[group(1), binding(0)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S; + [[group(0), binding(0)]] var<storage, read_write> sBufferDynamic : S; + [[group(1), binding(0)]] var<storage, read> sReadonlyBufferDynamic : S; [[stage(fragment)]] fn main() { var val : vec2<f32> = sBufferDynamic.value; @@ -1932,8 +1931,8 @@ value : vec2<f32>; }; - [[group(0), binding(0)]] var<storage> sBufferDynamic : [[access(read_write)]] S; - [[group(1), binding(0)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S; + [[group(0), binding(0)]] var<storage, read_write> sBufferDynamic : S; + [[group(1), binding(0)]] var<storage, read> sReadonlyBufferDynamic : S; [[stage(compute), workgroup_size(4, 4, 1)]] fn main() { var val : vec2<f32> = sBufferDynamic.value;
diff --git a/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp b/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp index de024b0..a0a531a 100644 --- a/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp +++ b/src/tests/unittests/validation/GetBindGroupLayoutValidationTests.cpp
@@ -66,7 +66,7 @@ [[block]] struct S3 { pos : mat4x4<f32>; }; - [[group(3), binding(0)]] var<storage> storage3 : [[access(read_write)]] S3; + [[group(3), binding(0)]] var<storage, read_write> storage3 : S3; [[stage(fragment)]] fn main() { var pos_u : vec4<f32> = uniform2.pos; @@ -202,7 +202,7 @@ [[block]] struct S { pos : vec4<f32>; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] S; + [[group(0), binding(0)]] var<storage, read_write> ssbo : S; [[stage(fragment)]] fn main() { var pos : vec4<f32> = ssbo.pos; @@ -229,7 +229,7 @@ [[block]] struct S { pos : vec4<f32>; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read)]] S; + [[group(0), binding(0)]] var<storage, read> ssbo : S; [[stage(fragment)]] fn main() { var pos : vec4<f32> = ssbo.pos; @@ -707,7 +707,7 @@ [[block]] struct S { pos : vec4<f32>; }; - [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] S; + [[group(0), binding(0)]] var<storage, read_write> ssbo : S; [[stage(fragment)]] fn main() { var pos : vec4<f32> = ssbo.pos; @@ -912,8 +912,8 @@ [[block]] struct Data { data : f32; }; - [[group(0), binding(0)]] var<storage> data0 : [[access(read_write)]] Data; - [[group(0), binding(1)]] var<storage> data1 : [[access(read_write)]] Data; + [[group(0), binding(0)]] var<storage, read_write> data0 : Data; + [[group(0), binding(1)]] var<storage, read_write> data1 : Data; [[stage(compute)]] fn compute0() { data0.data = 0.0;
diff --git a/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp b/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp index de4ec17..df88025 100644 --- a/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp +++ b/src/tests/unittests/validation/MinimumBufferSizeValidationTests.cpp
@@ -80,12 +80,10 @@ ostream << "var<uniform> b" << index << " : S" << index << ";\n"; break; case wgpu::BufferBindingType::Storage: - ostream << "var<storage> b" << index << " : [[access(read_write)]] S" << index - << ";\n"; + ostream << "var<storage, read_write> b" << index << " : S" << index << ";\n"; break; case wgpu::BufferBindingType::ReadOnlyStorage: - ostream << "var<storage> b" << index << " : [[access(read)]] S" << index - << ";\n"; + ostream << "var<storage, read> b" << index << " : S" << index << ";\n"; break; default: UNREACHABLE();
diff --git a/src/tests/unittests/validation/RenderBundleValidationTests.cpp b/src/tests/unittests/validation/RenderBundleValidationTests.cpp index 11277b2..b6da94e 100644 --- a/src/tests/unittests/validation/RenderBundleValidationTests.cpp +++ b/src/tests/unittests/validation/RenderBundleValidationTests.cpp
@@ -46,7 +46,7 @@ [[block]] struct Storage { dummy : array<f32>; }; - [[group(1), binding(1)]] var<storage> ssbo : [[access(read_write)]] Storage; + [[group(1), binding(1)]] var<storage, read_write> ssbo : Storage; [[stage(fragment)]] fn main() { })");
diff --git a/src/tests/unittests/validation/RenderPipelineValidationTests.cpp b/src/tests/unittests/validation/RenderPipelineValidationTests.cpp index 37f156f..593a365 100644 --- a/src/tests/unittests/validation/RenderPipelineValidationTests.cpp +++ b/src/tests/unittests/validation/RenderPipelineValidationTests.cpp
@@ -460,7 +460,7 @@ [[block]] struct Dst { data : array<u32, 100>; }; - [[group(0), binding(0)]] var<storage> dst : [[access(read_write)]] Dst; + [[group(0), binding(0)]] var<storage, read_write> dst : Dst; [[stage(vertex)]] fn main([[builtin(vertex_index)]] VertexIndex : u32) -> [[builtin(position)]] vec4<f32> { dst.data[VertexIndex] = 0x1234u; return vec4<f32>();
diff --git a/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp b/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp index 142ba56..d70ef9a 100644 --- a/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp +++ b/src/tests/unittests/validation/ResourceUsageTrackingTests.cpp
@@ -762,7 +762,7 @@ [[block]] struct RBuffer { value : f32; }; - [[group(0), binding(0)]] var<storage> rBuffer : [[access(read)]] RBuffer; + [[group(0), binding(0)]] var<storage, read> rBuffer : RBuffer; [[stage(fragment)]] fn main() { })"); utils::ComboRenderPipelineDescriptor pipelineDescriptor;
diff --git a/src/tests/unittests/validation/StorageTextureValidationTests.cpp b/src/tests/unittests/validation/StorageTextureValidationTests.cpp index 5294c0f..a212e10 100644 --- a/src/tests/unittests/validation/StorageTextureValidationTests.cpp +++ b/src/tests/unittests/validation/StorageTextureValidationTests.cpp
@@ -82,9 +82,8 @@ } std::ostringstream ostream; - ostream << "[[group(0), binding(0)]] var image0 : " - << "[[access(" << access << ")]] " << imageTypeDeclaration << "<" - << imageFormatQualifier + ostream << "[[group(0), binding(0)]] var image0 : " << imageTypeDeclaration << "<" + << imageFormatQualifier << ", " << access << ">;\n" "[[stage(compute)]] fn main() {\n" " textureDimensions(image0);\n" @@ -193,7 +192,7 @@ [[block]] struct Buf { data : f32; }; - [[group(0), binding(1)]] var<storage> buf : [[access(read_write)]] Buf; + [[group(0), binding(1)]] var<storage, read_write> buf : Buf; [[stage(compute)]] fn main([[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) { buf.data = textureLoad(image0, vec2<i32>(LocalInvocationID.xy)).x;