encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-12-20 10:25:28 +00:00
Code Issues Packages Projects Releases Wiki Activity

Update src/tint unittests to new @stage format.

Browse Source 
This CL updates all of the Tint unittests to the new @stage shorter
syntax. This also updates the WGSL writer to emit the new short forms
instead of using the deprecated form.

Bug: tint:1503
Change-Id: I8c49e5319a19cccb5b4b5078f3ab39c50f31a9a8
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/92483
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
dan sinclair
2022-06-07 13:55:34 +00:00
committed by Dawn LUCI CQ
parent e7099d4f3b
commit b29892be09
4586 changed files with 10823 additions and 10820 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
src/tint/transform/add_empty_entry_point_test.cc
View File

@@ -31,7 +31,7 @@ TEST_F(AddEmptyEntryPointTest, ShouldRunEmptyModule) {
TEST_F(AddEmptyEntryPointTest, ShouldRunExistingEntryPoint) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn existing() {}
)";
@@ -42,7 +42,7 @@ TEST_F(AddEmptyEntryPointTest, EmptyModule) {
auto* src = R"()";
auto* expect = R"(
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn unused_entry_point() {
}
)";
@@ -54,7 +54,7 @@ fn unused_entry_point() {
TEST_F(AddEmptyEntryPointTest, ExistingEntryPoint) {
auto* src = R"(
@stage(fragment)
@fragment
fn main() {
}
)";
@@ -70,7 +70,7 @@ TEST_F(AddEmptyEntryPointTest, NameClash) {
auto* src = R"(var<private> unused_entry_point : f32;)";
auto* expect = R"(
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn unused_entry_point_1() {
}

52
src/tint/transform/add_spirv_block_attribute_test.cc
View File

@@ -41,7 +41,7 @@ struct S {
var<private> p : S;
@stage(fragment)
@fragment
fn main() {
p.f = 1.0;
}
@@ -60,7 +60,7 @@ struct S {
f : f32,
}
@stage(fragment)
@fragment
fn main() -> S {
return S();
}
@@ -77,7 +77,7 @@ TEST_F(AddSpirvBlockAttributeTest, BasicScalar) {
@group(0) @binding(0)
var<uniform> u : f32;
@stage(fragment)
@fragment
fn main() {
let f = u;
}
@@ -90,7 +90,7 @@ struct u_block {
@group(0) @binding(0) var<uniform> u : u_block;
@stage(fragment)
@fragment
fn main() {
let f = u.inner;
}
@@ -106,7 +106,7 @@ TEST_F(AddSpirvBlockAttributeTest, BasicArray) {
@group(0) @binding(0)
var<uniform> u : array<vec4<f32>, 4u>;
@stage(fragment)
@fragment
fn main() {
let a = u;
}
@@ -119,7 +119,7 @@ struct u_block {
@group(0) @binding(0) var<uniform> u : u_block;
@stage(fragment)
@fragment
fn main() {
let a = u.inner;
}
@@ -137,7 +137,7 @@ type Numbers = array<vec4<f32>, 4u>;
@group(0) @binding(0)
var<uniform> u : Numbers;
@stage(fragment)
@fragment
fn main() {
let a = u;
}
@@ -152,7 +152,7 @@ struct u_block {
@group(0) @binding(0) var<uniform> u : u_block;
@stage(fragment)
@fragment
fn main() {
let a = u.inner;
}
@@ -172,7 +172,7 @@ struct S {
@group(0) @binding(0)
var<uniform> u : S;
@stage(fragment)
@fragment
fn main() {
let f = u.f;
}
@@ -185,7 +185,7 @@ struct S {
@group(0) @binding(0) var<uniform> u : S;
@stage(fragment)
@fragment
fn main() {
let f = u.f;
}
@@ -209,7 +209,7 @@ struct Outer {
@group(0) @binding(0)
var<uniform> u : Outer;
@stage(fragment)
@fragment
fn main() {
let f = u.i.f;
}
@@ -226,7 +226,7 @@ struct Outer {
@group(0) @binding(0) var<uniform> u : Outer;
@stage(fragment)
@fragment
fn main() {
let f = u.i.f;
}
@@ -253,7 +253,7 @@ var<uniform> u0 : Outer;
@group(0) @binding(1)
var<uniform> u1 : Inner;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.f;
@@ -278,7 +278,7 @@ struct u1_block {
@group(0) @binding(1) var<uniform> u1 : u1_block;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.inner.f;
@@ -305,7 +305,7 @@ var<private> p : Outer;
@group(0) @binding(1)
var<uniform> u : Inner;
@stage(fragment)
@fragment
fn main() {
let f0 = p.i.f;
let f1 = u.f;
@@ -329,7 +329,7 @@ struct u_block {
@group(0) @binding(1) var<uniform> u : u_block;
@stage(fragment)
@fragment
fn main() {
let f0 = p.i.f;
let f1 = u.inner.f;
@@ -360,7 +360,7 @@ var<uniform> u1 : Inner;
@group(0) @binding(2)
var<uniform> u2 : Inner;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.f;
@@ -388,7 +388,7 @@ struct u1_block {
@group(0) @binding(2) var<uniform> u2 : u1_block;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.inner.f;
@@ -410,7 +410,7 @@ struct S {
@group(0) @binding(0)
var<uniform> u : S;
@stage(fragment)
@fragment
fn main() {
let f = u.f;
let a = array<S, 4>();
@@ -428,7 +428,7 @@ struct u_block {
@group(0) @binding(0) var<uniform> u : u_block;
@stage(fragment)
@fragment
fn main() {
let f = u.inner.f;
let a = array<S, 4>();
@@ -452,7 +452,7 @@ var<uniform> u0 : S;
@group(0) @binding(1)
var<uniform> u1 : S;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.f;
let f1 = u1.f;
@@ -473,7 +473,7 @@ struct u0_block {
@group(0) @binding(1) var<uniform> u1 : u0_block;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.inner.f;
let f1 = u1.inner.f;
@@ -506,7 +506,7 @@ var<uniform> u0 : MyOuter;
@group(0) @binding(1)
var<uniform> u1 : MyInner;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.f;
@@ -535,7 +535,7 @@ struct u1_block {
@group(0) @binding(1) var<uniform> u1 : u1_block;
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.inner.f;
@@ -549,7 +549,7 @@ fn main() {
TEST_F(AddSpirvBlockAttributeTest, Aliases_Nested_OuterBuffer_InnerBuffer_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.f;
@@ -574,7 +574,7 @@ struct Inner {
};
)";
auto* expect = R"(
@stage(fragment)
@fragment
fn main() {
let f0 = u0.i.f;
let f1 = u1.inner.f;

32
src/tint/transform/array_length_from_uniform_test.cc
View File

@@ -40,7 +40,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
}
)";
@@ -57,7 +57,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb.arr);
}
@@ -75,7 +75,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb.arr);
}
@@ -94,7 +94,7 @@ TEST_F(ArrayLengthFromUniformTest, Basic) {
auto* src = R"(
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb);
}
@@ -109,7 +109,7 @@ struct tint_symbol {
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 4u);
}
@@ -137,7 +137,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb.arr);
}
@@ -157,7 +157,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
}
@@ -197,7 +197,7 @@ struct SB4 {
@group(3) @binding(2) var<storage, read> sb4 : SB4;
@group(4) @binding(2) var<storage, read> sb5 : array<vec4<f32>>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = arrayLength(&(sb1.arr1));
var len2 : u32 = arrayLength(&(sb2.arr2));
@@ -240,7 +240,7 @@ struct SB4 {
@group(4) @binding(2) var<storage, read> sb5 : array<vec4<f32>>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
var len2 : u32 = ((tint_symbol_1.buffer_size[0u][1u] - 16u) / 16u);
@@ -289,7 +289,7 @@ struct SB4 {
@group(3) @binding(2) var<storage, read> sb4 : SB4;
@group(4) @binding(2) var<storage, read> sb5 : array<vec4<f32>>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = arrayLength(&(sb1.arr1));
var len3 : u32 = arrayLength(&sb3);
@@ -329,7 +329,7 @@ struct SB4 {
@group(4) @binding(2) var<storage, read> sb5 : array<vec4<f32>>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
var len3 : u32 = (tint_symbol_1.buffer_size[0u][2u] / 16u);
@@ -363,7 +363,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = &(sb.arr);
}
@@ -397,7 +397,7 @@ struct SB2 {
@group(1) @binding(2) var<storage, read> sb2 : SB2;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = arrayLength(&(sb1.arr1));
var len2 : u32 = arrayLength(&(sb2.arr2));
@@ -426,7 +426,7 @@ struct SB2 {
@group(1) @binding(2) var<storage, read> sb2 : SB2;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
var len2 : u32 = arrayLength(&(sb2.arr2));
@@ -449,7 +449,7 @@ fn main() {
TEST_F(ArrayLengthFromUniformTest, OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb.arr);
}
@@ -469,7 +469,7 @@ struct tint_symbol {
@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
}

28
src/tint/transform/binding_remapper_test.cc
View File

@@ -74,7 +74,7 @@ struct S {
@group(3) @binding(2) var<storage, read> b : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -99,7 +99,7 @@ struct S {
@group(3) @binding(2) var<storage, read> b : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -113,7 +113,7 @@ struct S {
@group(3) @binding(2) var<storage, read> b : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -143,7 +143,7 @@ struct S {
@group(4) @binding(3) var<storage, read> c : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -159,7 +159,7 @@ struct S {
@group(4) @binding(3) var<storage, read> c : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -187,7 +187,7 @@ struct S {
@group(3) @binding(2) var<storage, read> b : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -201,7 +201,7 @@ struct S {
@group(6) @binding(7) var<storage, write> b : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -235,7 +235,7 @@ struct S {
@group(5) @binding(4) var<storage, read> d : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
let x : i32 = (((a.i + b.i) + c.i) + d.i);
}
@@ -254,7 +254,7 @@ struct S {
@internal(disable_validation__binding_point_collision) @group(5) @binding(4) var<storage, read> d : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
let x : i32 = (((a.i + b.i) + c.i) + d.i);
}
@@ -287,12 +287,12 @@ struct S {
@group(5) @binding(4) var<storage, read> d : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f1() {
let x : i32 = (a.i + c.i);
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f2() {
let x : i32 = (b.i + d.i);
}
@@ -311,12 +311,12 @@ struct S {
@group(5) @binding(4) var<storage, read> d : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f1() {
let x : i32 = (a.i + c.i);
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f2() {
let x : i32 = (b.i + d.i);
}
@@ -345,7 +345,7 @@ struct S {
@group(3) @binding(2) var<storage, read> b : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";

44
src/tint/transform/calculate_array_length_test.cc
View File

@@ -38,7 +38,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
}
)";
@@ -55,7 +55,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb.arr);
}
@@ -68,7 +68,7 @@ TEST_F(CalculateArrayLengthTest, BasicArray) {
auto* src = R"(
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb);
}
@@ -80,7 +80,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(sb, &(tint_symbol_1));
@@ -103,7 +103,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len : u32 = arrayLength(&sb.arr);
}
@@ -120,7 +120,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(sb, &(tint_symbol_1));
@@ -142,7 +142,7 @@ struct S {
@group(0) @binding(0) var<storage, read> arr : array<S>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let len = arrayLength(&arr);
}
@@ -157,7 +157,7 @@ struct S {
@group(0) @binding(0) var<storage, read> arr : array<S>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(arr, &(tint_symbol_1));
@@ -179,7 +179,7 @@ struct S {
@group(0) @binding(0) var<storage, read> arr : array<array<S, 4>>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let len = arrayLength(&arr);
}
@@ -194,7 +194,7 @@ struct S {
@group(0) @binding(0) var<storage, read> arr : array<array<S, 4>>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(arr, &(tint_symbol_1));
@@ -212,7 +212,7 @@ TEST_F(CalculateArrayLengthTest, InSameBlock) {
auto* src = R"(
@group(0) @binding(0) var<storage, read> sb : array<i32>;;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : u32 = arrayLength(&sb);
var b : u32 = arrayLength(&sb);
@@ -226,7 +226,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(sb, &(tint_symbol_1));
@@ -251,7 +251,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : u32 = arrayLength(&sb.arr);
var b : u32 = arrayLength(&sb.arr);
@@ -270,7 +270,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(sb, &(tint_symbol_1));
@@ -295,7 +295,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
if (true) {
var len : u32 = arrayLength(&sb.arr);
@@ -318,7 +318,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
if (true) {
var tint_symbol_1 : u32 = 0u;
@@ -359,7 +359,7 @@ struct SB2 {
@group(0) @binding(2) var<storage, read> sb3 : array<i32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = arrayLength(&(sb1.arr1));
var len2 : u32 = arrayLength(&(sb2.arr2));
@@ -394,7 +394,7 @@ struct SB2 {
@group(0) @binding(2) var<storage, read> sb3 : array<i32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(sb1, &(tint_symbol_1));
@@ -427,7 +427,7 @@ struct SB {
@group(0) @binding(0) var<storage, read> a : SB;
@group(0) @binding(1) var<storage, read> b : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x = &a;
var a : u32 = arrayLength(&a.arr);
@@ -451,7 +451,7 @@ struct SB {
@group(0) @binding(1) var<storage, read> b : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(a, &(tint_symbol_1));
@@ -473,7 +473,7 @@ fn main() {
TEST_F(CalculateArrayLengthTest, OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var len1 : u32 = arrayLength(&(sb1.arr1));
var len2 : u32 = arrayLength(&(sb2.arr2));
@@ -508,7 +508,7 @@ fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_par
@internal(intrinsic_buffer_size)
fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var tint_symbol_1 : u32 = 0u;
tint_symbol(sb1, &(tint_symbol_1));

336
src/tint/transform/canonicalize_entry_point_io_test.cc
View File

File diff suppressed because it is too large Load Diff

76
src/tint/transform/decompose_memory_access_test.cc
View File

@@ -78,7 +78,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = sb.a;
var b : u32 = sb.b;
@@ -213,7 +213,7 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa
return arr;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = tint_symbol(sb, 0u);
var b : u32 = tint_symbol_1(sb, 4u);
@@ -247,7 +247,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, SB_BasicLoad_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = sb.a;
var b : u32 = sb.b;
@@ -382,7 +382,7 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa
return arr;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = tint_symbol(sb, 0u);
var b : u32 = tint_symbol_1(sb, 4u);
@@ -470,7 +470,7 @@ struct UB {
@group(0) @binding(0) var<uniform> ub : UB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = ub.a;
var b : u32 = ub.b;
@@ -605,7 +605,7 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa
return arr;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = tint_symbol(ub, 0u);
var b : u32 = tint_symbol_1(ub, 4u);
@@ -639,7 +639,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, UB_BasicLoad_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = ub.a;
var b : u32 = ub.b;
@@ -774,7 +774,7 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa
return arr;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a : i32 = tint_symbol(ub, 0u);
var b : u32 = tint_symbol_1(ub, 4u);
@@ -862,7 +862,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
sb.a = i32();
sb.b = u32();
@@ -1014,7 +1014,7 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa
}
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
tint_symbol(sb, 0u, i32());
tint_symbol_1(sb, 4u, u32());
@@ -1048,7 +1048,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, SB_BasicStore_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
sb.a = i32();
sb.b = u32();
@@ -1200,7 +1200,7 @@ fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_pa
}
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
tint_symbol(sb, 0u, i32());
tint_symbol_1(sb, 4u, u32());
@@ -1288,7 +1288,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : SB = sb;
}
@@ -1406,7 +1406,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
return SB(tint_symbol_1(buffer, (offset + 0u)), tint_symbol_2(buffer, (offset + 4u)), tint_symbol_3(buffer, (offset + 8u)), tint_symbol_4(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)), tint_symbol_6(buffer, (offset + 32u)), tint_symbol_7(buffer, (offset + 48u)), tint_symbol_8(buffer, (offset + 64u)), tint_symbol_9(buffer, (offset + 80u)), tint_symbol_10(buffer, (offset + 96u)), tint_symbol_11(buffer, (offset + 112u)), tint_symbol_12(buffer, (offset + 128u)), tint_symbol_13(buffer, (offset + 144u)), tint_symbol_14(buffer, (offset + 160u)), tint_symbol_15(buffer, (offset + 192u)), tint_symbol_16(buffer, (offset + 224u)), tint_symbol_17(buffer, (offset + 256u)), tint_symbol_18(buffer, (offset + 304u)), tint_symbol_19(buffer, (offset + 352u)), tint_symbol_20(buffer, (offset + 384u)), tint_symbol_21(buffer, (offset + 448u)), tint_symbol_22(buffer, (offset + 512u)));
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : SB = tint_symbol(sb, 0u);
}
@@ -1419,7 +1419,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, LoadStructure_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : SB = sb;
}
@@ -1537,7 +1537,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
return SB(tint_symbol_1(buffer, (offset + 0u)), tint_symbol_2(buffer, (offset + 4u)), tint_symbol_3(buffer, (offset + 8u)), tint_symbol_4(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)), tint_symbol_6(buffer, (offset + 32u)), tint_symbol_7(buffer, (offset + 48u)), tint_symbol_8(buffer, (offset + 64u)), tint_symbol_9(buffer, (offset + 80u)), tint_symbol_10(buffer, (offset + 96u)), tint_symbol_11(buffer, (offset + 112u)), tint_symbol_12(buffer, (offset + 128u)), tint_symbol_13(buffer, (offset + 144u)), tint_symbol_14(buffer, (offset + 160u)), tint_symbol_15(buffer, (offset + 192u)), tint_symbol_16(buffer, (offset + 224u)), tint_symbol_17(buffer, (offset + 256u)), tint_symbol_18(buffer, (offset + 304u)), tint_symbol_19(buffer, (offset + 352u)), tint_symbol_20(buffer, (offset + 384u)), tint_symbol_21(buffer, (offset + 448u)), tint_symbol_22(buffer, (offset + 512u)));
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : SB = tint_symbol(sb, 0u);
}
@@ -1604,7 +1604,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
sb = SB();
}
@@ -1760,7 +1760,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
tint_symbol_22(buffer, (offset + 512u), value.v);
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
tint_symbol(sb, 0u, SB());
}
@@ -1773,7 +1773,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, StoreStructure_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
sb = SB();
}
@@ -1929,7 +1929,7 @@ fn tint_symbol(@internal(disable_validation__ignore_constructible_function_param
tint_symbol_22(buffer, (offset + 512u), value.v);
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
tint_symbol(sb, 0u, SB());
}
@@ -1993,7 +1993,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : f32 = sb.b[4].b[1].b.z;
}
@@ -2030,7 +2030,7 @@ struct SB {
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : f32 = tint_symbol(sb, 712u);
}
@@ -2043,7 +2043,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : f32 = sb.b[4].b[1].b.z;
}
@@ -2080,7 +2080,7 @@ struct S1 {
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var x : f32 = tint_symbol(sb, 712u);
}
@@ -2133,7 +2133,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2166,7 +2166,7 @@ struct SB {
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2182,7 +2182,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, ComplexDynamicAccessChain_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2215,7 +2215,7 @@ struct S1 {
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2279,7 +2279,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2320,7 +2320,7 @@ struct SB {
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2336,7 +2336,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, ComplexDynamicAccessChainWithAliases_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2377,7 +2377,7 @@ struct S1 {
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> f32
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var i : i32 = 4;
var j : u32 = 1u;
@@ -2429,7 +2429,7 @@ struct SB {
@group(0) @binding(0) var<storage, read_write> sb : SB;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
atomicStore(&sb.a, 123);
atomicLoad(&sb.a);
@@ -2542,7 +2542,7 @@ struct atomic_compare_exchange_weak_ret_type_1 {
@internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body)
fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
tint_atomicStore(sb, 16u, 123);
tint_atomicLoad(sb, 16u);
@@ -2576,7 +2576,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, StorageBufferAtomics_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
atomicStore(&sb.a, 123);
atomicLoad(&sb.a);
@@ -2689,7 +2689,7 @@ struct atomic_compare_exchange_weak_ret_type_1 {
@internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body)
fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
tint_atomicStore(sb, 16u, 123);
tint_atomicLoad(sb, 16u);
@@ -2739,7 +2739,7 @@ struct S {
var<workgroup> w : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
atomicStore(&(w.a), 123);
atomicLoad(&(w.a));
@@ -2775,7 +2775,7 @@ fn main() {
TEST_F(DecomposeMemoryAccessTest, WorkgroupBufferAtomics_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
atomicStore(&(w.a), 123);
atomicLoad(&(w.a));

44
src/tint/transform/decompose_strided_array_test.cc
View File

@@ -71,7 +71,7 @@ TEST_F(DecomposeStridedArrayTest, Empty) {
TEST_F(DecomposeStridedArrayTest, PrivateDefaultStridedArray) {
// var<private> arr : @stride(4) array<f32, 4u>
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(4) array<f32, 4u> = a;
// let b : f32 = arr[1];
@@ -92,7 +92,7 @@ TEST_F(DecomposeStridedArrayTest, PrivateDefaultStridedArray) {
auto* expect = R"(
var<private> arr : array<f32, 4u>;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : array<f32, 4u> = arr;
let b : f32 = arr[1i];
@@ -107,7 +107,7 @@ fn f() {
TEST_F(DecomposeStridedArrayTest, PrivateStridedArray) {
// var<private> arr : @stride(32) array<f32, 4u>
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(32) array<f32, 4u> = a;
// let b : f32 = arr[1];
@@ -133,7 +133,7 @@ struct strided_arr {
var<private> arr : array<strided_arr, 4u>;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : array<strided_arr, 4u> = arr;
let b : f32 = arr[1i].el;
@@ -151,7 +151,7 @@ TEST_F(DecomposeStridedArrayTest, ReadUniformStridedArray) {
// };
// @group(0) @binding(0) var<uniform> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(32) array<f32, 4u> = s.a;
// let b : f32 = s.a[1];
@@ -181,7 +181,7 @@ struct S {
@group(0) @binding(0) var<uniform> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : array<strided_arr, 4u> = s.a;
let b : f32 = s.a[1i].el;
@@ -199,7 +199,7 @@ TEST_F(DecomposeStridedArrayTest, ReadUniformDefaultStridedArray) {
// };
// @group(0) @binding(0) var<uniform> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(16) array<vec4<f32>, 4u> = s.a;
// let b : f32 = s.a[1][2];
@@ -227,7 +227,7 @@ struct S {
@group(0) @binding(0) var<uniform> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : array<vec4<f32>, 4u> = s.a;
let b : f32 = s.a[1i][2i];
@@ -245,7 +245,7 @@ TEST_F(DecomposeStridedArrayTest, ReadStorageStridedArray) {
// };
// @group(0) @binding(0) var<storage> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(32) array<f32, 4u> = s.a;
// let b : f32 = s.a[1];
@@ -275,7 +275,7 @@ struct S {
@group(0) @binding(0) var<storage> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : array<strided_arr, 4u> = s.a;
let b : f32 = s.a[1i].el;
@@ -293,7 +293,7 @@ TEST_F(DecomposeStridedArrayTest, ReadStorageDefaultStridedArray) {
// };
// @group(0) @binding(0) var<storage> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(4) array<f32, 4u> = s.a;
// let b : f32 = s.a[1];
@@ -318,7 +318,7 @@ struct S {
@group(0) @binding(0) var<storage> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : array<f32, 4u> = s.a;
let b : f32 = s.a[1i];
@@ -336,7 +336,7 @@ TEST_F(DecomposeStridedArrayTest, WriteStorageStridedArray) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// s.a = @stride(32) array<f32, 4u>();
// s.a = @stride(32) array<f32, 4u>(1.0, 2.0, 3.0, 4.0);
@@ -371,7 +371,7 @@ struct S {
@group(0) @binding(0) var<storage, read_write> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
s.a = array<strided_arr, 4u>();
s.a = array<strided_arr, 4u>(strided_arr(1.0f), strided_arr(2.0f), strided_arr(3.0f), strided_arr(4.0f));
@@ -390,7 +390,7 @@ TEST_F(DecomposeStridedArrayTest, WriteStorageDefaultStridedArray) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// s.a = @stride(4) array<f32, 4u>();
// s.a = @stride(4) array<f32, 4u>(1.0, 2.0, 3.0, 4.0);
@@ -420,7 +420,7 @@ struct S {
@group(0) @binding(0) var<storage, read_write> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
s.a = array<f32, 4u>();
s.a = array<f32, 4u>(1.0f, 2.0f, 3.0f, 4.0f);
@@ -439,7 +439,7 @@ TEST_F(DecomposeStridedArrayTest, ReadWriteViaPointerLets) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a = &s.a;
// let b = &*&*(a);
@@ -479,7 +479,7 @@ struct S {
@group(0) @binding(0) var<storage, read_write> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let c = s.a;
let d = s.a[1i].el;
@@ -500,7 +500,7 @@ TEST_F(DecomposeStridedArrayTest, PrivateAliasedStridedArray) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : ARR = s.a;
// let b : f32 = s.a[1];
@@ -541,7 +541,7 @@ struct S {
@group(0) @binding(0) var<storage, read_write> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : ARR = s.a;
let b : f32 = s.a[1i].el;
@@ -564,7 +564,7 @@ TEST_F(DecomposeStridedArrayTest, PrivateNestedStridedArray) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a : ARR_B = s.a;
// let b : array<@stride(8) array<f32, 2u>, 3u> = s.a[3];
@@ -641,7 +641,7 @@ struct S {
@group(0) @binding(0) var<storage, read_write> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let a : ARR_B = s.a;
let b : array<ARR_A, 3u> = s.a[3i].el;

40
src/tint/transform/decompose_strided_matrix_test.cc
View File

@@ -62,7 +62,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadUniformMatrix) {
// };
// @group(0) @binding(0) var<uniform> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let x : mat2x2<f32> = s.m;
// }
@@ -99,7 +99,7 @@ fn arr_to_mat2x2_stride_32(arr : @stride(32) array<vec2<f32>, 2u>) -> mat2x2<f32
return mat2x2<f32>(arr[0u], arr[1u]);
}
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let x : mat2x2<f32> = arr_to_mat2x2_stride_32(s.m);
}
@@ -118,7 +118,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadUniformColumn) {
// };
// @group(0) @binding(0) var<uniform> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let x : vec2<f32> = s.m[1];
// }
@@ -152,7 +152,7 @@ struct S {
@group(0) @binding(0) var<uniform> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let x : vec2<f32> = s.m[1i];
}
@@ -171,7 +171,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadUniformMatrix_DefaultStride) {
// };
// @group(0) @binding(0) var<uniform> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let x : mat2x2<f32> = s.m;
// }
@@ -205,7 +205,7 @@ struct S {
@group(0) @binding(0) var<uniform> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let x : mat2x2<f32> = s.m;
}
@@ -224,7 +224,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadStorageMatrix) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let x : mat2x2<f32> = s.m;
// }
@@ -262,7 +262,7 @@ fn arr_to_mat2x2_stride_32(arr : @stride(32) array<vec2<f32>, 2u>) -> mat2x2<f32
return mat2x2<f32>(arr[0u], arr[1u]);
}
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let x : mat2x2<f32> = arr_to_mat2x2_stride_32(s.m);
}
@@ -281,7 +281,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadStorageColumn) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let x : vec2<f32> = s.m[1];
// }
@@ -316,7 +316,7 @@ struct S {
@group(0) @binding(0) var<storage, read_write> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let x : vec2<f32> = s.m[1i];
}
@@ -335,7 +335,7 @@ TEST_F(DecomposeStridedMatrixTest, WriteStorageMatrix) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
// }
@@ -374,7 +374,7 @@ fn mat2x2_stride_32_to_arr(m : mat2x2<f32>) -> @stride(32) array<vec2<f32>, 2u>
return @stride(32) array<vec2<f32>, 2u>(m[0u], m[1u]);
}
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
s.m = mat2x2_stride_32_to_arr(mat2x2<f32>(vec2<f32>(1.0f, 2.0f), vec2<f32>(3.0f, 4.0f)));
}
@@ -393,7 +393,7 @@ TEST_F(DecomposeStridedMatrixTest, WriteStorageColumn) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// s.m[1] = vec2<f32>(1.0, 2.0);
// }
@@ -427,7 +427,7 @@ struct S {
@group(0) @binding(0) var<storage, read_write> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
s.m[1i] = vec2<f32>(1.0f, 2.0f);
}
@@ -446,7 +446,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadWriteViaPointerLets) {
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let a = &s.m;
// let b = &*&*(a);
@@ -500,7 +500,7 @@ fn mat2x2_stride_32_to_arr(m : mat2x2<f32>) -> @stride(32) array<vec2<f32>, 2u>
return @stride(32) array<vec2<f32>, 2u>(m[0u], m[1u]);
}
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let x = arr_to_mat2x2_stride_32(s.m);
let y = s.m[1i];
@@ -523,7 +523,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadPrivateMatrix) {
// };
// var<private> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// let x : mat2x2<f32> = s.m;
// }
@@ -557,7 +557,7 @@ struct S {
var<private> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
let x : mat2x2<f32> = s.m;
}
@@ -576,7 +576,7 @@ TEST_F(DecomposeStridedMatrixTest, WritePrivateMatrix) {
// };
// var<private> s : S;
//
// @stage(compute) @workgroup_size(1)
// @compute @workgroup_size(1)
// fn f() {
// s.m = mat2x2<f32>(vec2<f32>(1.0, 2.0), vec2<f32>(3.0, 4.0));
// }
@@ -611,7 +611,7 @@ struct S {
var<private> s : S;
@stage(compute) @workgroup_size(1i)
@compute @workgroup_size(1i)
fn f() {
s.m = mat2x2<f32>(vec2<f32>(1.0f, 2.0f), vec2<f32>(3.0f, 4.0f));
}

2
src/tint/transform/disable_uniformity_analysis_test.cc
View File

@@ -54,7 +54,7 @@ TEST_F(DisableUniformityAnalysisTest, NonEmptyModule) {
auto* src = R"(
@group(0) @binding(0) var<storage, read> global : i32;
@stage(compute) @workgroup_size(64)
@compute @workgroup_size(64)
fn main() {
if ((global == 42)) {
workgroupBarrier();

62
src/tint/transform/first_index_offset_test.cc
View File

@@ -33,7 +33,7 @@ TEST_F(FirstIndexOffsetTest, ShouldRunEmptyModule) {
TEST_F(FirstIndexOffsetTest, ShouldRunFragmentStage) {
auto* src = R"(
@stage(fragment)
@fragment
fn entry() {
return;
}
@@ -44,7 +44,7 @@ fn entry() {
TEST_F(FirstIndexOffsetTest, ShouldRunVertexStage) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@@ -70,7 +70,7 @@ TEST_F(FirstIndexOffsetTest, EmptyModule) {
TEST_F(FirstIndexOffsetTest, BasicVertexShader) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@@ -95,7 +95,7 @@ fn test(vert_idx : u32) -> u32 {
return vert_idx;
}
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
test(vert_idx);
return vec4<f32>();
@@ -114,7 +114,7 @@ fn test(vert_idx : u32) -> u32 {
return vert_idx;
}
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
test((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
@@ -135,7 +135,7 @@ fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32>
TEST_F(FirstIndexOffsetTest, BasicModuleVertexIndex_OutOfOrder) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
test(vert_idx);
return vec4<f32>();
@@ -154,7 +154,7 @@ struct tint_symbol {
@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
test((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
@@ -183,7 +183,7 @@ fn test(inst_idx : u32) -> u32 {
return inst_idx;
}
@stage(vertex)
@vertex
fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
test(inst_idx);
return vec4<f32>();
@@ -202,7 +202,7 @@ fn test(inst_idx : u32) -> u32 {
return inst_idx;
}
@stage(vertex)
@vertex
fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
test((inst_idx + tint_symbol_1.first_instance_index));
return vec4<f32>();
@@ -223,7 +223,7 @@ fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32
TEST_F(FirstIndexOffsetTest, BasicModuleInstanceIndex_OutOfOrder) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
test(inst_idx);
return vec4<f32>();
@@ -242,7 +242,7 @@ struct tint_symbol {
@binding(1) @group(7) var<uniform> tint_symbol_1 : tint_symbol;
@stage(vertex)
@vertex
fn entry(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
test((inst_idx + tint_symbol_1.first_instance_index));
return vec4<f32>();
@@ -276,7 +276,7 @@ struct Inputs {
@builtin(vertex_index) vert_idx : u32,
};
@stage(vertex)
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
test(inputs.instance_idx, inputs.vert_idx);
return vec4<f32>();
@@ -302,7 +302,7 @@ struct Inputs {
vert_idx : u32,
}
@stage(vertex)
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
test((inputs.instance_idx + tint_symbol_1.first_instance_index), (inputs.vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
@@ -323,7 +323,7 @@ fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
TEST_F(FirstIndexOffsetTest, BasicModuleBothIndex_OutOfOrder) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
test(inputs.instance_idx, inputs.vert_idx);
return vec4<f32>();
@@ -347,7 +347,7 @@ struct tint_symbol {
@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
@stage(vertex)
@vertex
fn entry(inputs : Inputs) -> @builtin(position) vec4<f32> {
test((inputs.instance_idx + tint_symbol_1.first_instance_index), (inputs.vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
@@ -387,7 +387,7 @@ fn func2(vert_idx : u32) -> u32 {
return func1(vert_idx);
}
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func2(vert_idx);
return vec4<f32>();
@@ -410,7 +410,7 @@ fn func2(vert_idx : u32) -> u32 {
return func1(vert_idx);
}
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func2((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
@@ -431,7 +431,7 @@ fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32>
TEST_F(FirstIndexOffsetTest, NestedCalls_OutOfOrder) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func2(vert_idx);
return vec4<f32>();
@@ -454,7 +454,7 @@ struct tint_symbol {
@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func2((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
@@ -487,19 +487,19 @@ fn func(i : u32) -> u32 {
return i;
}
@stage(vertex)
@vertex
fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func(vert_idx);
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func(vert_idx + inst_idx);
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func(inst_idx);
return vec4<f32>();
@@ -518,19 +518,19 @@ fn func(i : u32) -> u32 {
return i;
}
@stage(vertex)
@vertex
fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func(((vert_idx + tint_symbol_1.first_vertex_index) + (inst_idx + tint_symbol_1.first_instance_index)));
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func((inst_idx + tint_symbol_1.first_instance_index));
return vec4<f32>();
@@ -551,19 +551,19 @@ fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f
TEST_F(FirstIndexOffsetTest, MultipleEntryPoints_OutOfOrder) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func(vert_idx);
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func(vert_idx + inst_idx);
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func(inst_idx);
return vec4<f32>();
@@ -582,19 +582,19 @@ struct tint_symbol {
@binding(1) @group(2) var<uniform> tint_symbol_1 : tint_symbol;
@stage(vertex)
@vertex
fn entry_a(@builtin(vertex_index) vert_idx : u32) -> @builtin(position) vec4<f32> {
func((vert_idx + tint_symbol_1.first_vertex_index));
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_b(@builtin(vertex_index) vert_idx : u32, @builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func(((vert_idx + tint_symbol_1.first_vertex_index) + (inst_idx + tint_symbol_1.first_instance_index)));
return vec4<f32>();
}
@stage(vertex)
@vertex
fn entry_c(@builtin(instance_index) inst_idx : u32) -> @builtin(position) vec4<f32> {
func((inst_idx + tint_symbol_1.first_instance_index));
return vec4<f32>();

50
src/tint/transform/localize_struct_array_assignment_test.cc
View File

@@ -46,7 +46,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -69,7 +69,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -88,7 +88,7 @@ fn main() {
TEST_F(LocalizeStructArrayAssignmentTest, StructArray_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -111,7 +111,7 @@ struct Uniforms {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -162,7 +162,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -189,7 +189,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -208,7 +208,7 @@ fn main() {
TEST_F(LocalizeStructArrayAssignmentTest, StructStructArray_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -235,7 +235,7 @@ struct Uniforms {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -287,7 +287,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -311,7 +311,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -348,7 +348,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -375,7 +375,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -413,7 +413,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s : OuterS;
@@ -441,7 +441,7 @@ struct OuterS {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s : OuterS;
@@ -488,7 +488,7 @@ fn getNextIndex() -> u32 {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s : OuterS;
@@ -523,7 +523,7 @@ fn getNextIndex() -> u32 {
@group(1) @binding(4) var<uniform> uniforms : Uniforms;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s : OuterS;
@@ -545,7 +545,7 @@ fn main() {
TEST_F(LocalizeStructArrayAssignmentTest, IndexingWithSideEffectFunc_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s : OuterS;
@@ -579,7 +579,7 @@ struct InnerS {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s : OuterS;
@@ -643,7 +643,7 @@ fn f(p : ptr<function, OuterS>) {
(*p).a1[uniforms.i] = v;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var s1 : OuterS;
f(&s1);
@@ -675,7 +675,7 @@ fn f(p : ptr<function, OuterS>) {
}
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var s1 : OuterS;
f(&(s1));
@@ -688,7 +688,7 @@ fn main() {
TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var s1 : OuterS;
f(&s1);
@@ -714,7 +714,7 @@ struct Uniforms {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var s1 : OuterS;
f(&(s1));
@@ -769,7 +769,7 @@ fn f(p : ptr<function, InnerS>, v : InnerS) {
*(p) = v;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -797,7 +797,7 @@ fn f(p : ptr<function, InnerS>, v : InnerS) {
*(p) = v;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var v : InnerS;
var s1 : OuterS;
@@ -831,7 +831,7 @@ fn f(i : u32) -> u32 {
return (i + 1u);
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var s1 : OuterS;
var v : vec3<f32>;

120
src/tint/transform/module_scope_var_to_entry_point_param_test.cc
View File

@@ -42,14 +42,14 @@ TEST_F(ModuleScopeVarToEntryPointParamTest, Basic) {
var<private> p : f32;
var<workgroup> w : f32;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
w = p;
}
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol : f32;
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol_1 : f32;
@@ -64,7 +64,7 @@ fn main() {
TEST_F(ModuleScopeVarToEntryPointParamTest, Basic_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
w = p;
}
@@ -74,7 +74,7 @@ var<private> p : f32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol : f32;
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol_1 : f32;
@@ -106,7 +106,7 @@ fn foo(a : f32) {
no_uses();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo(1.0);
}
@@ -127,7 +127,7 @@ fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(di
no_uses();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol_4 : f32;
@internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol_5 : f32;
@@ -142,7 +142,7 @@ fn main() {
TEST_F(ModuleScopeVarToEntryPointParamTest, FunctionCalls_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo(1.0);
}
@@ -166,7 +166,7 @@ var<workgroup> w : f32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32;
@internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol_1 : f32;
@@ -198,14 +198,14 @@ TEST_F(ModuleScopeVarToEntryPointParamTest, Constructors) {
var<private> a : f32 = 1.0;
var<private> b : f32 = f32();
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x : f32 = a + b;
}
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32 = 1.0;
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol_1 : f32 = f32();
@@ -220,7 +220,7 @@ fn main() {
TEST_F(ModuleScopeVarToEntryPointParamTest, Constructors_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x : f32 = a + b;
}
@@ -230,7 +230,7 @@ var<private> a : f32 = 1.0;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32 = 1.0;
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol_1 : f32 = f32();
@@ -248,7 +248,7 @@ TEST_F(ModuleScopeVarToEntryPointParamTest, Pointers) {
var<private> p : f32;
var<workgroup> w : f32;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let p_ptr : ptr<private, f32> = &p;
let w_ptr : ptr<workgroup, f32> = &w;
@@ -258,7 +258,7 @@ fn main() {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32;
@internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol_1 : f32;
@@ -276,7 +276,7 @@ fn main() {
TEST_F(ModuleScopeVarToEntryPointParamTest, Pointers_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let p_ptr : ptr<private, f32> = &p;
let w_ptr : ptr<workgroup, f32> = &w;
@@ -289,7 +289,7 @@ var<private> p : f32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32;
@internal(disable_validation__ignore_storage_class) var<workgroup> tint_symbol_1 : f32;
@@ -317,7 +317,7 @@ fn foo() {
bar(&v);
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo();
}
@@ -332,7 +332,7 @@ fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_val
bar(tint_symbol);
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol_1 : f32;
foo(&(tint_symbol_1));
@@ -346,7 +346,7 @@ fn main() {
TEST_F(ModuleScopeVarToEntryPointParamTest, FoldAddressOfDeref_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo();
}
@@ -363,7 +363,7 @@ var<private> v : f32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_storage_class) var<private> tint_symbol : f32;
foo(&(tint_symbol));
@@ -394,7 +394,7 @@ var<uniform> u : S;
@group(0) @binding(1)
var<storage> s : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = u;
_ = s;
@@ -406,7 +406,7 @@ struct S {
a : f32,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr<storage, S>) {
_ = *(tint_symbol);
_ = *(tint_symbol_1);
@@ -420,7 +420,7 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_Basic_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = u;
_ = s;
@@ -436,7 +436,7 @@ struct S {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr<storage, S>) {
_ = *(tint_symbol);
_ = *(tint_symbol_1);
@@ -457,7 +457,7 @@ TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray) {
@group(0) @binding(0)
var<storage> buffer : array<f32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = buffer[0];
}
@@ -468,7 +468,7 @@ struct tint_symbol_1 {
arr : array<f32>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
_ = (*(tint_symbol)).arr[0];
}
@@ -481,7 +481,7 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = buffer[0];
}
@@ -495,7 +495,7 @@ struct tint_symbol_1 {
arr : array<f32>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
_ = (*(tint_symbol)).arr[0];
}
@@ -515,7 +515,7 @@ fn foo() {
_ = buffer[0];
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo();
}
@@ -530,7 +530,7 @@ fn foo(@internal(disable_validation__ignore_storage_class) @internal(disable_val
_ = (*(tint_symbol))[0];
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr<storage, tint_symbol_2>) {
foo(&((*(tint_symbol_1)).arr));
}
@@ -543,7 +543,7 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArrayInsideFunction_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo();
}
@@ -560,7 +560,7 @@ struct tint_symbol_1 {
arr : array<f32>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
foo(&((*(tint_symbol)).arr));
}
@@ -582,7 +582,7 @@ type myarray = array<f32>;
@group(0) @binding(0)
var<storage> buffer : myarray;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = buffer[0];
}
@@ -595,7 +595,7 @@ struct tint_symbol_1 {
type myarray = array<f32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
_ = (*(tint_symbol)).arr[0];
}
@@ -608,7 +608,7 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray_Alias_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = buffer[0];
}
@@ -623,7 +623,7 @@ struct tint_symbol_1 {
arr : array<f32>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
_ = (*(tint_symbol)).arr[0];
}
@@ -645,7 +645,7 @@ struct S {
@group(0) @binding(0)
var<storage> buffer : array<S>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = buffer[0];
}
@@ -660,7 +660,7 @@ struct tint_symbol_1 {
arr : array<S>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
_ = (*(tint_symbol)).arr[0];
}
@@ -673,7 +673,7 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_ArrayOfStruct_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = buffer[0];
}
@@ -694,7 +694,7 @@ struct tint_symbol_1 {
arr : array<S>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<storage, tint_symbol_1>) {
_ = (*(tint_symbol)).arr[0];
}
@@ -731,7 +731,7 @@ fn foo(a : f32) {
no_uses();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo(1.0);
}
@@ -757,7 +757,7 @@ fn foo(a : f32, @internal(disable_validation__ignore_storage_class) @internal(di
no_uses();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_4 : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_5 : ptr<storage, S>) {
foo(1.0, tint_symbol_4, tint_symbol_5);
}
@@ -770,7 +770,7 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_FunctionCalls_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo(1.0);
}
@@ -801,7 +801,7 @@ var<storage> s : S;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol : ptr<uniform, S>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) @internal(disable_validation__ignore_storage_class) tint_symbol_1 : ptr<storage, S>) {
foo(1.0, tint_symbol, tint_symbol_1);
}
@@ -836,7 +836,7 @@ TEST_F(ModuleScopeVarToEntryPointParamTest, HandleTypes_Basic) {
@group(0) @binding(0) var t : texture_2d<f32>;
@group(0) @binding(1) var s : sampler;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
_ = t;
_ = s;
@@ -844,7 +844,7 @@ fn main() {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) tint_symbol : texture_2d<f32>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) tint_symbol_1 : sampler) {
_ = tint_symbol;
_ = tint_symbol_1;
@@ -876,7 +876,7 @@ fn foo(a : f32) {
no_uses();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo(1.0);
}
@@ -898,7 +898,7 @@ fn foo(a : f32, tint_symbol_2 : texture_2d<f32>, tint_symbol_3 : sampler) {
no_uses();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) tint_symbol_4 : texture_2d<f32>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) tint_symbol_5 : sampler) {
foo(1.0, tint_symbol_4, tint_symbol_5);
}
@@ -911,7 +911,7 @@ fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_paramete
TEST_F(ModuleScopeVarToEntryPointParamTest, HandleTypes_FunctionCalls_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
foo(1.0);
}
@@ -936,7 +936,7 @@ fn bar(a : f32, b : f32) {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@group(0) @binding(0) @internal(disable_validation__entry_point_parameter) tint_symbol : texture_2d<f32>, @group(0) @binding(1) @internal(disable_validation__entry_point_parameter) tint_symbol_1 : sampler) {
foo(1.0, tint_symbol, tint_symbol_1);
}
@@ -966,7 +966,7 @@ TEST_F(ModuleScopeVarToEntryPointParamTest, Matrix) {
auto* src = R"(
var<workgroup> m : mat2x2<f32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x = m;
}
@@ -977,7 +977,7 @@ struct tint_symbol_2 {
m : mat2x2<f32>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
let tint_symbol : ptr<workgroup, mat2x2<f32>> = &((*(tint_symbol_1)).m);
let x = *(tint_symbol);
@@ -999,7 +999,7 @@ struct S2 {
};
var<workgroup> m : array<S2, 4>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x = m;
}
@@ -1018,7 +1018,7 @@ struct tint_symbol_2 {
m : array<S2, 4u>,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
let tint_symbol : ptr<workgroup, array<S2, 4u>> = &((*(tint_symbol_1)).m);
let x = *(tint_symbol);
@@ -1042,7 +1042,7 @@ var<workgroup> a : S;
var<workgroup> b : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x = a;
let y = b;
@@ -1059,7 +1059,7 @@ struct tint_symbol_3 {
b : S,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
let tint_symbol : ptr<workgroup, S> = &((*(tint_symbol_1)).a);
let tint_symbol_2 : ptr<workgroup, S> = &((*(tint_symbol_1)).b);
@@ -1077,7 +1077,7 @@ fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr
// variables that are promoted to threadgroup memory arguments.
TEST_F(ModuleScopeVarToEntryPointParamTest, DuplicateThreadgroupArgumentTypes_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x = a;
let y = b;
@@ -1101,7 +1101,7 @@ struct tint_symbol_3 {
b : S,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@internal(disable_validation__entry_point_parameter) tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
let tint_symbol : ptr<workgroup, S> = &((*(tint_symbol_1)).a);
let tint_symbol_2 : ptr<workgroup, S> = &((*(tint_symbol_1)).b);
@@ -1132,7 +1132,7 @@ var<storage> sb : S;
@group(0) @binding(2) var t : texture_2d<f32>;
@group(0) @binding(3) var s : sampler;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
}
)";
@@ -1142,7 +1142,7 @@ struct S {
a : f32,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
}
)";

76
src/tint/transform/multiplanar_external_texture_test.cc
View File

@@ -56,7 +56,7 @@ TEST_F(MultiplanarExternalTextureTest, ErrorNoPassedData) {
@group(0) @binding(0) var s : sampler;
@group(0) @binding(1) var ext_tex : texture_external;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleLevel(ext_tex, s, coord.xy);
}
@@ -74,7 +74,7 @@ TEST_F(MultiplanarExternalTextureTest, ErrorIncorrectBindingPont) {
@group(0) @binding(0) var s : sampler;
@group(0) @binding(1) var ext_tex : texture_external;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleLevel(ext_tex, s, coord.xy);
}
@@ -96,7 +96,7 @@ TEST_F(MultiplanarExternalTextureTest, Dimensions) {
auto* src = R"(
@group(0) @binding(0) var ext_tex : texture_external;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
dim = textureDimensions(ext_tex);
@@ -130,7 +130,7 @@ struct ExternalTextureParams {
@group(0) @binding(0) var ext_tex : texture_2d<f32>;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
dim = textureDimensions(ext_tex);
@@ -148,7 +148,7 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
// Tests that the transform works with a textureDimensions call.
TEST_F(MultiplanarExternalTextureTest, Dimensions_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
dim = textureDimensions(ext_tex);
@@ -182,7 +182,7 @@ struct ExternalTextureParams {
@group(0) @binding(2) var<uniform> ext_tex_params : ExternalTextureParams;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
var dim : vec2<i32>;
dim = textureDimensions(ext_tex);
@@ -205,7 +205,7 @@ TEST_F(MultiplanarExternalTextureTest, BasicTextureSampleLevel) {
@group(0) @binding(0) var s : sampler;
@group(0) @binding(1) var ext_tex : texture_external;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleLevel(ext_tex, s, coord.xy);
}
@@ -259,7 +259,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
return vec4<f32>(color, 1.0f);
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params);
}
@@ -275,7 +275,7 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
// Test that the transform works with a textureSampleLevel call.
TEST_F(MultiplanarExternalTextureTest, BasicTextureSampleLevel_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleLevel(ext_tex, s, coord.xy);
}
@@ -328,7 +328,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
return vec4<f32>(color, 1.0f);
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params);
}
@@ -350,7 +350,7 @@ TEST_F(MultiplanarExternalTextureTest, BasicTextureLoad) {
auto* src = R"(
@group(0) @binding(0) var ext_tex : texture_external;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureLoad(ext_tex, vec2<i32>(1, 1));
}
@@ -402,7 +402,7 @@ fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord
return vec4<f32>(color, 1.0f);
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params);
}
@@ -418,7 +418,7 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
// Tests that the transform works with a textureLoad call.
TEST_F(MultiplanarExternalTextureTest, BasicTextureLoad_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureLoad(ext_tex, vec2<i32>(1, 1));
}
@@ -470,7 +470,7 @@ fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord
return vec4<f32>(color, 1.0f);
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params);
}
@@ -492,7 +492,7 @@ TEST_F(MultiplanarExternalTextureTest, TextureSampleAndTextureLoad) {
@group(0) @binding(0) var s : sampler;
@group(0) @binding(1) var ext_tex : texture_external;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleLevel(ext_tex, s, coord.xy) + textureLoad(ext_tex, vec2<i32>(1, 1));
}
@@ -559,7 +559,7 @@ fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord
return vec4<f32>(color, 1.0f);
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return (textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params) + textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params));
}
@@ -576,7 +576,7 @@ fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
// call.
TEST_F(MultiplanarExternalTextureTest, TextureSampleAndTextureLoad_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleLevel(ext_tex, s, coord.xy) + textureLoad(ext_tex, vec2<i32>(1, 1));
}
@@ -642,7 +642,7 @@ fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord
return vec4<f32>(color, 1.0f);
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return (textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params) + textureLoadExternal(ext_tex, ext_tex_plane_1, vec2<i32>(1, 1), ext_tex_params));
}
@@ -668,7 +668,7 @@ TEST_F(MultiplanarExternalTextureTest, ManyTextureSampleLevel) {
@group(0) @binding(3) var ext_tex_2 : texture_external;
@group(1) @binding(0) var ext_tex_3 : texture_external;
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return textureSampleLevel(ext_tex, s, coord.xy) + textureSampleLevel(ext_tex_1, s, coord.xy) + textureSampleLevel(ext_tex_2, s, coord.xy) + textureSampleLevel(ext_tex_3, s, coord.xy);
}
@@ -740,7 +740,7 @@ fn textureSampleExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, smp
return vec4<f32>(color, 1.0f);
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord : vec4<f32>) -> @location(0) vec4<f32> {
return (((textureSampleExternal(ext_tex, ext_tex_plane_1, s, coord.xy, ext_tex_params) + textureSampleExternal(ext_tex_1, ext_tex_plane_1_1, s, coord.xy, ext_tex_params_1)) + textureSampleExternal(ext_tex_2, ext_tex_plane_1_2, s, coord.xy, ext_tex_params_2)) + textureSampleExternal(ext_tex_3, ext_tex_plane_1_3, s, coord.xy, ext_tex_params_3));
}
@@ -768,7 +768,7 @@ fn f(t : texture_external, s : sampler) {
@group(0) @binding(0) var ext_tex : texture_external;
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp);
}
@@ -826,7 +826,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
}
@@ -843,7 +843,7 @@ fn main() {
// correct output.
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParam_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp);
}
@@ -880,7 +880,7 @@ struct ExternalTextureParams {
@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
}
@@ -932,7 +932,7 @@ fn f(s : sampler, t : texture_external) {
@group(0) @binding(0) var ext_tex : texture_external;
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(smp, ext_tex);
}
@@ -990,7 +990,7 @@ fn f(s : sampler, t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(smp, ext_tex, ext_tex_plane_1, ext_tex_params);
}
@@ -1016,7 +1016,7 @@ fn f(t : texture_external, s : sampler, t2 : texture_external) {
@group(0) @binding(1) var smp : sampler;
@group(0) @binding(2) var ext_tex2 : texture_external;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp, ext_tex2);
}
@@ -1081,7 +1081,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2
@group(0) @binding(2) var ext_tex2 : texture_2d<f32>;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp, ext_tex2, ext_tex_plane_1_1, ext_tex_params_1);
}
@@ -1099,7 +1099,7 @@ fn main() {
// correct output.
TEST_F(MultiplanarExternalTextureTest, ExternalTexturePassedAsParamMultiple_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp, ext_tex2);
}
@@ -1143,7 +1143,7 @@ struct ExternalTextureParams {
@group(0) @binding(6) var<uniform> ext_tex_params_1 : ExternalTextureParams;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp, ext_tex2, ext_tex_plane_1_1, ext_tex_params_1);
}
@@ -1203,7 +1203,7 @@ fn f(t : texture_external, s : sampler) {
@group(0) @binding(0) var ext_tex : texture_external;
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp);
}
@@ -1265,7 +1265,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
}
@@ -1293,7 +1293,7 @@ fn f(t : texture_external, s : sampler) {
@group(0) @binding(0) var ext_tex : texture_external;
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp);
}
@@ -1355,7 +1355,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_2 : texture_2d<f32>, ext_tex_params_2
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
}
@@ -1421,7 +1421,7 @@ fn f(t : ET, s : sampler) {
@group(0) @binding(0) var ext_tex : ET;
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp);
}
@@ -1481,7 +1481,7 @@ fn f(t : texture_2d<f32>, ext_tex_plane_1_1 : texture_2d<f32>, ext_tex_params_1
@group(0) @binding(1) var smp : sampler;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
}
@@ -1497,7 +1497,7 @@ fn main() {
// Tests that the the transform handles aliases to external textures
TEST_F(MultiplanarExternalTextureTest, ExternalTextureAlias_OutOfOrder) {
auto* src = R"(
@stage(fragment)
@fragment
fn main() {
f(ext_tex, smp);
}
@@ -1536,7 +1536,7 @@ struct ExternalTextureParams {
@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
@stage(fragment)
@fragment
fn main() {
f(ext_tex, ext_tex_plane_1, ext_tex_params, smp);
}

64
src/tint/transform/num_workgroups_from_uniform_test.cc
View File

@@ -33,7 +33,7 @@ TEST_F(NumWorkgroupsFromUniformTest, ShouldRunEmptyModule) {
TEST_F(NumWorkgroupsFromUniformTest, ShouldRunHasNumWorkgroups) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
}
)";
@@ -43,7 +43,7 @@ fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
TEST_F(NumWorkgroupsFromUniformTest, Error_MissingTransformData) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
}
)";
@@ -61,7 +61,7 @@ fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
TEST_F(NumWorkgroupsFromUniformTest, Basic) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
let groups_x = num_wgs.x;
let groups_y = num_wgs.y;
@@ -82,7 +82,7 @@ fn main_inner(num_wgs : vec3<u32>) {
let groups_z = num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
main_inner(tint_symbol_3.num_workgroups);
}
@@ -101,7 +101,7 @@ struct Builtins {
@builtin(num_workgroups) num_wgs : vec3<u32>,
};
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(in : Builtins) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
@@ -126,7 +126,7 @@ fn main_inner(in : Builtins) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
main_inner(Builtins(tint_symbol_3.num_workgroups));
}
@@ -141,7 +141,7 @@ fn main() {
TEST_F(NumWorkgroupsFromUniformTest, StructOnlyMember_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(in : Builtins) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
@@ -166,7 +166,7 @@ fn main_inner(in : Builtins) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
main_inner(Builtins(tint_symbol_3.num_workgroups));
}
@@ -191,7 +191,7 @@ struct Builtins {
@builtin(workgroup_id) wgid : vec3<u32>,
};
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(in : Builtins) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
@@ -225,7 +225,7 @@ fn main_inner(in : Builtins) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(tint_symbol : tint_symbol_1) {
main_inner(Builtins(tint_symbol.gid, tint_symbol_3.num_workgroups, tint_symbol.wgid));
}
@@ -240,7 +240,7 @@ fn main(tint_symbol : tint_symbol_1) {
TEST_F(NumWorkgroupsFromUniformTest, StructMultipleMembers_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(in : Builtins) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
@@ -275,7 +275,7 @@ fn main_inner(in : Builtins) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(tint_symbol : tint_symbol_1) {
main_inner(Builtins(tint_symbol.gid, tint_symbol_3.num_workgroups, tint_symbol.wgid));
}
@@ -306,21 +306,21 @@ struct Builtins2 {
@builtin(workgroup_id) wgid : vec3<u32>,
};
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main1(in : Builtins1) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main2(in : Builtins2) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main3(@builtin(num_workgroups) num_wgs : vec3<u32>) {
let groups_x = num_wgs.x;
let groups_y = num_wgs.y;
@@ -351,7 +351,7 @@ fn main1_inner(in : Builtins1) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main1() {
main1_inner(Builtins1(tint_symbol_7.num_workgroups));
}
@@ -369,7 +369,7 @@ fn main2_inner(in : Builtins2) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main2(tint_symbol_2 : tint_symbol_3) {
main2_inner(Builtins2(tint_symbol_2.gid, tint_symbol_7.num_workgroups, tint_symbol_2.wgid));
}
@@ -380,7 +380,7 @@ fn main3_inner(num_wgs : vec3<u32>) {
let groups_z = num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main3() {
main3_inner(tint_symbol_7.num_workgroups);
}
@@ -400,7 +400,7 @@ struct Builtins {
@builtin(workgroup_id) wgid : vec3<u32>,
};
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(in : Builtins) {
}
)";
@@ -421,7 +421,7 @@ struct tint_symbol_1 {
fn main_inner(in : Builtins) {
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main(tint_symbol : tint_symbol_1) {
main_inner(Builtins(tint_symbol.gid, tint_symbol.wgid));
}
@@ -448,21 +448,21 @@ struct Builtins2 {
@builtin(workgroup_id) wgid : vec3<u32>,
};
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main1(in : Builtins1) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main2(in : Builtins2) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main3(@builtin(num_workgroups) num_wgs : vec3<u32>) {
let groups_x = num_wgs.x;
let groups_y = num_wgs.y;
@@ -493,7 +493,7 @@ fn main1_inner(in : Builtins1) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main1() {
main1_inner(Builtins1(tint_symbol_7.num_workgroups));
}
@@ -511,7 +511,7 @@ fn main2_inner(in : Builtins2) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main2(tint_symbol_2 : tint_symbol_3) {
main2_inner(Builtins2(tint_symbol_2.gid, tint_symbol_7.num_workgroups, tint_symbol_2.wgid));
}
@@ -522,7 +522,7 @@ fn main3_inner(num_wgs : vec3<u32>) {
let groups_z = num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main3() {
main3_inner(tint_symbol_7.num_workgroups);
}
@@ -572,7 +572,7 @@ struct S1 {
@group(1) @binding(3) var<storage, read> g9 : S0;
@group(3) @binding(2) var<storage, read_write> g10 : S0;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main1(in : Builtins1) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
@@ -580,14 +580,14 @@ fn main1(in : Builtins1) {
g8.m0 = 1u;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main2(in : Builtins2) {
let groups_x = in.num_wgs.x;
let groups_y = in.num_wgs.y;
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main3(@builtin(num_workgroups) num_wgs : vec3<u32>) {
let groups_x = num_wgs.x;
let groups_y = num_wgs.y;
@@ -647,7 +647,7 @@ fn main1_inner(in : Builtins1) {
g8.m0 = 1u;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main1() {
main1_inner(Builtins1(tint_symbol_7.num_workgroups));
}
@@ -665,7 +665,7 @@ fn main2_inner(in : Builtins2) {
let groups_z = in.num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main2(tint_symbol_2 : tint_symbol_3) {
main2_inner(Builtins2(tint_symbol_2.gid, tint_symbol_7.num_workgroups, tint_symbol_2.wgid));
}
@@ -676,7 +676,7 @@ fn main3_inner(num_wgs : vec3<u32>) {
let groups_z = num_wgs.z;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main3() {
main3_inner(tint_symbol_7.num_workgroups);
}

34
src/tint/transform/renamer_test.cc
View File

@@ -50,7 +50,7 @@ fn test(vert_idx : u32) -> u32 {
return vert_idx;
}
@stage(vertex)
@vertex
fn entry(@builtin(vertex_index) vert_idx : u32
) -> @builtin(position) vec4<f32> {
_ = test(vert_idx);
@@ -63,7 +63,7 @@ fn tint_symbol(tint_symbol_1 : u32) -> u32 {
return tint_symbol_1;
}
@stage(vertex)
@vertex
fn tint_symbol_2(@builtin(vertex_index) tint_symbol_1 : u32) -> @builtin(position) vec4<f32> {
_ = tint_symbol(tint_symbol_1);
return vec4<f32>();
@@ -87,7 +87,7 @@ fn tint_symbol_2(@builtin(vertex_index) tint_symbol_1 : u32) -> @builtin(positio
TEST_F(RenamerTest, PreserveSwizzles) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry() -> @builtin(position) vec4<f32> {
var v : vec4<f32>;
var rgba : f32;
@@ -97,7 +97,7 @@ fn entry() -> @builtin(position) vec4<f32> {
)";
auto* expect = R"(
@stage(vertex)
@vertex
fn tint_symbol() -> @builtin(position) vec4<f32> {
var tint_symbol_1 : vec4<f32>;
var tint_symbol_2 : f32;
@@ -124,7 +124,7 @@ fn tint_symbol() -> @builtin(position) vec4<f32> {
TEST_F(RenamerTest, PreserveBuiltins) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry() -> @builtin(position) vec4<f32> {
var blah : vec4<f32>;
return abs(blah);
@@ -132,7 +132,7 @@ fn entry() -> @builtin(position) vec4<f32> {
)";
auto* expect = R"(
@stage(vertex)
@vertex
fn tint_symbol() -> @builtin(position) vec4<f32> {
var tint_symbol_1 : vec4<f32>;
return abs(tint_symbol_1);
@@ -155,7 +155,7 @@ fn tint_symbol() -> @builtin(position) vec4<f32> {
TEST_F(RenamerTest, PreserveBuiltinTypes) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn entry() {
var a = modf(1.0).whole;
var b = modf(1.0).fract;
@@ -165,7 +165,7 @@ fn entry() {
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn tint_symbol() {
var tint_symbol_1 = modf(1.0).whole;
var tint_symbol_2 = modf(1.0).fract;
@@ -190,7 +190,7 @@ fn tint_symbol() {
TEST_F(RenamerTest, PreserveUnicode) {
auto src = R"(
@stage(fragment)
@fragment
fn frag_main() {
var )" + std::string(kUnicodeIdentifier) +
R"( : i32;
@@ -209,7 +209,7 @@ fn frag_main() {
TEST_F(RenamerTest, AttemptSymbolCollision) {
auto* src = R"(
@stage(vertex)
@vertex
fn entry() -> @builtin(position) vec4<f32> {
var tint_symbol : vec4<f32>;
var tint_symbol_2 : vec4<f32>;
@@ -219,7 +219,7 @@ fn entry() -> @builtin(position) vec4<f32> {
)";
auto* expect = R"(
@stage(vertex)
@vertex
fn tint_symbol() -> @builtin(position) vec4<f32> {
var tint_symbol_1 : vec4<f32>;
var tint_symbol_2 : vec4<f32>;
@@ -252,7 +252,7 @@ TEST_P(RenamerTestGlsl, Keywords) {
auto keyword = GetParam();
auto src = R"(
@stage(fragment)
@fragment
fn frag_main() {
var )" + keyword +
R"( : i32;
@@ -260,7 +260,7 @@ fn frag_main() {
)";
auto* expect = R"(
@stage(fragment)
@fragment
fn frag_main() {
var tint_symbol : i32;
}
@@ -278,7 +278,7 @@ TEST_P(RenamerTestHlsl, Keywords) {
auto keyword = GetParam();
auto src = R"(
@stage(fragment)
@fragment
fn frag_main() {
var )" + keyword +
R"( : i32;
@@ -286,7 +286,7 @@ fn frag_main() {
)";
auto* expect = R"(
@stage(fragment)
@fragment
fn frag_main() {
var tint_symbol : i32;
}
@@ -304,7 +304,7 @@ TEST_P(RenamerTestMsl, Keywords) {
auto keyword = GetParam();
auto src = R"(
@stage(fragment)
@fragment
fn frag_main() {
var )" + keyword +
R"( : i32;
@@ -312,7 +312,7 @@ fn frag_main() {
)";
auto* expect = R"(
@stage(fragment)
@fragment
fn frag_main() {
var tint_symbol : i32;
}

8
src/tint/transform/simplify_pointers_test.cc
View File

@@ -236,7 +236,7 @@ fn foo() -> i32 {
return 1;
}
@stage(fragment)
@fragment
fn main() {
var arr = array<f32, 4>();
for (let a = &arr[foo()]; ;) {
@@ -251,7 +251,7 @@ fn foo() -> i32 {
return 1;
}
@stage(fragment)
@fragment
fn main() {
var arr = array<f32, 4>();
let a_save = foo();
@@ -337,7 +337,7 @@ TEST_F(SimplifyPointersTest, ShadowPointer) {
auto* src = R"(
var<private> a : array<i32, 2>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
let x = &a;
var a : i32 = (*x)[0];
@@ -350,7 +350,7 @@ fn main() {
auto* expect = R"(
var<private> a : array<i32, 2>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
var a_1 : i32 = a[0];
{

70
src/tint/transform/single_entry_point_test.cc
View File

@@ -47,7 +47,7 @@ TEST_F(SingleEntryPointTest, Error_NoEntryPoints) {
TEST_F(SingleEntryPointTest, Error_InvalidEntryPoint) {
auto* src = R"(
@stage(vertex)
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@@ -68,7 +68,7 @@ TEST_F(SingleEntryPointTest, Error_NotAnEntryPoint) {
auto* src = R"(
fn foo() {}
@stage(fragment)
@fragment
fn main() {}
)";
@@ -85,7 +85,7 @@ fn main() {}
TEST_F(SingleEntryPointTest, SingleEntryPoint) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn main() {
}
)";
@@ -101,26 +101,26 @@ fn main() {
TEST_F(SingleEntryPointTest, MultipleEntryPoints) {
auto* src = R"(
@stage(vertex)
@vertex
fn vert_main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@stage(fragment)
@fragment
fn frag_main() {
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main2() {
}
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
}
)";
@@ -144,23 +144,23 @@ var<private> c : f32;
var<private> d : f32;
@stage(vertex)
@vertex
fn vert_main() -> @builtin(position) vec4<f32> {
a = 0.0;
return vec4<f32>();
}
@stage(fragment)
@fragment
fn frag_main() {
b = 0.0;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
c = 0.0;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main2() {
d = 0.0;
}
@@ -169,7 +169,7 @@ fn comp_main2() {
auto* expect = R"(
var<private> c : f32;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
c = 0.0;
}
@@ -194,23 +194,23 @@ let c : f32 = 1.0;
let d : f32 = 1.0;
@stage(vertex)
@vertex
fn vert_main() -> @builtin(position) vec4<f32> {
let local_a : f32 = a;
return vec4<f32>();
}
@stage(fragment)
@fragment
fn frag_main() {
let local_b : f32 = b;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
let local_c : f32 = c;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main2() {
let local_d : f32 = d;
}
@@ -219,7 +219,7 @@ fn comp_main2() {
auto* expect = R"(
let c : f32 = 1.0;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
let local_c : f32 = c;
}
@@ -238,7 +238,7 @@ TEST_F(SingleEntryPointTest, WorkgroupSizeLetPreserved) {
auto* src = R"(
let size : i32 = 1;
@stage(compute) @workgroup_size(size)
@compute @workgroup_size(size)
fn main() {
}
)";
@@ -261,27 +261,27 @@ TEST_F(SingleEntryPointTest, OverridableConstants) {
@id(0) override c3 : u32 = 1u;
@id(9999) override c4 : u32 = 1u;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
let local_d = c1;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main2() {
let local_d = c2;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main3() {
let local_d = c3;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main4() {
let local_d = c4;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main5() {
let local_d = 1u;
}
@@ -292,7 +292,7 @@ fn comp_main5() {
auto* expect = R"(
@id(1001) override c1 : u32 = 1u;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
let local_d = c1;
}
@@ -310,7 +310,7 @@ fn comp_main1() {
auto* expect = R"(
@id(1) override c2 : u32 = 1u;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main2() {
let local_d = c2;
}
@@ -326,7 +326,7 @@ fn comp_main2() {
auto* expect = R"(
@id(0) override c3 : u32 = 1u;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main3() {
let local_d = c3;
}
@@ -342,7 +342,7 @@ fn comp_main3() {
auto* expect = R"(
@id(9999) override c4 : u32 = 1u;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main4() {
let local_d = c4;
}
@@ -356,7 +356,7 @@ fn comp_main4() {
{
SingleEntryPoint::Config cfg("comp_main5");
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main5() {
let local_d = 1u;
}
@@ -389,12 +389,12 @@ fn outer2() {
inner_shared();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main2() {
outer2();
}
@@ -412,7 +412,7 @@ fn outer1() {
inner_shared();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}
@@ -463,12 +463,12 @@ fn outer2() {
outer2_var = 0.0;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main2() {
outer2();
}
@@ -495,7 +495,7 @@ fn outer1() {
outer1_var = 0.0;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}

82
src/tint/transform/unwind_discard_functions_test.cc
View File

@@ -102,7 +102,7 @@ fn f() {
let marker1 = 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
f();
let marker1 = 0;
@@ -122,7 +122,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
f();
if (tint_discard) {
@@ -156,7 +156,7 @@ fn f() -> S {
return s;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
f();
@@ -186,7 +186,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
f();
@@ -230,7 +230,7 @@ fn h() -> i32{
return 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
h();
@@ -275,7 +275,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
h();
@@ -311,7 +311,7 @@ fn h() {
let marker1 = 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
f();
@@ -348,7 +348,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
f();
@@ -381,7 +381,7 @@ fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
TEST_F(UnwindDiscardFunctionsTest, Call_DiscardFuncDeclaredBelow) {
auto* src = R"(
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
f();
let marker1 = 0;
@@ -400,7 +400,7 @@ fn tint_discard_func() {
var<private> tint_discard : bool = false;
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
f();
if (tint_discard) {
@@ -433,7 +433,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
if (f() == 42) {
let marker1 = 0;
@@ -456,7 +456,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let tint_symbol = f();
if (tint_discard) {
@@ -485,7 +485,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
if (true) {
let marker1 = 0;
@@ -512,7 +512,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
if (true) {
let marker1 = 0;
@@ -547,7 +547,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
var a = 0;
@@ -573,7 +573,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
var a = 0;
@@ -605,7 +605,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
for (f(); ; ) {
@@ -630,7 +630,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
var tint_symbol = f();
@@ -661,7 +661,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
for (let i = f(); ; ) {
@@ -686,7 +686,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
var tint_symbol = f();
@@ -717,7 +717,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
for (; f() == 42; ) {
@@ -742,7 +742,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
loop {
@@ -778,7 +778,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
for (; ; f()) {
@@ -809,7 +809,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
switch (f()) {
case 0: {
@@ -843,7 +843,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
var tint_symbol = f();
if (tint_discard) {
@@ -893,7 +893,7 @@ fn g() -> S {
return f();
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
g();
@@ -929,7 +929,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let marker1 = 0;
g();
@@ -956,7 +956,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
var a = f();
let marker1 = 0;
@@ -978,7 +978,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
var a = f();
if (tint_discard) {
@@ -1005,7 +1005,7 @@ fn f() -> i32 {
return 42;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
var a : i32;
a = f();
@@ -1028,7 +1028,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
var a : i32;
a = f();
@@ -1056,7 +1056,7 @@ fn f() -> i32 {
return 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
var b = array<i32, 10>();
b[f()] = 10;
@@ -1079,7 +1079,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
var b = array<i32, 10>();
let tint_symbol = f();
@@ -1115,7 +1115,7 @@ fn g() -> i32 {
return 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
var b = array<i32, 10>();
b[f()] = g();
@@ -1146,7 +1146,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
var b = array<i32, 10>();
let tint_symbol = g();
@@ -1194,7 +1194,7 @@ fn h() -> i32{
return 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
if ((f() + g() + h()) == 0) {
let marker1 = 0;
@@ -1233,7 +1233,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let tint_symbol = f();
if (tint_discard) {
@@ -1286,7 +1286,7 @@ fn h() -> i32{
return 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
if (f() == 1 && g() == 2 && h() == 3) {
let marker1 = 0;
@@ -1325,7 +1325,7 @@ fn tint_discard_func() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
let tint_symbol_2 = f();
if (tint_discard) {
@@ -1373,7 +1373,7 @@ fn f() {
let marker1 = 0;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
f();
let marker1 = 0;
@@ -1397,7 +1397,7 @@ fn tint_discard_func_1() {
discard;
}
@stage(fragment)
@fragment
fn main(@builtin(position) coord_in : vec4<f32>) -> @location(0) vec4<f32> {
f();
if (tint_discard_1) {

8
src/tint/transform/vectorize_scalar_matrix_constructors_test.cc
View File

@@ -53,7 +53,7 @@ TEST_P(VectorizeScalarMatrixConstructorsTest, SingleScalars) {
}
std::string src = R"(
@stage(fragment)
@fragment
fn main() {
let m = ${matrix}(42.0);
}
@@ -64,7 +64,7 @@ fn build_${matrix_no_type}(value : f32) -> ${matrix} {
return ${matrix}(${values});
}
@stage(fragment)
@fragment
fn main() {
let m = build_${matrix_no_type}(42.0);
}
@@ -107,7 +107,7 @@ TEST_P(VectorizeScalarMatrixConstructorsTest, MultipleScalars) {
}
std::string tmpl = R"(
@stage(fragment)
@fragment
fn main() {
let m = ${matrix}(${values});
}
@@ -137,7 +137,7 @@ TEST_P(VectorizeScalarMatrixConstructorsTest, NonScalarConstructors) {
}
std::string tmpl = R"(
@stage(fragment)
@fragment
fn main() {
let m = ${matrix}(${columns});
}

70
src/tint/transform/vertex_pulling_test.cc
View File

@@ -37,7 +37,7 @@ TEST_F(VertexPullingTest, Error_NoEntryPoint) {
TEST_F(VertexPullingTest, Error_InvalidEntryPoint) {
auto* src = R"(
@stage(vertex)
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@@ -57,7 +57,7 @@ fn main() -> @builtin(position) vec4<f32> {
TEST_F(VertexPullingTest, Error_EntryPointWrongStage) {
auto* src = R"(
@stage(fragment)
@fragment
fn main() {}
)";
@@ -75,7 +75,7 @@ fn main() {}
TEST_F(VertexPullingTest, Error_BadStride) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
return vec4<f32>(var_a, 0.0, 0.0, 1.0);
}
@@ -98,7 +98,7 @@ fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
TEST_F(VertexPullingTest, BasicModule) {
auto* src = R"(
@stage(vertex)
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@@ -109,7 +109,7 @@ struct TintVertexData {
tint_vertex_data : array<u32>,
}
@stage(vertex)
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@@ -127,7 +127,7 @@ fn main() -> @builtin(position) vec4<f32> {
TEST_F(VertexPullingTest, OneAttribute) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
return vec4<f32>(var_a, 0.0, 0.0, 1.0);
}
@@ -140,7 +140,7 @@ struct TintVertexData {
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var var_a : f32;
{
@@ -164,7 +164,7 @@ fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(posi
TEST_F(VertexPullingTest, OneInstancedAttribute) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
return vec4<f32>(var_a, 0.0, 0.0, 1.0);
}
@@ -177,7 +177,7 @@ struct TintVertexData {
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(instance_index) tint_pulling_instance_index : u32) -> @builtin(position) vec4<f32> {
var var_a : f32;
{
@@ -201,7 +201,7 @@ fn main(@builtin(instance_index) tint_pulling_instance_index : u32) -> @builtin(
TEST_F(VertexPullingTest, OneAttributeDifferentOutputSet) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : f32) -> @builtin(position) vec4<f32> {
return vec4<f32>(var_a, 0.0, 0.0, 1.0);
}
@@ -214,7 +214,7 @@ struct TintVertexData {
@binding(0) @group(5) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var var_a : f32;
{
@@ -243,7 +243,7 @@ struct Inputs {
@location(0) var_a : f32,
};
@stage(vertex)
@vertex
fn main(inputs : Inputs) -> @builtin(position) vec4<f32> {
return vec4<f32>(inputs.var_a, 0.0, 0.0, 1.0);
}
@@ -261,7 +261,7 @@ struct Inputs {
var_a : f32,
}
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var inputs : Inputs;
{
@@ -286,7 +286,7 @@ fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(posi
// We expect the transform to use an existing builtin variables if it finds them
TEST_F(VertexPullingTest, ExistingVertexIndexAndInstanceIndex) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : f32,
@location(1) var_b : f32,
@builtin(vertex_index) custom_vertex_index : u32,
@@ -305,7 +305,7 @@ struct TintVertexData {
@binding(1) @group(4) var<storage, read> tint_pulling_vertex_buffer_1 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) custom_vertex_index : u32, @builtin(instance_index) custom_instance_index : u32) -> @builtin(position) vec4<f32> {
var var_a : f32;
var var_b : f32;
@@ -350,7 +350,7 @@ struct Inputs {
@builtin(instance_index) custom_instance_index : u32,
};
@stage(vertex)
@vertex
fn main(inputs : Inputs) -> @builtin(position) vec4<f32> {
return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
}
@@ -383,7 +383,7 @@ struct Inputs {
custom_instance_index : u32,
}
@stage(vertex)
@vertex
fn main(tint_symbol_1 : tint_symbol) -> @builtin(position) vec4<f32> {
var inputs : Inputs;
inputs.custom_vertex_index = tint_symbol_1.custom_vertex_index;
@@ -422,7 +422,7 @@ fn main(tint_symbol_1 : tint_symbol) -> @builtin(position) vec4<f32> {
TEST_F(VertexPullingTest, ExistingVertexIndexAndInstanceIndex_Struct_OutOfOrder) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(inputs : Inputs) -> @builtin(position) vec4<f32> {
return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
}
@@ -451,7 +451,7 @@ struct tint_symbol {
custom_instance_index : u32,
}
@stage(vertex)
@vertex
fn main(tint_symbol_1 : tint_symbol) -> @builtin(position) vec4<f32> {
var inputs : Inputs;
inputs.custom_vertex_index = tint_symbol_1.custom_vertex_index;
@@ -511,7 +511,7 @@ struct Indices {
@builtin(instance_index) custom_instance_index : u32,
};
@stage(vertex)
@vertex
fn main(inputs : Inputs, indices : Indices) -> @builtin(position) vec4<f32> {
return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
}
@@ -540,7 +540,7 @@ struct Indices {
custom_instance_index : u32,
}
@stage(vertex)
@vertex
fn main(indices : Indices) -> @builtin(position) vec4<f32> {
var inputs : Inputs;
{
@@ -577,7 +577,7 @@ fn main(indices : Indices) -> @builtin(position) vec4<f32> {
TEST_F(VertexPullingTest, ExistingVertexIndexAndInstanceIndex_SeparateStruct_OutOfOrder) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(inputs : Inputs, indices : Indices) -> @builtin(position) vec4<f32> {
return vec4<f32>(inputs.var_a, inputs.var_b, 0.0, 1.0);
}
@@ -602,7 +602,7 @@ struct TintVertexData {
@binding(1) @group(4) var<storage, read> tint_pulling_vertex_buffer_1 : TintVertexData;
@stage(vertex)
@vertex
fn main(indices : Indices) -> @builtin(position) vec4<f32> {
var inputs : Inputs;
{
@@ -653,7 +653,7 @@ struct Indices {
TEST_F(VertexPullingTest, TwoAttributesSameBuffer) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : f32,
@location(1) var_b : vec4<f32>) -> @builtin(position) vec4<f32> {
return vec4<f32>();
@@ -667,7 +667,7 @@ struct TintVertexData {
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var var_a : f32;
var var_b : vec4<f32>;
@@ -695,7 +695,7 @@ fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(posi
TEST_F(VertexPullingTest, FloatVectorAttributes) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : vec2<f32>,
@location(1) var_b : vec3<f32>,
@location(2) var_c : vec4<f32>
@@ -715,7 +715,7 @@ struct TintVertexData {
@binding(2) @group(4) var<storage, read> tint_pulling_vertex_buffer_2 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var var_a : vec2<f32>;
var var_b : vec3<f32>;
@@ -749,7 +749,7 @@ fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(posi
TEST_F(VertexPullingTest, AttemptSymbolCollision) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(@location(0) var_a : f32,
@location(1) var_b : vec4<f32>) -> @builtin(position) vec4<f32> {
var tint_pulling_vertex_index : i32;
@@ -767,7 +767,7 @@ struct TintVertexData {
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0_1 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index_1 : u32) -> @builtin(position) vec4<f32> {
var var_a : f32;
var var_b : vec4<f32>;
@@ -799,7 +799,7 @@ fn main(@builtin(vertex_index) tint_pulling_vertex_index_1 : u32) -> @builtin(po
TEST_F(VertexPullingTest, FormatsAligned) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(
@location(0) uint8x2 : vec2<u32>,
@location(1) uint8x4 : vec4<u32>,
@@ -843,7 +843,7 @@ struct TintVertexData {
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var uint8x2 : vec2<u32>;
var uint8x4 : vec4<u32>;
@@ -944,7 +944,7 @@ fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(posi
TEST_F(VertexPullingTest, FormatsStrideUnaligned) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(
@location(0) uint8x2 : vec2<u32>,
@location(1) uint8x4 : vec4<u32>,
@@ -989,7 +989,7 @@ struct TintVertexData {
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var uint8x2 : vec2<u32>;
var uint8x4 : vec4<u32>;
@@ -1090,7 +1090,7 @@ fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(posi
TEST_F(VertexPullingTest, FormatsWithVectorsResized) {
auto* src = R"(
@stage(vertex)
@vertex
fn main(
@location(0) uint8x2 : vec3<u32>,
@location(1) uint8x4 : vec2<u32>,
@@ -1134,7 +1134,7 @@ struct TintVertexData {
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@stage(vertex)
@vertex
fn main(@builtin(vertex_index) tint_pulling_vertex_index : u32) -> @builtin(position) vec4<f32> {
var uint8x2 : vec3<u32>;
var uint8x4 : vec2<u32>;

120
src/tint/transform/zero_init_workgroup_memory_test.cc
View File

@@ -81,7 +81,7 @@ fn unreferenced() {
b = c;
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
)";
@@ -94,7 +94,7 @@ fn f() {
TEST_F(ZeroInitWorkgroupMemoryTest, UnreferencedWorkgroupVars_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
}
@@ -119,7 +119,7 @@ TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndex) {
auto* src = R"(
var<workgroup> v : i32;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
_ = v; // Initialization should be inserted above this statement
}
@@ -127,7 +127,7 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
auto* expect = R"(
var<workgroup> v : i32;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
{
v = i32();
@@ -144,7 +144,7 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndex_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
_ = v; // Initialization should be inserted above this statement
}
@@ -152,7 +152,7 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
var<workgroup> v : i32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
{
v = i32();
@@ -177,7 +177,7 @@ struct Params {
@builtin(local_invocation_index) local_idx : u32,
};
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(params : Params) {
_ = v; // Initialization should be inserted above this statement
}
@@ -190,7 +190,7 @@ struct Params {
local_idx : u32,
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(params : Params) {
{
v = i32();
@@ -207,7 +207,7 @@ fn f(params : Params) {
TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndexInStruct_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(params : Params) {
_ = v; // Initialization should be inserted above this statement
}
@@ -219,7 +219,7 @@ struct Params {
var<workgroup> v : i32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(params : Params) {
{
v = i32();
@@ -245,7 +245,7 @@ TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_InjectedLocalIndex) {
auto* src = R"(
var<workgroup> v : i32;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
_ = v; // Initialization should be inserted above this statement
}
@@ -253,7 +253,7 @@ fn f() {
auto* expect = R"(
var<workgroup> v : i32;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
{
v = i32();
@@ -270,7 +270,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_InjectedLocalIndex_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
_ = v; // Initialization should be inserted above this statement
}
@@ -278,7 +278,7 @@ fn f() {
var<workgroup> v : i32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
{
v = i32();
@@ -308,7 +308,7 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
_ = a; // Initialization should be inserted above this statement
_ = b;
@@ -327,7 +327,7 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
{
a = i32();
@@ -360,7 +360,7 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_ExistingLocalIndex_Size1_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
_ = a; // Initialization should be inserted above this statement
_ = b;
@@ -379,7 +379,7 @@ struct S {
};
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
{
a = i32();
@@ -434,7 +434,7 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(2, 3)
@compute @workgroup_size(2, 3)
fn f(@builtin(local_invocation_index) local_idx : u32) {
_ = a; // Initialization should be inserted above this statement
_ = b;
@@ -453,7 +453,7 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(2, 3)
@compute @workgroup_size(2, 3)
fn f(@builtin(local_invocation_index) local_idx : u32) {
if ((local_idx < 1u)) {
a = i32();
@@ -499,7 +499,7 @@ var<workgroup> c : array<S, 32>;
@id(1) override X : i32;
@stage(compute) @workgroup_size(2, 3, X)
@compute @workgroup_size(2, 3, X)
fn f(@builtin(local_invocation_index) local_idx : u32) {
_ = a; // Initialization should be inserted above this statement
_ = b;
@@ -521,7 +521,7 @@ var<workgroup> c : array<S, 32>;
@id(1) override X : i32;
@stage(compute) @workgroup_size(2, 3, X)
@compute @workgroup_size(2, 3, X)
fn f(@builtin(local_invocation_index) local_idx : u32) {
for(var idx : u32 = local_idx; (idx < 1u); idx = (idx + (u32(X) * 6u))) {
a = i32();
@@ -568,7 +568,7 @@ var<workgroup> c : array<S, 32>;
@id(1) override X : u32;
@stage(compute) @workgroup_size(5u, X, 10u)
@compute @workgroup_size(5u, X, 10u)
fn f(@builtin(local_invocation_index) local_idx : u32) {
_ = a; // Initialization should be inserted above this statement
_ = b;
@@ -591,7 +591,7 @@ var<workgroup> c : array<S, 32>;
@id(1) override X : u32;
@stage(compute) @workgroup_size(5u, X, 10u)
@compute @workgroup_size(5u, X, 10u)
fn f(@builtin(local_invocation_index) local_idx : u32) {
for(var idx : u32 = local_idx; (idx < 1u); idx = (idx + (X * 50u))) {
a = i32();
@@ -654,7 +654,7 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
_ = a; // Initialization should be inserted above this statement
_ = b;
@@ -673,7 +673,7 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index : u32) {
{
a = i32();
@@ -706,7 +706,7 @@ fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(loc
TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_InjectedLocalIndex_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
_ = a; // Initialization should be inserted above this statement
_ = b;
@@ -725,7 +725,7 @@ struct S {
};
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index : u32) {
{
a = i32();
@@ -780,18 +780,18 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f1() {
_ = a; // Initialization should be inserted above this statement
_ = c;
}
@stage(compute) @workgroup_size(1, 2, 3)
@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
_ = b; // Initialization should be inserted above this statement
}
@stage(compute) @workgroup_size(4, 5, 6)
@compute @workgroup_size(4, 5, 6)
fn f3() {
_ = c; // Initialization should be inserted above this statement
_ = a;
@@ -809,7 +809,7 @@ var<workgroup> b : S;
var<workgroup> c : array<S, 32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
{
a = i32();
@@ -828,7 +828,7 @@ fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
_ = c;
}
@stage(compute) @workgroup_size(1, 2, 3)
@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index_1 : u32) {
if ((local_invocation_index_1 < 1u)) {
b.x = i32();
@@ -841,7 +841,7 @@ fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(lo
_ = b;
}
@stage(compute) @workgroup_size(4, 5, 6)
@compute @workgroup_size(4, 5, 6)
fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
if ((local_invocation_index_2 < 1u)) {
a = i32();
@@ -868,18 +868,18 @@ fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_MultipleEntryPoints_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f1() {
_ = a; // Initialization should be inserted above this statement
_ = c;
}
@stage(compute) @workgroup_size(1, 2, 3)
@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
_ = b; // Initialization should be inserted above this statement
}
@stage(compute) @workgroup_size(4, 5, 6)
@compute @workgroup_size(4, 5, 6)
fn f3() {
_ = c; // Initialization should be inserted above this statement
_ = a;
@@ -897,7 +897,7 @@ struct S {
};
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
{
a = i32();
@@ -916,7 +916,7 @@ fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
_ = c;
}
@stage(compute) @workgroup_size(1, 2, 3)
@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index_1 : u32) {
if ((local_invocation_index_1 < 1u)) {
b.x = i32();
@@ -929,7 +929,7 @@ fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(lo
_ = b;
}
@stage(compute) @workgroup_size(4, 5, 6)
@compute @workgroup_size(4, 5, 6)
fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
if ((local_invocation_index_2 < 1u)) {
a = i32();
@@ -977,7 +977,7 @@ fn call_use_v() {
use_v();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
call_use_v(); // Initialization should be inserted above this statement
}
@@ -993,7 +993,7 @@ fn call_use_v() {
use_v();
}
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
{
v = i32();
@@ -1010,7 +1010,7 @@ fn f(@builtin(local_invocation_index) local_idx : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, TransitiveUsage_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
call_use_v(); // Initialization should be inserted above this statement
}
@@ -1026,7 +1026,7 @@ fn use_v() {
var<workgroup> v : i32;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
{
v = i32();
@@ -1056,7 +1056,7 @@ TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupAtomics) {
var<workgroup> i : atomic<i32>;
var<workgroup> u : atomic<u32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
atomicLoad(&(i)); // Initialization should be inserted above this statement
atomicLoad(&(u));
@@ -1067,7 +1067,7 @@ var<workgroup> i : atomic<i32>;
var<workgroup> u : atomic<u32>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
{
atomicStore(&(i), i32());
@@ -1086,7 +1086,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupAtomics_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
atomicLoad(&(i)); // Initialization should be inserted above this statement
atomicLoad(&(u));
@@ -1096,7 +1096,7 @@ var<workgroup> i : atomic<i32>;
var<workgroup> u : atomic<u32>;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
{
atomicStore(&(i), i32());
@@ -1129,7 +1129,7 @@ struct S {
var<workgroup> w : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
_ = w.a; // Initialization should be inserted above this statement
}
@@ -1145,7 +1145,7 @@ struct S {
var<workgroup> w : S;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
{
w.a = i32();
@@ -1166,7 +1166,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupStructOfAtomics_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
_ = w.a; // Initialization should be inserted above this statement
}
@@ -1182,7 +1182,7 @@ struct S {
};
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
{
w.a = i32();
@@ -1215,7 +1215,7 @@ TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfAtomics) {
auto* src = R"(
var<workgroup> w : array<atomic<u32>, 4>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
atomicLoad(&w[0]); // Initialization should be inserted above this statement
}
@@ -1223,7 +1223,7 @@ fn f() {
auto* expect = R"(
var<workgroup> w : array<atomic<u32>, 4>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
let i : u32 = idx;
@@ -1241,7 +1241,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfAtomics_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
atomicLoad(&w[0]); // Initialization should be inserted above this statement
}
@@ -1249,7 +1249,7 @@ fn f() {
var<workgroup> w : array<atomic<u32>, 4>;
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
let i : u32 = idx;
@@ -1279,7 +1279,7 @@ struct S {
var<workgroup> w : array<S, 4>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
_ = w[0].a; // Initialization should be inserted above this statement
}
@@ -1295,7 +1295,7 @@ struct S {
var<workgroup> w : array<S, 4>;
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
let i_1 : u32 = idx;
@@ -1317,7 +1317,7 @@ fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfStructOfAtomics_OutOfOrder) {
auto* src = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f() {
_ = w[0].a; // Initialization should be inserted above this statement
}
@@ -1333,7 +1333,7 @@ struct S {
};
)";
auto* expect = R"(
@stage(compute) @workgroup_size(1)
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
let i_1 : u32 = idx;