IntrinsicTable: De-duplicate returned Intrinsics
Much like sem::Type, it greatly simplifies downstream logic if we can compare sem::Intrinsic pointers to know if they refer to the same intrinsic overload.
Change-Id: If236247cd3979bbde821d9294f304ab85ba4938e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/58061
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/inspector/sampler_texture_pair.h b/src/inspector/sampler_texture_pair.h
index 12b4f7e..2af9fbe 100644
--- a/src/inspector/sampler_texture_pair.h
+++ b/src/inspector/sampler_texture_pair.h
@@ -51,8 +51,8 @@
namespace std {
-/// Custom std::hash specialization for ttint::inspector::SamplerTexturePair so
-/// SamplerTexturePairss be used as keys for std::unordered_map and
+/// Custom std::hash specialization for tint::inspector::SamplerTexturePair so
+/// SamplerTexturePairs be used as keys for std::unordered_map and
/// std::unordered_set.
template <>
class hash<tint::inspector::SamplerTexturePair> {
diff --git a/src/intrinsic_table.cc b/src/intrinsic_table.cc
index 0997eae..233aff4 100644
--- a/src/intrinsic_table.cc
+++ b/src/intrinsic_table.cc
@@ -27,6 +27,8 @@
#include "src/sem/pipeline_stage_set.h"
#include "src/sem/sampled_texture_type.h"
#include "src/sem/storage_texture_type.h"
+#include "src/utils/get_or_create.h"
+#include "src/utils/hash.h"
#include "src/utils/scoped_assignment.h"
namespace tint {
@@ -706,13 +708,13 @@
const sem::Intrinsic* Lookup(sem::IntrinsicType intrinsic_type,
const std::vector<const sem::Type*>& args,
- const Source& source) const override;
+ const Source& source) override;
private:
const sem::Intrinsic* Match(sem::IntrinsicType intrinsic_type,
const OverloadInfo& overload,
const std::vector<const sem::Type*>& args,
- int& match_score) const;
+ int& match_score);
MatchState Match(ClosedState& closed,
const OverloadInfo& overload,
@@ -724,6 +726,7 @@
ProgramBuilder& builder;
Matchers matchers;
+ std::unordered_map<sem::Intrinsic, sem::Intrinsic*> intrinsics;
};
/// @return a string representing a call to an intrinsic with the given argument
@@ -760,7 +763,7 @@
const sem::Intrinsic* Impl::Lookup(sem::IntrinsicType intrinsic_type,
const std::vector<const sem::Type*>& args,
- const Source& source) const {
+ const Source& source) {
// Candidate holds information about a mismatched overload that could be what
// the user intended to call.
struct Candidate {
@@ -809,7 +812,7 @@
const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
const OverloadInfo& overload,
const std::vector<const sem::Type*>& args,
- int& match_score) const {
+ int& match_score) {
// Score wait for argument <-> parameter count matches / mismatches
constexpr int kScorePerParamArgMismatch = -1;
constexpr int kScorePerMatchedParam = 2;
@@ -896,9 +899,14 @@
return_type = builder.create<sem::Void>();
}
- return builder.create<sem::Intrinsic>(
- intrinsic_type, const_cast<sem::Type*>(return_type),
- std::move(parameters), overload.supported_stages, overload.is_deprecated);
+ sem::Intrinsic intrinsic(intrinsic_type, const_cast<sem::Type*>(return_type),
+ std::move(parameters), overload.supported_stages,
+ overload.is_deprecated);
+
+ // De-duplicate intrinsics that are identical.
+ return utils::GetOrCreate(intrinsics, intrinsic, [&] {
+ return builder.create<sem::Intrinsic>(intrinsic);
+ });
}
MatchState Impl::Match(ClosedState& closed,
diff --git a/src/intrinsic_table.h b/src/intrinsic_table.h
index 6802e55..12532b6 100644
--- a/src/intrinsic_table.h
+++ b/src/intrinsic_table.h
@@ -45,7 +45,7 @@
virtual const sem::Intrinsic* Lookup(
sem::IntrinsicType type,
const std::vector<const sem::Type*>& args,
- const Source& source) const = 0;
+ const Source& source) = 0;
};
} // namespace tint
diff --git a/src/intrinsic_table_test.cc b/src/intrinsic_table_test.cc
index 9a45c24..5ff0654 100644
--- a/src/intrinsic_table_test.cc
+++ b/src/intrinsic_table_test.cc
@@ -548,5 +548,26 @@
)");
}
+TEST_F(IntrinsicTableTest, SameOverloadReturnsSameIntrinsicPointer) {
+ auto* f32 = create<sem::F32>();
+ auto* vec2_f32 = create<sem::Vector>(create<sem::F32>(), 2);
+ auto* bool_ = create<sem::Bool>();
+ auto* a = table->Lookup(IntrinsicType::kSelect, {f32, f32, bool_}, Source{});
+ ASSERT_NE(a, nullptr) << Diagnostics().str();
+
+ auto* b = table->Lookup(IntrinsicType::kSelect, {f32, f32, bool_}, Source{});
+ ASSERT_NE(b, nullptr) << Diagnostics().str();
+ ASSERT_EQ(Diagnostics().str(), "");
+
+ auto* c = table->Lookup(IntrinsicType::kSelect, {vec2_f32, vec2_f32, bool_},
+ Source{});
+ ASSERT_NE(c, nullptr) << Diagnostics().str();
+ ASSERT_EQ(Diagnostics().str(), "");
+
+ EXPECT_EQ(a, b);
+ EXPECT_NE(a, c);
+ EXPECT_NE(b, c);
+}
+
} // namespace
} // namespace tint
diff --git a/src/sem/call_target.cc b/src/sem/call_target.cc
index ce4c828..5306bf8 100644
--- a/src/sem/call_target.cc
+++ b/src/sem/call_target.cc
@@ -26,6 +26,8 @@
TINT_ASSERT(Semantic, return_type);
}
+CallTarget::CallTarget(const CallTarget&) = default;
+
CallTarget::~CallTarget() = default;
int IndexOf(const ParameterList& parameters, ParameterUsage usage) {
diff --git a/src/sem/call_target.h b/src/sem/call_target.h
index 520e22c..e994c80 100644
--- a/src/sem/call_target.h
+++ b/src/sem/call_target.h
@@ -20,6 +20,7 @@
#include "src/sem/node.h"
#include "src/sem/parameter_usage.h"
#include "src/sem/sampler_type.h"
+#include "src/utils/hash.h"
namespace tint {
@@ -37,11 +38,16 @@
std::ostream& operator<<(std::ostream& out, Parameter parameter);
-/// Comparison operator for Parameters
+/// Equality operator for Parameters
static inline bool operator==(const Parameter& a, const Parameter& b) {
return a.type == b.type && a.usage == b.usage;
}
+/// Inequality operator for Parameters
+static inline bool operator!=(const Parameter& a, const Parameter& b) {
+ return !(a == b);
+}
+
/// ParameterList is a list of Parameter
using ParameterList = std::vector<Parameter>;
@@ -59,6 +65,9 @@
/// @param parameters the parameters for the call target
CallTarget(sem::Type* return_type, const ParameterList& parameters);
+ /// Copy constructor
+ CallTarget(const CallTarget&);
+
/// @return the return type of the call target
sem::Type* ReturnType() const { return return_type_; }
@@ -76,4 +85,19 @@
} // namespace sem
} // namespace tint
+namespace std {
+
+/// Custom std::hash specialization for tint::sem::Parameter
+template <>
+class hash<tint::sem::Parameter> {
+ public:
+ /// @param p the tint::sem::Parameter to create a hash for
+ /// @return the hash value
+ inline std::size_t operator()(const tint::sem::Parameter& p) const {
+ return tint::utils::Hash(p.type, p.usage);
+ }
+};
+
+} // namespace std
+
#endif // SRC_SEM_CALL_TARGET_H_
diff --git a/src/sem/intrinsic.cc b/src/sem/intrinsic.cc
index be1a873..4be75e8 100644
--- a/src/sem/intrinsic.cc
+++ b/src/sem/intrinsic.cc
@@ -111,6 +111,8 @@
supported_stages_(supported_stages),
is_deprecated_(is_deprecated) {}
+Intrinsic::Intrinsic(const Intrinsic&) = default;
+
Intrinsic::~Intrinsic() = default;
bool Intrinsic::IsCoarseDerivative() const {
@@ -153,5 +155,25 @@
return IsAtomicIntrinsic(type_);
}
+bool operator==(const Intrinsic& a, const Intrinsic& b) {
+ static_assert(sizeof(Intrinsic(IntrinsicType::kNone, nullptr, ParameterList{},
+ PipelineStageSet{}, false)) > 0,
+ "don't forget to update the comparison below if you change the "
+ "constructor of Intrinsic!");
+
+ if (a.Type() != b.Type() || a.SupportedStages() != b.SupportedStages() ||
+ a.ReturnType() != b.ReturnType() ||
+ a.IsDeprecated() != b.IsDeprecated() ||
+ a.Parameters().size() != b.Parameters().size()) {
+ return false;
+ }
+ for (size_t i = 0; i < a.Parameters().size(); i++) {
+ if (a.Parameters()[i] != b.Parameters()[i]) {
+ return false;
+ }
+ }
+ return true;
+}
+
} // namespace sem
} // namespace tint
diff --git a/src/sem/intrinsic.h b/src/sem/intrinsic.h
index 11cea7b..25c0c59 100644
--- a/src/sem/intrinsic.h
+++ b/src/sem/intrinsic.h
@@ -20,6 +20,7 @@
#include "src/sem/call_target.h"
#include "src/sem/intrinsic_type.h"
#include "src/sem/pipeline_stage_set.h"
+#include "src/utils/hash.h"
namespace tint {
namespace sem {
@@ -91,6 +92,9 @@
PipelineStageSet supported_stages,
bool is_deprecated);
+ /// Copy constructor
+ Intrinsic(const Intrinsic&);
+
/// Destructor
~Intrinsic() override;
@@ -147,7 +151,31 @@
/// matches the name in the WGSL spec.
std::ostream& operator<<(std::ostream& out, IntrinsicType i);
+/// Equality operator for Intrinsics
+bool operator==(const Intrinsic& a, const Intrinsic& b);
+
+/// Inequality operator for Intrinsics
+static inline bool operator!=(const Intrinsic& a, const Intrinsic& b) {
+ return !(a == b);
+}
+
} // namespace sem
} // namespace tint
+namespace std {
+
+/// Custom std::hash specialization for tint::sem::Intrinsic
+template <>
+class hash<tint::sem::Intrinsic> {
+ public:
+ /// @param i the Intrinsic to create a hash for
+ /// @return the hash value
+ inline std::size_t operator()(const tint::sem::Intrinsic& i) const {
+ return tint::utils::Hash(i.Type(), i.SupportedStages(), i.ReturnType(),
+ i.Parameters(), i.IsDeprecated());
+ }
+};
+
+} // namespace std
+
#endif // SRC_SEM_INTRINSIC_H_
diff --git a/src/utils/enum_set.h b/src/utils/enum_set.h
index 396e911..46d93e1 100644
--- a/src/utils/enum_set.h
+++ b/src/utils/enum_set.h
@@ -16,6 +16,7 @@
#define SRC_UTILS_ENUM_SET_H_
#include <cstdint>
+#include <functional>
#include <type_traits>
namespace tint {
@@ -58,6 +59,19 @@
/// @return true if the set contains `e`
inline bool Contains(Enum e) { return (set & Bit(e)) != 0; }
+ /// Equality operator
+ /// @param rhs the other EnumSet to compare this to
+ /// @return true if this EnumSet is equal to rhs
+ inline bool operator==(const EnumSet& rhs) const { return set == rhs.set; }
+
+ /// Inequality operator
+ /// @param rhs the other EnumSet to compare this to
+ /// @return true if this EnumSet is not equal to rhs
+ inline bool operator!=(const EnumSet& rhs) const { return set != rhs.set; }
+
+ /// @return the underlying value for the EnumSet
+ inline uint64_t Value() const { return set; }
+
private:
static constexpr uint64_t Bit(Enum value) {
return static_cast<uint64_t>(1) << static_cast<uint64_t>(value);
@@ -76,4 +90,19 @@
} // namespace utils
} // namespace tint
+namespace std {
+
+/// Custom std::hash specialization for tint::utils::EnumSet<T>
+template <typename T>
+class hash<tint::utils::EnumSet<T>> {
+ public:
+ /// @param e the EnumSet to create a hash for
+ /// @return the hash value
+ inline std::size_t operator()(const tint::utils::EnumSet<T>& e) const {
+ return std::hash<uint64_t>()(e.Value());
+ }
+};
+
+} // namespace std
+
#endif // SRC_UTILS_ENUM_SET_H_
diff --git a/src/utils/enum_set_test.cc b/src/utils/enum_set_test.cc
index d53b834..80c3dcf 100644
--- a/src/utils/enum_set_test.cc
+++ b/src/utils/enum_set_test.cc
@@ -59,6 +59,30 @@
EXPECT_FALSE(set.Contains(E::C));
}
+TEST(EnumSetTest, Equality) {
+ EXPECT_TRUE(EnumSet<E>(E::A, E::B) == EnumSet<E>(E::A, E::B));
+ EXPECT_FALSE(EnumSet<E>(E::A, E::B) == EnumSet<E>(E::A, E::C));
+}
+
+TEST(EnumSetTest, Inequality) {
+ EXPECT_FALSE(EnumSet<E>(E::A, E::B) != EnumSet<E>(E::A, E::B));
+ EXPECT_TRUE(EnumSet<E>(E::A, E::B) != EnumSet<E>(E::A, E::C));
+}
+
+TEST(EnumSetTest, Hash) {
+ auto hash = [&](EnumSet<E> s) { return std::hash<EnumSet<E>>()(s); };
+ EXPECT_EQ(hash(EnumSet<E>(E::A, E::B)), hash(EnumSet<E>(E::A, E::B)));
+ EXPECT_NE(hash(EnumSet<E>(E::A, E::B)), hash(EnumSet<E>(E::A, E::C)));
+}
+
+TEST(EnumSetTest, Value) {
+ EXPECT_EQ(EnumSet<E>().Value(), 0u);
+ EXPECT_EQ(EnumSet<E>(E::A).Value(), 1u);
+ EXPECT_EQ(EnumSet<E>(E::B).Value(), 2u);
+ EXPECT_EQ(EnumSet<E>(E::C).Value(), 4u);
+ EXPECT_EQ(EnumSet<E>(E::A, E::C).Value(), 5u);
+}
+
} // namespace
} // namespace utils
} // namespace tint
diff --git a/src/utils/get_or_create.h b/src/utils/get_or_create.h
index b5b86ef..dd5e151 100644
--- a/src/utils/get_or_create.h
+++ b/src/utils/get_or_create.h
@@ -27,8 +27,10 @@
/// @param key the map key of the item to query or add
/// @param create a callable function-like object with the signature `V()`
/// @return the value of the item with the given key, or the newly created item
-template <typename K, typename V, typename CREATE, typename H>
-V GetOrCreate(std::unordered_map<K, V, H>& map, K key, CREATE&& create) {
+template <typename K, typename V, typename H, typename C, typename CREATE>
+V GetOrCreate(std::unordered_map<K, V, H, C>& map,
+ const K& key,
+ CREATE&& create) {
auto it = map.find(key);
if (it != map.end()) {
return it->second;
diff --git a/src/utils/hash.h b/src/utils/hash.h
index 44e1e49..bd78942 100644
--- a/src/utils/hash.h
+++ b/src/utils/hash.h
@@ -18,6 +18,7 @@
#include <stdint.h>
#include <cstdio>
#include <functional>
+#include <vector>
namespace tint {
namespace utils {
@@ -51,6 +52,15 @@
*hash ^= std::hash<T>()(value) + offset + (*hash << 6) + (*hash >> 2);
}
+// Helper for hashing vectors
+template <typename T>
+void HashCombine(size_t* hash, const std::vector<T>& vector) {
+ HashCombine(hash, vector.size());
+ for (auto& el : vector) {
+ HashCombine(hash, el);
+ }
+}
+
template <typename T, typename... ARGS>
void HashCombine(size_t* hash, const T& value, const ARGS&... args) {
HashCombine(hash, value);
diff --git a/src/utils/hash_test.cc b/src/utils/hash_test.cc
index 177c53b..dd0877d 100644
--- a/src/utils/hash_test.cc
+++ b/src/utils/hash_test.cc
@@ -23,14 +23,25 @@
namespace {
TEST(HashTests, Basic) {
+ EXPECT_EQ(Hash(123), Hash(123));
EXPECT_NE(Hash(123), Hash(321));
+ EXPECT_EQ(Hash(123, 456), Hash(123, 456));
+ EXPECT_NE(Hash(123, 456), Hash(456, 123));
EXPECT_NE(Hash(123, 456), Hash(123));
+ EXPECT_EQ(Hash(123, 456, false), Hash(123, 456, false));
EXPECT_NE(Hash(123, 456, false), Hash(123, 456));
+ EXPECT_EQ(Hash(std::string("hello")), Hash(std::string("hello")));
EXPECT_NE(Hash(std::string("hello")), Hash(std::string("world")));
}
-TEST(HashTests, Order) {
- EXPECT_NE(Hash(123, 456), Hash(456, 123));
+TEST(HashTests, Vector) {
+ EXPECT_EQ(Hash(std::vector<int>({})), Hash(std::vector<int>({})));
+ EXPECT_EQ(Hash(std::vector<int>({1, 2, 3})),
+ Hash(std::vector<int>({1, 2, 3})));
+ EXPECT_NE(Hash(std::vector<int>({1, 2, 3})),
+ Hash(std::vector<int>({1, 2, 4})));
+ EXPECT_NE(Hash(std::vector<int>({1, 2, 3})),
+ Hash(std::vector<int>({1, 2, 3, 4})));
}
} // namespace
diff --git a/test/intrinsics/repeated_use.wgsl b/test/intrinsics/repeated_use.wgsl
new file mode 100644
index 0000000..24a8370
--- /dev/null
+++ b/test/intrinsics/repeated_use.wgsl
@@ -0,0 +1,20 @@
+// Check that for backends that generate intrinsic helpers, repeated use of the
+// same intrinsic overload results in single helper being generated.
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ ignore(isNormal(vec4<f32>()));
+ ignore(isNormal(vec4<f32>(1.)));
+ ignore(isNormal(vec4<f32>(1., 2., 3., 4.)));
+
+ ignore(isNormal(vec3<f32>()));
+ ignore(isNormal(vec3<f32>(1.)));
+ ignore(isNormal(vec3<f32>(1., 2., 3.)));
+
+ ignore(isNormal(vec2<f32>()));
+ ignore(isNormal(vec2<f32>(1.)));
+ ignore(isNormal(vec2<f32>(1., 2.)));
+
+ ignore(isNormal(1.));
+ ignore(isNormal(2.));
+ ignore(isNormal(3.));
+}
diff --git a/test/intrinsics/repeated_use.wgsl.expected.hlsl b/test/intrinsics/repeated_use.wgsl.expected.hlsl
new file mode 100644
index 0000000..c89b9fb
--- /dev/null
+++ b/test/intrinsics/repeated_use.wgsl.expected.hlsl
@@ -0,0 +1,40 @@
+bool4 tint_isNormal(float4 param_0) {
+ uint4 exponent = asuint(param_0) & 0x7f80000;
+ uint4 clamped = clamp(exponent, 0x0080000, 0x7f00000);
+ return clamped == exponent;
+}
+
+bool3 tint_isNormal_1(float3 param_0) {
+ uint3 exponent = asuint(param_0) & 0x7f80000;
+ uint3 clamped = clamp(exponent, 0x0080000, 0x7f00000);
+ return clamped == exponent;
+}
+
+bool2 tint_isNormal_2(float2 param_0) {
+ uint2 exponent = asuint(param_0) & 0x7f80000;
+ uint2 clamped = clamp(exponent, 0x0080000, 0x7f00000);
+ return clamped == exponent;
+}
+
+bool tint_isNormal_3(float param_0) {
+ uint exponent = asuint(param_0) & 0x7f80000;
+ uint clamped = clamp(exponent, 0x0080000, 0x7f00000);
+ return clamped == exponent;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_isNormal(float4(0.0f, 0.0f, 0.0f, 0.0f));
+ tint_isNormal(float4((1.0f).xxxx));
+ tint_isNormal(float4(1.0f, 2.0f, 3.0f, 4.0f));
+ tint_isNormal_1(float3(0.0f, 0.0f, 0.0f));
+ tint_isNormal_1(float3((1.0f).xxx));
+ tint_isNormal_1(float3(1.0f, 2.0f, 3.0f));
+ tint_isNormal_2(float2(0.0f, 0.0f));
+ tint_isNormal_2(float2((1.0f).xx));
+ tint_isNormal_2(float2(1.0f, 2.0f));
+ tint_isNormal_3(1.0f);
+ tint_isNormal_3(2.0f);
+ tint_isNormal_3(3.0f);
+ return;
+}
diff --git a/test/intrinsics/repeated_use.wgsl.expected.msl b/test/intrinsics/repeated_use.wgsl.expected.msl
new file mode 100644
index 0000000..16b9698
--- /dev/null
+++ b/test/intrinsics/repeated_use.wgsl.expected.msl
@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ (void) isnormal(float4());
+ (void) isnormal(float4(1.0f));
+ (void) isnormal(float4(1.0f, 2.0f, 3.0f, 4.0f));
+ (void) isnormal(float3());
+ (void) isnormal(float3(1.0f));
+ (void) isnormal(float3(1.0f, 2.0f, 3.0f));
+ (void) isnormal(float2());
+ (void) isnormal(float2(1.0f));
+ (void) isnormal(float2(1.0f, 2.0f));
+ (void) isnormal(1.0f);
+ (void) isnormal(2.0f);
+ (void) isnormal(3.0f);
+ return;
+}
+
diff --git a/test/intrinsics/repeated_use.wgsl.expected.spvasm b/test/intrinsics/repeated_use.wgsl.expected.spvasm
new file mode 100644
index 0000000..42acf47
--- /dev/null
+++ b/test/intrinsics/repeated_use.wgsl.expected.spvasm
@@ -0,0 +1,120 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 121
+; Schema: 0
+ OpCapability Shader
+ %12 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %bool = OpTypeBool
+ %v4bool = OpTypeVector %bool 4
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %11 = OpConstantNull %v4float
+ %uint = OpTypeInt 32 0
+%uint_133693440 = OpConstant %uint 133693440
+%uint_524288 = OpConstant %uint 524288
+%uint_133169152 = OpConstant %uint 133169152
+ %v4uint = OpTypeVector %uint 4
+ %float_1 = OpConstant %float 1
+ %27 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+ %float_2 = OpConstant %float 2
+ %float_3 = OpConstant %float 3
+ %float_4 = OpConstant %float 4
+ %39 = OpConstantComposite %v4float %float_1 %float_2 %float_3 %float_4
+ %v3bool = OpTypeVector %bool 3
+ %v3float = OpTypeVector %float 3
+ %50 = OpConstantNull %v3float
+ %v3uint = OpTypeVector %uint 3
+ %60 = OpConstantComposite %v3float %float_1 %float_1 %float_1
+ %69 = OpConstantComposite %v3float %float_1 %float_2 %float_3
+ %v2bool = OpTypeVector %bool 2
+ %v2float = OpTypeVector %float 2
+ %80 = OpConstantNull %v2float
+ %v2uint = OpTypeVector %uint 2
+ %90 = OpConstantComposite %v2float %float_1 %float_1
+ %99 = OpConstantComposite %v2float %float_1 %float_2
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %18 = OpCompositeConstruct %v4uint %uint_133693440 %uint_133693440 %uint_133693440 %uint_133693440
+ %19 = OpCompositeConstruct %v4uint %uint_524288 %uint_524288 %uint_524288 %uint_524288
+ %20 = OpCompositeConstruct %v4uint %uint_133169152 %uint_133169152 %uint_133169152 %uint_133169152
+ %21 = OpBitcast %v4uint %11
+ %22 = OpBitwiseAnd %v4uint %21 %18
+ %23 = OpExtInst %v4uint %12 UClamp %22 %19 %20
+ %6 = OpIEqual %v4bool %22 %23
+ %28 = OpCompositeConstruct %v4uint %uint_133693440 %uint_133693440 %uint_133693440 %uint_133693440
+ %29 = OpCompositeConstruct %v4uint %uint_524288 %uint_524288 %uint_524288 %uint_524288
+ %30 = OpCompositeConstruct %v4uint %uint_133169152 %uint_133169152 %uint_133169152 %uint_133169152
+ %31 = OpBitcast %v4uint %27
+ %32 = OpBitwiseAnd %v4uint %31 %28
+ %33 = OpExtInst %v4uint %12 UClamp %32 %29 %30
+ %25 = OpIEqual %v4bool %32 %33
+ %40 = OpCompositeConstruct %v4uint %uint_133693440 %uint_133693440 %uint_133693440 %uint_133693440
+ %41 = OpCompositeConstruct %v4uint %uint_524288 %uint_524288 %uint_524288 %uint_524288
+ %42 = OpCompositeConstruct %v4uint %uint_133169152 %uint_133169152 %uint_133169152 %uint_133169152
+ %43 = OpBitcast %v4uint %39
+ %44 = OpBitwiseAnd %v4uint %43 %40
+ %45 = OpExtInst %v4uint %12 UClamp %44 %41 %42
+ %35 = OpIEqual %v4bool %44 %45
+ %52 = OpCompositeConstruct %v3uint %uint_133693440 %uint_133693440 %uint_133693440
+ %53 = OpCompositeConstruct %v3uint %uint_524288 %uint_524288 %uint_524288
+ %54 = OpCompositeConstruct %v3uint %uint_133169152 %uint_133169152 %uint_133169152
+ %55 = OpBitcast %v3uint %50
+ %56 = OpBitwiseAnd %v3uint %55 %52
+ %57 = OpExtInst %v3uint %12 UClamp %56 %53 %54
+ %47 = OpIEqual %v3bool %56 %57
+ %61 = OpCompositeConstruct %v3uint %uint_133693440 %uint_133693440 %uint_133693440
+ %62 = OpCompositeConstruct %v3uint %uint_524288 %uint_524288 %uint_524288
+ %63 = OpCompositeConstruct %v3uint %uint_133169152 %uint_133169152 %uint_133169152
+ %64 = OpBitcast %v3uint %60
+ %65 = OpBitwiseAnd %v3uint %64 %61
+ %66 = OpExtInst %v3uint %12 UClamp %65 %62 %63
+ %59 = OpIEqual %v3bool %65 %66
+ %70 = OpCompositeConstruct %v3uint %uint_133693440 %uint_133693440 %uint_133693440
+ %71 = OpCompositeConstruct %v3uint %uint_524288 %uint_524288 %uint_524288
+ %72 = OpCompositeConstruct %v3uint %uint_133169152 %uint_133169152 %uint_133169152
+ %73 = OpBitcast %v3uint %69
+ %74 = OpBitwiseAnd %v3uint %73 %70
+ %75 = OpExtInst %v3uint %12 UClamp %74 %71 %72
+ %68 = OpIEqual %v3bool %74 %75
+ %82 = OpCompositeConstruct %v2uint %uint_133693440 %uint_133693440
+ %83 = OpCompositeConstruct %v2uint %uint_524288 %uint_524288
+ %84 = OpCompositeConstruct %v2uint %uint_133169152 %uint_133169152
+ %85 = OpBitcast %v2uint %80
+ %86 = OpBitwiseAnd %v2uint %85 %82
+ %87 = OpExtInst %v2uint %12 UClamp %86 %83 %84
+ %77 = OpIEqual %v2bool %86 %87
+ %91 = OpCompositeConstruct %v2uint %uint_133693440 %uint_133693440
+ %92 = OpCompositeConstruct %v2uint %uint_524288 %uint_524288
+ %93 = OpCompositeConstruct %v2uint %uint_133169152 %uint_133169152
+ %94 = OpBitcast %v2uint %90
+ %95 = OpBitwiseAnd %v2uint %94 %91
+ %96 = OpExtInst %v2uint %12 UClamp %95 %92 %93
+ %89 = OpIEqual %v2bool %95 %96
+ %100 = OpCompositeConstruct %v2uint %uint_133693440 %uint_133693440
+ %101 = OpCompositeConstruct %v2uint %uint_524288 %uint_524288
+ %102 = OpCompositeConstruct %v2uint %uint_133169152 %uint_133169152
+ %103 = OpBitcast %v2uint %99
+ %104 = OpBitwiseAnd %v2uint %103 %100
+ %105 = OpExtInst %v2uint %12 UClamp %104 %101 %102
+ %98 = OpIEqual %v2bool %104 %105
+ %108 = OpBitcast %uint %float_1
+ %109 = OpBitwiseAnd %uint %108 %uint_133693440
+ %110 = OpExtInst %uint %12 UClamp %109 %uint_524288 %uint_133169152
+ %107 = OpIEqual %bool %109 %110
+ %113 = OpBitcast %uint %float_2
+ %114 = OpBitwiseAnd %uint %113 %uint_133693440
+ %115 = OpExtInst %uint %12 UClamp %114 %uint_524288 %uint_133169152
+ %112 = OpIEqual %bool %114 %115
+ %118 = OpBitcast %uint %float_3
+ %119 = OpBitwiseAnd %uint %118 %uint_133693440
+ %120 = OpExtInst %uint %12 UClamp %119 %uint_524288 %uint_133169152
+ %117 = OpIEqual %bool %119 %120
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/repeated_use.wgsl.expected.wgsl b/test/intrinsics/repeated_use.wgsl.expected.wgsl
new file mode 100644
index 0000000..e708833
--- /dev/null
+++ b/test/intrinsics/repeated_use.wgsl.expected.wgsl
@@ -0,0 +1,15 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ ignore(isNormal(vec4<f32>()));
+ ignore(isNormal(vec4<f32>(1.0)));
+ ignore(isNormal(vec4<f32>(1.0, 2.0, 3.0, 4.0)));
+ ignore(isNormal(vec3<f32>()));
+ ignore(isNormal(vec3<f32>(1.0)));
+ ignore(isNormal(vec3<f32>(1.0, 2.0, 3.0)));
+ ignore(isNormal(vec2<f32>()));
+ ignore(isNormal(vec2<f32>(1.0)));
+ ignore(isNormal(vec2<f32>(1.0, 2.0)));
+ ignore(isNormal(1.0));
+ ignore(isNormal(2.0));
+ ignore(isNormal(3.0));
+}
