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 <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index bc4bf28..a6792b1 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -118,6 +118,8 @@
type __frexp_result
[[display("__frexp_result_vec{N}")]] type __frexp_result_vec<N: num>
+type __atomic_compare_exchange_result<T>
+
////////////////////////////////////////////////////////////////////////////////
// Type matchers //
// //
@@ -603,7 +605,7 @@
[[stage("fragment", "compute")]] fn atomicOr<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
[[stage("fragment", "compute")]] fn atomicXor<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
[[stage("fragment", "compute")]] fn atomicExchange<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
-[[stage("fragment", "compute")]] fn atomicCompareExchangeWeak<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> vec2<T>
+[[stage("fragment", "compute")]] fn atomicCompareExchangeWeak<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T>
////////////////////////////////////////////////////////////////////////////////
// Type constructors //
diff --git a/src/tint/resolver/intrinsic_table.cc b/src/tint/resolver/intrinsic_table.cc
index 6c19f7c..8571176 100644
--- a/src/tint/resolver/intrinsic_table.cc
+++ b/src/tint/resolver/intrinsic_table.cc
@@ -722,6 +722,14 @@
return true;
}
+bool match_atomic_compare_exchange_result(const sem::Type* ty, const sem::Type*& T) {
+ if (ty->Is<Any>()) {
+ T = ty;
+ return true;
+ }
+ return false;
+}
+
struct NameAndType {
std::string name;
sem::Type* type;
@@ -779,6 +787,13 @@
{{"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<sem::Type*>(ty)},
+ {"exchanged", state.builder.create<sem::Bool>()}});
+}
+
/// 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 c518e58..7a422fb 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -1512,8 +1512,41 @@
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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
}
/// 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 @@
ModfResultVec ModfResultVec_;
FrexpResult FrexpResult_;
FrexpResultVec FrexpResultVec_;
+ AtomicCompareExchangeResult AtomicCompareExchangeResult_;
Fiu32 Fiu32_;
Fi32 Fi32_;
Iu32 Iu32_;
@@ -2233,7 +2267,7 @@
~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 @@
/* [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 @@
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [98] */
/* usage */ ParameterUsage::kSampler,
- /* matcher indices */ &kMatcherIndices[191],
+ /* matcher indices */ &kMatcherIndices[193],
},
{
/* [99] */
@@ -3038,12 +3075,12 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [232] */
/* usage */ ParameterUsage::kSampler,
- /* matcher indices */ &kMatcherIndices[191],
+ /* matcher indices */ &kMatcherIndices[193],
},
{
/* [233] */
@@ -3713,7 +3750,7 @@
{
/* [237] */
/* usage */ ParameterUsage::kSampler,
- /* matcher indices */ &kMatcherIndices[191],
+ /* matcher indices */ &kMatcherIndices[193],
},
{
/* [238] */
@@ -3728,12 +3765,12 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [257] */
/* usage */ ParameterUsage::kSampler,
- /* matcher indices */ &kMatcherIndices[191],
+ /* matcher indices */ &kMatcherIndices[193],
},
{
/* [258] */
@@ -3828,12 +3865,12 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [361] */
/* usage */ ParameterUsage::kSampler,
- /* matcher indices */ &kMatcherIndices[191],
+ /* matcher indices */ &kMatcherIndices[193],
},
{
/* [362] */
@@ -4343,12 +4380,12 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [407] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[195],
+ /* matcher indices */ &kMatcherIndices[197],
},
{
/* [408] */
@@ -4583,12 +4620,12 @@
{
/* [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 @@
{
/* [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 @@
{
/* [425] */
/* usage */ ParameterUsage::kSampler,
- /* matcher indices */ &kMatcherIndices[191],
+ /* matcher indices */ &kMatcherIndices[193],
},
{
/* [426] */
@@ -4763,12 +4800,12 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [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 @@
{
/* [579] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[196],
+ /* matcher indices */ &kMatcherIndices[198],
},
{
/* [580] */
@@ -5438,12 +5475,12 @@
{
/* [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 @@
{
/* [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 @@
{
/* [588] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[198],
+ /* matcher indices */ &kMatcherIndices[200],
},
{
/* [589] */
@@ -5498,12 +5535,12 @@
{
/* [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 @@
{
/* [637] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[196],
+ /* matcher indices */ &kMatcherIndices[198],
},
{
/* [638] */
@@ -5733,7 +5770,7 @@
{
/* [641] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[195],
+ /* matcher indices */ &kMatcherIndices[197],
},
{
/* [642] */
@@ -5753,7 +5790,7 @@
{
/* [645] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[194],
+ /* matcher indices */ &kMatcherIndices[196],
},
{
/* [646] */
@@ -5773,7 +5810,7 @@
{
/* [649] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[193],
+ /* matcher indices */ &kMatcherIndices[195],
},
{
/* [650] */
@@ -6163,7 +6200,7 @@
{
/* [727] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[197],
+ /* matcher indices */ &kMatcherIndices[199],
},
{
/* [728] */
@@ -6748,7 +6785,7 @@
{
/* [844] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[198],
+ /* matcher indices */ &kMatcherIndices[200],
},
{
/* [845] */
@@ -6758,7 +6795,7 @@
{
/* [846] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[193],
+ /* matcher indices */ &kMatcherIndices[195],
},
{
/* [847] */
@@ -6768,17 +6805,17 @@
{
/* [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 @@
{
/* [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 @@
{
/* [883] */
/* usage */ ParameterUsage::kTexture,
- /* matcher indices */ &kMatcherIndices[197],
+ /* matcher indices */ &kMatcherIndices[199],
},
{
/* [884] */
@@ -6968,27 +7005,27 @@
{
/* [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 @@
{
/* [919] */
/* usage */ ParameterUsage::kNone,
- /* matcher indices */ &kMatcherIndices[183],
+ /* matcher indices */ &kMatcherIndices[185],
},
{
/* [920] */
@@ -7183,7 +7220,7 @@
{
/* [931] */
/* usage */ ParameterUsage::kNone,
- /* matcher indices */ &kMatcherIndices[173],
+ /* matcher indices */ &kMatcherIndices[175],
},
{
/* [932] */
@@ -7486,7 +7523,7 @@
{
/* [1] */
/* name */ "U",
- /* matcher index */ 57,
+ /* matcher index */ 58,
},
{
/* [2] */
@@ -7496,7 +7533,7 @@
{
/* [3] */
/* name */ "U",
- /* matcher index */ 54,
+ /* matcher index */ 55,
},
{
/* [4] */
@@ -7506,7 +7543,7 @@
{
/* [5] */
/* name */ "U",
- /* matcher index */ 55,
+ /* matcher index */ 56,
},
{
/* [6] */
@@ -7516,12 +7553,12 @@
{
/* [7] */
/* name */ "U",
- /* matcher index */ 56,
+ /* matcher index */ 57,
},
{
/* [8] */
/* name */ "T",
- /* matcher index */ 48,
+ /* matcher index */ 49,
},
{
/* [9] */
@@ -7531,22 +7568,22 @@
{
/* [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 @@
{
/* [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 @@
/* 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 @@
/* 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 @@
/* 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 @@
/* 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 @@
/* 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 @@
/* 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 @@
/* 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 @@
/* 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 @@
/* 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 @@
/* 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 @@
},
{
/* [106] */
- /* fn atomicCompareExchangeWeak<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> vec2<T> */
+ /* fn atomicCompareExchangeWeak<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T> */
/* 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 775cc05..a90a6e2 100644
--- a/src/tint/transform/decompose_memory_access.cc
+++ b/src/tint/transform/decompose_memory_access.cc
@@ -644,14 +644,34 @@
<< el_ty->TypeInfo().name;
}
- auto* ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType());
- auto* func =
- b.create<ast::Function>(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<sem::Struct>();
+ 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<ast::Function>(
+ 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 @@
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 7a7b783..76cb23e 100644
--- a/src/tint/transform/decompose_memory_access.h
+++ b/src/tint/transform/decompose_memory_access.h
@@ -89,6 +89,9 @@
/// @return the newly cloned object
const Intrinsic* Clone(CloneContext* ctx) const override;
+ /// @return true if op is atomic
+ bool IsAtomic() const;
+
/// The op of the intrinsic
const Op op;
diff --git a/src/tint/transform/decompose_memory_access_test.cc b/src/tint/transform/decompose_memory_access_test.cc
index 22b5da4..19f8b2e 100644
--- a/src/tint/transform/decompose_memory_access_test.cc
+++ b/src/tint/transform/decompose_memory_access_test.cc
@@ -2467,95 +2467,105 @@
@group(0) @binding(0) var<storage, read_write> sb : SB;
@internal(intrinsic_atomic_store_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32)
+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<i32>
+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<u32>
+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 @@
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<i32>
+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<u32>
+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<storage, read_write> sb : SB;
diff --git a/src/tint/transform/manager.cc b/src/tint/transform/manager.cc
index 823474c..e5f7682 100644
--- a/src/tint/transform/manager.cc
+++ b/src/tint/transform/manager.cc
@@ -49,7 +49,7 @@
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 ccce06a..8a48156 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -911,39 +911,56 @@
return true;
}
case sem::BuiltinType::kAtomicCompareExchangeWeak: {
- return CallBuiltinHelper(
- out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& 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<sem::Struct>())) {
+ 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<sem::U32>()) {
- 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 72def34..1d566ca 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -174,14 +174,6 @@
/// @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 d6a5fa7..a2cac0f 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -968,7 +968,10 @@
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 @@
<< static_cast<int>(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 @@
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 @@
{
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<int>(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<int>(intrinsic->op);
+ return false;
}
bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
@@ -1788,6 +1759,12 @@
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<sem::Struct>())) {
+ return false;
+ }
+
auto* dest = expr->args[0];
auto* compare_value = expr->args[1];
auto* value = expr->args[2];
@@ -1807,7 +1784,7 @@
pre << ";";
}
- { // InterlockedCompareExchange(dst, compare, value, result.x);
+ { // InterlockedCompareExchange(dst, compare, value, result.old_value);
auto pre = line();
pre << "InterlockedCompareExchange";
{
@@ -1819,14 +1796,13 @@
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::EmitFunction(const ast::Function* func) {
auto* sem = builder_.Sem().Get(func);
+ // Emit storage atomic helpers
+ if (auto* intrinsic =
+ ast::GetAttribute<transform::DecomposeMemoryAccess::Intrinsic>(func->attributes)) {
+ if (intrinsic->storage_class == ast::StorageClass::kStorage && intrinsic->IsAtomic()) {
+ if (!EmitStorageAtomicIntrinsic(func, intrinsic)) {
+ return false;
+ }
+ }
+ return true;
+ }
+
if (ast::HasAttribute<ast::InternalAttribute>(func->attributes)) {
// An internal function. Do not emit.
return true;
@@ -3755,13 +3742,9 @@
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<ast::LocationAttribute>()) {
@@ -3826,7 +3809,6 @@
}
line(b) << "};";
-
return true;
}
diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h
index 86bbd7d..0e8ca4c 100644
--- a/src/tint/writer/hlsl/generator_impl.h
+++ b/src/tint/writer/hlsl/generator_impl.h
@@ -187,6 +187,12 @@
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 @@
TextBuffer helpers_; // Helper functions emitted at the top of the output
std::function<bool()> emit_continuing_;
- std::unordered_map<DMAIntrinsic, std::string, DMAIntrinsic::Hasher> dma_intrinsics_;
std::unordered_map<const sem::Matrix*, std::string> matrix_scalar_ctors_;
std::unordered_map<const sem::Builtin*, std::string> builtins_;
std::unordered_map<const sem::Struct*, std::string> structure_builders_;
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 54d9164..578e78d 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -806,6 +806,12 @@
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<sem::Struct>())) {
+ return false;
+ }
+
auto* ptr_ty = TypeOf(expr->args[0])->UnwrapRef()->As<sem::Pointer>();
auto sc = ptr_ty->StorageClass();
@@ -816,7 +822,8 @@
line(&buf) << "template <typename A, typename T>";
{
auto f = line(&buf);
- f << "vec<T, 2> " << name << "(";
+ auto str_name = StructName(builtin->ReturnType()->As<sem::Struct>());
+ f << str_name << " " << name << "(";
if (!EmitStorageClass(f, sc)) {
return "";
}
@@ -830,12 +837,12 @@
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 e2d1fac..eb8795f 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -3201,42 +3201,12 @@
return false;
}
- // zero := T(0)
- // one := T(1)
- uint32_t zero = 0;
- uint32_t one = 0;
- if (value_sem_type->Is<sem::I32>()) {
- zero = GenerateConstantIfNeeded(ScalarConstant::I32(0u));
- one = GenerateConstantIfNeeded(ScalarConstant::I32(1u));
- } else if (value_sem_type->Is<sem::U32>()) {
- 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<T>(original_value, xchg_success)
+ // result := __atomic_compare_exchange_result<T>(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 59a567a..7340074 100644
--- a/src/tint/writer/spirv/builder_builtin_test.cc
+++ b/src/tint/writer/spirv/builder_builtin_test.cc
@@ -2018,15 +2018,15 @@
TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) {
// struct S {
- // u : atomic<u32>;
- // i : atomic<i32>;
+ // u : atomic<u32>,
+ // i : atomic<i32>,
// }
//
// @binding(1) @group(2) var<storage, read_write> b : S;
//
// fn a_func() {
- // let u : vec2<u32> = atomicCompareExchangeWeak(&b.u, 10u);
- // let i : vec2<i32> = 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<u32>()),
@@ -2040,10 +2040,10 @@
Func("a_func", {}, ty.void_(),
ast::StatementList{
- Decl(Let("u", ty.vec2<u32>(),
+ Decl(Let("u", nullptr,
Call("atomicCompareExchangeWeak", AddressOf(MemberAccessor("b", "u")), 10_u,
20_u))),
- Decl(Let("i", ty.vec2<i32>(),
+ Decl(Let("i", nullptr,
Call("atomicCompareExchangeWeak", AddressOf(MemberAccessor("b", "i")), 10_i,
20_i))),
},
@@ -2062,33 +2062,29 @@
%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 5b6f132..05d4fe0 100644
--- a/test/tint/bug/chromium/1273230.wgsl.expected.hlsl
+++ b/test/tint/bug/chromium/1273230.wgsl.expected.hlsl
@@ -2,24 +2,6 @@
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 @@
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 @@
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 fa8e7e2..e201d4a 100644
--- a/test/tint/bug/tint/1113.wgsl.expected.hlsl
+++ b/test/tint/bug/tint/1113.wgsl.expected.hlsl
@@ -2,35 +2,6 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 a5ab12d..82a2ef6 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 @@
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 @@
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 9f7e997..f5ed4b6 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 @@
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 d1c5d71..6ec8aed 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 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 35e6ef3..f31fb5c 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 f87e491..62ae701 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 fa6c15f..281ec13 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 9d4eb2f..1c89d14 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 6aa8f5a..0000000
--- 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 9bd884c..0000000
--- 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 a7bb20c..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
+++ /dev/null
@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> 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.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm
deleted file mode 100644
index f4b46af..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm
+++ /dev/null
@@ -1,59 +0,0 @@
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 32
-; Schema: 0
- OpCapability Shader
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %fragment_main "fragment_main"
- OpEntryPoint GLCompute %compute_main "compute_main"
- OpExecutionMode %fragment_main OriginUpperLeft
- OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %SB_RW "SB_RW"
- OpMemberName %SB_RW 0 "arg_0"
- OpName %sb_rw "sb_rw"
- OpName %atomicCompareExchangeWeak_12871c "atomicCompareExchangeWeak_12871c"
- OpName %res "res"
- OpName %fragment_main "fragment_main"
- OpName %compute_main "compute_main"
- OpDecorate %SB_RW Block
- OpMemberDecorate %SB_RW 0 Offset 0
- OpDecorate %sb_rw DescriptorSet 0
- OpDecorate %sb_rw Binding 0
- %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
- %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
- %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
- %20 = OpIEqual %bool %19 %int_1
- %22 = OpSelect %int %20 %int_1 %int_0
- %9 = OpCompositeConstruct %v2int %19 %22
- OpStore %res %9
- OpReturn
- OpFunctionEnd
-%fragment_main = OpFunction %void None %5
- %27 = OpLabel
- %28 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c
- OpReturn
- OpFunctionEnd
-%compute_main = OpFunction %void None %5
- %30 = OpLabel
- %31 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl
deleted file mode 100644
index ba2ab03..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl
+++ /dev/null
@@ -1,19 +0,0 @@
-struct SB_RW {
- arg_0 : atomic<i32>,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_12871c() {
- var res : vec2<i32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
-}
-
-@stage(fragment)
-fn fragment_main() {
- atomicCompareExchangeWeak_12871c();
-}
-
-@stage(compute) @workgroup_size(1)
-fn compute_main() {
- atomicCompareExchangeWeak_12871c();
-}
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 1f1c2d1..f3c62f7 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl
@@ -27,17 +27,17 @@
};
@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
-// fn atomicCompareExchangeWeak(ptr<storage, atomic<i32>, read_write>, i32, i32) -> vec2<i32>
-fn atomicCompareExchangeWeak_12871c() {
- var res: vec2<i32> = atomicCompareExchangeWeak(&sb_rw.arg_0, 1, 1);
+// fn atomicCompareExchangeWeak(ptr<storage, atomic<i32>, read_write>, i32, i32) -> __atomic_compare_exchange_result<i32>
+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 0000000..001110f
--- /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 0000000..b1e30e0
--- /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 0000000..6cab275
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resulti32 {
+ int old_value;
+ bool exchanged;
+};
+template <typename A, typename T>
+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/1bd40a.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
new file mode 100644
index 0000000..454abb1
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %SB_RW "SB_RW"
+ OpMemberName %SB_RW 0 "arg_0"
+ OpName %sb_rw "sb_rw"
+ 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"
+ OpDecorate %SB_RW Block
+ 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
+ %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
+%_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___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
+ %9 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %19 %20
+ OpStore %res %9
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+ %25 = OpLabel
+ %26 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %5
+ %28 = OpLabel
+ %29 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl
new file mode 100644
index 0000000..7e22177
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+struct SB_RW {
+ arg_0 : atomic<i32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn atomicCompareExchangeWeak_1bd40a() {
+ var res = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
+}
+
+@stage(fragment)
+fn fragment_main() {
+ atomicCompareExchangeWeak_1bd40a();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+ atomicCompareExchangeWeak_1bd40a();
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl
similarity index 83%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl
index d447325..2fde0d8 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl
@@ -27,17 +27,17 @@
};
@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
-// fn atomicCompareExchangeWeak(ptr<storage, atomic<u32>, read_write>, u32, u32) -> vec2<u32>
-fn atomicCompareExchangeWeak_6673da() {
- var res: vec2<u32> = atomicCompareExchangeWeak(&sb_rw.arg_0, 1u, 1u);
+// fn atomicCompareExchangeWeak(ptr<storage, atomic<u32>, read_write>, u32, u32) -> __atomic_compare_exchange_result<u32>
+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 0000000..e5738b7
--- /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 0000000..9dc72b3
--- /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 0000000..94166ca
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resultu32 {
+ uint old_value;
+ bool exchanged;
+};
+template <typename A, typename T>
+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/63d8e6.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
new file mode 100644
index 0000000..09e83a0
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %fragment_main "fragment_main"
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %fragment_main OriginUpperLeft
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %SB_RW "SB_RW"
+ OpMemberName %SB_RW 0 "arg_0"
+ OpName %sb_rw "sb_rw"
+ 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"
+ OpDecorate %SB_RW Block
+ 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
+ %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
+%_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___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
+ %9 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18
+ OpStore %res %9
+ OpReturn
+ OpFunctionEnd
+%fragment_main = OpFunction %void None %5
+ %23 = OpLabel
+ %24 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %5
+ %26 = OpLabel
+ %27 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl
new file mode 100644
index 0000000..3ecac33
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+struct SB_RW {
+ arg_0 : atomic<u32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
+
+fn atomicCompareExchangeWeak_63d8e6() {
+ var res = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
+}
+
+@stage(fragment)
+fn fragment_main() {
+ atomicCompareExchangeWeak_63d8e6();
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+ 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 65d12ed..0000000
--- 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 430f132..0000000
--- 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 b3a827f..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
+++ /dev/null
@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> 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/6673da.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm
deleted file mode 100644
index 7edaed7..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm
+++ /dev/null
@@ -1,56 +0,0 @@
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 29
-; Schema: 0
- OpCapability Shader
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %fragment_main "fragment_main"
- OpEntryPoint GLCompute %compute_main "compute_main"
- OpExecutionMode %fragment_main OriginUpperLeft
- OpExecutionMode %compute_main LocalSize 1 1 1
- OpName %SB_RW "SB_RW"
- OpMemberName %SB_RW 0 "arg_0"
- OpName %sb_rw "sb_rw"
- OpName %atomicCompareExchangeWeak_6673da "atomicCompareExchangeWeak_6673da"
- OpName %res "res"
- OpName %fragment_main "fragment_main"
- OpName %compute_main "compute_main"
- OpDecorate %SB_RW Block
- OpMemberDecorate %SB_RW 0 Offset 0
- OpDecorate %sb_rw DescriptorSet 0
- OpDecorate %sb_rw Binding 0
- %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
- %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
- %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
- %18 = OpIEqual %bool %17 %uint_1
- %19 = OpSelect %uint %18 %uint_1 %uint_0
- %9 = OpCompositeConstruct %v2uint %17 %19
- OpStore %res %9
- OpReturn
- OpFunctionEnd
-%fragment_main = OpFunction %void None %5
- %24 = OpLabel
- %25 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da
- OpReturn
- OpFunctionEnd
-%compute_main = OpFunction %void None %5
- %27 = OpLabel
- %28 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl
deleted file mode 100644
index dff8dca..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl
+++ /dev/null
@@ -1,19 +0,0 @@
-struct SB_RW {
- arg_0 : atomic<u32>,
-}
-
-@group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_6673da() {
- var res : vec2<u32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
-}
-
-@stage(fragment)
-fn fragment_main() {
- atomicCompareExchangeWeak_6673da();
-}
-
-@stage(compute) @workgroup_size(1)
-fn compute_main() {
- atomicCompareExchangeWeak_6673da();
-}
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 d75b725..599f9be 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl
@@ -24,12 +24,12 @@
var<workgroup> arg_0: atomic<u32>;
-// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<u32>, read_write>, u32, u32) -> vec2<u32>
-fn atomicCompareExchangeWeak_b2ab2c() {
- var res: vec2<u32> = atomicCompareExchangeWeak(&arg_0, 1u, 1u);
+// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<u32>, read_write>, u32, u32) -> __atomic_compare_exchange_result<u32>
+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 0000000..589a98f
--- /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/83580d.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
new file mode 100644
index 0000000..3f46597
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
@@ -0,0 +1,32 @@
+struct atomic_compare_exchange_resultu32 {
+ uint old_value;
+ bool exchanged;
+};
+groupshared uint arg_0;
+
+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.old_value);
+ atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+ atomic_compare_exchange_resultu32 res = atomic_result;
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(arg_0, 0u, atomic_result_1);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ atomicCompareExchangeWeak_83580d();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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 0000000..7a3443b
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resultu32 {
+ uint old_value;
+ bool exchanged;
+};
+template <typename A, typename T>
+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/83580d.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
new file mode 100644
index 0000000..8264985
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ 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
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %bool = OpTypeBool
+%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool
+ %uint_2 = OpConstant %uint 2
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+%_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_83580d = OpFunction %void None %6
+ %9 = OpLabel
+ %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
+ %10 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18
+ OpStore %res %10
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %22
+%local_invocation_index = OpFunctionParameter %uint
+ %25 = OpLabel
+ OpAtomicStore %arg_0 %uint_2 %uint_0 %28
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %31 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %6
+ %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 0000000..77a8862
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : atomic<u32>;
+
+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 2f649cd..0000000
--- 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.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl
deleted file mode 100644
index 97cc6c4..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl
+++ /dev/null
@@ -1,28 +0,0 @@
-groupshared int arg_0;
-
-void atomicCompareExchangeWeak_89ea3b() {
- int2 atomic_result = int2(0, 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;
-}
-
-struct tint_symbol_1 {
- uint local_invocation_index : SV_GroupIndex;
-};
-
-void compute_main_inner(uint local_invocation_index) {
- {
- int atomic_result_1 = 0;
- InterlockedExchange(arg_0, 0, atomic_result_1);
- }
- GroupMemoryBarrierWithGroupSync();
- atomicCompareExchangeWeak_89ea3b();
-}
-
-[numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
- compute_main_inner(tint_symbol.local_invocation_index);
- return;
-}
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 b002e50..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
+++ /dev/null
@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> 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.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm
deleted file mode 100644
index 40f727b..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm
+++ /dev/null
@@ -1,60 +0,0 @@
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 39
-; Schema: 0
- OpCapability Shader
- OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
- 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 %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
- %uint = OpTypeInt 32 0
-%_ptr_Input_uint = OpTypePointer Input %uint
-%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
- %int = OpTypeInt 32 1
-%_ptr_Workgroup_int = OpTypePointer Workgroup %int
- %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
- %void = OpTypeVoid
- %7 = OpTypeFunction %void
- %v2int = OpTypeVector %int 2
- %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
- %uint_264 = OpConstant %uint 264
-%atomicCompareExchangeWeak_89ea3b = OpFunction %void None %7
- %10 = OpLabel
- %res = OpVariable %_ptr_Function_v2int Function %24
- %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
- OpStore %res %11
- OpReturn
- OpFunctionEnd
-%compute_main_inner = OpFunction %void None %25
-%local_invocation_index = OpFunctionParameter %uint
- %28 = OpLabel
- OpAtomicStore %arg_0 %uint_2 %uint_0 %31
- OpControlBarrier %uint_2 %uint_2 %uint_264
- %34 = OpFunctionCall %void %atomicCompareExchangeWeak_89ea3b
- OpReturn
- OpFunctionEnd
-%compute_main = OpFunction %void None %7
- %36 = OpLabel
- %38 = OpLoad %uint %local_invocation_index_1
- %37 = OpFunctionCall %void %compute_main_inner %38
- OpReturn
- OpFunctionEnd
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 4357511..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl
+++ /dev/null
@@ -1,10 +0,0 @@
-var<workgroup> arg_0 : atomic<i32>;
-
-fn atomicCompareExchangeWeak_89ea3b() {
- var res : vec2<i32> = 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 cbb201a..0000000
--- 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.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl
deleted file mode 100644
index 05cfbc2..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl
+++ /dev/null
@@ -1,28 +0,0 @@
-groupshared uint arg_0;
-
-void atomicCompareExchangeWeak_b2ab2c() {
- uint2 atomic_result = uint2(0u, 0u);
- 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;
-}
-
-struct tint_symbol_1 {
- uint local_invocation_index : SV_GroupIndex;
-};
-
-void compute_main_inner(uint local_invocation_index) {
- {
- uint atomic_result_1 = 0u;
- InterlockedExchange(arg_0, 0u, atomic_result_1);
- }
- GroupMemoryBarrierWithGroupSync();
- atomicCompareExchangeWeak_b2ab2c();
-}
-
-[numthreads(1, 1, 1)]
-void compute_main(tint_symbol_1 tint_symbol) {
- compute_main_inner(tint_symbol.local_invocation_index);
- return;
-}
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 6a94858..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
+++ /dev/null
@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> 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.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm
deleted file mode 100644
index 6098500..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm
+++ /dev/null
@@ -1,58 +0,0 @@
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 37
-; Schema: 0
- OpCapability Shader
- OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
- 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 %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
- %uint = OpTypeInt 32 0
-%_ptr_Input_uint = OpTypePointer Input %uint
-%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
-%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
- %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
- %void = OpTypeVoid
- %6 = OpTypeFunction %void
- %v2uint = OpTypeVector %uint 2
- %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
- %uint_264 = OpConstant %uint 264
-%atomicCompareExchangeWeak_b2ab2c = OpFunction %void None %6
- %9 = OpLabel
- %res = OpVariable %_ptr_Function_v2uint Function %22
- %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
- OpStore %res %10
- OpReturn
- OpFunctionEnd
-%compute_main_inner = OpFunction %void None %23
-%local_invocation_index = OpFunctionParameter %uint
- %26 = OpLabel
- OpAtomicStore %arg_0 %uint_2 %uint_0 %29
- OpControlBarrier %uint_2 %uint_2 %uint_264
- %32 = OpFunctionCall %void %atomicCompareExchangeWeak_b2ab2c
- OpReturn
- OpFunctionEnd
-%compute_main = OpFunction %void None %6
- %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/b2ab2c.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl
deleted file mode 100644
index d53e099..0000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl
+++ /dev/null
@@ -1,10 +0,0 @@
-var<workgroup> arg_0 : atomic<u32>;
-
-fn atomicCompareExchangeWeak_b2ab2c() {
- var res : vec2<u32> = 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 c81fc38..2147f98 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl
@@ -24,12 +24,12 @@
var<workgroup> arg_0: atomic<i32>;
-// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<i32>, read_write>, i32, i32) -> vec2<i32>
-fn atomicCompareExchangeWeak_89ea3b() {
- var res: vec2<i32> = atomicCompareExchangeWeak(&arg_0, 1, 1);
+// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<i32>, read_write>, i32, i32) -> __atomic_compare_exchange_result<i32>
+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 0000000..ff5e7a1
--- /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/e88938.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
new file mode 100644
index 0000000..4d201c4
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
@@ -0,0 +1,32 @@
+struct atomic_compare_exchange_resulti32 {
+ int old_value;
+ bool exchanged;
+};
+groupshared int arg_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.old_value);
+ atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+ atomic_compare_exchange_resulti32 res = atomic_result;
+}
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ int atomic_result_1 = 0;
+ InterlockedExchange(arg_0, 0, atomic_result_1);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ atomicCompareExchangeWeak_e88938();
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
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 0000000..9815b6d
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resulti32 {
+ int old_value;
+ bool exchanged;
+};
+template <typename A, typename T>
+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/e88938.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
new file mode 100644
index 0000000..a0f338c
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 37
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %arg_0 "arg_0"
+ 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
+ %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %bool = OpTypeBool
+%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool
+ %uint_2 = OpConstant %uint 2
+ %uint_0 = OpConstant %uint 0
+ %int_1 = OpConstant %int 1
+%_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_e88938 = OpFunction %void None %7
+ %10 = OpLabel
+ %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
+ %11 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %18 %19
+ OpStore %res %11
+ OpReturn
+ OpFunctionEnd
+%compute_main_inner = OpFunction %void None %23
+%local_invocation_index = OpFunctionParameter %uint
+ %26 = OpLabel
+ OpAtomicStore %arg_0 %uint_2 %uint_0 %29
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %32 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %7
+ %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 0000000..e882b01
--- /dev/null
+++ b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : atomic<i32>;
+
+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 506c6fe..b40e146 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 1ab0d6c..ea1abc5 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 8c58350..a278b2c 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 967d4c0..77edab3 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 621944f..c8f0893 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 4ac6cd8..f3d398c 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 5c55017..b8d48b8 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 4b350c3..3ae5176 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 21059e6..97883cb 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 1504dec..04be9d7 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 d78cdbe..cb8cfdd 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 afac632..599f575 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 26ad745..29d8e04 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 21e149f..e5d0027 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 9cc11db..9ed582e 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 7d483b7..bb38f42 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() {
