encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

tint/resolver: Resolve builtin structs 

Allow the resolver to understand builtin structures, like
__frexp_result_f16. This allows backend transforms to declare the types,
even if they're "untypable" by the user.

Bug: chromium:1430309
Change-Id: I392709118182a058f737ccf1b7b46fc6b0b7264d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/129482
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
Ben Clayton 2023-04-26 18:27:53 +00:00 committed by Dawn LUCI CQ
parent d3b09b90e3
commit 47dd30117d
104 changed files with 1970 additions and 1104 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
src/tint/BUILD.gn
View File

@ -1416,6 +1416,7 @@ if (tint_build_unittests) {
"resolver/attribute_validation_test.cc", "resolver/attribute_validation_test.cc",
"resolver/bitcast_validation_test.cc", "resolver/bitcast_validation_test.cc",
"resolver/builtin_enum_test.cc", "resolver/builtin_enum_test.cc",
"resolver/builtin_structs_test.cc",
"resolver/builtin_test.cc", "resolver/builtin_test.cc",
"resolver/builtin_validation_test.cc", "resolver/builtin_validation_test.cc",
"resolver/builtins_validation_test.cc", "resolver/builtins_validation_test.cc",

1
src/tint/CMakeLists.txt
View File

@ -913,6 +913,7 @@ if(TINT_BUILD_TESTS)
resolver/attribute_validation_test.cc resolver/attribute_validation_test.cc
resolver/bitcast_validation_test.cc resolver/bitcast_validation_test.cc
resolver/builtin_enum_test.cc resolver/builtin_enum_test.cc
resolver/builtin_structs_test.cc
resolver/builtin_test.cc resolver/builtin_test.cc
resolver/builtin_validation_test.cc resolver/builtin_validation_test.cc
resolver/builtins_validation_test.cc resolver/builtins_validation_test.cc

2
src/tint/builtin/access_bench.cc
View File

@ -43,7 +43,7 @@ void AccessParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(AccessParser); BENCHMARK(AccessParser);

2
src/tint/builtin/address_space_bench.cc
View File

@ -94,7 +94,7 @@ void AddressSpaceParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(AddressSpaceParser); BENCHMARK(AddressSpaceParser);

2
src/tint/builtin/attribute_bench.cc
View File

@ -143,7 +143,7 @@ void AttributeParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(AttributeParser); BENCHMARK(AttributeParser);

130
src/tint/builtin/builtin.cc
View File

@ -28,6 +28,84 @@ namespace tint::builtin {
/// @param str the string to parse /// @param str the string to parse
/// @returns the parsed enum, or Builtin::kUndefined if the string could not be parsed. /// @returns the parsed enum, or Builtin::kUndefined if the string could not be parsed.
Builtin ParseBuiltin(std::string_view str) { Builtin ParseBuiltin(std::string_view str) {
if (str == "__atomic_compare_exchange_result_i32") {
return Builtin::kAtomicCompareExchangeResultI32;
}
if (str == "__atomic_compare_exchange_result_u32") {
return Builtin::kAtomicCompareExchangeResultU32;
}
if (str == "__frexp_result_abstract") {
return Builtin::kFrexpResultAbstract;
}
if (str == "__frexp_result_f16") {
return Builtin::kFrexpResultF16;
}
if (str == "__frexp_result_f32") {
return Builtin::kFrexpResultF32;
}
if (str == "__frexp_result_vec2_abstract") {
return Builtin::kFrexpResultVec2Abstract;
}
if (str == "__frexp_result_vec2_f16") {
return Builtin::kFrexpResultVec2F16;
}
if (str == "__frexp_result_vec2_f32") {
return Builtin::kFrexpResultVec2F32;
}
if (str == "__frexp_result_vec3_abstract") {
return Builtin::kFrexpResultVec3Abstract;
}
if (str == "__frexp_result_vec3_f16") {
return Builtin::kFrexpResultVec3F16;
}
if (str == "__frexp_result_vec3_f32") {
return Builtin::kFrexpResultVec3F32;
}
if (str == "__frexp_result_vec4_abstract") {
return Builtin::kFrexpResultVec4Abstract;
}
if (str == "__frexp_result_vec4_f16") {
return Builtin::kFrexpResultVec4F16;
}
if (str == "__frexp_result_vec4_f32") {
return Builtin::kFrexpResultVec4F32;
}
if (str == "__modf_result_abstract") {
return Builtin::kModfResultAbstract;
}
if (str == "__modf_result_f16") {
return Builtin::kModfResultF16;
}
if (str == "__modf_result_f32") {
return Builtin::kModfResultF32;
}
if (str == "__modf_result_vec2_abstract") {
return Builtin::kModfResultVec2Abstract;
}
if (str == "__modf_result_vec2_f16") {
return Builtin::kModfResultVec2F16;
}
if (str == "__modf_result_vec2_f32") {
return Builtin::kModfResultVec2F32;
}
if (str == "__modf_result_vec3_abstract") {
return Builtin::kModfResultVec3Abstract;
}
if (str == "__modf_result_vec3_f16") {
return Builtin::kModfResultVec3F16;
}
if (str == "__modf_result_vec3_f32") {
return Builtin::kModfResultVec3F32;
}
if (str == "__modf_result_vec4_abstract") {
return Builtin::kModfResultVec4Abstract;
}
if (str == "__modf_result_vec4_f16") {
return Builtin::kModfResultVec4F16;
}
if (str == "__modf_result_vec4_f32") {
return Builtin::kModfResultVec4F32;
}
if (str == "__packed_vec3") { if (str == "__packed_vec3") {
return Builtin::kPackedVec3; return Builtin::kPackedVec3;
} }
@ -245,6 +323,58 @@ utils::StringStream& operator<<(utils::StringStream& out, Builtin value) {
switch (value) { switch (value) {
case Builtin::kUndefined: case Builtin::kUndefined:
return out << "undefined"; return out << "undefined";
case Builtin::kAtomicCompareExchangeResultI32:
return out << "__atomic_compare_exchange_result_i32";
case Builtin::kAtomicCompareExchangeResultU32:
return out << "__atomic_compare_exchange_result_u32";
case Builtin::kFrexpResultAbstract:
return out << "__frexp_result_abstract";
case Builtin::kFrexpResultF16:
return out << "__frexp_result_f16";
case Builtin::kFrexpResultF32:
return out << "__frexp_result_f32";
case Builtin::kFrexpResultVec2Abstract:
return out << "__frexp_result_vec2_abstract";
case Builtin::kFrexpResultVec2F16:
return out << "__frexp_result_vec2_f16";
case Builtin::kFrexpResultVec2F32:
return out << "__frexp_result_vec2_f32";
case Builtin::kFrexpResultVec3Abstract:
return out << "__frexp_result_vec3_abstract";
case Builtin::kFrexpResultVec3F16:
return out << "__frexp_result_vec3_f16";
case Builtin::kFrexpResultVec3F32:
return out << "__frexp_result_vec3_f32";
case Builtin::kFrexpResultVec4Abstract:
return out << "__frexp_result_vec4_abstract";
case Builtin::kFrexpResultVec4F16:
return out << "__frexp_result_vec4_f16";
case Builtin::kFrexpResultVec4F32:
return out << "__frexp_result_vec4_f32";
case Builtin::kModfResultAbstract:
return out << "__modf_result_abstract";
case Builtin::kModfResultF16:
return out << "__modf_result_f16";
case Builtin::kModfResultF32:
return out << "__modf_result_f32";
case Builtin::kModfResultVec2Abstract:
return out << "__modf_result_vec2_abstract";
case Builtin::kModfResultVec2F16:
return out << "__modf_result_vec2_f16";
case Builtin::kModfResultVec2F32:
return out << "__modf_result_vec2_f32";
case Builtin::kModfResultVec3Abstract:
return out << "__modf_result_vec3_abstract";
case Builtin::kModfResultVec3F16:
return out << "__modf_result_vec3_f16";
case Builtin::kModfResultVec3F32:
return out << "__modf_result_vec3_f32";
case Builtin::kModfResultVec4Abstract:
return out << "__modf_result_vec4_abstract";
case Builtin::kModfResultVec4F16:
return out << "__modf_result_vec4_f16";
case Builtin::kModfResultVec4F32:
return out << "__modf_result_vec4_f32";
case Builtin::kPackedVec3: case Builtin::kPackedVec3:
return out << "__packed_vec3"; return out << "__packed_vec3";
case Builtin::kArray: case Builtin::kArray:

52
src/tint/builtin/builtin.h
View File

@ -30,6 +30,32 @@ namespace tint::builtin {
/// An enumerator of builtin builtin. /// An enumerator of builtin builtin.
enum class Builtin { enum class Builtin {
kUndefined, kUndefined,
kAtomicCompareExchangeResultI32,
kAtomicCompareExchangeResultU32,
kFrexpResultAbstract,
kFrexpResultF16,
kFrexpResultF32,
kFrexpResultVec2Abstract,
kFrexpResultVec2F16,
kFrexpResultVec2F32,
kFrexpResultVec3Abstract,
kFrexpResultVec3F16,
kFrexpResultVec3F32,
kFrexpResultVec4Abstract,
kFrexpResultVec4F16,
kFrexpResultVec4F32,
kModfResultAbstract,
kModfResultF16,
kModfResultF32,
kModfResultVec2Abstract,
kModfResultVec2F16,
kModfResultVec2F32,
kModfResultVec3Abstract,
kModfResultVec3F16,
kModfResultVec3F32,
kModfResultVec4Abstract,
kModfResultVec4F16,
kModfResultVec4F32,
kPackedVec3, kPackedVec3,
kArray, kArray,
kAtomic, kAtomic,
@ -113,6 +139,32 @@ utils::StringStream& operator<<(utils::StringStream& out, Builtin value);
Builtin ParseBuiltin(std::string_view str); Builtin ParseBuiltin(std::string_view str);
constexpr const char* kBuiltinStrings[] = { constexpr const char* kBuiltinStrings[] = {
"__atomic_compare_exchange_result_i32",
"__atomic_compare_exchange_result_u32",
"__frexp_result_abstract",
"__frexp_result_f16",
"__frexp_result_f32",
"__frexp_result_vec2_abstract",
"__frexp_result_vec2_f16",
"__frexp_result_vec2_f32",
"__frexp_result_vec3_abstract",
"__frexp_result_vec3_f16",
"__frexp_result_vec3_f32",
"__frexp_result_vec4_abstract",
"__frexp_result_vec4_f16",
"__frexp_result_vec4_f32",
"__modf_result_abstract",
"__modf_result_f16",
"__modf_result_f32",
"__modf_result_vec2_abstract",
"__modf_result_vec2_f16",
"__modf_result_vec2_f32",
"__modf_result_vec3_abstract",
"__modf_result_vec3_f16",
"__modf_result_vec3_f32",
"__modf_result_vec4_abstract",
"__modf_result_vec4_f16",
"__modf_result_vec4_f32",
"__packed_vec3", "__packed_vec3",
"array", "array",
"atomic", "atomic",

1028
src/tint/builtin/builtin_bench.cc
View File

File diff suppressed because it is too large Load Diff

522
src/tint/builtin/builtin_test.cc
View File

@ -43,6 +43,32 @@ inline std::ostream& operator<<(std::ostream& out, Case c) {
} }
static constexpr Case kValidCases[] = { static constexpr Case kValidCases[] = {
{"__atomic_compare_exchange_result_i32", Builtin::kAtomicCompareExchangeResultI32},
{"__atomic_compare_exchange_result_u32", Builtin::kAtomicCompareExchangeResultU32},
{"__frexp_result_abstract", Builtin::kFrexpResultAbstract},
{"__frexp_result_f16", Builtin::kFrexpResultF16},
{"__frexp_result_f32", Builtin::kFrexpResultF32},
{"__frexp_result_vec2_abstract", Builtin::kFrexpResultVec2Abstract},
{"__frexp_result_vec2_f16", Builtin::kFrexpResultVec2F16},
{"__frexp_result_vec2_f32", Builtin::kFrexpResultVec2F32},
{"__frexp_result_vec3_abstract", Builtin::kFrexpResultVec3Abstract},
{"__frexp_result_vec3_f16", Builtin::kFrexpResultVec3F16},
{"__frexp_result_vec3_f32", Builtin::kFrexpResultVec3F32},
{"__frexp_result_vec4_abstract", Builtin::kFrexpResultVec4Abstract},
{"__frexp_result_vec4_f16", Builtin::kFrexpResultVec4F16},
{"__frexp_result_vec4_f32", Builtin::kFrexpResultVec4F32},
{"__modf_result_abstract", Builtin::kModfResultAbstract},
{"__modf_result_f16", Builtin::kModfResultF16},
{"__modf_result_f32", Builtin::kModfResultF32},
{"__modf_result_vec2_abstract", Builtin::kModfResultVec2Abstract},
{"__modf_result_vec2_f16", Builtin::kModfResultVec2F16},
{"__modf_result_vec2_f32", Builtin::kModfResultVec2F32},
{"__modf_result_vec3_abstract", Builtin::kModfResultVec3Abstract},
{"__modf_result_vec3_f16", Builtin::kModfResultVec3F16},
{"__modf_result_vec3_f32", Builtin::kModfResultVec3F32},
{"__modf_result_vec4_abstract", Builtin::kModfResultVec4Abstract},
{"__modf_result_vec4_f16", Builtin::kModfResultVec4F16},
{"__modf_result_vec4_f32", Builtin::kModfResultVec4F32},
{"__packed_vec3", Builtin::kPackedVec3}, {"__packed_vec3", Builtin::kPackedVec3},
{"array", Builtin::kArray}, {"array", Builtin::kArray},
{"atomic", Builtin::kAtomic}, {"atomic", Builtin::kAtomic},
@ -116,216 +142,294 @@ static constexpr Case kValidCases[] = {
}; };
static constexpr Case kInvalidCases[] = { static constexpr Case kInvalidCases[] = {
{"__acked_veccc", Builtin::kUndefined}, {"__atomic_compareexchangeccresult_i32", Builtin::kUndefined},
{"_pac3ed_v3", Builtin::kUndefined}, {"__atoml3_compare_exchane_resulti2", Builtin::kUndefined},
{"__packeV_vec3", Builtin::kUndefined}, {"__atomic_compare_Vxchange_result_i32", Builtin::kUndefined},
{"arra1", Builtin::kUndefined}, {"__atomic_com1are_exchange_result_u32", Builtin::kUndefined},
{"qqrJy", Builtin::kUndefined}, {"__atomic_qqompare_exchage_resulJ_u32", Builtin::kUndefined},
{"arrll7y", Builtin::kUndefined}, {"__atllmic_compare_exchange_result_u377", Builtin::kUndefined},
{"atppmHHc", Builtin::kUndefined}, {"qpp_frexp_resultHHbstract", Builtin::kUndefined},
{"cto", Builtin::kUndefined}, {"__fep_esulv_abstract", Builtin::kUndefined},
{"abGmi", Builtin::kUndefined}, {"__Gbexp_resul_abstract", Builtin::kUndefined},
{"bovii", Builtin::kUndefined}, {"_vfrexp_resiilt_f16", Builtin::kUndefined},
{"boWWl", Builtin::kUndefined}, {"__fr8xp_resultWWf16", Builtin::kUndefined},
{"Mxxl", Builtin::kUndefined}, {"__frxp_result_fMxx", Builtin::kUndefined},
{"fgg", Builtin::kUndefined}, {"gg_fXexp_reslt_f32", Builtin::kUndefined},
{"X", Builtin::kUndefined}, {"__frXxpresul_V32", Builtin::kUndefined},
{"316", Builtin::kUndefined}, {"__frexp_r3sult_f32", Builtin::kUndefined},
{"fE2", Builtin::kUndefined}, {"__frexpEresult_vec2_abstract", Builtin::kUndefined},
{"fPTT", Builtin::kUndefined}, {"__frex_rPPsult_vTTc2_abstract", Builtin::kUndefined},
{"dxx2", Builtin::kUndefined}, {"__frexp_resuddt_ec2_xxbstract", Builtin::kUndefined},
{"4432", Builtin::kUndefined}, {"__frexp_result_ve442_f16", Builtin::kUndefined},
{"iSVV2", Builtin::kUndefined}, {"_SSfrexp_resulVV_vec2_f16", Builtin::kUndefined},
{"RR2", Builtin::kUndefined}, {"__fRxpRr22sult_vec2_f16", Builtin::kUndefined},
{"at292", Builtin::kUndefined}, {"__frexp_res9lt_vec_fF2", Builtin::kUndefined},
{"mat2x", Builtin::kUndefined}, {"__frexp_result_ve2_f32", Builtin::kUndefined},
{"Vat2OR2", Builtin::kUndefined}, {"_OOfrexp_result_VeHRRf32", Builtin::kUndefined},
{"ma2xyf", Builtin::kUndefined}, {"__frexp_reyult_vec3_absract", Builtin::kUndefined},
{"llnarr2772f", Builtin::kUndefined}, {"__frexp_re77ulll_vecG_arrnstract", Builtin::kUndefined},
{"mat24200", Builtin::kUndefined}, {"__4rexp_result_vec3_00bstract", Builtin::kUndefined},
{"a2ooh", Builtin::kUndefined}, {"__oorxp_result_vec316", Builtin::kUndefined},
{"zz2x2h", Builtin::kUndefined}, {"zz_frexp_esult_ec3_f16", Builtin::kUndefined},
{"miitppx1", Builtin::kUndefined}, {"__iirex11_result_vp3_f16", Builtin::kUndefined},
{"maXX2x3", Builtin::kUndefined}, {"__frXXxp_result_vec3_f32", Builtin::kUndefined},
{"55IIt2nn99", Builtin::kUndefined}, {"__fnnexp99resIIlt_vec3_f355", Builtin::kUndefined},
{"aHHrrt2xSS", Builtin::kUndefined}, {"__faSSerrp_result_vHHc3_fY2", Builtin::kUndefined},
{"makkf", Builtin::kUndefined}, {"__freHp_resutve4_abstkkact", Builtin::kUndefined},
{"jatgRx", Builtin::kUndefined}, {"jfrexpgresult_veRR4_abstrac", Builtin::kUndefined},
{"mb2x3", Builtin::kUndefined}, {"__frexp_resul_vec4_absbrac", Builtin::kUndefined},
{"mat2xjh", Builtin::kUndefined}, {"_jfrexp_result_vec4_f16", Builtin::kUndefined},
{"at2x3h", Builtin::kUndefined}, {"__frexp_resultvec4_f16", Builtin::kUndefined},
{"q2x3h", Builtin::kUndefined}, {"__freqpresultvec4_f16", Builtin::kUndefined},
{"mNN2x4", Builtin::kUndefined}, {"__frexNN_result_vec_f32", Builtin::kUndefined},
{"mavv4", Builtin::kUndefined}, {"__frexp_resvvlt_vc4_f3", Builtin::kUndefined},
{"maQQx4", Builtin::kUndefined}, {"__frexp_esult_vec4_f3QQ", Builtin::kUndefined},
{"maffxr", Builtin::kUndefined}, {"rmodf_reffultabstract", Builtin::kUndefined},
{"mat2xjf", Builtin::kUndefined}, {"__jodf_result_abstract", Builtin::kUndefined},
{"mNNw2x48", Builtin::kUndefined}, {"_mNNwdf_r2sult8abstract", Builtin::kUndefined},
{"matx4h", Builtin::kUndefined}, {"__mdf_result_f16", Builtin::kUndefined},
{"mrrt2x4h", Builtin::kUndefined}, {"__modrr_result_f16", Builtin::kUndefined},
{"Gat2x4h", Builtin::kUndefined}, {"__mGdf_result_f16", Builtin::kUndefined},
{"matFFx2", Builtin::kUndefined}, {"__modf_resulFF_f32", Builtin::kUndefined},
{"mtx", Builtin::kUndefined}, {"__modf_eult_E3", Builtin::kUndefined},
{"mrrt3x", Builtin::kUndefined}, {"__odf_resurrt_f32", Builtin::kUndefined},
{"__modf_reslt_vec_abstract", Builtin::kUndefined},
{"__modfJJresuDt_Xc2_abstract", Builtin::kUndefined},
{"_modf_reslt_vec28abstrct", Builtin::kUndefined},
{"__odf_reult_vkc211f1", Builtin::kUndefined},
{"__mdf_result_vec2_f16", Builtin::kUndefined},
{"__modf_resuJt_vec2_f6", Builtin::kUndefined},
{"__modf_result_vec2cf32", Builtin::kUndefined},
{"__modf_result_vec2_fO2", Builtin::kUndefined},
{"KK_movvf_result_vec2_ftt__", Builtin::kUndefined},
{"xx_modf_r8sult_vec3_abtr5ct", Builtin::kUndefined},
{"__modf_resuFt_vec3_aqt__act", Builtin::kUndefined},
{"__modf_result_vec3_aqqstrac", Builtin::kUndefined},
{"__odf_33esult_vec3_f1O6", Builtin::kUndefined},
{"_ttm6df_resQQlt_ooec9_f16", Builtin::kUndefined},
{"_modf_resu66t_vec3_f16", Builtin::kUndefined},
{"__mdf_resultOvxc3_f36zz", Builtin::kUndefined},
{"__modf_resuyyt_vec3_f32", Builtin::kUndefined},
{"__mod_resul_vecZHHf32", Builtin::kUndefined},
{"__modf_reqult_44ec4WWbstract", Builtin::kUndefined},
{"__mof_result_vec4_abstrOOct", Builtin::kUndefined},
{"__modYooresult_vh4_bstract", Builtin::kUndefined},
{"__modf_relt_ve4_f16", Builtin::kUndefined},
{"__modf_result_ve4Ff16", Builtin::kUndefined},
{"__modf_result_wec4_f1", Builtin::kUndefined},
{"__Kdff_rGsult_vec4_f2", Builtin::kUndefined},
{"__modf_reKKulq_vec4_f32", Builtin::kUndefined},
{"__modf_resummt3vec4_f3F", Builtin::kUndefined},
{"__packed_ec3", Builtin::kUndefined},
{"__packed_ecq", Builtin::kUndefined},
{"_backed_bbec3", Builtin::kUndefined},
{"iira", Builtin::kUndefined},
{"aqOOy", Builtin::kUndefined},
{"arvvTTy", Builtin::kUndefined},
{"atomFFc", Builtin::kUndefined},
{"aoQ00P", Builtin::kUndefined},
{"atPmic", Builtin::kUndefined},
{"bos77", Builtin::kUndefined},
{"CoRbbl", Builtin::kUndefined},
{"booXX", Builtin::kUndefined},
{"qOOO6", Builtin::kUndefined},
{"fs", Builtin::kUndefined},
{"f1X", Builtin::kUndefined},
{"f3", Builtin::kUndefined},
{"q", Builtin::kUndefined},
{"f322", Builtin::kUndefined},
{"0yz2", Builtin::kUndefined},
{"iVP", Builtin::kUndefined},
{"Cnn", Builtin::kUndefined},
{"AtqqHH2", Builtin::kUndefined},
{"at2x2", Builtin::kUndefined},
{"mafKK", Builtin::kUndefined},
{"ltgg2f", Builtin::kUndefined},
{"mat2xf", Builtin::kUndefined},
{"NTTtcx4f", Builtin::kUndefined},
{"ma7ppl2h", Builtin::kUndefined},
{"mNNt2xg", Builtin::kUndefined},
{"uub2XX2h", Builtin::kUndefined},
{"mt2x3", Builtin::kUndefined},
{"m88xK", Builtin::kUndefined},
{"maqx3", Builtin::kUndefined},
{"m11t2x3f", Builtin::kUndefined},
{"22at2iif", Builtin::kUndefined},
{"at2x377", Builtin::kUndefined},
{"m2t2xNh", Builtin::kUndefined},
{"mVVt2x3h", Builtin::kUndefined},
{"FaWW2w11h", Builtin::kUndefined},
{"matww4", Builtin::kUndefined},
{"mat2D4", Builtin::kUndefined},
{"maKx4", Builtin::kUndefined},
{"mat21PPhf", Builtin::kUndefined},
{"mat24f", Builtin::kUndefined},
{"mYYt2x4f", Builtin::kUndefined},
{"mttHH4kk", Builtin::kUndefined},
{"mat2rr4h", Builtin::kUndefined},
{"WWas2x4h", Builtin::kUndefined},
{"maYx2", Builtin::kUndefined},
{"mq3f2", Builtin::kUndefined},
{"vvafu222", Builtin::kUndefined},
{"t3x2f", Builtin::kUndefined}, {"t3x2f", Builtin::kUndefined},
{"Da3xJJf", Builtin::kUndefined}, {"YYat3f", Builtin::kUndefined},
{"ma82", Builtin::kUndefined}, {"may3x2EYY", Builtin::kUndefined},
{"m11k2", Builtin::kUndefined}, {"da3xMoh", Builtin::kUndefined},
{"matx2h", Builtin::kUndefined}, {"matMMx2", Builtin::kUndefined},
{"maJx2h", Builtin::kUndefined}, {"mat3x55h", Builtin::kUndefined},
{"cat3x3", Builtin::kUndefined}, {"maN3", Builtin::kUndefined},
{"mat3O3", Builtin::kUndefined}, {"ma33x3", Builtin::kUndefined},
{"ttKavv3x__", Builtin::kUndefined}, {"mt3x3", Builtin::kUndefined},
{"xx83x3f", Builtin::kUndefined}, {"mm66Issf", Builtin::kUndefined},
{"__qatF3", Builtin::kUndefined}, {"mat3x1f", Builtin::kUndefined},
{"matqx3f", Builtin::kUndefined}, {"Xt3x3", Builtin::kUndefined},
{"33atOx3h", Builtin::kUndefined}, {"LatIx3h", Builtin::kUndefined},
{"mtt63x9oQQ", Builtin::kUndefined}, {"at3fh", Builtin::kUndefined},
{"ma3x66h", Builtin::kUndefined}, {"mYtURD3", Builtin::kUndefined},
{"66aOzx4", Builtin::kUndefined}, {"mah3x4", Builtin::kUndefined},
{"myyt3x4", Builtin::kUndefined}, {"muqII4", Builtin::kUndefined},
{"HHZx4", Builtin::kUndefined}, {"mat3xH", Builtin::kUndefined},
{"4WWt3q4f", Builtin::kUndefined}, {"at3QQvv", Builtin::kUndefined},
{"mOO3x4f", Builtin::kUndefined}, {"at66ef", Builtin::kUndefined},
{"oatY4f", Builtin::kUndefined}, {"ma7O4f", Builtin::kUndefined},
{"matx", Builtin::kUndefined}, {"m55t3x0DD", Builtin::kUndefined},
{"ma3xFh", Builtin::kUndefined}, {"maH3x4II", Builtin::kUndefined},
{"at3x4w", Builtin::kUndefined}, {"at3x4", Builtin::kUndefined},
{"ma4Gf", Builtin::kUndefined}, {"ma994x2", Builtin::kUndefined},
{"qatKKx2", Builtin::kUndefined}, {"mWWt4Gt2", Builtin::kUndefined},
{"mmmt4x2", Builtin::kUndefined}, {"ay42", Builtin::kUndefined},
{"at4x2f", Builtin::kUndefined}, {"mt4x2f", Builtin::kUndefined},
{"mt4x2q", Builtin::kUndefined}, {"IIaBB4x2f", Builtin::kUndefined},
{"mat4xbb", Builtin::kUndefined}, {"TTat4x833", Builtin::kUndefined},
{"mi4x2h", Builtin::kUndefined}, {"ddUUnntYYx2h", Builtin::kUndefined},
{"maOO4xq", Builtin::kUndefined}, {"m5CCxxdZ", Builtin::kUndefined},
{"matTvvx2h", Builtin::kUndefined}, {"matkkq2h", Builtin::kUndefined},
{"mat4FF3", Builtin::kUndefined}, {"5iitp00", Builtin::kUndefined},
{"mtQ00P", Builtin::kUndefined}, {"mnntIIx3", Builtin::kUndefined},
{"maP4x3", Builtin::kUndefined}, {"ccaKx", Builtin::kUndefined},
{"ma774xss", Builtin::kUndefined}, {"m43KK", Builtin::kUndefined},
{"RRCbb4x3f", Builtin::kUndefined}, {"mat66x3f", Builtin::kUndefined},
{"mXXt4x3f", Builtin::kUndefined}, {"Et4PP3K", Builtin::kUndefined},
{"qaCC4xOOh", Builtin::kUndefined}, {"xxatx3h", Builtin::kUndefined},
{"ma4s3L", Builtin::kUndefined}, {"qat4x3h", Builtin::kUndefined},
{"mXt4x3h", Builtin::kUndefined}, {"MMayySrxh", Builtin::kUndefined},
{"mat4x", Builtin::kUndefined}, {"uat4", Builtin::kUndefined},
{"qqt4", Builtin::kUndefined}, {"tx4", Builtin::kUndefined},
{"mat4x22", Builtin::kUndefined}, {"ma54FF4", Builtin::kUndefined},
{"myzz40XX", Builtin::kUndefined}, {"rra444z4f", Builtin::kUndefined},
{"matVViP", Builtin::kUndefined}, {"matWW", Builtin::kUndefined},
{"mannC4f", Builtin::kUndefined}, {"CatZJXx4f", Builtin::kUndefined},
{"atx4AHHq", Builtin::kUndefined}, {"maPPx4h", Builtin::kUndefined},
{"may4x4", Builtin::kUndefined}, {"mat4c4h", Builtin::kUndefined},
{"aOOOZZh", Builtin::kUndefined}, {"matPPll6h", Builtin::kUndefined},
{"V", Builtin::kUndefined}, {"9tyy", Builtin::kUndefined},
{"ptf__", Builtin::kUndefined}, {"ptKK", Builtin::kUndefined},
{"4lMT", Builtin::kUndefined}, {"x_", Builtin::kUndefined},
{"sNNmplg", Builtin::kUndefined}, {"ayKer", Builtin::kUndefined},
{"uubpXXer", Builtin::kUndefined}, {"szmpVek", Builtin::kUndefined},
{"samler", Builtin::kUndefined}, {"sampqeK", Builtin::kUndefined},
{"m88ler_cQmparisoK", Builtin::kUndefined}, {"sampler_comparisn", Builtin::kUndefined},
{"qa9ler_comparison", Builtin::kUndefined}, {"sapler_comparisVVn", Builtin::kUndefined},
{"sampler_comparis11n", Builtin::kUndefined}, {"samplerIcompaAUison", Builtin::kUndefined},
{"teiiu22eF1d", Builtin::kUndefined}, {"jexurbRd", Builtin::kUndefined},
{"tex77ur_1d", Builtin::kUndefined}, {"exure_YYd", Builtin::kUndefined},
{"te2urNN_1d", Builtin::kUndefined}, {"exture_1d", Builtin::kUndefined},
{"texturVV_2d", Builtin::kUndefined}, {"texxxur_1d", Builtin::kUndefined},
{"WWFxtu11e_wd", Builtin::kUndefined}, {"tJxucce_2d", Builtin::kUndefined},
{"txture_2ww", Builtin::kUndefined}, {"texure_JJd", Builtin::kUndefined},
{"texture_2d_arrDy", Builtin::kUndefined}, {"lDexture_fCC_arraU", Builtin::kUndefined},
{"teKtre_2d_array", Builtin::kUndefined}, {"tegture_2d_array", Builtin::kUndefined},
{"texhure_2fra11raPP", Builtin::kUndefined}, {"teCCure2d_arra", Builtin::kUndefined},
{"texture3d", Builtin::kUndefined}, {"textue_3d", Builtin::kUndefined},
{"texture_3YY", Builtin::kUndefined}, {"tIx__ure_3d", Builtin::kUndefined},
{"HHtxtrkk_3d", Builtin::kUndefined}, {"texurettPP", Builtin::kUndefined},
{"texrrure_cube", Builtin::kUndefined}, {"tddx3ure_cube", Builtin::kUndefined},
{"tssxturWW_cue", Builtin::kUndefined}, {"teKyyur_cube", Builtin::kUndefined},
{"teYure_cube", Builtin::kUndefined}, {"teturecub", Builtin::kUndefined},
{"txture_Lufe_arraq", Builtin::kUndefined}, {"textinne_c03e_array", Builtin::kUndefined},
{"te22ture_uuubevvfray", Builtin::kUndefined}, {"nextCCruuvcubK_array", Builtin::kUndefined},
{"texturecube_aray", Builtin::kUndefined}, {"tXxturellcbe_array", Builtin::kUndefined},
{"texture_Yepth_2", Builtin::kUndefined}, {"tppxture_depth_2d", Builtin::kUndefined},
{"teytYYEe_77epth_2d", Builtin::kUndefined}, {"txture_deptww_2d", Builtin::kUndefined},
{"teMture_deootd2d", Builtin::kUndefined}, {"gexturedemmthuu2", Builtin::kUndefined},
{"texMMre_depth_2d_array", Builtin::kUndefined}, {"texmmre_depthaa2daray", Builtin::kUndefined},
{"texture_depth_2d_arra55", Builtin::kUndefined}, {"texture_RRepth_Td_ccZray", Builtin::kUndefined},
{"texture_deh_2d_aNray", Builtin::kUndefined}, {"text88re_depthTOd_array", Builtin::kUndefined},
{"te3ture_dpth_cO3be", Builtin::kUndefined}, {"texture_depth_cm00e", Builtin::kUndefined},
{"texture_depth_cub3", Builtin::kUndefined}, {"texture_Bmepth_cube", Builtin::kUndefined},
{"Iexturedepth_cume", Builtin::kUndefined}, {"Mextre_ppeph_cube", Builtin::kUndefined},
{"texture_depthnncube_Krrry", Builtin::kUndefined}, {"texturOO_depth_cub_array", Builtin::kUndefined},
{"texture_dth_XXube_rra", Builtin::kUndefined}, {"GeGGture_depthcube_array", Builtin::kUndefined},
{"textIre_depph_ubeLLarray", Builtin::kUndefined}, {"texture11Hdepth_cube_array", Builtin::kUndefined},
{"txtfre_depthmultisampled_2d", Builtin::kUndefined}, {"textu6e_FFepth_multiameeled_2d", Builtin::kUndefined},
{"texURuYe_Depthmultisampled_2d", Builtin::kUndefined}, {"texture_epth_mltisampled_2d", Builtin::kUndefined},
{"texture_depth_multisamphed_2d", Builtin::kUndefined}, {"texture_depth_mullKsaiipled_2d", Builtin::kUndefined},
{"teqtureuIIextnal", Builtin::kUndefined}, {"texture_extenal", Builtin::kUndefined},
{"texture_externaH", Builtin::kUndefined}, {"IIext99reexvvernal", Builtin::kUndefined},
{"texre_externaQvv", Builtin::kUndefined}, {"texture_externl", Builtin::kUndefined},
{"textureemultismp66ed_d", Builtin::kUndefined}, {"texture_mhltisampled_2d", Builtin::kUndefined},
{"tW7trO_multisampled_2d", Builtin::kUndefined}, {"texturemPllltisampzzed_2d", Builtin::kUndefined},
{"texture_mult550ampled_2DD", Builtin::kUndefined}, {"exture_mltisamed_2d", Builtin::kUndefined},
{"teHture_storIIge_1d", Builtin::kUndefined}, {"texture_qqtoragff_1", Builtin::kUndefined},
{"textue_storage_1d", Builtin::kUndefined}, {"textre_JJddorage_1W", Builtin::kUndefined},
{"rexture_storae_1d", Builtin::kUndefined}, {"XXrxture_storae1zz", Builtin::kUndefined},
{"texture_stolage_2d", Builtin::kUndefined}, {"texturestorag2_2d", Builtin::kUndefined},
{"txture_JJtGrgtt_2d", Builtin::kUndefined}, {"yyNxture_storage_2d", Builtin::kUndefined},
{"yexture_storage2d", Builtin::kUndefined}, {"etue_storage_2OO", Builtin::kUndefined},
{"texture_storage_2d_rray", Builtin::kUndefined}, {"reutuPe_storZgeE2d_array", Builtin::kUndefined},
{"texture_IItorage_2d_BBrray", Builtin::kUndefined}, {"texlure_storddeee_d_22rray", Builtin::kUndefined},
{"33exture_TTtorge_Kd_ar88ay", Builtin::kUndefined}, {"texture_mtorage_2V_a9ra", Builtin::kUndefined},
{"texSnnYUUure_storage_3d", Builtin::kUndefined}, {"teII1re_storage_3d", Builtin::kUndefined},
{"textuxe_5torCCdZ_3d", Builtin::kUndefined}, {"texture_storagb_3d", Builtin::kUndefined},
{"tkkxture_storaqe_3d", Builtin::kUndefined}, {"texizrestorge73d", Builtin::kUndefined},
{"5i00", Builtin::kUndefined}, {"u3oi", Builtin::kUndefined},
{"unII2", Builtin::kUndefined}, {"3", Builtin::kUndefined},
{"cc", Builtin::kUndefined}, {"S2", Builtin::kUndefined},
{"KK", Builtin::kUndefined}, {"e22", Builtin::kUndefined},
{"66ec2", Builtin::kUndefined}, {"1eC2", Builtin::kUndefined},
{"PPEK", Builtin::kUndefined}, {"vf8c2", Builtin::kUndefined},
{"vexxf", Builtin::kUndefined}, {"c2f", Builtin::kUndefined},
{"qec2f", Builtin::kUndefined}, {"JJecSSf", Builtin::kUndefined},
{"veSyMMr", Builtin::kUndefined}, {"92f", Builtin::kUndefined},
{"v2u", Builtin::kUndefined}, {"vbbJJ2TT", Builtin::kUndefined},
{"ec", Builtin::kUndefined}, {"e66h", Builtin::kUndefined},
{"5eFF2h", Builtin::kUndefined}, {"u662h", Builtin::kUndefined},
{"rrecz44", Builtin::kUndefined}, {"vW2i", Builtin::kUndefined},
{"vWW", Builtin::kUndefined}, {"v2i", Builtin::kUndefined},
{"XJecCZZ", Builtin::kUndefined}, {"veci", Builtin::kUndefined},
{"vePP2", Builtin::kUndefined}, {"rec2u", Builtin::kUndefined},
{"vec2c", Builtin::kUndefined}, {"2ec2B", Builtin::kUndefined},
{"ve6ll2u", Builtin::kUndefined}, {"vcBBu", Builtin::kUndefined},
{"vey99", Builtin::kUndefined}, {"veRR", Builtin::kUndefined},
{"vKKc3", Builtin::kUndefined}, {"VLL0", Builtin::kUndefined},
{"x_3", Builtin::kUndefined}, {"KOe3", Builtin::kUndefined},
{"Ky3", Builtin::kUndefined}, {"vgwcf", Builtin::kUndefined},
{"zek3f", Builtin::kUndefined}, {"vLphf", Builtin::kUndefined},
{"veKSf", Builtin::kUndefined}, {"eiiEf", Builtin::kUndefined},
{"vc3h", Builtin::kUndefined}, {"ec3h", Builtin::kUndefined},
{"ec3VV", Builtin::kUndefined}, {"UU883", Builtin::kUndefined},
{"IAAc3h", Builtin::kUndefined}, {"rrecvvh", Builtin::kUndefined},
{"jbR", Builtin::kUndefined}, {"ecmm", Builtin::kUndefined},
{"veY4", Builtin::kUndefined}, {"vec4j", Builtin::kUndefined},
{"ec3i", Builtin::kUndefined}, {"vec3X", Builtin::kUndefined},
{"vc911", Builtin::kUndefined}, {"vec38", Builtin::kUndefined},
{"mmccu", Builtin::kUndefined}, {"vecvEE", Builtin::kUndefined},
{"vJJcu", Builtin::kUndefined}, {"z99ci", Builtin::kUndefined},
{"lDCfcU", Builtin::kUndefined}, {"JJGeQQ4", Builtin::kUndefined},
{"veg4", Builtin::kUndefined}, {"ssec4", Builtin::kUndefined},
{"CC", Builtin::kUndefined}, {"PecK", Builtin::kUndefined},
{"ec4f", Builtin::kUndefined}, {"tpc4f", Builtin::kUndefined},
{"vIc__f", Builtin::kUndefined}, {"vec", Builtin::kUndefined},
{"ePPtt", Builtin::kUndefined}, {"MMec4f", Builtin::kUndefined},
{"v3dc4h", Builtin::kUndefined}, {"vJJc40", Builtin::kUndefined},
{"vcyyh", Builtin::kUndefined}, {"8c", Builtin::kUndefined},
{"u4", Builtin::kUndefined}, {"vecggKh", Builtin::kUndefined},
{"v03nni", Builtin::kUndefined}, {"vecfi", Builtin::kUndefined},
{"Cuuecnv", Builtin::kUndefined}, {"vec47Q", Builtin::kUndefined},
{"vX4ll", Builtin::kUndefined}, {"veY4i", Builtin::kUndefined},
{"vocppu", Builtin::kUndefined}, {"keSu", Builtin::kUndefined},
{"vwwc4", Builtin::kUndefined}, {"n422", Builtin::kUndefined},
{"veuug", Builtin::kUndefined}, {"vFFu", Builtin::kUndefined},
}; };
using BuiltinParseTest = testing::TestWithParam<Case>; using BuiltinParseTest = testing::TestWithParam<Case>;

2
src/tint/builtin/builtin_value_bench.cc
View File

@ -129,7 +129,7 @@ void BuiltinValueParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(BuiltinValueParser); BENCHMARK(BuiltinValueParser);

4
src/tint/builtin/diagnostic_rule_bench.cc
View File

@ -41,7 +41,7 @@ void CoreDiagnosticRuleParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(CoreDiagnosticRuleParser); BENCHMARK(CoreDiagnosticRuleParser);
@ -56,7 +56,7 @@ void ChromiumDiagnosticRuleParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(ChromiumDiagnosticRuleParser); BENCHMARK(ChromiumDiagnosticRuleParser);

2
src/tint/builtin/diagnostic_severity_bench.cc
View File

@ -42,7 +42,7 @@ void DiagnosticSeverityParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(DiagnosticSeverityParser); BENCHMARK(DiagnosticSeverityParser);

2
src/tint/builtin/extension_bench.cc
View File

@ -80,7 +80,7 @@ void ExtensionParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(ExtensionParser); BENCHMARK(ExtensionParser);

2
src/tint/builtin/interpolation_sampling_bench.cc
View File

@ -40,7 +40,7 @@ void InterpolationSamplingParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(InterpolationSamplingParser); BENCHMARK(InterpolationSamplingParser);

2
src/tint/builtin/interpolation_type_bench.cc
View File

@ -43,7 +43,7 @@ void InterpolationTypeParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(InterpolationTypeParser); BENCHMARK(InterpolationTypeParser);

2
src/tint/builtin/texel_format_bench.cc
View File

@ -62,7 +62,7 @@ void TexelFormatParser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK(TexelFormatParser); BENCHMARK(TexelFormatParser);

26
src/tint/intrinsics.def
View File

@ -218,6 +218,32 @@ enum builtin_type {
// Internal types. // Internal types.
__packed_vec3 __packed_vec3
__atomic_compare_exchange_result_i32
__atomic_compare_exchange_result_u32
__frexp_result_abstract
__frexp_result_f16
__frexp_result_f32
__frexp_result_vec2_abstract
__frexp_result_vec2_f16
__frexp_result_vec2_f32
__frexp_result_vec3_abstract
__frexp_result_vec3_f16
__frexp_result_vec3_f32
__frexp_result_vec4_abstract
__frexp_result_vec4_f16
__frexp_result_vec4_f32
__modf_result_abstract
__modf_result_f16
__modf_result_f32
__modf_result_vec2_abstract
__modf_result_vec2_f16
__modf_result_vec2_f32
__modf_result_vec3_abstract
__modf_result_vec3_f16
__modf_result_vec3_f32
__modf_result_vec4_abstract
__modf_result_vec4_f16
__modf_result_vec4_f32
} }
// https://gpuweb.github.io/gpuweb/wgsl/#attributes // https://gpuweb.github.io/gpuweb/wgsl/#attributes

2
src/tint/program_builder.h
View File

@ -1470,7 +1470,7 @@ class ProgramBuilder {
/// @param name the symbol string /// @param name the symbol string
/// @return a Symbol with the given name /// @return a Symbol with the given name
Symbol Sym(const std::string& name) { return Symbols().Register(name); } Symbol Sym(std::string_view name) { return Symbols().Register(name); }
/// @param enumerator the enumerator /// @param enumerator the enumerator
/// @return a Symbol with the given enum value /// @return a Symbol with the given enum value

111
src/tint/resolver/builtin_structs.cc
View File

@ -29,12 +29,12 @@ namespace tint::resolver {
namespace { namespace {
struct NameAndType { struct NameAndType {
std::string name; std::string_view name;
const type::Type* type; const type::Type* type;
}; };
sem::Struct* BuildStruct(ProgramBuilder& b, sem::Struct* BuildStruct(ProgramBuilder& b,
std::string name, builtin::Builtin name,
std::initializer_list<NameAndType> member_names_and_types) { std::initializer_list<NameAndType> member_names_and_types) {
uint32_t offset = 0; uint32_t offset = 0;
uint32_t max_align = 0; uint32_t max_align = 0;
@ -69,45 +69,64 @@ sem::Struct* BuildStruct(ProgramBuilder& b,
} }
} // namespace } // namespace
constexpr std::array kModfVecF32Names{
builtin::Builtin::kModfResultVec2F32,
builtin::Builtin::kModfResultVec3F32,
builtin::Builtin::kModfResultVec4F32,
};
constexpr std::array kModfVecF16Names{
builtin::Builtin::kModfResultVec2F16,
builtin::Builtin::kModfResultVec3F16,
builtin::Builtin::kModfResultVec4F16,
};
constexpr std::array kModfVecAbstractNames{
builtin::Builtin::kModfResultVec2Abstract,
builtin::Builtin::kModfResultVec3Abstract,
builtin::Builtin::kModfResultVec4Abstract,
};
sem::Struct* CreateModfResult(ProgramBuilder& b, const type::Type* ty) { sem::Struct* CreateModfResult(ProgramBuilder& b, const type::Type* ty) {
return Switch( return Switch(
ty, ty,
[&](const type::F32*) { [&](const type::F32*) {
return BuildStruct(b, "__modf_result_f32", {{"fract", ty}, {"whole", ty}}); return BuildStruct(b, builtin::Builtin::kModfResultF32, {{"fract", ty}, {"whole", ty}});
}, // }, //
[&](const type::F16*) { [&](const type::F16*) {
return BuildStruct(b, "__modf_result_f16", {{"fract", ty}, {"whole", ty}}); return BuildStruct(b, builtin::Builtin::kModfResultF16, {{"fract", ty}, {"whole", ty}});
}, },
[&](const type::AbstractFloat*) { [&](const type::AbstractFloat*) {
auto* abstract = auto* abstract = BuildStruct(b, builtin::Builtin::kModfResultAbstract,
BuildStruct(b, "__modf_result_abstract", {{"fract", ty}, {"whole", ty}}); {{"fract", ty}, {"whole", ty}});
auto* f32 = b.create<type::F32>(); auto* f32 = b.create<type::F32>();
auto* f16 = b.create<type::F16>(); auto* f16 = b.create<type::F16>();
abstract->SetConcreteTypes(utils::Vector{ abstract->SetConcreteTypes(utils::Vector{
BuildStruct(b, "__modf_result_f32", {{"fract", f32}, {"whole", f32}}), BuildStruct(b, builtin::Builtin::kModfResultF32, {{"fract", f32}, {"whole", f32}}),
BuildStruct(b, "__modf_result_f16", {{"fract", f16}, {"whole", f16}}), BuildStruct(b, builtin::Builtin::kModfResultF16, {{"fract", f16}, {"whole", f16}}),
}); });
return abstract; return abstract;
}, },
[&](const type::Vector* vec) { [&](const type::Vector* vec) {
auto width = vec->Width(); auto width = vec->Width();
auto prefix = "__modf_result_vec" + std::to_string(width);
return Switch( return Switch(
vec->type(), // vec->type(), //
[&](const type::F32*) { [&](const type::F32*) {
return BuildStruct(b, prefix + "_f32", {{"fract", vec}, {"whole", vec}}); return BuildStruct(b, kModfVecF32Names[width - 2],
{{"fract", vec}, {"whole", vec}});
}, },
[&](const type::F16*) { [&](const type::F16*) {
return BuildStruct(b, prefix + "_f16", {{"fract", vec}, {"whole", vec}}); return BuildStruct(b, kModfVecF16Names[width - 2],
{{"fract", vec}, {"whole", vec}});
}, },
[&](const type::AbstractFloat*) { [&](const type::AbstractFloat*) {
auto* vec_f32 = b.create<type::Vector>(b.create<type::F32>(), width); auto* vec_f32 = b.create<type::Vector>(b.create<type::F32>(), width);
auto* vec_f16 = b.create<type::Vector>(b.create<type::F16>(), width); auto* vec_f16 = b.create<type::Vector>(b.create<type::F16>(), width);
auto* abstract = auto* abstract = BuildStruct(b, kModfVecAbstractNames[width - 2],
BuildStruct(b, prefix + "_abstract", {{"fract", vec}, {"whole", vec}}); {{"fract", vec}, {"whole", vec}});
abstract->SetConcreteTypes(utils::Vector{ abstract->SetConcreteTypes(utils::Vector{
BuildStruct(b, prefix + "_f32", {{"fract", vec_f32}, {"whole", vec_f32}}), BuildStruct(b, kModfVecF32Names[width - 2],
BuildStruct(b, prefix + "_f16", {{"fract", vec_f16}, {"whole", vec_f16}}), {{"fract", vec_f32}, {"whole", vec_f32}}),
BuildStruct(b, kModfVecF16Names[width - 2],
{{"fract", vec_f16}, {"whole", vec_f16}}),
}); });
return abstract; return abstract;
}, },
@ -123,53 +142,71 @@ sem::Struct* CreateModfResult(ProgramBuilder& b, const type::Type* ty) {
}); });
} }
constexpr std::array kFrexpVecF32Names{
builtin::Builtin::kFrexpResultVec2F32,
builtin::Builtin::kFrexpResultVec3F32,
builtin::Builtin::kFrexpResultVec4F32,
};
constexpr std::array kFrexpVecF16Names{
builtin::Builtin::kFrexpResultVec2F16,
builtin::Builtin::kFrexpResultVec3F16,
builtin::Builtin::kFrexpResultVec4F16,
};
constexpr std::array kFrexpVecAbstractNames{
builtin::Builtin::kFrexpResultVec2Abstract,
builtin::Builtin::kFrexpResultVec3Abstract,
builtin::Builtin::kFrexpResultVec4Abstract,
};
sem::Struct* CreateFrexpResult(ProgramBuilder& b, const type::Type* ty) { sem::Struct* CreateFrexpResult(ProgramBuilder& b, const type::Type* ty) {
return Switch( return Switch(
ty, // ty, //
[&](const type::F32*) { [&](const type::F32*) {
auto* i32 = b.create<type::I32>(); auto* i32 = b.create<type::I32>();
return BuildStruct(b, "__frexp_result_f32", {{"fract", ty}, {"exp", i32}}); return BuildStruct(b, builtin::Builtin::kFrexpResultF32, {{"fract", ty}, {"exp", i32}});
}, },
[&](const type::F16*) { [&](const type::F16*) {
auto* i32 = b.create<type::I32>(); auto* i32 = b.create<type::I32>();
return BuildStruct(b, "__frexp_result_f16", {{"fract", ty}, {"exp", i32}}); return BuildStruct(b, builtin::Builtin::kFrexpResultF16, {{"fract", ty}, {"exp", i32}});
}, },
[&](const type::AbstractFloat*) { [&](const type::AbstractFloat*) {
auto* f32 = b.create<type::F32>(); auto* f32 = b.create<type::F32>();
auto* f16 = b.create<type::F16>(); auto* f16 = b.create<type::F16>();
auto* i32 = b.create<type::I32>(); auto* i32 = b.create<type::I32>();
auto* ai = b.create<type::AbstractInt>(); auto* ai = b.create<type::AbstractInt>();
auto* abstract = auto* abstract = BuildStruct(b, builtin::Builtin::kFrexpResultAbstract,
BuildStruct(b, "__frexp_result_abstract", {{"fract", ty}, {"exp", ai}}); {{"fract", ty}, {"exp", ai}});
abstract->SetConcreteTypes(utils::Vector{ abstract->SetConcreteTypes(utils::Vector{
BuildStruct(b, "__frexp_result_f32", {{"fract", f32}, {"exp", i32}}), BuildStruct(b, builtin::Builtin::kFrexpResultF32, {{"fract", f32}, {"exp", i32}}),
BuildStruct(b, "__frexp_result_f16", {{"fract", f16}, {"exp", i32}}), BuildStruct(b, builtin::Builtin::kFrexpResultF16, {{"fract", f16}, {"exp", i32}}),
}); });
return abstract; return abstract;
}, },
[&](const type::Vector* vec) { [&](const type::Vector* vec) {
auto width = vec->Width(); auto width = vec->Width();
auto prefix = "__frexp_result_vec" + std::to_string(width);
return Switch( return Switch(
vec->type(), // vec->type(), //
[&](const type::F32*) { [&](const type::F32*) {
auto* vec_i32 = b.create<type::Vector>(b.create<type::I32>(), width); auto* vec_i32 = b.create<type::Vector>(b.create<type::I32>(), width);
return BuildStruct(b, prefix + "_f32", {{"fract", ty}, {"exp", vec_i32}}); return BuildStruct(b, kFrexpVecF32Names[width - 2],
{{"fract", ty}, {"exp", vec_i32}});
}, },
[&](const type::F16*) { [&](const type::F16*) {
auto* vec_i32 = b.create<type::Vector>(b.create<type::I32>(), width); auto* vec_i32 = b.create<type::Vector>(b.create<type::I32>(), width);
return BuildStruct(b, prefix + "_f16", {{"fract", ty}, {"exp", vec_i32}}); return BuildStruct(b, kFrexpVecF16Names[width - 2],
{{"fract", ty}, {"exp", vec_i32}});
}, },
[&](const type::AbstractFloat*) { [&](const type::AbstractFloat*) {
auto* vec_f32 = b.create<type::Vector>(b.create<type::F32>(), width); auto* vec_f32 = b.create<type::Vector>(b.create<type::F32>(), width);
auto* vec_f16 = b.create<type::Vector>(b.create<type::F16>(), width); auto* vec_f16 = b.create<type::Vector>(b.create<type::F16>(), width);
auto* vec_i32 = b.create<type::Vector>(b.create<type::I32>(), width); auto* vec_i32 = b.create<type::Vector>(b.create<type::I32>(), width);
auto* vec_ai = b.create<type::Vector>(b.create<type::AbstractInt>(), width); auto* vec_ai = b.create<type::Vector>(b.create<type::AbstractInt>(), width);
auto* abstract = auto* abstract = BuildStruct(b, kFrexpVecAbstractNames[width - 2],
BuildStruct(b, prefix + "_abstract", {{"fract", ty}, {"exp", vec_ai}}); {{"fract", ty}, {"exp", vec_ai}});
abstract->SetConcreteTypes(utils::Vector{ abstract->SetConcreteTypes(utils::Vector{
BuildStruct(b, prefix + "_f32", {{"fract", vec_f32}, {"exp", vec_i32}}), BuildStruct(b, kFrexpVecF32Names[width - 2],
BuildStruct(b, prefix + "_f16", {{"fract", vec_f16}, {"exp", vec_i32}}), {{"fract", vec_f32}, {"exp", vec_i32}}),
BuildStruct(b, kFrexpVecF16Names[width - 2],
{{"fract", vec_f16}, {"exp", vec_i32}}),
}); });
return abstract; return abstract;
}, },
@ -186,9 +223,21 @@ sem::Struct* CreateFrexpResult(ProgramBuilder& b, const type::Type* ty) {
} }
sem::Struct* CreateAtomicCompareExchangeResult(ProgramBuilder& b, const type::Type* ty) { sem::Struct* CreateAtomicCompareExchangeResult(ProgramBuilder& b, const type::Type* ty) {
return BuildStruct( return Switch(
b, "__atomic_compare_exchange_result" + ty->FriendlyName(), ty, //
{{"old_value", const_cast<type::Type*>(ty)}, {"exchanged", b.create<type::Bool>()}}); [&](const type::I32*) {
return BuildStruct(b, builtin::Builtin::kAtomicCompareExchangeResultI32,
{{"old_value", ty}, {"exchanged", b.create<type::Bool>()}});
},
[&](const type::U32*) {
return BuildStruct(b, builtin::Builtin::kAtomicCompareExchangeResultU32,
{{"old_value", ty}, {"exchanged", b.create<type::Bool>()}});
},
[&](Default) {
TINT_ICE(Resolver, b.Diagnostics())
<< "unhandled atomic_compare_exchange type: " << b.FriendlyName(ty);
return nullptr;
});
} }
} // namespace tint::resolver } // namespace tint::resolver

73
src/tint/resolver/builtin_structs_test.cc Normal file
View File

@ -0,0 +1,73 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/resolver/resolver.h"
#include "src/tint/resolver/resolver_test_helper.h"
#include "gmock/gmock.h"
using namespace tint::number_suffixes; // NOLINT
namespace tint::resolver {
namespace {
////////////////////////////////////////////////////////////////////////////////
// access
////////////////////////////////////////////////////////////////////////////////
using ResolverBuiltinStructs = ResolverTestWithParam<builtin::Builtin>;
TEST_P(ResolverBuiltinStructs, Resolve) {
Enable(builtin::Extension::kF16);
// var<private> p : NAME;
auto* var = GlobalVar("p", ty(GetParam()), builtin::AddressSpace::kPrivate);
ASSERT_TRUE(r()->Resolve()) << r()->error();
auto* str = As<sem::Struct>(TypeOf(var)->UnwrapRef());
ASSERT_NE(str, nullptr);
EXPECT_EQ(str->Name().Name(), utils::ToString(GetParam()));
EXPECT_EQ(str->Declaration(), nullptr);
}
INSTANTIATE_TEST_SUITE_P(,
ResolverBuiltinStructs,
testing::Values(builtin::Builtin::kAtomicCompareExchangeResultI32,
builtin::Builtin::kAtomicCompareExchangeResultU32,
builtin::Builtin::kFrexpResultAbstract,
builtin::Builtin::kFrexpResultF16,
builtin::Builtin::kFrexpResultF32,
builtin::Builtin::kFrexpResultVec2Abstract,
builtin::Builtin::kFrexpResultVec2F16,
builtin::Builtin::kFrexpResultVec2F32,
builtin::Builtin::kFrexpResultVec3Abstract,
builtin::Builtin::kFrexpResultVec3F16,
builtin::Builtin::kFrexpResultVec3F32,
builtin::Builtin::kFrexpResultVec4Abstract,
builtin::Builtin::kFrexpResultVec4F16,
builtin::Builtin::kFrexpResultVec4F32,
builtin::Builtin::kModfResultAbstract,
builtin::Builtin::kModfResultF16,
builtin::Builtin::kModfResultF32,
builtin::Builtin::kModfResultVec2Abstract,
builtin::Builtin::kModfResultVec2F16,
builtin::Builtin::kModfResultVec2F32,
builtin::Builtin::kModfResultVec3Abstract,
builtin::Builtin::kModfResultVec3F16,
builtin::Builtin::kModfResultVec3F32,
builtin::Builtin::kModfResultVec4Abstract,
builtin::Builtin::kModfResultVec4F16,
builtin::Builtin::kModfResultVec4F32));
} // namespace
} // namespace tint::resolver

57
src/tint/resolver/resolver.cc
View File

@ -43,6 +43,7 @@
#include "src/tint/ast/while_statement.h" #include "src/tint/ast/while_statement.h"
#include "src/tint/ast/workgroup_attribute.h" #include "src/tint/ast/workgroup_attribute.h"
#include "src/tint/builtin/builtin.h" #include "src/tint/builtin/builtin.h"
#include "src/tint/resolver/builtin_structs.h"
#include "src/tint/resolver/uniformity.h" #include "src/tint/resolver/uniformity.h"
#include "src/tint/sem/break_if_statement.h" #include "src/tint/sem/break_if_statement.h"
#include "src/tint/sem/builtin_enum_expression.h" #include "src/tint/sem/builtin_enum_expression.h"
@ -2336,6 +2337,7 @@ type::Type* Resolver::BuiltinType(builtin::Builtin builtin_ty, const ast::Identi
auto check_no_tmpl_args = [&](type::Type* ty) -> type::Type* { auto check_no_tmpl_args = [&](type::Type* ty) -> type::Type* {
return TINT_LIKELY(CheckNotTemplated("type", ident)) ? ty : nullptr; return TINT_LIKELY(CheckNotTemplated("type", ident)) ? ty : nullptr;
}; };
auto af = [&] { return b.create<type::AbstractFloat>(); };
auto f32 = [&] { return b.create<type::F32>(); }; auto f32 = [&] { return b.create<type::F32>(); };
auto i32 = [&] { return b.create<type::I32>(); }; auto i32 = [&] { return b.create<type::I32>(); };
auto u32 = [&] { return b.create<type::U32>(); }; auto u32 = [&] { return b.create<type::U32>(); };
@ -2737,9 +2739,60 @@ type::Type* Resolver::BuiltinType(builtin::Builtin builtin_ty, const ast::Identi
return storage_texture(type::TextureDimension::k2dArray); return storage_texture(type::TextureDimension::k2dArray);
case builtin::Builtin::kTextureStorage3D: case builtin::Builtin::kTextureStorage3D:
return storage_texture(type::TextureDimension::k3d); return storage_texture(type::TextureDimension::k3d);
case builtin::Builtin::kPackedVec3: { case builtin::Builtin::kPackedVec3:
return packed_vec3_t(); return packed_vec3_t();
} case builtin::Builtin::kAtomicCompareExchangeResultI32:
return CreateAtomicCompareExchangeResult(*builder_, i32());
case builtin::Builtin::kAtomicCompareExchangeResultU32:
return CreateAtomicCompareExchangeResult(*builder_, u32());
case builtin::Builtin::kFrexpResultAbstract:
return CreateFrexpResult(*builder_, af());
case builtin::Builtin::kFrexpResultF16:
return CreateFrexpResult(*builder_, f16());
case builtin::Builtin::kFrexpResultF32:
return CreateFrexpResult(*builder_, f32());
case builtin::Builtin::kFrexpResultVec2Abstract:
return CreateFrexpResult(*builder_, vec(af(), 2));
case builtin::Builtin::kFrexpResultVec2F16:
return CreateFrexpResult(*builder_, vec(f16(), 2));
case builtin::Builtin::kFrexpResultVec2F32:
return CreateFrexpResult(*builder_, vec(f32(), 2));
case builtin::Builtin::kFrexpResultVec3Abstract:
return CreateFrexpResult(*builder_, vec(af(), 3));
case builtin::Builtin::kFrexpResultVec3F16:
return CreateFrexpResult(*builder_, vec(f16(), 3));
case builtin::Builtin::kFrexpResultVec3F32:
return CreateFrexpResult(*builder_, vec(f32(), 3));
case builtin::Builtin::kFrexpResultVec4Abstract:
return CreateFrexpResult(*builder_, vec(af(), 4));
case builtin::Builtin::kFrexpResultVec4F16:
return CreateFrexpResult(*builder_, vec(f16(), 4));
case builtin::Builtin::kFrexpResultVec4F32:
return CreateFrexpResult(*builder_, vec(f32(), 4));
case builtin::Builtin::kModfResultAbstract:
return CreateModfResult(*builder_, af());
case builtin::Builtin::kModfResultF16:
return CreateModfResult(*builder_, f16());
case builtin::Builtin::kModfResultF32:
return CreateModfResult(*builder_, f32());
case builtin::Builtin::kModfResultVec2Abstract:
return CreateModfResult(*builder_, vec(af(), 2));
case builtin::Builtin::kModfResultVec2F16:
return CreateModfResult(*builder_, vec(f16(), 2));
case builtin::Builtin::kModfResultVec2F32:
return CreateModfResult(*builder_, vec(f32(), 2));
case builtin::Builtin::kModfResultVec3Abstract:
return CreateModfResult(*builder_, vec(af(), 3));
case builtin::Builtin::kModfResultVec3F16:
return CreateModfResult(*builder_, vec(f16(), 3));
case builtin::Builtin::kModfResultVec3F32:
return CreateModfResult(*builder_, vec(f32(), 3));
case builtin::Builtin::kModfResultVec4Abstract:
return CreateModfResult(*builder_, vec(af(), 4));
case builtin::Builtin::kModfResultVec4F16:
return CreateModfResult(*builder_, vec(f16(), 4));
case builtin::Builtin::kModfResultVec4F32:
return CreateModfResult(*builder_, vec(f32(), 4));
case builtin::Builtin::kUndefined: case builtin::Builtin::kUndefined:
break; break;
} }

2
src/tint/templates/enums.tmpl.inc
View File

@ -191,7 +191,7 @@ void {{$enum}}Parser(::benchmark::State& state) {
benchmark::DoNotOptimize(result); benchmark::DoNotOptimize(result);
} }
} }
} } // NOLINT(readability/fn_size)
BENCHMARK({{$enum}}Parser); BENCHMARK({{$enum}}Parser);
{{- end -}} {{- end -}}

11
test/tint/bug/chromium/1430309.wgsl Normal file
View File

@ -0,0 +1,11 @@
struct frexp_result_f32 {
f : f32,
}
var<private> a : frexp_result_f32;
var<private> b = frexp(1);
@fragment
fn main() -> @location(0) vec4f {
return vec4f(a.f, b.fract, 0, 0);
}

26
test/tint/bug/chromium/1430309.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,26 @@
struct frexp_result_f32_1 {
float fract;
int exp;
};
struct frexp_result_f32 {
float f;
};
static frexp_result_f32 a = (frexp_result_f32)0;
const frexp_result_f32_1 c = {0.5f, 1};
static frexp_result_f32_1 b = c;
struct tint_symbol {
float4 value : SV_Target0;
};
float4 main_inner() {
return float4(a.f, b.fract, 0.0f, 0.0f);
}
tint_symbol main() {
const float4 inner_result = main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}

26
test/tint/bug/chromium/1430309.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,26 @@
struct frexp_result_f32_1 {
float fract;
int exp;
};
struct frexp_result_f32 {
float f;
};
static frexp_result_f32 a = (frexp_result_f32)0;
const frexp_result_f32_1 c = {0.5f, 1};
static frexp_result_f32_1 b = c;
struct tint_symbol {
float4 value : SV_Target0;
};
float4 main_inner() {
return float4(a.f, b.fract, 0.0f, 0.0f);
}
tint_symbol main() {
const float4 inner_result = main_inner();
tint_symbol wrapper_result = (tint_symbol)0;
wrapper_result.value = inner_result;
return wrapper_result;
}

25
test/tint/bug/chromium/1430309.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,25 @@
#version 310 es
precision highp float;
struct frexp_result_f32_1 {
float fract;
int exp;
};
layout(location = 0) out vec4 value;
struct frexp_result_f32 {
float f;
};
frexp_result_f32 a = frexp_result_f32(0.0f);
frexp_result_f32_1 b = frexp_result_f32_1(0.5f, 1);
vec4 tint_symbol() {
return vec4(a.f, b.fract, 0.0f, 0.0f);
}
void main() {
vec4 inner_result = tint_symbol();
value = inner_result;
return;
}

34
test/tint/bug/chromium/1430309.wgsl.expected.msl Normal file
View File

@ -0,0 +1,34 @@
#include <metal_stdlib>
using namespace metal;
struct frexp_result_f32_1 {
float fract;
int exp;
};
struct frexp_result_f32 {
float f;
};
struct tint_private_vars_struct {
frexp_result_f32 a;
frexp_result_f32_1 b;
};
struct tint_symbol_1 {
float4 value [[color(0)]];
};
float4 tint_symbol_inner(thread tint_private_vars_struct* const tint_private_vars) {
return float4((*(tint_private_vars)).a.f, (*(tint_private_vars)).b.fract, 0.0f, 0.0f);
}
fragment tint_symbol_1 tint_symbol() {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.b = frexp_result_f32_1{.fract=0.5f, .exp=1};
float4 const inner_result = tint_symbol_inner(&(tint_private_vars));
tint_symbol_1 wrapper_result = {};
wrapper_result.value = inner_result;
return wrapper_result;
}

61
test/tint/bug/chromium/1430309.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,61 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 34
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %main "main" %value
OpExecutionMode %main OriginUpperLeft
OpName %value "value"
OpName %frexp_result_f32 "frexp_result_f32"
OpMemberName %frexp_result_f32 0 "f"
OpName %a "a"
OpName %__frexp_result_f32 "__frexp_result_f32"
OpMemberName %__frexp_result_f32 0 "fract"
OpMemberName %__frexp_result_f32 1 "exp"
OpName %b "b"
OpName %main_inner "main_inner"
OpName %main "main"
OpDecorate %value Location 0
OpMemberDecorate %frexp_result_f32 0 Offset 0
OpMemberDecorate %__frexp_result_f32 0 Offset 0
OpMemberDecorate %__frexp_result_f32 1 Offset 4
%float = OpTypeFloat 32
%v4float = OpTypeVector %float 4
%_ptr_Output_v4float = OpTypePointer Output %v4float
%5 = OpConstantNull %v4float
%value = OpVariable %_ptr_Output_v4float Output %5
%frexp_result_f32 = OpTypeStruct %float
%_ptr_Private_frexp_result_f32 = OpTypePointer Private %frexp_result_f32
%9 = OpConstantNull %frexp_result_f32
%a = OpVariable %_ptr_Private_frexp_result_f32 Private %9
%int = OpTypeInt 32 1
%__frexp_result_f32 = OpTypeStruct %float %int
%float_0_5 = OpConstant %float 0.5
%int_1 = OpConstant %int 1
%14 = OpConstantComposite %__frexp_result_f32 %float_0_5 %int_1
%_ptr_Private___frexp_result_f32 = OpTypePointer Private %__frexp_result_f32
%b = OpVariable %_ptr_Private___frexp_result_f32 Private %14
%17 = OpTypeFunction %v4float
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%_ptr_Private_float = OpTypePointer Private %float
%27 = OpConstantNull %float
%void = OpTypeVoid
%29 = OpTypeFunction %void
%main_inner = OpFunction %v4float None %17
%19 = OpLabel
%23 = OpAccessChain %_ptr_Private_float %a %uint_0
%24 = OpLoad %float %23
%25 = OpAccessChain %_ptr_Private_float %b %uint_0
%26 = OpLoad %float %25
%28 = OpCompositeConstruct %v4float %24 %26 %27 %27
OpReturnValue %28
OpFunctionEnd
%main = OpFunction %void None %29
%32 = OpLabel
%33 = OpFunctionCall %v4float %main_inner
OpStore %value %33
OpReturn
OpFunctionEnd

12
test/tint/bug/chromium/1430309.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,12 @@
struct frexp_result_f32 {
f : f32,
}
var<private> a : frexp_result_f32;
var<private> b = frexp(1);
@fragment
fn main() -> @location(0) vec4f {
return vec4f(a.f, b.fract, 0, 0);
}

6
test/tint/bug/tint/1573.wgsl.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -12,10 +12,10 @@ layout(binding = 0, std430) buffer a_block_ssbo {
void tint_symbol() { void tint_symbol() {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(a.inner, 0u, value); atomic_compare_result.old_value = atomicCompSwap(a.inner, 0u, value);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u;
atomic_compare_exchange_resultu32 result = atomic_compare_result; atomic_compare_exchange_result_u32 result = atomic_compare_result;
} }
layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in; layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;

6
test/tint/bug/tint/1573.wgsl.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -14,7 +14,7 @@ atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint
kernel void tint_symbol(device atomic_uint* tint_symbol_1 [[buffer(0)]]) { kernel void tint_symbol(device atomic_uint* tint_symbol_1 [[buffer(0)]]) {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 const result = atomicCompareExchangeWeak_1(tint_symbol_1, 0u, value); atomic_compare_exchange_result_u32 const result = atomicCompareExchangeWeak_1(tint_symbol_1, 0u, value);
return; return;
} }

14
test/tint/bug/tint/1573.wgsl.expected.spvasm
View File

@ -12,15 +12,15 @@
OpName %a "a" OpName %a "a"
OpName %main "main" OpName %main "main"
OpName %value "value" OpName %value "value"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpDecorate %a_block Block OpDecorate %a_block Block
OpMemberDecorate %a_block 0 Offset 0 OpMemberDecorate %a_block 0 Offset 0
OpDecorate %a DescriptorSet 0 OpDecorate %a DescriptorSet 0
OpDecorate %a Binding 0 OpDecorate %a Binding 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%a_block = OpTypeStruct %uint %a_block = OpTypeStruct %uint
%_ptr_StorageBuffer_a_block = OpTypePointer StorageBuffer %a_block %_ptr_StorageBuffer_a_block = OpTypePointer StorageBuffer %a_block
@ -31,7 +31,7 @@
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%12 = OpConstantNull %uint %12 = OpConstantNull %uint
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
@ -43,6 +43,6 @@
%21 = OpLoad %uint %value %21 = OpLoad %uint %value
%22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %12 %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %12
%23 = OpIEqual %bool %22 %12 %23 = OpIEqual %bool %22 %12
%13 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23 %13 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %22 %23
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

28
test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl
View File

@ -1,8 +1,8 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -62,39 +62,39 @@ void main_inner(uint local_invocation_index) {
} }
{ {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 atomic_result_2 = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result_2 = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = 0u; uint atomic_compare_value = 0u;
InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value); InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value);
atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value; atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value;
const atomic_compare_exchange_resultu32 r1 = atomic_result_2; const atomic_compare_exchange_result_u32 r1 = atomic_result_2;
atomic_compare_exchange_resultu32 atomic_result_3 = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result_3 = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value_1 = 0u; uint atomic_compare_value_1 = 0u;
InterlockedCompareExchange(b_u32, atomic_compare_value_1, value, atomic_result_3.old_value); InterlockedCompareExchange(b_u32, atomic_compare_value_1, value, atomic_result_3.old_value);
atomic_result_3.exchanged = atomic_result_3.old_value == atomic_compare_value_1; atomic_result_3.exchanged = atomic_result_3.old_value == atomic_compare_value_1;
const atomic_compare_exchange_resultu32 r2 = atomic_result_3; const atomic_compare_exchange_result_u32 r2 = atomic_result_3;
atomic_compare_exchange_resultu32 atomic_result_4 = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result_4 = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value_2 = 0u; uint atomic_compare_value_2 = 0u;
InterlockedCompareExchange(b_u32, atomic_compare_value_2, value, atomic_result_4.old_value); InterlockedCompareExchange(b_u32, atomic_compare_value_2, value, atomic_result_4.old_value);
atomic_result_4.exchanged = atomic_result_4.old_value == atomic_compare_value_2; atomic_result_4.exchanged = atomic_result_4.old_value == atomic_compare_value_2;
const atomic_compare_exchange_resultu32 r3 = atomic_result_4; const atomic_compare_exchange_result_u32 r3 = atomic_result_4;
} }
{ {
int value = 42; int value = 42;
atomic_compare_exchange_resulti32 atomic_result_5 = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result_5 = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value_3 = 0; int atomic_compare_value_3 = 0;
InterlockedCompareExchange(b_i32, atomic_compare_value_3, value, atomic_result_5.old_value); InterlockedCompareExchange(b_i32, atomic_compare_value_3, value, atomic_result_5.old_value);
atomic_result_5.exchanged = atomic_result_5.old_value == atomic_compare_value_3; atomic_result_5.exchanged = atomic_result_5.old_value == atomic_compare_value_3;
const atomic_compare_exchange_resulti32 r1 = atomic_result_5; const atomic_compare_exchange_result_i32 r1 = atomic_result_5;
atomic_compare_exchange_resulti32 atomic_result_6 = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result_6 = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value_4 = 0; int atomic_compare_value_4 = 0;
InterlockedCompareExchange(b_i32, atomic_compare_value_4, value, atomic_result_6.old_value); InterlockedCompareExchange(b_i32, atomic_compare_value_4, value, atomic_result_6.old_value);
atomic_result_6.exchanged = atomic_result_6.old_value == atomic_compare_value_4; atomic_result_6.exchanged = atomic_result_6.old_value == atomic_compare_value_4;
const atomic_compare_exchange_resulti32 r2 = atomic_result_6; const atomic_compare_exchange_result_i32 r2 = atomic_result_6;
atomic_compare_exchange_resulti32 atomic_result_7 = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result_7 = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value_5 = 0; int atomic_compare_value_5 = 0;
InterlockedCompareExchange(b_i32, atomic_compare_value_5, value, atomic_result_7.old_value); InterlockedCompareExchange(b_i32, atomic_compare_value_5, value, atomic_result_7.old_value);
atomic_result_7.exchanged = atomic_result_7.old_value == atomic_compare_value_5; atomic_result_7.exchanged = atomic_result_7.old_value == atomic_compare_value_5;
const atomic_compare_exchange_resulti32 r3 = atomic_result_7; const atomic_compare_exchange_result_i32 r3 = atomic_result_7;
} }
} }

28
test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl
View File

@ -1,8 +1,8 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -62,39 +62,39 @@ void main_inner(uint local_invocation_index) {
} }
{ {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 atomic_result_2 = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result_2 = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = 0u; uint atomic_compare_value = 0u;
InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value); InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value);
atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value; atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value;
const atomic_compare_exchange_resultu32 r1 = atomic_result_2; const atomic_compare_exchange_result_u32 r1 = atomic_result_2;
atomic_compare_exchange_resultu32 atomic_result_3 = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result_3 = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value_1 = 0u; uint atomic_compare_value_1 = 0u;
InterlockedCompareExchange(b_u32, atomic_compare_value_1, value, atomic_result_3.old_value); InterlockedCompareExchange(b_u32, atomic_compare_value_1, value, atomic_result_3.old_value);
atomic_result_3.exchanged = atomic_result_3.old_value == atomic_compare_value_1; atomic_result_3.exchanged = atomic_result_3.old_value == atomic_compare_value_1;
const atomic_compare_exchange_resultu32 r2 = atomic_result_3; const atomic_compare_exchange_result_u32 r2 = atomic_result_3;
atomic_compare_exchange_resultu32 atomic_result_4 = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result_4 = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value_2 = 0u; uint atomic_compare_value_2 = 0u;
InterlockedCompareExchange(b_u32, atomic_compare_value_2, value, atomic_result_4.old_value); InterlockedCompareExchange(b_u32, atomic_compare_value_2, value, atomic_result_4.old_value);
atomic_result_4.exchanged = atomic_result_4.old_value == atomic_compare_value_2; atomic_result_4.exchanged = atomic_result_4.old_value == atomic_compare_value_2;
const atomic_compare_exchange_resultu32 r3 = atomic_result_4; const atomic_compare_exchange_result_u32 r3 = atomic_result_4;
} }
{ {
int value = 42; int value = 42;
atomic_compare_exchange_resulti32 atomic_result_5 = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result_5 = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value_3 = 0; int atomic_compare_value_3 = 0;
InterlockedCompareExchange(b_i32, atomic_compare_value_3, value, atomic_result_5.old_value); InterlockedCompareExchange(b_i32, atomic_compare_value_3, value, atomic_result_5.old_value);
atomic_result_5.exchanged = atomic_result_5.old_value == atomic_compare_value_3; atomic_result_5.exchanged = atomic_result_5.old_value == atomic_compare_value_3;
const atomic_compare_exchange_resulti32 r1 = atomic_result_5; const atomic_compare_exchange_result_i32 r1 = atomic_result_5;
atomic_compare_exchange_resulti32 atomic_result_6 = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result_6 = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value_4 = 0; int atomic_compare_value_4 = 0;
InterlockedCompareExchange(b_i32, atomic_compare_value_4, value, atomic_result_6.old_value); InterlockedCompareExchange(b_i32, atomic_compare_value_4, value, atomic_result_6.old_value);
atomic_result_6.exchanged = atomic_result_6.old_value == atomic_compare_value_4; atomic_result_6.exchanged = atomic_result_6.old_value == atomic_compare_value_4;
const atomic_compare_exchange_resulti32 r2 = atomic_result_6; const atomic_compare_exchange_result_i32 r2 = atomic_result_6;
atomic_compare_exchange_resulti32 atomic_result_7 = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result_7 = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value_5 = 0; int atomic_compare_value_5 = 0;
InterlockedCompareExchange(b_i32, atomic_compare_value_5, value, atomic_result_7.old_value); InterlockedCompareExchange(b_i32, atomic_compare_value_5, value, atomic_result_7.old_value);
atomic_result_7.exchanged = atomic_result_7.old_value == atomic_compare_value_5; atomic_result_7.exchanged = atomic_result_7.old_value == atomic_compare_value_5;
const atomic_compare_exchange_resulti32 r3 = atomic_result_7; const atomic_compare_exchange_result_i32 r3 = atomic_result_7;
} }
} }

52
test/tint/bug/tint/1574.wgsl.expected.glsl
View File

@ -1,11 +1,11 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -29,63 +29,63 @@ void tint_symbol(uint local_invocation_index) {
barrier(); barrier();
{ {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(a_u32.inner, 0u, value); atomic_compare_result.old_value = atomicCompSwap(a_u32.inner, 0u, value);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u;
atomic_compare_exchange_resultu32 r1 = atomic_compare_result; atomic_compare_exchange_result_u32 r1 = atomic_compare_result;
atomic_compare_exchange_resultu32 atomic_compare_result_1; atomic_compare_exchange_result_u32 atomic_compare_result_1;
atomic_compare_result_1.old_value = atomicCompSwap(a_u32.inner, 0u, value); atomic_compare_result_1.old_value = atomicCompSwap(a_u32.inner, 0u, value);
atomic_compare_result_1.exchanged = atomic_compare_result_1.old_value == 0u; atomic_compare_result_1.exchanged = atomic_compare_result_1.old_value == 0u;
atomic_compare_exchange_resultu32 r2 = atomic_compare_result_1; atomic_compare_exchange_result_u32 r2 = atomic_compare_result_1;
atomic_compare_exchange_resultu32 atomic_compare_result_2; atomic_compare_exchange_result_u32 atomic_compare_result_2;
atomic_compare_result_2.old_value = atomicCompSwap(a_u32.inner, 0u, value); atomic_compare_result_2.old_value = atomicCompSwap(a_u32.inner, 0u, value);
atomic_compare_result_2.exchanged = atomic_compare_result_2.old_value == 0u; atomic_compare_result_2.exchanged = atomic_compare_result_2.old_value == 0u;
atomic_compare_exchange_resultu32 r3 = atomic_compare_result_2; atomic_compare_exchange_result_u32 r3 = atomic_compare_result_2;
} }
{ {
int value = 42; int value = 42;
atomic_compare_exchange_resulti32 atomic_compare_result_3; atomic_compare_exchange_result_i32 atomic_compare_result_3;
atomic_compare_result_3.old_value = atomicCompSwap(a_i32.inner, 0, value); atomic_compare_result_3.old_value = atomicCompSwap(a_i32.inner, 0, value);
atomic_compare_result_3.exchanged = atomic_compare_result_3.old_value == 0; atomic_compare_result_3.exchanged = atomic_compare_result_3.old_value == 0;
atomic_compare_exchange_resulti32 r1 = atomic_compare_result_3; atomic_compare_exchange_result_i32 r1 = atomic_compare_result_3;
atomic_compare_exchange_resulti32 atomic_compare_result_4; atomic_compare_exchange_result_i32 atomic_compare_result_4;
atomic_compare_result_4.old_value = atomicCompSwap(a_i32.inner, 0, value); atomic_compare_result_4.old_value = atomicCompSwap(a_i32.inner, 0, value);
atomic_compare_result_4.exchanged = atomic_compare_result_4.old_value == 0; atomic_compare_result_4.exchanged = atomic_compare_result_4.old_value == 0;
atomic_compare_exchange_resulti32 r2 = atomic_compare_result_4; atomic_compare_exchange_result_i32 r2 = atomic_compare_result_4;
atomic_compare_exchange_resulti32 atomic_compare_result_5; atomic_compare_exchange_result_i32 atomic_compare_result_5;
atomic_compare_result_5.old_value = atomicCompSwap(a_i32.inner, 0, value); atomic_compare_result_5.old_value = atomicCompSwap(a_i32.inner, 0, value);
atomic_compare_result_5.exchanged = atomic_compare_result_5.old_value == 0; atomic_compare_result_5.exchanged = atomic_compare_result_5.old_value == 0;
atomic_compare_exchange_resulti32 r3 = atomic_compare_result_5; atomic_compare_exchange_result_i32 r3 = atomic_compare_result_5;
} }
{ {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 atomic_compare_result_6; atomic_compare_exchange_result_u32 atomic_compare_result_6;
atomic_compare_result_6.old_value = atomicCompSwap(b_u32, 0u, value); atomic_compare_result_6.old_value = atomicCompSwap(b_u32, 0u, value);
atomic_compare_result_6.exchanged = atomic_compare_result_6.old_value == 0u; atomic_compare_result_6.exchanged = atomic_compare_result_6.old_value == 0u;
atomic_compare_exchange_resultu32 r1 = atomic_compare_result_6; atomic_compare_exchange_result_u32 r1 = atomic_compare_result_6;
atomic_compare_exchange_resultu32 atomic_compare_result_7; atomic_compare_exchange_result_u32 atomic_compare_result_7;
atomic_compare_result_7.old_value = atomicCompSwap(b_u32, 0u, value); atomic_compare_result_7.old_value = atomicCompSwap(b_u32, 0u, value);
atomic_compare_result_7.exchanged = atomic_compare_result_7.old_value == 0u; atomic_compare_result_7.exchanged = atomic_compare_result_7.old_value == 0u;
atomic_compare_exchange_resultu32 r2 = atomic_compare_result_7; atomic_compare_exchange_result_u32 r2 = atomic_compare_result_7;
atomic_compare_exchange_resultu32 atomic_compare_result_8; atomic_compare_exchange_result_u32 atomic_compare_result_8;
atomic_compare_result_8.old_value = atomicCompSwap(b_u32, 0u, value); atomic_compare_result_8.old_value = atomicCompSwap(b_u32, 0u, value);
atomic_compare_result_8.exchanged = atomic_compare_result_8.old_value == 0u; atomic_compare_result_8.exchanged = atomic_compare_result_8.old_value == 0u;
atomic_compare_exchange_resultu32 r3 = atomic_compare_result_8; atomic_compare_exchange_result_u32 r3 = atomic_compare_result_8;
} }
{ {
int value = 42; int value = 42;
atomic_compare_exchange_resulti32 atomic_compare_result_9; atomic_compare_exchange_result_i32 atomic_compare_result_9;
atomic_compare_result_9.old_value = atomicCompSwap(b_i32, 0, value); atomic_compare_result_9.old_value = atomicCompSwap(b_i32, 0, value);
atomic_compare_result_9.exchanged = atomic_compare_result_9.old_value == 0; atomic_compare_result_9.exchanged = atomic_compare_result_9.old_value == 0;
atomic_compare_exchange_resulti32 r1 = atomic_compare_result_9; atomic_compare_exchange_result_i32 r1 = atomic_compare_result_9;
atomic_compare_exchange_resulti32 atomic_compare_result_10; atomic_compare_exchange_result_i32 atomic_compare_result_10;
atomic_compare_result_10.old_value = atomicCompSwap(b_i32, 0, value); atomic_compare_result_10.old_value = atomicCompSwap(b_i32, 0, value);
atomic_compare_result_10.exchanged = atomic_compare_result_10.old_value == 0; atomic_compare_result_10.exchanged = atomic_compare_result_10.old_value == 0;
atomic_compare_exchange_resulti32 r2 = atomic_compare_result_10; atomic_compare_exchange_result_i32 r2 = atomic_compare_result_10;
atomic_compare_exchange_resulti32 atomic_compare_result_11; atomic_compare_exchange_result_i32 atomic_compare_result_11;
atomic_compare_result_11.old_value = atomicCompSwap(b_i32, 0, value); atomic_compare_result_11.old_value = atomicCompSwap(b_i32, 0, value);
atomic_compare_result_11.exchanged = atomic_compare_result_11.old_value == 0; atomic_compare_result_11.exchanged = atomic_compare_result_11.old_value == 0;
atomic_compare_exchange_resulti32 r3 = atomic_compare_result_11; atomic_compare_exchange_result_i32 r3 = atomic_compare_result_11;
} }
} }

36
test/tint/bug/tint/1574.wgsl.expected.msl
View File

@ -2,33 +2,33 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
} }
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_2(device atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_2(device atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
} }
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_3(threadgroup atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_3(threadgroup atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
} }
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_4(threadgroup atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_4(threadgroup atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -42,27 +42,27 @@ void tint_symbol_inner(uint local_invocation_index, threadgroup atomic_uint* con
threadgroup_barrier(mem_flags::mem_threadgroup); threadgroup_barrier(mem_flags::mem_threadgroup);
{ {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 const r1 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); atomic_compare_exchange_result_u32 const r1 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value);
atomic_compare_exchange_resultu32 const r2 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); atomic_compare_exchange_result_u32 const r2 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value);
atomic_compare_exchange_resultu32 const r3 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); atomic_compare_exchange_result_u32 const r3 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value);
} }
{ {
int value = 42; int value = 42;
atomic_compare_exchange_resulti32 const r1 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); atomic_compare_exchange_result_i32 const r1 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value);
atomic_compare_exchange_resulti32 const r2 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); atomic_compare_exchange_result_i32 const r2 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value);
atomic_compare_exchange_resulti32 const r3 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); atomic_compare_exchange_result_i32 const r3 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value);
} }
{ {
uint value = 42u; uint value = 42u;
atomic_compare_exchange_resultu32 const r1 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); atomic_compare_exchange_result_u32 const r1 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value);
atomic_compare_exchange_resultu32 const r2 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); atomic_compare_exchange_result_u32 const r2 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value);
atomic_compare_exchange_resultu32 const r3 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); atomic_compare_exchange_result_u32 const r3 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value);
} }
{ {
int value = 42; int value = 42;
atomic_compare_exchange_resulti32 const r1 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); atomic_compare_exchange_result_i32 const r1 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value);
atomic_compare_exchange_resulti32 const r2 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); atomic_compare_exchange_result_i32 const r2 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value);
atomic_compare_exchange_resulti32 const r3 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); atomic_compare_exchange_result_i32 const r3 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value);
} }
} }

