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

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>
This commit is contained in:
Ben Clayton 2021-06-08 15:36:44 +00:00 committed by Dawn LUCI CQ
parent 3a07f7410f
commit 15eba9a048
31 changed files with 103 additions and 110 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
examples/ComputeBoids.cpp
View File

@ -166,8 +166,8 @@ void initSim() {
particles : array<Particle>;
};
[[binding(0), group(0)]] var<uniform> params : SimParams;
[[binding(1), group(0)]] var<storage> particlesA : [[access(read)]] Particles;
[[binding(2), group(0)]] var<storage> particlesB : [[access(read_write)]] Particles;
[[binding(1), group(0)]] var<storage, read> particlesA : Particles;
[[binding(2), group(0)]] var<storage, read_write> particlesB : Particles;
// https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
[[stage(compute)]]

4
src/dawn_native/QueryHelper.cpp
View File

@ -55,9 +55,9 @@ namespace dawn_native {
};
[[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;

2
src/dawn_native/d3d12/ShaderModuleD3D12.cpp
View File

@ -234,7 +234,7 @@ namespace dawn_native { namespace d3d12 {
bgl->GetBindingInfo(bindingIndex).buffer.type ==
wgpu::BufferBindingType::Storage);
if (forceStorageBufferAsUAV) {
accessControls.emplace(srcBindingPoint, tint::ast::AccessControl::kReadWrite);
accessControls.emplace(srcBindingPoint, tint::ast::Access::kReadWrite);
}
}
}

2
src/tests/DawnTest.cpp
View File

@ -1098,7 +1098,7 @@ std::ostringstream& DawnTestBase::ExpectSampledDepthData(wgpu::Texture texture,
};
[[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>

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

@ -72,8 +72,8 @@ class BindGroupTests : public DawnTest {
<< " : 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 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) {
};
[[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 @@ TEST_P(BindGroupTests, ReadonlyStorage) {
[[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 @@ TEST_P(BindGroupTests, ReallyLargeBindGroup) {
};
)";
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 @@ TEST_P(BindGroupTests, ReallyLargeBindGroup) {
};
)";
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";

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

@ -1035,7 +1035,7 @@ TEST_P(BufferZeroInitTest, BoundAsReadonlyStorageBuffer) {
[[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 @@ TEST_P(BufferZeroInitTest, BoundAsStorageBuffer) {
[[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() {

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

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

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

@ -37,7 +37,7 @@ class ComputeDispatchTests : public DawnTest {
};
[[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>) {

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

@ -78,7 +78,7 @@ TEST_P(ComputeSharedMemoryTests, Basic) {
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)]]

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

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

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

@ -146,7 +146,7 @@ class CopyTextureForBrowserTests : public DawnTest {
};
[[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 @@ class CopyTextureForBrowserTests : public DawnTest {
// 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, VerifyCopyOnYDirection) {
}
TEST_P(CopyTextureForBrowserTests, VerifyCopyFromLargeTexture) {
constexpr uint32_t kWidth = 899;
constexpr uint32_t kHeight = 999;

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

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

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

@ -211,7 +211,7 @@ TEST_P(D3D12CachingTests, ReuseShaderWithMultipleEntryPoints) {
[[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;

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

@ -141,8 +141,8 @@ class DepthStencilSamplingTest : public DawnTest {
<< " : 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 @@ class DepthStencilSamplingTest : public DawnTest {
<< " : 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 @@ class DepthStencilSamplingTest : public DawnTest {
[[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);

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

@ -125,9 +125,9 @@ class DynamicBufferOffsetTests : public DawnTest {
};
[[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 @@ class DynamicBufferOffsetTests : public DawnTest {
};
[[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) {

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

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

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

@ -132,7 +132,7 @@ struct VertexOutputs {
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"(

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

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

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

@ -100,7 +100,7 @@ class MultisampledSamplingTest : public DawnTest {
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) {

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

@ -63,10 +63,10 @@ class OpArrayLengthTest : public DawnTest {
};
// 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 @@ class OpArrayLengthTest : public DawnTest {
[[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 @@ TEST_P(OpArrayLengthTest, Compute) {
[[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);

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

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

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

@ -167,12 +167,12 @@ class StorageTextureTests : public DawnTest {
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 @@ TEST_P(StorageTextureTests, ReadonlyStorageTextureInComputeShader) {
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 @@ TEST_P(StorageTextureTests, Readonly2DArrayStorageTexture) {
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 @@ TEST_P(StorageTextureZeroInitTests, ReadonlyStorageTextureClearsToZeroInComputeP
};
[[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()) {

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

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

12
src/tests/perf_tests/ShaderRobustnessPerf.cpp
View File

@ -29,9 +29,9 @@ namespace {
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 @@ namespace {
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> {

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

@ -1278,8 +1278,8 @@ class SetBindGroupValidationTest : public ValidationTest {
[[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 @@ class SetBindGroupValidationTest : public ValidationTest {
[[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 @@ class SetBindGroupPersistenceValidationTest : public ValidationTest {
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 @@ class BindGroupLayoutCompatibilityTest : public ValidationTest {
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 @@ class BindGroupLayoutCompatibilityTest : public ValidationTest {
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;

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

@ -66,7 +66,7 @@ TEST_F(GetBindGroupLayoutTests, SameObject) {
[[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 @@ TEST_F(GetBindGroupLayoutTests, BindingType) {
[[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 @@ TEST_F(GetBindGroupLayoutTests, BindingType) {
[[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 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingType) {
[[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 @@ TEST_F(GetBindGroupLayoutTests, FromCorrectEntryPoint) {
[[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;

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

@ -80,12 +80,10 @@ namespace {
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();

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

@ -46,7 +46,7 @@ namespace {
[[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() {
})");

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

@ -460,7 +460,7 @@ TEST_F(RenderPipelineValidationTest, StorageBufferInVertexShaderNoLayout) {
[[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>();

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

@ -762,7 +762,7 @@ namespace {
[[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;

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

@ -82,9 +82,8 @@ class StorageTextureValidationTests : public ValidationTest {
}
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 @@ TEST_F(StorageTextureValidationTests, ComputePipeline) {
[[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;