[tint] Add sem::Load for swizzle from a pointer
Without the implicit load, conversion to IR fails (found by fuzzer in
the linked bug). This also causes the alias analysis and uniformity
analysis to miss certain errors.
Unlike all other cases where the load rule is invoked, the type of the
source expression here is a pointer, since there is also an implicit
dereference. This change allows `sem::Load` to have a pointer
expression as its source type to support this.
Bug: 353039528
Change-Id: I59513509468bbddb6e638d4c75bf566f257c8c0a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/199176
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc b/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
index 816e014..941b42d 100644
--- a/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
+++ b/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
@@ -467,7 +467,7 @@
[&](const sem::Load* load) {
// Unpack loads of types that contain vec3s in host-shareable address spaces.
if (ContainsVec3(load->Type()) &&
- core::IsHostShareable(load->ReferenceType()->AddressSpace())) {
+ core::IsHostShareable(load->MemoryView()->AddressSpace())) {
to_unpack.Add(load);
}
},
diff --git a/src/tint/lang/spirv/reader/ast_lower/atomics.cc b/src/tint/lang/spirv/reader/ast_lower/atomics.cc
index 8b74df7..ea2b778 100644
--- a/src/tint/lang/spirv/reader/ast_lower/atomics.cc
+++ b/src/tint/lang/spirv/reader/ast_lower/atomics.cc
@@ -261,9 +261,9 @@
// with atomicLoad and atomicStore.
for (auto* node : ctx.src->ASTNodes().Objects()) {
if (auto* load = ctx.src->Sem().Get<sem::Load>(node)) {
- if (is_ref_to_atomic_var(load->Reference())) {
- ctx.Replace(load->Reference()->Declaration(), [=] {
- auto* expr = ctx.CloneWithoutTransform(load->Reference()->Declaration());
+ if (is_ref_to_atomic_var(load->Source())) {
+ ctx.Replace(load->Source()->Declaration(), [=] {
+ auto* expr = ctx.CloneWithoutTransform(load->Source()->Declaration());
return b.Call(wgsl::BuiltinFn::kAtomicLoad, b.AddressOf(expr));
});
}
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index cda00e9..4c2c6fd 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -520,11 +520,11 @@
}
}
if (auto* load = expr->As<sem::Load>()) {
- auto ref_id = GenerateExpression(load->Reference());
+ auto ref_id = GenerateExpression(load->Source());
if (ref_id == 0) {
return 0;
}
- return GenerateLoad(load->ReferenceType(), ref_id);
+ return GenerateLoad(load->MemoryView(), ref_id);
}
return Switch(
expr->Declaration(), //
@@ -1120,7 +1120,7 @@
"' does not resolve to a variable";
}
-uint32_t Builder::GenerateLoad(const core::type::Reference* type, uint32_t id) {
+uint32_t Builder::GenerateLoad(const core::type::MemoryView* type, uint32_t id) {
auto type_id = GenerateTypeIfNeeded(type->StoreType());
auto result = result_op();
auto result_id = std::get<uint32_t>(result);
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.h b/src/tint/lang/spirv/writer/ast_printer/builder.h
index 14e03d4..b54fd21 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.h
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.h
@@ -386,7 +386,7 @@
/// @param type the reference type of the expression
/// @param id the SPIR-V id of the expression
/// @returns the ID of the loaded value or 0 on failure.
- uint32_t GenerateLoad(const core::type::Reference* type, uint32_t id);
+ uint32_t GenerateLoad(const core::type::MemoryView* type, uint32_t id);
/// Generates an OpLoad on the given ID if it has reference type in WGSL, otherwise return the
/// ID itself.
/// @param type the type of the expression
diff --git a/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc b/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc
index ad4d713..f47fce6 100644
--- a/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc
+++ b/src/tint/lang/wgsl/ast/transform/simplify_pointers.cc
@@ -28,6 +28,7 @@
#include "src/tint/lang/wgsl/ast/transform/simplify_pointers.h"
#include <unordered_set>
+#include "src/tint/lang/wgsl/sem/load.h"
#include "src/tint/utils/containers/hashset.h"
#include "src/tint/lang/wgsl/ast/transform/unshadow.h"
@@ -134,13 +135,17 @@
break;
}
}
- if (auto* user = ctx.src->Sem().Get<sem::VariableUser>(op.expr)) {
- auto* var = user->Variable();
- if (var->Is<sem::LocalVariable>() && //
- var->Declaration()->Is<Let>() && //
- var->Type()->Is<core::type::Pointer>()) {
- op.expr = var->Declaration()->initializer;
- continue;
+
+ if (auto* sem = ctx.src->Sem().Get<sem::ValueExpression>(op.expr)) {
+ // There may be an implicit load before the identifier due to a swizzle on pointer.
+ if (auto* user = sem->UnwrapLoad()->As<sem::VariableUser>()) {
+ auto* var = user->Variable();
+ if (var->Is<sem::LocalVariable>() && //
+ var->Declaration()->Is<Let>() && //
+ var->Type()->Is<core::type::Pointer>()) {
+ op.expr = var->Declaration()->initializer;
+ continue;
+ }
}
}
return op;
@@ -236,7 +241,8 @@
},
[&](const AccessorExpression* accessor) {
if (auto* a = ctx.src->Sem().Get<sem::ValueExpression>(accessor->object)) {
- if (a->Type()->Is<core::type::Pointer>()) {
+ // There may be an implicit load if this is a swizzle.
+ if (a->UnwrapLoad()->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 1c92b18..9af6f5b 100644
--- a/src/tint/lang/wgsl/ast/transform/simplify_pointers_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/simplify_pointers_test.cc
@@ -513,5 +513,26 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(SimplifyPointersTest, SwizzleFromPointer) {
+ auto* src = R"(
+fn f() {
+ var a : vec4f;
+ let p : ptr<function, vec4f> = &a;
+ let v : vec2f = p.yw;
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var a : vec4f;
+ let v : vec2f = a.yw;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace tint::ast::transform
diff --git a/src/tint/lang/wgsl/resolver/alias_analysis_test.cc b/src/tint/lang/wgsl/resolver/alias_analysis_test.cc
index e545670..047d9e2 100644
--- a/src/tint/lang/wgsl/resolver/alias_analysis_test.cc
+++ b/src/tint/lang/wgsl/resolver/alias_analysis_test.cc
@@ -850,6 +850,37 @@
12:34 note: aliases with another argument passed here)");
}
+TEST_F(ResolverAliasAnalysisTest, Read_MultiComponentSwizzle_FromPointer) {
+ // fn f2(p1 : ptr<function, vec4<f32>, p2 : ptr<function, vec4<f32>) {
+ // _ = p2.zy;
+ // *p1 = vec4<f32>();
+ // }
+ // fn f1() {
+ // var v : vec4<f32>;
+ // f2(&v, &v);
+ // }
+ Structure("S", Vector{Member("a", ty.i32())});
+ Func("f2",
+ Vector{
+ Param("p1", ty.ptr<function, vec4<f32>>()),
+ Param("p2", ty.ptr<function, vec4<f32>>()),
+ },
+ ty.void_(),
+ Vector{
+ Assign(Phony(), MemberAccessor("p2", "zy")),
+ Assign(Deref("p1"), Call<vec4<f32>>()),
+ });
+ Func("f1", tint::Empty, ty.void_(),
+ Vector{
+ Decl(Var("v", ty.vec4<f32>())),
+ CallStmt(
+ Call("f2", AddressOf(Source{{12, 34}}, "v"), AddressOf(Source{{56, 76}}, "v"))),
+ });
+ EXPECT_FALSE(r()->Resolve()) << r()->error();
+ EXPECT_EQ(r()->error(), R"(56:76 error: invalid aliased pointer argument
+12:34 note: aliases with another argument passed here)");
+}
+
TEST_F(ResolverAliasAnalysisTest, SinglePointerReadWrite) {
// Test that we can both read and write from a single pointer parameter.
//
diff --git a/src/tint/lang/wgsl/resolver/load_test.cc b/src/tint/lang/wgsl/resolver/load_test.cc
index b4234a0..4eeee7b 100644
--- a/src/tint/lang/wgsl/resolver/load_test.cc
+++ b/src/tint/lang/wgsl/resolver/load_test.cc
@@ -52,8 +52,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, LetInitializer) {
@@ -67,8 +67,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, Assignment) {
@@ -84,8 +84,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, CompoundAssignment) {
@@ -101,8 +101,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, UnaryOp) {
@@ -116,8 +116,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, UnaryOp_NoLoad) {
@@ -145,8 +145,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, Index) {
@@ -160,8 +160,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, MultiComponentSwizzle) {
@@ -175,8 +175,25 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::Vector>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::Vector>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::Vector>());
+}
+
+TEST_F(ResolverLoadTest, MultiComponentSwizzle_FromPointer) {
+ // var ref = vec4(1);
+ // let ptr = &ref;
+ // var v = ptr.xyz;
+ auto* ident = Expr("ptr");
+ WrapInFunction(Var("ref", Call<vec4<i32>>(1_i)), //
+ Let("ptr", AddressOf("ref")), //
+ Var("v", MemberAccessor(ident, "xyz")));
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+ auto* load = Sem().Get<sem::Load>(ident);
+ ASSERT_NE(load, nullptr);
+ EXPECT_TRUE(load->Type()->Is<core::type::Vector>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Pointer>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapPtr()->Is<core::type::Vector>());
}
TEST_F(ResolverLoadTest, Bitcast) {
@@ -190,8 +207,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::F32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::F32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::F32>());
}
TEST_F(ResolverLoadTest, BuiltinArg) {
@@ -205,8 +222,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::F32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::F32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::F32>());
}
TEST_F(ResolverLoadTest, FunctionArg) {
@@ -222,8 +239,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::F32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::F32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::F32>());
}
TEST_F(ResolverLoadTest, FunctionArg_Handles) {
@@ -255,15 +272,15 @@
auto* load = Sem().Get<sem::Load>(t_ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::SampledTexture>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::SampledTexture>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::SampledTexture>());
}
{
auto* load = Sem().Get<sem::Load>(s_ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::Sampler>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::Sampler>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::Sampler>());
}
}
@@ -281,8 +298,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::F32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::F32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::F32>());
}
TEST_F(ResolverLoadTest, IfCond) {
@@ -296,8 +313,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::Bool>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::Bool>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::Bool>());
}
TEST_F(ResolverLoadTest, Switch) {
@@ -313,8 +330,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::I32>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::I32>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::I32>());
}
TEST_F(ResolverLoadTest, BreakIfCond) {
@@ -332,8 +349,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::Bool>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::Bool>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::Bool>());
}
TEST_F(ResolverLoadTest, ForCond) {
@@ -347,8 +364,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::Bool>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::Bool>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::Bool>());
}
TEST_F(ResolverLoadTest, WhileCond) {
@@ -362,8 +379,8 @@
auto* load = Sem().Get<sem::Load>(ident);
ASSERT_NE(load, nullptr);
EXPECT_TRUE(load->Type()->Is<core::type::Bool>());
- EXPECT_TRUE(load->Reference()->Type()->Is<core::type::Reference>());
- EXPECT_TRUE(load->Reference()->Type()->UnwrapRef()->Is<core::type::Bool>());
+ EXPECT_TRUE(load->Source()->Type()->Is<core::type::Reference>());
+ EXPECT_TRUE(load->Source()->Type()->UnwrapRef()->Is<core::type::Bool>());
}
TEST_F(ResolverLoadTest, AddressOf) {
diff --git a/src/tint/lang/wgsl/resolver/ptr_ref_test.cc b/src/tint/lang/wgsl/resolver/ptr_ref_test.cc
index 892dd90..ba1dd8e 100644
--- a/src/tint/lang/wgsl/resolver/ptr_ref_test.cc
+++ b/src/tint/lang/wgsl/resolver/ptr_ref_test.cc
@@ -71,7 +71,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -139,7 +139,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -160,7 +160,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -198,7 +198,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -219,7 +219,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -257,7 +257,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -278,7 +278,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -316,7 +316,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -340,7 +340,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -385,7 +385,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
@@ -410,7 +410,7 @@
auto* load = Sem().Get<sem::Load>(expr);
ASSERT_NE(load, nullptr);
- auto* ref = load->Reference();
+ auto* ref = load->Source();
ASSERT_NE(ref, nullptr);
ASSERT_TRUE(ref->Type()->Is<core::type::Reference>());
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index aeba08d..a5e89d8 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -3529,8 +3529,25 @@
// the swizzle.
ty = b.create<core::type::Vector>(vec->type(), static_cast<uint32_t>(size));
- // The load rule is invoked before the swizzle, if necessary.
- obj_expr = Load(object);
+ if (obj_expr->Type()->Is<core::type::Pointer>()) {
+ // If the LHS is a pointer, the load rule is invoked. We special case this
+ // because our usual handling of implicit loads assumes the expression has
+ // reference type. This expression also has an implicit dereference before the
+ // load, but we have no way of representing that, so we create the load directly
+ // from the pointer expression.
+ auto* load =
+ b.create<sem::Load>(obj_expr, current_statement_, obj_expr->Stage());
+ load->Behaviors() = obj_expr->Behaviors();
+ b.Sem().Replace(obj_expr->Declaration(), load);
+
+ // Register the load for the alias analysis.
+ RegisterLoad(obj_expr);
+
+ obj_expr = load;
+ } else {
+ // The load rule is invoked before the swizzle, if necessary.
+ obj_expr = Load(obj_expr);
+ }
}
const core::constant::Value* val = nullptr;
if (auto* obj_val = object->ConstantValue()) {
diff --git a/src/tint/lang/wgsl/resolver/uniformity_test.cc b/src/tint/lang/wgsl/resolver/uniformity_test.cc
index f4f7614..022998c 100644
--- a/src/tint/lang/wgsl/resolver/uniformity_test.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity_test.cc
@@ -5486,7 +5486,7 @@
TEST_F(UniformityAnalysisTest, VectorElement_NonUniform) {
std::string src = R"(
-@group(0) @binding(0) var<storage, read_write> v : array<i32>;
+@group(0) @binding(0) var<storage, read_write> v : vec4<i32>;
fn foo() {
if (v[2] == 0) {
@@ -5511,6 +5511,60 @@
)");
}
+TEST_F(UniformityAnalysisTest, VectorSwizzle_NonUniform) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> v : vec4<i32>;
+
+fn foo() {
+ if (any(v.xy == vec2())) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:5:3 note: control flow depends on possibly non-uniform value
+ if (any(v.xy == vec2())) {
+ ^^
+
+test:5:11 note: reading from read_write storage buffer 'v' may result in a non-uniform value
+ if (any(v.xy == vec2())) {
+ ^
+)");
+}
+
+TEST_F(UniformityAnalysisTest, VectorSwizzleFromPointer_NonUniform) {
+ std::string src = R"(
+@group(0) @binding(0) var<storage, read_write> v : vec4<i32>;
+
+fn foo() {
+ if (any((&v).xy == vec2())) {
+ workgroupBarrier();
+ }
+}
+)";
+
+ RunTest(src, false);
+ EXPECT_EQ(error_,
+ R"(test:6:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ workgroupBarrier();
+ ^^^^^^^^^^^^^^^^
+
+test:5:3 note: control flow depends on possibly non-uniform value
+ if (any((&v).xy == vec2())) {
+ ^^
+
+test:5:13 note: reading from read_write storage buffer 'v' may result in a non-uniform value
+ if (any((&v).xy == vec2())) {
+ ^
+)");
+}
+
TEST_F(UniformityAnalysisTest, VectorElement_BecomesNonUniform_BeforeCondition) {
std::string src = R"(
@group(0) @binding(0) var<storage, read_write> rw : i32;
diff --git a/src/tint/lang/wgsl/sem/load.cc b/src/tint/lang/wgsl/sem/load.cc
index 636c24e..41eeba9 100644
--- a/src/tint/lang/wgsl/sem/load.cc
+++ b/src/tint/lang/wgsl/sem/load.cc
@@ -33,16 +33,16 @@
TINT_INSTANTIATE_TYPEINFO(tint::sem::Load);
namespace tint::sem {
-Load::Load(const ValueExpression* ref, const Statement* statement, core::EvaluationStage stage)
- : Base(/* declaration */ ref->Declaration(),
- /* type */ ref->Type()->UnwrapRef(),
+Load::Load(const ValueExpression* src, const Statement* statement, core::EvaluationStage stage)
+ : Base(/* declaration */ src->Declaration(),
+ /* type */ src->Type()->UnwrapPtrOrRef(),
/* stage */ stage,
/* statement */ statement,
/* constant */ nullptr,
- /* has_side_effects */ ref->HasSideEffects(),
- /* root_ident */ ref->RootIdentifier()),
- reference_(ref) {
- TINT_ASSERT(ref->Type()->Is<core::type::Reference>());
+ /* has_side_effects */ src->HasSideEffects(),
+ /* root_ident */ src->RootIdentifier()),
+ source_(src) {
+ TINT_ASSERT(src->Type()->Is<core::type::MemoryView>());
}
Load::~Load() = default;
diff --git a/src/tint/lang/wgsl/sem/load.h b/src/tint/lang/wgsl/sem/load.h
index 1711a25..f8a41e2 100644
--- a/src/tint/lang/wgsl/sem/load.h
+++ b/src/tint/lang/wgsl/sem/load.h
@@ -33,30 +33,30 @@
namespace tint::sem {
-/// Load is a semantic expression which represents the load of a reference to a non-reference value.
+/// Load is a semantic expression which represents the load of a memory view to a value.
/// Loads from reference types are implicit in WGSL, so the Load semantic node shares the same AST
/// node as the inner semantic node.
class Load final : public Castable<Load, ValueExpression> {
public:
/// Constructor
- /// @param reference the reference expression being loaded
+ /// @param source the source expression being loaded from
/// @param statement the statement that owns this expression
/// @param stage the earliest evaluation stage for the expression
- Load(const ValueExpression* reference, const Statement* statement, core::EvaluationStage stage);
+ Load(const ValueExpression* source, const Statement* statement, core::EvaluationStage stage);
/// Destructor
~Load() override;
- /// @return the reference being loaded
- const ValueExpression* Reference() const { return reference_; }
+ /// @return the source object being loaded
+ const ValueExpression* Source() const { return source_; }
- /// @returns the type of the loaded reference.
- const core::type::Reference* ReferenceType() const {
- return static_cast<const core::type::Reference*>(reference_->Type());
+ /// @returns the type of the memory view being loaded from.
+ const core::type::MemoryView* MemoryView() const {
+ return static_cast<const core::type::MemoryView*>(source_->Type());
}
private:
- ValueExpression const* const reference_;
+ ValueExpression const* const source_;
};
} // namespace tint::sem
diff --git a/src/tint/lang/wgsl/sem/value_expression.cc b/src/tint/lang/wgsl/sem/value_expression.cc
index 274eeaf..44ba0a4 100644
--- a/src/tint/lang/wgsl/sem/value_expression.cc
+++ b/src/tint/lang/wgsl/sem/value_expression.cc
@@ -68,7 +68,7 @@
const ValueExpression* ValueExpression::UnwrapLoad() const {
if (auto* l = As<Load>()) {
- return l->Reference();
+ return l->Source();
}
return this;
}
@@ -76,7 +76,7 @@
const ValueExpression* ValueExpression::Unwrap() const {
return Switch(
this, // note: An expression can only be wrapped by a Load or Materialize, not both.
- [&](const Load* load) { return load->Reference(); },
+ [&](const Load* load) { return load->Source(); },
[&](const Materialize* materialize) { return materialize->Expr(); },
[&](Default) { return this; });
}
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl b/test/tint/ptr_sugar/vector_swizzle.wgsl
new file mode 100644
index 0000000..5f349de
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl
@@ -0,0 +1,17 @@
+@group(0) @binding(0) var<storage, read_write> buffer : vec4i;
+
+fn deref() {
+ let p = &buffer;
+ buffer = (*p).wzyx;
+}
+
+fn no_deref() {
+ let p = &buffer;
+ buffer = p.wzyx;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.dxc.hlsl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b56140a
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.dxc.hlsl
@@ -0,0 +1,16 @@
+RWByteAddressBuffer buffer : register(u0);
+
+void deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+void no_deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.fxc.hlsl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b56140a
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.fxc.hlsl
@@ -0,0 +1,16 @@
+RWByteAddressBuffer buffer : register(u0);
+
+void deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+void no_deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.glsl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.glsl
new file mode 100644
index 0000000..8389673
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ ivec4 inner;
+} tint_symbol;
+
+void deref() {
+ tint_symbol.inner = tint_symbol.inner.wzyx;
+}
+
+void no_deref() {
+ tint_symbol.inner = tint_symbol.inner.wzyx;
+}
+
+void tint_symbol_1() {
+ deref();
+ no_deref();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_1();
+ return;
+}
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.dxc.hlsl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..787f342
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+RWByteAddressBuffer buffer : register(u0);
+void deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+void no_deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+}
+
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.fxc.hlsl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..787f342
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+RWByteAddressBuffer buffer : register(u0);
+void deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+void no_deref() {
+ buffer.Store4(0u, asuint(asint(buffer.Load4(0u)).wzyx));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ deref();
+ no_deref();
+}
+
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.msl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.msl
new file mode 100644
index 0000000..5b58134
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.msl
@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int4* tint_symbol;
+};
+
+void deref(tint_module_vars_struct tint_module_vars) {
+ device int4* const p = tint_module_vars.tint_symbol;
+ (*tint_module_vars.tint_symbol) = (*p).wzyx;
+}
+
+void no_deref(tint_module_vars_struct tint_module_vars) {
+ device int4* const p = tint_module_vars.tint_symbol;
+ (*tint_module_vars.tint_symbol) = (*p).wzyx;
+}
+
+kernel void tint_symbol_1(device int4* tint_symbol [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_symbol=tint_symbol};
+ deref(tint_module_vars);
+ no_deref(tint_module_vars);
+}
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.spvasm b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..1253649
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.ir.spvasm
@@ -0,0 +1,9 @@
+SKIP: FAILED
+
+../../src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc:1164 internal compiler error: expression did not resolve to a value
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.msl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.msl
new file mode 100644
index 0000000..2b7b722
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.msl
@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+void deref(device int4* const tint_symbol_2) {
+ *(tint_symbol_2) = (*(tint_symbol_2)).wzyx;
+}
+
+void no_deref(device int4* const tint_symbol_3) {
+ *(tint_symbol_3) = (*(tint_symbol_3)).wzyx;
+}
+
+kernel void tint_symbol_1(device int4* tint_symbol_4 [[buffer(0)]]) {
+ deref(tint_symbol_4);
+ no_deref(tint_symbol_4);
+ return;
+}
+
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.spvasm b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.spvasm
new file mode 100644
index 0000000..df4cbbc
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.spvasm
@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 27
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %buffer_block "buffer_block"
+ OpMemberName %buffer_block 0 "inner"
+ OpName %buffer "buffer"
+ OpName %deref "deref"
+ OpName %no_deref "no_deref"
+ OpName %main "main"
+ OpDecorate %buffer_block Block
+ OpMemberDecorate %buffer_block 0 Offset 0
+ OpDecorate %buffer DescriptorSet 0
+ OpDecorate %buffer Binding 0
+ %int = OpTypeInt 32 1
+ %v4int = OpTypeVector %int 4
+%buffer_block = OpTypeStruct %v4int
+%_ptr_StorageBuffer_buffer_block = OpTypePointer StorageBuffer %buffer_block
+ %buffer = OpVariable %_ptr_StorageBuffer_buffer_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4int = OpTypePointer StorageBuffer %v4int
+ %deref = OpFunction %void None %6
+ %9 = OpLabel
+ %13 = OpAccessChain %_ptr_StorageBuffer_v4int %buffer %uint_0
+ %14 = OpAccessChain %_ptr_StorageBuffer_v4int %buffer %uint_0
+ %15 = OpLoad %v4int %14
+ %16 = OpVectorShuffle %v4int %15 %15 3 2 1 0
+ OpStore %13 %16
+ OpReturn
+ OpFunctionEnd
+ %no_deref = OpFunction %void None %6
+ %18 = OpLabel
+ %19 = OpAccessChain %_ptr_StorageBuffer_v4int %buffer %uint_0
+ %20 = OpAccessChain %_ptr_StorageBuffer_v4int %buffer %uint_0
+ %21 = OpLoad %v4int %20
+ %22 = OpVectorShuffle %v4int %21 %21 3 2 1 0
+ OpStore %19 %22
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %6
+ %24 = OpLabel
+ %25 = OpFunctionCall %void %deref
+ %26 = OpFunctionCall %void %no_deref
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.wgsl b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.wgsl
new file mode 100644
index 0000000..316ee2d
--- /dev/null
+++ b/test/tint/ptr_sugar/vector_swizzle.wgsl.expected.wgsl
@@ -0,0 +1,17 @@
+@group(0) @binding(0) var<storage, read_write> buffer : vec4i;
+
+fn deref() {
+ let p = &(buffer);
+ buffer = (*(p)).wzyx;
+}
+
+fn no_deref() {
+ let p = &(buffer);
+ buffer = p.wzyx;
+}
+
+@compute @workgroup_size(1)
+fn main() {
+ deref();
+ no_deref();
+}