48
test/tint/bug/tint/1574.wgsl.expected.spvasm
View File

@ -19,13 +19,13 @@
OpName %main_inner "main_inner" OpName %main_inner "main_inner"
OpName %local_invocation_index "local_invocation_index" OpName %local_invocation_index "local_invocation_index"
OpName %value "value" OpName %value "value"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %value_0 "value" OpName %value_0 "value"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %value_1 "value" OpName %value_1 "value"
OpName %value_2 "value" OpName %value_2 "value"
OpName %main "main" OpName %main "main"
@ -38,10 +38,10 @@
OpMemberDecorate %a_i32_block 0 Offset 0 OpMemberDecorate %a_i32_block 0 Offset 0
OpDecorate %a_i32 DescriptorSet 0 OpDecorate %a_i32 DescriptorSet 0
OpDecorate %a_i32 Binding 1 OpDecorate %a_i32 Binding 1
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
@ -67,11 +67,11 @@
%uint_264 = OpConstant %uint 264 %uint_264 = OpConstant %uint 264
%uint_42 = OpConstant %uint 42 %uint_42 = OpConstant %uint 42
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%int_42 = OpConstant %int 42 %int_42 = OpConstant %int 42
%_ptr_Function_int = OpTypePointer Function %int %_ptr_Function_int = OpTypePointer Function %int
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%113 = OpTypeFunction %void %113 = OpTypeFunction %void
%main_inner = OpFunction %void None %15 %main_inner = OpFunction %void None %15
@ -95,59 +95,59 @@
%43 = OpLoad %uint %value %43 = OpLoad %uint %value
%44 = OpAtomicCompareExchange %uint %42 %uint_1 %uint_0 %uint_0 %43 %29 %44 = OpAtomicCompareExchange %uint %42 %uint_1 %uint_0 %uint_0 %43 %29
%45 = OpIEqual %bool %44 %29 %45 = OpIEqual %bool %44 %29
%38 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %44 %45 %38 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %44 %45
%48 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0 %48 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0
%49 = OpLoad %uint %value %49 = OpLoad %uint %value
%50 = OpAtomicCompareExchange %uint %48 %uint_1 %uint_0 %uint_0 %49 %29 %50 = OpAtomicCompareExchange %uint %48 %uint_1 %uint_0 %uint_0 %49 %29
%51 = OpIEqual %bool %50 %29 %51 = OpIEqual %bool %50 %29
%46 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %50 %51 %46 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %50 %51
%54 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0 %54 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0
%55 = OpLoad %uint %value %55 = OpLoad %uint %value
%56 = OpAtomicCompareExchange %uint %54 %uint_1 %uint_0 %uint_0 %55 %29 %56 = OpAtomicCompareExchange %uint %54 %uint_1 %uint_0 %uint_0 %55 %29
%57 = OpIEqual %bool %56 %29 %57 = OpIEqual %bool %56 %29
%52 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %56 %57 %52 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %56 %57
OpStore %value_0 %int_42 OpStore %value_0 %int_42
%65 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0 %65 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0
%66 = OpLoad %int %value_0 %66 = OpLoad %int %value_0
%67 = OpAtomicCompareExchange %int %65 %uint_1 %uint_0 %uint_0 %66 %32 %67 = OpAtomicCompareExchange %int %65 %uint_1 %uint_0 %uint_0 %66 %32
%68 = OpIEqual %bool %67 %32 %68 = OpIEqual %bool %67 %32
%61 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %67 %68 %61 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %67 %68
%71 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0 %71 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0
%72 = OpLoad %int %value_0 %72 = OpLoad %int %value_0
%73 = OpAtomicCompareExchange %int %71 %uint_1 %uint_0 %uint_0 %72 %32 %73 = OpAtomicCompareExchange %int %71 %uint_1 %uint_0 %uint_0 %72 %32
%74 = OpIEqual %bool %73 %32 %74 = OpIEqual %bool %73 %32
%69 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %73 %74 %69 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %73 %74
%77 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0 %77 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0
%78 = OpLoad %int %value_0 %78 = OpLoad %int %value_0
%79 = OpAtomicCompareExchange %int %77 %uint_1 %uint_0 %uint_0 %78 %32 %79 = OpAtomicCompareExchange %int %77 %uint_1 %uint_0 %uint_0 %78 %32
%80 = OpIEqual %bool %79 %32 %80 = OpIEqual %bool %79 %32
%75 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %79 %80 %75 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %79 %80
OpStore %value_1 %uint_42 OpStore %value_1 %uint_42
%84 = OpLoad %uint %value_1 %84 = OpLoad %uint %value_1
%85 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %84 %29 %85 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %84 %29
%86 = OpIEqual %bool %85 %29 %86 = OpIEqual %bool %85 %29
%82 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %85 %86 %82 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %85 %86
%89 = OpLoad %uint %value_1 %89 = OpLoad %uint %value_1
%90 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %89 %29 %90 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %89 %29
%91 = OpIEqual %bool %90 %29 %91 = OpIEqual %bool %90 %29
%87 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %90 %91 %87 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %90 %91
%94 = OpLoad %uint %value_1 %94 = OpLoad %uint %value_1
%95 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %94 %29 %95 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %94 %29
%96 = OpIEqual %bool %95 %29 %96 = OpIEqual %bool %95 %29
%92 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %95 %96 %92 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %95 %96
OpStore %value_2 %int_42 OpStore %value_2 %int_42
%100 = OpLoad %int %value_2 %100 = OpLoad %int %value_2
%101 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %100 %32 %101 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %100 %32
%102 = OpIEqual %bool %101 %32 %102 = OpIEqual %bool %101 %32
%98 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %101 %102 %98 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %101 %102
%105 = OpLoad %int %value_2 %105 = OpLoad %int %value_2
%106 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %105 %32 %106 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %105 %32
%107 = OpIEqual %bool %106 %32 %107 = OpIEqual %bool %106 %32
%103 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %106 %107 %103 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %106 %107
%110 = OpLoad %int %value_2 %110 = OpLoad %int %value_2
%111 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %110 %32 %111 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %110 %32
%112 = OpIEqual %bool %111 %32 %112 = OpIEqual %bool %111 %32
%108 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %111 %112 %108 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %111 %112
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%main = OpFunction %void None %113 %main = OpFunction %void None %113

