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

Fix Storage Buffers in WGSL tests & examples 

Converts var<storage_buffer> to var<storage>

Adds in [[access(read_write)]] when no access qualifier was supplied.

BUG=dawn:699

Change-Id: I9da95366d0bb8734f1d134ee2dc764a7c16ef0cf
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/45300
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Austin Eng <enga@chromium.org>
This commit is contained in:
Ryan Harrison 2021-03-18 17:20:48 +00:00 committed by Commit Bot service account
parent 9895c273d6
commit bbabda3590
26 changed files with 83 additions and 83 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
examples/ComputeBoids.cpp
View File

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

4
src/dawn_native/QueryHelper.cpp
View File

@ -53,9 +53,9 @@ namespace dawn_native {
}; };
[[group(0), binding(0)]] [[group(0), binding(0)]]
var<storage_buffer> timestamps : [[access(read_write)]] TimestampArr; var<storage> timestamps : [[access(read_write)]] TimestampArr;
[[group(0), binding(1)]] [[group(0), binding(1)]]
var<storage_buffer> availability : [[access(read)]] AvailabilityArr; var<storage> availability : [[access(read)]] AvailabilityArr;
[[group(0), binding(2)]] var<uniform> params : TimestampParams; [[group(0), binding(2)]] var<uniform> params : TimestampParams;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;

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

@ -76,7 +76,7 @@ class BindGroupTests : public DawnTest {
<< " : Buffer" << i << ";"; << " : Buffer" << i << ";";
break; break;
case wgpu::BufferBindingType::Storage: case wgpu::BufferBindingType::Storage:
fs << "\n[[group(" << i << "), binding(0)]] var<storage_buffer> buffer" << i fs << "\n[[group(" << i << "), binding(0)]] var<storage> buffer" << i
<< " : [[access(read)]] Buffer" << i << ";"; << " : [[access(read)]] Buffer" << i << ";";
break; break;
default: default:
@ -854,9 +854,9 @@ TEST_P(BindGroupTests, DynamicOffsetOrder) {
}; };
[[group(0), binding(2)]] var<uniform> buffer2 : Buffer2; [[group(0), binding(2)]] var<uniform> buffer2 : Buffer2;
[[group(0), binding(3)]] var<storage_buffer> buffer3 : [[access(read)]] Buffer3; [[group(0), binding(3)]] var<storage> buffer3 : [[access(read)]] Buffer3;
[[group(0), binding(0)]] var<storage_buffer> buffer0 : [[access(read)]] Buffer0; [[group(0), binding(0)]] var<storage> buffer0 : [[access(read)]] Buffer0;
[[group(0), binding(4)]] var<storage_buffer> outputBuffer : [[access(read_write)]] OutputBuffer; [[group(0), binding(4)]] var<storage> outputBuffer : [[access(read_write)]] OutputBuffer;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
outputBuffer.value = vec3<u32>(buffer0.value, buffer2.value, buffer3.value); outputBuffer.value = vec3<u32>(buffer0.value, buffer2.value, buffer3.value);
@ -1123,7 +1123,7 @@ TEST_P(BindGroupTests, ReadonlyStorage) {
[[block]] struct Buffer0 { [[block]] struct Buffer0 {
color : vec4<f32>; color : vec4<f32>;
}; };
[[group(0), binding(0)]] var<storage_buffer> buffer0 : [[access(read)]] Buffer0; [[group(0), binding(0)]] var<storage> buffer0 : [[access(read)]] Buffer0;
[[location(0)]] var<out> fragColor : vec4<f32>; [[location(0)]] var<out> fragColor : vec4<f32>;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
@ -1261,8 +1261,8 @@ TEST_P(BindGroupTests, ReallyLargeBindGroup) {
}; };
)"; )";
interface << "[[group(0), binding(" << binding++ << ")]] " interface << "[[group(0), binding(" << binding++ << ")]] "
<< "var<storage_buffer> sbuf" << i << " : [[access(read)]] ReadOnlyStorageBuffer" << "var<storage> sbuf" << i << " : [[access(read)]] ReadOnlyStorageBuffer" << i
<< i << ";\n"; << ";\n";
body << "if (sbuf" << i << ".value != " << expectedValue++ << "u) {\n"; body << "if (sbuf" << i << ".value != " << expectedValue++ << "u) {\n";
body << " return;\n"; body << " return;\n";
@ -1278,7 +1278,7 @@ TEST_P(BindGroupTests, ReallyLargeBindGroup) {
}; };
)"; )";
interface << "[[group(0), binding(" << binding++ << ")]] " interface << "[[group(0), binding(" << binding++ << ")]] "
<< "var<storage_buffer> result : [[access(read_write)]] ReadWriteStorageBuffer;\n"; << "var<storage> result : [[access(read_write)]] ReadWriteStorageBuffer;\n";
body << "result.value = 1u;\n"; body << "result.value = 1u;\n";

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

