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/dawn_native/QueryHelper.cpp b/src/dawn_native/QueryHelper.cpp
index 70d9a9c..eb3d0bf 100644
--- a/src/dawn_native/QueryHelper.cpp
+++ b/src/dawn_native/QueryHelper.cpp
@@ -55,9 +55,9 @@
};
[[group(0), binding(0)]]
- var<storage> timestamps : [[access(read_write)]] TimestampArr;
+ var<storage, read_write> timestamps : TimestampArr;
[[group(0), binding(1)]]
- var<storage> availability : [[access(read)]] AvailabilityArr;
+ var<storage, read> availability : AvailabilityArr;
[[group(0), binding(2)]] var<uniform> params : TimestampParams;
diff --git a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
index ed4727d..f7a5a91 100644
--- a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp
@@ -234,7 +234,7 @@
bgl->GetBindingInfo(bindingIndex).buffer.type ==
wgpu::BufferBindingType::Storage);
if (forceStorageBufferAsUAV) {
- accessControls.emplace(srcBindingPoint, tint::ast::AccessControl::kReadWrite);
+ accessControls.emplace(srcBindingPoint, tint::ast::Access::kReadWrite);
}
}
}
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;