12
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -26,10 +26,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_1bd40a() { void atomicCompareExchangeWeak_1bd40a() {
x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32(0, false); x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32(0, false);
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
atomic_compare_exchange_resulti32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_i32 tint_symbol = atomic_compare_result;
int old_value_1 = tint_symbol.old_value; int old_value_1 = tint_symbol.old_value;
int x_19 = old_value_1; int x_19 = old_value_1;
x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_19, (x_19 == 1)); x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_19, (x_19 == 1));
@ -52,7 +52,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -77,10 +77,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_1bd40a() { void atomicCompareExchangeWeak_1bd40a() {
x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32(0, false); x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32(0, false);
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
atomic_compare_exchange_resulti32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_i32 tint_symbol = atomic_compare_result;
int old_value_1 = tint_symbol.old_value; int old_value_1 = tint_symbol.old_value;
int x_19 = old_value_1; int x_19 = old_value_1;
x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_19, (x_19 == 1)); x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_19, (x_19 == 1));

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -27,7 +27,7 @@ struct x__atomic_compare_exchange_resulti32 {
void atomicCompareExchangeWeak_1bd40a(device SB_RW_atomic* const tint_symbol_2) { void atomicCompareExchangeWeak_1bd40a(device SB_RW_atomic* const tint_symbol_2) {
x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32{}; x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32{};
atomic_compare_exchange_resulti32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), 1, 1); atomic_compare_exchange_result_i32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), 1, 1);
int const old_value_1 = tint_symbol.old_value; int const old_value_1 = tint_symbol.old_value;
int const x_19 = old_value_1; int const x_19 = old_value_1;
x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_19, .exchanged=(x_19 == 1)}; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_19, .exchanged=(x_19 == 1)};