@ -1037,7 +1037,7 @@ TEST_P(BufferZeroInitTest, BoundAsReadonlyStorageBuffer) {
[[block]] struct SSBO { [[block]] struct SSBO {
value : vec4<u32>; value : vec4<u32>;
}; };
[[group(0), binding(0)]] var<storage> ssbo : SSBO; [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
[[group(0), binding(1)]] var outImage : [[access(write)]] texture_storage_2d<rgba8unorm>; [[group(0), binding(1)]] var outImage : [[access(write)]] texture_storage_2d<rgba8unorm>;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {

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 // TODO(crbug.com/tint/386): Use the same struct type
[[set(0), binding(0)]] var<storage> src : Buf1; [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1;
[[set(0), binding(1)]] var<storage> dst : Buf2; [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
@ -124,8 +124,8 @@ TEST_P(ComputeCopyStorageBufferTests, SizedArrayOfStruct) {
}; };
// TODO(crbug.com/tint/386): Use the same struct type // TODO(crbug.com/tint/386): Use the same struct type
[[set(0), binding(0)]] var<storage> src : Buf1; [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1;
[[set(0), binding(1)]] var<storage> dst : Buf2; [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
@ -147,8 +147,8 @@ TEST_P(ComputeCopyStorageBufferTests, UnsizedArrayOfBasic) {
}; };
// TODO(crbug.com/tint/386): Use the same struct type // TODO(crbug.com/tint/386): Use the same struct type
[[set(0), binding(0)]] var<storage> src : Buf1; [[set(0), binding(0)]] var<storage> src : [[access(read_write)]] Buf1;
[[set(0), binding(1)]] var<storage> dst : Buf2; [[set(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf2;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> 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(0)]] var<uniform> input : InputBuf;
[[group(0), binding(1)]] var<storage_buffer> output : OutputBuf; [[group(0), binding(1)]] var<storage> output : [[access(read_write)]] OutputBuf;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;

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

@ -36,7 +36,7 @@ TEST_P(ComputeStorageBufferBarrierTests, AddIncrement) {
data : array<u32, 100>; data : array<u32, 100>;
}; };
[[group(0), binding(0)]] var<storage_buffer> buf : [[access(read_write)]] Buf; [[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
@ -92,8 +92,8 @@ TEST_P(ComputeStorageBufferBarrierTests, AddPingPong) {
data : array<u32, 100>; data : array<u32, 100>;
}; };
[[group(0), binding(0)]] var<storage_buffer> src : [[access(read_write)]] Src; [[group(0), binding(0)]] var<storage> src : [[access(read_write)]] Src;
[[group(0), binding(1)]] var<storage_buffer> dst : [[access(read_write)]] Dst; [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Dst;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
@ -163,8 +163,8 @@ TEST_P(ComputeStorageBufferBarrierTests, StorageAndReadonlyStoragePingPongInOneP
data : array<u32, 100>; data : array<u32, 100>;
}; };
[[group(0), binding(0)]] var<storage_buffer> src : [[access(read)]] Src; [[group(0), binding(0)]] var<storage> src : [[access(read)]] Src;
[[group(0), binding(1)]] var<storage_buffer> dst : [[access(read_write)]] Dst; [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Dst;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
@ -233,7 +233,7 @@ TEST_P(ComputeStorageBufferBarrierTests, UniformToStorageAddPingPong) {
}; };
[[group(0), binding(0)]] var<uniform> src : Buf; [[group(0), binding(0)]] var<uniform> src : Buf;
[[group(0), binding(1)]] var<storage_buffer> dst : [[access(read_write)]] Buf; [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
@ -301,7 +301,7 @@ TEST_P(ComputeStorageBufferBarrierTests, UniformToStorageAddPingPongInOnePass) {
}; };
[[group(0), binding(0)]] var<uniform> src : Buf; [[group(0), binding(0)]] var<uniform> src : Buf;
[[group(0), binding(1)]] var<storage_buffer> dst : [[access(read_write)]] Buf; [[group(0), binding(1)]] var<storage> dst : [[access(read_write)]] Buf;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {

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

@ -80,7 +80,7 @@ class CopyTextureForBrowserTests : public DawnTest {
}; };
[[group(0), binding(0)]] var src : texture_2d<f32>; [[group(0), binding(0)]] var src : texture_2d<f32>;
[[group(0), binding(1)]] var dst : texture_2d<f32>; [[group(0), binding(1)]] var dst : texture_2d<f32>;
[[group(0), binding(2)]] var<storage_buffer> output : OutputBuf; [[group(0), binding(2)]] var<storage> output : [[access(read_write)]] OutputBuf;
[[group(0), binding(3)]] var<uniform> uniforms : Uniforms; [[group(0), binding(3)]] var<uniform> uniforms : Uniforms;
[[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>; [[builtin(global_invocation_id)]] var<in> GlobalInvocationID : vec3<u32>;
[[stage(compute), workgroup_size(1, 1, 1)]] [[stage(compute), workgroup_size(1, 1, 1)]]

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

@ -38,7 +38,7 @@ TEST_P(CreatePipelineAsyncTest, BasicUseOfCreateComputePipelineAsync) {
[[block]] struct SSBO { [[block]] struct SSBO {
value : u32; value : u32;
}; };
[[group(0), binding(0)]] var<storage_buffer> ssbo : SSBO; [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
ssbo.value = 1u; ssbo.value = 1u;
@ -105,7 +105,7 @@ TEST_P(CreatePipelineAsyncTest, CreateComputePipelineFailed) {
[[block]] struct SSBO { [[block]] struct SSBO {
value : u32; value : u32;
}; };
[[group(0), binding(0)]] var<storage_buffer> ssbo : SSBO; [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] SSBO;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
ssbo.value = 1u; ssbo.value = 1u;

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

@ -229,7 +229,7 @@ TEST_P(D3D12CachingTests, ReuseShaderWithMultipleEntryPoints) {
[[block]] struct Data { [[block]] struct Data {
data : u32; data : u32;
}; };
[[binding(0), group(0)]] var<storage_buffer> data : Data; [[binding(0), group(0)]] var<storage> data : [[access(read_write)]] Data;
[[stage(compute)]] fn write1() -> void { [[stage(compute)]] fn write1() -> void {
data.data = 1u; data.data = 1u;

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

@ -139,8 +139,8 @@ class DepthStencilSamplingTest : public DawnTest {
<< " : texture_2d<f32>;\n"; << " : texture_2d<f32>;\n";
shaderSource << "[[group(0), binding(" << 2 * index + 1 shaderSource << "[[group(0), binding(" << 2 * index + 1
<< ")]] var<storage_buffer> result" << index << ")]] var<storage> result" << index
<< " : DepthResult;\n"; << " : [[access(read_write)]] DepthResult;\n";
shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index
<< ", vec2<i32>(0, 0), 0)[" << componentIndex << "];"; << ", vec2<i32>(0, 0), 0)[" << componentIndex << "];";
@ -150,8 +150,8 @@ class DepthStencilSamplingTest : public DawnTest {
<< " : texture_2d<u32>;\n"; << " : texture_2d<u32>;\n";
shaderSource << "[[group(0), binding(" << 2 * index + 1 shaderSource << "[[group(0), binding(" << 2 * index + 1
<< ")]] var<storage_buffer> result" << index << ")]] var<storage> result" << index
<< " : StencilResult;\n"; << " : [[access(read_write)]] StencilResult;\n";
shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index shaderBody << "\nresult" << index << ".value = textureLoad(tex" << index
<< ", vec2<i32>(0, 0), 0)[" << componentIndex << "];"; << ", vec2<i32>(0, 0), 0)[" << componentIndex << "];";
@ -223,7 +223,7 @@ class DepthStencilSamplingTest : public DawnTest {
[[block]] struct SamplerResult { [[block]] struct SamplerResult {
value : f32; value : f32;
}; };
[[group(0), binding(3)]] var<storage_buffer> samplerResult : SamplerResult; [[group(0), binding(3)]] var<storage> samplerResult : [[access(read_write)]] SamplerResult;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
samplerResult.value = textureSampleCompare(tex, samp, vec2<f32>(0.5, 0.5), uniforms.compareRef); samplerResult.value = textureSampleCompare(tex, samp, vec2<f32>(0.5, 0.5), uniforms.compareRef);

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

@ -126,9 +126,9 @@ class DynamicBufferOffsetTests : public DawnTest {
}; };
[[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1; [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1;
[[group(0), binding(1)]] var<storage_buffer> sBufferNotDynamic : [[access(read_write)]] Buffer2; [[group(0), binding(1)]] var<storage> sBufferNotDynamic : [[access(read_write)]] Buffer2;
[[group(0), binding(3)]] var<uniform> uBuffer : Buffer3; [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3;
[[group(0), binding(4)]] var<storage_buffer> sBuffer : [[access(read_write)]] Buffer4; [[group(0), binding(4)]] var<storage> sBuffer : [[access(read_write)]] Buffer4;
)"; )";
if (isInheritedPipeline) { if (isInheritedPipeline) {
@ -195,9 +195,9 @@ class DynamicBufferOffsetTests : public DawnTest {
}; };
[[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1; [[group(0), binding(0)]] var<uniform> uBufferNotDynamic : Buffer1;
[[group(0), binding(1)]] var<storage_buffer> sBufferNotDynamic : [[access(read_write)]] Buffer2; [[group(0), binding(1)]] var<storage> sBufferNotDynamic : [[access(read_write)]] Buffer2;
[[group(0), binding(3)]] var<uniform> uBuffer : Buffer3; [[group(0), binding(3)]] var<uniform> uBuffer : Buffer3;
[[group(0), binding(4)]] var<storage_buffer> sBuffer : [[access(read_write)]] Buffer4; [[group(0), binding(4)]] var<storage> sBuffer : [[access(read_write)]] Buffer4;
)"; )";
if (isInheritedPipeline) { if (isInheritedPipeline) {

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

@ -74,7 +74,7 @@ TEST_P(EntryPointTests, TwoComputeInModule) {
[[block]] struct Data { [[block]] struct Data {
data : u32; data : u32;
}; };
[[binding(0), group(0)]] var<storage_buffer> data : Data; [[binding(0), group(0)]] var<storage> data : [[access(read_write)]] Data;
[[stage(compute)]] fn write1() -> void { [[stage(compute)]] fn write1() -> void {
data.data = 1u; data.data = 1u;

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

@ -119,7 +119,7 @@ void FirstIndexOffsetTests::TestImpl(DrawMode mode,
instance_index : u32; instance_index : u32;
}; };
[[group(0), binding(0)]] var<storage_buffer> idx_vals : [[access(read_write)]] IndexVals; [[group(0), binding(0)]] var<storage> idx_vals : [[access(read_write)]] IndexVals;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
)"; )";

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

@ -39,7 +39,7 @@ class GpuMemorySyncTests : public DawnTest {
[[block]] struct Data { [[block]] struct Data {
a : i32; a : i32;
}; };
[[group(0), binding(0)]] var<storage_buffer> data : [[access(read_write)]] Data; [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
data.a = data.a + 1; data.a = data.a + 1;
})"); })");
@ -67,7 +67,7 @@ class GpuMemorySyncTests : public DawnTest {
[[block]] struct Data { [[block]] struct Data {
i : i32; i : i32;
}; };
[[group(0), binding(0)]] var<storage_buffer> data : [[access(read_write)]] Data; [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data;
[[location(0)]] var<out> fragColor : vec4<f32>; [[location(0)]] var<out> fragColor : vec4<f32>;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
data.i = data.i + 1; data.i = data.i + 1;
@ -257,7 +257,7 @@ TEST_P(GpuMemorySyncTests, SampledAndROStorageTextureInComputePass) {
sampledOut: u32; sampledOut: u32;
storageOut: u32; storageOut: u32;
}; };
[[group(0), binding(0)]] var<storage_buffer> output : [[access(write)]] Output; [[group(0), binding(0)]] var<storage> output : [[access(write)]] Output;
[[group(0), binding(1)]] var sampledTex : texture_2d<u32>; [[group(0), binding(1)]] var sampledTex : texture_2d<u32>;
[[group(0), binding(2)]] var storageTex : [[access(read)]] texture_storage_2d<r32uint>; [[group(0), binding(2)]] var storageTex : [[access(read)]] texture_storage_2d<r32uint>;
@ -317,7 +317,7 @@ class StorageToUniformSyncTests : public DawnTest {
[[block]] struct Data { [[block]] struct Data {
a : f32; a : f32;
}; };
[[group(0), binding(0)]] var<storage_buffer> data : [[access(read_write)]] Data; [[group(0), binding(0)]] var<storage> data : [[access(read_write)]] Data;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
data.a = 1.0; data.a = 1.0;
})"); })");
@ -516,12 +516,12 @@ TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) {
[[block]] struct VBContents { [[block]] struct VBContents {
pos : array<vec4<f32>, 4>; pos : array<vec4<f32>, 4>;
}; };
[[group(0), binding(0)]] var<storage_buffer> vbContents : [[access(read_write)]] VBContents; [[group(0), binding(0)]] var<storage> vbContents : [[access(read_write)]] VBContents;
[[block]] struct IBContents { [[block]] struct IBContents {
indices : array<vec4<i32>, 2>; indices : array<vec4<i32>, 2>;
}; };
[[group(0), binding(1)]] var<storage_buffer> ibContents : [[access(read_write)]] IBContents; [[group(0), binding(1)]] var<storage> ibContents : [[access(read_write)]] IBContents;
// TODO(crbug.com/tint/386): Use the same struct. // TODO(crbug.com/tint/386): Use the same struct.
[[block]] struct ColorContents1 { [[block]] struct ColorContents1 {
@ -530,8 +530,8 @@ TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) {
[[block]] struct ColorContents2 { [[block]] struct ColorContents2 {
color : f32; color : f32;
}; };
[[group(0), binding(2)]] var<storage_buffer> uniformContents : [[access(read_write)]] ColorContents1; [[group(0), binding(2)]] var<storage> uniformContents : [[access(read_write)]] ColorContents1;
[[group(0), binding(3)]] var<storage_buffer> storageContents : [[access(read_write)]] ColorContents2; [[group(0), binding(3)]] var<storage> storageContents : [[access(read_write)]] ColorContents2;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
vbContents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0); vbContents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0);
@ -586,7 +586,7 @@ TEST_P(MultipleWriteThenMultipleReadTests, SeparateBuffers) {
}; };
[[group(0), binding(0)]] var<uniform> uniformBuffer : Buf; [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf;
[[group(0), binding(1)]] var<storage_buffer> storageBuffer : [[access(read)]] Buf; [[group(0), binding(1)]] var<storage> storageBuffer : [[access(read)]] Buf;
[[location(0)]] var<out> fragColor : vec4<f32>; [[location(0)]] var<out> fragColor : vec4<f32>;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
@ -648,7 +648,7 @@ TEST_P(MultipleWriteThenMultipleReadTests, OneBuffer) {
[[align(256)]] color1 : f32; [[align(256)]] color1 : f32;
}; };
[[group(0), binding(0)]] var<storage_buffer> contents : [[access(read_write)]] Contents; [[group(0), binding(0)]] var<storage> contents : [[access(read_write)]] Contents;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
contents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0); contents.pos[0] = vec4<f32>(-1.0, 1.0, 0.0, 1.0);
@ -703,7 +703,7 @@ TEST_P(MultipleWriteThenMultipleReadTests, OneBuffer) {
color : f32; color : f32;
}; };
[[group(0), binding(0)]] var<uniform> uniformBuffer : Buf; [[group(0), binding(0)]] var<uniform> uniformBuffer : Buf;
[[group(0), binding(1)]] var<storage_buffer> storageBuffer : [[access(read)]] Buf; [[group(0), binding(1)]] var<storage> storageBuffer : [[access(read)]] Buf;
[[location(0)]] var<out> fragColor : vec4<f32>; [[location(0)]] var<out> fragColor : vec4<f32>;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {

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

@ -96,7 +96,7 @@ class MultisampledSamplingTest : public DawnTest {
colorSamples : array<f32, 4>; colorSamples : array<f32, 4>;
depthSamples : array<f32, 4>; depthSamples : array<f32, 4>;
}; };
[[group(0), binding(2)]] var<storage_buffer> results : [[access(read_write)]] Results; [[group(0), binding(2)]] var<storage> results : [[access(read_write)]] Results;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
for (var i : i32 = 0; i < 4; i = i + 1) { 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. // The length should be 1 because the buffer is 4-byte long.
[[group(0), binding(0)]] var<storage_buffer> buffer1 : [[access(read)]] DataBuffer1; [[group(0), binding(0)]] var<storage> buffer1 : [[access(read)]] DataBuffer1;
// The length should be 64 because the buffer is 256 bytes long. // The length should be 64 because the buffer is 256 bytes long.
[[group(0), binding(1)]] var<storage_buffer> buffer2 : [[access(read)]] DataBuffer2; [[group(0), binding(1)]] var<storage> buffer2 : [[access(read)]] DataBuffer2;
// The length should be (512 - 16*4) / 8 = 56 because the buffer is 512 bytes long // The length should be (512 - 16*4) / 8 = 56 because the buffer is 512 bytes long
// and the structure is 8 bytes big. // and the structure is 8 bytes big.
@ -79,7 +79,7 @@ class OpArrayLengthTest : public DawnTest {
[[size(64)]] garbage : mat4x4<f32>; [[size(64)]] garbage : mat4x4<f32>;
data : [[stride(8)]] array<Buffer3Data>; data : [[stride(8)]] array<Buffer3Data>;
}; };
[[group(0), binding(2)]] var<storage_buffer> buffer3 : [[access(read)]] Buffer3; [[group(0), binding(2)]] var<storage> buffer3 : [[access(read)]] Buffer3;
)"; )";
// See comments in the shader for an explanation of these values // See comments in the shader for an explanation of these values
@ -128,7 +128,7 @@ TEST_P(OpArrayLengthTest, Compute) {
[[block]] struct ResultBuffer { [[block]] struct ResultBuffer {
data : [[stride(4)]] array<u32, 3>; data : [[stride(4)]] array<u32, 3>;
}; };
[[group(1), binding(0)]] var<storage_buffer> result : [[access(read_write)]] ResultBuffer; [[group(1), binding(0)]] var<storage> result : [[access(read_write)]] ResultBuffer;
)" + mShaderInterface + R"( )" + mShaderInterface + R"(
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
result.data[0] = arrayLength(buffer1.data); 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>; data : array<u32, 19>;
}; };
[[group(0), binding(0)]] var<storage_buffer> buf : [[access(read_write)]] Buf; [[group(0), binding(0)]] var<storage> buf : [[access(read_write)]] Buf;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
const factor : f32 = 1.0001; const factor : f32 = 1.0001;

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

@ -712,7 +712,7 @@ TEST_P(StorageTextureTests, ReadonlyStorageTextureInComputeShader) {
result : u32; result : u32;
}; };
[[group(0), binding(1)]] var<storage_buffer> dstBuffer : DstBuffer; [[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer;
)" << CommonReadOnlyTestCode(format) )" << CommonReadOnlyTestCode(format)
<< R"( << R"(
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
@ -937,7 +937,7 @@ TEST_P(StorageTextureTests, Readonly2DArrayStorageTexture) {
result : u32; result : u32;
}; };
[[group(0), binding(1)]] var<storage_buffer> dstBuffer : DstBuffer; [[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer;
)" << CommonReadOnlyTestCode(kTextureFormat, true) )" << CommonReadOnlyTestCode(kTextureFormat, true)
<< R"( << R"(
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
@ -1207,7 +1207,7 @@ TEST_P(StorageTextureZeroInitTests, ReadonlyStorageTextureClearsToZeroInComputeP
}; };
[[group(0), binding(0)]] var srcImage : [[access(read)]] texture_storage_2d<r32uint>; [[group(0), binding(0)]] var srcImage : [[access(read)]] texture_storage_2d<r32uint>;
[[group(0), binding(1)]] var<storage_buffer> dstBuffer : DstBuffer; [[group(0), binding(1)]] var<storage> dstBuffer : [[access(read_write)]] DstBuffer;
)") + kCommonReadOnlyZeroInitTestCode + R"( )") + kCommonReadOnlyZeroInitTestCode + R"(
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
if (doTest()) { if (doTest()) {

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

@ -977,7 +977,7 @@ TEST_P(TextureZeroInitTest, ComputePassSampledTextureClear) {
[[block]] struct Result { [[block]] struct Result {
value : vec4<f32>; value : vec4<f32>;
}; };
[[group(0), binding(1)]] var<storage> result : Result; [[group(0), binding(1)]] var<storage> result : [[access(read_write)]] Result;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
result.value = textureLoad(tex, vec2<i32>(0,0), 0); result.value = textureLoad(tex, vec2<i32>(0,0), 0);
} }

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

@ -1094,8 +1094,8 @@ class SetBindGroupValidationTest : public ValidationTest {
[[group(0), binding(0)]] var<uniform> uBufferDynamic : S; [[group(0), binding(0)]] var<uniform> uBufferDynamic : S;
[[group(0), binding(1)]] var<uniform> uBuffer : S; [[group(0), binding(1)]] var<uniform> uBuffer : S;
[[group(0), binding(2)]] var<storage_buffer> sBufferDynamic : [[access(read_write)]] S; [[group(0), binding(2)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
[[group(0), binding(3)]] var<storage_buffer> sReadonlyBufferDynamic : [[access(read)]] S; [[group(0), binding(3)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})"); })");
@ -1117,8 +1117,8 @@ class SetBindGroupValidationTest : public ValidationTest {
[[group(0), binding(0)]] var<uniform> uBufferDynamic : S; [[group(0), binding(0)]] var<uniform> uBufferDynamic : S;
[[group(0), binding(1)]] var<uniform> uBuffer : S; [[group(0), binding(1)]] var<uniform> uBuffer : S;
[[group(0), binding(2)]] var<storage_buffer> sBufferDynamic : [[access(read_write)]] S; [[group(0), binding(2)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
[[group(0), binding(3)]] var<storage_buffer> sReadonlyBufferDynamic : [[access(read)]] S; [[group(0), binding(3)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
[[stage(compute), workgroup_size(4, 4, 1)]] fn main() -> void { [[stage(compute), workgroup_size(4, 4, 1)]] fn main() -> void {
})"); })");
@ -1543,7 +1543,7 @@ class SetBindGroupPersistenceValidationTest : public ValidationTest {
ss << "[[group(" << l << "), binding(" << b << ")]] "; ss << "[[group(" << l << "), binding(" << b << ")]] ";
switch (binding) { switch (binding) {
case wgpu::BufferBindingType::Storage: case wgpu::BufferBindingType::Storage:
ss << "var<storage_buffer> set" << l << "_binding" << b ss << "var<storage> set" << l << "_binding" << b
<< " : [[access(read_write)]] S;"; << " : [[access(read_write)]] S;";
break; break;
case wgpu::BufferBindingType::Uniform: case wgpu::BufferBindingType::Uniform:
@ -1711,8 +1711,8 @@ class BindGroupLayoutCompatibilityTest : public ValidationTest {
value : vec2<f32>; value : vec2<f32>;
}; };
[[group(0), binding(0)]] var<storage_buffer> sBufferDynamic : [[access(read_write)]] S; [[group(0), binding(0)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
[[group(1), binding(0)]] var<storage_buffer> sReadonlyBufferDynamic : [[access(read)]] S; [[group(1), binding(0)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})", })",
@ -1744,8 +1744,8 @@ class BindGroupLayoutCompatibilityTest : public ValidationTest {
value : vec2<f32>; value : vec2<f32>;
}; };
[[group(0), binding(0)]] var<storage_buffer> sBufferDynamic : [[access(read_write)]] S; [[group(0), binding(0)]] var<storage> sBufferDynamic : [[access(read_write)]] S;
[[group(1), binding(0)]] var<storage_buffer> sReadonlyBufferDynamic : [[access(read)]] S; [[group(1), binding(0)]] var<storage> sReadonlyBufferDynamic : [[access(read)]] S;
[[stage(compute), workgroup_size(4, 4, 1)]] fn main() -> void { [[stage(compute), workgroup_size(4, 4, 1)]] fn main() -> void {
})", })",

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

@ -62,7 +62,7 @@ TEST_F(GetBindGroupLayoutTests, SameObject) {
[[block]] struct S3 { [[block]] struct S3 {
pos : mat4x4<f32>; pos : mat4x4<f32>;
}; };
[[group(3), binding(0)]] var<storage_buffer> storage3 : S3; [[group(3), binding(0)]] var<storage> storage3 : [[access(read_write)]] S3;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})"); })");
@ -194,7 +194,7 @@ TEST_F(GetBindGroupLayoutTests, BindingType) {
[[block]] struct S { [[block]] struct S {
pos : vec4<f32>; pos : vec4<f32>;
}; };
[[group(0), binding(0)]] var<storage_buffer> ssbo : S; [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] S;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})"); })");
@ -219,7 +219,7 @@ TEST_F(GetBindGroupLayoutTests, BindingType) {
[[block]] struct S { [[block]] struct S {
pos : vec4<f32>; pos : vec4<f32>;
}; };
[[group(0), binding(0)]] var<storage_buffer> ssbo : [[access(read)]] S; [[group(0), binding(0)]] var<storage> ssbo : [[access(read)]] S;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})"); })");
@ -638,7 +638,7 @@ TEST_F(GetBindGroupLayoutTests, ConflictingBindingType) {
[[block]] struct S { [[block]] struct S {
pos : vec4<f32>; pos : vec4<f32>;
}; };
[[group(0), binding(0)]] var<storage_buffer> ssbo : S; [[group(0), binding(0)]] var<storage> ssbo : [[access(read_write)]] S;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})"); })");
@ -827,8 +827,8 @@ TEST_F(GetBindGroupLayoutTests, DISABLED_FromCorrectEntryPoint) {
[[block]] struct Data { [[block]] struct Data {
data : f32; data : f32;
}; };
[[binding 0, set 0]] var<storage_buffer> data0 : Data; [[binding 0, set 0]] var<storage> data0 : [[access(read_write)]] Data;
[[binding 1, set 0]] var<storage_buffer> data1 : Data; [[binding 1, set 0]] var<storage> data1 : [[access(read_write)]] Data;
fn compute0() -> void { fn compute0() -> void {
data0.data = 0.0; data0.data = 0.0;

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

@ -47,7 +47,7 @@ namespace {
[[block]] struct Storage { [[block]] struct Storage {
dummy : array<f32>; dummy : array<f32>;
}; };
[[group(1), binding(1)]] var<storage_buffer> ssbo : [[access(read_write)]] Storage; [[group(1), binding(1)]] var<storage> ssbo : [[access(read_write)]] Storage;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})"); })");

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

@ -486,7 +486,7 @@ TEST_F(RenderPipelineValidationTest, StorageBufferInVertexShaderNoLayout) {
[[block]] struct Dst { [[block]] struct Dst {
data : array<u32, 100>; data : array<u32, 100>;
}; };
[[group(0), binding(0)]] var<storage_buffer> dst : [[access(read_write)]] Dst; [[group(0), binding(0)]] var<storage> dst : [[access(read_write)]] Dst;
[[builtin(vertex_index)]] var<in> VertexIndex : u32; [[builtin(vertex_index)]] var<in> VertexIndex : u32;
[[stage(vertex)]] fn main() -> void { [[stage(vertex)]] fn main() -> void {
dst.data[VertexIndex] = 0x1234u; dst.data[VertexIndex] = 0x1234u;

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

@ -779,7 +779,7 @@ namespace {
[[block]] struct RBuffer { [[block]] struct RBuffer {
value : f32; value : f32;
}; };
[[group(0), binding(0)]] var<storage_buffer> rBuffer : [[access(read)]] RBuffer; [[group(0), binding(0)]] var<storage> rBuffer : [[access(read)]] RBuffer;
[[stage(fragment)]] fn main() -> void { [[stage(fragment)]] fn main() -> void {
})"); })");
utils::ComboRenderPipelineDescriptor2 pipelineDescriptor; utils::ComboRenderPipelineDescriptor2 pipelineDescriptor;
@ -819,7 +819,7 @@ namespace {
[[block]] struct RBuffer { [[block]] struct RBuffer {
value : f32; value : f32;
}; };
[[group(0), binding(0)]] var<storage_buffer> rBuffer : [[access(read)]] RBuffer; [[group(0), binding(0)]] var<storage> rBuffer : [[access(read)]] RBuffer;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
})"); })");
wgpu::ComputePipelineDescriptor pipelineDescriptor; wgpu::ComputePipelineDescriptor pipelineDescriptor;

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

@ -197,7 +197,7 @@ TEST_F(StorageTextureValidationTests, ComputePipeline) {
[[block]] struct Buf { [[block]] struct Buf {
data : f32; data : f32;
}; };
[[group(0), binding(1)]] var<storage_buffer> buf : [[access(read_write)]] Buf; [[group(0), binding(1)]] var<storage> buf : [[access(read_write)]] Buf;
[[stage(compute)]] fn main() -> void { [[stage(compute)]] fn main() -> void {
buf.data = textureLoad(image0, vec2<i32>(LocalInvocationID.xy)).x; buf.data = textureLoad(image0, vec2<i32>(LocalInvocationID.xy)).x;