mirror of
https://github.com/encounter/dawn-cmake.git
synced 2025-06-05 14:13:39 +00:00
4cc4315d6c
Change ast::Array to use an ast::Expression for its `size` field. The WGSL frontend now parses the array size as an `primary_expression`, and the Resolver is responsible for validating the expression is a signed or unsigned integer, and either a literal or a non-overridable module-scope constant. The Resolver evaluates the constant value of the size expression, and so the resolved sem::Array type still has a constant size as before. Fixed: tint:1068 Fixed: tint:1117 Change-Id: Icfa141482ea1e47ea8c21a25e9eb48221f176e9a Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63061 Auto-Submit: James Price <jrprice@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Commit-Queue: James Price <jrprice@google.com> Reviewed-by: Ben Clayton <bclayton@google.com>
1418 lines
54 KiB
C++
1418 lines
54 KiB
C++
// Copyright 2021 The Tint Authors.
|
|
//
|
|
// Licensed under the Apache License, Version 2.0 (the "License");
|
|
// you may not use this file except in compliance with the License.
|
|
// You may obtain a copy of the License at
|
|
//
|
|
// http://www.apache.org/licenses/LICENSE-2.0
|
|
//
|
|
// Unless required by applicable law or agreed to in writing, software
|
|
// distributed under the License is distributed on an "AS IS" BASIS,
|
|
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
// See the License for the specific language governing permissions and
|
|
// limitations under the License.
|
|
|
|
#include "src/transform/decompose_memory_access.h"
|
|
|
|
#include "src/transform/test_helper.h"
|
|
|
|
namespace tint {
|
|
namespace transform {
|
|
namespace {
|
|
|
|
using DecomposeMemoryAccessTest = TransformTest;
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, SB_BasicLoad) {
|
|
auto* src = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var a : i32 = sb.a;
|
|
var b : u32 = sb.b;
|
|
var c : f32 = sb.c;
|
|
var d : vec2<i32> = sb.d;
|
|
var e : vec2<u32> = sb.e;
|
|
var f : vec2<f32> = sb.f;
|
|
var g : vec3<i32> = sb.g;
|
|
var h : vec3<u32> = sb.h;
|
|
var i : vec3<f32> = sb.i;
|
|
var j : vec4<i32> = sb.j;
|
|
var k : vec4<u32> = sb.k;
|
|
var l : vec4<f32> = sb.l;
|
|
var m : mat2x2<f32> = sb.m;
|
|
var n : mat2x3<f32> = sb.n;
|
|
var o : mat2x4<f32> = sb.o;
|
|
var p : mat3x2<f32> = sb.p;
|
|
var q : mat3x3<f32> = sb.q;
|
|
var r : mat3x4<f32> = sb.r;
|
|
var s : mat4x2<f32> = sb.s;
|
|
var t : mat4x3<f32> = sb.t;
|
|
var u : mat4x4<f32> = sb.u;
|
|
var v : array<vec3<f32>, 2> = sb.v;
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
|
|
|
|
[[internal(intrinsic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
|
|
|
|
[[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
|
|
|
|
[[internal(intrinsic_load_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
|
|
|
|
[[internal(intrinsic_load_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
|
|
|
|
[[internal(intrinsic_load_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
|
|
|
|
[[internal(intrinsic_load_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
|
|
|
|
[[internal(intrinsic_load_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
|
|
|
|
[[internal(intrinsic_load_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
|
|
|
|
[[internal(intrinsic_load_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
|
|
|
|
[[internal(intrinsic_load_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
|
|
|
|
[[internal(intrinsic_load_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
|
|
|
|
fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
|
|
return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
|
|
}
|
|
|
|
fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
|
|
return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
|
|
return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
|
|
return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
|
|
return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
|
|
}
|
|
|
|
fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
|
|
return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
|
|
}
|
|
|
|
fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
|
|
return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
|
|
}
|
|
|
|
fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
|
|
return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
|
|
}
|
|
|
|
fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
|
|
return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
|
|
}
|
|
|
|
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
|
|
var arr : array<vec3<f32>, 2u>;
|
|
for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
|
|
arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
|
|
}
|
|
return arr;
|
|
}
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var a : i32 = tint_symbol(sb, 0u);
|
|
var b : u32 = tint_symbol_1(sb, 4u);
|
|
var c : f32 = tint_symbol_2(sb, 8u);
|
|
var d : vec2<i32> = tint_symbol_3(sb, 16u);
|
|
var e : vec2<u32> = tint_symbol_4(sb, 24u);
|
|
var f : vec2<f32> = tint_symbol_5(sb, 32u);
|
|
var g : vec3<i32> = tint_symbol_6(sb, 48u);
|
|
var h : vec3<u32> = tint_symbol_7(sb, 64u);
|
|
var i : vec3<f32> = tint_symbol_8(sb, 80u);
|
|
var j : vec4<i32> = tint_symbol_9(sb, 96u);
|
|
var k : vec4<u32> = tint_symbol_10(sb, 112u);
|
|
var l : vec4<f32> = tint_symbol_11(sb, 128u);
|
|
var m : mat2x2<f32> = tint_symbol_12(sb, 144u);
|
|
var n : mat2x3<f32> = tint_symbol_13(sb, 160u);
|
|
var o : mat2x4<f32> = tint_symbol_14(sb, 192u);
|
|
var p : mat3x2<f32> = tint_symbol_15(sb, 224u);
|
|
var q : mat3x3<f32> = tint_symbol_16(sb, 256u);
|
|
var r : mat3x4<f32> = tint_symbol_17(sb, 304u);
|
|
var s : mat4x2<f32> = tint_symbol_18(sb, 352u);
|
|
var t : mat4x3<f32> = tint_symbol_19(sb, 384u);
|
|
var u : mat4x4<f32> = tint_symbol_20(sb, 448u);
|
|
var v : array<vec3<f32>, 2> = tint_symbol_21(sb, 512u);
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, UB_BasicLoad) {
|
|
auto* src = R"(
|
|
[[block]]
|
|
struct UB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<uniform> ub : UB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var a : i32 = ub.a;
|
|
var b : u32 = ub.b;
|
|
var c : f32 = ub.c;
|
|
var d : vec2<i32> = ub.d;
|
|
var e : vec2<u32> = ub.e;
|
|
var f : vec2<f32> = ub.f;
|
|
var g : vec3<i32> = ub.g;
|
|
var h : vec3<u32> = ub.h;
|
|
var i : vec3<f32> = ub.i;
|
|
var j : vec4<i32> = ub.j;
|
|
var k : vec4<u32> = ub.k;
|
|
var l : vec4<f32> = ub.l;
|
|
var m : mat2x2<f32> = ub.m;
|
|
var n : mat2x3<f32> = ub.n;
|
|
var o : mat2x4<f32> = ub.o;
|
|
var p : mat3x2<f32> = ub.p;
|
|
var q : mat3x3<f32> = ub.q;
|
|
var r : mat3x4<f32> = ub.r;
|
|
var s : mat4x2<f32> = ub.s;
|
|
var t : mat4x3<f32> = ub.t;
|
|
var u : mat4x4<f32> = ub.u;
|
|
var v : array<vec3<f32>, 2> = ub.v;
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
[[block]]
|
|
struct UB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<uniform> ub : UB;
|
|
|
|
[[internal(intrinsic_load_uniform_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> i32
|
|
|
|
[[internal(intrinsic_load_uniform_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> u32
|
|
|
|
[[internal(intrinsic_load_uniform_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> f32
|
|
|
|
[[internal(intrinsic_load_uniform_vec2_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec2<i32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec2_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec2<u32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec2_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec2<f32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec3_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec3<i32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec3_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec3<u32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec3_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec3<f32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec4_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec4<i32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec4_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec4<u32>
|
|
|
|
[[internal(intrinsic_load_uniform_vec4_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> vec4<f32>
|
|
|
|
fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat2x2<f32> {
|
|
return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
|
|
}
|
|
|
|
fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat2x3<f32> {
|
|
return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat2x4<f32> {
|
|
return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat3x2<f32> {
|
|
return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat3x3<f32> {
|
|
return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
|
|
}
|
|
|
|
fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat3x4<f32> {
|
|
return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
|
|
}
|
|
|
|
fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat4x2<f32> {
|
|
return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
|
|
}
|
|
|
|
fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat4x3<f32> {
|
|
return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
|
|
}
|
|
|
|
fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> mat4x4<f32> {
|
|
return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
|
|
}
|
|
|
|
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> array<vec3<f32>, 2u> {
|
|
var arr : array<vec3<f32>, 2u>;
|
|
for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
|
|
arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
|
|
}
|
|
return arr;
|
|
}
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var a : i32 = tint_symbol(ub, 0u);
|
|
var b : u32 = tint_symbol_1(ub, 4u);
|
|
var c : f32 = tint_symbol_2(ub, 8u);
|
|
var d : vec2<i32> = tint_symbol_3(ub, 16u);
|
|
var e : vec2<u32> = tint_symbol_4(ub, 24u);
|
|
var f : vec2<f32> = tint_symbol_5(ub, 32u);
|
|
var g : vec3<i32> = tint_symbol_6(ub, 48u);
|
|
var h : vec3<u32> = tint_symbol_7(ub, 64u);
|
|
var i : vec3<f32> = tint_symbol_8(ub, 80u);
|
|
var j : vec4<i32> = tint_symbol_9(ub, 96u);
|
|
var k : vec4<u32> = tint_symbol_10(ub, 112u);
|
|
var l : vec4<f32> = tint_symbol_11(ub, 128u);
|
|
var m : mat2x2<f32> = tint_symbol_12(ub, 144u);
|
|
var n : mat2x3<f32> = tint_symbol_13(ub, 160u);
|
|
var o : mat2x4<f32> = tint_symbol_14(ub, 192u);
|
|
var p : mat3x2<f32> = tint_symbol_15(ub, 224u);
|
|
var q : mat3x3<f32> = tint_symbol_16(ub, 256u);
|
|
var r : mat3x4<f32> = tint_symbol_17(ub, 304u);
|
|
var s : mat4x2<f32> = tint_symbol_18(ub, 352u);
|
|
var t : mat4x3<f32> = tint_symbol_19(ub, 384u);
|
|
var u : mat4x4<f32> = tint_symbol_20(ub, 448u);
|
|
var v : array<vec3<f32>, 2> = tint_symbol_21(ub, 512u);
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, SB_BasicStore) {
|
|
auto* src = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
sb.a = i32();
|
|
sb.b = u32();
|
|
sb.c = f32();
|
|
sb.d = vec2<i32>();
|
|
sb.e = vec2<u32>();
|
|
sb.f = vec2<f32>();
|
|
sb.g = vec3<i32>();
|
|
sb.h = vec3<u32>();
|
|
sb.i = vec3<f32>();
|
|
sb.j = vec4<i32>();
|
|
sb.k = vec4<u32>();
|
|
sb.l = vec4<f32>();
|
|
sb.m = mat2x2<f32>();
|
|
sb.n = mat2x3<f32>();
|
|
sb.o = mat2x4<f32>();
|
|
sb.p = mat3x2<f32>();
|
|
sb.q = mat3x3<f32>();
|
|
sb.r = mat3x4<f32>();
|
|
sb.s = mat4x2<f32>();
|
|
sb.t = mat4x3<f32>();
|
|
sb.u = mat4x4<f32>();
|
|
sb.v = array<vec3<f32>, 2>();
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
|
|
|
|
[[internal(intrinsic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
|
|
|
|
[[internal(intrinsic_store_storage_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
|
|
|
|
[[internal(intrinsic_store_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
|
|
|
|
fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
|
|
tint_symbol_5(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_5(buffer, (offset + 8u), value[1u]);
|
|
}
|
|
|
|
fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
|
|
tint_symbol_8(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_8(buffer, (offset + 16u), value[1u]);
|
|
}
|
|
|
|
fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
|
|
tint_symbol_11(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_11(buffer, (offset + 16u), value[1u]);
|
|
}
|
|
|
|
fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
|
|
tint_symbol_5(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_5(buffer, (offset + 8u), value[1u]);
|
|
tint_symbol_5(buffer, (offset + 16u), value[2u]);
|
|
}
|
|
|
|
fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
|
|
tint_symbol_8(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_8(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_8(buffer, (offset + 32u), value[2u]);
|
|
}
|
|
|
|
fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
|
|
tint_symbol_11(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_11(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_11(buffer, (offset + 32u), value[2u]);
|
|
}
|
|
|
|
fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
|
|
tint_symbol_5(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_5(buffer, (offset + 8u), value[1u]);
|
|
tint_symbol_5(buffer, (offset + 16u), value[2u]);
|
|
tint_symbol_5(buffer, (offset + 24u), value[3u]);
|
|
}
|
|
|
|
fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
|
|
tint_symbol_8(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_8(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_8(buffer, (offset + 32u), value[2u]);
|
|
tint_symbol_8(buffer, (offset + 48u), value[3u]);
|
|
}
|
|
|
|
fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
|
|
tint_symbol_11(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_11(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_11(buffer, (offset + 32u), value[2u]);
|
|
tint_symbol_11(buffer, (offset + 48u), value[3u]);
|
|
}
|
|
|
|
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
|
|
var array = value;
|
|
for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
|
|
tint_symbol_8(buffer, (offset + (i_1 * 16u)), array[i_1]);
|
|
}
|
|
}
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
tint_symbol(sb, 0u, i32());
|
|
tint_symbol_1(sb, 4u, u32());
|
|
tint_symbol_2(sb, 8u, f32());
|
|
tint_symbol_3(sb, 16u, vec2<i32>());
|
|
tint_symbol_4(sb, 24u, vec2<u32>());
|
|
tint_symbol_5(sb, 32u, vec2<f32>());
|
|
tint_symbol_6(sb, 48u, vec3<i32>());
|
|
tint_symbol_7(sb, 64u, vec3<u32>());
|
|
tint_symbol_8(sb, 80u, vec3<f32>());
|
|
tint_symbol_9(sb, 96u, vec4<i32>());
|
|
tint_symbol_10(sb, 112u, vec4<u32>());
|
|
tint_symbol_11(sb, 128u, vec4<f32>());
|
|
tint_symbol_12(sb, 144u, mat2x2<f32>());
|
|
tint_symbol_13(sb, 160u, mat2x3<f32>());
|
|
tint_symbol_14(sb, 192u, mat2x4<f32>());
|
|
tint_symbol_15(sb, 224u, mat3x2<f32>());
|
|
tint_symbol_16(sb, 256u, mat3x3<f32>());
|
|
tint_symbol_17(sb, 304u, mat3x4<f32>());
|
|
tint_symbol_18(sb, 352u, mat4x2<f32>());
|
|
tint_symbol_19(sb, 384u, mat4x3<f32>());
|
|
tint_symbol_20(sb, 448u, mat4x4<f32>());
|
|
tint_symbol_21(sb, 512u, array<vec3<f32>, 2>());
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, LoadStructure) {
|
|
auto* src = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var x : SB = sb;
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
|
|
|
|
[[internal(intrinsic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
|
|
|
|
[[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
|
|
|
|
[[internal(intrinsic_load_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
|
|
|
|
[[internal(intrinsic_load_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
|
|
|
|
[[internal(intrinsic_load_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
|
|
|
|
[[internal(intrinsic_load_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
|
|
|
|
[[internal(intrinsic_load_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
|
|
|
|
[[internal(intrinsic_load_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
|
|
|
|
[[internal(intrinsic_load_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
|
|
|
|
[[internal(intrinsic_load_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
|
|
|
|
[[internal(intrinsic_load_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
|
|
|
|
fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
|
|
return mat2x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)));
|
|
}
|
|
|
|
fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
|
|
return mat2x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
|
|
return mat2x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
|
|
return mat3x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)));
|
|
}
|
|
|
|
fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
|
|
return mat3x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)));
|
|
}
|
|
|
|
fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
|
|
return mat3x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)));
|
|
}
|
|
|
|
fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
|
|
return mat4x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)), tint_symbol_6(buffer, (offset + 24u)));
|
|
}
|
|
|
|
fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
|
|
return mat4x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)), tint_symbol_9(buffer, (offset + 48u)));
|
|
}
|
|
|
|
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
|
|
return mat4x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)), tint_symbol_12(buffer, (offset + 48u)));
|
|
}
|
|
|
|
fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
|
|
var arr : array<vec3<f32>, 2u>;
|
|
for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
|
|
arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
|
|
}
|
|
return arr;
|
|
}
|
|
|
|
fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> SB {
|
|
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)]]
|
|
fn main() {
|
|
var x : SB = tint_symbol(sb, 0u);
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, StoreStructure) {
|
|
auto* src = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
sb = SB();
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
[[block]]
|
|
struct SB {
|
|
a : i32;
|
|
b : u32;
|
|
c : f32;
|
|
d : vec2<i32>;
|
|
e : vec2<u32>;
|
|
f : vec2<f32>;
|
|
g : vec3<i32>;
|
|
h : vec3<u32>;
|
|
i : vec3<f32>;
|
|
j : vec4<i32>;
|
|
k : vec4<u32>;
|
|
l : vec4<f32>;
|
|
m : mat2x2<f32>;
|
|
n : mat2x3<f32>;
|
|
o : mat2x4<f32>;
|
|
p : mat3x2<f32>;
|
|
q : mat3x3<f32>;
|
|
r : mat3x4<f32>;
|
|
s : mat4x2<f32>;
|
|
t : mat4x3<f32>;
|
|
u : mat4x4<f32>;
|
|
v : array<vec3<f32>, 2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
|
|
|
|
[[internal(intrinsic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
|
|
|
|
[[internal(intrinsic_store_storage_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
|
|
|
|
[[internal(intrinsic_store_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
|
|
|
|
[[internal(intrinsic_store_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
|
|
|
|
fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
|
|
tint_symbol_6(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_6(buffer, (offset + 8u), value[1u]);
|
|
}
|
|
|
|
fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
|
|
tint_symbol_9(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_9(buffer, (offset + 16u), value[1u]);
|
|
}
|
|
|
|
fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
|
|
tint_symbol_12(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_12(buffer, (offset + 16u), value[1u]);
|
|
}
|
|
|
|
fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
|
|
tint_symbol_6(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_6(buffer, (offset + 8u), value[1u]);
|
|
tint_symbol_6(buffer, (offset + 16u), value[2u]);
|
|
}
|
|
|
|
fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
|
|
tint_symbol_9(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_9(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_9(buffer, (offset + 32u), value[2u]);
|
|
}
|
|
|
|
fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
|
|
tint_symbol_12(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_12(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_12(buffer, (offset + 32u), value[2u]);
|
|
}
|
|
|
|
fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
|
|
tint_symbol_6(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_6(buffer, (offset + 8u), value[1u]);
|
|
tint_symbol_6(buffer, (offset + 16u), value[2u]);
|
|
tint_symbol_6(buffer, (offset + 24u), value[3u]);
|
|
}
|
|
|
|
fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
|
|
tint_symbol_9(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_9(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_9(buffer, (offset + 32u), value[2u]);
|
|
tint_symbol_9(buffer, (offset + 48u), value[3u]);
|
|
}
|
|
|
|
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
|
|
tint_symbol_12(buffer, (offset + 0u), value[0u]);
|
|
tint_symbol_12(buffer, (offset + 16u), value[1u]);
|
|
tint_symbol_12(buffer, (offset + 32u), value[2u]);
|
|
tint_symbol_12(buffer, (offset + 48u), value[3u]);
|
|
}
|
|
|
|
fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
|
|
var array = value;
|
|
for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
|
|
tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1]);
|
|
}
|
|
}
|
|
|
|
fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : SB) {
|
|
tint_symbol_1(buffer, (offset + 0u), value.a);
|
|
tint_symbol_2(buffer, (offset + 4u), value.b);
|
|
tint_symbol_3(buffer, (offset + 8u), value.c);
|
|
tint_symbol_4(buffer, (offset + 16u), value.d);
|
|
tint_symbol_5(buffer, (offset + 24u), value.e);
|
|
tint_symbol_6(buffer, (offset + 32u), value.f);
|
|
tint_symbol_7(buffer, (offset + 48u), value.g);
|
|
tint_symbol_8(buffer, (offset + 64u), value.h);
|
|
tint_symbol_9(buffer, (offset + 80u), value.i);
|
|
tint_symbol_10(buffer, (offset + 96u), value.j);
|
|
tint_symbol_11(buffer, (offset + 112u), value.k);
|
|
tint_symbol_12(buffer, (offset + 128u), value.l);
|
|
tint_symbol_13(buffer, (offset + 144u), value.m);
|
|
tint_symbol_14(buffer, (offset + 160u), value.n);
|
|
tint_symbol_15(buffer, (offset + 192u), value.o);
|
|
tint_symbol_16(buffer, (offset + 224u), value.p);
|
|
tint_symbol_17(buffer, (offset + 256u), value.q);
|
|
tint_symbol_18(buffer, (offset + 304u), value.r);
|
|
tint_symbol_19(buffer, (offset + 352u), value.s);
|
|
tint_symbol_20(buffer, (offset + 384u), value.t);
|
|
tint_symbol_21(buffer, (offset + 448u), value.u);
|
|
tint_symbol_22(buffer, (offset + 512u), value.v);
|
|
}
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
tint_symbol(sb, 0u, SB());
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain) {
|
|
auto* src = R"(
|
|
struct S1 {
|
|
a : i32;
|
|
b : vec3<f32>;
|
|
c : i32;
|
|
};
|
|
|
|
struct S2 {
|
|
a : i32;
|
|
b : [[stride(32)]] array<S1, 3>;
|
|
c : i32;
|
|
};
|
|
|
|
[[block]]
|
|
struct SB {
|
|
[[size(128)]]
|
|
a : i32;
|
|
b : [[stride(256)]] array<S2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var x : f32 = sb.b[4].b[1].b.z;
|
|
}
|
|
)";
|
|
|
|
// sb.b[4].b[1].b.z
|
|
// ^ ^ ^ ^ ^ ^
|
|
// | | | | | |
|
|
// 128 | |1200| 1224
|
|
// | | |
|
|
// 1152 1168 1216
|
|
|
|
auto* expect = R"(
|
|
struct S1 {
|
|
a : i32;
|
|
b : vec3<f32>;
|
|
c : i32;
|
|
};
|
|
|
|
struct S2 {
|
|
a : i32;
|
|
b : [[stride(32)]] array<S1, 3>;
|
|
c : i32;
|
|
};
|
|
|
|
[[block]]
|
|
struct SB {
|
|
[[size(128)]]
|
|
a : i32;
|
|
b : [[stride(256)]] array<S2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : 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)]]
|
|
fn main() {
|
|
var x : f32 = tint_symbol(sb, 1224u);
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, ComplexDynamicAccessChain) {
|
|
auto* src = R"(
|
|
struct S1 {
|
|
a : i32;
|
|
b : vec3<f32>;
|
|
c : i32;
|
|
};
|
|
|
|
struct S2 {
|
|
a : i32;
|
|
b : [[stride(32)]] array<S1, 3>;
|
|
c : i32;
|
|
};
|
|
|
|
[[block]]
|
|
struct SB {
|
|
[[size(128)]]
|
|
a : i32;
|
|
b : [[stride(256)]] array<S2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var i : i32 = 4;
|
|
var j : u32 = 1u;
|
|
var k : i32 = 2;
|
|
var x : f32 = sb.b[i].b[j].b[k];
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
struct S1 {
|
|
a : i32;
|
|
b : vec3<f32>;
|
|
c : i32;
|
|
};
|
|
|
|
struct S2 {
|
|
a : i32;
|
|
b : [[stride(32)]] array<S1, 3>;
|
|
c : i32;
|
|
};
|
|
|
|
[[block]]
|
|
struct SB {
|
|
[[size(128)]]
|
|
a : i32;
|
|
b : [[stride(256)]] array<S2>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : 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)]]
|
|
fn main() {
|
|
var i : i32 = 4;
|
|
var j : u32 = 1u;
|
|
var k : i32 = 2;
|
|
var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, ComplexDynamicAccessChainWithAliases) {
|
|
auto* src = R"(
|
|
struct S1 {
|
|
a : i32;
|
|
b : vec3<f32>;
|
|
c : i32;
|
|
};
|
|
|
|
type A1 = S1;
|
|
|
|
type A1_Array = [[stride(32)]] array<S1, 3>;
|
|
|
|
struct S2 {
|
|
a : i32;
|
|
b : A1_Array;
|
|
c : i32;
|
|
};
|
|
|
|
type A2 = S2;
|
|
|
|
type A2_Array = [[stride(256)]] array<S2>;
|
|
|
|
[[block]]
|
|
struct SB {
|
|
[[size(128)]]
|
|
a : i32;
|
|
b : A2_Array;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
var i : i32 = 4;
|
|
var j : u32 = 1u;
|
|
var k : i32 = 2;
|
|
var x : f32 = sb.b[i].b[j].b[k];
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
struct S1 {
|
|
a : i32;
|
|
b : vec3<f32>;
|
|
c : i32;
|
|
};
|
|
|
|
type A1 = S1;
|
|
|
|
type A1_Array = [[stride(32)]] array<S1, 3>;
|
|
|
|
struct S2 {
|
|
a : i32;
|
|
b : A1_Array;
|
|
c : i32;
|
|
};
|
|
|
|
type A2 = S2;
|
|
|
|
type A2_Array = [[stride(256)]] array<S2>;
|
|
|
|
[[block]]
|
|
struct SB {
|
|
[[size(128)]]
|
|
a : i32;
|
|
b : A2_Array;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : 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)]]
|
|
fn main() {
|
|
var i : i32 = 4;
|
|
var j : u32 = 1u;
|
|
var k : i32 = 2;
|
|
var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, StorageBufferAtomics) {
|
|
auto* src = R"(
|
|
[[block]]
|
|
struct SB {
|
|
padding : vec4<f32>;
|
|
a : atomic<i32>;
|
|
b : atomic<u32>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
atomicStore(&sb.a, 123);
|
|
ignore(atomicLoad(&sb.a));
|
|
ignore(atomicAdd(&sb.a, 123));
|
|
ignore(atomicSub(&sb.a, 123));
|
|
ignore(atomicMax(&sb.a, 123));
|
|
ignore(atomicMin(&sb.a, 123));
|
|
ignore(atomicAnd(&sb.a, 123));
|
|
ignore(atomicOr(&sb.a, 123));
|
|
ignore(atomicXor(&sb.a, 123));
|
|
ignore(atomicExchange(&sb.a, 123));
|
|
ignore(atomicCompareExchangeWeak(&sb.a, 123, 345));
|
|
|
|
atomicStore(&sb.b, 123u);
|
|
ignore(atomicLoad(&sb.b));
|
|
ignore(atomicAdd(&sb.b, 123u));
|
|
ignore(atomicSub(&sb.b, 123u));
|
|
ignore(atomicMax(&sb.b, 123u));
|
|
ignore(atomicMin(&sb.b, 123u));
|
|
ignore(atomicAnd(&sb.b, 123u));
|
|
ignore(atomicOr(&sb.b, 123u));
|
|
ignore(atomicXor(&sb.b, 123u));
|
|
ignore(atomicExchange(&sb.b, 123u));
|
|
ignore(atomicCompareExchangeWeak(&sb.b, 123u, 345u));
|
|
}
|
|
)";
|
|
|
|
auto* expect = R"(
|
|
[[block]]
|
|
struct SB {
|
|
padding : vec4<f32>;
|
|
a : atomic<i32>;
|
|
b : atomic<u32>;
|
|
};
|
|
|
|
[[group(0), binding(0)]] var<storage, read_write> sb : SB;
|
|
|
|
[[internal(intrinsic_atomic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32)
|
|
|
|
[[internal(intrinsic_atomic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_add_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_sub_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_max_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_min_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_and_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_or_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_xor_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_exchange_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32) -> i32
|
|
|
|
[[internal(intrinsic_atomic_compare_exchange_weak_storage_i32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2<i32>
|
|
|
|
[[internal(intrinsic_atomic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32)
|
|
|
|
[[internal(intrinsic_atomic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_add_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_sub_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_max_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_min_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_and_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_or_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_xor_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_exchange_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32) -> u32
|
|
|
|
[[internal(intrinsic_atomic_compare_exchange_weak_storage_u32), internal(disable_validation__function_has_no_body)]]
|
|
fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2<u32>
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
tint_symbol(sb, 16u, 123);
|
|
ignore(tint_symbol_1(sb, 16u));
|
|
ignore(tint_symbol_2(sb, 16u, 123));
|
|
ignore(tint_symbol_3(sb, 16u, 123));
|
|
ignore(tint_symbol_4(sb, 16u, 123));
|
|
ignore(tint_symbol_5(sb, 16u, 123));
|
|
ignore(tint_symbol_6(sb, 16u, 123));
|
|
ignore(tint_symbol_7(sb, 16u, 123));
|
|
ignore(tint_symbol_8(sb, 16u, 123));
|
|
ignore(tint_symbol_9(sb, 16u, 123));
|
|
ignore(tint_symbol_10(sb, 16u, 123, 345));
|
|
tint_symbol_11(sb, 20u, 123u);
|
|
ignore(tint_symbol_12(sb, 20u));
|
|
ignore(tint_symbol_13(sb, 20u, 123u));
|
|
ignore(tint_symbol_14(sb, 20u, 123u));
|
|
ignore(tint_symbol_15(sb, 20u, 123u));
|
|
ignore(tint_symbol_16(sb, 20u, 123u));
|
|
ignore(tint_symbol_17(sb, 20u, 123u));
|
|
ignore(tint_symbol_18(sb, 20u, 123u));
|
|
ignore(tint_symbol_19(sb, 20u, 123u));
|
|
ignore(tint_symbol_20(sb, 20u, 123u));
|
|
ignore(tint_symbol_21(sb, 20u, 123u, 345u));
|
|
}
|
|
)";
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
TEST_F(DecomposeMemoryAccessTest, WorkgroupBufferAtomics) {
|
|
auto* src = R"(
|
|
struct S {
|
|
padding : vec4<f32>;
|
|
a : atomic<i32>;
|
|
b : atomic<u32>;
|
|
};
|
|
|
|
var<workgroup> w : S;
|
|
|
|
[[stage(compute), workgroup_size(1)]]
|
|
fn main() {
|
|
atomicStore(&(w.a), 123);
|
|
ignore(atomicLoad(&(w.a)));
|
|
ignore(atomicAdd(&(w.a), 123));
|
|
ignore(atomicSub(&(w.a), 123));
|
|
ignore(atomicMax(&(w.a), 123));
|
|
ignore(atomicMin(&(w.a), 123));
|
|
ignore(atomicAnd(&(w.a), 123));
|
|
ignore(atomicOr(&(w.a), 123));
|
|
ignore(atomicXor(&(w.a), 123));
|
|
ignore(atomicExchange(&(w.a), 123));
|
|
ignore(atomicCompareExchangeWeak(&(w.a), 123, 345));
|
|
atomicStore(&(w.b), 123u);
|
|
ignore(atomicLoad(&(w.b)));
|
|
ignore(atomicAdd(&(w.b), 123u));
|
|
ignore(atomicSub(&(w.b), 123u));
|
|
ignore(atomicMax(&(w.b), 123u));
|
|
ignore(atomicMin(&(w.b), 123u));
|
|
ignore(atomicAnd(&(w.b), 123u));
|
|
ignore(atomicOr(&(w.b), 123u));
|
|
ignore(atomicXor(&(w.b), 123u));
|
|
ignore(atomicExchange(&(w.b), 123u));
|
|
ignore(atomicCompareExchangeWeak(&(w.b), 123u, 345u));
|
|
}
|
|
)";
|
|
|
|
auto* expect = src;
|
|
|
|
auto got = Run<DecomposeMemoryAccess>(src);
|
|
|
|
EXPECT_EQ(expect, str(got));
|
|
}
|
|
|
|
} // namespace
|
|
} // namespace transform
|
|
} // namespace tint
|