14
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm
View File

@ -19,9 +19,9 @@
OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main_1 "fragment_main_1"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -33,8 +33,8 @@
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%SB_RW_atomic = OpTypeStruct %int %SB_RW_atomic = OpTypeStruct %int
%sb_rw_block = OpTypeStruct %SB_RW_atomic %sb_rw_block = OpTypeStruct %SB_RW_atomic
@ -46,7 +46,7 @@
%x__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %x__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool
%12 = OpConstantNull %x__atomic_compare_exchange_resulti32 %12 = OpConstantNull %x__atomic_compare_exchange_resulti32
%_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32 %_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
@ -59,7 +59,7 @@
%22 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0 %22 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0
%24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %int_1 %int_1 %24 = OpAtomicCompareExchange %int %22 %uint_1 %uint_0 %uint_0 %int_1 %int_1
%25 = OpIEqual %bool %24 %int_1 %25 = OpIEqual %bool %24 %int_1
%15 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25 %15 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %24 %25
%26 = OpCompositeExtract %int %15 0 %26 = OpCompositeExtract %int %15 0
%27 = OpIEqual %bool %26 %int_1 %27 = OpIEqual %bool %26 %int_1
%28 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %26 %27 %28 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %26 %27

12
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -26,10 +26,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_63d8e6() { void atomicCompareExchangeWeak_63d8e6() {
x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32(0u, false); x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32(0u, false);
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
atomic_compare_exchange_resultu32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_u32 tint_symbol = atomic_compare_result;
uint old_value_1 = tint_symbol.old_value; uint old_value_1 = tint_symbol.old_value;
uint x_17 = old_value_1; uint x_17 = old_value_1;
x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u)); x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u));
@ -52,7 +52,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -77,10 +77,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_63d8e6() { void atomicCompareExchangeWeak_63d8e6() {
x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32(0u, false); x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32(0u, false);
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
atomic_compare_exchange_resultu32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_u32 tint_symbol = atomic_compare_result;
uint old_value_1 = tint_symbol.old_value; uint old_value_1 = tint_symbol.old_value;
uint x_17 = old_value_1; uint x_17 = old_value_1;
x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u)); x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u));

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -27,7 +27,7 @@ struct x__atomic_compare_exchange_resultu32 {
void atomicCompareExchangeWeak_63d8e6(device SB_RW_atomic* const tint_symbol_2) { void atomicCompareExchangeWeak_63d8e6(device SB_RW_atomic* const tint_symbol_2) {
x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32{}; x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32{};
atomic_compare_exchange_resultu32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), 1u, 1u); atomic_compare_exchange_result_u32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), 1u, 1u);
uint const old_value_1 = tint_symbol.old_value; uint const old_value_1 = tint_symbol.old_value;
uint const x_17 = old_value_1; uint const x_17 = old_value_1;
x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_17, .exchanged=(x_17 == 1u)}; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_17, .exchanged=(x_17 == 1u)};

