tint: add syntax sugar for dereferencing pointers
Adds support to Tint for syntax sugar for dereferencing pointers for
member or index access as per:
https://github.com/gpuweb/gpuweb/pull/4311
- Resolver: when the lhs of a accessor expression is a pointer, it is
now resolved to a sem::Reference.
- Added "pointer_composite_access" feature as experimental, hooked up
validation in Resolver, and added tests.
- Added resolver tests for the new syntax to resolver/ptr_ref_test.cc.
- Fixed multiple transforms to deal with the fact that the lhs of
accessor expressions can now be pointers, including: Robustness,
Renamer, Std140, and SimplifyPointers.
- In transforms that rely on other transforms, such as
SimplifyPointers, to remove/inline pointers, I added asserts that
the type is not a pointer.
- Added unit tests for transforms that use pointer-dot/index for
accessor expressions.
- Fixed uniformity analysis code so that ProcessLValueExpression
correctly deals with accessor expressions where the object is a
pointer, in the same way we do for UnaryOp::kIndirection, including
partial pointer checks. Added many tests for these new cases.
- Fixed ProgramToIR so that EmitAccess handles the new syntax. Added
multiple tests.
- Added end2end tests under test/tint/ptr_sugar
For Googlers, see my work log at
go/add-syntax-sugar-for-dereferencing-composites for more details.
Bug: tint:2113
Change-Id: I7a0093f52ca2237be598e44245b45049f21d056c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/164900
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/dawn/common/WGSLFeatureMapping.h b/src/dawn/common/WGSLFeatureMapping.h
index 14ba2b6..4810cee 100644
--- a/src/dawn/common/WGSLFeatureMapping.h
+++ b/src/dawn/common/WGSLFeatureMapping.h
@@ -32,6 +32,7 @@
X(kUndefined, Undefined) \
X(kReadonlyAndReadwriteStorageTextures, ReadonlyAndReadwriteStorageTextures) \
X(kPacked4X8IntegerDotProduct, Packed4x8IntegerDotProduct) \
+ X(kPointerCompositeAccess, PointerCompositeAccess) \
X(kChromiumTestingUnimplemented, ChromiumTestingUnimplemented) \
X(kChromiumTestingUnsafeExperimental, ChromiumTestingUnsafeExperimental) \
X(kChromiumTestingExperimental, ChromiumTestingExperimental) \
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index 452abc5..38219bd 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -652,7 +652,7 @@
break;
default: {
auto* vec = TypeOf(lhs_row_access->object)
- ->UnwrapRef()
+ ->UnwrapPtrOrRef()
->As<core::type::Vector>();
TINT_UNREACHABLE() << "invalid vector size " << vec->Width();
break;
@@ -908,10 +908,24 @@
bool ASTPrinter::EmitAssign(const ast::AssignmentStatement* stmt) {
if (auto* lhs_access = stmt->lhs->As<ast::IndexAccessorExpression>()) {
+ auto validate_obj_not_pointer = [&](const core::type::Type* object_ty) {
+ if (TINT_UNLIKELY(object_ty->Is<core::type::Pointer>())) {
+ TINT_ICE() << "lhs of index accessor should not be a pointer. These should have "
+ "been removed by transforms such as SimplifyPointers, "
+ "DecomposeMemoryAccess, and DirectVariableAccess";
+ return false;
+ }
+ return true;
+ };
+
// BUG(crbug.com/tint/1333): work around assignment of scalar to matrices
// with at least one dynamic index
if (auto* lhs_sub_access = lhs_access->object->As<ast::IndexAccessorExpression>()) {
- if (auto* mat = TypeOf(lhs_sub_access->object)->UnwrapRef()->As<core::type::Matrix>()) {
+ const auto* lhs_sub_access_type = TypeOf(lhs_sub_access->object);
+ if (!validate_obj_not_pointer(lhs_sub_access_type)) {
+ return false;
+ }
+ if (auto* mat = lhs_sub_access_type->UnwrapRef()->As<core::type::Matrix>()) {
auto* rhs_row_idx_sem = builder_.Sem().GetVal(lhs_access->index);
auto* rhs_col_idx_sem = builder_.Sem().GetVal(lhs_sub_access->index);
if (!rhs_row_idx_sem->ConstantValue() || !rhs_col_idx_sem->ConstantValue()) {
@@ -921,8 +935,11 @@
}
// BUG(crbug.com/tint/1333): work around assignment of vector to matrices
// with dynamic indices
- const auto* lhs_access_type = TypeOf(lhs_access->object)->UnwrapRef();
- if (auto* mat = lhs_access_type->As<core::type::Matrix>()) {
+ const auto* lhs_access_type = TypeOf(lhs_access->object);
+ if (!validate_obj_not_pointer(lhs_access_type)) {
+ return false;
+ }
+ if (auto* mat = lhs_access_type->UnwrapRef()->As<core::type::Matrix>()) {
auto* lhs_index_sem = builder_.Sem().GetVal(lhs_access->index);
if (!lhs_index_sem->ConstantValue()) {
return EmitDynamicMatrixVectorAssignment(stmt, mat);
@@ -930,7 +947,7 @@
}
// BUG(crbug.com/tint/534): work around assignment to vectors with dynamic
// indices
- if (auto* vec = lhs_access_type->As<core::type::Vector>()) {
+ if (auto* vec = lhs_access_type->UnwrapRef()->As<core::type::Vector>()) {
auto* rhs_sem = builder_.Sem().GetVal(lhs_access->index);
if (!rhs_sem->ConstantValue()) {
return EmitDynamicVectorAssignment(stmt, vec);
diff --git a/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length.cc b/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length.cc
index 0dacf5e..3dd8c8a 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length.cc
@@ -178,6 +178,12 @@
"&struct_var.array_member";
break;
}
+ if (TINT_UNLIKELY(storage_buffer_sem->Type()->Is<core::type::Pointer>())) {
+ TINT_ICE()
+ << "storage buffer variable should not be a pointer. These should have "
+ "been removed by the SimplifyPointers transform";
+ break;
+ }
auto* storage_buffer_var = storage_buffer_sem->Variable();
auto* storage_buffer_type =
storage_buffer_sem->Type()->As<core::type::Reference>();
diff --git a/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length_test.cc b/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length_test.cc
index 4768b3a..60e769b 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/calculate_array_length_test.cc
@@ -149,6 +149,47 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(CalculateArrayLengthTest, BasicInStruct_ViaPointerDot) {
+ auto* src = R"(
+struct SB {
+ x : i32,
+ arr : array<i32>,
+};
+
+@group(0) @binding(0) var<storage, read> sb : SB;
+
+@compute @workgroup_size(1)
+fn main() {
+ let p = &sb;
+ var len : u32 = arrayLength(&p.arr);
+}
+)";
+
+ auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__function_parameter) buffer : ptr<storage, SB, read>, result : ptr<function, u32>)
+
+struct SB {
+ x : i32,
+ arr : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read> sb : SB;
+
+@compute @workgroup_size(1)
+fn main() {
+ var tint_symbol_1 : u32 = 0u;
+ tint_symbol(&(sb), &(tint_symbol_1));
+ let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u);
+ var len : u32 = tint_symbol_2;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(CalculateArrayLengthTest, ArrayOfStruct) {
auto* src = R"(
struct S {
diff --git a/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access_test.cc b/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access_test.cc
index 32414a1..d9addc3 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access_test.cc
@@ -26,8 +26,8 @@
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
#include "src/tint/lang/hlsl/writer/ast_raise/decompose_memory_access.h"
-
#include "src/tint/lang/wgsl/ast/transform/helper_test.h"
+#include "src/tint/lang/wgsl/ast/transform/simplify_pointers.h"
namespace tint::hlsl::writer {
namespace {
@@ -3198,6 +3198,81 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain_ViaPointerDot) {
+ auto* src = R"(
+// sizeof(S1) == 32
+// alignof(S1) == 16
+struct S1 {
+ a : i32,
+ b : vec3<f32>,
+ c : i32,
+};
+
+// sizeof(S2) == 116
+// alignof(S2) == 16
+struct S2 {
+ a : i32,
+ b : array<S1, 3>,
+ c : i32,
+};
+
+struct SB {
+ @size(128)
+ a : i32,
+ b : array<S2>,
+};
+
+@group(0) @binding(0) var<storage, read_write> sb : SB;
+
+@compute @workgroup_size(1)
+fn main() {
+ let p = &sb;
+ var x : f32 = (*p).b[4].b[1].b.z;
+}
+)";
+
+ // sb.b[4].b[1].b.z
+ // ^ ^ ^ ^ ^ ^
+ // | | | | | |
+ // 128 | |688 | 712
+ // | | |
+ // 640 656 704
+
+ auto* expect = R"(
+struct S1 {
+ a : i32,
+ b : vec3<f32>,
+ c : i32,
+}
+
+struct S2 {
+ a : i32,
+ b : array<S1, 3>,
+ c : i32,
+}
+
+struct SB {
+ @size(128)
+ a : i32,
+ b : 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 sb_load(offset : u32) -> f32
+
+@compute @workgroup_size(1)
+fn main() {
+ var x : f32 = sb_load(712u);
+}
+)";
+
+ auto got = Run<ast::transform::SimplifyPointers, DecomposeMemoryAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain_OutOfOrder) {
auto* src = R"(
@compute @workgroup_size(1)
diff --git a/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment.cc b/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment.cc
index c52041d..84574c8 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment.cc
@@ -183,8 +183,15 @@
if (!idx_sem->ConstantValue()) {
// Indexing a member access expr?
if (auto* ma = ia->object->As<ast::MemberAccessorExpression>()) {
+ const auto* ma_ty = src.TypeOf(ma);
+ if (TINT_UNLIKELY(ma_ty->Is<core::type::Pointer>())) {
+ TINT_ICE()
+ << "lhs of index accessor expression should not be a pointer. These "
+ "should have been removed by the SimplifyPointers transform";
+ return ast::TraverseAction::Stop;
+ }
// That accesses an array?
- if (src.TypeOf(ma)->UnwrapRef()->Is<core::type::Array>()) {
+ if (ma_ty->UnwrapRef()->Is<core::type::Array>()) {
result = true;
return ast::TraverseAction::Stop;
}
diff --git a/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment_test.cc b/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment_test.cc
index 0ad39aa..c6a6f6b 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment_test.cc
@@ -699,6 +699,67 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg_PointerDot) {
+ auto* src = R"(
+struct Uniforms {
+ i : u32,
+};
+struct InnerS {
+ v : i32,
+};
+struct OuterS {
+ a1 : array<InnerS, 8>,
+};
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+fn f(p : ptr<function, OuterS>) {
+ var v : InnerS;
+ p.a1[uniforms.i] = v;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var s1 : OuterS;
+ f(&s1);
+}
+)";
+
+ auto* expect = R"(
+struct Uniforms {
+ i : u32,
+}
+
+struct InnerS {
+ v : i32,
+}
+
+struct OuterS {
+ a1 : array<InnerS, 8>,
+}
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+fn f(p : ptr<function, OuterS>) {
+ var v : InnerS;
+ {
+ let tint_symbol = &((*(p)).a1);
+ var tint_symbol_1 = *(tint_symbol);
+ tint_symbol_1[uniforms.i] = v;
+ *(tint_symbol) = tint_symbol_1;
+ }
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ var s1 : OuterS;
+ f(&(s1));
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg_OutOfOrder) {
auto* src = R"(
@compute @workgroup_size(1)
@@ -858,5 +919,91 @@
EXPECT_FALSE(ShouldRun<LocalizeStructArrayAssignment>(src));
}
+TEST_F(LocalizeStructArrayAssignmentTest, ArrayStructArray) {
+ auto* src = R"(
+struct Uniforms {
+ i : u32,
+};
+
+struct InnerS {
+ v : i32,
+};
+
+struct OuterS {
+ a1 : array<InnerS, 8>,
+};
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+@compute @workgroup_size(1)
+fn main() {
+ var v : InnerS;
+ var s1 : array<OuterS, 2>;
+ s1[uniforms.i].a1[uniforms.i] = v;
+}
+)";
+
+ // Transform does nothing as the struct-of-array is in an array, which FXC has no problem with.
+ EXPECT_FALSE(ShouldRun<LocalizeStructArrayAssignment>(src));
+}
+
+TEST_F(LocalizeStructArrayAssignmentTest, ArrayStructArray_ViaPointerDerefIndex) {
+ auto* src = R"(
+struct Uniforms {
+ i : u32,
+};
+
+struct InnerS {
+ v : i32,
+};
+
+struct OuterS {
+ a1 : array<InnerS, 8>,
+};
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+@compute @workgroup_size(1)
+fn main() {
+ var v : InnerS;
+ var s1 : array<OuterS, 2>;
+ let p = &s1;
+ (*p)[uniforms.i].a1[uniforms.i] = v;
+}
+)";
+
+ // Transform does nothing as the struct-of-array is in an array, which FXC has no problem with.
+ EXPECT_FALSE(ShouldRun<LocalizeStructArrayAssignment>(src));
+}
+
+TEST_F(LocalizeStructArrayAssignmentTest, ArrayStructArray_ViaPointerIndex) {
+ auto* src = R"(
+struct Uniforms {
+ i : u32,
+};
+
+struct InnerS {
+ v : i32,
+};
+
+struct OuterS {
+ a1 : array<InnerS, 8>,
+};
+
+@group(1) @binding(4) var<uniform> uniforms : Uniforms;
+
+@compute @workgroup_size(1)
+fn main() {
+ var v : InnerS;
+ var s1 : array<OuterS, 2>;
+ let p = &s1;
+ p[uniforms.i].a1[uniforms.i] = v;
+}
+)";
+
+ // Transform does nothing as the struct-of-array is in an array, which FXC has no problem with.
+ EXPECT_FALSE(ShouldRun<LocalizeStructArrayAssignment>(src));
+}
+
} // namespace
} // namespace tint::hlsl::writer
diff --git a/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform_test.cc b/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform_test.cc
index a8fed98..d27b132 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform_test.cc
@@ -111,6 +111,124 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(NumWorkgroupsFromUniformTest, Basic_VarCopy) {
+ auto* src = R"(
+@compute @workgroup_size(1)
+fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
+ var a = num_wgs;
+ let groups_x = a.x;
+ let groups_y = a.y;
+ let groups_z = a.z;
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_2 {
+ num_workgroups : vec3<u32>,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_3 : tint_symbol_2;
+
+fn main_inner(num_wgs : vec3<u32>) {
+ var a = num_wgs;
+ let groups_x = a.x;
+ let groups_y = a.y;
+ let groups_z = a.z;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ main_inner(tint_symbol_3.num_workgroups);
+}
+)";
+
+ ast::transform::DataMap data;
+ data.Add<CanonicalizeEntryPointIO::Config>(CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
+ data.Add<NumWorkgroupsFromUniform::Config>(BindingPoint{0, 30u});
+ auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(NumWorkgroupsFromUniformTest, Basic_VarCopy_ViaPointerDerefDot) {
+ auto* src = R"(
+@compute @workgroup_size(1)
+fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
+ var a = num_wgs;
+ let p = &a;
+ let groups_x = (*p).x;
+ let groups_y = (*p).y;
+ let groups_z = (*p).z;
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_2 {
+ num_workgroups : vec3<u32>,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_3 : tint_symbol_2;
+
+fn main_inner(num_wgs : vec3<u32>) {
+ var a = num_wgs;
+ let p = &(a);
+ let groups_x = (*(p)).x;
+ let groups_y = (*(p)).y;
+ let groups_z = (*(p)).z;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ main_inner(tint_symbol_3.num_workgroups);
+}
+)";
+
+ ast::transform::DataMap data;
+ data.Add<CanonicalizeEntryPointIO::Config>(CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
+ data.Add<NumWorkgroupsFromUniform::Config>(BindingPoint{0, 30u});
+ auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(NumWorkgroupsFromUniformTest, Basic_VarCopy_ViaPointerDot) {
+ auto* src = R"(
+@compute @workgroup_size(1)
+fn main(@builtin(num_workgroups) num_wgs : vec3<u32>) {
+ var a = num_wgs;
+ let p = &a;
+ let groups_x = p.x;
+ let groups_y = p.y;
+ let groups_z = p.z;
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol_2 {
+ num_workgroups : vec3<u32>,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_3 : tint_symbol_2;
+
+fn main_inner(num_wgs : vec3<u32>) {
+ var a = num_wgs;
+ let p = &(a);
+ let groups_x = p.x;
+ let groups_y = p.y;
+ let groups_z = p.z;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ main_inner(tint_symbol_3.num_workgroups);
+}
+)";
+
+ ast::transform::DataMap data;
+ data.Add<CanonicalizeEntryPointIO::Config>(CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
+ data.Add<NumWorkgroupsFromUniform::Config>(BindingPoint{0, 30u});
+ auto got = Run<Unshadow, CanonicalizeEntryPointIO, NumWorkgroupsFromUniform>(src, data);
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(NumWorkgroupsFromUniformTest, StructOnlyMember) {
auto* src = R"(
struct Builtins {
diff --git a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
index b3f3fb1..325f528 100644
--- a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
+++ b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
@@ -291,6 +291,60 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, Vec3_ReadComponent_IndexAccessor_ViaDerefPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage> v : vec3<f32>;
+
+fn f() {
+ let p = &v;
+ let x = (*p)[1];
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+@group(0) @binding(0) var<storage> v : __packed_vec3<f32>;
+
+fn f() {
+ let p = &(v);
+ let x = (*(p))[1];
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, Vec3_ReadComponent_IndexAccessor_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage> v : vec3<f32>;
+
+fn f() {
+ let p = &v;
+ let x = p[1];
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+@group(0) @binding(0) var<storage> v : __packed_vec3<f32>;
+
+fn f() {
+ let p = &(v);
+ let x = p[1];
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, Vec3_WriteVector_ValueRHS) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> v : vec3<f32>;
@@ -369,6 +423,60 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PackedVec3Test, Vec3_WriteComponent_MemberAccessor_ViaDerefPointerDot) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : vec3<f32>;
+
+fn f() {
+ let p = &v;
+ (*p).y = 1.23;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+@group(0) @binding(0) var<storage, read_write> v : __packed_vec3<f32>;
+
+fn f() {
+ let p = &(v);
+ (*(p)).y = 1.22999999999999998224;
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, Vec3_WriteComponent_MemberAccessor_ViaPointerDot) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> v : vec3<f32>;
+
+fn f() {
+ let p = &v;
+ p.y = 1.23;
+}
+)";
+
+ auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+@group(0) @binding(0) var<storage, read_write> v : __packed_vec3<f32>;
+
+fn f() {
+ let p = &(v);
+ p.y = 1.22999999999999998224;
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<PackedVec3>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PackedVec3Test, Vec3_WriteComponent_IndexAccessor) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> v : vec3<f32>;
diff --git a/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc b/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc
index 54a1f9d..818297c 100644
--- a/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc
+++ b/src/tint/lang/spirv/reader/ast_lower/atomics_test.cc
@@ -200,6 +200,54 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(AtomicsTest, ArrayOfU32_ViaDerefPointerIndex) {
+ auto* src = R"(
+var<workgroup> wg : array<u32, 4>;
+
+fn f() {
+ let p = &wg;
+ stub_atomicStore_u32((*p)[1], 1u);
+}
+)";
+
+ auto* expect = R"(
+var<workgroup> wg : array<atomic<u32>, 4u>;
+
+fn f() {
+ let p = &(wg);
+ atomicStore(&((*(p))[1]), 1u);
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AtomicsTest, ArrayOfU32_ViaPointerIndex) {
+ auto* src = R"(
+var<workgroup> wg : array<u32, 4>;
+
+fn f() {
+ let p = &wg;
+ stub_atomicStore_u32(p[1], 1u);
+}
+)";
+
+ auto* expect = R"(
+var<workgroup> wg : array<atomic<u32>, 4u>;
+
+fn f() {
+ let p = &(wg);
+ atomicStore(&(p[1]), 1u);
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(AtomicsTest, ArraysOfU32) {
auto* src = R"(
var<workgroup> wg : array<array<array<u32, 1>, 2>, 3>;
@@ -294,6 +342,78 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(AtomicsTest, FlatStructSingleAtomic_ViaDerefPointerDot) {
+ auto* src = R"(
+struct S {
+ a : u32,
+}
+
+var<workgroup> wg : S;
+
+fn f() {
+ let p = &wg;
+ stub_atomicStore_u32((*p).a, 1u);
+}
+)";
+
+ auto* expect = R"(
+struct S_atomic {
+ a : atomic<u32>,
+}
+
+struct S {
+ a : u32,
+}
+
+var<workgroup> wg : S_atomic;
+
+fn f() {
+ let p = &(wg);
+ atomicStore(&((*(p)).a), 1u);
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AtomicsTest, FlatStructSingleAtomic_ViaPointerDot) {
+ auto* src = R"(
+struct S {
+ a : u32,
+}
+
+var<workgroup> wg : S;
+
+fn f() {
+ let p = &wg;
+ stub_atomicStore_u32(p.a, 1u);
+}
+)";
+
+ auto* expect = R"(
+struct S_atomic {
+ a : atomic<u32>,
+}
+
+struct S {
+ a : u32,
+}
+
+var<workgroup> wg : S_atomic;
+
+fn f() {
+ let p = &(wg);
+ atomicStore(&(p.a), 1u);
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(AtomicsTest, FlatStructMultipleAtomic) {
auto* src = R"(
struct S {
@@ -455,6 +575,90 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(AtomicsTest, ArrayOfStruct_ViaDerefPointerIndex) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : i32,
+ c : u32,
+}
+
+@group(0) @binding(1) var<storage, read_write> arr : array<S>;
+
+fn f() {
+ let p = &arr;
+ stub_atomicStore_i32((*p)[4].b, 1i);
+}
+)";
+
+ auto* expect = R"(
+struct S_atomic {
+ a : u32,
+ b : atomic<i32>,
+ c : u32,
+}
+
+struct S {
+ a : u32,
+ b : i32,
+ c : u32,
+}
+
+@group(0) @binding(1) var<storage, read_write> arr : array<S_atomic>;
+
+fn f() {
+ let p = &(arr);
+ atomicStore(&((*(p))[4].b), 1i);
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AtomicsTest, ArrayOfStruct_ViaPointerIndex) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : i32,
+ c : u32,
+}
+
+@group(0) @binding(1) var<storage, read_write> arr : array<S>;
+
+fn f() {
+ let p = &arr;
+ stub_atomicStore_i32(p[4].b, 1i);
+}
+)";
+
+ auto* expect = R"(
+struct S_atomic {
+ a : u32,
+ b : atomic<i32>,
+ c : u32,
+}
+
+struct S {
+ a : u32,
+ b : i32,
+ c : u32,
+}
+
+@group(0) @binding(1) var<storage, read_write> arr : array<S_atomic>;
+
+fn f() {
+ let p = &(arr);
+ atomicStore(&(p[4].b), 1i);
+}
+)";
+
+ auto got = Run(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(AtomicsTest, StructOfArray) {
auto* src = R"(
struct S {
diff --git a/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array.cc b/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array.cc
index e26ac38..47e242b8 100644
--- a/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array.cc
+++ b/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array.cc
@@ -143,6 +143,11 @@
// Example: `arr[i]` -> `arr[i].el`
ctx.ReplaceAll([&](const ast::IndexAccessorExpression* idx) -> const ast::Expression* {
if (auto* ty = src.TypeOf(idx->object)) {
+ if (TINT_UNLIKELY(ty->Is<core::type::Pointer>())) {
+ TINT_ICE() << "lhs of index accessor expression should not be a pointer. These "
+ "should have been removed by the SimplifyPointers transform";
+ return nullptr;
+ }
if (auto* arr = ty->UnwrapRef()->As<core::type::Array>()) {
if (!arr->IsStrideImplicit()) {
auto* expr = ctx.CloneWithoutTransform(idx);
diff --git a/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array_test.cc b/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array_test.cc
index 87d7ceb..84103ce 100644
--- a/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array_test.cc
+++ b/src/tint/lang/spirv/reader/ast_lower/decompose_strided_array_test.cc
@@ -100,7 +100,7 @@
//
// @compute @workgroup_size(1)
// fn f() {
- // let a : @stride(4) array<f32, 4u> = a;
+ // let a : @stride(4) array<f32, 4u> = arr;
// let b : f32 = arr[1];
// }
@@ -139,6 +139,52 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DecomposeStridedArrayTest, PrivateDefaultStridedArray_ViaPointerIndex) {
+ // var<private> arr : @stride(4) array<f32, 4u>
+ //
+ // @compute @workgroup_size(1)
+ // fn f() {
+ // let a : @stride(4) array<f32, 4u> = arr;
+ // let p = &arr;
+ // let b : f32 = p[1];
+ // }
+
+ ProgramBuilder b;
+ b.GlobalVar("arr",
+ b.ty.array<f32, 4u>(Vector{
+ b.Stride(4),
+ }),
+ core::AddressSpace::kPrivate);
+ b.Func("f", tint::Empty, b.ty.void_(),
+ Vector{
+ b.Decl(b.Let("a",
+ b.ty.array<f32, 4u>(Vector{
+ b.Stride(4),
+ }),
+ b.Expr("arr"))),
+ b.Decl(b.Let("p", b.AddressOf(b.Expr("arr")))),
+ b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor("p", 1_i))),
+ },
+ Vector{
+ b.Stage(ast::PipelineStage::kCompute),
+ b.WorkgroupSize(1_i),
+ });
+
+ auto* expect = R"(
+var<private> arr : array<f32, 4u>;
+
+@compute @workgroup_size(1i)
+fn f() {
+ let a : array<f32, 4u> = arr;
+ let b : f32 = arr[1i];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(resolver::Resolve(b));
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(DecomposeStridedArrayTest, PrivateStridedArray) {
// var<private> arr : @stride(32) array<f32, 4u>
//
diff --git a/src/tint/lang/spirv/writer/ast_printer/accessor_expression_test.cc b/src/tint/lang/spirv/writer/ast_printer/accessor_expression_test.cc
index f09814b..1dde2c8 100644
--- a/src/tint/lang/spirv/writer/ast_printer/accessor_expression_test.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/accessor_expression_test.cc
@@ -138,6 +138,84 @@
Validate(b);
}
+TEST_F(SpirvASTPrinterTest, Runtime_IndexAccessor_Vector_ViaDerefPointerIndex) {
+ // var ary : vec3<u32>;
+ // let p = &ary;
+ // var x = (*p)[1i];
+
+ auto* ary = Var("ary", ty.vec3<u32>());
+ auto* p = Let("p", AddressOf("ary"));
+ auto* x = Var("x", IndexAccessor(Deref(p), 1_i));
+ WrapInFunction(ary, p, x);
+
+ Builder& b = SanitizeAndBuild();
+
+ ASSERT_TRUE(b.Build()) << b.Diagnostics();
+
+ EXPECT_EQ(DumpInstructions(b.Module().Types()), R"(%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeInt 32 0
+%7 = OpTypeVector %8 3
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%10 = OpTypeInt 32 1
+%11 = OpConstant %10 1
+%12 = OpTypePointer Function %8
+%16 = OpConstantNull %8
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].variables()),
+ R"(%5 = OpVariable %6 Function %9
+%15 = OpVariable %12 Function %16
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].instructions()),
+ R"(%13 = OpAccessChain %12 %5 %11
+%14 = OpLoad %8 %13
+OpStore %15 %14
+OpReturn
+)");
+
+ Validate(b);
+}
+
+TEST_F(SpirvASTPrinterTest, Runtime_IndexAccessor_Vector_ViaPointerIndex) {
+ // var ary : vec3<u32>;
+ // let p = &ary;
+ // var x = p[1i];
+
+ auto* ary = Var("ary", ty.vec3<u32>());
+ auto* p = Let("p", AddressOf("ary"));
+ auto* x = Var("x", IndexAccessor(p, 1_i));
+ WrapInFunction(ary, p, x);
+
+ Builder& b = SanitizeAndBuild();
+
+ ASSERT_TRUE(b.Build()) << b.Diagnostics();
+
+ EXPECT_EQ(DumpInstructions(b.Module().Types()), R"(%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeInt 32 0
+%7 = OpTypeVector %8 3
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%10 = OpTypeInt 32 1
+%11 = OpConstant %10 1
+%12 = OpTypePointer Function %8
+%16 = OpConstantNull %8
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].variables()),
+ R"(%5 = OpVariable %6 Function %9
+%15 = OpVariable %12 Function %16
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].instructions()),
+ R"(%13 = OpAccessChain %12 %5 %11
+%14 = OpLoad %8 %13
+OpStore %15 %14
+OpReturn
+)");
+
+ Validate(b);
+}
+
TEST_F(SpirvASTPrinterTest, Dynamic_IndexAccessor_Vector) {
// var ary : vec3<f32>;
// var idx : i32;
@@ -978,6 +1056,96 @@
Validate(b);
}
+TEST_F(SpirvASTPrinterTest, MemberAccessor_ViaDerefPointerDot) {
+ // my_struct {
+ // a : f32
+ // b : f32
+ // }
+ // var ident : my_struct
+ // let p = &ident;
+ // (*p).b
+
+ auto* s = Structure("my_struct", Vector{
+ Member("a", ty.f32()),
+ Member("b", ty.f32()),
+ });
+
+ auto* var = Var("ident", ty.Of(s));
+ auto* p = Let("p", AddressOf(var));
+ auto* expr = MemberAccessor(Deref(p), "b");
+ WrapInFunction(var, p, expr);
+
+ Builder& b = SanitizeAndBuild();
+
+ ASSERT_TRUE(b.Build()) << b.Diagnostics();
+
+ EXPECT_EQ(DumpInstructions(b.Module().Types()), R"(%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeFloat 32
+%7 = OpTypeStruct %8 %8
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%10 = OpTypeInt 32 0
+%11 = OpConstant %10 1
+%12 = OpTypePointer Function %8
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].variables()),
+ R"(%5 = OpVariable %6 Function %9
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].instructions()),
+ R"(%13 = OpAccessChain %12 %5 %11
+%14 = OpLoad %8 %13
+OpReturn
+)");
+
+ Validate(b);
+}
+
+TEST_F(SpirvASTPrinterTest, MemberAccessor_ViaPointerDot) {
+ // my_struct {
+ // a : f32
+ // b : f32
+ // }
+ // var ident : my_struct
+ // let p = &ident;
+ // p.b
+
+ auto* s = Structure("my_struct", Vector{
+ Member("a", ty.f32()),
+ Member("b", ty.f32()),
+ });
+
+ auto* var = Var("ident", ty.Of(s));
+ auto* p = Let("p", AddressOf(var));
+ auto* expr = MemberAccessor(p, "b");
+ WrapInFunction(var, p, expr);
+
+ Builder& b = SanitizeAndBuild();
+
+ ASSERT_TRUE(b.Build()) << b.Diagnostics();
+
+ EXPECT_EQ(DumpInstructions(b.Module().Types()), R"(%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%8 = OpTypeFloat 32
+%7 = OpTypeStruct %8 %8
+%6 = OpTypePointer Function %7
+%9 = OpConstantNull %7
+%10 = OpTypeInt 32 0
+%11 = OpConstant %10 1
+%12 = OpTypePointer Function %8
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].variables()),
+ R"(%5 = OpVariable %6 Function %9
+)");
+ EXPECT_EQ(DumpInstructions(b.Module().Functions()[0].instructions()),
+ R"(%13 = OpAccessChain %12 %5 %11
+%14 = OpLoad %8 %13
+OpReturn
+)");
+
+ Validate(b);
+}
+
TEST_F(SpirvASTPrinterTest, MemberAccessor_Nested) {
// inner_struct {
// a : f32
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index 19a5662..b736b28 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -866,12 +866,14 @@
}
// If the source is a reference, we access chain into it.
- // In the future, pointers may support access-chaining.
- // See https://github.com/gpuweb/gpuweb/pull/1580
if (info->source_type->Is<core::type::Reference>()) {
info->access_chain_indices.push_back(idx_id);
info->source_type = builder_.Sem().Get(expr)->UnwrapLoad()->Type();
return true;
+ } else if (TINT_UNLIKELY(info->source_type->Is<core::type::Pointer>())) {
+ TINT_ICE() << "lhs of index accesor expression should not be a pointer. These should have "
+ "been removed by the SimplifyPointers transform";
+ return false;
}
auto result_type_id = GenerateTypeIfNeeded(TypeOf(expr));
diff --git a/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc b/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc
index 91654cd..5d95703 100644
--- a/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc
+++ b/src/tint/lang/spirv/writer/ast_raise/var_for_dynamic_index_test.cc
@@ -470,6 +470,76 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(VarForDynamicIndexTest, VarArrayIndexDynamic) {
+ auto* src = R"(
+fn f() {
+ var i : i32;
+ var p = array<i32, 4>(1, 2, 3, 4);
+ let x = p[i];
+}
+)";
+
+ auto* expect = src;
+
+ ast::transform::DataMap data;
+ auto got = Run<VarForDynamicIndex>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(VarForDynamicIndexTest, VarMatrixIndexDynamic) {
+ auto* src = R"(
+fn f() {
+ var i : i32;
+ var p = mat2x2(1.0, 2.0, 3.0, 4.0);
+ let x = p[i];
+}
+)";
+
+ auto* expect = src;
+
+ ast::transform::DataMap data;
+ auto got = Run<VarForDynamicIndex>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(VarForDynamicIndexTest, VarArrayIndexDynamic_ViaPointerIndex) {
+ auto* src = R"(
+fn f() {
+ var i : i32;
+ var v = array<i32, 4>(1, 2, 3, 4);
+ let p = &(v);
+ let x = p[i];
+}
+)";
+
+ auto* expect = src;
+
+ ast::transform::DataMap data;
+ auto got = Run<VarForDynamicIndex>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(VarForDynamicIndexTest, VarMatrixIndexDynamic_ViaPointerIndex) {
+ auto* src = R"(
+fn f() {
+ var i : i32;
+ var v = mat2x2(1.0, 2.0, 3.0, 4.0);
+ let p = &(v);
+ let x = p[i];
+}
+)";
+
+ auto* expect = src;
+
+ ast::transform::DataMap data;
+ auto got = Run<VarForDynamicIndex>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(VarForDynamicIndexTest, ArrayIndexLiteral) {
auto* src = R"(
fn f() {
diff --git a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
index d98068d..27b354c 100644
--- a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
+++ b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform.cc
@@ -162,6 +162,11 @@
// array_length = ----------------------------------------
// array_stride
const Expression* total_size = total_storage_buffer_size;
+ if (TINT_UNLIKELY(storage_buffer_sem->Type()->Is<core::type::Pointer>())) {
+ TINT_ICE() << "storage buffer variable should not be a pointer. These should have "
+ "been removed by the SimplifyPointers transform";
+ return;
+ }
auto* storage_buffer_type = storage_buffer_sem->Type()->UnwrapRef();
const core::type::Array* array_type = nullptr;
if (auto* str = storage_buffer_type->As<core::type::Struct>()) {
diff --git a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc
index b6b4ab2..2a9c6ac 100644
--- a/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/array_length_from_uniform_test.cc
@@ -206,6 +206,57 @@
got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
}
+// Should output the same as BasicInStruct because SimplifyPointers outputs the same AST for
+// explicit and implicit pointer dereference.
+TEST_F(ArrayLengthFromUniformTest, BasicInStruct_ViaPointerDot) {
+ auto* src = R"(
+struct SB {
+ x : i32,
+ arr : array<i32>,
+};
+
+@group(0) @binding(0) var<storage, read> sb : SB;
+
+@compute @workgroup_size(1)
+fn main() {
+ let p = &sb;
+ var len : u32 = arrayLength(&p.arr);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ buffer_size : array<vec4<u32>, 1u>,
+}
+
+@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
+
+struct SB {
+ x : i32,
+ arr : array<i32>,
+}
+
+@group(0) @binding(0) var<storage, read> sb : SB;
+
+@compute @workgroup_size(1)
+fn main() {
+ var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
+}
+)";
+
+ ArrayLengthFromUniform::Config cfg({0, 30u});
+ cfg.bindpoint_to_size_index.emplace(BindingPoint{0, 0}, 0);
+
+ DataMap data;
+ data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(std::unordered_set<uint32_t>({0}),
+ got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
+}
+
TEST_F(ArrayLengthFromUniformTest, MultipleStorageBuffers) {
auto* src = R"(
struct SB1 {
diff --git a/src/tint/lang/wgsl/ast/transform/demote_to_helper_test.cc b/src/tint/lang/wgsl/ast/transform/demote_to_helper_test.cc
index 181a010..de590a4 100644
--- a/src/tint/lang/wgsl/ast/transform/demote_to_helper_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/demote_to_helper_test.cc
@@ -334,6 +334,56 @@
EXPECT_EQ(expect, str(got));
}
+// Test that write via sugared pointer also discards
+TEST_F(DemoteToHelperTest, WriteInEntryPoint_DiscardInEntryPoint_ViaPointerDot) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+@group(0) @binding(2) var<storage, read_write> v : vec4<f32>;
+
+@fragment
+fn foo(@location(0) in : f32, @location(1) coord : vec2<f32>) {
+ if (in == 0.0) {
+ discard;
+ }
+ let ret = textureSample(t, s, coord);
+ let p = &v;
+ p.x = ret.x;
+}
+)";
+
+ auto* expect = R"(
+var<private> tint_discarded = false;
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+@group(0) @binding(2) var<storage, read_write> v : vec4<f32>;
+
+@fragment
+fn foo(@location(0) in : f32, @location(1) coord : vec2<f32>) {
+ if ((in == 0.0)) {
+ tint_discarded = true;
+ }
+ let ret = textureSample(t, s, coord);
+ let p = &(v);
+ if (!(tint_discarded)) {
+ p.x = ret.x;
+ }
+ if (tint_discarded) {
+ discard;
+ }
+}
+)";
+
+ auto got = Run<DemoteToHelper>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
// Test that no additional discards are inserted when the function unconditionally returns in a
// nested block.
TEST_F(DemoteToHelperTest, EntryPointReturn_NestedInBlock) {
@@ -667,6 +717,58 @@
EXPECT_EQ(expect, str(got));
}
+// Test that we do not mask writes to invocation-private address spaces via a sugared pointer write
+TEST_F(DemoteToHelperTest, InvocationPrivateWritesViaPointerDot) {
+ auto* src = R"(
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+var<private> vp : vec4<f32>;
+
+@fragment
+fn foo(@location(0) in : f32, @location(1) coord : vec2<f32>) {
+ if (in == 0.0) {
+ discard;
+ }
+ let ret = textureSample(t, s, coord);
+ var vf : f32;
+ vf = ret.x;
+ let p = &vp;
+ p.x = ret.x;
+}
+)";
+
+ auto* expect = R"(
+var<private> tint_discarded = false;
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var s : sampler;
+
+var<private> vp : vec4<f32>;
+
+@fragment
+fn foo(@location(0) in : f32, @location(1) coord : vec2<f32>) {
+ if ((in == 0.0)) {
+ tint_discarded = true;
+ }
+ let ret = textureSample(t, s, coord);
+ var vf : f32;
+ vf = ret.x;
+ let p = &(vp);
+ p.x = ret.x;
+ if (tint_discarded) {
+ discard;
+ }
+}
+)";
+
+ auto got = Run<DemoteToHelper>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(DemoteToHelperTest, TextureStoreInEntryPoint) {
auto* src = R"(
@group(0) @binding(0) var t : texture_2d<f32>;
diff --git a/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc b/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
index 33c68c1..cb5598f 100644
--- a/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
@@ -287,6 +287,74 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessPtrChainsTest, ConstantIndices_ViaPointerIndex) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
+
+fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
+ return *p;
+}
+
+fn b() {
+ let p0 = &U;
+ let p1 = &p0[1];
+ let p2 = &p1[1+1];
+ let p3 = &p2[2*2 - 1];
+ a(10, p3, 20);
+}
+
+fn c(p : ptr<uniform, array<array<array<vec4<i32>, 8>, 8>, 8>>) {
+ let p0 = p;
+ let p1 = &p0[1];
+ let p2 = &p1[1+1];
+ let p3 = &p2[2*2 - 1];
+ a(10, p3, 20);
+}
+
+fn d() {
+ c(&U);
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
+
+alias U_X_X_X = array<u32, 3u>;
+
+fn a_U_X_X_X(pre : i32, p : U_X_X_X, post : i32) -> vec4<i32> {
+ return U[p[0]][p[1]][p[2]];
+}
+
+fn b() {
+ let p0 = &(U);
+ let p1 = &(p0[1]);
+ let p2 = &(p1[(1 + 1)]);
+ let p3 = &(p2[((2 * 2) - 1)]);
+ a_U_X_X_X(10, U_X_X_X(1, 2, 3), 20);
+}
+
+fn c_U() {
+ let p0 = &(U);
+ let p1 = &(U[1]);
+ let p2 = &(U[1][2]);
+ let p3 = &(U[1][2][3]);
+ a_U_X_X_X(10, U_X_X_X(1, 2, 3), 20);
+}
+
+fn d() {
+ c_U();
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(DirectVariableAccessPtrChainsTest, HoistIndices) {
auto* src = R"(
enable chromium_experimental_full_ptr_parameters;
@@ -839,6 +907,46 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessUniformASTest, Param_ptr_vec4i32_Via_array_DynamicRead_ViaPointerDot) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8>;
+
+fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
+ return *p;
+}
+
+fn b() {
+ var I = vec2<i32>(3, 3);
+ let p = &I;
+ a(10, &U[p.x], 20);
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8>;
+
+alias U_X = array<u32, 1u>;
+
+fn a_U_X(pre : i32, p : U_X, post : i32) -> vec4<i32> {
+ return U[p[0]];
+}
+
+fn b() {
+ var I = vec2<i32>(3, 3);
+ let p = &(I);
+ a_U_X(10, U_X(u32(p.x)), 20);
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(DirectVariableAccessUniformASTest, CallChaining) {
auto* src = R"(
enable chromium_experimental_full_ptr_parameters;
@@ -1002,6 +1110,169 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessUniformASTest, CallChaining_ViaPointerDotOrIndex) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+};
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+};
+
+@group(0) @binding(0) var<uniform> U : Outer;
+
+fn f0(p : ptr<uniform, vec4<f32>>) -> f32 {
+ return p.x;
+}
+
+fn f1(p : ptr<uniform, mat3x4<f32>>) -> f32 {
+ var res : f32;
+ {
+ // call f0() with inline usage of p
+ res += f0(&p[1]);
+ }
+ {
+ // call f0() with pointer-let usage of p
+ let p_vec = &p[1];
+ res += f0(p_vec);
+ }
+ {
+ // call f0() with inline usage of U
+ res += f0(&U.arr[2].mat[1]);
+ }
+ {
+ // call f0() with pointer-let usage of U
+ let p_vec = &U.arr[2].mat[1];
+ res += f0(p_vec);
+ }
+ return res;
+}
+
+fn f2(p : ptr<uniform, Inner>) -> f32 {
+ let p_mat = &p.mat;
+ return f1(p_mat);
+}
+
+fn f3(p0 : ptr<uniform, InnerArr>, p1 : ptr<uniform, mat3x4<f32>>) -> f32 {
+ let p0_inner = &(*p0)[3];
+ return f2(p0_inner) + f1(p1);
+}
+
+fn f4(p : ptr<uniform, Outer>) -> f32 {
+ return f3(&p.arr, &U.mat);
+}
+
+fn b() {
+ f4(&U);
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+}
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+}
+
+@group(0) @binding(0) var<uniform> U : Outer;
+
+alias U_mat_X = array<u32, 1u>;
+
+fn f0_U_mat_X(p : U_mat_X) -> f32 {
+ return (&(U.mat[p[0]])).x;
+}
+
+alias U_arr_X_mat_X = array<u32, 2u>;
+
+fn f0_U_arr_X_mat_X(p : U_arr_X_mat_X) -> f32 {
+ return (&(U.arr[p[0]].mat[p[0]])).x;
+}
+
+alias U_arr_X_mat_X_1 = array<u32, 2u>;
+
+fn f0_U_arr_X_mat_X_1(p : U_arr_X_mat_X_1) -> f32 {
+ return (&(U.arr[p[0]].mat[p[1]])).x;
+}
+
+fn f1_U_mat() -> f32 {
+ var res : f32;
+ {
+ res += f0_U_mat_X(U_mat_X(1));
+ }
+ {
+ let p_vec = &(U.mat[1]);
+ res += f0_U_mat_X(U_mat_X(1));
+ }
+ {
+ res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ }
+ {
+ let p_vec = &(U.arr[2].mat[1]);
+ res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ }
+ return res;
+}
+
+alias U_arr_X_mat = array<u32, 1u>;
+
+fn f1_U_arr_X_mat(p : U_arr_X_mat) -> f32 {
+ var res : f32;
+ {
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(p[0u], 1));
+ }
+ {
+ let p_vec = &(U.arr[p[0]].mat[1]);
+ res += f0_U_arr_X_mat_X(U_arr_X_mat_X(p[0u], 1));
+ }
+ {
+ res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ }
+ {
+ let p_vec = &(U.arr[2].mat[1]);
+ res += f0_U_arr_X_mat_X_1(U_arr_X_mat_X_1(2, 1));
+ }
+ return res;
+}
+
+alias U_arr_X = array<u32, 1u>;
+
+fn f2_U_arr_X(p : U_arr_X) -> f32 {
+ let p_mat = &(U.arr[p[0]].mat);
+ return f1_U_arr_X_mat(U_arr_X_mat(p[0u]));
+}
+
+fn f3_U_arr_U_mat() -> f32 {
+ let p0_inner = &(U.arr[3]);
+ return (f2_U_arr_X(U_arr_X(3)) + f1_U_mat());
+}
+
+fn f4_U() -> f32 {
+ return f3_U_arr_U_mat();
+}
+
+fn b() {
+ f4_U();
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace uniform_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -1296,6 +1567,169 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessStorageASTest, CallChaining_ViaPointerDotOrIndex) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+};
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+};
+
+@group(0) @binding(0) var<storage> S : Outer;
+
+fn f0(p : ptr<storage, vec4<f32>>) -> f32 {
+ return p.x;
+}
+
+fn f1(p : ptr<storage, mat3x4<f32>>) -> f32 {
+ var res : f32;
+ {
+ // call f0() with inline usage of p
+ res += f0(&p[1]);
+ }
+ {
+ // call f0() with pointer-let usage of p
+ let p_vec = &p[1];
+ res += f0(p_vec);
+ }
+ {
+ // call f0() with inline usage of S
+ res += f0(&S.arr[2].mat[1]);
+ }
+ {
+ // call f0() with pointer-let usage of S
+ let p_vec = &S.arr[2].mat[1];
+ res += f0(p_vec);
+ }
+ return res;
+}
+
+fn f2(p : ptr<storage, Inner>) -> f32 {
+ let p_mat = &p.mat;
+ return f1(p_mat);
+}
+
+fn f3(p0 : ptr<storage, InnerArr>, p1 : ptr<storage, mat3x4<f32>>) -> f32 {
+ let p0_inner = &p0[3];
+ return f2(p0_inner) + f1(p1);
+}
+
+fn f4(p : ptr<storage, Outer>) -> f32 {
+ return f3(&p.arr, &S.mat);
+}
+
+fn b() {
+ f4(&S);
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+}
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+}
+
+@group(0) @binding(0) var<storage> S : Outer;
+
+alias S_mat_X = array<u32, 1u>;
+
+fn f0_S_mat_X(p : S_mat_X) -> f32 {
+ return (&(S.mat[p[0]])).x;
+}
+
+alias S_arr_X_mat_X = array<u32, 2u>;
+
+fn f0_S_arr_X_mat_X(p : S_arr_X_mat_X) -> f32 {
+ return (&(S.arr[p[0]].mat[p[0]])).x;
+}
+
+alias S_arr_X_mat_X_1 = array<u32, 2u>;
+
+fn f0_S_arr_X_mat_X_1(p : S_arr_X_mat_X_1) -> f32 {
+ return (&(S.arr[p[0]].mat[p[1]])).x;
+}
+
+fn f1_S_mat() -> f32 {
+ var res : f32;
+ {
+ res += f0_S_mat_X(S_mat_X(1));
+ }
+ {
+ let p_vec = &(S.mat[1]);
+ res += f0_S_mat_X(S_mat_X(1));
+ }
+ {
+ res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ }
+ {
+ let p_vec = &(S.arr[2].mat[1]);
+ res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ }
+ return res;
+}
+
+alias S_arr_X_mat = array<u32, 1u>;
+
+fn f1_S_arr_X_mat(p : S_arr_X_mat) -> f32 {
+ var res : f32;
+ {
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(p[0u], 1));
+ }
+ {
+ let p_vec = &(S.arr[p[0]].mat[1]);
+ res += f0_S_arr_X_mat_X(S_arr_X_mat_X(p[0u], 1));
+ }
+ {
+ res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ }
+ {
+ let p_vec = &(S.arr[2].mat[1]);
+ res += f0_S_arr_X_mat_X_1(S_arr_X_mat_X_1(2, 1));
+ }
+ return res;
+}
+
+alias S_arr_X = array<u32, 1u>;
+
+fn f2_S_arr_X(p : S_arr_X) -> f32 {
+ let p_mat = &(S.arr[p[0]].mat);
+ return f1_S_arr_X_mat(S_arr_X_mat(p[0u]));
+}
+
+fn f3_S_arr_S_mat() -> f32 {
+ let p0_inner = &(S.arr[3]);
+ return (f2_S_arr_X(S_arr_X(3)) + f1_S_mat());
+}
+
+fn f4_S() -> f32 {
+ return f3_S_arr_S_mat();
+}
+
+fn b() {
+ f4_S();
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace storage_as_tests
////////////////////////////////////////////////////////////////////////////////
@@ -1540,6 +1974,169 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(DirectVariableAccessWorkgroupASTest, CallChaining_ViaPointerDotOrIndex) {
+ auto* src = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+};
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+};
+
+var<workgroup> W : Outer;
+
+fn f0(p : ptr<workgroup, vec4<f32>>) -> f32 {
+ return p.x;
+}
+
+fn f1(p : ptr<workgroup, mat3x4<f32>>) -> f32 {
+ var res : f32;
+ {
+ // call f0() with inline usage of p
+ res += f0(&p[1]);
+ }
+ {
+ // call f0() with pointer-let usage of p
+ let p_vec = &p[1];
+ res += f0(p_vec);
+ }
+ {
+ // call f0() with inline usage of W
+ res += f0(&W.arr[2].mat[1]);
+ }
+ {
+ // call f0() with pointer-let usage of W
+ let p_vec = &W.arr[2].mat[1];
+ res += f0(p_vec);
+ }
+ return res;
+}
+
+fn f2(p : ptr<workgroup, Inner>) -> f32 {
+ let p_mat = &p.mat;
+ return f1(p_mat);
+}
+
+fn f3(p0 : ptr<workgroup, InnerArr>, p1 : ptr<workgroup, mat3x4<f32>>) -> f32 {
+ let p0_inner = &p0[3];
+ return f2(p0_inner) + f1(p1);
+}
+
+fn f4(p : ptr<workgroup, Outer>) -> f32 {
+ return f3(&p.arr, &W.mat);
+}
+
+fn b() {
+ f4(&W);
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct Inner {
+ mat : mat3x4<f32>,
+}
+
+alias InnerArr = array<Inner, 4>;
+
+struct Outer {
+ arr : InnerArr,
+ mat : mat3x4<f32>,
+}
+
+var<workgroup> W : Outer;
+
+alias W_mat_X = array<u32, 1u>;
+
+fn f0_W_mat_X(p : W_mat_X) -> f32 {
+ return (&(W.mat[p[0]])).x;
+}
+
+alias W_arr_X_mat_X = array<u32, 2u>;
+
+fn f0_W_arr_X_mat_X(p : W_arr_X_mat_X) -> f32 {
+ return (&(W.arr[p[0]].mat[p[0]])).x;
+}
+
+alias W_arr_X_mat_X_1 = array<u32, 2u>;
+
+fn f0_W_arr_X_mat_X_1(p : W_arr_X_mat_X_1) -> f32 {
+ return (&(W.arr[p[0]].mat[p[1]])).x;
+}
+
+fn f1_W_mat() -> f32 {
+ var res : f32;
+ {
+ res += f0_W_mat_X(W_mat_X(1));
+ }
+ {
+ let p_vec = &(W.mat[1]);
+ res += f0_W_mat_X(W_mat_X(1));
+ }
+ {
+ res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ }
+ {
+ let p_vec = &(W.arr[2].mat[1]);
+ res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ }
+ return res;
+}
+
+alias W_arr_X_mat = array<u32, 1u>;
+
+fn f1_W_arr_X_mat(p : W_arr_X_mat) -> f32 {
+ var res : f32;
+ {
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(p[0u], 1));
+ }
+ {
+ let p_vec = &(W.arr[p[0]].mat[1]);
+ res += f0_W_arr_X_mat_X(W_arr_X_mat_X(p[0u], 1));
+ }
+ {
+ res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ }
+ {
+ let p_vec = &(W.arr[2].mat[1]);
+ res += f0_W_arr_X_mat_X_1(W_arr_X_mat_X_1(2, 1));
+ }
+ return res;
+}
+
+alias W_arr_X = array<u32, 1u>;
+
+fn f2_W_arr_X(p : W_arr_X) -> f32 {
+ let p_mat = &(W.arr[p[0]].mat);
+ return f1_W_arr_X_mat(W_arr_X_mat(p[0u]));
+}
+
+fn f3_W_arr_W_mat() -> f32 {
+ let p0_inner = &(W.arr[3]);
+ return (f2_W_arr_X(W_arr_X(3)) + f1_W_mat());
+}
+
+fn f4_W() -> f32 {
+ return f3_W_arr_W_mat();
+}
+
+fn b() {
+ f4_W();
+}
+)";
+
+ auto got = Run<DirectVariableAccess>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace workgroup_as_tests
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/wgsl/ast/transform/expand_compound_assignment.cc b/src/tint/lang/wgsl/ast/transform/expand_compound_assignment.cc
index 06da9de..d4c66c1 100644
--- a/src/tint/lang/wgsl/ast/transform/expand_compound_assignment.cc
+++ b/src/tint/lang/wgsl/ast/transform/expand_compound_assignment.cc
@@ -85,8 +85,10 @@
// Helper function to create a variable that is a pointer to `expr`.
auto hoist_pointer_to = [&](const Expression* expr) {
+ // Lhs may already be a pointer, in which case we don't take it's address
+ bool is_pointer = ctx.src->Sem().GetVal(expr)->Type()->Is<core::type::Pointer>();
auto name = b.Sym();
- auto* ptr = b.AddressOf(ctx.Clone(expr));
+ auto* ptr = is_pointer ? ctx.Clone(expr) : b.AddressOf(ctx.Clone(expr));
auto* decl = b.Decl(b.Let(name, ptr));
hoist_to_decl_before.InsertBefore(ctx.src->Sem().Get(stmt), decl);
return name;
@@ -103,7 +105,7 @@
// Helper function that returns `true` if the type of `expr` is a vector.
auto is_vec = [&](const Expression* expr) {
if (auto* val_expr = ctx.src->Sem().GetVal(expr)) {
- return val_expr->Type()->UnwrapRef()->Is<core::type::Vector>();
+ return val_expr->Type()->UnwrapPtrOrRef()->Is<core::type::Vector>();
}
return false;
};
@@ -116,9 +118,11 @@
auto* member_accessor = lhs->As<MemberAccessorExpression>();
if (lhs->Is<IdentifierExpression>() ||
(member_accessor && member_accessor->object->Is<IdentifierExpression>())) {
- // This is the simple case with no side effects, so we can just use the
- // original LHS expression directly.
- // Before:
+ // TODO(crbug.com/tint/2115): This branch should also handle (recursive) deref'd
+ // identifiers (e.g. (*p).bar += rhs)).
+
+ // This is the simple case with no side effects, so we can just use
+ // the original LHS expression directly. Before:
// foo.bar += rhs;
// After:
// foo.bar = foo.bar + rhs;
diff --git a/src/tint/lang/wgsl/ast/transform/expand_compound_assignment_test.cc b/src/tint/lang/wgsl/ast/transform/expand_compound_assignment_test.cc
index 83d154c..8ffc045 100644
--- a/src/tint/lang/wgsl/ast/transform/expand_compound_assignment_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/expand_compound_assignment_test.cc
@@ -222,6 +222,122 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(ExpandCompoundAssignmentTest, LhsArrayOfVectorComponent_MemberAccessor_ViaArrayIndex) {
+ auto* src = R"(
+fn main() {
+ var v : array<vec4<i32>, 3>;
+ v[0].y += 1;
+}
+)";
+
+ auto* expect = R"(
+fn main() {
+ var v : array<vec4<i32>, 3>;
+ let tint_symbol = &(v[0]);
+ (*(tint_symbol)).y = ((*(tint_symbol)).y + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ExpandCompoundAssignmentTest, LhsVectorComponent_MemberAccessor_ViaDerefPointerDot) {
+ auto* src = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &v;
+ (*p).y += 1;
+}
+)";
+
+ // TODO(crbug.com/tint/2115): we currently needlessly hoist pointer-deref to another pointer.
+ auto* expect = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &(v);
+ let tint_symbol = &(*(p));
+ (*(tint_symbol)).y = ((*(tint_symbol)).y + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ExpandCompoundAssignmentTest, LhsVectorComponent_MemberAccessor_ViaPointerDot) {
+ auto* src = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &v;
+ p.y += 1;
+}
+)";
+
+ auto* expect = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &(v);
+ p.y = (p.y + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ExpandCompoundAssignmentTest, LhsVectorComponent_MemberAccessor_ViaDerefPointerIndex) {
+ auto* src = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &v;
+ (*p)[0] += 1;
+}
+)";
+
+ // TODO(crbug.com/tint/2115): we currently needlessly hoist pointer-deref to another pointer.
+ auto* expect = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &(v);
+ let tint_symbol = &(*(p));
+ let tint_symbol_1 = 0;
+ (*(tint_symbol))[tint_symbol_1] = ((*(tint_symbol))[tint_symbol_1] + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ExpandCompoundAssignmentTest, LhsVectorComponent_MemberAccessor_ViaPointerIndex) {
+ auto* src = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &v;
+ p[0] += 1;
+}
+)";
+
+ auto* expect = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &(v);
+ let tint_symbol = p;
+ let tint_symbol_1 = 0;
+ (*(tint_symbol))[tint_symbol_1] = ((*(tint_symbol))[tint_symbol_1] + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(ExpandCompoundAssignmentTest, LhsMatrixColumn) {
auto* src = R"(
var<private> m : mat4x4<f32>;
@@ -678,6 +794,79 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(ExpandCompoundAssignmentTest,
+ Increment_LhsVectorComponent_ArrayAccessor_ViaDerefPointerIndex) {
+ auto* src = R"(
+var<private> v : vec4<i32>;
+
+fn idx() -> i32 {
+ v.y = 42;
+ return 1;
+}
+
+fn main() {
+ let p = &v;
+ (*p)[idx()]++;
+}
+)";
+
+ auto* expect = R"(
+var<private> v : vec4<i32>;
+
+fn idx() -> i32 {
+ v.y = 42;
+ return 1;
+}
+
+fn main() {
+ let p = &(v);
+ let tint_symbol = &(*(p));
+ let tint_symbol_1 = idx();
+ (*(tint_symbol))[tint_symbol_1] = ((*(tint_symbol))[tint_symbol_1] + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ExpandCompoundAssignmentTest, Increment_LhsVectorComponent_ArrayAccessor_ViaPointerIndex) {
+ auto* src = R"(
+var<private> v : vec4<i32>;
+
+fn idx() -> i32 {
+ v.y = 42;
+ return 1;
+}
+
+fn main() {
+ let p = &v;
+ p[idx()]++;
+}
+)";
+
+ auto* expect = R"(
+var<private> v : vec4<i32>;
+
+fn idx() -> i32 {
+ v.y = 42;
+ return 1;
+}
+
+fn main() {
+ let p = &(v);
+ let tint_symbol = p;
+ let tint_symbol_1 = idx();
+ (*(tint_symbol))[tint_symbol_1] = ((*(tint_symbol))[tint_symbol_1] + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(ExpandCompoundAssignmentTest, Increment_LhsVectorComponent_MemberAccessor) {
auto* src = R"(
fn main() {
@@ -698,6 +887,53 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(ExpandCompoundAssignmentTest,
+ Increment_LhsVectorComponent_MemberAccessor_ViaDerefPointerDot) {
+ auto* src = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &v;
+ (*p).y++;
+}
+)";
+
+ // TODO(crbug.com/tint/2115): we currently needlessly hoist pointer-deref to another pointer.
+ auto* expect = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &(v);
+ let tint_symbol = &(*(p));
+ (*(tint_symbol)).y = ((*(tint_symbol)).y + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ExpandCompoundAssignmentTest, Increment_LhsVectorComponent_MemberAccessor_ViaPointerDot) {
+ auto* src = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &v;
+ p.y++;
+}
+)";
+
+ auto* expect = R"(
+fn main() {
+ var v : vec4<i32>;
+ let p = &(v);
+ p.y = (p.y + 1);
+}
+)";
+
+ auto got = Run<ExpandCompoundAssignment>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(ExpandCompoundAssignmentTest, Increment_ForLoopCont) {
auto* src = R"(
var<private> a : array<vec4<i32>, 4>;
diff --git a/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc b/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc
index c5f43f3..469448b 100644
--- a/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc
@@ -150,6 +150,67 @@
EXPECT_EQ(expect, str(got));
}
+// Same should happen via a sugared pointer write
+TEST_F(PreservePaddingTest, StructTrailingPadding_ViaPointerDot) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : u32,
+ c : u32,
+ d : u32,
+ e : vec3<u32>,
+}
+
+struct Outer {
+ s : S,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : Outer;
+
+@compute @workgroup_size(1)
+fn foo() {
+ let p = &v;
+ p.s = S();
+}
+)";
+
+ auto* expect = R"(
+enable chromium_experimental_full_ptr_parameters;
+
+struct S {
+ a : u32,
+ b : u32,
+ c : u32,
+ d : u32,
+ e : vec3<u32>,
+}
+
+struct Outer {
+ s : S,
+}
+
+@group(0) @binding(0) var<storage, read_write> v : Outer;
+
+fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) {
+ (*(dest)).a = value.a;
+ (*(dest)).b = value.b;
+ (*(dest)).c = value.c;
+ (*(dest)).d = value.d;
+ (*(dest)).e = value.e;
+}
+
+@compute @workgroup_size(1)
+fn foo() {
+ let p = &(v);
+ assign_and_preserve_padding(&(p.s), S());
+}
+)";
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PreservePaddingTest, StructInternalPadding) {
auto* src = R"(
struct S {
@@ -745,6 +806,34 @@
EXPECT_EQ(expect, str(got));
}
+// Same should happen via a sugared pointer write.
+TEST_F(PreservePaddingTest, NoModify_Workgroup_ViaPointerDot) {
+ auto* src = R"(
+struct S {
+ a : u32,
+ b : vec3<u32>,
+}
+
+struct Outer {
+ s : S,
+}
+
+var<workgroup> v : Outer;
+
+@compute @workgroup_size(1)
+fn foo() {
+ let p = &(v);
+ p.s = S();
+}
+)";
+
+ auto* expect = src;
+
+ auto got = Run<PreservePadding>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PreservePaddingTest, NoModify_Private) {
auto* src = R"(
struct S {
diff --git a/src/tint/lang/wgsl/ast/transform/promote_side_effects_to_decl_test.cc b/src/tint/lang/wgsl/ast/transform/promote_side_effects_to_decl_test.cc
index ea506c0..f3238ee 100644
--- a/src/tint/lang/wgsl/ast/transform/promote_side_effects_to_decl_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/promote_side_effects_to_decl_test.cc
@@ -2725,6 +2725,107 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PromoteSideEffectsToDeclTest, IndexAccessor_2D_LeftSE_ViaPointerIndex) {
+ auto* src = R"(
+fn a(i : i32) -> i32 {
+ return 1;
+}
+
+fn f() {
+ var b = array<array<i32, 10>, 10>();
+ let p = &b;
+ var c = 1;
+ var r = p[a(0)][c];
+}
+)";
+
+ auto* expect = R"(
+fn a(i : i32) -> i32 {
+ return 1;
+}
+
+fn f() {
+ var b = array<array<i32, 10>, 10>();
+ let p = &(b);
+ var c = 1;
+ let tint_symbol : i32 = a(0);
+ var r = p[tint_symbol][c];
+}
+)";
+
+ auto got = Run<PromoteSideEffectsToDecl>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PromoteSideEffectsToDeclTest, IndexAccessor_2D_RightSE_ViaPointerIndex) {
+ auto* src = R"(
+fn a(i : i32) -> i32 {
+ return 1;
+}
+
+fn f() {
+ var b = array<array<i32, 10>, 10>();
+ let p = &b;
+ var c = 1;
+ let tint_symbol = c;
+ let tint_symbol_1 = a(0);
+ var r = p[tint_symbol][tint_symbol_1];
+}
+)";
+
+ auto* expect = R"(
+fn a(i : i32) -> i32 {
+ return 1;
+}
+
+fn f() {
+ var b = array<array<i32, 10>, 10>();
+ let p = &(b);
+ var c = 1;
+ let tint_symbol = c;
+ let tint_symbol_1 = a(0);
+ var r = p[tint_symbol][tint_symbol_1];
+}
+)";
+
+ auto got = Run<PromoteSideEffectsToDecl>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PromoteSideEffectsToDeclTest, IndexAccessor_2D_BothSE_ViaPointerIndex) {
+ auto* src = R"(
+fn a(i : i32) -> i32 {
+ return 1;
+}
+
+fn f() {
+ var b = array<array<i32, 10>, 10>();
+ let p = &b;
+ var r = p[a(0)][a(1)];
+}
+)";
+
+ auto* expect = R"(
+fn a(i : i32) -> i32 {
+ return 1;
+}
+
+fn f() {
+ var b = array<array<i32, 10>, 10>();
+ let p = &(b);
+ let tint_symbol : i32 = a(0);
+ let tint_symbol_1 : i32 = a(1);
+ var r = p[tint_symbol][tint_symbol_1];
+}
+)";
+
+ auto got = Run<PromoteSideEffectsToDecl>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PromoteSideEffectsToDeclTest, Assignment_ToPhony) {
auto* src = R"(
fn a(i : i32) -> i32 {
@@ -3643,6 +3744,43 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(PromoteSideEffectsToDeclTest, BinaryMemberAccessorPlusSE_ViaPointerDot) {
+ // bclayton@'s example:
+ // https://dawn-review.googlesource.com/c/tint/+/78620/6..8/src/transform/promote_side_effects_to_decl.cc#b490
+ auto* src = R"(
+fn modify_vec(p : ptr<function, vec4<i32>>) -> i32 {
+ (*p).x = 42;
+ return 0;
+}
+
+fn f() {
+ var v = vec4<i32>();
+ let p = &v;
+ let l = p.x + modify_vec(&v);
+ // l should be 0, not 42
+}
+)";
+
+ auto* expect = R"(
+fn modify_vec(p : ptr<function, vec4<i32>>) -> i32 {
+ (*(p)).x = 42;
+ return 0;
+}
+
+fn f() {
+ var v = vec4<i32>();
+ let p = &(v);
+ let tint_symbol : i32 = p.x;
+ let tint_symbol_1 : i32 = modify_vec(&(v));
+ let l = (tint_symbol + tint_symbol_1);
+}
+)";
+
+ auto got = Run<PromoteSideEffectsToDecl>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(PromoteSideEffectsToDeclTest, Call_ReadOnlyArgAndSE) {
// Make sure that read-only args don't get hoisted (tex and samp)
auto* src = R"(
diff --git a/src/tint/lang/wgsl/ast/transform/renamer.cc b/src/tint/lang/wgsl/ast/transform/renamer.cc
index 0183f76..7233498 100644
--- a/src/tint/lang/wgsl/ast/transform/renamer.cc
+++ b/src/tint/lang/wgsl/ast/transform/renamer.cc
@@ -1298,7 +1298,7 @@
if (sem->Is<sem::Swizzle>()) {
preserved_identifiers.Add(accessor->member);
} else if (auto* str_expr = src.Sem().GetVal(accessor->object)) {
- if (auto* ty = str_expr->Type()->UnwrapRef()->As<core::type::Struct>()) {
+ if (auto* ty = str_expr->Type()->UnwrapPtrOrRef()->As<core::type::Struct>()) {
if (!ty->Is<sem::Struct>()) { // Builtin structure
preserved_identifiers.Add(accessor->member);
}
diff --git a/src/tint/lang/wgsl/ast/transform/renamer_test.cc b/src/tint/lang/wgsl/ast/transform/renamer_test.cc
index fca0b22..b24a440 100644
--- a/src/tint/lang/wgsl/ast/transform/renamer_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/renamer_test.cc
@@ -274,6 +274,51 @@
EXPECT_THAT(data->remappings, ContainerEq(expected_remappings));
}
+TEST_F(RenamerTest, PreserveBuiltinTypes_ViaPointerDot) {
+ auto* src = R"(
+@compute @workgroup_size(1)
+fn entry() {
+ var m = modf(1.0);
+ let p1 = &m;
+ var f = frexp(1.0);
+ let p2 = &f;
+
+ var a = p1.whole;
+ var b = p1.fract;
+ var c = p2.fract;
+ var d = p2.exp;
+}
+)";
+
+ auto* expect = R"(
+@compute @workgroup_size(1)
+fn tint_symbol() {
+ var tint_symbol_1 = modf(1.0);
+ let tint_symbol_2 = &(tint_symbol_1);
+ var tint_symbol_3 = frexp(1.0);
+ let tint_symbol_4 = &(tint_symbol_3);
+ var tint_symbol_5 = tint_symbol_2.whole;
+ var tint_symbol_6 = tint_symbol_2.fract;
+ var tint_symbol_7 = tint_symbol_4.fract;
+ var tint_symbol_8 = tint_symbol_4.exp;
+}
+)";
+
+ auto got = Run<Renamer>(src);
+
+ EXPECT_EQ(expect, str(got));
+
+ auto* data = got.data.Get<Renamer::Data>();
+
+ ASSERT_NE(data, nullptr);
+ Renamer::Remappings expected_remappings = {
+ {"entry", "tint_symbol"}, {"m", "tint_symbol_1"}, {"p1", "tint_symbol_2"},
+ {"f", "tint_symbol_3"}, {"p2", "tint_symbol_4"}, {"a", "tint_symbol_5"},
+ {"b", "tint_symbol_6"}, {"c", "tint_symbol_7"}, {"d", "tint_symbol_8"},
+ };
+ EXPECT_THAT(data->remappings, ContainerEq(expected_remappings));
+}
+
TEST_F(RenamerTest, PreserveCoreDiagnosticRuleName) {
auto* src = R"(
diagnostic(off, chromium.unreachable_code);
diff --git a/src/tint/lang/wgsl/ast/transform/robustness.cc b/src/tint/lang/wgsl/ast/transform/robustness.cc
index 0bc1575..8169d28 100644
--- a/src/tint/lang/wgsl/ast/transform/robustness.cc
+++ b/src/tint/lang/wgsl/ast/transform/robustness.cc
@@ -31,6 +31,7 @@
#include <limits>
#include <utility>
+#include "src/tint/lang/core/type/memory_view.h"
#include "src/tint/lang/core/type/reference.h"
#include "src/tint/lang/wgsl/ast/transform/hoist_to_decl_before.h"
#include "src/tint/lang/wgsl/program/clone_context.h"
@@ -194,7 +195,7 @@
if (auto pred = predicates.Get(expr)) {
// Expression is predicated
auto* sem_expr = sem.GetVal(expr);
- if (!sem_expr->Type()->IsAnyOf<core::type::Reference, core::type::Pointer>()) {
+ if (!sem_expr->Type()->Is<core::type::MemoryView>()) {
auto pred_load = b.Symbols().New("predicated_expr");
auto ty = CreateASTTypeFor(ctx, sem_expr->Type());
hoist.InsertBefore(sem_expr->Stmt(), b.Decl(b.Var(pred_load, ty)));
@@ -236,7 +237,7 @@
const Expression* DynamicLimitFor(const sem::IndexAccessorExpression* expr) {
auto* obj_type = expr->Object()->Type();
return Switch(
- obj_type->UnwrapRef(), //
+ obj_type->UnwrapPtrOrRef(), //
[&](const core::type::Vector* vec) -> const Expression* {
if (expr->Index()->ConstantValue() || expr->Index()->Is<sem::Swizzle>()) {
// Index and size is constant.
@@ -708,7 +709,7 @@
/// @returns true if expr is an IndexAccessorExpression whose object is a runtime-sized array.
bool IsIndexAccessingRuntimeSizedArray(const sem::IndexAccessorExpression* expr) {
- auto* array_type = expr->Object()->Type()->UnwrapRef()->As<core::type::Array>();
+ auto* array_type = expr->Object()->Type()->UnwrapPtrOrRef()->As<core::type::Array>();
return array_type != nullptr && array_type->Count()->Is<core::type::RuntimeArrayCount>();
}
diff --git a/src/tint/lang/wgsl/ast/transform/robustness_test.cc b/src/tint/lang/wgsl/ast/transform/robustness_test.cc
index f1e733b..647e669 100644
--- a/src/tint/lang/wgsl/ast/transform/robustness_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/robustness_test.cc
@@ -464,6 +464,67 @@
EXPECT_EQ(expect, str(got));
}
+TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeArrayIndexViaPointerIndex) {
+ auto* src = R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let p1 = &(a);
+ let p2 = &(b);
+ var c : f32 = p1[p2[i]];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let p1 = &(a);
+ let p2 = &(b);
+ var c : f32 = p1[min(u32(p2[min(i, 4u)]), 2u)];
+}
+)",
+ /* predicate */ R"(
+var<private> a : array<f32, 3>;
+
+var<private> b : array<i32, 5>;
+
+var<private> i : u32;
+
+fn f() {
+ let p1 = &(a);
+ let p2 = &(b);
+ let index = i;
+ let predicate = (u32(index) <= 4u);
+ var predicated_expr : i32;
+ if (predicate) {
+ predicated_expr = p2[index];
+ }
+ let index_1 = predicated_expr;
+ let predicate_1 = (u32(index_1) <= 2u);
+ var predicated_expr_1 : f32;
+ if (predicate_1) {
+ predicated_expr_1 = p1[index_1];
+ }
+ var c : f32 = predicated_expr_1;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeExpression) {
auto* src = R"(
var<private> a : array<f32, 3>;
@@ -1379,6 +1440,52 @@
EXPECT_EQ(expect, str(got));
}
+TEST_P(RobustnessTest, Read_Vector_IndexWithRuntimeExpression_ViaPointerIndex) {
+ auto* src = R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let p = &(a);
+ var b : f32 = p[((c + 2) - 3)];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let p = &(a);
+ var b : f32 = p[min(u32(((c + 2) - 3)), 2u)];
+}
+)",
+ /* predicate */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let p = &(a);
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 2u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = p[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithGlobalVar) {
auto* src = R"(
var<private> a : vec3<f32>;
@@ -1465,6 +1572,52 @@
EXPECT_EQ(expect, str(got));
}
+TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithRuntimeExpression_ViaPointerDot) {
+ auto* src = R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let p = &(a);
+ var b : f32 = p.xy[((c + 2) - 3)];
+}
+)";
+
+ auto* expect = Expect(GetParam(),
+ /* ignore */ src,
+ /* clamp */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let p = &(a);
+ var b : f32 = p.xy[min(u32(((c + 2) - 3)), 1u)];
+}
+)",
+ /* predicate */ R"(
+var<private> a : vec3<f32>;
+
+var<private> c : i32;
+
+fn f() {
+ let p = &(a);
+ let index = ((c + 2) - 3);
+ let predicate = (u32(index) <= 1u);
+ var predicated_expr : f32;
+ if (predicate) {
+ predicated_expr = p.xy[index];
+ }
+ var b : f32 = predicated_expr;
+}
+)");
+
+ auto got = Run<Robustness>(src, Config(GetParam()));
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_P(RobustnessTest, Read_Vector_IndexWithOverride) {
auto* src = R"(
@id(1300) override idx : i32;
@@ -5571,6 +5724,30 @@
EXPECT_EQ(expect, str(got));
}
+TEST_P(RobustnessTest, Read_disable_unsized_array_index_clamping_abstract_int_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read> s : array<f32>;
+
+fn f() {
+ let p = &(s);
+ var d : f32 = p[25];
+}
+)";
+
+ auto* expect = R"(
+@group(0) @binding(0) var<storage, read> s : array<f32>;
+
+fn f() {
+ let p = &(s);
+ var d : f32 = p[u32(25)];
+}
+)";
+
+ auto got = Run<Robustness>(src, Config(GetParam(), true));
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_P(RobustnessTest, Assign_disable_unsized_array_index_clamping_i32) {
auto* src = R"(
@group(0) @binding(0) var<storage, read_write> s : array<f32>;
@@ -5634,6 +5811,30 @@
EXPECT_EQ(expect, str(got));
}
+TEST_P(RobustnessTest, Assign_disable_unsized_array_index_clamping_abstract_int_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> s : array<f32>;
+
+fn f() {
+ let p = &(s);
+ p[25] = 0.5f;
+}
+)";
+
+ auto* expect = R"(
+@group(0) @binding(0) var<storage, read_write> s : array<f32>;
+
+fn f() {
+ let p = &(s);
+ p[u32(25)] = 0.5f;
+}
+)";
+
+ auto got = Run<Robustness>(src, Config(GetParam(), true));
+
+ EXPECT_EQ(expect, str(got));
+}
+
INSTANTIATE_TEST_SUITE_P(,
RobustnessTest,
testing::Values(Robustness::Action::kIgnore,
diff --git a/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc b/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc
index 771b954..3095252 100644
--- a/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc
+++ b/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc
@@ -27,15 +27,14 @@
#include "src/tint/lang/wgsl/ast/transform/simplify_pointers.h"
-#include <memory>
-#include <unordered_map>
-#include <utility>
-#include <vector>
+#include <unordered_set>
+#include "src/tint/utils/containers/hashset.h"
#include "src/tint/lang/wgsl/ast/transform/unshadow.h"
#include "src/tint/lang/wgsl/program/clone_context.h"
#include "src/tint/lang/wgsl/program/program_builder.h"
#include "src/tint/lang/wgsl/resolver/resolve.h"
+#include "src/tint/lang/wgsl/sem/accessor_expression.h"
#include "src/tint/lang/wgsl/sem/block_statement.h"
#include "src/tint/lang/wgsl/sem/function.h"
#include "src/tint/lang/wgsl/sem/statement.h"
@@ -69,6 +68,9 @@
ProgramBuilder b;
/// The clone context
program::CloneContext ctx = {&b, &src, /* auto_clone_symbols */ true};
+ /// Set of accessor expression objects that are pointers, used to handle
+ /// pointer-index/dot sugar syntax.
+ Hashset<const Expression*, 4> is_accessor_object_pointer;
/// Constructor
/// @param program the source program
@@ -110,9 +112,14 @@
/// indirection ops into a PointerOp.
/// @param in the expression to walk
/// @returns the reduced PointerOp
- PointerOp Reduce(const Expression* in) const {
+ PointerOp Reduce(const Expression* in) {
PointerOp op{0, in};
while (true) {
+ if (is_accessor_object_pointer.Contains(op.expr)) {
+ // Object is an implicitly dereferenced pointer (i.e. syntax sugar).
+ op.indirections++;
+ }
+
if (auto* unary = op.expr->As<UnaryOpExpression>()) {
switch (unary->op) {
case core::UnaryOp::kIndirection:
@@ -226,6 +233,14 @@
// will be no pointers that can be inlined.
needs_transform = true;
}
+ },
+ [&](const AccessorExpression* accessor) {
+ if (auto* a = ctx.src->Sem().Get<sem::ValueExpression>(accessor->object)) {
+ if (a->Type()->Is<core::type::Pointer>()) {
+ // Object is an implicitly dereferenced pointer (i.e. syntax sugar).
+ is_accessor_object_pointer.Add(accessor->object);
+ }
+ }
});
}
diff --git a/src/tint/lang/wgsl/ast/transform/simplify_pointers_test.cc b/src/tint/lang/wgsl/ast/transform/simplify_pointers_test.cc
index e2c7efd..1c92b18 100644
--- a/src/tint/lang/wgsl/ast/transform/simplify_pointers_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/simplify_pointers_test.cc
@@ -144,6 +144,92 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(SimplifyPointersTest, PointerDerefIndex) {
+ auto* src = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let p = &a;
+ let v = (*p)[1];
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let v = a[1];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(SimplifyPointersTest, PointerIndex) {
+ auto* src = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let p = &a;
+ let v = p[1];
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let v = a[1];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(SimplifyPointersTest, SimpleChain) {
+ auto* src = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let ap : ptr<function, array<f32, 2>> = &a;
+ let vp : ptr<function, f32> = &(*ap)[1];
+ let v : f32 = *vp;
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let v : f32 = a[1];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(SimplifyPointersTest, SimpleChain_ViaPointerIndex) {
+ auto* src = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let ap : ptr<function, array<f32, 2>> = &a;
+ let vp : ptr<function, f32> = &ap[1];
+ let v : f32 = *vp;
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var a : array<f32, 2>;
+ let v : f32 = a[1];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(SimplifyPointersTest, ComplexChain) {
auto* src = R"(
fn f() {
@@ -167,6 +253,29 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(SimplifyPointersTest, ComplexChain_ViaPointerIndex) {
+ auto* src = R"(
+fn f() {
+ var a : array<mat4x4<f32>, 4>;
+ let ap : ptr<function, array<mat4x4<f32>, 4>> = &a;
+ let mp : ptr<function, mat4x4<f32>> = &ap[3];
+ let vp : ptr<function, vec4<f32>> = &mp[2];
+ let v : vec4<f32> = *vp;
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var a : array<mat4x4<f32>, 4>;
+ let v : vec4<f32> = a[3][2];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(SimplifyPointersTest, SavedVars) {
auto* src = R"(
struct S {
diff --git a/src/tint/lang/wgsl/ast/transform/std140.cc b/src/tint/lang/wgsl/ast/transform/std140.cc
index d933ec6..b28e5a7 100644
--- a/src/tint/lang/wgsl/ast/transform/std140.cc
+++ b/src/tint/lang/wgsl/ast/transform/std140.cc
@@ -551,7 +551,7 @@
expr = a->Object();
// Is the object a std140 decomposed matrix?
- if (auto* mat = expr->Type()->UnwrapRef()->As<core::type::Matrix>()) {
+ if (auto* mat = expr->Type()->UnwrapPtrOrRef()->As<core::type::Matrix>()) {
if (std140_mats.Contains(mat)) {
// Record this on the access.
access.std140_mat_idx = access.indices.Length();
diff --git a/src/tint/lang/wgsl/ast/transform/std140_exhaustive_test.cc b/src/tint/lang/wgsl/ast/transform/std140_exhaustive_test.cc
index 2e9c0a1..f9dca89 100644
--- a/src/tint/lang/wgsl/ast/transform/std140_exhaustive_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/std140_exhaustive_test.cc
@@ -2458,6 +2458,217 @@
EXPECT_EQ(expect, str(got));
}
+TEST_P(Std140Test_Matrix, ArrayStructArrayStructMatUniform_LoadsViaImplicitDerefPtrs) {
+ auto matrix = GetParam();
+
+ std::string src = R"(
+enable f16;
+
+struct Inner {
+ @size(64)
+ m : ${mat},
+}
+
+struct Outer {
+ a : array<Inner, 4>,
+}
+
+@group(0) @binding(0) var<uniform> a : array<Outer, 4>;
+
+fn f() {
+ let I = 1;
+ let J = 2;
+ let K = 0;
+ let p_a = &(a);
+ let p_a_3 = &(p_a[3]);
+ let p_a_I = &(p_a[I]);
+ let p_a_3_a = &((*(p_a_3)).a);
+ let p_a_I_a = &((*(p_a_I)).a);
+ let p_a_3_a_2 = &(p_a_3_a[2]);
+ let p_a_3_a_I = &(p_a_3_a[I]);
+ let p_a_I_a_2 = &(p_a_I_a[2]);
+ let p_a_I_a_J = &(p_a_I_a[J]);
+ let p_a_3_a_2_m = &((*(p_a_3_a_2)).m);
+ let p_a_3_a_I_m = &((*(p_a_3_a_I)).m);
+ let p_a_I_a_2_m = &((*(p_a_I_a_2)).m);
+ let p_a_I_a_J_m = &((*(p_a_I_a_J)).m);
+ let p_a_3_a_2_m_1 = &(p_a_3_a_2_m[1]);
+ let p_a_I_a_J_m_K = &(p_a_I_a_J_m[K]);
+ let l_a : array<Outer, 4> = *(p_a);
+ let l_a_3 : Outer = *(p_a_3);
+ let l_a_I : Outer = *(p_a_I);
+ let l_a_3_a : array<Inner, 4> = *(p_a_3_a);
+ let l_a_I_a : array<Inner, 4> = *(p_a_I_a);
+ let l_a_3_a_2 : Inner = *(p_a_3_a_2);
+ let l_a_3_a_I : Inner = *(p_a_3_a_I);
+ let l_a_I_a_2 : Inner = *(p_a_I_a_2);
+ let l_a_I_a_J : Inner = *(p_a_I_a_J);
+ let l_a_3_a_2_m : ${mat} = *(p_a_3_a_2_m);
+ let l_a_3_a_I_m : ${mat} = *(p_a_3_a_I_m);
+ let l_a_I_a_2_m : ${mat} = *(p_a_I_a_2_m);
+ let l_a_I_a_J_m : ${mat} = *(p_a_I_a_J_m);
+ let l_a_3_a_2_m_1 : ${col_vector_type} = *(p_a_3_a_2_m_1);
+ let l_a_I_a_J_m_K : ${col_vector_type} = *(p_a_I_a_J_m_K);
+ let l_a_2_a_0_m_1_0 : ${elem_type} = p_a_3_a_2_m_1[0];
+ let l_a_I_a_J_m_K_I : ${elem_type} = p_a_I_a_J_m_K[I];
+}
+)";
+ src = matrix.ReplaceFieldsInString(src);
+
+ std::string expect;
+ if (matrix.NotStd140Compatible()) {
+ expect = R"(
+enable f16;
+
+struct Inner {
+ @size(64)
+ m : ${mat},
+}
+
+struct Inner_std140 {
+${col_vectors}
+}
+
+struct Outer {
+ a : array<Inner, 4>,
+}
+
+struct Outer_std140 {
+ a : array<Inner_std140, 4u>,
+}
+
+@group(0) @binding(0) var<uniform> a : array<Outer_std140, 4u>;
+
+fn conv_Inner(val : Inner_std140) -> Inner {
+ return Inner(${mat}(${col_vectors_inline_conv_Inner}));
+}
+
+fn conv_arr4_Inner(val : array<Inner_std140, 4u>) -> array<Inner, 4u> {
+ var arr : array<Inner, 4u>;
+ for(var i : u32; (i < 4u); i = (i + 1)) {
+ arr[i] = conv_Inner(val[i]);
+ }
+ return arr;
+}
+
+fn conv_Outer(val : Outer_std140) -> Outer {
+ return Outer(conv_arr4_Inner(val.a));
+}
+
+fn conv_arr4_Outer(val : array<Outer_std140, 4u>) -> array<Outer, 4u> {
+ var arr : array<Outer, 4u>;
+ for(var i : u32; (i < 4u); i = (i + 1)) {
+ arr[i] = conv_Outer(val[i]);
+ }
+ return arr;
+}
+
+fn load_a_3_a_2_m() -> ${mat} {
+ let s = &(a[3u].a[2u]);
+ return ${mat}(${col_vectors_inline_load_matrix});
+}
+
+fn load_a_3_a_p0_m(p0 : u32) -> ${mat} {
+ let s = &(a[3u].a[p0]);
+ return ${mat}(${col_vectors_inline_load_matrix});
+}
+
+fn load_a_p0_a_2_m(p0 : u32) -> ${mat} {
+ let s = &(a[p0].a[2u]);
+ return ${mat}(${col_vectors_inline_load_matrix});
+}
+
+fn load_a_p0_a_p1_m(p0 : u32, p1 : u32) -> ${mat} {
+ let s = &(a[p0].a[p1]);
+ return ${mat}(${col_vectors_inline_load_matrix});
+}
+
+fn load_a_p0_a_p1_m_p2(p0 : u32, p1 : u32, p2 : u32) -> ${col_vector_type} {
+ switch(p2) {
+${col_table_load_column}
+ default: {
+ return ${col_vector_type}();
+ }
+ }
+}
+
+fn load_a_p0_a_p1_m_p2_p3(p0 : u32, p1 : u32, p2 : u32, p3 : u32) -> ${elem_type} {
+ switch(p2) {
+${col_table_load_element}
+ default: {
+ return ${elem_type}();
+ }
+ }
+}
+
+fn f() {
+ let I = 1;
+ let J = 2;
+ let K = 0;
+ let p_a = conv_arr4_Outer(a);
+ let p_a_3 = conv_Outer(a[3u]);
+ let p_a_I = conv_Outer(a[I]);
+ let p_a_3_a = conv_arr4_Inner(a[3u].a);
+ let p_a_I_a = conv_arr4_Inner(a[I].a);
+ let p_a_3_a_2 = conv_Inner(a[3u].a[2u]);
+ let p_a_3_a_I = conv_Inner(a[3u].a[I]);
+ let p_a_I_a_2 = conv_Inner(a[I].a[2u]);
+ let p_a_I_a_J = conv_Inner(a[I].a[J]);
+ let p_a_3_a_2_m = load_a_3_a_2_m();
+ let p_a_3_a_I_m = load_a_3_a_p0_m(u32(I));
+ let p_a_I_a_2_m = load_a_p0_a_2_m(u32(I));
+ let p_a_I_a_J_m = load_a_p0_a_p1_m(u32(I), u32(J));
+ let p_a_3_a_2_m_1 = a[3u].a[2u].m_1;
+ let p_a_I_a_J_m_K = load_a_p0_a_p1_m_p2(u32(I), u32(J), u32(K));
+ let l_a : array<Outer, 4> = conv_arr4_Outer(a);
+ let l_a_3 : Outer = conv_Outer(a[3u]);
+ let l_a_I : Outer = conv_Outer(a[I]);
+ let l_a_3_a : array<Inner, 4> = conv_arr4_Inner(a[3u].a);
+ let l_a_I_a : array<Inner, 4> = conv_arr4_Inner(a[I].a);
+ let l_a_3_a_2 : Inner = conv_Inner(a[3u].a[2u]);
+ let l_a_3_a_I : Inner = conv_Inner(a[3u].a[I]);
+ let l_a_I_a_2 : Inner = conv_Inner(a[I].a[2u]);
+ let l_a_I_a_J : Inner = conv_Inner(a[I].a[J]);
+ let l_a_3_a_2_m : ${mat} = load_a_3_a_2_m();
+ let l_a_3_a_I_m : ${mat} = load_a_3_a_p0_m(u32(I));
+ let l_a_I_a_2_m : ${mat} = load_a_p0_a_2_m(u32(I));
+ let l_a_I_a_J_m : ${mat} = load_a_p0_a_p1_m(u32(I), u32(J));
+ let l_a_3_a_2_m_1 : ${col_vector_type} = a[3u].a[2u].m_1;
+ let l_a_I_a_J_m_K : ${col_vector_type} = load_a_p0_a_p1_m_p2(u32(I), u32(J), u32(K));
+ let l_a_2_a_0_m_1_0 : ${elem_type} = a[3u].a[2u].m_1[0u];
+ let l_a_I_a_J_m_K_I : ${elem_type} = load_a_p0_a_p1_m_p2_p3(u32(I), u32(J), u32(K), u32(I));
+}
+)";
+ std::string col_tableLoadColumn = matrix.JoinTemplatedStringForEachMatrixColumn( //
+ R"( case ${col_id_for_tmpl}u: {
+ return a[p0].a[p1].m_${col_id_for_tmpl};
+ })",
+ "\n");
+ std::string col_tableLoadElement = matrix.JoinTemplatedStringForEachMatrixColumn( //
+ R"( case ${col_id_for_tmpl}u: {
+ return a[p0].a[p1].m_${col_id_for_tmpl}[p3];
+ })",
+ "\n");
+ uint32_t last_size =
+ 64 - static_cast<uint32_t>(matrix.ColumnVectorAlign() * (matrix.columns - 1));
+ expect = matrix.ReplaceFieldsInString(
+ expect,
+ {{"${col_vectors}", matrix.ExpendedColumnVectorsWithLastSize(2, "m_", last_size)},
+ {"${col_vectors_inline_conv_Inner}",
+ matrix.ExpendedColumnVectorsInline("val.m_", ", ")},
+ {"${col_vectors_inline_load_matrix}",
+ matrix.ExpendedColumnVectorsInline("(*(s)).m_", ", ")},
+ {"${col_table_load_column}", col_tableLoadColumn},
+ {"${col_table_load_element}", col_tableLoadElement}});
+ } else {
+ expect = src;
+ }
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_P(Std140Test_Matrix, ArrayStructMatUniform_CopyArray_UniformToStorage) {
auto matrix = GetParam();
diff --git a/src/tint/lang/wgsl/ast/transform/std140_f32_test.cc b/src/tint/lang/wgsl/ast/transform/std140_f32_test.cc
index f5a7d19..ed81103 100644
--- a/src/tint/lang/wgsl/ast/transform/std140_f32_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/std140_f32_test.cc
@@ -693,6 +693,360 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(Std140Test_F32, MatUniform_LoadColumn_ConstIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let p = &a;
+ let l = p[1];
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn f() {
+ let p = conv_mat3x2_f32(a);
+ let l = a.col1;
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(Std140Test_F32, MatUniform_LoadColumn_VariableIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let I = 1;
+ let p = &a;
+ let l = p[I];
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn load_a_p0(p0 : u32) -> vec2<f32> {
+ switch(p0) {
+ case 0u: {
+ return a.col0;
+ }
+ case 1u: {
+ return a.col1;
+ }
+ case 2u: {
+ return a.col2;
+ }
+ default: {
+ return vec2<f32>();
+ }
+ }
+}
+
+fn f() {
+ let I = 1;
+ let p = conv_mat3x2_f32(a);
+ let l = load_a_p0(u32(I));
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(Std140Test_F32, MatUniform_LoadColumnSwizzle_ConstIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let p = &a;
+ let l = p[1].yx;
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn f() {
+ let p = conv_mat3x2_f32(a);
+ let l = a.col1.yx;
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(Std140Test_F32, MatUniform_LoadColumnSwizzle_VariableIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let I = 1;
+ let p = &a;
+ let l = p[I].yx;
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn load_a_p0_yx(p0 : u32) -> vec2<f32> {
+ switch(p0) {
+ case 0u: {
+ return a.col0.yx;
+ }
+ case 1u: {
+ return a.col1.yx;
+ }
+ case 2u: {
+ return a.col2.yx;
+ }
+ default: {
+ return vec2<f32>();
+ }
+ }
+}
+
+fn f() {
+ let I = 1;
+ let p = conv_mat3x2_f32(a);
+ let l = load_a_p0_yx(u32(I));
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(Std140Test_F32,
+ MatUniform_LoadScalar_ConstColumnIndex_ConstRowIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let p = &a;
+ let l = p[1][0];
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn f() {
+ let p = conv_mat3x2_f32(a);
+ let l = a.col1[0u];
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(Std140Test_F32,
+ MatUniform_LoadScalar_VariableColumnIndex_ConstRowIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let I = 0;
+ let p = &a;
+ let l = p[I][0];
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn load_a_p0_0(p0 : u32) -> f32 {
+ switch(p0) {
+ case 0u: {
+ return a.col0[0u];
+ }
+ case 1u: {
+ return a.col1[0u];
+ }
+ case 2u: {
+ return a.col2[0u];
+ }
+ default: {
+ return f32();
+ }
+ }
+}
+
+fn f() {
+ let I = 0;
+ let p = conv_mat3x2_f32(a);
+ let l = load_a_p0_0(u32(I));
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(Std140Test_F32,
+ MatUniform_LoadScalar_ConstColumnIndex_VariableRowIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let I = 0;
+ let p = &a;
+ let l = p[1][I];
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn f() {
+ let I = 0;
+ let p = conv_mat3x2_f32(a);
+ let l = a.col1[I];
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(Std140Test_F32,
+ MatUniform_LoadScalar_VariableColumnIndex_VariableRowIndex_Mat3x2F32_ViaPointerIndex) {
+ auto* src = R"(
+@group(0) @binding(0) var<uniform> a : mat3x2<f32>;
+
+fn f() {
+ let I = 0;
+ let p = &a;
+ let l = p[I][I];
+}
+)";
+
+ auto* expect = R"(
+struct mat3x2_f32 {
+ col0 : vec2<f32>,
+ col1 : vec2<f32>,
+ col2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> a : mat3x2_f32;
+
+fn conv_mat3x2_f32(val : mat3x2_f32) -> mat3x2<f32> {
+ return mat3x2<f32>(val.col0, val.col1, val.col2);
+}
+
+fn load_a_p0_p1(p0 : u32, p1 : u32) -> f32 {
+ switch(p0) {
+ case 0u: {
+ return a.col0[p1];
+ }
+ case 1u: {
+ return a.col1[p1];
+ }
+ case 2u: {
+ return a.col2[p1];
+ }
+ default: {
+ return f32();
+ }
+ }
+}
+
+fn f() {
+ let I = 0;
+ let p = conv_mat3x2_f32(a);
+ let l = load_a_p0_p1(u32(I), u32(I));
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(Std140Test_F32, StructMatUniform_NameCollision_Mat3x2F32) {
auto* src = R"(
struct S {
@@ -841,6 +1195,48 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(Std140Test_F32, StructMatUniform_LoadColumn_ConstIndex_Mat3x2F32_ViaPointerDot) {
+ auto* src = R"(
+struct S {
+ m : mat3x2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> s : S;
+
+fn f() {
+ let p = &s;
+ let l = p.m[1];
+}
+)";
+
+ auto* expect = R"(
+struct S {
+ m : mat3x2<f32>,
+}
+
+struct S_std140 {
+ m_0 : vec2<f32>,
+ m_1 : vec2<f32>,
+ m_2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> s : S_std140;
+
+fn conv_S(val : S_std140) -> S {
+ return S(mat3x2<f32>(val.m_0, val.m_1, val.m_2));
+}
+
+fn f() {
+ let p = conv_S(s);
+ let l = s.m_1;
+}
+)";
+
+ auto got = Run<Std140>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(Std140Test_F32, StructMatUniform_LoadColumn_VariableIndex_Mat3x2F32) {
auto* src = R"(
struct S {
diff --git a/src/tint/lang/wgsl/ast/transform/vectorize_scalar_matrix_initializers_test.cc b/src/tint/lang/wgsl/ast/transform/vectorize_scalar_matrix_initializers_test.cc
index 88e3eb6..fd69553 100644
--- a/src/tint/lang/wgsl/ast/transform/vectorize_scalar_matrix_initializers_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/vectorize_scalar_matrix_initializers_test.cc
@@ -86,7 +86,7 @@
EXPECT_EQ(expect, str(got));
}
-TEST_P(VectorizeScalarMatrixInitializersTest, MultipleScalarsReference) {
+TEST_P(VectorizeScalarMatrixInitializersTest, MultipleScalarsViaReference) {
uint32_t cols = GetParam().first;
uint32_t rows = GetParam().second;
std::string mat_type = "mat" + std::to_string(cols) + "x" + std::to_string(rows) + "<f32>";
@@ -129,6 +129,50 @@
EXPECT_EQ(expect, str(got));
}
+TEST_P(VectorizeScalarMatrixInitializersTest, MultipleScalarsViaPointerIndex) {
+ uint32_t cols = GetParam().first;
+ uint32_t rows = GetParam().second;
+ std::string mat_type = "mat" + std::to_string(cols) + "x" + std::to_string(rows) + "<f32>";
+ std::string vec_type = "vec" + std::to_string(rows) + "<f32>";
+ std::string scalar_values;
+ std::string vector_values;
+ for (uint32_t c = 0; c < cols; c++) {
+ if (c > 0) {
+ vector_values += ", ";
+ scalar_values += ", ";
+ }
+ vector_values += vec_type + "(";
+ for (uint32_t r = 0; r < rows; r++) {
+ if (r > 0) {
+ scalar_values += ", ";
+ vector_values += ", ";
+ }
+ auto value = "p[" + std::to_string((c * rows + r) % 4) + "]";
+ scalar_values += value;
+ vector_values += value;
+ }
+ vector_values += ")";
+ }
+
+ std::string tmpl = R"(
+@fragment
+fn main() {
+ var v = vec4<f32>(1.0, 2.0, 3.0, 8.0);
+ let p = &(v);
+ let m = ${matrix}(${values});
+}
+)";
+ tmpl = tint::ReplaceAll(tmpl, "${matrix}", mat_type);
+ auto src = tint::ReplaceAll(tmpl, "${values}", scalar_values);
+ auto expect = tint::ReplaceAll(tmpl, "${values}", vector_values);
+
+ EXPECT_TRUE(ShouldRun<VectorizeScalarMatrixInitializers>(src));
+
+ auto got = Run<VectorizeScalarMatrixInitializers>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_P(VectorizeScalarMatrixInitializersTest, NonScalarInitializers) {
uint32_t cols = GetParam().first;
uint32_t rows = GetParam().second;
diff --git a/src/tint/lang/wgsl/features/language_feature.cc b/src/tint/lang/wgsl/features/language_feature.cc
index 0dc4b3a..4561eb0 100644
--- a/src/tint/lang/wgsl/features/language_feature.cc
+++ b/src/tint/lang/wgsl/features/language_feature.cc
@@ -60,6 +60,9 @@
if (str == "packed_4x8_integer_dot_product") {
return LanguageFeature::kPacked4X8IntegerDotProduct;
}
+ if (str == "pointer_composite_access") {
+ return LanguageFeature::kPointerCompositeAccess;
+ }
if (str == "readonly_and_readwrite_storage_textures") {
return LanguageFeature::kReadonlyAndReadwriteStorageTextures;
}
@@ -82,6 +85,8 @@
return "chromium_testing_unsafe_experimental";
case LanguageFeature::kPacked4X8IntegerDotProduct:
return "packed_4x8_integer_dot_product";
+ case LanguageFeature::kPointerCompositeAccess:
+ return "pointer_composite_access";
case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
return "readonly_and_readwrite_storage_textures";
}
diff --git a/src/tint/lang/wgsl/features/language_feature.h b/src/tint/lang/wgsl/features/language_feature.h
index 0f1d2c9..aa7315b 100644
--- a/src/tint/lang/wgsl/features/language_feature.h
+++ b/src/tint/lang/wgsl/features/language_feature.h
@@ -52,6 +52,7 @@
kChromiumTestingUnimplemented,
kChromiumTestingUnsafeExperimental,
kPacked4X8IntegerDotProduct,
+ kPointerCompositeAccess,
kReadonlyAndReadwriteStorageTextures,
};
@@ -71,6 +72,7 @@
"chromium_testing_unimplemented",
"chromium_testing_unsafe_experimental",
"packed_4x8_integer_dot_product",
+ "pointer_composite_access",
"readonly_and_readwrite_storage_textures",
};
@@ -82,6 +84,7 @@
LanguageFeature::kChromiumTestingUnimplemented,
LanguageFeature::kChromiumTestingUnsafeExperimental,
LanguageFeature::kPacked4X8IntegerDotProduct,
+ LanguageFeature::kPointerCompositeAccess,
LanguageFeature::kReadonlyAndReadwriteStorageTextures,
};
diff --git a/src/tint/lang/wgsl/features/status.cc b/src/tint/lang/wgsl/features/status.cc
index b0c4bb3..f8297ea 100644
--- a/src/tint/lang/wgsl/features/status.cc
+++ b/src/tint/lang/wgsl/features/status.cc
@@ -35,6 +35,7 @@
switch (f) {
case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
case LanguageFeature::kPacked4X8IntegerDotProduct:
+ case LanguageFeature::kPointerCompositeAccess:
return FeatureStatus::kExperimental;
case LanguageFeature::kUndefined:
return FeatureStatus::kUnknown;
diff --git a/src/tint/lang/wgsl/language_feature_test.cc b/src/tint/lang/wgsl/language_feature_test.cc
index 6a103ad..991f5db 100644
--- a/src/tint/lang/wgsl/language_feature_test.cc
+++ b/src/tint/lang/wgsl/language_feature_test.cc
@@ -64,6 +64,7 @@
{"chromium_testing_unimplemented", LanguageFeature::kChromiumTestingUnimplemented},
{"chromium_testing_unsafe_experimental", LanguageFeature::kChromiumTestingUnsafeExperimental},
{"packed_4x8_integer_dot_product", LanguageFeature::kPacked4X8IntegerDotProduct},
+ {"pointer_composite_access", LanguageFeature::kPointerCompositeAccess},
{"readonly_and_readwrite_storage_textures",
LanguageFeature::kReadonlyAndReadwriteStorageTextures},
};
@@ -87,9 +88,12 @@
{"packed_4x8_integer_Eot_product", LanguageFeature::kUndefined},
{"paked_4x8_integePP_dTTt_product", LanguageFeature::kUndefined},
{"packed_4x8_integxxrdot_pddoduct", LanguageFeature::kUndefined},
- {"readon44y_and_readwrite_storage_textures", LanguageFeature::kUndefined},
- {"readonly_and_readwrite_storageVVSSextures", LanguageFeature::kUndefined},
- {"rRadonly_an_rea22write_storRge_textures", LanguageFeature::kUndefined},
+ {"44ointer_composite_access", LanguageFeature::kUndefined},
+ {"VVSSinter_composite_access", LanguageFeature::kUndefined},
+ {"poinR2er_compRsite_acess", LanguageFeature::kUndefined},
+ {"readonlF_and_readwrite_st9rage_textues", LanguageFeature::kUndefined},
+ {"readonly_and_radwrite_storage_textures", LanguageFeature::kUndefined},
+ {"readonly_and_readwrite_sOOrage_tVxRRures", LanguageFeature::kUndefined},
};
using LanguageFeatureParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
index 9e0ffa9..b763ef9 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/accessor_test.cc
@@ -66,6 +66,60 @@
)");
}
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_ArraySingleIndex_ViaDerefPointerIndex) {
+ // var a: array<u32, 3>
+ // let p = &a;
+ // let b = (*p)[2]
+
+ auto* a = Var("a", ty.array<u32, 3>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", IndexAccessor(Deref(p), 2_u)));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, array<u32, 3>, read_write> = var
+ %p:ptr<function, array<u32, 3>, read_write> = let %a
+ %4:ptr<function, u32, read_write> = access %p, 2u
+ %5:u32 = load %4
+ %b:u32 = let %5
+ ret
+ }
+}
+)");
+}
+
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_ArraySingleIndex_ViaPointerIndex) {
+ // var a: array<u32, 3>
+ // let p = &a;
+ // let b = p[2]
+
+ auto* a = Var("a", ty.array<u32, 3>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", IndexAccessor(p, 2_u)));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, array<u32, 3>, read_write> = var
+ %p:ptr<function, array<u32, 3>, read_write> = let %a
+ %4:ptr<function, u32, read_write> = access %p, 2u
+ %5:u32 = load %4
+ %b:u32 = let %5
+ ret
+ }
+}
+)");
+}
+
TEST_F(ProgramToIRAccessorTest, Accessor_Multiple) {
// let a: vec4<u32> = vec4();
// let b = a[2]
@@ -117,6 +171,58 @@
)");
}
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_VectorSingleIndex_ViaDerefPointerIndex) {
+ // var a: vec3<u32>
+ // let p = &a;
+ // let b = (*p)[2]
+
+ auto* a = Var("a", ty.vec3<u32>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", IndexAccessor(Deref(p), 2_u)));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, vec3<u32>, read_write> = var
+ %p:ptr<function, vec3<u32>, read_write> = let %a
+ %4:u32 = load_vector_element %p, 2u
+ %b:u32 = let %4
+ ret
+ }
+}
+)");
+}
+
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_VectorSingleIndex_ViaPointerIndex) {
+ // var a: vec3<u32>
+ // let p = &a;
+ // let b = p[2]
+
+ auto* a = Var("a", ty.vec3<u32>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", IndexAccessor(p, 2_u)));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, vec3<u32>, read_write> = var
+ %p:ptr<function, vec3<u32>, read_write> = let %a
+ %4:u32 = load_vector_element %p, 2u
+ %b:u32 = let %4
+ ret
+ }
+}
+)");
+}
+
TEST_F(ProgramToIRAccessorTest, Accessor_Var_ArraysMultiIndex) {
// var a: array<array<f32, 4>, 3>
// let b = a[2][3]
@@ -141,6 +247,60 @@
)");
}
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_ArraysMultiIndex_ViaDerefPointerIndex) {
+ // var a: array<array<f32, 4>, 3>
+ // let p = &a;
+ // let b = (*p)[2][3]
+
+ auto* a = Var("a", ty.array<array<f32, 4>, 3>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", IndexAccessor(IndexAccessor(Deref(p), 2_u), 3_u)));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, array<array<f32, 4>, 3>, read_write> = var
+ %p:ptr<function, array<array<f32, 4>, 3>, read_write> = let %a
+ %4:ptr<function, f32, read_write> = access %p, 2u, 3u
+ %5:f32 = load %4
+ %b:f32 = let %5
+ ret
+ }
+}
+)");
+}
+
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_ArraysMultiIndex_ViaPointerIndex) {
+ // var a: array<array<f32, 4>, 3>
+ // let p = &a;
+ // let b = p[2][3]
+
+ auto* a = Var("a", ty.array<array<f32, 4>, 3>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", IndexAccessor(IndexAccessor(p, 2_u), 3_u)));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, array<array<f32, 4>, 3>, read_write> = var
+ %p:ptr<function, array<array<f32, 4>, 3>, read_write> = let %a
+ %4:ptr<function, f32, read_write> = access %p, 2u, 3u
+ %5:f32 = load %4
+ %b:f32 = let %5
+ ret
+ }
+}
+)");
+}
+
TEST_F(ProgramToIRAccessorTest, Accessor_Var_MatrixMultiIndex) {
// var a: mat3x4<f32>
// let b = a[2][3]
@@ -197,6 +357,76 @@
)");
}
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_SingleMember_ViaDerefPointerIndex) {
+ // struct MyStruct { foo: i32 }
+ // var a: MyStruct;
+ // let p = &a;
+ // let b = (*p).foo
+
+ auto* s = Structure("MyStruct", Vector{
+ Member("foo", ty.i32()),
+ });
+ auto* a = Var("a", ty.Of(s), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", MemberAccessor(Deref(p), "foo")));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(MyStruct = struct @align(4) {
+ foo:i32 @offset(0)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, MyStruct, read_write> = var
+ %p:ptr<function, MyStruct, read_write> = let %a
+ %4:ptr<function, i32, read_write> = access %p, 0u
+ %5:i32 = load %4
+ %b:i32 = let %5
+ ret
+ }
+}
+)");
+}
+
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_SingleMember_ViaPointerIndex) {
+ // struct MyStruct { foo: i32 }
+ // var a: MyStruct;
+ // let p = &a;
+ // let b = p.foo
+
+ auto* s = Structure("MyStruct", Vector{
+ Member("foo", ty.i32()),
+ });
+ auto* a = Var("a", ty.Of(s), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", MemberAccessor(p, "foo")));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(MyStruct = struct @align(4) {
+ foo:i32 @offset(0)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, MyStruct, read_write> = var
+ %p:ptr<function, MyStruct, read_write> = let %a
+ %4:ptr<function, i32, read_write> = access %p, 0u
+ %5:i32 = load %4
+ %b:i32 = let %5
+ ret
+ }
+}
+)");
+}
+
TEST_F(ProgramToIRAccessorTest, Accessor_Var_MultiMember) {
// struct Inner { bar: f32 }
// struct Outer { a: i32, foo: Inner }
@@ -333,6 +563,58 @@
)");
}
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_VectorElementSwizzle_ViaDerefPointerIndex) {
+ // var a: vec2<f32>
+ // let p = &a;
+ // let b = (*p).y
+
+ auto* a = Var("a", ty.vec2<f32>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", MemberAccessor(Deref(p), "y")));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, vec2<f32>, read_write> = var
+ %p:ptr<function, vec2<f32>, read_write> = let %a
+ %4:f32 = load_vector_element %p, 1u
+ %b:f32 = let %4
+ ret
+ }
+}
+)");
+}
+
+TEST_F(ProgramToIRAccessorTest, Accessor_Var_VectorElementSwizzle_ViaPointerIndex) {
+ // var a: vec2<f32>
+ // let p = &a;
+ // let b = p.y
+
+ auto* a = Var("a", ty.vec2<f32>(), core::AddressSpace::kFunction);
+ auto* p = Let("p", AddressOf(a));
+ auto* expr = Decl(Let("b", MemberAccessor(p, "y")));
+ WrapInFunction(Decl(a), Decl(p), expr);
+
+ auto m = Build();
+ ASSERT_EQ(m, Success);
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %a:ptr<function, vec2<f32>, read_write> = var
+ %p:ptr<function, vec2<f32>, read_write> = let %a
+ %4:f32 = load_vector_element %p, 1u
+ %b:f32 = let %4
+ ret
+ }
+}
+)");
+}
+
TEST_F(ProgramToIRAccessorTest, Accessor_Var_MultiElementSwizzle) {
// var a: vec3<f32>
// let b = a.zyxz
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index 740c918..cc5a137 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -1098,12 +1098,12 @@
return std::nullopt;
}
- auto* ref = access->Object()->Type()->As<core::type::Reference>();
- if (!ref) {
+ auto* memory_view = access->Object()->Type()->As<core::type::MemoryView>();
+ if (!memory_view) {
return std::nullopt;
}
- if (!ref->StoreType()->Is<core::type::Vector>()) {
+ if (!memory_view->StoreType()->Is<core::type::Vector>()) {
return std::nullopt;
}
return tint::Switch(
diff --git a/src/tint/lang/wgsl/resolver/array_accessor_test.cc b/src/tint/lang/wgsl/resolver/array_accessor_test.cc
index aece64f..937568d 100644
--- a/src/tint/lang/wgsl/resolver/array_accessor_test.cc
+++ b/src/tint/lang/wgsl/resolver/array_accessor_test.cc
@@ -41,7 +41,7 @@
using ResolverIndexAccessorTest = ResolverTest;
-TEST_F(ResolverIndexAccessorTest, Matrix_Dynamic_F32) {
+TEST_F(ResolverIndexAccessorTest, Matrix_F32) {
GlobalVar("my_var", ty.mat2x3<f32>(), core::AddressSpace::kPrivate);
auto* acc = IndexAccessor("my_var", Expr(Source{{12, 34}}, 1_f));
WrapInFunction(acc);
@@ -82,7 +82,7 @@
TEST_F(ResolverIndexAccessorTest, Matrix_Dynamic) {
GlobalConst("my_const", ty.mat2x3<f32>(), Call<mat2x3<f32>>());
auto* idx = Var("idx", ty.i32(), Call<i32>());
- auto* acc = IndexAccessor("my_const", Expr(Source{{12, 34}}, idx));
+ auto* acc = IndexAccessor("my_const", idx);
WrapInFunction(Decl(idx), acc);
EXPECT_TRUE(r()->Resolve());
@@ -97,7 +97,7 @@
TEST_F(ResolverIndexAccessorTest, Matrix_XDimension_Dynamic) {
GlobalConst("my_const", ty.mat4x4<f32>(), Call<mat4x4<f32>>());
auto* idx = Var("idx", ty.u32(), Expr(3_u));
- auto* acc = IndexAccessor("my_const", Expr(Source{{12, 34}}, idx));
+ auto* acc = IndexAccessor("my_const", idx);
WrapInFunction(Decl(idx), acc);
EXPECT_TRUE(r()->Resolve());
@@ -106,9 +106,10 @@
TEST_F(ResolverIndexAccessorTest, Matrix_BothDimension_Dynamic) {
GlobalConst("my_const", ty.mat4x4<f32>(), Call<mat4x4<f32>>());
- auto* idx = Var("idy", ty.u32(), Expr(2_u));
- auto* acc = IndexAccessor(IndexAccessor("my_const", Expr(Source{{12, 34}}, idx)), 1_i);
- WrapInFunction(Decl(idx), acc);
+ auto* idx = Var("idx", ty.u32(), Expr(3_u));
+ auto* idy = Var("idy", ty.u32(), Expr(2_u));
+ auto* acc = IndexAccessor(IndexAccessor("my_const", idx), idy);
+ WrapInFunction(Decl(idx), Decl(idy), acc);
EXPECT_TRUE(r()->Resolve());
EXPECT_EQ(r()->error(), "");
@@ -175,16 +176,16 @@
TEST_F(ResolverIndexAccessorTest, Vector_Dynamic) {
GlobalConst("my_const", ty.vec3<f32>(), Call<vec3<f32>>());
auto* idx = Var("idx", ty.i32(), Expr(2_i));
- auto* acc = IndexAccessor("my_const", Expr(Source{{12, 34}}, idx));
+ auto* acc = IndexAccessor("my_const", idx);
WrapInFunction(Decl(idx), acc);
EXPECT_TRUE(r()->Resolve());
}
TEST_F(ResolverIndexAccessorTest, Vector) {
- GlobalVar("my_var", ty.vec3<f32>(), core::AddressSpace::kPrivate);
+ GlobalConst("my_const", ty.vec3<f32>(), Call<vec3<f32>>());
- auto* acc = IndexAccessor("my_var", 2_i);
+ auto* acc = IndexAccessor("my_const", 2_i);
WrapInFunction(acc);
EXPECT_TRUE(r()->Resolve()) << r()->error();
@@ -346,6 +347,26 @@
EXPECT_EQ(idx_sem->Object()->Declaration(), acc->object);
}
+TEST_F(ResolverIndexAccessorTest, Expr_ImplicitDeref_FuncGoodParent) {
+ // fn func(p: ptr<function, vec4<f32>>) -> f32 {
+ // let idx: u32 = u32();
+ // let x: f32 = p[idx];
+ // return x;
+ // }
+ auto* p = Param("p", ty.ptr<function, vec4<f32>>());
+ auto* idx = Let("idx", ty.u32(), Call<u32>());
+ auto* acc = IndexAccessor(Source{{12, 34}}, p, idx);
+ auto* x = Var("x", ty.f32(), acc);
+ Func("func", Vector{p}, ty.f32(), Vector{Decl(idx), Decl(x), Return(x)});
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto idx_sem = Sem().Get(acc)->UnwrapLoad()->As<sem::IndexAccessorExpression>();
+ ASSERT_NE(idx_sem, nullptr);
+ EXPECT_EQ(idx_sem->Index()->Declaration(), acc->index);
+ EXPECT_EQ(idx_sem->Object()->Declaration(), acc->object);
+}
+
TEST_F(ResolverIndexAccessorTest, Expr_Deref_FuncBadParent) {
// fn func(p: ptr<function, vec4<f32>>) -> f32 {
// let idx: u32 = u32();
@@ -360,11 +381,10 @@
Func("func", Vector{p}, ty.f32(), Vector{Decl(idx), Decl(x), Return(x)});
EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- "12:34 error: cannot index type 'ptr<function, vec4<f32>, read_write>'");
+ EXPECT_EQ(r()->error(), "12:34 error: cannot dereference expression of type 'f32'");
}
-TEST_F(ResolverIndexAccessorTest, Exr_Deref_BadParent) {
+TEST_F(ResolverIndexAccessorTest, Expr_Deref_BadParent) {
// var param: vec4<f32>
// let x: f32 = *(¶m)[0];
auto* param = Var("param", ty.vec4<f32>());
@@ -376,8 +396,7 @@
WrapInFunction(param, idx, x);
EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- "12:34 error: cannot index type 'ptr<function, vec4<f32>, read_write>'");
+ EXPECT_EQ(r()->error(), "12:34 error: cannot dereference expression of type 'f32'");
}
} // namespace
diff --git a/src/tint/lang/wgsl/resolver/ptr_ref_test.cc b/src/tint/lang/wgsl/resolver/ptr_ref_test.cc
index 70c5b7b..892dd90 100644
--- a/src/tint/lang/wgsl/resolver/ptr_ref_test.cc
+++ b/src/tint/lang/wgsl/resolver/ptr_ref_test.cc
@@ -35,6 +35,7 @@
namespace tint::resolver {
namespace {
+using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
struct ResolverPtrRefTest : public resolver::TestHelper, public testing::Test {};
@@ -122,5 +123,320 @@
EXPECT_EQ(TypeOf(storage_ptr)->As<core::type::Pointer>()->Access(), core::Access::kRead);
}
+TEST_F(ResolverPtrRefTest, ArrayIndexAccessorViaDerefPointer) {
+ // var a : array<i32, 3>;
+ // let p = &a;
+ // (*p)[0]
+
+ auto* v = Var("v", ty.array<i32, 3>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(Deref(p), 0_i);
+
+ WrapInFunction(v, p, expr);
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, ArrayIndexAccessorViaPointer) {
+ // var a : array<i32, 3>;
+ // let p = &a;
+ // p[0]
+
+ auto* v = Var("v", ty.array<i32, 3>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(p, 0_i);
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, ArrayIndexAccessorViaPointer_FeatureDisallowed) {
+ // var a : array<i32, 3>;
+ // let p = &a;
+ // p[0]
+
+ auto* v = Var("v", ty.array<i32, 3>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(Source{Source::Location{12, 34}}, p, 0_i);
+ WrapInFunction(v, p, expr);
+
+ auto resolver = Resolver(this, {});
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "12:34 error: pointer composite access requires the pointer_composite_access "
+ "language feature, which is not allowed in the current environment");
+}
+
+TEST_F(ResolverPtrRefTest, VectorIndexAccessorViaDerefPointer) {
+ // var a : vec3<i32>;
+ // let p = &a;
+ // (*p)[0]
+
+ auto* v = Var("v", ty.vec3<i32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(Deref(p), 0_i);
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, VectorIndexAccessorViaPointer) {
+ // var a : vec3<i32>;
+ // let p = &a;
+ // p[0]
+
+ auto* v = Var("v", ty.vec3<i32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(p, 0_i);
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, VectorIndexAccessorViaPointer_FeatureDisallowed) {
+ // var a : vec3<i32>;
+ // let p = &a;
+ // p[0]
+
+ auto* v = Var("v", ty.vec3<i32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(Source{Source::Location{12, 34}}, p, 0_i);
+ WrapInFunction(v, p, expr);
+
+ auto resolver = Resolver(this, {});
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "12:34 error: pointer composite access requires the pointer_composite_access "
+ "language feature, which is not allowed in the current environment");
+}
+
+TEST_F(ResolverPtrRefTest, VectorMemberAccessorViaDerefPointer) {
+ // var a : vec3<i32>;
+ // let p = &a;
+ // (*p).x
+
+ auto* v = Var("v", ty.vec3<i32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = MemberAccessor(Deref(p), "x");
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, VectorMemberAccessorViaPointer) {
+ // var a : vec3<i32>;
+ // let p = &a;
+ // p.x
+
+ auto* v = Var("v", ty.vec3<i32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = MemberAccessor(p, "x");
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, VectorMemberAccessorViaPointer_FeatureDisallowed) {
+ // var a : vec3<i32>;
+ // let p = &a;
+ // p.x
+
+ auto* v = Var("v", ty.vec3<i32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = MemberAccessor(Source{Source::Location{12, 34}}, p, "x");
+ WrapInFunction(v, p, expr);
+
+ auto resolver = Resolver(this, {});
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "12:34 error: pointer composite access requires the pointer_composite_access "
+ "language feature, which is not allowed in the current environment");
+}
+
+TEST_F(ResolverPtrRefTest, MatrixIndexAccessorViaDerefPointer) {
+ // var a : mat2x3<f32>;
+ // let p = &a;
+ // (*p)[0]
+
+ auto* v = Var("v", ty.mat2x3<f32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(Deref(p), 0_i);
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ auto* vec = ref->Type()->As<core::type::Reference>()->StoreType()->As<core::type::Vector>();
+ ASSERT_TRUE(vec);
+ EXPECT_EQ(vec->Elements().count, 3u);
+ EXPECT_TRUE(vec->Elements().type->Is<core::type::F32>());
+}
+
+TEST_F(ResolverPtrRefTest, MatrixIndexAccessorViaPointer) {
+ // var a : mat2x3<f32>;
+ // let p = &a;
+ // p[0]
+
+ auto* v = Var("v", ty.mat2x3<f32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(p, 0_i);
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ auto* vec = ref->Type()->As<core::type::Reference>()->StoreType()->As<core::type::Vector>();
+ ASSERT_TRUE(vec);
+ EXPECT_EQ(vec->Elements().count, 3u);
+ EXPECT_TRUE(vec->Elements().type->Is<core::type::F32>());
+}
+
+TEST_F(ResolverPtrRefTest, MatrixIndexAccessorViaPointer_FeatureDisallowed) {
+ // var a : mat2x3<f32>;
+ // let p = &a;
+ // p[0]
+
+ auto* v = Var("v", ty.mat2x3<f32>());
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = IndexAccessor(Source{Source::Location{12, 34}}, p, 0_i);
+ WrapInFunction(v, p, expr);
+
+ auto resolver = Resolver(this, {});
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "12:34 error: pointer composite access requires the pointer_composite_access "
+ "language feature, which is not allowed in the current environment");
+}
+
+TEST_F(ResolverPtrRefTest, StructMemberAccessorViaDerefPointer) {
+ // struct S { a : i32, }
+ // var a : S;
+ // let p = &a;
+ // (*p).a
+
+ auto* s = Structure("S", Vector{
+ Member("a", ty.i32()),
+ });
+ auto* v = Var("v", ty.Of(s));
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = MemberAccessor(Deref(p), "a");
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, StructMemberAccessorViaPointer) {
+ // struct S { a : i32, }
+ // var a : S;
+ // let p = &a;
+ // p.a
+
+ auto* s = Structure("S", Vector{
+ Member("a", ty.i32()),
+ });
+ auto* v = Var("v", ty.Of(s));
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = MemberAccessor(p, "a");
+ WrapInFunction(v, p, expr);
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* load = Sem().Get<sem::Load>(expr);
+ ASSERT_NE(load, nullptr);
+
+ auto* ref = load->Reference();
+ ASSERT_NE(ref, nullptr);
+
+ ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(ref->Type()->As<core::type::Reference>()->StoreType()->Is<core::type::I32>());
+}
+
+TEST_F(ResolverPtrRefTest, StructMemberAccessorViaPointer_FeatureDisallowed) {
+ // struct S { a : i32, }
+ // var a : S;
+ // let p = &a;
+ // p.a
+
+ auto* s = Structure("S", Vector{
+ Member("a", ty.i32()),
+ });
+ auto* v = Var("v", ty.Of(s));
+ auto* p = Let("p", AddressOf(v));
+ auto* expr = MemberAccessor(Source{Source::Location{12, 34}}, p, "a");
+ WrapInFunction(v, p, expr);
+
+ auto resolver = Resolver(this, {});
+ EXPECT_FALSE(resolver.Resolve());
+ EXPECT_EQ(resolver.error(),
+ "12:34 error: pointer composite access requires the pointer_composite_access "
+ "language feature, which is not allowed in the current environment");
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index a94a59f..454e928 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -44,6 +44,7 @@
#include "src/tint/lang/core/type/depth_multisampled_texture.h"
#include "src/tint/lang/core/type/depth_texture.h"
#include "src/tint/lang/core/type/external_texture.h"
+#include "src/tint/lang/core/type/memory_view.h"
#include "src/tint/lang/core/type/multisampled_texture.h"
#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/lang/core/type/reference.h"
@@ -1979,17 +1980,30 @@
if (!obj) {
return nullptr;
}
- auto* obj_raw_ty = obj->Type();
- auto* obj_ty = obj_raw_ty->UnwrapRef();
+ auto* object_ty = obj->Type();
+ auto* const memory_view = object_ty->As<core::type::MemoryView>();
+ const core::type::Type* storage_ty = object_ty->UnwrapRef();
+ if (memory_view) {
+ if (memory_view->Is<core::type::Pointer>() &&
+ !allowed_features_.features.count(wgsl::LanguageFeature::kPointerCompositeAccess)) {
+ AddError(
+ "pointer composite access requires the pointer_composite_access language feature, "
+ "which is not allowed in the current environment",
+ expr->source);
+ return nullptr;
+ }
+ storage_ty = memory_view->StoreType();
+ }
+
auto* ty = Switch(
- obj_ty, //
+ storage_ty, //
[&](const sem::Array* arr) { return arr->ElemType(); },
[&](const core::type::Vector* vec) { return vec->type(); },
[&](const core::type::Matrix* mat) {
return b.create<core::type::Vector>(mat->type(), mat->rows());
},
[&](Default) {
- AddError("cannot index type '" + sem_.TypeNameOf(obj_ty) + "'", expr->source);
+ AddError("cannot index type '" + sem_.TypeNameOf(storage_ty) + "'", expr->source);
return nullptr;
});
if (ty == nullptr) {
@@ -2003,9 +2017,10 @@
return nullptr;
}
- // If we're extracting from a reference, we return a reference.
- if (auto* ref = obj_raw_ty->As<core::type::Reference>()) {
- ty = b.create<core::type::Reference>(ref->AddressSpace(), ty, ref->Access());
+ // If we're extracting from a memory view, we return a reference.
+ if (memory_view) {
+ ty =
+ b.create<core::type::Reference>(memory_view->AddressSpace(), ty, memory_view->Access());
}
const core::constant::Value* val = nullptr;
@@ -3398,7 +3413,20 @@
}
auto* object_ty = object->Type();
- auto* storage_ty = object_ty->UnwrapRef();
+
+ auto* const memory_view = object_ty->As<core::type::MemoryView>();
+ const core::type::Type* storage_ty = object_ty->UnwrapRef();
+ if (memory_view) {
+ if (memory_view->Is<core::type::Pointer>() &&
+ !allowed_features_.features.count(wgsl::LanguageFeature::kPointerCompositeAccess)) {
+ AddError(
+ "pointer composite access requires the pointer_composite_access language feature, "
+ "which is not allowed in the current environment",
+ expr->source);
+ return nullptr;
+ }
+ storage_ty = memory_view->StoreType();
+ }
auto* root_ident = object->RootIdentifier();
@@ -3429,9 +3457,10 @@
ty = member->Type();
- // If we're extracting from a reference, we return a reference.
- if (auto* ref = object_ty->As<core::type::Reference>()) {
- ty = b.create<core::type::Reference>(ref->AddressSpace(), ty, ref->Access());
+ // If we're extracting from a memory view, we return a reference.
+ if (memory_view) {
+ ty = b.create<core::type::Reference>(memory_view->AddressSpace(), ty,
+ memory_view->Access());
}
const core::constant::Value* val = nullptr;
@@ -3497,9 +3526,10 @@
if (size == 1) {
// A single element swizzle is just the type of the vector.
ty = vec->type();
- // If we're extracting from a reference, we return a reference.
- if (auto* ref = object_ty->As<core::type::Reference>()) {
- ty = b.create<core::type::Reference>(ref->AddressSpace(), ty, ref->Access());
+ // If we're extracting from a memory view, we return a reference.
+ if (memory_view) {
+ ty = b.create<core::type::Reference>(memory_view->AddressSpace(), ty,
+ memory_view->Access());
}
} else {
// The vector will have a number of components equal to the length of
diff --git a/src/tint/lang/wgsl/resolver/uniformity.cc b/src/tint/lang/wgsl/resolver/uniformity.cc
index 0cea0be..236c458 100644
--- a/src/tint/lang/wgsl/resolver/uniformity.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity.cc
@@ -1445,14 +1445,43 @@
},
[&](const ast::IndexAccessorExpression* i) {
- auto [cf1, l1, root_ident] =
- ProcessLValueExpression(cf, i->object, /*is_partial_reference*/ true);
+ auto* sem_object = sem_.GetVal(i->object);
+
+ LValue object_result;
+ if (sem_object->Type()->Is<core::type::Pointer>()) {
+ // Sugared pointer access, treat as indirection
+ auto* root_ident = sem_object->RootIdentifier();
+ auto* deref = CreateNode({NameFor(root_ident), "_deref"});
+ if (auto* old_value = current_function_->variables.Get(root_ident)) {
+ // We're dereferecing a partial pointer, so link back to the variable's
+ // previous value.
+ deref->AddEdge(old_value);
+ }
+ object_result = LValue{cf, deref, root_ident};
+ } else {
+ object_result =
+ ProcessLValueExpression(cf, i->object, /*is_partial_reference*/ true);
+ }
+ auto [cf1, l1, root_ident] = object_result;
auto [cf2, v2] = ProcessExpression(cf1, i->index);
l1->AddEdge(v2);
return LValue{cf2, l1, root_ident};
},
[&](const ast::MemberAccessorExpression* m) {
+ auto* sem_object = sem_.GetVal(m->object);
+ if (sem_object->Type()->Is<core::type::Pointer>()) {
+ // Sugared pointer access, treat as indirection
+ auto* root_ident = sem_object->RootIdentifier();
+ auto* deref = CreateNode({NameFor(root_ident), "_deref"});
+ if (auto* old_value = current_function_->variables.Get(root_ident)) {
+ // We're dereferecing a partial pointer, so link back to the variable's
+ // previous value.
+ deref->AddEdge(old_value);
+ }
+ return LValue{cf, deref, root_ident};
+ }
+
return ProcessLValueExpression(cf, m->object, /*is_partial_reference*/ true);
},
diff --git a/src/tint/lang/wgsl/resolver/uniformity_test.cc b/src/tint/lang/wgsl/resolver/uniformity_test.cc
index 769ab82..fe1389e 100644
--- a/src/tint/lang/wgsl/resolver/uniformity_test.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity_test.cc
@@ -5631,6 +5631,70 @@
)");
}
+TEST_F(UniformityAnalysisTest,
+ VectorElement_VectorBecomesUniform_PartialAssignment_ViaPointerDerefIndex) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ let p = &v;
+ (*p)[1] = rw;
+ v = vec4(1, 2, 3, v[3]);
+ if (v[1] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (v[1] == 0) {
+ ^^
+
+test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ (*p)[1] = rw;
+ ^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ VectorElement_VectorBecomesUniform_PartialAssignment_ViaPointerIndex) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ let p = &v;
+ p[1] = rw;
+ v = vec4(1, 2, 3, v[3]);
+ if (v[1] == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (v[1] == 0) {
+ ^^
+
+test:7:10 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ p[1] = rw;
+ ^^
+)");
+}
+
TEST_F(UniformityAnalysisTest, VectorElementViaMember_VectorBecomesUniform_PartialAssignment) {
std::string src = R"(
@group(0) @binding(0) var<storage, read_write> rw : i32;
@@ -5661,6 +5725,70 @@
)");
}
+TEST_F(UniformityAnalysisTest,
+ VectorElementViaMember_VectorBecomesUniform_PartialAssignment_ViaPointerDerefDot) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ let p = &v;
+ (*p).y = rw;
+ v = vec4(1, 2, 3, v.w);
+ if (v.y == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (v.y == 0) {
+ ^^
+
+test:7:12 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ (*p).y = rw;
+ ^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ VectorElementViaMember_VectorBecomesUniform_PartialAssignment_ViaPointerDot) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : i32;
+
+fn foo() {
+ var v : vec4<i32>;
+ let p = &v;
+ p.y = rw;
+ v = vec4(1, 2, 3, v.w);
+ if (v.y == 0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (v.y == 0) {
+ ^^
+
+test:7:9 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ p.y = rw;
+ ^^
+)");
+}
+
TEST_F(UniformityAnalysisTest, VectorElement_DifferentElementBecomesUniform) {
// For aggregate types, we conservatively consider them to be non-uniform once they
// become non-uniform. Test that after assigning a uniform value to an element, the whole vector
@@ -5961,6 +6089,146 @@
}
TEST_F(UniformityAnalysisTest,
+ MatrixElement_ColumnBecomesUniform_ThroughCapturedPartialPointer_PointerDerefIndex) {
+ // For aggregate types, we conservatively consider them to be non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, that element is
+ // still considered to be non-uniform.
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : f32;
+
+fn foo() {
+ var m : mat3x3<f32>;
+ let p = &m[1];
+ m[1][1] = rw;
+ (*p)[0] = 0.0;
+ if (m[1][1] == 0.0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (m[1][1] == 0.0) {
+ ^^
+
+test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ m[1][1] = rw;
+ ^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ MatrixElement_ColumnBecomesUniform_ThroughCapturedPartialPointer_PointerIndex) {
+ // For aggregate types, we conservatively consider them to be non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, that element is
+ // still considered to be non-uniform.
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : f32;
+
+fn foo() {
+ var m : mat3x3<f32>;
+ let p = &m[1];
+ m[1][1] = rw;
+ p[0] = 0.0;
+ if (m[1][1] == 0.0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (m[1][1] == 0.0) {
+ ^^
+
+test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ m[1][1] = rw;
+ ^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ MatrixElement_ColumnBecomesUniform_ThroughCapturedPartialPointer_PointerDerefDot) {
+ // For aggregate types, we conservatively consider them to be non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, that element is
+ // still considered to be non-uniform.
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : f32;
+
+fn foo() {
+ var m : mat3x3<f32>;
+ let p = &m[1];
+ m[1][1] = rw;
+ (*p).x = 0.0;
+ if (m[1][1] == 0.0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (m[1][1] == 0.0) {
+ ^^
+
+test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ m[1][1] = rw;
+ ^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
+ MatrixElement_ColumnBecomesUniform_ThroughCapturedPartialPointer_PointerDot) {
+ // For aggregate types, we conservatively consider them to be non-uniform once they
+ // become non-uniform. Test that after assigning a uniform value to an element, that element is
+ // still considered to be non-uniform.
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> rw : f32;
+
+fn foo() {
+ var m : mat3x3<f32>;
+ let p = &m[1];
+ m[1][1] = rw;
+ p.x = 0.0;
+ if (m[1][1] == 0.0) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:9:3 note: control flow depends on possibly non-uniform value
+ if (m[1][1] == 0.0) {
+ ^^
+
+test:7:13 note: reading from read_write storage buffer 'rw' may result in a non-uniform value
+ m[1][1] = rw;
+ ^^
+)");
+}
+
+TEST_F(UniformityAnalysisTest,
MatrixElement_ColumnBecomesUniform_ThroughCapturedPartialPointerChain) {
// For aggregate types, we conservatively consider them to be non-uniform once they
// become non-uniform. Test that after assigning a uniform value to an element, that element is
diff --git a/src/tint/lang/wgsl/resolver/validation_test.cc b/src/tint/lang/wgsl/resolver/validation_test.cc
index 9d369b8..6bc6fc2 100644
--- a/src/tint/lang/wgsl/resolver/validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/validation_test.cc
@@ -383,12 +383,10 @@
WrapInFunction(Decl(param), Decl(ret));
EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- "12:34 error: invalid member accessor expression. Expected vector or struct, got "
- "'ptr<function, vec4<f32>, read_write>'");
+ EXPECT_EQ(r()->error(), "error: cannot dereference expression of type 'f32'");
}
-TEST_F(ResolverValidationTest, EXpr_MemberAccessor_FuncGoodParent) {
+TEST_F(ResolverValidationTest, Expr_MemberAccessor_FuncGoodParent) {
// fn func(p: ptr<function, vec4<f32>>) -> f32 {
// let x: f32 = (*p).z;
// return x;
@@ -405,7 +403,7 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
-TEST_F(ResolverValidationTest, EXpr_MemberAccessor_FuncBadParent) {
+TEST_F(ResolverValidationTest, Expr_MemberAccessor_FuncBadParent) {
// fn func(p: ptr<function, vec4<f32>>) -> f32 {
// let x: f32 = *p.z;
// return x;
@@ -421,9 +419,7 @@
});
EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- "12:34 error: invalid member accessor expression. Expected vector or struct, got "
- "'ptr<function, vec4<f32>, read_write>'");
+ EXPECT_EQ(r()->error(), "error: cannot dereference expression of type 'f32'");
}
TEST_F(ResolverValidationTest,
@@ -1276,24 +1272,6 @@
EXPECT_EQ(r()->error(), R"(12:24 error: value 4294967296 cannot be represented as 'u32')");
}
-// var a: array<i32,2>;
-// *&a[0] = 1;
-TEST_F(ResolverTest, PointerIndexing_Fail) {
- // var a: array<i32,2>;
- // let p = &a;
- // *p[0] = 0;
-
- auto* a = Var("a", ty.array<i32, 2>());
- auto* p = AddressOf("a");
- auto* idx = Assign(Deref(IndexAccessor(p, 0_u)), 0_u);
-
- WrapInFunction(a, idx);
-
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- R"(error: cannot index type 'ptr<function, array<i32, 2>, read_write>')");
-}
-
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 8b40557..41cd202 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -93,6 +93,7 @@
enum language_feature {
readonly_and_readwrite_storage_textures
packed_4x8_integer_dot_product
+ pointer_composite_access
// Language features used only for testing whose status will never change.
chromium_testing_unimplemented
diff --git a/test/tint/ptr_sugar/array.wgsl b/test/tint/ptr_sugar/array.wgsl
new file mode 100644
index 0000000..d48a62b
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var b = (*p)[0];
+ (*p)[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : array<i32, 10>;
+ let p = &a;
+ let i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : array<i32, 10>;
+ let p = &a;
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : array<i32, 10>;
+ let p = &a;
+ var i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/array.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..c876ea4
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.dxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/array.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..c876ea4
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.fxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int a[10] = (int[10])0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int a[10] = (int[10])0;
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int a[10] = (int[10])0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.glsl b/test/tint/ptr_sugar/array.wgsl.expected.glsl
new file mode 100644
index 0000000..48419fc
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.glsl
@@ -0,0 +1,56 @@
+#version 310 es
+
+void deref_const() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int a[10] = int[10](0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.msl b/test/tint/ptr_sugar/array.wgsl.expected.msl
new file mode 100644
index 0000000..9bd0861
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.msl
@@ -0,0 +1,66 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+void deref_const() {
+ tint_array<int, 10> a = {};
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ tint_array<int, 10> a = {};
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ tint_array<int, 10> a = {};
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ tint_array<int, 10> a = {};
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ tint_array<int, 10> a = {};
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ tint_array<int, 10> a = {};
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+kernel void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.spvasm b/test/tint/ptr_sugar/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..fd7b149
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.spvasm
@@ -0,0 +1,124 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 66
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref_const "deref_const"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref_const "no_deref_const"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %deref_let "deref_let"
+ OpName %a_1 "a"
+ OpName %b_1 "b"
+ OpName %no_deref_let "no_deref_let"
+ OpName %a_2 "a"
+ OpName %b_2 "b"
+ OpName %deref_var "deref_var"
+ OpName %a_3 "a"
+ OpName %i "i"
+ OpName %b_3 "b"
+ OpName %no_deref_var "no_deref_var"
+ OpName %a_4 "a"
+ OpName %i_0 "i"
+ OpName %b_4 "b"
+ OpName %main "main"
+ OpDecorate %_arr_int_uint_10 ArrayStride 4
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %uint_10 = OpConstant %uint 10
+%_arr_int_uint_10 = OpTypeArray %int %uint_10
+%_ptr_Function__arr_int_uint_10 = OpTypePointer Function %_arr_int_uint_10
+ %11 = OpConstantNull %_arr_int_uint_10
+ %12 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+%deref_const = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b = OpVariable %_ptr_Function_int Function %12
+ %14 = OpAccessChain %_ptr_Function_int %a %12
+ %15 = OpLoad %int %14
+ OpStore %b %15
+ %17 = OpAccessChain %_ptr_Function_int %a %12
+ OpStore %17 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_const = OpFunction %void None %1
+ %20 = OpLabel
+ %a_0 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b_0 = OpVariable %_ptr_Function_int Function %12
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %12
+ %23 = OpLoad %int %22
+ OpStore %b_0 %23
+ %25 = OpAccessChain %_ptr_Function_int %a_0 %12
+ OpStore %25 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_let = OpFunction %void None %1
+ %27 = OpLabel
+ %a_1 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b_1 = OpVariable %_ptr_Function_int Function %12
+ %29 = OpAccessChain %_ptr_Function_int %a_1 %12
+ %30 = OpLoad %int %29
+ OpStore %b_1 %30
+ %32 = OpAccessChain %_ptr_Function_int %a_1 %12
+ OpStore %32 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_let = OpFunction %void None %1
+ %34 = OpLabel
+ %a_2 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %b_2 = OpVariable %_ptr_Function_int Function %12
+ %36 = OpAccessChain %_ptr_Function_int %a_2 %12
+ %37 = OpLoad %int %36
+ OpStore %b_2 %37
+ %39 = OpAccessChain %_ptr_Function_int %a_2 %12
+ OpStore %39 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_var = OpFunction %void None %1
+ %41 = OpLabel
+ %a_3 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %i = OpVariable %_ptr_Function_int Function %12
+ %b_3 = OpVariable %_ptr_Function_int Function %12
+ OpStore %i %12
+ %44 = OpLoad %int %i
+ %45 = OpAccessChain %_ptr_Function_int %a_3 %44
+ %46 = OpLoad %int %45
+ OpStore %b_3 %46
+ %48 = OpAccessChain %_ptr_Function_int %a_3 %12
+ OpStore %48 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_var = OpFunction %void None %1
+ %50 = OpLabel
+ %a_4 = OpVariable %_ptr_Function__arr_int_uint_10 Function %11
+ %i_0 = OpVariable %_ptr_Function_int Function %12
+ %b_4 = OpVariable %_ptr_Function_int Function %12
+ OpStore %i_0 %12
+ %53 = OpLoad %int %i_0
+ %54 = OpAccessChain %_ptr_Function_int %a_4 %53
+ %55 = OpLoad %int %54
+ OpStore %b_4 %55
+ %57 = OpAccessChain %_ptr_Function_int %a_4 %12
+ OpStore %57 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %59 = OpLabel
+ %60 = OpFunctionCall %void %deref_const
+ %61 = OpFunctionCall %void %no_deref_const
+ %62 = OpFunctionCall %void %deref_let
+ %63 = OpFunctionCall %void %no_deref_let
+ %64 = OpFunctionCall %void %deref_var
+ %65 = OpFunctionCall %void %no_deref_var
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/array.wgsl.expected.wgsl b/test/tint/ptr_sugar/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..da29411
--- /dev/null
+++ b/test/tint/ptr_sugar/array.wgsl.expected.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var b = (*(p))[0];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ let i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : array<i32, 10>;
+ let p = &(a);
+ var i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl b/test/tint/ptr_sugar/builtin_struct.wgsl
new file mode 100644
index 0000000..b44cbc4
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl
@@ -0,0 +1,35 @@
+fn deref_modf() {
+ var a = modf(1.5);
+ let p = &a;
+ var fract = (*p).fract;
+ var whole = (*p).whole;
+}
+
+fn no_deref_modf() {
+ var a = modf(1.5);
+ let p = &a;
+ var fract = p.fract;
+ var whole = p.whole;
+}
+
+fn deref_frexp() {
+ var a = frexp(1.5);
+ let p = &a;
+ var fract = (*p).fract;
+ var exp = (*p).exp;
+}
+
+fn no_deref_frexp() {
+ var a = frexp(1.5);
+ let p = &a;
+ var fract = p.fract;
+ var exp = p.exp;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..ec96b03
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.dxc.hlsl
@@ -0,0 +1,40 @@
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+void deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+ return;
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..ec96b03
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.fxc.hlsl
@@ -0,0 +1,40 @@
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+void deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = {0.5f, 1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = {0.75f, 1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+ return;
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.glsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.glsl
new file mode 100644
index 0000000..e852152
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+
+
+void deref_modf() {
+ modf_result_f32 a = modf_result_f32(0.5f, 1.0f);
+ float tint_symbol = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = modf_result_f32(0.5f, 1.0f);
+ float tint_symbol = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32(0.75f, 1);
+ float tint_symbol = a.fract;
+ int tint_symbol_1 = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32(0.75f, 1);
+ float tint_symbol = a.fract;
+ int tint_symbol_1 = a.exp;
+}
+
+void tint_symbol_2() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_2();
+ return;
+}
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.msl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.msl
new file mode 100644
index 0000000..13aa15f
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.msl
@@ -0,0 +1,44 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct modf_result_f32 {
+ float fract;
+ float whole;
+};
+struct frexp_result_f32 {
+ float fract;
+ int exp;
+};
+void deref_modf() {
+ modf_result_f32 a = modf_result_f32{.fract=0.5f, .whole=1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void no_deref_modf() {
+ modf_result_f32 a = modf_result_f32{.fract=0.5f, .whole=1.0f};
+ float fract = a.fract;
+ float whole = a.whole;
+}
+
+void deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32{.fract=0.75f, .exp=1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+void no_deref_frexp() {
+ frexp_result_f32 a = frexp_result_f32{.fract=0.75f, .exp=1};
+ float fract = a.fract;
+ int exp = a.exp;
+}
+
+kernel void tint_symbol() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.spvasm b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..5459222
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.spvasm
@@ -0,0 +1,123 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 66
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref_modf "deref_modf"
+ OpName %__modf_result_f32 "__modf_result_f32"
+ OpMemberName %__modf_result_f32 0 "fract"
+ OpMemberName %__modf_result_f32 1 "whole"
+ OpName %a "a"
+ OpName %fract "fract"
+ OpName %whole "whole"
+ OpName %no_deref_modf "no_deref_modf"
+ OpName %a_0 "a"
+ OpName %fract_0 "fract"
+ OpName %whole_0 "whole"
+ OpName %deref_frexp "deref_frexp"
+ OpName %__frexp_result_f32 "__frexp_result_f32"
+ OpMemberName %__frexp_result_f32 0 "fract"
+ OpMemberName %__frexp_result_f32 1 "exp"
+ OpName %a_1 "a"
+ OpName %fract_1 "fract"
+ OpName %exp "exp"
+ OpName %no_deref_frexp "no_deref_frexp"
+ OpName %a_2 "a"
+ OpName %fract_2 "fract"
+ OpName %exp_0 "exp"
+ OpName %main "main"
+ OpMemberDecorate %__modf_result_f32 0 Offset 0
+ OpMemberDecorate %__modf_result_f32 1 Offset 4
+ OpMemberDecorate %__frexp_result_f32 0 Offset 0
+ OpMemberDecorate %__frexp_result_f32 1 Offset 4
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+%__modf_result_f32 = OpTypeStruct %float %float
+ %float_0_5 = OpConstant %float 0.5
+ %float_1 = OpConstant %float 1
+ %9 = OpConstantComposite %__modf_result_f32 %float_0_5 %float_1
+%_ptr_Function___modf_result_f32 = OpTypePointer Function %__modf_result_f32
+ %12 = OpConstantNull %__modf_result_f32
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_float = OpTypePointer Function %float
+ %19 = OpConstantNull %float
+ %uint_1 = OpConstant %uint 1
+ %int = OpTypeInt 32 1
+%__frexp_result_f32 = OpTypeStruct %float %int
+ %float_0_75 = OpConstant %float 0.75
+ %int_1 = OpConstant %int 1
+ %39 = OpConstantComposite %__frexp_result_f32 %float_0_75 %int_1
+%_ptr_Function___frexp_result_f32 = OpTypePointer Function %__frexp_result_f32
+ %42 = OpConstantNull %__frexp_result_f32
+%_ptr_Function_int = OpTypePointer Function %int
+ %50 = OpConstantNull %int
+ %deref_modf = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function___modf_result_f32 Function %12
+ %fract = OpVariable %_ptr_Function_float Function %19
+ %whole = OpVariable %_ptr_Function_float Function %19
+ OpStore %a %9
+ %16 = OpAccessChain %_ptr_Function_float %a %uint_0
+ %17 = OpLoad %float %16
+ OpStore %fract %17
+ %21 = OpAccessChain %_ptr_Function_float %a %uint_1
+ %22 = OpLoad %float %21
+ OpStore %whole %22
+ OpReturn
+ OpFunctionEnd
+%no_deref_modf = OpFunction %void None %1
+ %25 = OpLabel
+ %a_0 = OpVariable %_ptr_Function___modf_result_f32 Function %12
+ %fract_0 = OpVariable %_ptr_Function_float Function %19
+ %whole_0 = OpVariable %_ptr_Function_float Function %19
+ OpStore %a_0 %9
+ %27 = OpAccessChain %_ptr_Function_float %a_0 %uint_0
+ %28 = OpLoad %float %27
+ OpStore %fract_0 %28
+ %30 = OpAccessChain %_ptr_Function_float %a_0 %uint_1
+ %31 = OpLoad %float %30
+ OpStore %whole_0 %31
+ OpReturn
+ OpFunctionEnd
+%deref_frexp = OpFunction %void None %1
+ %34 = OpLabel
+ %a_1 = OpVariable %_ptr_Function___frexp_result_f32 Function %42
+ %fract_1 = OpVariable %_ptr_Function_float Function %19
+ %exp = OpVariable %_ptr_Function_int Function %50
+ OpStore %a_1 %39
+ %43 = OpAccessChain %_ptr_Function_float %a_1 %uint_0
+ %44 = OpLoad %float %43
+ OpStore %fract_1 %44
+ %47 = OpAccessChain %_ptr_Function_int %a_1 %uint_1
+ %48 = OpLoad %int %47
+ OpStore %exp %48
+ OpReturn
+ OpFunctionEnd
+%no_deref_frexp = OpFunction %void None %1
+ %52 = OpLabel
+ %a_2 = OpVariable %_ptr_Function___frexp_result_f32 Function %42
+ %fract_2 = OpVariable %_ptr_Function_float Function %19
+ %exp_0 = OpVariable %_ptr_Function_int Function %50
+ OpStore %a_2 %39
+ %54 = OpAccessChain %_ptr_Function_float %a_2 %uint_0
+ %55 = OpLoad %float %54
+ OpStore %fract_2 %55
+ %57 = OpAccessChain %_ptr_Function_int %a_2 %uint_1
+ %58 = OpLoad %int %57
+ OpStore %exp_0 %58
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %61 = OpLabel
+ %62 = OpFunctionCall %void %deref_modf
+ %63 = OpFunctionCall %void %no_deref_modf
+ %64 = OpFunctionCall %void %deref_frexp
+ %65 = OpFunctionCall %void %no_deref_frexp
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/builtin_struct.wgsl.expected.wgsl b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..7349c16
--- /dev/null
+++ b/test/tint/ptr_sugar/builtin_struct.wgsl.expected.wgsl
@@ -0,0 +1,35 @@
+fn deref_modf() {
+ var a = modf(1.5);
+ let p = &(a);
+ var fract = (*(p)).fract;
+ var whole = (*(p)).whole;
+}
+
+fn no_deref_modf() {
+ var a = modf(1.5);
+ let p = &(a);
+ var fract = p.fract;
+ var whole = p.whole;
+}
+
+fn deref_frexp() {
+ var a = frexp(1.5);
+ let p = &(a);
+ var fract = (*(p)).fract;
+ var exp = (*(p)).exp;
+}
+
+fn no_deref_frexp() {
+ var a = frexp(1.5);
+ let p = &(a);
+ var fract = p.fract;
+ var exp = p.exp;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_modf();
+ no_deref_modf();
+ deref_frexp();
+ no_deref_frexp();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl b/test/tint/ptr_sugar/compound_assign_index.wgsl
new file mode 100644
index 0000000..ff01d4c
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl
@@ -0,0 +1,31 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ (*p)[0] += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ p[0] += 42;
+}
+
+fn deref_inc() {
+ var a : vec3<i32>;
+ let p = &a;
+ (*p)[0]++;
+}
+
+fn no_deref_inc() {
+ var a : vec3<i32>;
+ let p = &a;
+ p[0]++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..25853e3
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.dxc.hlsl
@@ -0,0 +1,36 @@
+void set_int3(inout int3 vec, int idx, int val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+void deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_1 = 0;
+ set_int3(a, tint_symbol_1, (a[tint_symbol_1] + 42));
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_3 = 0;
+ set_int3(a, tint_symbol_3, (a[tint_symbol_3] + 42));
+}
+
+void deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_5 = 0;
+ set_int3(a, tint_symbol_5, (a[tint_symbol_5] + 1));
+}
+
+void no_deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_7 = 0;
+ set_int3(a, tint_symbol_7, (a[tint_symbol_7] + 1));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..25853e3
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.fxc.hlsl
@@ -0,0 +1,36 @@
+void set_int3(inout int3 vec, int idx, int val) {
+ vec = (idx.xxx == int3(0, 1, 2)) ? val.xxx : vec;
+}
+
+void deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_1 = 0;
+ set_int3(a, tint_symbol_1, (a[tint_symbol_1] + 42));
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_3 = 0;
+ set_int3(a, tint_symbol_3, (a[tint_symbol_3] + 42));
+}
+
+void deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_5 = 0;
+ set_int3(a, tint_symbol_5, (a[tint_symbol_5] + 1));
+}
+
+void no_deref_inc() {
+ int3 a = int3(0, 0, 0);
+ const int tint_symbol_7 = 0;
+ set_int3(a, tint_symbol_7, (a[tint_symbol_7] + 1));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.glsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.glsl
new file mode 100644
index 0000000..dd04e0f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.glsl
@@ -0,0 +1,38 @@
+#version 310 es
+
+void deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_2 = 0;
+ a[tint_symbol_2] = (a[tint_symbol_2] + 42);
+}
+
+void no_deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_4 = 0;
+ a[tint_symbol_4] = (a[tint_symbol_4] + 42);
+}
+
+void deref_inc() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_6 = 0;
+ a[tint_symbol_6] = (a[tint_symbol_6] + 1);
+}
+
+void no_deref_inc() {
+ ivec3 a = ivec3(0, 0, 0);
+ int tint_symbol_8 = 0;
+ a[tint_symbol_8] = (a[tint_symbol_8] + 1);
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.msl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.msl
new file mode 100644
index 0000000..0a18439
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ int3 a = 0;
+ int const tint_symbol_2 = 0;
+ a[tint_symbol_2] = as_type<int>((as_type<uint>(a[tint_symbol_2]) + as_type<uint>(42)));
+}
+
+void no_deref() {
+ int3 a = 0;
+ int const tint_symbol_4 = 0;
+ a[tint_symbol_4] = as_type<int>((as_type<uint>(a[tint_symbol_4]) + as_type<uint>(42)));
+}
+
+void deref_inc() {
+ int3 a = 0;
+ int const tint_symbol_6 = 0;
+ a[tint_symbol_6] = as_type<int>((as_type<uint>(a[tint_symbol_6]) + as_type<uint>(1)));
+}
+
+void no_deref_inc() {
+ int3 a = 0;
+ int const tint_symbol_8 = 0;
+ a[tint_symbol_8] = as_type<int>((as_type<uint>(a[tint_symbol_8]) + as_type<uint>(1)));
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.spvasm b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.spvasm
new file mode 100644
index 0000000..237c95a
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.spvasm
@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 45
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %deref_inc "deref_inc"
+ OpName %a_1 "a"
+ OpName %no_deref_inc "no_deref_inc"
+ OpName %a_2 "a"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %10 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+ %int_1 = OpConstant %int 1
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %12 = OpAccessChain %_ptr_Function_int %a %10
+ %13 = OpAccessChain %_ptr_Function_int %a %10
+ %14 = OpLoad %int %13
+ %16 = OpIAdd %int %14 %int_42
+ OpStore %12 %16
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %18 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %20 = OpAccessChain %_ptr_Function_int %a_0 %10
+ %21 = OpAccessChain %_ptr_Function_int %a_0 %10
+ %22 = OpLoad %int %21
+ %23 = OpIAdd %int %22 %int_42
+ OpStore %20 %23
+ OpReturn
+ OpFunctionEnd
+ %deref_inc = OpFunction %void None %1
+ %25 = OpLabel
+ %a_1 = OpVariable %_ptr_Function_v3int Function %9
+ %27 = OpAccessChain %_ptr_Function_int %a_1 %10
+ %28 = OpAccessChain %_ptr_Function_int %a_1 %10
+ %29 = OpLoad %int %28
+ %31 = OpIAdd %int %29 %int_1
+ OpStore %27 %31
+ OpReturn
+ OpFunctionEnd
+%no_deref_inc = OpFunction %void None %1
+ %33 = OpLabel
+ %a_2 = OpVariable %_ptr_Function_v3int Function %9
+ %35 = OpAccessChain %_ptr_Function_int %a_2 %10
+ %36 = OpAccessChain %_ptr_Function_int %a_2 %10
+ %37 = OpLoad %int %36
+ %38 = OpIAdd %int %37 %int_1
+ OpStore %35 %38
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %deref
+ %42 = OpFunctionCall %void %no_deref
+ %43 = OpFunctionCall %void %deref_inc
+ %44 = OpFunctionCall %void %no_deref_inc
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.wgsl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.wgsl
new file mode 100644
index 0000000..9fad9aa
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.wgsl
@@ -0,0 +1,31 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ (*(p))[0] += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ p[0] += 42;
+}
+
+fn deref_inc() {
+ var a : vec3<i32>;
+ let p = &(a);
+ (*(p))[0]++;
+}
+
+fn no_deref_inc() {
+ var a : vec3<i32>;
+ let p = &(a);
+ p[0]++;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+ deref_inc();
+ no_deref_inc();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl b/test/tint/ptr_sugar/compound_assign_member.wgsl
new file mode 100644
index 0000000..2785bd5
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl
@@ -0,0 +1,17 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ (*p).x += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ p.x += 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..630d44f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.dxc.hlsl
@@ -0,0 +1,16 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..630d44f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.fxc.hlsl
@@ -0,0 +1,16 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.glsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.glsl
new file mode 100644
index 0000000..62a3b1f
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.glsl
@@ -0,0 +1,22 @@
+#version 310 es
+
+void deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void no_deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ a.x = (a.x + 42);
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.msl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.msl
new file mode 100644
index 0000000..65e3c9e
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.msl
@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ int3 a = 0;
+ a[0] = as_type<int>((as_type<uint>(a[0]) + as_type<uint>(42)));
+}
+
+void no_deref() {
+ int3 a = 0;
+ a[0] = as_type<int>((as_type<uint>(a[0]) + as_type<uint>(42)));
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.spvasm b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.spvasm
new file mode 100644
index 0000000..7e3b8ff
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.spvasm
@@ -0,0 +1,50 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %13 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %14 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %15 = OpLoad %int %14
+ %17 = OpIAdd %int %15 %int_42
+ OpStore %13 %17
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %19 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %21 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %23 = OpLoad %int %22
+ %24 = OpIAdd %int %23 %int_42
+ OpStore %21 %24
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %26 = OpLabel
+ %27 = OpFunctionCall %void %deref
+ %28 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.wgsl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.wgsl
new file mode 100644
index 0000000..3a01d23
--- /dev/null
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.wgsl
@@ -0,0 +1,17 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ (*(p)).x += 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ p.x += 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl b/test/tint/ptr_sugar/matrix.wgsl
new file mode 100644
index 0000000..5a768d9
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : mat2x3<f32>;
+ let p = &a;
+ var b = (*p)[0];
+ (*p)[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+fn no_deref() {
+ var a : mat2x3<f32>;
+ let p = &a;
+ var b = p[0];
+ p[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/matrix.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..7790b20
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/matrix.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7790b20
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ float2x3 a = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.glsl b/test/tint/ptr_sugar/matrix.wgsl.expected.glsl
new file mode 100644
index 0000000..25cae38
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+void deref() {
+ mat2x3 a = mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ vec3 b = a[0];
+ a[0] = vec3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ mat2x3 a = mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+ vec3 b = a[0];
+ a[0] = vec3(1.0f, 2.0f, 3.0f);
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.msl b/test/tint/ptr_sugar/matrix.wgsl.expected.msl
new file mode 100644
index 0000000..65ab860
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ float2x3 a = float2x3(0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+void no_deref() {
+ float2x3 a = float2x3(0.0f);
+ float3 b = a[0];
+ a[0] = float3(1.0f, 2.0f, 3.0f);
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.spvasm b/test/tint/ptr_sugar/matrix.wgsl.expected.spvasm
new file mode 100644
index 0000000..4f1e62d
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.spvasm
@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 34
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %float = OpTypeFloat 32
+ %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Function_mat2v3float = OpTypePointer Function %mat2v3float
+ %10 = OpConstantNull %mat2v3float
+ %int = OpTypeInt 32 1
+ %12 = OpConstantNull %int
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+ %17 = OpConstantNull %v3float
+ %float_1 = OpConstant %float 1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %22 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_mat2v3float Function %10
+ %b = OpVariable %_ptr_Function_v3float Function %17
+ %14 = OpAccessChain %_ptr_Function_v3float %a %12
+ %15 = OpLoad %v3float %14
+ OpStore %b %15
+ %18 = OpAccessChain %_ptr_Function_v3float %a %12
+ OpStore %18 %22
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %24 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_mat2v3float Function %10
+ %b_0 = OpVariable %_ptr_Function_v3float Function %17
+ %26 = OpAccessChain %_ptr_Function_v3float %a_0 %12
+ %27 = OpLoad %v3float %26
+ OpStore %b_0 %27
+ %29 = OpAccessChain %_ptr_Function_v3float %a_0 %12
+ OpStore %29 %22
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %31 = OpLabel
+ %32 = OpFunctionCall %void %deref
+ %33 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/matrix.wgsl.expected.wgsl b/test/tint/ptr_sugar/matrix.wgsl.expected.wgsl
new file mode 100644
index 0000000..3b18d24
--- /dev/null
+++ b/test/tint/ptr_sugar/matrix.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : mat2x3<f32>;
+ let p = &(a);
+ var b = (*(p))[0];
+ (*(p))[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+fn no_deref() {
+ var a : mat2x3<f32>;
+ let p = &(a);
+ var b = p[0];
+ p[0] = vec3<f32>(1.0, 2.0, 3.0);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl b/test/tint/ptr_sugar/struct.wgsl
new file mode 100644
index 0000000..1c6c7fe
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl
@@ -0,0 +1,23 @@
+struct S {
+ x : i32,
+}
+
+fn deref() {
+ var a : S;
+ let p = &a;
+ var b = (*p).x;
+ (*p).x = 42;
+}
+
+fn no_deref() {
+ var a : S;
+ let p = &a;
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/struct.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..628808f
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.dxc.hlsl
@@ -0,0 +1,22 @@
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/struct.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..628808f
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.fxc.hlsl
@@ -0,0 +1,22 @@
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = (S)0;
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.glsl b/test/tint/ptr_sugar/struct.wgsl.expected.glsl
new file mode 100644
index 0000000..dceeb2c
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.glsl
@@ -0,0 +1,28 @@
+#version 310 es
+
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = S(0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = S(0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.msl b/test/tint/ptr_sugar/struct.wgsl.expected.msl
new file mode 100644
index 0000000..1b6ebcd
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ int x;
+};
+
+void deref() {
+ S a = {};
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ S a = {};
+ int b = a.x;
+ a.x = 42;
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.spvasm b/test/tint/ptr_sugar/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..d4c38c9
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.spvasm
@@ -0,0 +1,58 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %S "S"
+ OpMemberName %S 0 "x"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %main "main"
+ OpMemberDecorate %S 0 Offset 0
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int
+%_ptr_Function_S = OpTypePointer Function %S
+ %9 = OpConstantNull %S
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+ %16 = OpConstantNull %int
+ %int_42 = OpConstant %int 42
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_S Function %9
+ %b = OpVariable %_ptr_Function_int Function %16
+ %13 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %14 = OpLoad %int %13
+ OpStore %b %14
+ %17 = OpAccessChain %_ptr_Function_int %a %uint_0
+ OpStore %17 %int_42
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %20 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_S Function %9
+ %b_0 = OpVariable %_ptr_Function_int Function %16
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %23 = OpLoad %int %22
+ OpStore %b_0 %23
+ %25 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ OpStore %25 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %27 = OpLabel
+ %28 = OpFunctionCall %void %deref
+ %29 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/struct.wgsl.expected.wgsl b/test/tint/ptr_sugar/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..e01aecd
--- /dev/null
+++ b/test/tint/ptr_sugar/struct.wgsl.expected.wgsl
@@ -0,0 +1,23 @@
+struct S {
+ x : i32,
+}
+
+fn deref() {
+ var a : S;
+ let p = &(a);
+ var b = (*(p)).x;
+ (*(p)).x = 42;
+}
+
+fn no_deref() {
+ var a : S;
+ let p = &(a);
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl b/test/tint/ptr_sugar/vector_index.wgsl
new file mode 100644
index 0000000..c4a549a
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = (*p)[0];
+ (*p)[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : vec3<i32>;
+ let p = &a;
+ let i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : vec3<i32>;
+ let p = &a;
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : vec3<i32>;
+ let p = &a;
+ var i = 0;
+ var b = (*p)[i];
+ (*p)[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : vec3<i32>;
+ let p = &a;
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..7d0de2e
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.dxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int3 a = int3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7d0de2e
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.fxc.hlsl
@@ -0,0 +1,50 @@
+void deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int3 a = int3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int3 a = int3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int3 a = int3(0, 0, 0);
+ const int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.glsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.glsl
new file mode 100644
index 0000000..b78430b
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.glsl
@@ -0,0 +1,56 @@
+#version 310 es
+
+void deref_const() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ ivec3 a = ivec3(0, 0, 0);
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.msl b/test/tint/ptr_sugar/vector_index.wgsl.expected.msl
new file mode 100644
index 0000000..2d99f38
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.msl
@@ -0,0 +1,53 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref_const() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref_const() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void deref_let() {
+ int3 a = 0;
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_let() {
+ int3 a = 0;
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void deref_var() {
+ int3 a = 0;
+ int i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+void no_deref_var() {
+ int3 a = 0;
+ int const i = 0;
+ int b = a[i];
+ a[0] = 42;
+}
+
+kernel void tint_symbol() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.spvasm b/test/tint/ptr_sugar/vector_index.wgsl.expected.spvasm
new file mode 100644
index 0000000..7706621
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.spvasm
@@ -0,0 +1,117 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 62
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref_const "deref_const"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref_const "no_deref_const"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %deref_let "deref_let"
+ OpName %a_1 "a"
+ OpName %b_1 "b"
+ OpName %no_deref_let "no_deref_let"
+ OpName %a_2 "a"
+ OpName %b_2 "b"
+ OpName %deref_var "deref_var"
+ OpName %a_3 "a"
+ OpName %i "i"
+ OpName %b_3 "b"
+ OpName %no_deref_var "no_deref_var"
+ OpName %a_4 "a"
+ OpName %b_4 "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %10 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_42 = OpConstant %int 42
+%deref_const = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %b = OpVariable %_ptr_Function_int Function %10
+ %12 = OpAccessChain %_ptr_Function_int %a %10
+ %13 = OpLoad %int %12
+ OpStore %b %13
+ %15 = OpAccessChain %_ptr_Function_int %a %10
+ OpStore %15 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_const = OpFunction %void None %1
+ %18 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %b_0 = OpVariable %_ptr_Function_int Function %10
+ %20 = OpAccessChain %_ptr_Function_int %a_0 %10
+ %21 = OpLoad %int %20
+ OpStore %b_0 %21
+ %23 = OpAccessChain %_ptr_Function_int %a_0 %10
+ OpStore %23 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_let = OpFunction %void None %1
+ %25 = OpLabel
+ %a_1 = OpVariable %_ptr_Function_v3int Function %9
+ %b_1 = OpVariable %_ptr_Function_int Function %10
+ %27 = OpAccessChain %_ptr_Function_int %a_1 %10
+ %28 = OpLoad %int %27
+ OpStore %b_1 %28
+ %30 = OpAccessChain %_ptr_Function_int %a_1 %10
+ OpStore %30 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_let = OpFunction %void None %1
+ %32 = OpLabel
+ %a_2 = OpVariable %_ptr_Function_v3int Function %9
+ %b_2 = OpVariable %_ptr_Function_int Function %10
+ %34 = OpAccessChain %_ptr_Function_int %a_2 %10
+ %35 = OpLoad %int %34
+ OpStore %b_2 %35
+ %37 = OpAccessChain %_ptr_Function_int %a_2 %10
+ OpStore %37 %int_42
+ OpReturn
+ OpFunctionEnd
+ %deref_var = OpFunction %void None %1
+ %39 = OpLabel
+ %a_3 = OpVariable %_ptr_Function_v3int Function %9
+ %i = OpVariable %_ptr_Function_int Function %10
+ %b_3 = OpVariable %_ptr_Function_int Function %10
+ OpStore %i %10
+ %42 = OpLoad %int %i
+ %43 = OpAccessChain %_ptr_Function_int %a_3 %42
+ %44 = OpLoad %int %43
+ OpStore %b_3 %44
+ %46 = OpAccessChain %_ptr_Function_int %a_3 %10
+ OpStore %46 %int_42
+ OpReturn
+ OpFunctionEnd
+%no_deref_var = OpFunction %void None %1
+ %48 = OpLabel
+ %a_4 = OpVariable %_ptr_Function_v3int Function %9
+ %b_4 = OpVariable %_ptr_Function_int Function %10
+ %50 = OpAccessChain %_ptr_Function_int %a_4 %10
+ %51 = OpLoad %int %50
+ OpStore %b_4 %51
+ %53 = OpAccessChain %_ptr_Function_int %a_4 %10
+ OpStore %53 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %55 = OpLabel
+ %56 = OpFunctionCall %void %deref_const
+ %57 = OpFunctionCall %void %no_deref_const
+ %58 = OpFunctionCall %void %deref_let
+ %59 = OpFunctionCall %void %no_deref_let
+ %60 = OpFunctionCall %void %deref_var
+ %61 = OpFunctionCall %void %no_deref_var
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/vector_index.wgsl.expected.wgsl b/test/tint/ptr_sugar/vector_index.wgsl.expected.wgsl
new file mode 100644
index 0000000..28f45ae
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_index.wgsl.expected.wgsl
@@ -0,0 +1,55 @@
+fn deref_const() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = (*(p))[0];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_const() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = p[0];
+ p[0] = 42;
+}
+
+fn deref_let() {
+ var a : vec3<i32>;
+ let p = &(a);
+ let i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_let() {
+ var a : vec3<i32>;
+ let p = &(a);
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+fn deref_var() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var i = 0;
+ var b = (*(p))[i];
+ (*(p))[0] = 42;
+}
+
+fn no_deref_var() {
+ var a : vec3<i32>;
+ let p = &(a);
+ let i = 0;
+ var b = p[i];
+ p[0] = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref_const();
+ no_deref_const();
+ deref_let();
+ no_deref_let();
+ deref_var();
+ no_deref_var();
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl b/test/tint/ptr_sugar/vector_member.wgsl
new file mode 100644
index 0000000..cd2eab8
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = (*p).x;
+ (*p).x = 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &a;
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b9ff0dc
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.dxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b9ff0dc
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.fxc.hlsl
@@ -0,0 +1,18 @@
+void deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ int3 a = int3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.glsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.glsl
new file mode 100644
index 0000000..4380434
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+void deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void no_deref() {
+ ivec3 a = ivec3(0, 0, 0);
+ int b = a.x;
+ a.x = 42;
+}
+
+void tint_symbol() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.msl b/test/tint/ptr_sugar/vector_member.wgsl.expected.msl
new file mode 100644
index 0000000..b4b88f1
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.msl
@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+void no_deref() {
+ int3 a = 0;
+ int b = a[0];
+ a[0] = 42;
+}
+
+kernel void tint_symbol() {
+ deref();
+ no_deref();
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.spvasm b/test/tint/ptr_sugar/vector_member.wgsl.expected.spvasm
new file mode 100644
index 0000000..90d445f
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.spvasm
@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %deref "deref"
+ OpName %a "a"
+ OpName %b "b"
+ OpName %no_deref "no_deref"
+ OpName %a_0 "a"
+ OpName %b_0 "b"
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+ %9 = OpConstantNull %v3int
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+ %16 = OpConstantNull %int
+ %int_42 = OpConstant %int 42
+ %deref = OpFunction %void None %1
+ %4 = OpLabel
+ %a = OpVariable %_ptr_Function_v3int Function %9
+ %b = OpVariable %_ptr_Function_int Function %16
+ %13 = OpAccessChain %_ptr_Function_int %a %uint_0
+ %14 = OpLoad %int %13
+ OpStore %b %14
+ %17 = OpAccessChain %_ptr_Function_int %a %uint_0
+ OpStore %17 %int_42
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %1
+ %20 = OpLabel
+ %a_0 = OpVariable %_ptr_Function_v3int Function %9
+ %b_0 = OpVariable %_ptr_Function_int Function %16
+ %22 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ %23 = OpLoad %int %22
+ OpStore %b_0 %23
+ %25 = OpAccessChain %_ptr_Function_int %a_0 %uint_0
+ OpStore %25 %int_42
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %1
+ %27 = OpLabel
+ %28 = OpFunctionCall %void %deref
+ %29 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/vector_member.wgsl.expected.wgsl b/test/tint/ptr_sugar/vector_member.wgsl.expected.wgsl
new file mode 100644
index 0000000..4c48f2b
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_member.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+fn deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = (*(p)).x;
+ (*(p)).x = 42;
+}
+
+fn no_deref() {
+ var a : vec3<i32>;
+ let p = &(a);
+ var b = p.x;
+ p.x = 42;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
