From 08f4b557fcf03e7fa6fea0342fb47b7c194f27be Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Tue, 31 May 2022 13:20:28 +0000 Subject: [PATCH] Implement atomicCompareExchangeWeak returning struct instead of vec2 Also fixed implementation of this atomic in GLSL. It was emitting code that would not compile because, as for HLSL, we must pass in the variable directly to atomic funcs, not via an in/out arg to a function. Bug: tint:1185 Change-Id: Id0e9f99d6368717511ef3a94473634c512e10cb8 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/91881 Kokoro: Kokoro Reviewed-by: Ben Clayton Commit-Queue: Antonio Maiorano --- src/tint/intrinsics.def | 4 +- src/tint/resolver/intrinsic_table.cc | 15 + src/tint/resolver/intrinsic_table.inl | 521 ++++++++++-------- src/tint/transform/decompose_memory_access.cc | 40 +- src/tint/transform/decompose_memory_access.h | 3 + .../transform/decompose_memory_access_test.cc | 196 ++++--- src/tint/transform/manager.cc | 2 +- src/tint/writer/glsl/generator_impl.cc | 77 ++- src/tint/writer/glsl/generator_impl.h | 8 - src/tint/writer/hlsl/generator_impl.cc | 408 +++++++------- src/tint/writer/hlsl/generator_impl.h | 7 +- src/tint/writer/msl/generator_impl.cc | 17 +- src/tint/writer/spirv/builder.cc | 34 +- src/tint/writer/spirv/builder_builtin_test.cc | 54 +- .../bug/chromium/1273230.wgsl.expected.hlsl | 45 +- test/tint/bug/tint/1113.wgsl.expected.hlsl | 84 +-- test/tint/bug/tint/1121.wgsl.expected.hlsl | 15 +- test/tint/bug/tint/926.wgsl.expected.hlsl | 15 +- test/tint/bug/tint/993.wgsl.expected.hlsl | 15 +- .../gen/atomicAdd/8a199a.wgsl.expected.hlsl | 7 +- .../gen/atomicAdd/d32fe4.wgsl.expected.hlsl | 7 +- .../gen/atomicAnd/152966.wgsl.expected.hlsl | 7 +- .../gen/atomicAnd/85a8d9.wgsl.expected.hlsl | 7 +- .../12871c.wgsl.expected.glsl | 76 --- .../12871c.wgsl.expected.hlsl | 23 - .../12871c.wgsl.expected.msl | 29 - .../{12871c.wgsl => 1bd40a.wgsl} | 10 +- .../1bd40a.wgsl.expected.glsl | 62 +++ .../1bd40a.wgsl.expected.hlsl | 29 + .../1bd40a.wgsl.expected.msl | 33 ++ ...ted.spvasm => 1bd40a.wgsl.expected.spvasm} | 37 +- ...xpected.wgsl => 1bd40a.wgsl.expected.wgsl} | 8 +- .../{6673da.wgsl => 63d8e6.wgsl} | 10 +- .../63d8e6.wgsl.expected.glsl | 62 +++ .../63d8e6.wgsl.expected.hlsl | 29 + .../63d8e6.wgsl.expected.msl | 33 ++ ...ted.spvasm => 63d8e6.wgsl.expected.spvasm} | 36 +- ...xpected.wgsl => 63d8e6.wgsl.expected.wgsl} | 8 +- .../6673da.wgsl.expected.glsl | 76 --- .../6673da.wgsl.expected.hlsl | 23 - .../6673da.wgsl.expected.msl | 29 - .../{b2ab2c.wgsl => 83580d.wgsl} | 8 +- .../83580d.wgsl.expected.glsl | 29 + ...xpected.hlsl => 83580d.wgsl.expected.hlsl} | 16 +- .../83580d.wgsl.expected.msl | 33 ++ ...ted.spvasm => 83580d.wgsl.expected.spvasm} | 42 +- .../83580d.wgsl.expected.wgsl | 10 + .../89ea3b.wgsl.expected.glsl | 37 -- .../89ea3b.wgsl.expected.msl | 29 - .../89ea3b.wgsl.expected.wgsl | 10 - .../b2ab2c.wgsl.expected.glsl | 37 -- .../b2ab2c.wgsl.expected.msl | 29 - .../b2ab2c.wgsl.expected.wgsl | 10 - .../{89ea3b.wgsl => e88938.wgsl} | 8 +- .../e88938.wgsl.expected.glsl | 29 + ...xpected.hlsl => e88938.wgsl.expected.hlsl} | 16 +- .../e88938.wgsl.expected.msl | 33 ++ ...ted.spvasm => e88938.wgsl.expected.spvasm} | 43 +- .../e88938.wgsl.expected.wgsl | 10 + .../atomicExchange/d59712.wgsl.expected.hlsl | 7 +- .../atomicExchange/f2e22f.wgsl.expected.hlsl | 7 +- .../gen/atomicLoad/0806ad.wgsl.expected.hlsl | 7 +- .../gen/atomicLoad/fe6cc3.wgsl.expected.hlsl | 7 +- .../gen/atomicMax/51b9be.wgsl.expected.hlsl | 7 +- .../gen/atomicMax/92aa72.wgsl.expected.hlsl | 7 +- .../gen/atomicMin/8e38dc.wgsl.expected.hlsl | 7 +- .../gen/atomicMin/c67a74.wgsl.expected.hlsl | 7 +- .../gen/atomicOr/5e95d4.wgsl.expected.hlsl | 7 +- .../gen/atomicOr/8d96a0.wgsl.expected.hlsl | 7 +- .../gen/atomicStore/cdc29e.wgsl.expected.hlsl | 7 +- .../gen/atomicStore/d1e9a6.wgsl.expected.hlsl | 7 +- .../gen/atomicSub/051100.wgsl.expected.hlsl | 7 +- .../gen/atomicSub/15bfc9.wgsl.expected.hlsl | 7 +- .../gen/atomicXor/54510e.wgsl.expected.hlsl | 7 +- .../gen/atomicXor/c1b78c.wgsl.expected.hlsl | 7 +- 75 files changed, 1428 insertions(+), 1314 deletions(-) delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{12871c.wgsl => 1bd40a.wgsl} (84%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.hlsl create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{12871c.wgsl.expected.spvasm => 1bd40a.wgsl.expected.spvasm} (56%) rename test/tint/builtins/gen/atomicCompareExchangeWeak/{12871c.wgsl.expected.wgsl => 1bd40a.wgsl.expected.wgsl} (52%) rename test/tint/builtins/gen/atomicCompareExchangeWeak/{6673da.wgsl => 63d8e6.wgsl} (84%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.hlsl create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{6673da.wgsl.expected.spvasm => 63d8e6.wgsl.expected.spvasm} (55%) rename test/tint/builtins/gen/atomicCompareExchangeWeak/{6673da.wgsl.expected.wgsl => 63d8e6.wgsl.expected.wgsl} (51%) delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{b2ab2c.wgsl => 83580d.wgsl} (85%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{b2ab2c.wgsl.expected.hlsl => 83580d.wgsl.expected.hlsl} (56%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.msl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{b2ab2c.wgsl.expected.spvasm => 83580d.wgsl.expected.spvasm} (54%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.wgsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl delete mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{89ea3b.wgsl => e88938.wgsl} (85%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{89ea3b.wgsl.expected.hlsl => e88938.wgsl.expected.hlsl} (56%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.msl rename test/tint/builtins/gen/atomicCompareExchangeWeak/{89ea3b.wgsl.expected.spvasm => e88938.wgsl.expected.spvasm} (54%) create mode 100644 test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.wgsl diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def index bc4bf28dfd..a6792b13ae 100644 --- a/src/tint/intrinsics.def +++ b/src/tint/intrinsics.def @@ -118,6 +118,8 @@ type __modf_result type __frexp_result [[display("__frexp_result_vec{N}")]] type __frexp_result_vec +type __atomic_compare_exchange_result + //////////////////////////////////////////////////////////////////////////////// // Type matchers // // // @@ -603,7 +605,7 @@ fn textureLoad(texture: texture_external, coords: vec2) -> vec4 [[stage("fragment", "compute")]] fn atomicOr(ptr, read_write>, T) -> T [[stage("fragment", "compute")]] fn atomicXor(ptr, read_write>, T) -> T [[stage("fragment", "compute")]] fn atomicExchange(ptr, read_write>, T) -> T -[[stage("fragment", "compute")]] fn atomicCompareExchangeWeak(ptr, read_write>, T, T) -> vec2 +[[stage("fragment", "compute")]] fn atomicCompareExchangeWeak(ptr, read_write>, T, T) -> __atomic_compare_exchange_result //////////////////////////////////////////////////////////////////////////////// // Type constructors // diff --git a/src/tint/resolver/intrinsic_table.cc b/src/tint/resolver/intrinsic_table.cc index 6c19f7cd94..85711769f1 100644 --- a/src/tint/resolver/intrinsic_table.cc +++ b/src/tint/resolver/intrinsic_table.cc @@ -722,6 +722,14 @@ bool match_frexp_result_vec(const sem::Type* ty, Number& N) { return true; } +bool match_atomic_compare_exchange_result(const sem::Type* ty, const sem::Type*& T) { + if (ty->Is()) { + T = ty; + return true; + } + return false; +} + struct NameAndType { std::string name; sem::Type* type; @@ -779,6 +787,13 @@ const sem::Struct* build_frexp_result_vec(MatchState& state, Number& n) { {{"sig", vec_f32}, {"exp", vec_i32}}); } +const sem::Struct* build_atomic_compare_exchange_result(MatchState& state, const sem::Type* ty) { + return build_struct( + state, "__atomic_compare_exchange_result" + ty->FriendlyName(state.builder.Symbols()), + {{"old_value", const_cast(ty)}, + {"exchanged", state.builder.create()}}); +} + /// ParameterInfo describes a parameter struct ParameterInfo { /// The parameter usage (parameter name in definition file) diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl index c518e58f7d..7a422fb6f7 100644 --- a/src/tint/resolver/intrinsic_table.inl +++ b/src/tint/resolver/intrinsic_table.inl @@ -1512,8 +1512,41 @@ std::string FrexpResultVec::String(MatchState* state) const { return ss.str(); } +/// TypeMatcher for 'type __atomic_compare_exchange_result' +/// @see src/tint/intrinsics.def:121:6 +class AtomicCompareExchangeResult : public TypeMatcher { + public: + /// Checks whether the given type matches the matcher rules. + /// Match may define and refine the template types and numbers in state. + /// @param state the MatchState + /// @param type the type to match + /// @returns the canonicalized type on match, otherwise nullptr + const sem::Type* Match(MatchState& state, + const sem::Type* type) const override; + /// @param state the MatchState + /// @return a string representation of the matcher. + std::string String(MatchState* state) const override; +}; + +const sem::Type* AtomicCompareExchangeResult::Match(MatchState& state, const sem::Type* ty) const { + const sem::Type* T = nullptr; + if (!match_atomic_compare_exchange_result(ty, T)) { + return nullptr; + } + T = state.Type(T); + if (T == nullptr) { + return nullptr; + } + return build_atomic_compare_exchange_result(state, T); +} + +std::string AtomicCompareExchangeResult::String(MatchState* state) const { + const std::string T = state->TypeName(); + return "__atomic_compare_exchange_result<" + T + ">"; +} + /// TypeMatcher for 'match fiu32' -/// @see src/tint/intrinsics.def:127:7 +/// @see src/tint/intrinsics.def:129:7 class Fiu32 : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1551,7 +1584,7 @@ std::string Fiu32::String(MatchState*) const { } /// TypeMatcher for 'match fi32' -/// @see src/tint/intrinsics.def:128:7 +/// @see src/tint/intrinsics.def:130:7 class Fi32 : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1586,7 +1619,7 @@ std::string Fi32::String(MatchState*) const { } /// TypeMatcher for 'match iu32' -/// @see src/tint/intrinsics.def:129:7 +/// @see src/tint/intrinsics.def:131:7 class Iu32 : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1621,7 +1654,7 @@ std::string Iu32::String(MatchState*) const { } /// TypeMatcher for 'match scalar' -/// @see src/tint/intrinsics.def:130:7 +/// @see src/tint/intrinsics.def:132:7 class Scalar : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1662,7 +1695,7 @@ std::string Scalar::String(MatchState*) const { } /// TypeMatcher for 'match abstract_or_scalar' -/// @see src/tint/intrinsics.def:131:7 +/// @see src/tint/intrinsics.def:133:7 class AbstractOrScalar : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1709,7 +1742,7 @@ std::string AbstractOrScalar::String(MatchState*) const { } /// TypeMatcher for 'match af_f32' -/// @see src/tint/intrinsics.def:132:7 +/// @see src/tint/intrinsics.def:134:7 class AfF32 : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1744,7 +1777,7 @@ std::string AfF32::String(MatchState*) const { } /// TypeMatcher for 'match scalar_no_f32' -/// @see src/tint/intrinsics.def:133:7 +/// @see src/tint/intrinsics.def:135:7 class ScalarNoF32 : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1782,7 +1815,7 @@ std::string ScalarNoF32::String(MatchState*) const { } /// TypeMatcher for 'match scalar_no_i32' -/// @see src/tint/intrinsics.def:134:7 +/// @see src/tint/intrinsics.def:136:7 class ScalarNoI32 : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1820,7 +1853,7 @@ std::string ScalarNoI32::String(MatchState*) const { } /// TypeMatcher for 'match scalar_no_u32' -/// @see src/tint/intrinsics.def:135:7 +/// @see src/tint/intrinsics.def:137:7 class ScalarNoU32 : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1858,7 +1891,7 @@ std::string ScalarNoU32::String(MatchState*) const { } /// TypeMatcher for 'match scalar_no_bool' -/// @see src/tint/intrinsics.def:136:7 +/// @see src/tint/intrinsics.def:138:7 class ScalarNoBool : public TypeMatcher { public: /// Checks whether the given type matches the matcher rules, and returns the @@ -1896,7 +1929,7 @@ std::string ScalarNoBool::String(MatchState*) const { } /// EnumMatcher for 'match f32_texel_format' -/// @see src/tint/intrinsics.def:147:7 +/// @see src/tint/intrinsics.def:149:7 class F32TexelFormat : public NumberMatcher { public: /// Checks whether the given number matches the enum matcher rules. @@ -1929,7 +1962,7 @@ std::string F32TexelFormat::String(MatchState*) const { } /// EnumMatcher for 'match i32_texel_format' -/// @see src/tint/intrinsics.def:149:7 +/// @see src/tint/intrinsics.def:151:7 class I32TexelFormat : public NumberMatcher { public: /// Checks whether the given number matches the enum matcher rules. @@ -1961,7 +1994,7 @@ std::string I32TexelFormat::String(MatchState*) const { } /// EnumMatcher for 'match u32_texel_format' -/// @see src/tint/intrinsics.def:151:7 +/// @see src/tint/intrinsics.def:153:7 class U32TexelFormat : public NumberMatcher { public: /// Checks whether the given number matches the enum matcher rules. @@ -1993,7 +2026,7 @@ std::string U32TexelFormat::String(MatchState*) const { } /// EnumMatcher for 'match write_only' -/// @see src/tint/intrinsics.def:154:7 +/// @see src/tint/intrinsics.def:156:7 class WriteOnly : public NumberMatcher { public: /// Checks whether the given number matches the enum matcher rules. @@ -2019,7 +2052,7 @@ std::string WriteOnly::String(MatchState*) const { } /// EnumMatcher for 'match function_private_workgroup' -/// @see src/tint/intrinsics.def:156:7 +/// @see src/tint/intrinsics.def:158:7 class FunctionPrivateWorkgroup : public NumberMatcher { public: /// Checks whether the given number matches the enum matcher rules. @@ -2049,7 +2082,7 @@ std::string FunctionPrivateWorkgroup::String(MatchState*) const { } /// EnumMatcher for 'match workgroup_or_storage' -/// @see src/tint/intrinsics.def:157:7 +/// @see src/tint/intrinsics.def:159:7 class WorkgroupOrStorage : public NumberMatcher { public: /// Checks whether the given number matches the enum matcher rules. @@ -2206,6 +2239,7 @@ class Matchers { ModfResultVec ModfResultVec_; FrexpResult FrexpResult_; FrexpResultVec FrexpResultVec_; + AtomicCompareExchangeResult AtomicCompareExchangeResult_; Fiu32 Fiu32_; Fi32 Fi32_; Iu32 Iu32_; @@ -2233,7 +2267,7 @@ class Matchers { ~Matchers(); /// The template types, types, and type matchers - TypeMatcher const* const type[58] = { + TypeMatcher const* const type[59] = { /* [0] */ &template_type_0_, /* [1] */ &template_type_1_, /* [2] */ &Bool_, @@ -2282,16 +2316,17 @@ class Matchers { /* [45] */ &ModfResultVec_, /* [46] */ &FrexpResult_, /* [47] */ &FrexpResultVec_, - /* [48] */ &Fiu32_, - /* [49] */ &Fi32_, - /* [50] */ &Iu32_, - /* [51] */ &Scalar_, - /* [52] */ &AbstractOrScalar_, - /* [53] */ &AfF32_, - /* [54] */ &ScalarNoF32_, - /* [55] */ &ScalarNoI32_, - /* [56] */ &ScalarNoU32_, - /* [57] */ &ScalarNoBool_, + /* [48] */ &AtomicCompareExchangeResult_, + /* [49] */ &Fiu32_, + /* [50] */ &Fi32_, + /* [51] */ &Iu32_, + /* [52] */ &Scalar_, + /* [53] */ &AbstractOrScalar_, + /* [54] */ &AfF32_, + /* [55] */ &ScalarNoF32_, + /* [56] */ &ScalarNoI32_, + /* [57] */ &ScalarNoU32_, + /* [58] */ &ScalarNoBool_, }; /// The template numbers, and number matchers @@ -2488,34 +2523,36 @@ constexpr MatcherIndex kMatcherIndices[] = { /* [170] */ 7, /* [171] */ 17, /* [172] */ 0, - /* [173] */ 18, - /* [174] */ 7, + /* [173] */ 48, + /* [174] */ 0, /* [175] */ 18, - /* [176] */ 0, - /* [177] */ 27, - /* [178] */ 7, - /* [179] */ 28, + /* [176] */ 7, + /* [177] */ 18, + /* [178] */ 0, + /* [179] */ 27, /* [180] */ 7, - /* [181] */ 29, + /* [181] */ 28, /* [182] */ 7, - /* [183] */ 19, + /* [183] */ 29, /* [184] */ 7, - /* [185] */ 30, + /* [185] */ 19, /* [186] */ 7, - /* [187] */ 31, + /* [187] */ 30, /* [188] */ 7, - /* [189] */ 32, + /* [189] */ 31, /* [190] */ 7, - /* [191] */ 25, - /* [192] */ 26, - /* [193] */ 37, - /* [194] */ 36, - /* [195] */ 35, - /* [196] */ 34, - /* [197] */ 43, - /* [198] */ 38, - /* [199] */ 44, - /* [200] */ 46, + /* [191] */ 32, + /* [192] */ 7, + /* [193] */ 25, + /* [194] */ 26, + /* [195] */ 37, + /* [196] */ 36, + /* [197] */ 35, + /* [198] */ 34, + /* [199] */ 43, + /* [200] */ 38, + /* [201] */ 44, + /* [202] */ 46, }; // Assert that the MatcherIndex is big enough to index all the matchers, plus @@ -2853,12 +2890,12 @@ constexpr ParameterInfo kParameters[] = { { /* [65] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [66] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [67] */ @@ -2888,12 +2925,12 @@ constexpr ParameterInfo kParameters[] = { { /* [72] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [73] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [74] */ @@ -2948,12 +2985,12 @@ constexpr ParameterInfo kParameters[] = { { /* [84] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [85] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [86] */ @@ -3018,7 +3055,7 @@ constexpr ParameterInfo kParameters[] = { { /* [98] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [99] */ @@ -3038,12 +3075,12 @@ constexpr ParameterInfo kParameters[] = { { /* [102] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [103] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [104] */ @@ -3068,12 +3105,12 @@ constexpr ParameterInfo kParameters[] = { { /* [108] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [109] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [110] */ @@ -3098,12 +3135,12 @@ constexpr ParameterInfo kParameters[] = { { /* [114] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [115] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [116] */ @@ -3128,12 +3165,12 @@ constexpr ParameterInfo kParameters[] = { { /* [120] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[189], + /* matcher indices */ &kMatcherIndices[191], }, { /* [121] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [122] */ @@ -3158,12 +3195,12 @@ constexpr ParameterInfo kParameters[] = { { /* [126] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [127] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [128] */ @@ -3188,12 +3225,12 @@ constexpr ParameterInfo kParameters[] = { { /* [132] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [133] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [134] */ @@ -3218,12 +3255,12 @@ constexpr ParameterInfo kParameters[] = { { /* [138] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [139] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [140] */ @@ -3248,12 +3285,12 @@ constexpr ParameterInfo kParameters[] = { { /* [144] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [145] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [146] */ @@ -3278,12 +3315,12 @@ constexpr ParameterInfo kParameters[] = { { /* [150] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [151] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [152] */ @@ -3303,12 +3340,12 @@ constexpr ParameterInfo kParameters[] = { { /* [155] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [156] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [157] */ @@ -3328,12 +3365,12 @@ constexpr ParameterInfo kParameters[] = { { /* [160] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [161] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [162] */ @@ -3353,12 +3390,12 @@ constexpr ParameterInfo kParameters[] = { { /* [165] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[189], + /* matcher indices */ &kMatcherIndices[191], }, { /* [166] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [167] */ @@ -3378,12 +3415,12 @@ constexpr ParameterInfo kParameters[] = { { /* [170] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [171] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [172] */ @@ -3403,12 +3440,12 @@ constexpr ParameterInfo kParameters[] = { { /* [175] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [176] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [177] */ @@ -3428,12 +3465,12 @@ constexpr ParameterInfo kParameters[] = { { /* [180] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [181] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [182] */ @@ -3453,12 +3490,12 @@ constexpr ParameterInfo kParameters[] = { { /* [185] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [186] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [187] */ @@ -3478,12 +3515,12 @@ constexpr ParameterInfo kParameters[] = { { /* [190] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [191] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [192] */ @@ -3503,12 +3540,12 @@ constexpr ParameterInfo kParameters[] = { { /* [195] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [196] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [197] */ @@ -3528,12 +3565,12 @@ constexpr ParameterInfo kParameters[] = { { /* [200] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [201] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [202] */ @@ -3553,12 +3590,12 @@ constexpr ParameterInfo kParameters[] = { { /* [205] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [206] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [207] */ @@ -3578,12 +3615,12 @@ constexpr ParameterInfo kParameters[] = { { /* [210] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[187], + /* matcher indices */ &kMatcherIndices[189], }, { /* [211] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [212] */ @@ -3603,12 +3640,12 @@ constexpr ParameterInfo kParameters[] = { { /* [215] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [216] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [217] */ @@ -3628,12 +3665,12 @@ constexpr ParameterInfo kParameters[] = { { /* [220] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [221] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [222] */ @@ -3653,12 +3690,12 @@ constexpr ParameterInfo kParameters[] = { { /* [225] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [226] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [227] */ @@ -3688,7 +3725,7 @@ constexpr ParameterInfo kParameters[] = { { /* [232] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [233] */ @@ -3713,7 +3750,7 @@ constexpr ParameterInfo kParameters[] = { { /* [237] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [238] */ @@ -3728,12 +3765,12 @@ constexpr ParameterInfo kParameters[] = { { /* [240] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [241] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [242] */ @@ -3753,12 +3790,12 @@ constexpr ParameterInfo kParameters[] = { { /* [245] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [246] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [247] */ @@ -3778,12 +3815,12 @@ constexpr ParameterInfo kParameters[] = { { /* [250] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [251] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [252] */ @@ -3813,7 +3850,7 @@ constexpr ParameterInfo kParameters[] = { { /* [257] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [258] */ @@ -3828,12 +3865,12 @@ constexpr ParameterInfo kParameters[] = { { /* [260] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [261] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [262] */ @@ -3853,12 +3890,12 @@ constexpr ParameterInfo kParameters[] = { { /* [265] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [266] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [267] */ @@ -3878,12 +3915,12 @@ constexpr ParameterInfo kParameters[] = { { /* [270] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[189], + /* matcher indices */ &kMatcherIndices[191], }, { /* [271] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [272] */ @@ -3903,12 +3940,12 @@ constexpr ParameterInfo kParameters[] = { { /* [275] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [276] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [277] */ @@ -3928,12 +3965,12 @@ constexpr ParameterInfo kParameters[] = { { /* [280] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [281] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [282] */ @@ -3953,12 +3990,12 @@ constexpr ParameterInfo kParameters[] = { { /* [285] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [286] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [287] */ @@ -3978,12 +4015,12 @@ constexpr ParameterInfo kParameters[] = { { /* [290] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [291] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [292] */ @@ -4003,12 +4040,12 @@ constexpr ParameterInfo kParameters[] = { { /* [295] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [296] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [297] */ @@ -4023,12 +4060,12 @@ constexpr ParameterInfo kParameters[] = { { /* [299] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [300] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [301] */ @@ -4043,12 +4080,12 @@ constexpr ParameterInfo kParameters[] = { { /* [303] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [304] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [305] */ @@ -4063,12 +4100,12 @@ constexpr ParameterInfo kParameters[] = { { /* [307] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[187], + /* matcher indices */ &kMatcherIndices[189], }, { /* [308] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [309] */ @@ -4083,12 +4120,12 @@ constexpr ParameterInfo kParameters[] = { { /* [311] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [312] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [313] */ @@ -4103,12 +4140,12 @@ constexpr ParameterInfo kParameters[] = { { /* [315] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [316] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [317] */ @@ -4123,12 +4160,12 @@ constexpr ParameterInfo kParameters[] = { { /* [319] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [320] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [321] */ @@ -4163,12 +4200,12 @@ constexpr ParameterInfo kParameters[] = { { /* [327] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [328] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [329] */ @@ -4183,12 +4220,12 @@ constexpr ParameterInfo kParameters[] = { { /* [331] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [332] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [333] */ @@ -4203,12 +4240,12 @@ constexpr ParameterInfo kParameters[] = { { /* [335] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [336] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [337] */ @@ -4223,12 +4260,12 @@ constexpr ParameterInfo kParameters[] = { { /* [339] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[187], + /* matcher indices */ &kMatcherIndices[189], }, { /* [340] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [341] */ @@ -4243,12 +4280,12 @@ constexpr ParameterInfo kParameters[] = { { /* [343] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [344] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [345] */ @@ -4263,12 +4300,12 @@ constexpr ParameterInfo kParameters[] = { { /* [347] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [348] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [349] */ @@ -4283,12 +4320,12 @@ constexpr ParameterInfo kParameters[] = { { /* [351] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[189], + /* matcher indices */ &kMatcherIndices[191], }, { /* [352] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [353] */ @@ -4303,12 +4340,12 @@ constexpr ParameterInfo kParameters[] = { { /* [355] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [356] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [357] */ @@ -4333,7 +4370,7 @@ constexpr ParameterInfo kParameters[] = { { /* [361] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [362] */ @@ -4343,12 +4380,12 @@ constexpr ParameterInfo kParameters[] = { { /* [363] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[181], + /* matcher indices */ &kMatcherIndices[183], }, { /* [364] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [365] */ @@ -4363,12 +4400,12 @@ constexpr ParameterInfo kParameters[] = { { /* [367] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [368] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [369] */ @@ -4383,12 +4420,12 @@ constexpr ParameterInfo kParameters[] = { { /* [371] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [372] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [373] */ @@ -4403,12 +4440,12 @@ constexpr ParameterInfo kParameters[] = { { /* [375] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [376] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [377] */ @@ -4423,12 +4460,12 @@ constexpr ParameterInfo kParameters[] = { { /* [379] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [380] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [381] */ @@ -4483,12 +4520,12 @@ constexpr ParameterInfo kParameters[] = { { /* [391] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [392] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[192], + /* matcher indices */ &kMatcherIndices[194], }, { /* [393] */ @@ -4503,12 +4540,12 @@ constexpr ParameterInfo kParameters[] = { { /* [395] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [396] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [397] */ @@ -4563,7 +4600,7 @@ constexpr ParameterInfo kParameters[] = { { /* [407] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [408] */ @@ -4583,12 +4620,12 @@ constexpr ParameterInfo kParameters[] = { { /* [411] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [412] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [413] */ @@ -4603,12 +4640,12 @@ constexpr ParameterInfo kParameters[] = { { /* [415] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [416] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [417] */ @@ -4653,7 +4690,7 @@ constexpr ParameterInfo kParameters[] = { { /* [425] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [426] */ @@ -4763,12 +4800,12 @@ constexpr ParameterInfo kParameters[] = { { /* [447] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [448] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [449] */ @@ -4808,12 +4845,12 @@ constexpr ParameterInfo kParameters[] = { { /* [456] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[197], + /* matcher indices */ &kMatcherIndices[199], }, { /* [457] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [458] */ @@ -5198,12 +5235,12 @@ constexpr ParameterInfo kParameters[] = { { /* [534] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [535] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [536] */ @@ -5258,12 +5295,12 @@ constexpr ParameterInfo kParameters[] = { { /* [546] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [547] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [548] */ @@ -5348,12 +5385,12 @@ constexpr ParameterInfo kParameters[] = { { /* [564] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[187], + /* matcher indices */ &kMatcherIndices[189], }, { /* [565] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [566] */ @@ -5378,12 +5415,12 @@ constexpr ParameterInfo kParameters[] = { { /* [570] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[185], + /* matcher indices */ &kMatcherIndices[187], }, { /* [571] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [572] */ @@ -5423,7 +5460,7 @@ constexpr ParameterInfo kParameters[] = { { /* [579] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [580] */ @@ -5438,12 +5475,12 @@ constexpr ParameterInfo kParameters[] = { { /* [582] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[179], + /* matcher indices */ &kMatcherIndices[181], }, { /* [583] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [584] */ @@ -5453,12 +5490,12 @@ constexpr ParameterInfo kParameters[] = { { /* [585] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[177], + /* matcher indices */ &kMatcherIndices[179], }, { /* [586] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [587] */ @@ -5468,7 +5505,7 @@ constexpr ParameterInfo kParameters[] = { { /* [588] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[198], + /* matcher indices */ &kMatcherIndices[200], }, { /* [589] */ @@ -5498,12 +5535,12 @@ constexpr ParameterInfo kParameters[] = { { /* [594] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [595] */ /* usage */ ParameterUsage::kSampler, - /* matcher indices */ &kMatcherIndices[191], + /* matcher indices */ &kMatcherIndices[193], }, { /* [596] */ @@ -5713,7 +5750,7 @@ constexpr ParameterInfo kParameters[] = { { /* [637] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [638] */ @@ -5733,7 +5770,7 @@ constexpr ParameterInfo kParameters[] = { { /* [641] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [642] */ @@ -5753,7 +5790,7 @@ constexpr ParameterInfo kParameters[] = { { /* [645] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [646] */ @@ -5773,7 +5810,7 @@ constexpr ParameterInfo kParameters[] = { { /* [649] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [650] */ @@ -6163,7 +6200,7 @@ constexpr ParameterInfo kParameters[] = { { /* [727] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[197], + /* matcher indices */ &kMatcherIndices[199], }, { /* [728] */ @@ -6748,7 +6785,7 @@ constexpr ParameterInfo kParameters[] = { { /* [844] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[198], + /* matcher indices */ &kMatcherIndices[200], }, { /* [845] */ @@ -6758,7 +6795,7 @@ constexpr ParameterInfo kParameters[] = { { /* [846] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [847] */ @@ -6768,17 +6805,17 @@ constexpr ParameterInfo kParameters[] = { { /* [848] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [849] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [850] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [851] */ @@ -6848,12 +6885,12 @@ constexpr ParameterInfo kParameters[] = { { /* [864] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [865] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [866] */ @@ -6943,7 +6980,7 @@ constexpr ParameterInfo kParameters[] = { { /* [883] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[197], + /* matcher indices */ &kMatcherIndices[199], }, { /* [884] */ @@ -6968,27 +7005,27 @@ constexpr ParameterInfo kParameters[] = { { /* [888] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[198], + /* matcher indices */ &kMatcherIndices[200], }, { /* [889] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[193], + /* matcher indices */ &kMatcherIndices[195], }, { /* [890] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[194], + /* matcher indices */ &kMatcherIndices[196], }, { /* [891] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[195], + /* matcher indices */ &kMatcherIndices[197], }, { /* [892] */ /* usage */ ParameterUsage::kTexture, - /* matcher indices */ &kMatcherIndices[196], + /* matcher indices */ &kMatcherIndices[198], }, { /* [893] */ @@ -7123,7 +7160,7 @@ constexpr ParameterInfo kParameters[] = { { /* [919] */ /* usage */ ParameterUsage::kNone, - /* matcher indices */ &kMatcherIndices[183], + /* matcher indices */ &kMatcherIndices[185], }, { /* [920] */ @@ -7183,7 +7220,7 @@ constexpr ParameterInfo kParameters[] = { { /* [931] */ /* usage */ ParameterUsage::kNone, - /* matcher indices */ &kMatcherIndices[173], + /* matcher indices */ &kMatcherIndices[175], }, { /* [932] */ @@ -7486,7 +7523,7 @@ constexpr TemplateTypeInfo kTemplateTypes[] = { { /* [1] */ /* name */ "U", - /* matcher index */ 57, + /* matcher index */ 58, }, { /* [2] */ @@ -7496,7 +7533,7 @@ constexpr TemplateTypeInfo kTemplateTypes[] = { { /* [3] */ /* name */ "U", - /* matcher index */ 54, + /* matcher index */ 55, }, { /* [4] */ @@ -7506,7 +7543,7 @@ constexpr TemplateTypeInfo kTemplateTypes[] = { { /* [5] */ /* name */ "U", - /* matcher index */ 55, + /* matcher index */ 56, }, { /* [6] */ @@ -7516,12 +7553,12 @@ constexpr TemplateTypeInfo kTemplateTypes[] = { { /* [7] */ /* name */ "U", - /* matcher index */ 56, + /* matcher index */ 57, }, { /* [8] */ /* name */ "T", - /* matcher index */ 48, + /* matcher index */ 49, }, { /* [9] */ @@ -7531,22 +7568,22 @@ constexpr TemplateTypeInfo kTemplateTypes[] = { { /* [10] */ /* name */ "T", - /* matcher index */ 53, + /* matcher index */ 54, }, { /* [11] */ /* name */ "T", - /* matcher index */ 50, + /* matcher index */ 51, }, { /* [12] */ /* name */ "T", - /* matcher index */ 52, + /* matcher index */ 53, }, { /* [13] */ /* name */ "T", - /* matcher index */ 51, + /* matcher index */ 52, }, { /* [14] */ @@ -7556,27 +7593,27 @@ constexpr TemplateTypeInfo kTemplateTypes[] = { { /* [15] */ /* name */ "T", - /* matcher index */ 57, + /* matcher index */ 58, }, { /* [16] */ /* name */ "T", - /* matcher index */ 54, + /* matcher index */ 55, }, { /* [17] */ /* name */ "T", - /* matcher index */ 56, + /* matcher index */ 57, }, { /* [18] */ /* name */ "T", - /* matcher index */ 55, + /* matcher index */ 56, }, { /* [19] */ /* name */ "T", - /* matcher index */ 49, + /* matcher index */ 50, }, }; @@ -9952,7 +9989,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[20], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[990], - /* return matcher indices */ &kMatcherIndices[173], + /* return matcher indices */ &kMatcherIndices[175], /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -9963,7 +10000,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[9], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[931], - /* return matcher indices */ &kMatcherIndices[173], + /* return matcher indices */ &kMatcherIndices[175], /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -9974,7 +10011,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[10], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[928], - /* return matcher indices */ &kMatcherIndices[175], + /* return matcher indices */ &kMatcherIndices[177], /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -9985,7 +10022,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[10], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[28], - /* return matcher indices */ &kMatcherIndices[175], + /* return matcher indices */ &kMatcherIndices[177], /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -9996,7 +10033,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[10], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[435], - /* return matcher indices */ &kMatcherIndices[175], + /* return matcher indices */ &kMatcherIndices[177], /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -10117,7 +10154,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[20], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[990], - /* return matcher indices */ &kMatcherIndices[183], + /* return matcher indices */ &kMatcherIndices[185], /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -10128,7 +10165,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[9], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[919], - /* return matcher indices */ &kMatcherIndices[183], + /* return matcher indices */ &kMatcherIndices[185], /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -10843,7 +10880,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[20], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[843], - /* return matcher indices */ &kMatcherIndices[200], + /* return matcher indices */ &kMatcherIndices[202], /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -10865,7 +10902,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[20], /* template numbers */ &kTemplateNumbers[10], /* parameters */ &kParameters[973], - /* return matcher indices */ &kMatcherIndices[199], + /* return matcher indices */ &kMatcherIndices[201], /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, { @@ -12526,7 +12563,7 @@ constexpr OverloadInfo kOverloads[] = { /* template types */ &kTemplateTypes[11], /* template numbers */ &kTemplateNumbers[9], /* parameters */ &kParameters[591], - /* return matcher indices */ &kMatcherIndices[105], + /* return matcher indices */ &kMatcherIndices[173], /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline), }, }; @@ -13358,7 +13395,7 @@ constexpr IntrinsicInfo kBuiltins[] = { }, { /* [106] */ - /* fn atomicCompareExchangeWeak(ptr, read_write>, T, T) -> vec2 */ + /* fn atomicCompareExchangeWeak(ptr, read_write>, T, T) -> __atomic_compare_exchange_result */ /* num overloads */ 1, /* overloads */ &kOverloads[444], }, diff --git a/src/tint/transform/decompose_memory_access.cc b/src/tint/transform/decompose_memory_access.cc index 775cc05f63..a90a6e2711 100644 --- a/src/tint/transform/decompose_memory_access.cc +++ b/src/tint/transform/decompose_memory_access.cc @@ -644,14 +644,34 @@ struct DecomposeMemoryAccess::State { << el_ty->TypeInfo().name; } - auto* ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType()); - auto* func = - b.create(b.Sym(), params, ret_ty, nullptr, - ast::AttributeList{ - atomic, - b.Disable(ast::DisabledValidation::kFunctionHasNoBody), - }, - ast::AttributeList{}); + const ast::Type* ret_ty = nullptr; + + // For intrinsics that return a struct, there is no AST node for it, so create one now. + if (intrinsic->Type() == sem::BuiltinType::kAtomicCompareExchangeWeak) { + auto* str = intrinsic->ReturnType()->As(); + TINT_ASSERT(Transform, str && str->Declaration() == nullptr); + + ast::StructMemberList ast_members; + ast_members.reserve(str->Members().size()); + for (auto& m : str->Members()) { + ast_members.push_back( + b.Member(ctx.Clone(m->Name()), CreateASTTypeFor(ctx, m->Type()))); + } + + auto name = b.Symbols().New("atomic_compare_exchange_weak_ret_type"); + auto* new_str = b.Structure(name, std::move(ast_members)); + ret_ty = b.ty.Of(new_str); + } else { + ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType()); + } + + auto* func = b.create( + b.Symbols().New(std::string{"tint_"} + intrinsic->str()), params, ret_ty, nullptr, + ast::AttributeList{ + atomic, + b.Disable(ast::DisabledValidation::kFunctionHasNoBody), + }, + ast::AttributeList{}); b.AST().AddFunction(func); return func->symbol; @@ -753,6 +773,10 @@ const DecomposeMemoryAccess::Intrinsic* DecomposeMemoryAccess::Intrinsic::Clone( storage_class, type); } +bool DecomposeMemoryAccess::Intrinsic::IsAtomic() const { + return op != Op::kLoad && op != Op::kStore; +} + DecomposeMemoryAccess::DecomposeMemoryAccess() = default; DecomposeMemoryAccess::~DecomposeMemoryAccess() = default; diff --git a/src/tint/transform/decompose_memory_access.h b/src/tint/transform/decompose_memory_access.h index 7a7b783d83..76cb23e2ff 100644 --- a/src/tint/transform/decompose_memory_access.h +++ b/src/tint/transform/decompose_memory_access.h @@ -89,6 +89,9 @@ class DecomposeMemoryAccess final : public Castable sb : SB; @internal(intrinsic_atomic_store_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) +fn tint_atomicStore(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) @internal(intrinsic_atomic_load_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 +fn tint_atomicLoad(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 @internal(intrinsic_atomic_add_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAdd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_sub_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicSub(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_max_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMax(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_min_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMin(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_and_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAnd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_or_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicOr(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_xor_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicXor(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicExchange(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 + +struct atomic_compare_exchange_weak_ret_type { + old_value : i32, + exchanged : bool, +} @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2 +fn tint_atomicCompareExchangeWeak(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) +fn tint_atomicStore_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) @internal(intrinsic_atomic_load_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 +fn tint_atomicLoad_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 @internal(intrinsic_atomic_add_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAdd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_sub_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicSub_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_max_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMax_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_min_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMin_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_and_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAnd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_or_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicOr_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_xor_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicXor_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicExchange_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 + +struct atomic_compare_exchange_weak_ret_type_1 { + old_value : u32, + exchanged : bool, +} @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2 +fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 @stage(compute) @workgroup_size(1) fn main() { - tint_symbol(sb, 16u, 123); - tint_symbol_1(sb, 16u); - tint_symbol_2(sb, 16u, 123); - tint_symbol_3(sb, 16u, 123); - tint_symbol_4(sb, 16u, 123); - tint_symbol_5(sb, 16u, 123); - tint_symbol_6(sb, 16u, 123); - tint_symbol_7(sb, 16u, 123); - tint_symbol_8(sb, 16u, 123); - tint_symbol_9(sb, 16u, 123); - tint_symbol_10(sb, 16u, 123, 345); - tint_symbol_11(sb, 20u, 123u); - tint_symbol_12(sb, 20u); - tint_symbol_13(sb, 20u, 123u); - tint_symbol_14(sb, 20u, 123u); - tint_symbol_15(sb, 20u, 123u); - tint_symbol_16(sb, 20u, 123u); - tint_symbol_17(sb, 20u, 123u); - tint_symbol_18(sb, 20u, 123u); - tint_symbol_19(sb, 20u, 123u); - tint_symbol_20(sb, 20u, 123u); - tint_symbol_21(sb, 20u, 123u, 345u); + tint_atomicStore(sb, 16u, 123); + tint_atomicLoad(sb, 16u); + tint_atomicAdd(sb, 16u, 123); + tint_atomicSub(sb, 16u, 123); + tint_atomicMax(sb, 16u, 123); + tint_atomicMin(sb, 16u, 123); + tint_atomicAnd(sb, 16u, 123); + tint_atomicOr(sb, 16u, 123); + tint_atomicXor(sb, 16u, 123); + tint_atomicExchange(sb, 16u, 123); + tint_atomicCompareExchangeWeak(sb, 16u, 123, 345); + tint_atomicStore_1(sb, 20u, 123u); + tint_atomicLoad_1(sb, 20u); + tint_atomicAdd_1(sb, 20u, 123u); + tint_atomicSub_1(sb, 20u, 123u); + tint_atomicMax_1(sb, 20u, 123u); + tint_atomicMin_1(sb, 20u, 123u); + tint_atomicAnd_1(sb, 20u, 123u); + tint_atomicOr_1(sb, 20u, 123u); + tint_atomicXor_1(sb, 20u, 123u); + tint_atomicExchange_1(sb, 20u, 123u); + tint_atomicCompareExchangeWeak_1(sb, 20u, 123u, 345u); } )"; @@ -2604,95 +2614,105 @@ struct SB { auto* expect = R"( @internal(intrinsic_atomic_store_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) +fn tint_atomicStore(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) @internal(intrinsic_atomic_load_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 +fn tint_atomicLoad(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32 @internal(intrinsic_atomic_add_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAdd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_sub_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicSub(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_max_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMax(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_min_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicMin(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_and_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicAnd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_or_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicOr(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_xor_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicXor(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 +fn tint_atomicExchange(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32 + +struct atomic_compare_exchange_weak_ret_type { + old_value : i32, + exchanged : bool, +} @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2 +fn tint_atomicCompareExchangeWeak(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) +fn tint_atomicStore_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) @internal(intrinsic_atomic_load_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 +fn tint_atomicLoad_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32 @internal(intrinsic_atomic_add_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAdd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_sub_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicSub_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_max_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMax_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_min_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicMin_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_and_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicAnd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_or_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicOr_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_xor_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicXor_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 +fn tint_atomicExchange_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32 + +struct atomic_compare_exchange_weak_ret_type_1 { + old_value : u32, + exchanged : bool, +} @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body) -fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2 +fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1 @stage(compute) @workgroup_size(1) fn main() { - tint_symbol(sb, 16u, 123); - tint_symbol_1(sb, 16u); - tint_symbol_2(sb, 16u, 123); - tint_symbol_3(sb, 16u, 123); - tint_symbol_4(sb, 16u, 123); - tint_symbol_5(sb, 16u, 123); - tint_symbol_6(sb, 16u, 123); - tint_symbol_7(sb, 16u, 123); - tint_symbol_8(sb, 16u, 123); - tint_symbol_9(sb, 16u, 123); - tint_symbol_10(sb, 16u, 123, 345); - tint_symbol_11(sb, 20u, 123u); - tint_symbol_12(sb, 20u); - tint_symbol_13(sb, 20u, 123u); - tint_symbol_14(sb, 20u, 123u); - tint_symbol_15(sb, 20u, 123u); - tint_symbol_16(sb, 20u, 123u); - tint_symbol_17(sb, 20u, 123u); - tint_symbol_18(sb, 20u, 123u); - tint_symbol_19(sb, 20u, 123u); - tint_symbol_20(sb, 20u, 123u); - tint_symbol_21(sb, 20u, 123u, 345u); + tint_atomicStore(sb, 16u, 123); + tint_atomicLoad(sb, 16u); + tint_atomicAdd(sb, 16u, 123); + tint_atomicSub(sb, 16u, 123); + tint_atomicMax(sb, 16u, 123); + tint_atomicMin(sb, 16u, 123); + tint_atomicAnd(sb, 16u, 123); + tint_atomicOr(sb, 16u, 123); + tint_atomicXor(sb, 16u, 123); + tint_atomicExchange(sb, 16u, 123); + tint_atomicCompareExchangeWeak(sb, 16u, 123, 345); + tint_atomicStore_1(sb, 20u, 123u); + tint_atomicLoad_1(sb, 20u); + tint_atomicAdd_1(sb, 20u, 123u); + tint_atomicSub_1(sb, 20u, 123u); + tint_atomicMax_1(sb, 20u, 123u); + tint_atomicMin_1(sb, 20u, 123u); + tint_atomicAnd_1(sb, 20u, 123u); + tint_atomicOr_1(sb, 20u, 123u); + tint_atomicXor_1(sb, 20u, 123u); + tint_atomicExchange_1(sb, 20u, 123u); + tint_atomicCompareExchangeWeak_1(sb, 20u, 123u, 345u); } @group(0) @binding(0) var sb : SB; diff --git a/src/tint/transform/manager.cc b/src/tint/transform/manager.cc index 823474c63b..e5f7682ead 100644 --- a/src/tint/transform/manager.cc +++ b/src/tint/transform/manager.cc @@ -49,7 +49,7 @@ Output Manager::Run(const Program* program, const DataMap& data) const { Output out; for (const auto& transform : transforms_) { if (!transform->ShouldRun(in, data)) { - TINT_IF_PRINT_PROGRAM(std::cout << "Skipping " << transform->TypeInfo().name); + TINT_IF_PRINT_PROGRAM(std::cout << "Skipping " << transform->TypeInfo().name << std::endl); continue; } TINT_IF_PRINT_PROGRAM(print_program("Input to", transform.get())); diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc index ccce06a89a..8a48156fe0 100644 --- a/src/tint/writer/glsl/generator_impl.cc +++ b/src/tint/writer/glsl/generator_impl.cc @@ -911,39 +911,56 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, return true; } case sem::BuiltinType::kAtomicCompareExchangeWeak: { - return CallBuiltinHelper( - out, expr, builtin, [&](TextBuffer* b, const std::vector& params) { - { - auto pre = line(b); - if (!EmitTypeAndName(pre, builtin->ReturnType(), ast::StorageClass::kNone, - ast::Access::kUndefined, "result")) { - return false; - } - pre << ";"; + // Emit the builtin return type unique to this overload. This does not + // exist in the AST, so it will not be generated in Generate(). + if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { + return false; + } + + auto* dest = expr->args[0]; + auto* compare_value = expr->args[1]; + auto* value = expr->args[2]; + + std::string result = UniqueIdentifier("atomic_compare_result"); + + { + auto pre = line(); + if (!EmitTypeAndName(pre, builtin->ReturnType(), ast::StorageClass::kNone, + ast::Access::kUndefined, result)) { + return false; + } + pre << ";"; + } + { + auto pre = line(); + pre << result << ".old_value = atomicCompSwap"; + { + ScopedParen sp(pre); + if (!EmitExpression(pre, dest)) { + return false; } - { - auto pre = line(b); - pre << "result.x = atomicCompSwap"; - { - ScopedParen sp(pre); - pre << params[0]; - pre << ", " << params[1]; - pre << ", " << params[2]; - } - pre << ";"; + pre << ", "; + if (!EmitExpression(pre, compare_value)) { + return false; } - { - auto pre = line(b); - pre << "result.y = result.x == " << params[2] << " ? "; - if (TypeOf(expr->args[2])->Is()) { - pre << "1u : 0u;"; - } else { - pre << "1 : 0;"; - } + pre << ", "; + if (!EmitExpression(pre, value)) { + return false; } - line(b) << "return result;"; - return true; - }); + } + pre << ";"; + } + { + auto pre = line(); + pre << result << ".exchanged = " << result << ".old_value == "; + if (!EmitExpression(pre, compare_value)) { + return false; + } + pre << ";"; + } + + out << result; + return true; } case sem::BuiltinType::kAtomicAdd: diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h index 72def34c62..1d566ca509 100644 --- a/src/tint/writer/glsl/generator_impl.h +++ b/src/tint/writer/glsl/generator_impl.h @@ -174,14 +174,6 @@ class GeneratorImpl : public TextGenerator { /// @param builtin the semantic information for the barrier builtin /// @returns true if the call expression is emitted bool EmitBarrierCall(std::ostream& out, const sem::Builtin* builtin); - /// Handles generating an atomic intrinsic call for a storage buffer variable - /// @param out the output of the expression stream - /// @param expr the call expression - /// @param intrinsic the atomic intrinsic - /// @returns true if the call expression is emitted - bool EmitStorageAtomicCall(std::ostream& out, - const ast::CallExpression* expr, - const transform::DecomposeMemoryAccess::Intrinsic* intrinsic); /// Handles generating an atomic builtin call for a workgroup variable /// @param out the output of the expression stream /// @param expr the call expression diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc index d6a5fa7549..a2cac0f8a9 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc @@ -968,7 +968,10 @@ bool GeneratorImpl::EmitFunctionCall(std::ostream& out, case ast::StorageClass::kUniform: return EmitUniformBufferAccess(out, expr, intrinsic); case ast::StorageClass::kStorage: - return EmitStorageBufferAccess(out, expr, intrinsic); + if (!intrinsic->IsAtomic()) { + return EmitStorageBufferAccess(out, expr, intrinsic); + } + break; default: TINT_UNREACHABLE(Writer, diagnostics_) << "unsupported DecomposeMemoryAccess::Intrinsic storage class:" @@ -1445,19 +1448,10 @@ bool GeneratorImpl::EmitStorageBufferAccess( << static_cast(intrinsic->type); return false; } - - case Op::kAtomicLoad: - case Op::kAtomicStore: - case Op::kAtomicAdd: - case Op::kAtomicSub: - case Op::kAtomicMax: - case Op::kAtomicMin: - case Op::kAtomicAnd: - case Op::kAtomicOr: - case Op::kAtomicXor: - case Op::kAtomicExchange: - case Op::kAtomicCompareExchangeWeak: - return EmitStorageAtomicCall(out, expr, intrinsic); + default: + // Break out to error case below/ + // Note that atomic intrinsics are generated as functions. + break; } TINT_UNREACHABLE(Writer, diagnostics_) @@ -1465,32 +1459,127 @@ bool GeneratorImpl::EmitStorageBufferAccess( return false; } -bool GeneratorImpl::EmitStorageAtomicCall( - std::ostream& out, - const ast::CallExpression* expr, +bool GeneratorImpl::EmitStorageAtomicIntrinsic( + const ast::Function* func, const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) { using Op = transform::DecomposeMemoryAccess::Intrinsic::Op; - auto* result_ty = TypeOf(expr); + const sem::Function* sem_func = builder_.Sem().Get(func); + auto* result_ty = sem_func->ReturnType(); + const auto& params = sem_func->Parameters(); + const auto name = builder_.Symbols().NameFor(func->symbol); + auto& buf = *current_buffer_; - auto& buf = helpers_; + auto rmw = [&](const char* hlsl) -> bool { + { + auto fn = line(&buf); + if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined, + name)) { + return false; + } + fn << "(RWByteAddressBuffer buffer, uint offset, "; + if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined, + "value")) { + return false; + } + fn << ") {"; + } - // generate_helper() generates a helper function that translates the - // DecomposeMemoryAccess::Intrinsic call into the corresponding HLSL - // atomic intrinsic function. - auto generate_helper = [&]() -> std::string { - auto rmw = [&](const char* wgsl, const char* hlsl) -> std::string { - auto name = UniqueIdentifier(wgsl); + buf.IncrementIndent(); + TINT_DEFER({ + buf.DecrementIndent(); + line(&buf) << "}"; + line(&buf); + }); + + { + auto l = line(&buf); + if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined, + "original_value")) { + return false; + } + l << " = 0;"; + } + { + auto l = line(&buf); + l << "buffer." << hlsl << "(offset, "; + if (intrinsic->op == Op::kAtomicSub) { + l << "-"; + } + l << "value, original_value);"; + } + line(&buf) << "return original_value;"; + return true; + }; + + switch (intrinsic->op) { + case Op::kAtomicAdd: + return rmw("InterlockedAdd"); + + case Op::kAtomicSub: + // Use add with the operand negated. + return rmw("InterlockedAdd"); + + case Op::kAtomicMax: + return rmw("InterlockedMax"); + + case Op::kAtomicMin: + return rmw("InterlockedMin"); + + case Op::kAtomicAnd: + return rmw("InterlockedAnd"); + + case Op::kAtomicOr: + return rmw("InterlockedOr"); + + case Op::kAtomicXor: + return rmw("InterlockedXor"); + + case Op::kAtomicExchange: + return rmw("InterlockedExchange"); + + case Op::kAtomicLoad: { + // HLSL does not have an InterlockedLoad, so we emulate it with + // InterlockedOr using 0 as the OR value { auto fn = line(&buf); if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined, name)) { - return ""; + return false; } - fn << "(RWByteAddressBuffer buffer, uint offset, "; - if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, + fn << "(RWByteAddressBuffer buffer, uint offset) {"; + } + + buf.IncrementIndent(); + TINT_DEFER({ + buf.DecrementIndent(); + line(&buf) << "}"; + line(&buf); + }); + + { + auto l = line(&buf); + if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined, "value")) { - return ""; + return false; + } + l << " = 0;"; + } + + line(&buf) << "buffer.InterlockedOr(offset, 0, value);"; + line(&buf) << "return value;"; + return true; + } + case Op::kAtomicStore: { + // HLSL does not have an InterlockedStore, so we emulate it with + // InterlockedExchange and discard the returned value + auto* value_ty = params[2]->Type()->UnwrapRef(); + { + auto fn = line(&buf); + fn << "void " << name << "(RWByteAddressBuffer buffer, uint offset, "; + if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone, + ast::Access::kUndefined, "value")) { + return false; } fn << ") {"; } @@ -1504,191 +1593,73 @@ bool GeneratorImpl::EmitStorageAtomicCall( { auto l = line(&buf); - if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, "original_value")) { - return ""; + if (!EmitTypeAndName(l, value_ty, ast::StorageClass::kNone, ast::Access::kUndefined, + "ignored")) { + return false; } - l << " = 0;"; + l << ";"; } + line(&buf) << "buffer.InterlockedExchange(offset, value, ignored);"; + return true; + } + case Op::kAtomicCompareExchangeWeak: { + // NOTE: We don't need to emit the return type struct here as DecomposeMemoryAccess + // already added it to the AST, and it should have already been emitted by now. + auto* value_ty = params[2]->Type()->UnwrapRef(); { + auto fn = line(&buf); + if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, + ast::Access::kUndefined, name)) { + return false; + } + fn << "(RWByteAddressBuffer buffer, uint offset, "; + if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone, + ast::Access::kUndefined, "compare")) { + return false; + } + fn << ", "; + if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone, + ast::Access::kUndefined, "value")) { + return false; + } + fn << ") {"; + } + + buf.IncrementIndent(); + TINT_DEFER({ + buf.DecrementIndent(); + line(&buf) << "}"; + line(&buf); + }); + + { // T result = {0}; auto l = line(&buf); - l << "buffer." << hlsl << "(offset, "; - if (intrinsic->op == Op::kAtomicSub) { - l << "-"; + if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone, + ast::Access::kUndefined, "result")) { + return false; } - l << "value, original_value);"; + l << "="; + if (!EmitZeroValue(l, result_ty)) { + return false; + } + l << ";"; } - line(&buf) << "return original_value;"; - return name; - }; - switch (intrinsic->op) { - case Op::kAtomicAdd: - return rmw("atomicAdd", "InterlockedAdd"); + line(&buf) << "buffer.InterlockedCompareExchange(offset, compare, value, " + "result.old_value);"; + line(&buf) << "result.exchanged = result.old_value == compare;"; + line(&buf) << "return result;"; - case Op::kAtomicSub: - // Use add with the operand negated. - return rmw("atomicSub", "InterlockedAdd"); - - case Op::kAtomicMax: - return rmw("atomicMax", "InterlockedMax"); - - case Op::kAtomicMin: - return rmw("atomicMin", "InterlockedMin"); - - case Op::kAtomicAnd: - return rmw("atomicAnd", "InterlockedAnd"); - - case Op::kAtomicOr: - return rmw("atomicOr", "InterlockedOr"); - - case Op::kAtomicXor: - return rmw("atomicXor", "InterlockedXor"); - - case Op::kAtomicExchange: - return rmw("atomicExchange", "InterlockedExchange"); - - case Op::kAtomicLoad: { - // HLSL does not have an InterlockedLoad, so we emulate it with - // InterlockedOr using 0 as the OR value - auto name = UniqueIdentifier("atomicLoad"); - { - auto fn = line(&buf); - if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, name)) { - return ""; - } - fn << "(RWByteAddressBuffer buffer, uint offset) {"; - } - - buf.IncrementIndent(); - TINT_DEFER({ - buf.DecrementIndent(); - line(&buf) << "}"; - line(&buf); - }); - - { - auto l = line(&buf); - if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, "value")) { - return ""; - } - l << " = 0;"; - } - - line(&buf) << "buffer.InterlockedOr(offset, 0, value);"; - line(&buf) << "return value;"; - return name; - } - case Op::kAtomicStore: { - // HLSL does not have an InterlockedStore, so we emulate it with - // InterlockedExchange and discard the returned value - auto* value_ty = TypeOf(expr->args[2])->UnwrapRef(); - auto name = UniqueIdentifier("atomicStore"); - { - auto fn = line(&buf); - fn << "void " << name << "(RWByteAddressBuffer buffer, uint offset, "; - if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, "value")) { - return ""; - } - fn << ") {"; - } - - buf.IncrementIndent(); - TINT_DEFER({ - buf.DecrementIndent(); - line(&buf) << "}"; - line(&buf); - }); - - { - auto l = line(&buf); - if (!EmitTypeAndName(l, value_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, "ignored")) { - return ""; - } - l << ";"; - } - line(&buf) << "buffer.InterlockedExchange(offset, value, ignored);"; - return name; - } - case Op::kAtomicCompareExchangeWeak: { - auto* value_ty = TypeOf(expr->args[2])->UnwrapRef(); - - auto name = UniqueIdentifier("atomicCompareExchangeWeak"); - { - auto fn = line(&buf); - if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, name)) { - return ""; - } - fn << "(RWByteAddressBuffer buffer, uint offset, "; - if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, "compare")) { - return ""; - } - fn << ", "; - if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, "value")) { - return ""; - } - fn << ") {"; - } - - buf.IncrementIndent(); - TINT_DEFER({ - buf.DecrementIndent(); - line(&buf) << "}"; - line(&buf); - }); - - { // T result = {0, 0}; - auto l = line(&buf); - if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone, - ast::Access::kUndefined, "result")) { - return ""; - } - l << " = {0, 0};"; - } - line(&buf) << "buffer.InterlockedCompareExchange(offset, compare, " - "value, result.x);"; - line(&buf) << "result.y = result.x == compare;"; - line(&buf) << "return result;"; - return name; - } - default: - break; + return true; } - TINT_UNREACHABLE(Writer, diagnostics_) - << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: " - << static_cast(intrinsic->op); - return ""; - }; - - auto func = utils::GetOrCreate(dma_intrinsics_, DMAIntrinsic{intrinsic->op, intrinsic->type}, - generate_helper); - if (func.empty()) { - return false; + default: + break; } - out << func; - { - ScopedParen sp(out); - bool first = true; - for (auto* arg : expr->args) { - if (!first) { - out << ", "; - } - first = false; - if (!EmitExpression(out, arg)) { - return false; - } - } - } - - return true; + TINT_UNREACHABLE(Writer, diagnostics_) + << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: " + << static_cast(intrinsic->op); + return false; } bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, @@ -1788,6 +1759,12 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, return true; } case sem::BuiltinType::kAtomicCompareExchangeWeak: { + // Emit the builtin return type unique to this overload. This does not + // exist in the AST, so it will not be generated in Generate(). + if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { + return false; + } + auto* dest = expr->args[0]; auto* compare_value = expr->args[1]; auto* value = expr->args[2]; @@ -1807,7 +1784,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, pre << ";"; } - { // InterlockedCompareExchange(dst, compare, value, result.x); + { // InterlockedCompareExchange(dst, compare, value, result.old_value); auto pre = line(); pre << "InterlockedCompareExchange"; { @@ -1819,14 +1796,13 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out, if (!EmitExpression(pre, value)) { return false; } - pre << ", " << result << ".x"; + pre << ", " << result << ".old_value"; } pre << ";"; } - { // result.y = result.x == compare; - line() << result << ".y = " << result << ".x == " << compare << ";"; - } + // result.exchanged = result.old_value == compare; + line() << result << ".exchanged = " << result << ".old_value == " << compare << ";"; out << result; return true; @@ -2740,6 +2716,17 @@ bool GeneratorImpl::EmitIf(const ast::IfStatement* stmt) { bool GeneratorImpl::EmitFunction(const ast::Function* func) { auto* sem = builder_.Sem().Get(func); + // Emit storage atomic helpers + if (auto* intrinsic = + ast::GetAttribute(func->attributes)) { + if (intrinsic->storage_class == ast::StorageClass::kStorage && intrinsic->IsAtomic()) { + if (!EmitStorageAtomicIntrinsic(func, intrinsic)) { + return false; + } + } + return true; + } + if (ast::HasAttribute(func->attributes)) { // An internal function. Do not emit. return true; @@ -3755,13 +3742,9 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { ScopedIndent si(b); for (auto* mem : str->Members()) { auto mem_name = builder_.Symbols().NameFor(mem->Name()); - auto* ty = mem->Type(); - auto out = line(b); - std::string pre, post; - if (auto* decl = mem->Declaration()) { for (auto* attr : decl->attributes) { if (auto* location = attr->As()) { @@ -3826,7 +3809,6 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) { } line(b) << "};"; - return true; } diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h index 86bbd7dedc..0e8ca4c2be 100644 --- a/src/tint/writer/hlsl/generator_impl.h +++ b/src/tint/writer/hlsl/generator_impl.h @@ -187,6 +187,12 @@ class GeneratorImpl : public TextGenerator { bool EmitStorageAtomicCall(std::ostream& out, const ast::CallExpression* expr, const transform::DecomposeMemoryAccess::Intrinsic* intrinsic); + /// Handles generating the helper function for the atomic intrinsic function + /// @param func the function + /// @param intrinsic the atomic intrinsic + /// @returns true if the function is emitted + bool EmitStorageAtomicIntrinsic(const ast::Function* func, + const transform::DecomposeMemoryAccess::Intrinsic* intrinsic); /// Handles generating an atomic intrinsic call for a workgroup variable /// @param out the output of the expression stream /// @param expr the call expression @@ -511,7 +517,6 @@ class GeneratorImpl : public TextGenerator { TextBuffer helpers_; // Helper functions emitted at the top of the output std::function emit_continuing_; - std::unordered_map dma_intrinsics_; std::unordered_map matrix_scalar_ctors_; std::unordered_map builtins_; std::unordered_map structure_builders_; diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index 54d9164823..578e78df8d 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -806,6 +806,12 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out, return call("atomic_exchange_explicit", true); case sem::BuiltinType::kAtomicCompareExchangeWeak: { + // Emit the builtin return type unique to this overload. This does not + // exist in the AST, so it will not be generated in Generate(). + if (!EmitStructType(&helpers_, builtin->ReturnType()->As())) { + return false; + } + auto* ptr_ty = TypeOf(expr->args[0])->UnwrapRef()->As(); auto sc = ptr_ty->StorageClass(); @@ -816,7 +822,8 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out, line(&buf) << "template "; { auto f = line(&buf); - f << "vec " << name << "("; + auto str_name = StructName(builtin->ReturnType()->As()); + f << str_name << " " << name << "("; if (!EmitStorageClass(f, sc)) { return ""; } @@ -830,12 +837,12 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out, line(&buf); }); - line(&buf) << "T prev_value = compare;"; - line(&buf) << "bool matched = " + line(&buf) << "T old_value = compare;"; + line(&buf) << "bool exchanged = " "atomic_compare_exchange_weak_explicit(atomic, " - "&prev_value, value, memory_order_relaxed, " + "&old_value, value, memory_order_relaxed, " "memory_order_relaxed);"; - line(&buf) << "return {prev_value, matched};"; + line(&buf) << "return {old_value, exchanged};"; return name; }); diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc index e2d1fac831..eb8795f714 100644 --- a/src/tint/writer/spirv/builder.cc +++ b/src/tint/writer/spirv/builder.cc @@ -3201,42 +3201,12 @@ bool Builder::GenerateAtomicBuiltin(const sem::Call* call, return false; } - // zero := T(0) - // one := T(1) - uint32_t zero = 0; - uint32_t one = 0; - if (value_sem_type->Is()) { - zero = GenerateConstantIfNeeded(ScalarConstant::I32(0u)); - one = GenerateConstantIfNeeded(ScalarConstant::I32(1u)); - } else if (value_sem_type->Is()) { - zero = GenerateConstantIfNeeded(ScalarConstant::U32(0u)); - one = GenerateConstantIfNeeded(ScalarConstant::U32(1u)); - } else { - TINT_UNREACHABLE(Writer, builder_.Diagnostics()) - << "unsupported atomic type " << value_sem_type->TypeInfo().name; - } - if (zero == 0 || one == 0) { - return false; - } - - // xchg_success := values_equal ? one : zero - auto xchg_success = result_op(); - if (!push_function_inst(spv::Op::OpSelect, { - Operand(value_type), - xchg_success, - values_equal, - Operand(one), - Operand(zero), - })) { - return false; - } - - // result := vec2(original_value, xchg_success) + // result := __atomic_compare_exchange_result(original_value, values_equal) return push_function_inst(spv::Op::OpCompositeConstruct, { result_type, result_id, original_value, - xchg_success, + values_equal, }); } default: diff --git a/src/tint/writer/spirv/builder_builtin_test.cc b/src/tint/writer/spirv/builder_builtin_test.cc index 59a567aae1..734007429c 100644 --- a/src/tint/writer/spirv/builder_builtin_test.cc +++ b/src/tint/writer/spirv/builder_builtin_test.cc @@ -2018,15 +2018,15 @@ OpReturn TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) { // struct S { - // u : atomic; - // i : atomic; + // u : atomic, + // i : atomic, // } // // @binding(1) @group(2) var b : S; // // fn a_func() { - // let u : vec2 = atomicCompareExchangeWeak(&b.u, 10u); - // let i : vec2 = atomicCompareExchangeWeak(&b.i, 10); + // let u = atomicCompareExchangeWeak(&b.u, 10u, 20u); + // let i = atomicCompareExchangeWeak(&b.i, 10, 10); // } auto* s = Structure("S", { Member("u", ty.atomic()), @@ -2040,10 +2040,10 @@ TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) { Func("a_func", {}, ty.void_(), ast::StatementList{ - Decl(Let("u", ty.vec2(), + Decl(Let("u", nullptr, Call("atomicCompareExchangeWeak", AddressOf(MemberAccessor("b", "u")), 10_u, 20_u))), - Decl(Let("i", ty.vec2(), + Decl(Let("i", nullptr, Call("atomicCompareExchangeWeak", AddressOf(MemberAccessor("b", "i")), 10_i, 20_i))), }, @@ -2062,33 +2062,29 @@ TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) { %1 = OpVariable %2 StorageBuffer %7 = OpTypeVoid %6 = OpTypeFunction %7 -%11 = OpTypeVector %4 2 -%12 = OpConstant %4 1 -%13 = OpConstant %4 0 -%15 = OpTypePointer StorageBuffer %4 -%17 = OpConstant %4 20 -%18 = OpConstant %4 10 -%19 = OpTypeBool -%24 = OpTypeVector %5 2 -%26 = OpTypePointer StorageBuffer %5 -%28 = OpConstant %5 20 -%29 = OpConstant %5 10 -%32 = OpConstant %5 0 -%33 = OpConstant %5 1 +%12 = OpTypeBool +%11 = OpTypeStruct %4 %12 +%13 = OpConstant %4 1 +%14 = OpConstant %4 0 +%16 = OpTypePointer StorageBuffer %4 +%18 = OpConstant %4 20 +%19 = OpConstant %4 10 +%23 = OpTypeStruct %5 %12 +%25 = OpTypePointer StorageBuffer %5 +%27 = OpConstant %5 20 +%28 = OpConstant %5 10 )"; auto got_types = DumpInstructions(b.types()); EXPECT_EQ(expected_types, got_types); - auto* expected_instructions = R"(%16 = OpAccessChain %15 %1 %13 -%20 = OpAtomicCompareExchange %4 %16 %12 %13 %13 %17 %18 -%21 = OpIEqual %19 %20 %17 -%22 = OpSelect %4 %21 %12 %13 -%10 = OpCompositeConstruct %11 %20 %22 -%27 = OpAccessChain %26 %1 %12 -%30 = OpAtomicCompareExchange %5 %27 %12 %13 %13 %28 %29 -%31 = OpIEqual %19 %30 %28 -%34 = OpSelect %5 %31 %33 %32 -%23 = OpCompositeConstruct %24 %30 %34 + auto* expected_instructions = R"(%17 = OpAccessChain %16 %1 %14 +%20 = OpAtomicCompareExchange %4 %17 %13 %14 %14 %18 %19 +%21 = OpIEqual %12 %20 %18 +%10 = OpCompositeConstruct %11 %20 %21 +%26 = OpAccessChain %25 %1 %13 +%29 = OpAtomicCompareExchange %5 %26 %13 %14 %14 %27 %28 +%30 = OpIEqual %12 %29 %27 +%22 = OpCompositeConstruct %23 %29 %30 OpReturn )"; auto got_instructions = DumpInstructions(b.functions()[0].instructions()); diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.hlsl b/test/tint/bug/chromium/1273230.wgsl.expected.hlsl index 5b6f132b18..05d4fe035b 100644 --- a/test/tint/bug/chromium/1273230.wgsl.expected.hlsl +++ b/test/tint/bug/chromium/1273230.wgsl.expected.hlsl @@ -2,24 +2,6 @@ uint value_or_one_if_zero_uint(uint value) { return value == 0u ? 1u : value; } -uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { - uint value = 0; - buffer.InterlockedOr(offset, 0, value); - return value; -} - -int atomicLoad_2(RWByteAddressBuffer buffer, uint offset) { - int value = 0; - buffer.InterlockedOr(offset, 0, value); - return value; -} - -int atomicAdd_1(RWByteAddressBuffer buffer, uint offset, int value) { - int original_value = 0; - buffer.InterlockedAdd(offset, value, original_value); - return original_value; -} - void marg8uintin() { } @@ -61,19 +43,40 @@ float3 loadPosition(uint vertexIndex) { return position; } +uint tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) { + uint value = 0; + buffer.InterlockedOr(offset, 0, value); + return value; +} + + +int tint_atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { + int value = 0; + buffer.InterlockedOr(offset, 0, value); + return value; +} + + void doIgnore() { uint g43 = uniforms[0].x; uint kj6 = dbg.Load(20u); - uint b53 = atomicLoad_1(counters, (4u * uint(0))); + uint b53 = tint_atomicLoad(counters, (4u * uint(0))); uint rwg = indices.Load((4u * uint(0))); float rb5 = asfloat(positions.Load((4u * uint(0)))); - int g55 = atomicLoad_2(LUT, (4u * uint(0))); + int g55 = tint_atomicLoad_1(LUT, (4u * uint(0))); } struct tint_symbol_1 { uint3 GlobalInvocationID : SV_DispatchThreadID; }; +int tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, int value) { + int original_value = 0; + buffer.InterlockedAdd(offset, value, original_value); + return original_value; +} + + void main_count_inner(uint3 GlobalInvocationID) { uint triangleIndex = GlobalInvocationID.x; if ((triangleIndex >= uniforms[0].x)) { @@ -89,7 +92,7 @@ void main_count_inner(uint3 GlobalInvocationID) { float3 center = (((p0 + p2) + p1) / 3.0f); float3 voxelPos = toVoxelPos(p1); uint lIndex = toIndex1D(uniforms[0].y, p0); - int triangleOffset = atomicAdd_1(LUT, (4u * i1), 1); + int triangleOffset = tint_atomicAdd(LUT, (4u * i1), 1); } [numthreads(128, 1, 1)] diff --git a/test/tint/bug/tint/1113.wgsl.expected.hlsl b/test/tint/bug/tint/1113.wgsl.expected.hlsl index fa8e7e281e..e201d4a83c 100644 --- a/test/tint/bug/tint/1113.wgsl.expected.hlsl +++ b/test/tint/bug/tint/1113.wgsl.expected.hlsl @@ -2,35 +2,6 @@ uint value_or_one_if_zero_uint(uint value) { return value == 0u ? 1u : value; } -uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { - uint value = 0; - buffer.InterlockedOr(offset, 0, value); - return value; -} - -int atomicLoad_2(RWByteAddressBuffer buffer, uint offset) { - int value = 0; - buffer.InterlockedOr(offset, 0, value); - return value; -} - -uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) { - uint original_value = 0; - buffer.InterlockedAdd(offset, value, original_value); - return original_value; -} - -void atomicStore_1(RWByteAddressBuffer buffer, uint offset, int value) { - int ignored; - buffer.InterlockedExchange(offset, value, ignored); -} - -int atomicAdd_2(RWByteAddressBuffer buffer, uint offset, int value) { - int original_value = 0; - buffer.InterlockedAdd(offset, value, original_value); - return original_value; -} - cbuffer cbuffer_uniforms : register(b0, space0) { uint4 uniforms[3]; }; @@ -69,19 +40,40 @@ float3 loadPosition(uint vertexIndex) { return position; } +uint tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) { + uint value = 0; + buffer.InterlockedOr(offset, 0, value); + return value; +} + + +int tint_atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { + int value = 0; + buffer.InterlockedOr(offset, 0, value); + return value; +} + + void doIgnore() { uint g42 = uniforms[0].x; uint kj6 = dbg.Load(20u); - uint b53 = atomicLoad_1(counters, (4u * uint(0))); + uint b53 = tint_atomicLoad(counters, (4u * uint(0))); uint rwg = indices.Load((4u * uint(0))); float rb5 = asfloat(positions.Load((4u * uint(0)))); - int g55 = atomicLoad_2(LUT, (4u * uint(0))); + int g55 = tint_atomicLoad_1(LUT, (4u * uint(0))); } struct tint_symbol_1 { uint3 GlobalInvocationID : SV_DispatchThreadID; }; +uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) { + uint original_value = 0; + buffer.InterlockedAdd(offset, value, original_value); + return original_value; +} + + void main_count_inner(uint3 GlobalInvocationID) { uint triangleIndex = GlobalInvocationID.x; if ((triangleIndex >= uniforms[0].x)) { @@ -97,7 +89,7 @@ void main_count_inner(uint3 GlobalInvocationID) { float3 center = (((p0 + p1) + p2) / 3.0f); float3 voxelPos = toVoxelPos(center); uint voxelIndex = toIndex1D(uniforms[0].y, voxelPos); - uint acefg = atomicAdd_1(counters, (4u * voxelIndex), 1u); + uint acefg = tint_atomicAdd(counters, (4u * voxelIndex), 1u); if ((triangleIndex == 0u)) { dbg.Store(16u, asuint(uniforms[0].y)); dbg.Store(32u, asuint(center.x)); @@ -116,6 +108,19 @@ struct tint_symbol_3 { uint3 GlobalInvocationID : SV_DispatchThreadID; }; +uint tint_atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) { + uint original_value = 0; + buffer.InterlockedAdd(offset, value, original_value); + return original_value; +} + + +void tint_atomicStore(RWByteAddressBuffer buffer, uint offset, int value) { + int ignored; + buffer.InterlockedExchange(offset, value, ignored); +} + + void main_create_lut_inner(uint3 GlobalInvocationID) { uint voxelIndex = GlobalInvocationID.x; doIgnore(); @@ -123,13 +128,13 @@ void main_create_lut_inner(uint3 GlobalInvocationID) { if ((voxelIndex >= maxVoxels)) { return; } - uint numTriangles = atomicLoad_1(counters, (4u * voxelIndex)); + uint numTriangles = tint_atomicLoad(counters, (4u * voxelIndex)); int offset = -1; if ((numTriangles > 0u)) { - const uint tint_symbol_6 = atomicAdd_1(dbg, 0u, numTriangles); + const uint tint_symbol_6 = tint_atomicAdd_1(dbg, 0u, numTriangles); offset = int(tint_symbol_6); } - atomicStore_1(LUT, (4u * voxelIndex), offset); + tint_atomicStore(LUT, (4u * voxelIndex), offset); } [numthreads(128, 1, 1)] @@ -142,6 +147,13 @@ struct tint_symbol_5 { uint3 GlobalInvocationID : SV_DispatchThreadID; }; +int tint_atomicAdd_2(RWByteAddressBuffer buffer, uint offset, int value) { + int original_value = 0; + buffer.InterlockedAdd(offset, value, original_value); + return original_value; +} + + void main_sort_triangles_inner(uint3 GlobalInvocationID) { uint triangleIndex = GlobalInvocationID.x; doIgnore(); @@ -157,7 +169,7 @@ void main_sort_triangles_inner(uint3 GlobalInvocationID) { float3 center = (((p0 + p1) + p2) / 3.0f); float3 voxelPos = toVoxelPos(center); uint voxelIndex = toIndex1D(uniforms[0].y, voxelPos); - int triangleOffset = atomicAdd_2(LUT, (4u * voxelIndex), 1); + int triangleOffset = tint_atomicAdd_2(LUT, (4u * voxelIndex), 1); } [numthreads(128, 1, 1)] diff --git a/test/tint/bug/tint/1121.wgsl.expected.hlsl b/test/tint/bug/tint/1121.wgsl.expected.hlsl index a5ab12d742..82a2ef64ea 100644 --- a/test/tint/bug/tint/1121.wgsl.expected.hlsl +++ b/test/tint/bug/tint/1121.wgsl.expected.hlsl @@ -1,9 +1,3 @@ -uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) { - uint original_value = 0; - buffer.InterlockedAdd(offset, value, original_value); - return original_value; -} - RWByteAddressBuffer lightsBuffer : register(u0, space0); RWByteAddressBuffer tileLightId : register(u0, space1); @@ -28,6 +22,13 @@ float4x4 tint_symbol_6(uint4 buffer[11], uint offset) { return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4])); } +uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) { + uint original_value = 0; + buffer.InterlockedAdd(offset, value, original_value); + return original_value; +} + + void main_inner(uint3 GlobalInvocationID) { uint index = GlobalInvocationID.x; if ((index >= config[0].x)) { @@ -96,7 +97,7 @@ void main_inner(uint3 GlobalInvocationID) { if ((tint_tmp)) { continue; } - uint offset = atomicAdd_1(tileLightId, (260u * tileId), 1u); + uint offset = tint_atomicAdd(tileLightId, (260u * tileId), 1u); if ((offset >= config[1].x)) { continue; } diff --git a/test/tint/bug/tint/926.wgsl.expected.hlsl b/test/tint/bug/tint/926.wgsl.expected.hlsl index 9f7e997c6f..f5ed4b6d4a 100644 --- a/test/tint/bug/tint/926.wgsl.expected.hlsl +++ b/test/tint/bug/tint/926.wgsl.expected.hlsl @@ -1,9 +1,3 @@ -uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) { - uint original_value = 0; - buffer.InterlockedAdd(offset, value, original_value); - return original_value; -} - RWByteAddressBuffer drawOut : register(u5, space0); static uint cubeVerts = 0u; @@ -11,8 +5,15 @@ struct tint_symbol_1 { uint3 global_id : SV_DispatchThreadID; }; +uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) { + uint original_value = 0; + buffer.InterlockedAdd(offset, value, original_value); + return original_value; +} + + void computeMain_inner(uint3 global_id) { - const uint firstVertex = atomicAdd_1(drawOut, 0u, cubeVerts); + const uint firstVertex = tint_atomicAdd(drawOut, 0u, cubeVerts); } [numthreads(1, 1, 1)] diff --git a/test/tint/bug/tint/993.wgsl.expected.hlsl b/test/tint/bug/tint/993.wgsl.expected.hlsl index d1c5d71291..6ec8aed1a5 100644 --- a/test/tint/bug/tint/993.wgsl.expected.hlsl +++ b/test/tint/bug/tint/993.wgsl.expected.hlsl @@ -1,9 +1,3 @@ -int atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { - int value = 0; - buffer.InterlockedOr(offset, 0, value); - return value; -} - cbuffer cbuffer_constants : register(b0, space1) { uint4 constants[1]; }; @@ -12,8 +6,15 @@ RWByteAddressBuffer result : register(u1, space1); RWByteAddressBuffer s : register(u0, space0); +int tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) { + int value = 0; + buffer.InterlockedOr(offset, 0, value); + return value; +} + + int runTest() { - return atomicLoad_1(s, (4u * (0u + uint(constants[0].x)))); + return tint_atomicLoad(s, (4u * (0u + uint(constants[0].x)))); } [numthreads(1, 1, 1)] diff --git a/test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl index 35e6ef35c5..f31fb5cf7f 100644 --- a/test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedAdd(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicAdd_8a199a() { - uint res = atomicAdd_1(sb_rw, 0u, 1u); + uint res = tint_atomicAdd(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl index f87e491c2c..62ae701abd 100644 --- a/test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicAdd_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedAdd(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicAdd_d32fe4() { - int res = atomicAdd_1(sb_rw, 0u, 1); + int res = tint_atomicAdd(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl index fa6c15f837..281ec13b3d 100644 --- a/test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicAnd_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicAnd(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedAnd(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicAnd_152966() { - int res = atomicAnd_1(sb_rw, 0u, 1); + int res = tint_atomicAnd(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl index 9d4eb2f1a1..1c89d149e9 100644 --- a/test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicAnd_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicAnd(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedAnd(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicAnd_85a8d9() { - uint res = atomicAnd_1(sb_rw, 0u, 1u); + uint res = tint_atomicAnd(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl deleted file mode 100644 index 6aa8f5a9f7..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl +++ /dev/null @@ -1,76 +0,0 @@ -SKIP: FAILED - -#version 310 es -precision mediump float; - -ivec2 tint_atomicCompareExchangeWeak(inout int param_0, int param_1, int param_2) { - ivec2 result; - result.x = atomicCompSwap(param_0, param_1, param_2); - result.y = result.x == param_2 ? 1 : 0; - return result; -} - - -struct SB_RW { - int arg_0; -}; - -layout(binding = 0, std430) buffer SB_RW_1 { - int arg_0; -} sb_rw; -void atomicCompareExchangeWeak_12871c() { - ivec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1, 1); -} - -void fragment_main() { - atomicCompareExchangeWeak_12871c(); -} - -void main() { - fragment_main(); - return; -} -Error parsing GLSL shader: -ERROR: 0:6: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. -ERROR: 0:6: '' : compilation terminated -ERROR: 2 compilation errors. No code generated. - - - -#version 310 es - -ivec2 tint_atomicCompareExchangeWeak(inout int param_0, int param_1, int param_2) { - ivec2 result; - result.x = atomicCompSwap(param_0, param_1, param_2); - result.y = result.x == param_2 ? 1 : 0; - return result; -} - - -struct SB_RW { - int arg_0; -}; - -layout(binding = 0, std430) buffer SB_RW_1 { - int arg_0; -} sb_rw; -void atomicCompareExchangeWeak_12871c() { - ivec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1, 1); -} - -void compute_main() { - atomicCompareExchangeWeak_12871c(); -} - -layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; -void main() { - compute_main(); - return; -} -Error parsing GLSL shader: -ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. -ERROR: 0:5: '' : compilation terminated -ERROR: 2 compilation errors. No code generated. - - - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl deleted file mode 100644 index 9bd884cb28..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl +++ /dev/null @@ -1,23 +0,0 @@ -int2 atomicCompareExchangeWeak_1(RWByteAddressBuffer buffer, uint offset, int compare, int value) { - int2 result = {0, 0}; - buffer.InterlockedCompareExchange(offset, compare, value, result.x); - result.y = result.x == compare; - return result; -} - -RWByteAddressBuffer sb_rw : register(u0, space0); - -void atomicCompareExchangeWeak_12871c() { - int2 res = atomicCompareExchangeWeak_1(sb_rw, 0u, 1, 1); -} - -void fragment_main() { - atomicCompareExchangeWeak_12871c(); - return; -} - -[numthreads(1, 1, 1)] -void compute_main() { - atomicCompareExchangeWeak_12871c(); - return; -} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl deleted file mode 100644 index a7bb20c970..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl +++ /dev/null @@ -1,29 +0,0 @@ -#include - -using namespace metal; - -template -vec atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) { - T prev_value = compare; - bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed); - return {prev_value, matched}; -} - -struct SB_RW { - /* 0x0000 */ atomic_int arg_0; -}; - -void atomicCompareExchangeWeak_12871c(device SB_RW* const tint_symbol) { - int2 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1, 1); -} - -fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { - atomicCompareExchangeWeak_12871c(tint_symbol_1); - return; -} - -kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) { - atomicCompareExchangeWeak_12871c(tint_symbol_2); - return; -} - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl similarity index 84% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl index 1f1c2d1e47..f3c62f75cb 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl @@ -27,17 +27,17 @@ struct SB_RW { }; @group(0) @binding(0) var sb_rw : SB_RW; -// fn atomicCompareExchangeWeak(ptr, read_write>, i32, i32) -> vec2 -fn atomicCompareExchangeWeak_12871c() { - var res: vec2 = atomicCompareExchangeWeak(&sb_rw.arg_0, 1, 1); +// fn atomicCompareExchangeWeak(ptr, read_write>, i32, i32) -> __atomic_compare_exchange_result +fn atomicCompareExchangeWeak_1bd40a() { + var res = atomicCompareExchangeWeak(&sb_rw.arg_0, 1, 1); } @stage(fragment) fn fragment_main() { - atomicCompareExchangeWeak_12871c(); + atomicCompareExchangeWeak_1bd40a(); } @stage(compute) @workgroup_size(1) fn compute_main() { - atomicCompareExchangeWeak_12871c(); + atomicCompareExchangeWeak_1bd40a(); } diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl new file mode 100644 index 0000000000..001110fd65 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl @@ -0,0 +1,62 @@ +#version 310 es +precision mediump float; + +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; + + +struct SB_RW { + int arg_0; +}; + +layout(binding = 0, std430) buffer SB_RW_1 { + int arg_0; +} sb_rw; +void atomicCompareExchangeWeak_1bd40a() { + atomic_compare_exchange_resulti32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1, 1); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; + atomic_compare_exchange_resulti32 res = atomic_compare_result; +} + +void fragment_main() { + atomicCompareExchangeWeak_1bd40a(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; + + +struct SB_RW { + int arg_0; +}; + +layout(binding = 0, std430) buffer SB_RW_1 { + int arg_0; +} sb_rw; +void atomicCompareExchangeWeak_1bd40a() { + atomic_compare_exchange_resulti32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1, 1); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; + atomic_compare_exchange_resulti32 res = atomic_compare_result; +} + +void compute_main() { + atomicCompareExchangeWeak_1bd40a(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.hlsl new file mode 100644 index 0000000000..b1e30e0e79 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.hlsl @@ -0,0 +1,29 @@ +RWByteAddressBuffer sb_rw : register(u0, space0); + +struct atomic_compare_exchange_weak_ret_type { + int old_value; + bool exchanged; +}; + +atomic_compare_exchange_weak_ret_type tint_atomicCompareExchangeWeak(RWByteAddressBuffer buffer, uint offset, int compare, int value) { + atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; + buffer.InterlockedCompareExchange(offset, compare, value, result.old_value); + result.exchanged = result.old_value == compare; + return result; +} + + +void atomicCompareExchangeWeak_1bd40a() { + atomic_compare_exchange_weak_ret_type res = tint_atomicCompareExchangeWeak(sb_rw, 0u, 1, 1); +} + +void fragment_main() { + atomicCompareExchangeWeak_1bd40a(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + atomicCompareExchangeWeak_1bd40a(); + return; +} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl new file mode 100644 index 0000000000..6cab275aa3 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; + +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; +template +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) { + T old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +struct SB_RW { + /* 0x0000 */ atomic_int arg_0; +}; + +void atomicCompareExchangeWeak_1bd40a(device SB_RW* const tint_symbol) { + atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1, 1); +} + +fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { + atomicCompareExchangeWeak_1bd40a(tint_symbol_1); + return; +} + +kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) { + atomicCompareExchangeWeak_1bd40a(tint_symbol_2); + return; +} + diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm similarity index 56% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm rename to test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm index f4b46af60c..454abb1a57 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 32 +; Bound: 30 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -12,7 +12,10 @@ OpName %SB_RW "SB_RW" OpMemberName %SB_RW 0 "arg_0" OpName %sb_rw "sb_rw" - OpName %atomicCompareExchangeWeak_12871c "atomicCompareExchangeWeak_12871c" + OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a" + OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" + OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" + OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpName %res "res" OpName %fragment_main "fragment_main" OpName %compute_main "compute_main" @@ -20,40 +23,40 @@ OpMemberDecorate %SB_RW 0 Offset 0 OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw Binding 0 + OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 %int = OpTypeInt 32 1 %SB_RW = OpTypeStruct %int %_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer %void = OpTypeVoid %5 = OpTypeFunction %void - %v2int = OpTypeVector %int 2 + %bool = OpTypeBool +%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %uint = OpTypeInt 32 0 %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %int_1 = OpConstant %int 1 - %bool = OpTypeBool - %int_0 = OpConstant %int 0 -%_ptr_Function_v2int = OpTypePointer Function %v2int - %25 = OpConstantNull %v2int -%atomicCompareExchangeWeak_12871c = OpFunction %void None %5 +%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 + %23 = OpConstantNull %__atomic_compare_exchange_resulti32 +%atomicCompareExchangeWeak_1bd40a = OpFunction %void None %5 %8 = OpLabel - %res = OpVariable %_ptr_Function_v2int Function %25 - %16 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 - %19 = OpAtomicCompareExchange %int %16 %uint_1 %uint_0 %uint_0 %int_1 %int_1 + %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %23 + %17 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 + %19 = OpAtomicCompareExchange %int %17 %uint_1 %uint_0 %uint_0 %int_1 %int_1 %20 = OpIEqual %bool %19 %int_1 - %22 = OpSelect %int %20 %int_1 %int_0 - %9 = OpCompositeConstruct %v2int %19 %22 + %9 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %19 %20 OpStore %res %9 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %5 - %27 = OpLabel - %28 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c + %25 = OpLabel + %26 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a OpReturn OpFunctionEnd %compute_main = OpFunction %void None %5 - %30 = OpLabel - %31 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c + %28 = OpLabel + %29 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl similarity index 52% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl index ba2ab03c7e..7e22177963 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl @@ -4,16 +4,16 @@ struct SB_RW { @group(0) @binding(0) var sb_rw : SB_RW; -fn atomicCompareExchangeWeak_12871c() { - var res : vec2 = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1); +fn atomicCompareExchangeWeak_1bd40a() { + var res = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1); } @stage(fragment) fn fragment_main() { - atomicCompareExchangeWeak_12871c(); + atomicCompareExchangeWeak_1bd40a(); } @stage(compute) @workgroup_size(1) fn compute_main() { - atomicCompareExchangeWeak_12871c(); + atomicCompareExchangeWeak_1bd40a(); } diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl similarity index 84% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl index d4473253e9..2fde0d8537 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl @@ -27,17 +27,17 @@ struct SB_RW { }; @group(0) @binding(0) var sb_rw : SB_RW; -// fn atomicCompareExchangeWeak(ptr, read_write>, u32, u32) -> vec2 -fn atomicCompareExchangeWeak_6673da() { - var res: vec2 = atomicCompareExchangeWeak(&sb_rw.arg_0, 1u, 1u); +// fn atomicCompareExchangeWeak(ptr, read_write>, u32, u32) -> __atomic_compare_exchange_result +fn atomicCompareExchangeWeak_63d8e6() { + var res = atomicCompareExchangeWeak(&sb_rw.arg_0, 1u, 1u); } @stage(fragment) fn fragment_main() { - atomicCompareExchangeWeak_6673da(); + atomicCompareExchangeWeak_63d8e6(); } @stage(compute) @workgroup_size(1) fn compute_main() { - atomicCompareExchangeWeak_6673da(); + atomicCompareExchangeWeak_63d8e6(); } diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl new file mode 100644 index 0000000000..e5738b7cb4 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl @@ -0,0 +1,62 @@ +#version 310 es +precision mediump float; + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; + + +struct SB_RW { + uint arg_0; +}; + +layout(binding = 0, std430) buffer SB_RW_1 { + uint arg_0; +} sb_rw; +void atomicCompareExchangeWeak_63d8e6() { + atomic_compare_exchange_resultu32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1u, 1u); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; + atomic_compare_exchange_resultu32 res = atomic_compare_result; +} + +void fragment_main() { + atomicCompareExchangeWeak_63d8e6(); +} + +void main() { + fragment_main(); + return; +} +#version 310 es + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; + + +struct SB_RW { + uint arg_0; +}; + +layout(binding = 0, std430) buffer SB_RW_1 { + uint arg_0; +} sb_rw; +void atomicCompareExchangeWeak_63d8e6() { + atomic_compare_exchange_resultu32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1u, 1u); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; + atomic_compare_exchange_resultu32 res = atomic_compare_result; +} + +void compute_main() { + atomicCompareExchangeWeak_63d8e6(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(); + return; +} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.hlsl new file mode 100644 index 0000000000..9dc72b3c46 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.hlsl @@ -0,0 +1,29 @@ +RWByteAddressBuffer sb_rw : register(u0, space0); + +struct atomic_compare_exchange_weak_ret_type { + uint old_value; + bool exchanged; +}; + +atomic_compare_exchange_weak_ret_type tint_atomicCompareExchangeWeak(RWByteAddressBuffer buffer, uint offset, uint compare, uint value) { + atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0; + buffer.InterlockedCompareExchange(offset, compare, value, result.old_value); + result.exchanged = result.old_value == compare; + return result; +} + + +void atomicCompareExchangeWeak_63d8e6() { + atomic_compare_exchange_weak_ret_type res = tint_atomicCompareExchangeWeak(sb_rw, 0u, 1u, 1u); +} + +void fragment_main() { + atomicCompareExchangeWeak_63d8e6(); + return; +} + +[numthreads(1, 1, 1)] +void compute_main() { + atomicCompareExchangeWeak_63d8e6(); + return; +} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl new file mode 100644 index 0000000000..94166ca79b --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; +template +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) { + T old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +struct SB_RW { + /* 0x0000 */ atomic_uint arg_0; +}; + +void atomicCompareExchangeWeak_63d8e6(device SB_RW* const tint_symbol) { + atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1u, 1u); +} + +fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { + atomicCompareExchangeWeak_63d8e6(tint_symbol_1); + return; +} + +kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) { + atomicCompareExchangeWeak_63d8e6(tint_symbol_2); + return; +} + diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm similarity index 55% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm rename to test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm index 7edaed7b1f..09e83a0ae1 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 29 +; Bound: 28 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -12,7 +12,10 @@ OpName %SB_RW "SB_RW" OpMemberName %SB_RW 0 "arg_0" OpName %sb_rw "sb_rw" - OpName %atomicCompareExchangeWeak_6673da "atomicCompareExchangeWeak_6673da" + OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6" + OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" + OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" + OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpName %res "res" OpName %fragment_main "fragment_main" OpName %compute_main "compute_main" @@ -20,37 +23,38 @@ OpMemberDecorate %SB_RW 0 Offset 0 OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw Binding 0 + OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 %uint = OpTypeInt 32 0 %SB_RW = OpTypeStruct %uint %_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer %void = OpTypeVoid %5 = OpTypeFunction %void - %v2uint = OpTypeVector %uint 2 + %bool = OpTypeBool +%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint - %bool = OpTypeBool -%_ptr_Function_v2uint = OpTypePointer Function %v2uint - %22 = OpConstantNull %v2uint -%atomicCompareExchangeWeak_6673da = OpFunction %void None %5 +%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 + %21 = OpConstantNull %__atomic_compare_exchange_resultu32 +%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %5 %8 = OpLabel - %res = OpVariable %_ptr_Function_v2uint Function %22 - %15 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 - %17 = OpAtomicCompareExchange %uint %15 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1 + %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %21 + %16 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 + %17 = OpAtomicCompareExchange %uint %16 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1 %18 = OpIEqual %bool %17 %uint_1 - %19 = OpSelect %uint %18 %uint_1 %uint_0 - %9 = OpCompositeConstruct %v2uint %17 %19 + %9 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18 OpStore %res %9 OpReturn OpFunctionEnd %fragment_main = OpFunction %void None %5 - %24 = OpLabel - %25 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da + %23 = OpLabel + %24 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %5 - %27 = OpLabel - %28 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da + %26 = OpLabel + %27 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl similarity index 51% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl index dff8dcad30..3ecac33dbf 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl @@ -4,16 +4,16 @@ struct SB_RW { @group(0) @binding(0) var sb_rw : SB_RW; -fn atomicCompareExchangeWeak_6673da() { - var res : vec2 = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u); +fn atomicCompareExchangeWeak_63d8e6() { + var res = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u); } @stage(fragment) fn fragment_main() { - atomicCompareExchangeWeak_6673da(); + atomicCompareExchangeWeak_63d8e6(); } @stage(compute) @workgroup_size(1) fn compute_main() { - atomicCompareExchangeWeak_6673da(); + atomicCompareExchangeWeak_63d8e6(); } diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl deleted file mode 100644 index 65d12ed736..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl +++ /dev/null @@ -1,76 +0,0 @@ -SKIP: FAILED - -#version 310 es -precision mediump float; - -uvec2 tint_atomicCompareExchangeWeak(inout uint param_0, uint param_1, uint param_2) { - uvec2 result; - result.x = atomicCompSwap(param_0, param_1, param_2); - result.y = result.x == param_2 ? 1u : 0u; - return result; -} - - -struct SB_RW { - uint arg_0; -}; - -layout(binding = 0, std430) buffer SB_RW_1 { - uint arg_0; -} sb_rw; -void atomicCompareExchangeWeak_6673da() { - uvec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1u, 1u); -} - -void fragment_main() { - atomicCompareExchangeWeak_6673da(); -} - -void main() { - fragment_main(); - return; -} -Error parsing GLSL shader: -ERROR: 0:6: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. -ERROR: 0:6: '' : compilation terminated -ERROR: 2 compilation errors. No code generated. - - - -#version 310 es - -uvec2 tint_atomicCompareExchangeWeak(inout uint param_0, uint param_1, uint param_2) { - uvec2 result; - result.x = atomicCompSwap(param_0, param_1, param_2); - result.y = result.x == param_2 ? 1u : 0u; - return result; -} - - -struct SB_RW { - uint arg_0; -}; - -layout(binding = 0, std430) buffer SB_RW_1 { - uint arg_0; -} sb_rw; -void atomicCompareExchangeWeak_6673da() { - uvec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1u, 1u); -} - -void compute_main() { - atomicCompareExchangeWeak_6673da(); -} - -layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; -void main() { - compute_main(); - return; -} -Error parsing GLSL shader: -ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. -ERROR: 0:5: '' : compilation terminated -ERROR: 2 compilation errors. No code generated. - - - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl deleted file mode 100644 index 430f1323cb..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl +++ /dev/null @@ -1,23 +0,0 @@ -uint2 atomicCompareExchangeWeak_1(RWByteAddressBuffer buffer, uint offset, uint compare, uint value) { - uint2 result = {0, 0}; - buffer.InterlockedCompareExchange(offset, compare, value, result.x); - result.y = result.x == compare; - return result; -} - -RWByteAddressBuffer sb_rw : register(u0, space0); - -void atomicCompareExchangeWeak_6673da() { - uint2 res = atomicCompareExchangeWeak_1(sb_rw, 0u, 1u, 1u); -} - -void fragment_main() { - atomicCompareExchangeWeak_6673da(); - return; -} - -[numthreads(1, 1, 1)] -void compute_main() { - atomicCompareExchangeWeak_6673da(); - return; -} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl deleted file mode 100644 index b3a827f788..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl +++ /dev/null @@ -1,29 +0,0 @@ -#include - -using namespace metal; - -template -vec atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) { - T prev_value = compare; - bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed); - return {prev_value, matched}; -} - -struct SB_RW { - /* 0x0000 */ atomic_uint arg_0; -}; - -void atomicCompareExchangeWeak_6673da(device SB_RW* const tint_symbol) { - uint2 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1u, 1u); -} - -fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { - atomicCompareExchangeWeak_6673da(tint_symbol_1); - return; -} - -kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) { - atomicCompareExchangeWeak_6673da(tint_symbol_2); - return; -} - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl similarity index 85% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl index d75b725e85..599f9be635 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl @@ -24,12 +24,12 @@ var arg_0: atomic; -// fn atomicCompareExchangeWeak(ptr, read_write>, u32, u32) -> vec2 -fn atomicCompareExchangeWeak_b2ab2c() { - var res: vec2 = atomicCompareExchangeWeak(&arg_0, 1u, 1u); +// fn atomicCompareExchangeWeak(ptr, read_write>, u32, u32) -> __atomic_compare_exchange_result +fn atomicCompareExchangeWeak_83580d() { + var res = atomicCompareExchangeWeak(&arg_0, 1u, 1u); } @stage(compute) @workgroup_size(1) fn compute_main() { - atomicCompareExchangeWeak_b2ab2c(); + atomicCompareExchangeWeak_83580d(); } diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl new file mode 100644 index 0000000000..589a98f416 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl @@ -0,0 +1,29 @@ +#version 310 es + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; + + +shared uint arg_0; +void atomicCompareExchangeWeak_83580d() { + atomic_compare_exchange_resultu32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(arg_0, 1u, 1u); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; + atomic_compare_exchange_resultu32 res = atomic_compare_result; +} + +void compute_main(uint local_invocation_index) { + { + atomicExchange(arg_0, 0u); + } + barrier(); + atomicCompareExchangeWeak_83580d(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(gl_LocalInvocationIndex); + return; +} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl similarity index 56% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl index 05cfbc2e74..3f46597640 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl @@ -1,11 +1,15 @@ +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; groupshared uint arg_0; -void atomicCompareExchangeWeak_b2ab2c() { - uint2 atomic_result = uint2(0u, 0u); +void atomicCompareExchangeWeak_83580d() { + atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; uint atomic_compare_value = 1u; - InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.x); - atomic_result.y = atomic_result.x == atomic_compare_value; - uint2 res = atomic_result; + InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value); + atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; + atomic_compare_exchange_resultu32 res = atomic_result; } struct tint_symbol_1 { @@ -18,7 +22,7 @@ void compute_main_inner(uint local_invocation_index) { InterlockedExchange(arg_0, 0u, atomic_result_1); } GroupMemoryBarrierWithGroupSync(); - atomicCompareExchangeWeak_b2ab2c(); + atomicCompareExchangeWeak_83580d(); } [numthreads(1, 1, 1)] diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.msl new file mode 100644 index 0000000000..7a3443b90d --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; + +struct atomic_compare_exchange_resultu32 { + uint old_value; + bool exchanged; +}; +template +atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { + T old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) { + atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u); +} + +void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) { + { + atomic_store_explicit(tint_symbol_1, uint(), memory_order_relaxed); + } + threadgroup_barrier(mem_flags::mem_threadgroup); + atomicCompareExchangeWeak_83580d(tint_symbol_1); +} + +kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) { + threadgroup atomic_uint tint_symbol_2; + compute_main_inner(local_invocation_index, &(tint_symbol_2)); + return; +} + diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm similarity index 54% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm rename to test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm index 6098500d00..826498596a 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 37 +; Bound: 36 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -9,12 +9,17 @@ OpExecutionMode %compute_main LocalSize 1 1 1 OpName %local_invocation_index_1 "local_invocation_index_1" OpName %arg_0 "arg_0" - OpName %atomicCompareExchangeWeak_b2ab2c "atomicCompareExchangeWeak_b2ab2c" + OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d" + OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" + OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" + OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpName %res "res" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index "local_invocation_index" OpName %compute_main "compute_main" OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex + OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input @@ -22,37 +27,36 @@ %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup %void = OpTypeVoid %6 = OpTypeFunction %void - %v2uint = OpTypeVector %uint 2 + %bool = OpTypeBool +%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 - %bool = OpTypeBool -%_ptr_Function_v2uint = OpTypePointer Function %v2uint - %22 = OpConstantNull %v2uint - %23 = OpTypeFunction %void %uint - %29 = OpConstantNull %uint +%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 + %21 = OpConstantNull %__atomic_compare_exchange_resultu32 + %22 = OpTypeFunction %void %uint + %28 = OpConstantNull %uint %uint_264 = OpConstant %uint 264 -%atomicCompareExchangeWeak_b2ab2c = OpFunction %void None %6 +%atomicCompareExchangeWeak_83580d = OpFunction %void None %6 %9 = OpLabel - %res = OpVariable %_ptr_Function_v2uint Function %22 + %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %21 %17 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1 %18 = OpIEqual %bool %17 %uint_1 - %19 = OpSelect %uint %18 %uint_1 %uint_0 - %10 = OpCompositeConstruct %v2uint %17 %19 + %10 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18 OpStore %res %10 OpReturn OpFunctionEnd -%compute_main_inner = OpFunction %void None %23 +%compute_main_inner = OpFunction %void None %22 %local_invocation_index = OpFunctionParameter %uint - %26 = OpLabel - OpAtomicStore %arg_0 %uint_2 %uint_0 %29 + %25 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %28 OpControlBarrier %uint_2 %uint_2 %uint_264 - %32 = OpFunctionCall %void %atomicCompareExchangeWeak_b2ab2c + %31 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d OpReturn OpFunctionEnd %compute_main = OpFunction %void None %6 - %34 = OpLabel - %36 = OpLoad %uint %local_invocation_index_1 - %35 = OpFunctionCall %void %compute_main_inner %36 + %33 = OpLabel + %35 = OpLoad %uint %local_invocation_index_1 + %34 = OpFunctionCall %void %compute_main_inner %35 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.wgsl new file mode 100644 index 0000000000..77a8862a5c --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.wgsl @@ -0,0 +1,10 @@ +var arg_0 : atomic; + +fn atomicCompareExchangeWeak_83580d() { + var res = atomicCompareExchangeWeak(&(arg_0), 1u, 1u); +} + +@stage(compute) @workgroup_size(1) +fn compute_main() { + atomicCompareExchangeWeak_83580d(); +} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl deleted file mode 100644 index 2f649cd762..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl +++ /dev/null @@ -1,37 +0,0 @@ -SKIP: FAILED - -#version 310 es - -ivec2 tint_atomicCompareExchangeWeak(inout int param_0, int param_1, int param_2) { - ivec2 result; - result.x = atomicCompSwap(param_0, param_1, param_2); - result.y = result.x == param_2 ? 1 : 0; - return result; -} - - -shared int arg_0; -void atomicCompareExchangeWeak_89ea3b() { - ivec2 res = tint_atomicCompareExchangeWeak(arg_0, 1, 1); -} - -void compute_main(uint local_invocation_index) { - { - atomicExchange(arg_0, 0); - } - barrier(); - atomicCompareExchangeWeak_89ea3b(); -} - -layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; -void main() { - compute_main(gl_LocalInvocationIndex); - return; -} -Error parsing GLSL shader: -ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. -ERROR: 0:5: '' : compilation terminated -ERROR: 2 compilation errors. No code generated. - - - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl deleted file mode 100644 index b002e501d4..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl +++ /dev/null @@ -1,29 +0,0 @@ -#include - -using namespace metal; - -template -vec atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { - T prev_value = compare; - bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed); - return {prev_value, matched}; -} - -void atomicCompareExchangeWeak_89ea3b(threadgroup atomic_int* const tint_symbol) { - int2 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1); -} - -void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) { - { - atomic_store_explicit(tint_symbol_1, int(), memory_order_relaxed); - } - threadgroup_barrier(mem_flags::mem_threadgroup); - atomicCompareExchangeWeak_89ea3b(tint_symbol_1); -} - -kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) { - threadgroup atomic_int tint_symbol_2; - compute_main_inner(local_invocation_index, &(tint_symbol_2)); - return; -} - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl deleted file mode 100644 index 4357511711..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl +++ /dev/null @@ -1,10 +0,0 @@ -var arg_0 : atomic; - -fn atomicCompareExchangeWeak_89ea3b() { - var res : vec2 = atomicCompareExchangeWeak(&(arg_0), 1, 1); -} - -@stage(compute) @workgroup_size(1) -fn compute_main() { - atomicCompareExchangeWeak_89ea3b(); -} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl deleted file mode 100644 index cbb201a116..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl +++ /dev/null @@ -1,37 +0,0 @@ -SKIP: FAILED - -#version 310 es - -uvec2 tint_atomicCompareExchangeWeak(inout uint param_0, uint param_1, uint param_2) { - uvec2 result; - result.x = atomicCompSwap(param_0, param_1, param_2); - result.y = result.x == param_2 ? 1u : 0u; - return result; -} - - -shared uint arg_0; -void atomicCompareExchangeWeak_b2ab2c() { - uvec2 res = tint_atomicCompareExchangeWeak(arg_0, 1u, 1u); -} - -void compute_main(uint local_invocation_index) { - { - atomicExchange(arg_0, 0u); - } - barrier(); - atomicCompareExchangeWeak_b2ab2c(); -} - -layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; -void main() { - compute_main(gl_LocalInvocationIndex); - return; -} -Error parsing GLSL shader: -ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. -ERROR: 0:5: '' : compilation terminated -ERROR: 2 compilation errors. No code generated. - - - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl deleted file mode 100644 index 6a9485892d..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl +++ /dev/null @@ -1,29 +0,0 @@ -#include - -using namespace metal; - -template -vec atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { - T prev_value = compare; - bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed); - return {prev_value, matched}; -} - -void atomicCompareExchangeWeak_b2ab2c(threadgroup atomic_uint* const tint_symbol) { - uint2 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u); -} - -void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) { - { - atomic_store_explicit(tint_symbol_1, uint(), memory_order_relaxed); - } - threadgroup_barrier(mem_flags::mem_threadgroup); - atomicCompareExchangeWeak_b2ab2c(tint_symbol_1); -} - -kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) { - threadgroup atomic_uint tint_symbol_2; - compute_main_inner(local_invocation_index, &(tint_symbol_2)); - return; -} - diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl deleted file mode 100644 index d53e099785..0000000000 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl +++ /dev/null @@ -1,10 +0,0 @@ -var arg_0 : atomic; - -fn atomicCompareExchangeWeak_b2ab2c() { - var res : vec2 = atomicCompareExchangeWeak(&(arg_0), 1u, 1u); -} - -@stage(compute) @workgroup_size(1) -fn compute_main() { - atomicCompareExchangeWeak_b2ab2c(); -} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl similarity index 85% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl index c81fc384cd..2147f98f2f 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl @@ -24,12 +24,12 @@ var arg_0: atomic; -// fn atomicCompareExchangeWeak(ptr, read_write>, i32, i32) -> vec2 -fn atomicCompareExchangeWeak_89ea3b() { - var res: vec2 = atomicCompareExchangeWeak(&arg_0, 1, 1); +// fn atomicCompareExchangeWeak(ptr, read_write>, i32, i32) -> __atomic_compare_exchange_result +fn atomicCompareExchangeWeak_e88938() { + var res = atomicCompareExchangeWeak(&arg_0, 1, 1); } @stage(compute) @workgroup_size(1) fn compute_main() { - atomicCompareExchangeWeak_89ea3b(); + atomicCompareExchangeWeak_e88938(); } diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl new file mode 100644 index 0000000000..ff5e7a1b5e --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl @@ -0,0 +1,29 @@ +#version 310 es + +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; + + +shared int arg_0; +void atomicCompareExchangeWeak_e88938() { + atomic_compare_exchange_resulti32 atomic_compare_result; + atomic_compare_result.old_value = atomicCompSwap(arg_0, 1, 1); + atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; + atomic_compare_exchange_resulti32 res = atomic_compare_result; +} + +void compute_main(uint local_invocation_index) { + { + atomicExchange(arg_0, 0); + } + barrier(); + atomicCompareExchangeWeak_e88938(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + compute_main(gl_LocalInvocationIndex); + return; +} diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl similarity index 56% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl rename to test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl index 97cc6c4c1e..4d201c49d5 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl @@ -1,11 +1,15 @@ +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; groupshared int arg_0; -void atomicCompareExchangeWeak_89ea3b() { - int2 atomic_result = int2(0, 0); +void atomicCompareExchangeWeak_e88938() { + atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; int atomic_compare_value = 1; - InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.x); - atomic_result.y = atomic_result.x == atomic_compare_value; - int2 res = atomic_result; + InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value); + atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; + atomic_compare_exchange_resulti32 res = atomic_result; } struct tint_symbol_1 { @@ -18,7 +22,7 @@ void compute_main_inner(uint local_invocation_index) { InterlockedExchange(arg_0, 0, atomic_result_1); } GroupMemoryBarrierWithGroupSync(); - atomicCompareExchangeWeak_89ea3b(); + atomicCompareExchangeWeak_e88938(); } [numthreads(1, 1, 1)] diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.msl new file mode 100644 index 0000000000..9815b6d0d6 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; + +struct atomic_compare_exchange_resulti32 { + int old_value; + bool exchanged; +}; +template +atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) { + T old_value = compare; + bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); + return {old_value, exchanged}; +} + +void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol) { + atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1); +} + +void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) { + { + atomic_store_explicit(tint_symbol_1, int(), memory_order_relaxed); + } + threadgroup_barrier(mem_flags::mem_threadgroup); + atomicCompareExchangeWeak_e88938(tint_symbol_1); +} + +kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) { + threadgroup atomic_int tint_symbol_2; + compute_main_inner(local_invocation_index, &(tint_symbol_2)); + return; +} + diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm similarity index 54% rename from test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm rename to test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm index 40f727b99d..a0f338c664 100644 --- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 39 +; Bound: 37 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -9,12 +9,17 @@ OpExecutionMode %compute_main LocalSize 1 1 1 OpName %local_invocation_index_1 "local_invocation_index_1" OpName %arg_0 "arg_0" - OpName %atomicCompareExchangeWeak_89ea3b "atomicCompareExchangeWeak_89ea3b" + OpName %atomicCompareExchangeWeak_e88938 "atomicCompareExchangeWeak_e88938" + OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" + OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" + OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpName %res "res" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index "local_invocation_index" OpName %compute_main "compute_main" OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex + OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input @@ -23,38 +28,36 @@ %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup %void = OpTypeVoid %7 = OpTypeFunction %void - %v2int = OpTypeVector %int 2 + %bool = OpTypeBool +%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 - %bool = OpTypeBool - %int_0 = OpConstant %int 0 -%_ptr_Function_v2int = OpTypePointer Function %v2int - %24 = OpConstantNull %v2int - %25 = OpTypeFunction %void %uint - %31 = OpConstantNull %int +%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 + %22 = OpConstantNull %__atomic_compare_exchange_resulti32 + %23 = OpTypeFunction %void %uint + %29 = OpConstantNull %int %uint_264 = OpConstant %uint 264 -%atomicCompareExchangeWeak_89ea3b = OpFunction %void None %7 +%atomicCompareExchangeWeak_e88938 = OpFunction %void None %7 %10 = OpLabel - %res = OpVariable %_ptr_Function_v2int Function %24 + %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %22 %18 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1 %19 = OpIEqual %bool %18 %int_1 - %21 = OpSelect %int %19 %int_1 %int_0 - %11 = OpCompositeConstruct %v2int %18 %21 + %11 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %18 %19 OpStore %res %11 OpReturn OpFunctionEnd -%compute_main_inner = OpFunction %void None %25 +%compute_main_inner = OpFunction %void None %23 %local_invocation_index = OpFunctionParameter %uint - %28 = OpLabel - OpAtomicStore %arg_0 %uint_2 %uint_0 %31 + %26 = OpLabel + OpAtomicStore %arg_0 %uint_2 %uint_0 %29 OpControlBarrier %uint_2 %uint_2 %uint_264 - %34 = OpFunctionCall %void %atomicCompareExchangeWeak_89ea3b + %32 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938 OpReturn OpFunctionEnd %compute_main = OpFunction %void None %7 - %36 = OpLabel - %38 = OpLoad %uint %local_invocation_index_1 - %37 = OpFunctionCall %void %compute_main_inner %38 + %34 = OpLabel + %36 = OpLoad %uint %local_invocation_index_1 + %35 = OpFunctionCall %void %compute_main_inner %36 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.wgsl new file mode 100644 index 0000000000..e882b01620 --- /dev/null +++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.wgsl @@ -0,0 +1,10 @@ +var arg_0 : atomic; + +fn atomicCompareExchangeWeak_e88938() { + var res = atomicCompareExchangeWeak(&(arg_0), 1, 1); +} + +@stage(compute) @workgroup_size(1) +fn compute_main() { + atomicCompareExchangeWeak_e88938(); +} diff --git a/test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl index 506c6fee25..b40e1463d9 100644 --- a/test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicExchange_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicExchange(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedExchange(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicExchange_d59712() { - uint res = atomicExchange_1(sb_rw, 0u, 1u); + uint res = tint_atomicExchange(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl index 1ab0d6cf4d..ea1abc5d03 100644 --- a/test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicExchange_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicExchange(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedExchange(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicExchange_f2e22f() { - int res = atomicExchange_1(sb_rw, 0u, 1); + int res = tint_atomicExchange(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl index 8c58350d1c..a278b2c16c 100644 --- a/test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) { int value = 0; buffer.InterlockedOr(offset, 0, value); return value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicLoad_0806ad() { - int res = atomicLoad_1(sb_rw, 0u); + int res = tint_atomicLoad(sb_rw, 0u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl index 967d4c0cad..77edab3db2 100644 --- a/test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) { uint value = 0; buffer.InterlockedOr(offset, 0, value); return value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicLoad_fe6cc3() { - uint res = atomicLoad_1(sb_rw, 0u); + uint res = tint_atomicLoad(sb_rw, 0u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl index 621944f603..c8f0893154 100644 --- a/test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicMax_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicMax(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedMax(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicMax_51b9be() { - uint res = atomicMax_1(sb_rw, 0u, 1u); + uint res = tint_atomicMax(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl index 4ac6cd8f3c..f3d398cc33 100644 --- a/test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicMax_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicMax(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedMax(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicMax_92aa72() { - int res = atomicMax_1(sb_rw, 0u, 1); + int res = tint_atomicMax(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl index 5c550171a3..b8d48b8247 100644 --- a/test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicMin_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicMin(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedMin(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicMin_8e38dc() { - int res = atomicMin_1(sb_rw, 0u, 1); + int res = tint_atomicMin(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl index 4b350c308f..3ae5176288 100644 --- a/test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicMin_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicMin(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedMin(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicMin_c67a74() { - uint res = atomicMin_1(sb_rw, 0u, 1u); + uint res = tint_atomicMin(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl index 21059e60ee..97883cbd8c 100644 --- a/test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicOr_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicOr(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedOr(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicOr_5e95d4() { - uint res = atomicOr_1(sb_rw, 0u, 1u); + uint res = tint_atomicOr(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl index 1504decac9..04be9d7304 100644 --- a/test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicOr_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicOr(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedOr(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicOr_8d96a0() { - int res = atomicOr_1(sb_rw, 0u, 1); + int res = tint_atomicOr(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl index d78cdbec5e..cb8cfddaab 100644 --- a/test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl @@ -1,12 +1,13 @@ -void atomicStore_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +void tint_atomicStore(RWByteAddressBuffer buffer, uint offset, uint value) { uint ignored; buffer.InterlockedExchange(offset, value, ignored); } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicStore_cdc29e() { - atomicStore_1(sb_rw, 0u, 1u); + tint_atomicStore(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl index afac632ba9..599f575c7f 100644 --- a/test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl @@ -1,12 +1,13 @@ -void atomicStore_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +void tint_atomicStore(RWByteAddressBuffer buffer, uint offset, int value) { int ignored; buffer.InterlockedExchange(offset, value, ignored); } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicStore_d1e9a6() { - atomicStore_1(sb_rw, 0u, 1); + tint_atomicStore(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl index 26ad745c6d..29d8e0413c 100644 --- a/test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicSub_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicSub(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedAdd(offset, -value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicSub_051100() { - int res = atomicSub_1(sb_rw, 0u, 1); + int res = tint_atomicSub(sb_rw, 0u, 1); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl index 21e149f65d..e5d0027b60 100644 --- a/test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicSub_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicSub(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedAdd(offset, -value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicSub_15bfc9() { - uint res = atomicSub_1(sb_rw, 0u, 1u); + uint res = tint_atomicSub(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl index 9cc11db1f3..9ed582e64e 100644 --- a/test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -uint atomicXor_1(RWByteAddressBuffer buffer, uint offset, uint value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +uint tint_atomicXor(RWByteAddressBuffer buffer, uint offset, uint value) { uint original_value = 0; buffer.InterlockedXor(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicXor_54510e() { - uint res = atomicXor_1(sb_rw, 0u, 1u); + uint res = tint_atomicXor(sb_rw, 0u, 1u); } void fragment_main() { diff --git a/test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl index 7d483b70f8..bb38f4211e 100644 --- a/test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl +++ b/test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl @@ -1,13 +1,14 @@ -int atomicXor_1(RWByteAddressBuffer buffer, uint offset, int value) { +RWByteAddressBuffer sb_rw : register(u0, space0); + +int tint_atomicXor(RWByteAddressBuffer buffer, uint offset, int value) { int original_value = 0; buffer.InterlockedXor(offset, value, original_value); return original_value; } -RWByteAddressBuffer sb_rw : register(u0, space0); void atomicXor_c1b78c() { - int res = atomicXor_1(sb_rw, 0u, 1); + int res = tint_atomicXor(sb_rw, 0u, 1); } void fragment_main() {