14
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm
View File

@ -19,9 +19,9 @@
OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main_1 "fragment_main_1"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -33,8 +33,8 @@
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%SB_RW_atomic = OpTypeStruct %uint %SB_RW_atomic = OpTypeStruct %uint
%sb_rw_block = OpTypeStruct %SB_RW_atomic %sb_rw_block = OpTypeStruct %SB_RW_atomic
@ -46,7 +46,7 @@
%x__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %x__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool
%12 = OpConstantNull %x__atomic_compare_exchange_resultu32 %12 = OpConstantNull %x__atomic_compare_exchange_resultu32
%_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32 %_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
@ -57,7 +57,7 @@
%21 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0 %21 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0
%22 = OpAtomicCompareExchange %uint %21 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1 %22 = OpAtomicCompareExchange %uint %21 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1
%23 = OpIEqual %bool %22 %uint_1 %23 = OpIEqual %bool %22 %uint_1
%15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23 %15 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %22 %23
%24 = OpCompositeExtract %uint %15 0 %24 = OpCompositeExtract %uint %15 0
%25 = OpIEqual %bool %24 %uint_1 %25 = OpIEqual %bool %24 %uint_1
%26 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %24 %25 %26 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %24 %25

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -12,11 +12,11 @@ groupshared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
x__atomic_compare_exchange_resulti32 res = (x__atomic_compare_exchange_resulti32)0; x__atomic_compare_exchange_resulti32 res = (x__atomic_compare_exchange_resulti32)0;
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = 1; int atomic_compare_value = 1;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resulti32 tint_symbol = atomic_result; const atomic_compare_exchange_result_i32 tint_symbol = atomic_result;
const int old_value_1 = tint_symbol.old_value; const int old_value_1 = tint_symbol.old_value;
const int x_18 = old_value_1; const int x_18 = old_value_1;
const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)}; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)};

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -12,11 +12,11 @@ groupshared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
x__atomic_compare_exchange_resulti32 res = (x__atomic_compare_exchange_resulti32)0; x__atomic_compare_exchange_resulti32 res = (x__atomic_compare_exchange_resulti32)0;
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = 1; int atomic_compare_value = 1;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resulti32 tint_symbol = atomic_result; const atomic_compare_exchange_result_i32 tint_symbol = atomic_result;
const int old_value_1 = tint_symbol.old_value; const int old_value_1 = tint_symbol.old_value;
const int x_18 = old_value_1; const int x_18 = old_value_1;
const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)}; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)};

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -15,10 +15,10 @@ uint local_invocation_index_1 = 0u;
shared int arg_0; shared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32(0, false); x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32(0, false);
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, 1, 1); atomic_compare_result.old_value = atomicCompSwap(arg_0, 1, 1);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
atomic_compare_exchange_resulti32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_i32 tint_symbol = atomic_compare_result;
int old_value_1 = tint_symbol.old_value; int old_value_1 = tint_symbol.old_value;
int x_18 = old_value_1; int x_18 = old_value_1;
x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_18, (x_18 == 1)); x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_18, (x_18 == 1));

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -23,7 +23,7 @@ struct x__atomic_compare_exchange_resulti32 {
void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol_2) { void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol_2) {
x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32{}; x__atomic_compare_exchange_resulti32 res = x__atomic_compare_exchange_resulti32{};
atomic_compare_exchange_resulti32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, 1, 1); atomic_compare_exchange_result_i32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, 1, 1);
int const old_value_1 = tint_symbol.old_value; int const old_value_1 = tint_symbol.old_value;
int const x_18 = old_value_1; int const x_18 = old_value_1;
x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_18, .exchanged=(x_18 == 1)}; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_18, .exchanged=(x_18 == 1)};

14
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm
View File

@ -15,9 +15,9 @@
OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %compute_main_inner "compute_main_inner" OpName %compute_main_inner "compute_main_inner"
OpName %local_invocation_index_2 "local_invocation_index_2" OpName %local_invocation_index_2 "local_invocation_index_2"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -27,8 +27,8 @@
OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
@ -44,7 +44,7 @@
%x__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %x__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool
%16 = OpConstantNull %x__atomic_compare_exchange_resulti32 %16 = OpConstantNull %x__atomic_compare_exchange_resulti32
%_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32 %_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
@ -57,7 +57,7 @@
OpStore %res %16 OpStore %res %16
%25 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1 %25 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1
%26 = OpIEqual %bool %25 %int_1 %26 = OpIEqual %bool %25 %int_1
%19 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %25 %26 %19 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %25 %26
%27 = OpCompositeExtract %int %19 0 %27 = OpCompositeExtract %int %19 0
%28 = OpIEqual %bool %27 %int_1 %28 = OpIEqual %bool %27 %int_1
%29 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %27 %28 %29 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %27 %28

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -12,11 +12,11 @@ groupshared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
x__atomic_compare_exchange_resultu32 res = (x__atomic_compare_exchange_resultu32)0; x__atomic_compare_exchange_resultu32 res = (x__atomic_compare_exchange_resultu32)0;
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = 1u; uint atomic_compare_value = 1u;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resultu32 tint_symbol = atomic_result; const atomic_compare_exchange_result_u32 tint_symbol = atomic_result;
const uint old_value_1 = tint_symbol.old_value; const uint old_value_1 = tint_symbol.old_value;
const uint x_17 = old_value_1; const uint x_17 = old_value_1;
const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)}; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)};

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -12,11 +12,11 @@ groupshared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
x__atomic_compare_exchange_resultu32 res = (x__atomic_compare_exchange_resultu32)0; x__atomic_compare_exchange_resultu32 res = (x__atomic_compare_exchange_resultu32)0;
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = 1u; uint atomic_compare_value = 1u;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resultu32 tint_symbol = atomic_result; const atomic_compare_exchange_result_u32 tint_symbol = atomic_result;
const uint old_value_1 = tint_symbol.old_value; const uint old_value_1 = tint_symbol.old_value;
const uint x_17 = old_value_1; const uint x_17 = old_value_1;
const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)}; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)};

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -15,10 +15,10 @@ uint local_invocation_index_1 = 0u;
shared uint arg_0; shared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32(0u, false); x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32(0u, false);
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, 1u, 1u); atomic_compare_result.old_value = atomicCompSwap(arg_0, 1u, 1u);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
atomic_compare_exchange_resultu32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_u32 tint_symbol = atomic_compare_result;
uint old_value_1 = tint_symbol.old_value; uint old_value_1 = tint_symbol.old_value;
uint x_17 = old_value_1; uint x_17 = old_value_1;
x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u)); x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u));

6
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -23,7 +23,7 @@ struct x__atomic_compare_exchange_resultu32 {
void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol_2) { void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol_2) {
x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32{}; x__atomic_compare_exchange_resultu32 res = x__atomic_compare_exchange_resultu32{};
atomic_compare_exchange_resultu32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, 1u, 1u); atomic_compare_exchange_result_u32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, 1u, 1u);
uint const old_value_1 = tint_symbol.old_value; uint const old_value_1 = tint_symbol.old_value;
uint const x_17 = old_value_1; uint const x_17 = old_value_1;
x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_17, .exchanged=(x_17 == 1u)}; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_17, .exchanged=(x_17 == 1u)};

14
test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm
View File

@ -15,9 +15,9 @@
OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %compute_main_inner "compute_main_inner" OpName %compute_main_inner "compute_main_inner"
OpName %local_invocation_index_2 "local_invocation_index_2" OpName %local_invocation_index_2 "local_invocation_index_2"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -27,8 +27,8 @@
OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
@ -43,7 +43,7 @@
%x__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %x__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool
%15 = OpConstantNull %x__atomic_compare_exchange_resultu32 %15 = OpConstantNull %x__atomic_compare_exchange_resultu32
%_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32 %_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
@ -55,7 +55,7 @@
OpStore %res %15 OpStore %res %15
%24 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1 %24 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1
%25 = OpIEqual %bool %24 %uint_1 %25 = OpIEqual %bool %24 %uint_1
%18 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %24 %25 %18 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %24 %25
%26 = OpCompositeExtract %uint %18 0 %26 = OpCompositeExtract %uint %18 0
%27 = OpIEqual %bool %26 %uint_1 %27 = OpIEqual %bool %26 %uint_1
%28 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %26 %27 %28 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %26 %27

12
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -32,10 +32,10 @@ void atomicCompareExchangeWeak_1bd40a() {
arg_2 = 1; arg_2 = 1;
int x_23 = arg_2; int x_23 = arg_2;
int x_24 = arg_1; int x_24 = arg_1;
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_24, x_23); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_24, x_23);
atomic_compare_result.exchanged = atomic_compare_result.old_value == x_24; atomic_compare_result.exchanged = atomic_compare_result.old_value == x_24;
atomic_compare_exchange_resulti32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_i32 tint_symbol = atomic_compare_result;
int old_value_1 = tint_symbol.old_value; int old_value_1 = tint_symbol.old_value;
int x_25 = old_value_1; int x_25 = old_value_1;
x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_25, (x_25 == x_23)); x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_25, (x_25 == x_23));
@ -58,7 +58,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -89,10 +89,10 @@ void atomicCompareExchangeWeak_1bd40a() {
arg_2 = 1; arg_2 = 1;
int x_23 = arg_2; int x_23 = arg_2;
int x_24 = arg_1; int x_24 = arg_1;
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_24, x_23); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_24, x_23);
atomic_compare_result.exchanged = atomic_compare_result.old_value == x_24; atomic_compare_result.exchanged = atomic_compare_result.old_value == x_24;
atomic_compare_exchange_resulti32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_i32 tint_symbol = atomic_compare_result;
int old_value_1 = tint_symbol.old_value; int old_value_1 = tint_symbol.old_value;
int x_25 = old_value_1; int x_25 = old_value_1;
x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_25, (x_25 == x_23)); x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_25, (x_25 == x_23));

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -33,7 +33,7 @@ void atomicCompareExchangeWeak_1bd40a(device SB_RW_atomic* const tint_symbol_2)
arg_2 = 1; arg_2 = 1;
int const x_23 = arg_2; int const x_23 = arg_2;
int const x_24 = arg_1; int const x_24 = arg_1;
atomic_compare_exchange_resulti32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), x_24, x_23); atomic_compare_exchange_result_i32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), x_24, x_23);
int const old_value_1 = tint_symbol.old_value; int const old_value_1 = tint_symbol.old_value;
int const x_25 = old_value_1; int const x_25 = old_value_1;
x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_25, .exchanged=(x_25 == x_23)}; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_25, .exchanged=(x_25 == x_23)};

14
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm
View File

@ -21,9 +21,9 @@
OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main_1 "fragment_main_1"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -35,8 +35,8 @@
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%SB_RW_atomic = OpTypeStruct %int %SB_RW_atomic = OpTypeStruct %int
%sb_rw_block = OpTypeStruct %SB_RW_atomic %sb_rw_block = OpTypeStruct %SB_RW_atomic
@ -51,7 +51,7 @@
%16 = OpConstantNull %x__atomic_compare_exchange_resulti32 %16 = OpConstantNull %x__atomic_compare_exchange_resulti32
%_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32 %_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
@ -71,7 +71,7 @@
%29 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0 %29 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0
%30 = OpAtomicCompareExchange %int %29 %uint_1 %uint_0 %uint_0 %20 %21 %30 = OpAtomicCompareExchange %int %29 %uint_1 %uint_0 %uint_0 %20 %21
%31 = OpIEqual %bool %30 %21 %31 = OpIEqual %bool %30 %21
%22 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %30 %31 %22 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %30 %31
%32 = OpCompositeExtract %int %22 0 %32 = OpCompositeExtract %int %22 0
%33 = OpIEqual %bool %32 %20 %33 = OpIEqual %bool %32 %20
%34 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %32 %33 %34 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %32 %33

12
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -32,10 +32,10 @@ void atomicCompareExchangeWeak_63d8e6() {
arg_2 = 1u; arg_2 = 1u;
uint x_21 = arg_2; uint x_21 = arg_2;
uint x_22 = arg_1; uint x_22 = arg_1;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_22, x_21); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_22, x_21);
atomic_compare_result.exchanged = atomic_compare_result.old_value == x_22; atomic_compare_result.exchanged = atomic_compare_result.old_value == x_22;
atomic_compare_exchange_resultu32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_u32 tint_symbol = atomic_compare_result;
uint old_value_1 = tint_symbol.old_value; uint old_value_1 = tint_symbol.old_value;
uint x_23 = old_value_1; uint x_23 = old_value_1;
x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21)); x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21));
@ -58,7 +58,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -89,10 +89,10 @@ void atomicCompareExchangeWeak_63d8e6() {
arg_2 = 1u; arg_2 = 1u;
uint x_21 = arg_2; uint x_21 = arg_2;
uint x_22 = arg_1; uint x_22 = arg_1;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_22, x_21); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, x_22, x_21);
atomic_compare_result.exchanged = atomic_compare_result.old_value == x_22; atomic_compare_result.exchanged = atomic_compare_result.old_value == x_22;
atomic_compare_exchange_resultu32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_u32 tint_symbol = atomic_compare_result;
uint old_value_1 = tint_symbol.old_value; uint old_value_1 = tint_symbol.old_value;
uint x_23 = old_value_1; uint x_23 = old_value_1;
x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21)); x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21));

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -33,7 +33,7 @@ void atomicCompareExchangeWeak_63d8e6(device SB_RW_atomic* const tint_symbol_2)
arg_2 = 1u; arg_2 = 1u;
uint const x_21 = arg_2; uint const x_21 = arg_2;
uint const x_22 = arg_1; uint const x_22 = arg_1;
atomic_compare_exchange_resultu32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), x_22, x_21); atomic_compare_exchange_result_u32 const tint_symbol = atomicCompareExchangeWeak_1(&((*(tint_symbol_2)).arg_0), x_22, x_21);
uint const old_value_1 = tint_symbol.old_value; uint const old_value_1 = tint_symbol.old_value;
uint const x_23 = old_value_1; uint const x_23 = old_value_1;
x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_23, .exchanged=(x_23 == x_21)}; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_23, .exchanged=(x_23 == x_21)};

14
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm
View File

@ -21,9 +21,9 @@
OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main_1 "fragment_main_1"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -35,8 +35,8 @@
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%SB_RW_atomic = OpTypeStruct %uint %SB_RW_atomic = OpTypeStruct %uint
%sb_rw_block = OpTypeStruct %SB_RW_atomic %sb_rw_block = OpTypeStruct %SB_RW_atomic
@ -51,7 +51,7 @@
%16 = OpConstantNull %x__atomic_compare_exchange_resultu32 %16 = OpConstantNull %x__atomic_compare_exchange_resultu32
%_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32 %_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6 %atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6
@ -69,7 +69,7 @@
%27 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0 %27 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0
%28 = OpAtomicCompareExchange %uint %27 %uint_1 %uint_0 %uint_0 %20 %21 %28 = OpAtomicCompareExchange %uint %27 %uint_1 %uint_0 %uint_0 %20 %21
%29 = OpIEqual %bool %28 %21 %29 = OpIEqual %bool %28 %21
%22 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %28 %29 %22 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %28 %29
%30 = OpCompositeExtract %uint %22 0 %30 = OpCompositeExtract %uint %22 0
%31 = OpIEqual %bool %30 %20 %31 = OpIEqual %bool %30 %20
%32 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %30 %31 %32 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %30 %31

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_e88938() {
arg_2 = 1; arg_2 = 1;
const int x_22 = arg_2; const int x_22 = arg_2;
const int x_23 = arg_1; const int x_23 = arg_1;
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = x_23; int atomic_compare_value = x_23;
InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resulti32 tint_symbol = atomic_result; const atomic_compare_exchange_result_i32 tint_symbol = atomic_result;
const int old_value_1 = tint_symbol.old_value; const int old_value_1 = tint_symbol.old_value;
const int x_24 = old_value_1; const int x_24 = old_value_1;
const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)}; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)};

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_e88938() {
arg_2 = 1; arg_2 = 1;
const int x_22 = arg_2; const int x_22 = arg_2;
const int x_23 = arg_1; const int x_23 = arg_1;
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = x_23; int atomic_compare_value = x_23;
InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resulti32 tint_symbol = atomic_result; const atomic_compare_exchange_result_i32 tint_symbol = atomic_result;
const int old_value_1 = tint_symbol.old_value; const int old_value_1 = tint_symbol.old_value;
const int x_24 = old_value_1; const int x_24 = old_value_1;
const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)}; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)};

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -21,10 +21,10 @@ void atomicCompareExchangeWeak_e88938() {
arg_2 = 1; arg_2 = 1;
int x_22 = arg_2; int x_22 = arg_2;
int x_23 = arg_1; int x_23 = arg_1;
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, x_23, x_22); atomic_compare_result.old_value = atomicCompSwap(arg_0, x_23, x_22);
atomic_compare_result.exchanged = atomic_compare_result.old_value == x_23; atomic_compare_result.exchanged = atomic_compare_result.old_value == x_23;
atomic_compare_exchange_resulti32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_i32 tint_symbol = atomic_compare_result;
int old_value_1 = tint_symbol.old_value; int old_value_1 = tint_symbol.old_value;
int x_24 = old_value_1; int x_24 = old_value_1;
x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_24, (x_24 == x_22)); x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_24, (x_24 == x_22));

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -29,7 +29,7 @@ void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol_
arg_2 = 1; arg_2 = 1;
int const x_22 = arg_2; int const x_22 = arg_2;
int const x_23 = arg_1; int const x_23 = arg_1;
atomic_compare_exchange_resulti32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, x_23, x_22); atomic_compare_exchange_result_i32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, x_23, x_22);
int const old_value_1 = tint_symbol.old_value; int const old_value_1 = tint_symbol.old_value;
int const x_24 = old_value_1; int const x_24 = old_value_1;
x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_24, .exchanged=(x_24 == x_22)}; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_24, .exchanged=(x_24 == x_22)};

14
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm
View File

@ -17,9 +17,9 @@
OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %compute_main_inner "compute_main_inner" OpName %compute_main_inner "compute_main_inner"
OpName %local_invocation_index_2 "local_invocation_index_2" OpName %local_invocation_index_2 "local_invocation_index_2"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -29,8 +29,8 @@
OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
@ -49,7 +49,7 @@
%20 = OpConstantNull %x__atomic_compare_exchange_resulti32 %20 = OpConstantNull %x__atomic_compare_exchange_resulti32
%_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32 %_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%36 = OpTypeFunction %void %uint %36 = OpTypeFunction %void %uint
@ -68,7 +68,7 @@
%25 = OpLoad %int %arg_1 %25 = OpLoad %int %arg_1
%31 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %24 %25 %31 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %24 %25
%32 = OpIEqual %bool %31 %25 %32 = OpIEqual %bool %31 %25
%26 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %31 %32 %26 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %31 %32
%33 = OpCompositeExtract %int %26 0 %33 = OpCompositeExtract %int %26 0
%34 = OpIEqual %bool %33 %24 %34 = OpIEqual %bool %33 %24
%35 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %33 %34 %35 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %33 %34

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_83580d() {
arg_2 = 1u; arg_2 = 1u;
const uint x_21 = arg_2; const uint x_21 = arg_2;
const uint x_22 = arg_1; const uint x_22 = arg_1;
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = x_22; uint atomic_compare_value = x_22;
InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resultu32 tint_symbol = atomic_result; const atomic_compare_exchange_result_u32 tint_symbol = atomic_result;
const uint old_value_1 = tint_symbol.old_value; const uint old_value_1 = tint_symbol.old_value;
const uint x_23 = old_value_1; const uint x_23 = old_value_1;
const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)}; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)};

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_83580d() {
arg_2 = 1u; arg_2 = 1u;
const uint x_21 = arg_2; const uint x_21 = arg_2;
const uint x_22 = arg_1; const uint x_22 = arg_1;
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = x_22; uint atomic_compare_value = x_22;
InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
const atomic_compare_exchange_resultu32 tint_symbol = atomic_result; const atomic_compare_exchange_result_u32 tint_symbol = atomic_result;
const uint old_value_1 = tint_symbol.old_value; const uint old_value_1 = tint_symbol.old_value;
const uint x_23 = old_value_1; const uint x_23 = old_value_1;
const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)}; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)};

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -21,10 +21,10 @@ void atomicCompareExchangeWeak_83580d() {
arg_2 = 1u; arg_2 = 1u;
uint x_21 = arg_2; uint x_21 = arg_2;
uint x_22 = arg_1; uint x_22 = arg_1;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, x_22, x_21); atomic_compare_result.old_value = atomicCompSwap(arg_0, x_22, x_21);
atomic_compare_result.exchanged = atomic_compare_result.old_value == x_22; atomic_compare_result.exchanged = atomic_compare_result.old_value == x_22;
atomic_compare_exchange_resultu32 tint_symbol = atomic_compare_result; atomic_compare_exchange_result_u32 tint_symbol = atomic_compare_result;
uint old_value_1 = tint_symbol.old_value; uint old_value_1 = tint_symbol.old_value;
uint x_23 = old_value_1; uint x_23 = old_value_1;
x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21)); x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21));

6
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -29,7 +29,7 @@ void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol
arg_2 = 1u; arg_2 = 1u;
uint const x_21 = arg_2; uint const x_21 = arg_2;
uint const x_22 = arg_1; uint const x_22 = arg_1;
atomic_compare_exchange_resultu32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, x_22, x_21); atomic_compare_exchange_result_u32 const tint_symbol = atomicCompareExchangeWeak_1(tint_symbol_2, x_22, x_21);
uint const old_value_1 = tint_symbol.old_value; uint const old_value_1 = tint_symbol.old_value;
uint const x_23 = old_value_1; uint const x_23 = old_value_1;
x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_23, .exchanged=(x_23 == x_21)}; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_23, .exchanged=(x_23 == x_21)};

14
test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm
View File

@ -17,9 +17,9 @@
OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value"
OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %compute_main_inner "compute_main_inner" OpName %compute_main_inner "compute_main_inner"
OpName %local_invocation_index_2 "local_invocation_index_2" OpName %local_invocation_index_2 "local_invocation_index_2"
OpName %compute_main_1 "compute_main_1" OpName %compute_main_1 "compute_main_1"
@ -29,8 +29,8 @@
OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex
OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0
OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input
@ -47,7 +47,7 @@
%18 = OpConstantNull %x__atomic_compare_exchange_resultu32 %18 = OpConstantNull %x__atomic_compare_exchange_resultu32
%_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32 %_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%34 = OpTypeFunction %void %uint %34 = OpTypeFunction %void %uint
@ -66,7 +66,7 @@
%23 = OpLoad %uint %arg_1 %23 = OpLoad %uint %arg_1
%29 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %22 %23 %29 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %22 %23
%30 = OpIEqual %bool %29 %23 %30 = OpIEqual %bool %29 %23
%24 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %29 %30 %24 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %29 %30
%31 = OpCompositeExtract %uint %24 0 %31 = OpCompositeExtract %uint %24 0
%32 = OpIEqual %bool %31 %22 %32 = OpIEqual %bool %31 %22
%33 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %31 %32 %33 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %31 %32

12
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -16,10 +16,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
} sb_rw; } sb_rw;
void atomicCompareExchangeWeak_1bd40a() { void atomicCompareExchangeWeak_1bd40a() {
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
atomic_compare_exchange_resulti32 res = atomic_compare_result; atomic_compare_exchange_result_i32 res = atomic_compare_result;
} }
void fragment_main() { void fragment_main() {
@ -32,7 +32,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -47,10 +47,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
} sb_rw; } sb_rw;
void atomicCompareExchangeWeak_1bd40a() { void atomicCompareExchangeWeak_1bd40a() {
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1, 1);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
atomic_compare_exchange_resulti32 res = atomic_compare_result; atomic_compare_exchange_result_i32 res = atomic_compare_result;
} }
void compute_main() { void compute_main() {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -17,7 +17,7 @@ struct SB_RW {
}; };
void atomicCompareExchangeWeak_1bd40a(device SB_RW* const tint_symbol) { void atomicCompareExchangeWeak_1bd40a(device SB_RW* const tint_symbol) {
atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1, 1); atomic_compare_exchange_result_i32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1, 1);
} }
fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {

20
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
View File

@ -15,9 +15,9 @@
OpMemberName %SB_RW 0 "arg_0" OpMemberName %SB_RW 0 "arg_0"
OpName %sb_rw "sb_rw" OpName %sb_rw "sb_rw"
OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a" OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main" OpName %compute_main "compute_main"
@ -26,8 +26,8 @@
OpMemberDecorate %SB_RW 0 Offset 0 OpMemberDecorate %SB_RW 0 Offset 0
OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw DescriptorSet 0
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%SB_RW = OpTypeStruct %int %SB_RW = OpTypeStruct %int
%sb_rw_block = OpTypeStruct %SB_RW %sb_rw_block = OpTypeStruct %SB_RW
@ -36,21 +36,21 @@
%void = OpTypeVoid %void = OpTypeVoid
%6 = OpTypeFunction %void %6 = OpTypeFunction %void
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 %_ptr_Function___atomic_compare_exchange_result_i32 = OpTypePointer Function %__atomic_compare_exchange_result_i32
%24 = OpConstantNull %__atomic_compare_exchange_resulti32 %24 = OpConstantNull %__atomic_compare_exchange_result_i32
%atomicCompareExchangeWeak_1bd40a = OpFunction %void None %6 %atomicCompareExchangeWeak_1bd40a = OpFunction %void None %6
%9 = OpLabel %9 = OpLabel
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %24 %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_i32 Function %24
%18 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0 %18 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0
%20 = OpAtomicCompareExchange %int %18 %uint_1 %uint_0 %uint_0 %int_1 %int_1 %20 = OpAtomicCompareExchange %int %18 %uint_1 %uint_0 %uint_0 %int_1 %int_1
%21 = OpIEqual %bool %20 %int_1 %21 = OpIEqual %bool %20 %int_1
%10 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %20 %21 %10 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %20 %21
OpStore %res %10 OpStore %res %10
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

12
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -16,10 +16,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
} sb_rw; } sb_rw;
void atomicCompareExchangeWeak_63d8e6() { void atomicCompareExchangeWeak_63d8e6() {
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
atomic_compare_exchange_resultu32 res = atomic_compare_result; atomic_compare_exchange_result_u32 res = atomic_compare_result;
} }
void fragment_main() { void fragment_main() {
@ -32,7 +32,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -47,10 +47,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
} sb_rw; } sb_rw;
void atomicCompareExchangeWeak_63d8e6() { void atomicCompareExchangeWeak_63d8e6() {
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, 1u, 1u);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
atomic_compare_exchange_resultu32 res = atomic_compare_result; atomic_compare_exchange_result_u32 res = atomic_compare_result;
} }
void compute_main() { void compute_main() {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -17,7 +17,7 @@ struct SB_RW {
}; };
void atomicCompareExchangeWeak_63d8e6(device SB_RW* const tint_symbol) { void atomicCompareExchangeWeak_63d8e6(device SB_RW* const tint_symbol) {
atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1u, 1u); atomic_compare_exchange_result_u32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1u, 1u);
} }
fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {

20
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
View File

@ -15,9 +15,9 @@
OpMemberName %SB_RW 0 "arg_0" OpMemberName %SB_RW 0 "arg_0"
OpName %sb_rw "sb_rw" OpName %sb_rw "sb_rw"
OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6" OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main" OpName %compute_main "compute_main"
@ -26,8 +26,8 @@
OpMemberDecorate %SB_RW 0 Offset 0 OpMemberDecorate %SB_RW 0 Offset 0
OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw DescriptorSet 0
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%SB_RW = OpTypeStruct %uint %SB_RW = OpTypeStruct %uint
%sb_rw_block = OpTypeStruct %SB_RW %sb_rw_block = OpTypeStruct %SB_RW
@ -36,19 +36,19 @@
%void = OpTypeVoid %void = OpTypeVoid
%6 = OpTypeFunction %void %6 = OpTypeFunction %void
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 %_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32
%22 = OpConstantNull %__atomic_compare_exchange_resultu32 %22 = OpConstantNull %__atomic_compare_exchange_result_u32
%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6 %atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6
%9 = OpLabel %9 = OpLabel
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %22 %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_u32 Function %22
%17 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0 %17 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0
%18 = OpAtomicCompareExchange %uint %17 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1 %18 = OpAtomicCompareExchange %uint %17 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1
%19 = OpIEqual %bool %18 %uint_1 %19 = OpIEqual %bool %18 %uint_1
%10 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %18 %19 %10 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %18 %19
OpStore %res %10 OpStore %res %10
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl
View File

@ -1,15 +1,15 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
groupshared uint arg_0; groupshared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = 1u; uint atomic_compare_value = 1u;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resultu32 res = atomic_result; atomic_compare_exchange_result_u32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl
View File

@ -1,15 +1,15 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
groupshared uint arg_0; groupshared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = 1u; uint atomic_compare_value = 1u;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resultu32 res = atomic_result; atomic_compare_exchange_result_u32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -8,10 +8,10 @@ struct atomic_compare_exchange_resultu32 {
shared uint arg_0; shared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, 1u, 1u); atomic_compare_result.old_value = atomicCompSwap(arg_0, 1u, 1u);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
atomic_compare_exchange_resultu32 res = atomic_compare_result; atomic_compare_exchange_result_u32 res = atomic_compare_result;
} }
void compute_main(uint local_invocation_index) { void compute_main(uint local_invocation_index) {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl
View File

@ -2,18 +2,18 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
} }
void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) { void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) {
atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u); atomic_compare_exchange_result_u32 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u);
} }
void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) { void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {

20
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
View File

@ -10,16 +10,16 @@
OpName %local_invocation_index_1 "local_invocation_index_1" OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %arg_0 "arg_0" OpName %arg_0 "arg_0"
OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d" OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %compute_main_inner "compute_main_inner" OpName %compute_main_inner "compute_main_inner"
OpName %local_invocation_index "local_invocation_index" OpName %local_invocation_index "local_invocation_index"
OpName %compute_main "compute_main" OpName %compute_main "compute_main"
OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
@ -28,21 +28,21 @@
%void = OpTypeVoid %void = OpTypeVoid
%6 = OpTypeFunction %void %6 = OpTypeFunction %void
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 %_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32
%21 = OpConstantNull %__atomic_compare_exchange_resultu32 %21 = OpConstantNull %__atomic_compare_exchange_result_u32
%22 = OpTypeFunction %void %uint %22 = OpTypeFunction %void %uint
%28 = OpConstantNull %uint %28 = OpConstantNull %uint
%uint_264 = OpConstant %uint 264 %uint_264 = OpConstant %uint 264
%atomicCompareExchangeWeak_83580d = OpFunction %void None %6 %atomicCompareExchangeWeak_83580d = OpFunction %void None %6
%9 = OpLabel %9 = OpLabel
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %21 %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_u32 Function %21
%17 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1 %17 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1
%18 = OpIEqual %bool %17 %uint_1 %18 = OpIEqual %bool %17 %uint_1
%10 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18 %10 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %17 %18
OpStore %res %10 OpStore %res %10
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl
View File

@ -1,15 +1,15 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
groupshared int arg_0; groupshared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = 1; int atomic_compare_value = 1;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resulti32 res = atomic_result; atomic_compare_exchange_result_i32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl
View File

@ -1,15 +1,15 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
groupshared int arg_0; groupshared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = 1; int atomic_compare_value = 1;
InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resulti32 res = atomic_result; atomic_compare_exchange_result_i32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -8,10 +8,10 @@ struct atomic_compare_exchange_resulti32 {
shared int arg_0; shared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, 1, 1); atomic_compare_result.old_value = atomicCompSwap(arg_0, 1, 1);
atomic_compare_result.exchanged = atomic_compare_result.old_value == 1; atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
atomic_compare_exchange_resulti32 res = atomic_compare_result; atomic_compare_exchange_result_i32 res = atomic_compare_result;
} }
void compute_main(uint local_invocation_index) { void compute_main(uint local_invocation_index) {

6
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl
View File

@ -2,18 +2,18 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
} }
void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol) { void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol) {
atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1); atomic_compare_exchange_result_i32 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1);
} }
void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) { void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {

20
test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
View File

@ -10,16 +10,16 @@
OpName %local_invocation_index_1 "local_invocation_index_1" OpName %local_invocation_index_1 "local_invocation_index_1"
OpName %arg_0 "arg_0" OpName %arg_0 "arg_0"
OpName %atomicCompareExchangeWeak_e88938 "atomicCompareExchangeWeak_e88938" OpName %atomicCompareExchangeWeak_e88938 "atomicCompareExchangeWeak_e88938"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %compute_main_inner "compute_main_inner" OpName %compute_main_inner "compute_main_inner"
OpName %local_invocation_index "local_invocation_index" OpName %local_invocation_index "local_invocation_index"
OpName %compute_main "compute_main" OpName %compute_main "compute_main"
OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
@ -29,21 +29,21 @@
%void = OpTypeVoid %void = OpTypeVoid
%7 = OpTypeFunction %void %7 = OpTypeFunction %void
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%int_1 = OpConstant %int 1 %int_1 = OpConstant %int 1
%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 %_ptr_Function___atomic_compare_exchange_result_i32 = OpTypePointer Function %__atomic_compare_exchange_result_i32
%22 = OpConstantNull %__atomic_compare_exchange_resulti32 %22 = OpConstantNull %__atomic_compare_exchange_result_i32
%23 = OpTypeFunction %void %uint %23 = OpTypeFunction %void %uint
%29 = OpConstantNull %int %29 = OpConstantNull %int
%uint_264 = OpConstant %uint 264 %uint_264 = OpConstant %uint 264
%atomicCompareExchangeWeak_e88938 = OpFunction %void None %7 %atomicCompareExchangeWeak_e88938 = OpFunction %void None %7
%10 = OpLabel %10 = OpLabel
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %22 %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_i32 Function %22
%18 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1 %18 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1
%19 = OpIEqual %bool %18 %int_1 %19 = OpIEqual %bool %18 %int_1
%11 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %18 %19 %11 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %18 %19
OpStore %res %11 OpStore %res %11
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

12
test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -18,10 +18,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_1bd40a() { void atomicCompareExchangeWeak_1bd40a() {
int arg_1 = 1; int arg_1 = 1;
int arg_2 = 1; int arg_2 = 1;
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2);
atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1; atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1;
atomic_compare_exchange_resulti32 res = atomic_compare_result; atomic_compare_exchange_result_i32 res = atomic_compare_result;
} }
void fragment_main() { void fragment_main() {
@ -34,7 +34,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -51,10 +51,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_1bd40a() { void atomicCompareExchangeWeak_1bd40a() {
int arg_1 = 1; int arg_1 = 1;
int arg_2 = 1; int arg_2 = 1;
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2);
atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1; atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1;
atomic_compare_exchange_resulti32 res = atomic_compare_result; atomic_compare_exchange_result_i32 res = atomic_compare_result;
} }
void compute_main() { void compute_main() {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -19,7 +19,7 @@ struct SB_RW {
void atomicCompareExchangeWeak_1bd40a(device SB_RW* const tint_symbol) { void atomicCompareExchangeWeak_1bd40a(device SB_RW* const tint_symbol) {
int arg_1 = 1; int arg_1 = 1;
int arg_2 = 1; int arg_2 = 1;
atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), arg_1, arg_2); atomic_compare_exchange_result_i32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), arg_1, arg_2);
} }
fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {

20
test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
View File

@ -17,9 +17,9 @@
OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a" OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a"
OpName %arg_1 "arg_1" OpName %arg_1 "arg_1"
OpName %arg_2 "arg_2" OpName %arg_2 "arg_2"
OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32"
OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main" OpName %compute_main "compute_main"
@ -28,8 +28,8 @@
OpMemberDecorate %SB_RW 0 Offset 0 OpMemberDecorate %SB_RW 0 Offset 0
OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw DescriptorSet 0
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4
%int = OpTypeInt 32 1 %int = OpTypeInt 32 1
%SB_RW = OpTypeStruct %int %SB_RW = OpTypeStruct %int
%sb_rw_block = OpTypeStruct %SB_RW %sb_rw_block = OpTypeStruct %SB_RW
@ -41,18 +41,18 @@
%_ptr_Function_int = OpTypePointer Function %int %_ptr_Function_int = OpTypePointer Function %int
%13 = OpConstantNull %int %13 = OpConstantNull %int
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%uint_1 = OpConstant %uint 1 %uint_1 = OpConstant %uint 1
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 %_ptr_Function___atomic_compare_exchange_result_i32 = OpTypePointer Function %__atomic_compare_exchange_result_i32
%30 = OpConstantNull %__atomic_compare_exchange_resulti32 %30 = OpConstantNull %__atomic_compare_exchange_result_i32
%atomicCompareExchangeWeak_1bd40a = OpFunction %void None %6 %atomicCompareExchangeWeak_1bd40a = OpFunction %void None %6
%9 = OpLabel %9 = OpLabel
%arg_1 = OpVariable %_ptr_Function_int Function %13 %arg_1 = OpVariable %_ptr_Function_int Function %13
%arg_2 = OpVariable %_ptr_Function_int Function %13 %arg_2 = OpVariable %_ptr_Function_int Function %13
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %30 %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_i32 Function %30
OpStore %arg_1 %int_1 OpStore %arg_1 %int_1
OpStore %arg_2 %int_1 OpStore %arg_2 %int_1
%23 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0 %23 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0
@ -60,7 +60,7 @@
%25 = OpLoad %int %arg_1 %25 = OpLoad %int %arg_1
%26 = OpAtomicCompareExchange %int %23 %uint_1 %uint_0 %uint_0 %24 %25 %26 = OpAtomicCompareExchange %int %23 %uint_1 %uint_0 %uint_0 %24 %25
%27 = OpIEqual %bool %26 %25 %27 = OpIEqual %bool %26 %25
%15 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %26 %27 %15 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %26 %27
OpStore %res %15 OpStore %res %15
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

12
test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl
View File

@ -1,7 +1,7 @@
#version 310 es #version 310 es
precision highp float; precision highp float;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -18,10 +18,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_63d8e6() { void atomicCompareExchangeWeak_63d8e6() {
uint arg_1 = 1u; uint arg_1 = 1u;
uint arg_2 = 1u; uint arg_2 = 1u;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2);
atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1; atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1;
atomic_compare_exchange_resultu32 res = atomic_compare_result; atomic_compare_exchange_result_u32 res = atomic_compare_result;
} }
void fragment_main() { void fragment_main() {
@ -34,7 +34,7 @@ void main() {
} }
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -51,10 +51,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo {
void atomicCompareExchangeWeak_63d8e6() { void atomicCompareExchangeWeak_63d8e6() {
uint arg_1 = 1u; uint arg_1 = 1u;
uint arg_2 = 1u; uint arg_2 = 1u;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2); atomic_compare_result.old_value = atomicCompSwap(sb_rw.inner.arg_0, arg_1, arg_2);
atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1; atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1;
atomic_compare_exchange_resultu32 res = atomic_compare_result; atomic_compare_exchange_result_u32 res = atomic_compare_result;
} }
void compute_main() { void compute_main() {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -19,7 +19,7 @@ struct SB_RW {
void atomicCompareExchangeWeak_63d8e6(device SB_RW* const tint_symbol) { void atomicCompareExchangeWeak_63d8e6(device SB_RW* const tint_symbol) {
uint arg_1 = 1u; uint arg_1 = 1u;
uint arg_2 = 1u; uint arg_2 = 1u;
atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), arg_1, arg_2); atomic_compare_exchange_result_u32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), arg_1, arg_2);
} }
fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) { fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {

20
test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
View File

@ -17,9 +17,9 @@
OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6" OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6"
OpName %arg_1 "arg_1" OpName %arg_1 "arg_1"
OpName %arg_2 "arg_2" OpName %arg_2 "arg_2"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %fragment_main "fragment_main" OpName %fragment_main "fragment_main"
OpName %compute_main "compute_main" OpName %compute_main "compute_main"
@ -28,8 +28,8 @@
OpMemberDecorate %SB_RW 0 Offset 0 OpMemberDecorate %SB_RW 0 Offset 0
OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw DescriptorSet 0
OpDecorate %sb_rw Binding 0 OpDecorate %sb_rw Binding 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%SB_RW = OpTypeStruct %uint %SB_RW = OpTypeStruct %uint
%sb_rw_block = OpTypeStruct %SB_RW %sb_rw_block = OpTypeStruct %SB_RW
@ -41,16 +41,16 @@
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%13 = OpConstantNull %uint %13 = OpConstantNull %uint
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 %_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32
%28 = OpConstantNull %__atomic_compare_exchange_resultu32 %28 = OpConstantNull %__atomic_compare_exchange_result_u32
%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6 %atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6
%9 = OpLabel %9 = OpLabel
%arg_1 = OpVariable %_ptr_Function_uint Function %13 %arg_1 = OpVariable %_ptr_Function_uint Function %13
%arg_2 = OpVariable %_ptr_Function_uint Function %13 %arg_2 = OpVariable %_ptr_Function_uint Function %13
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %28 %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_u32 Function %28
OpStore %arg_1 %uint_1 OpStore %arg_1 %uint_1
OpStore %arg_2 %uint_1 OpStore %arg_2 %uint_1
%21 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0 %21 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0
@ -58,7 +58,7 @@
%23 = OpLoad %uint %arg_1 %23 = OpLoad %uint %arg_1
%24 = OpAtomicCompareExchange %uint %21 %uint_1 %uint_0 %uint_0 %22 %23 %24 = OpAtomicCompareExchange %uint %21 %uint_1 %uint_0 %uint_0 %22 %23
%25 = OpIEqual %bool %24 %23 %25 = OpIEqual %bool %24 %23
%15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %24 %25 %15 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %24 %25
OpStore %res %15 OpStore %res %15
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -7,11 +7,11 @@ groupshared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
uint arg_1 = 1u; uint arg_1 = 1u;
uint arg_2 = 1u; uint arg_2 = 1u;
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = arg_1; uint atomic_compare_value = arg_1;
InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resultu32 res = atomic_result; atomic_compare_exchange_result_u32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -7,11 +7,11 @@ groupshared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
uint arg_1 = 1u; uint arg_1 = 1u;
uint arg_2 = 1u; uint arg_2 = 1u;
atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0; atomic_compare_exchange_result_u32 atomic_result = (atomic_compare_exchange_result_u32)0;
uint atomic_compare_value = arg_1; uint atomic_compare_value = arg_1;
InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resultu32 res = atomic_result; atomic_compare_exchange_result_u32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
@ -10,10 +10,10 @@ shared uint arg_0;
void atomicCompareExchangeWeak_83580d() { void atomicCompareExchangeWeak_83580d() {
uint arg_1 = 1u; uint arg_1 = 1u;
uint arg_2 = 1u; uint arg_2 = 1u;
atomic_compare_exchange_resultu32 atomic_compare_result; atomic_compare_exchange_result_u32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, arg_1, arg_2); atomic_compare_result.old_value = atomicCompSwap(arg_0, arg_1, arg_2);
atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1; atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1;
atomic_compare_exchange_resultu32 res = atomic_compare_result; atomic_compare_exchange_result_u32 res = atomic_compare_result;
} }
void compute_main(uint local_invocation_index) { void compute_main(uint local_invocation_index) {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resultu32 { struct atomic_compare_exchange_result_u32 {
uint old_value; uint old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) { atomic_compare_exchange_result_u32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) {
uint old_value = compare; uint old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -15,7 +15,7 @@ atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic
void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) { void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) {
uint arg_1 = 1u; uint arg_1 = 1u;
uint arg_2 = 1u; uint arg_2 = 1u;
atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(tint_symbol, arg_1, arg_2); atomic_compare_exchange_result_u32 res = atomicCompareExchangeWeak_1(tint_symbol, arg_1, arg_2);
} }
void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) { void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {

20
test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
View File

@ -12,16 +12,16 @@
OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d" OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d"
OpName %arg_1 "arg_1" OpName %arg_1 "arg_1"
OpName %arg_2 "arg_2" OpName %arg_2 "arg_2"
OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32"
OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value"
OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged"
OpName %res "res" OpName %res "res"
OpName %compute_main_inner "compute_main_inner" OpName %compute_main_inner "compute_main_inner"
OpName %local_invocation_index "local_invocation_index" OpName %local_invocation_index "local_invocation_index"
OpName %compute_main "compute_main" OpName %compute_main "compute_main"
OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0
OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4
%uint = OpTypeInt 32 0 %uint = OpTypeInt 32 0
%_ptr_Input_uint = OpTypePointer Input %uint %_ptr_Input_uint = OpTypePointer Input %uint
%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
@ -33,25 +33,25 @@
%_ptr_Function_uint = OpTypePointer Function %uint %_ptr_Function_uint = OpTypePointer Function %uint
%13 = OpConstantNull %uint %13 = OpConstantNull %uint
%bool = OpTypeBool %bool = OpTypeBool
%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool
%uint_2 = OpConstant %uint 2 %uint_2 = OpConstant %uint 2
%uint_0 = OpConstant %uint 0 %uint_0 = OpConstant %uint 0
%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 %_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32
%27 = OpConstantNull %__atomic_compare_exchange_resultu32 %27 = OpConstantNull %__atomic_compare_exchange_result_u32
%28 = OpTypeFunction %void %uint %28 = OpTypeFunction %void %uint
%uint_264 = OpConstant %uint 264 %uint_264 = OpConstant %uint 264
%atomicCompareExchangeWeak_83580d = OpFunction %void None %6 %atomicCompareExchangeWeak_83580d = OpFunction %void None %6
%9 = OpLabel %9 = OpLabel
%arg_1 = OpVariable %_ptr_Function_uint Function %13 %arg_1 = OpVariable %_ptr_Function_uint Function %13
%arg_2 = OpVariable %_ptr_Function_uint Function %13 %arg_2 = OpVariable %_ptr_Function_uint Function %13
%res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %27 %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_u32 Function %27
OpStore %arg_1 %uint_1 OpStore %arg_1 %uint_1
OpStore %arg_2 %uint_1 OpStore %arg_2 %uint_1
%21 = OpLoad %uint %arg_2 %21 = OpLoad %uint %arg_2
%22 = OpLoad %uint %arg_1 %22 = OpLoad %uint %arg_1
%23 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %22 %23 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %22
%24 = OpIEqual %bool %23 %22 %24 = OpIEqual %bool %23 %22
%15 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %23 %24 %15 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %23 %24
OpStore %res %15 OpStore %res %15
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -7,11 +7,11 @@ groupshared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
int arg_1 = 1; int arg_1 = 1;
int arg_2 = 1; int arg_2 = 1;
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = arg_1; int atomic_compare_value = arg_1;
InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resulti32 res = atomic_result; atomic_compare_exchange_result_i32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl
View File

@ -1,4 +1,4 @@
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -7,11 +7,11 @@ groupshared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
int arg_1 = 1; int arg_1 = 1;
int arg_2 = 1; int arg_2 = 1;
atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0; atomic_compare_exchange_result_i32 atomic_result = (atomic_compare_exchange_result_i32)0;
int atomic_compare_value = arg_1; int atomic_compare_value = arg_1;
InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value); InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_value);
atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
atomic_compare_exchange_resulti32 res = atomic_result; atomic_compare_exchange_result_i32 res = atomic_result;
} }
struct tint_symbol_1 { struct tint_symbol_1 {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl
View File

@ -1,6 +1,6 @@
#version 310 es #version 310 es
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
@ -10,10 +10,10 @@ shared int arg_0;
void atomicCompareExchangeWeak_e88938() { void atomicCompareExchangeWeak_e88938() {
int arg_1 = 1; int arg_1 = 1;
int arg_2 = 1; int arg_2 = 1;
atomic_compare_exchange_resulti32 atomic_compare_result; atomic_compare_exchange_result_i32 atomic_compare_result;
atomic_compare_result.old_value = atomicCompSwap(arg_0, arg_1, arg_2); atomic_compare_result.old_value = atomicCompSwap(arg_0, arg_1, arg_2);
atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1; atomic_compare_result.exchanged = atomic_compare_result.old_value == arg_1;
atomic_compare_exchange_resulti32 res = atomic_compare_result; atomic_compare_exchange_result_i32 res = atomic_compare_result;
} }
void compute_main(uint local_invocation_index) { void compute_main(uint local_invocation_index) {

6
test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl
View File

@ -2,11 +2,11 @@
using namespace metal; using namespace metal;
struct atomic_compare_exchange_resulti32 { struct atomic_compare_exchange_result_i32 {
int old_value; int old_value;
bool exchanged; bool exchanged;
}; };
atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) { atomic_compare_exchange_result_i32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) {
int old_value = compare; int old_value = compare;
bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
return {old_value, exchanged}; return {old_value, exchanged};
@ -15,7 +15,7 @@ atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic
void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol) { void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol) {
int arg_1 = 1; int arg_1 = 1;
int arg_2 = 1; int arg_2 = 1;
atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(tint_symbol, arg_1, arg_2); atomic_compare_exchange_result_i32 res = atomicCompareExchangeWeak_1(tint_symbol, arg_1, arg_2);
} }
void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) { void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {

Some files were not shown because too many files have changed in this diff Show More