diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 8f9a688e29..ee57d81110 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -1416,6 +1416,7 @@ if (tint_build_unittests) { "resolver/attribute_validation_test.cc", "resolver/bitcast_validation_test.cc", "resolver/builtin_enum_test.cc", + "resolver/builtin_structs_test.cc", "resolver/builtin_test.cc", "resolver/builtin_validation_test.cc", "resolver/builtins_validation_test.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index c5bcfb0793..64e0939a8e 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -913,6 +913,7 @@ if(TINT_BUILD_TESTS) resolver/attribute_validation_test.cc resolver/bitcast_validation_test.cc resolver/builtin_enum_test.cc + resolver/builtin_structs_test.cc resolver/builtin_test.cc resolver/builtin_validation_test.cc resolver/builtins_validation_test.cc diff --git a/src/tint/builtin/access_bench.cc b/src/tint/builtin/access_bench.cc index 91f808f1b1..58a5f47a91 100644 --- a/src/tint/builtin/access_bench.cc +++ b/src/tint/builtin/access_bench.cc @@ -43,7 +43,7 @@ void AccessParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(AccessParser); diff --git a/src/tint/builtin/address_space_bench.cc b/src/tint/builtin/address_space_bench.cc index 18cc9d241b..81b27eb285 100644 --- a/src/tint/builtin/address_space_bench.cc +++ b/src/tint/builtin/address_space_bench.cc @@ -94,7 +94,7 @@ void AddressSpaceParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(AddressSpaceParser); diff --git a/src/tint/builtin/attribute_bench.cc b/src/tint/builtin/attribute_bench.cc index 5d12ab3200..a25fbf12d1 100644 --- a/src/tint/builtin/attribute_bench.cc +++ b/src/tint/builtin/attribute_bench.cc @@ -143,7 +143,7 @@ void AttributeParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(AttributeParser); diff --git a/src/tint/builtin/builtin.cc b/src/tint/builtin/builtin.cc index 085a2661a7..8f26436153 100644 --- a/src/tint/builtin/builtin.cc +++ b/src/tint/builtin/builtin.cc @@ -28,6 +28,84 @@ namespace tint::builtin { /// @param str the string to parse /// @returns the parsed enum, or Builtin::kUndefined if the string could not be parsed. 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") { return Builtin::kPackedVec3; } @@ -245,6 +323,58 @@ utils::StringStream& operator<<(utils::StringStream& out, Builtin value) { switch (value) { case Builtin::kUndefined: 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: return out << "__packed_vec3"; case Builtin::kArray: diff --git a/src/tint/builtin/builtin.h b/src/tint/builtin/builtin.h index da5226df7f..d891294486 100644 --- a/src/tint/builtin/builtin.h +++ b/src/tint/builtin/builtin.h @@ -30,6 +30,32 @@ namespace tint::builtin { /// An enumerator of builtin builtin. enum class Builtin { 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, kArray, kAtomic, @@ -113,6 +139,32 @@ utils::StringStream& operator<<(utils::StringStream& out, Builtin value); Builtin ParseBuiltin(std::string_view str); 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", "array", "atomic", diff --git a/src/tint/builtin/builtin_bench.cc b/src/tint/builtin/builtin_bench.cc index 1f309f5677..4804ff7b70 100644 --- a/src/tint/builtin/builtin_bench.cc +++ b/src/tint/builtin/builtin_bench.cc @@ -31,496 +31,678 @@ namespace { void BuiltinParser(::benchmark::State& state) { const char* kStrings[] = { - "__acked_veccc", - "_pac3ed_v3", - "__packeV_vec3", + "__atomic_compareexchangeccresult_i32", + "__atoml3_compare_exchane_resulti2", + "__atomic_compare_Vxchange_result_i32", + "__atomic_compare_exchange_result_i32", + "__atomic_com1are_exchange_result_i32", + "__atomic_qqompare_exchage_resulJ_i32", + "__atllmic_compare_exchange_result_i377", + "__atomicppcompareqqexchange_reslt_uHH2", + "__atomi_compare_exchavge_cesult_3", + "__atomic_copare_eGbhange_result_u32", + "__atomic_compare_exchange_result_u32", + "__atomic_coiipare_exvhange_result_u32", + "__atomic_compaWWe_excha8ge_result_u32", + "__atomic_comparxxMexchage_result_u32", + "__fXexp_resgglt_bstract", + "V_frexp_resul_abuXrct", + "__frexp_result_abstra3t", + "__frexp_result_abstract", + "__frexp_resElt_abstract", + "__frexTT_Pesult_abstract", + "__frexp_resulxxddabstrct", + "44_frexp_result_f16", + "_VVfrexp_resulSS_f16", + "__frexp_reRult_fR6", + "__frexp_result_f16", + "__frFxp_re9ut_f16", + "__frep_result_f16", + "__frRRVH_rOOsultf16", + "__frepyresult_f32", + "_nrr77rexp_result_fGll", + "__4rex00_result_f32", + "__frexp_result_f32", + "__oorep_reult_f2", + "__fzzexp_result_3", + "__iir11x_respplt_f3", + "__frexp_resuXXt_vec2_abstract", + "55n99frexp_result_vec2_abstraIIt", + "__fHHexpSSaresrrlt_Yec2_abstract", + "__frexp_result_vec2_abstract", + "__freHp_resutve2_abstkkact", + "jfrexpgresult_veRR2_abstrac", + "__frexp_resul_vec2_absbrac", + "_jfrexp_result_vec2_f16", + "__frexp_resultvec2_f16", + "__freqpresultvec2_f16", + "__frexp_result_vec2_f16", + "__frexNN_result_vec_f16", + "__frexp_resvvlt_vc2_f1", + "__frexp_esult_vec2_f1QQ", + "__rerp_result_ffec2_f2", + "__frexp_result_vjc2_f32", + "__frewwp_reul2_vec2_NN82", + "__frexp_result_vec2_f32", + "__frexpresult_vec2_f32", + "__frexp_result_vec2_frr2", + "_Gfrexp_result_vec2_f32", + "__frexp_resulFF_vec3_abstract", + "_frexp_resultvec3_Estract", + "__fexp_result_vec3_abstrract", + "__frexp_result_vec3_abstract", + "frexp_result_vec3_abstract", + "D_rexp_resXlt_veJJ3_abstract", + "_frexp_resut_v8c_abstract", + "_frexp_rsl1k_vec3_f16", + "__frexp_reslt_vec3_f16", + "__frexJ_reult_vec3_f16", + "__frexp_result_vec3_f16", + "c_frexp_result_vec3_f16", + "__frexp_result_vec3Of16", + "___frexp_reKKultvvvec3_f1tt", + "8_frexp_reult_vxxc3_f32", + "_frexp_resul___veFqqf32", + "_qqfrexp_result_vec_f32", + "__frexp_result_vec3_f32", + "33_fOexp_result_ve3_6632", + "__oorexQQ_rttsult_ve639f32", + "__rexp_result_vec3_f662", + "__frexp_reszzlt_Oc4xabstrac66", + "__frexp_resyylt_vec4_abstract", + "__frexp_resut_vecHH_aZsracZ", + "__frexp_result_vec4_abstract", + "_WWfrex44_resulq_vec4_astract", + "__frexp_rsult_veOO4_abstract", + "__frexp_resultoovc4_abstYct", + "_frexp_esultvec4_f16", + "__Frexp_result_ec4_f16", + "__frewp_resut_vec4_f16", + "__frexp_result_vec4_f16", + "__frexp_reslt_veK4fG16", + "__fqexp_result_veKK4_f16", + "_F3rexp_result_vec4_f1mm", + "__frexp_result_ec4_f32", + "__frexp_result_qe4_f32", + "__frbbxp_result_vec4_b2", + "__frexp_result_vec4_f32", + "__frexp_reslt_iiec4_f2", + "__frexO_resulq_vec4_f32", + "__frexp_resulTT_vec4vvf32", + "__modf_resulFF_abstract", + "fm00df_rePult_abstraQt", + "__modf_result_abstPact", + "__modf_result_abstract", + "_modf_result_abstssac77", + "__modf_resulC_bbRbstract", + "__modf_result_abstracXX", + "__OOofCCresuOOt_f16", + "_smodf_resuutfL6", + "__modX_result_f16", + "__modf_result_f16", + "__modf_reult_f16", + "__modf_resqqO16", + "__modf22result_f16", + "__modf_X0eszzlt_fy", + "_VVmPf_result_f3i", + "__monnfCresultf32", + "__modf_result_f32", + "_HHAmodf_resqltf32", + "__modf_resut_f32", + "__modresuft_f3KK", + "__modlPrsultggvec2_astract", + "__odf_result_vec2_abstract", + "__mocTf_result_vNc2_abstra4t", + "__modf_result_vec2_abstract", + "__modf77result_vec2_plbtract", + "__mdf_resultNNvec2zabstgact", + "_modf_bbesult_vuuc2_abtraXXt", + "__modf_esult_vec2_f16", + "__mQdf_esuKt_vec_8816", + "q_m9dfresult_vec2_f16", + "__modf_result_vec2_f16", + "__11odf_result_vec2_f16", + "_iimodf_result_vF222f16", + "_77modf_result_vec2f16", + "__odfNNr2sult_vec2_f32", + "__modf_rVVsult_vec2_f32", + "__modf_Fesult_vewW2_f311", + "__modf_result_vec2_f32", + "__modf_rwwsult_vec_f32", + "__modf_result_Dec2_f32", + "__modf_result_ec2_f3K", + "__modf_resul1PP_vech_abstfact", + "__modf_result_vec_abstract", + "__YYodf_result_vec3_abstract", + "__modf_result_vec3_abstract", + "__mHHdfresult_kkec3_abstract", + "__modf_result_vec3rrabstract", + "__modf_ssesulWW_vec_abstract", + "__mYdf_reslt_vec3_f16", + "q_modLrfsult_vec3_f16", + "uu_vvo22f_rfsult_ec3_f16", + "__modf_result_vec3_f16", + "__mdf_reslt_vec3_f16", + "__modfYYresult_ve3f16", + "__modfEr77sult_vec3_yY16", + "__odf_desuMot_vec3_f32", + "__mMMf_result_vec3_f32", + "__modf_result_vec3_f355", + "__modf_result_vec3_f32", + "__modf_rest_vec3Nf32", + "_m33df_result_Oec3_f32", + "__modf_re3ult_vec3_f32", + "__momf_esult_Iec4_abstract", + "__modf_resultrvec4_absnnracK", + "__modf_eslt_ve4_absXXact", + "__modf_result_vec4_abstract", + "__modf_rsult_pLLI4_abstract", + "_modf_resflt_vec4_bstract", + "_Ymodf_resultURDec4_abtract", + "__hodf_result_vec4_f16", + "__moquu_rslt_vec4_f1II", + "__modf_result_vecH_f16", + "__modf_result_vec4_f16", + "__oQQf_resultvvvc4_f16", + "__modf_eeult66ec4_f16", + "_Omodf_r7sut_vec4_W16", + "__modf_DDes0lt_v55c4_f32", + "__modf_rIIsult_Hec4_f32", + "_modf_result_vec4_f32", + "__modf_result_vec4_f32", + "_modf_result_vrc4_f32", + "_lmodf_result_vec4_f32", + "tt_modfGeslt_vec4_fJJ2", + "__paked_vyc3", + "_packed_vec3", + "__pIIckedBBvec3", "__packed_vec3", - "__pa1ked_vec3", - "_qqJcked_vec3", - "__pack77d_vllc3", - "arqHapp", - "vy", - "Grby", + "__8aTTked_v33c3", + "dnnUUpackeSSY_vec3", + "xC_5ackedZvec3", + "kkrraq", + "a005iy", + "anIIray", "array", - "arviay", - "ar8WWy", - "Mxxra", - "atXggi", - "Xoic", - "ato3ic", + "ccrW", + "rKK", + "arr66y", + "aKKoPi", + "atxxmc", + "atoqic", "atomic", - "aEomic", - "toTTiPP", - "ddtoxxi", - "44ool", - "VVSSol", - "RoRl", + "rMoyyiSS", + "utom", + "oic", + "5oFFl", + "borz4l", + "WW", "bool", - "oFl", - "boo", - "ORVHl", - "y1", - "l77rrn6", - "4016", + "ZJJCoX", + "boPP", + "bocl", + "fll66", + "91yy", + "f1KK", "f16", - "5", - "u16", - "f", - "f3kk", - "fi", - "f3XX", + "x_", + "K", + "kVz", + "K3S", + "f2", + "fVV", "f32", - "55399II", - "frSSHHa", - "U", - "jV3", - "", - "GG", + "IAU2", + "j", + "Y4", + "i2", + "1xx", + "ccm", "i32", - "2", - "", - "jj", - "a2xrf", - "mat2j2", - "m82wNN2", - "mat2x2", + "iJJ", + "UfCDD", + "i3g", + "CCtx", "mt2x2", - "rrat2x2", - "mGt2x2", - "mat2x2FF", - "at2f", - "marrx2f", + "mat2x__", + "mat2x2", + "attxPP", + "mdd32x2", + "yyK2x2", + "m2uu", + "ma0nnx2i", + "KanuuCC2f", "mat2x2f", - "t2x2f", - "Da2xJJf", - "ma82", - "m11k2", - "matx2h", - "maJx2h", + "mlX2x2f", + "oat2pp2f", + "wwat22f", + "matguum", + "mt2ma2", + "Tat2xZRRh", "mat2x2h", - "mat2c2h", - "mat2x2O", - "KK_atvvtt2h", - "5txxx8", - "a__xqq", - "maqq2x", + "ma8T2xOh", + "m0at2x2h", + "mBBt2x2h", + "Matpp", + "Oat2x3", + "GGG2x3", "mat2x3", - "ma32x66", - "mttQQo2x3", - "mat66x", - "mtOxzz66", - "mat2yy3f", - "ZaHH3Z", - "mat2x3f", - "4WWt2q3f", - "mOO2x3f", - "oatY3f", + "mHHt2113", + "mateF63", "matx", - "ma2xFh", - "at2x3w", + "mat2ii3l", + "mt2x3f", + "IIvvt2x39", + "mat2x3f", + "mat23f", + "mat2h3f", + "mllt2xPzz", + "t3h", + "mtffxqqh", + "mtJJx3dd", "mat2x3h", - "fGtxKh", - "matqKx3h", - "matmmxFh", - "at2x4", - "matxq", - "mb2bb4", + "mzz2X3h", + "matx32", + "maN2yy3h", + "atxO", + "rauExP", + "meet22dd", "mat2x4", - "it2x4", - "mOO2xq", - "mat2Tvv4", - "maFF2x4f", - "Pa00xQf", - "mPt2x4f", + "maV92", + "maI2x1", + "mab2x4", + "matzf", + "mao2ii4f", + "mat45", "mat2x4f", - "ma772xss", - "RRCbb2x4f", - "mXXt2x4f", - "qaCC2xOOh", - "ma2s4L", - "mXt2x4h", + "at2xSf", + "mat22f", + "maG1C4f", + "maff284h", + "t2x4h", + "SSatJJx4h", "mat2x4h", - "mat24h", - "qa2O4", - "mat2x22h", - "mzzyt3x", - "atiVP2", - "mt3Cnn", + "atx9h", + "maJJbbTT4h", + "66a2xh", + "ma663u", + "Wa3x2", + "ma32", "mat3x2", - "AtqqHH2", - "at3x2", - "mafKK", - "ltgg2f", - "mat3xf", - "NTTtcx4f", + "ma3x2", + "rat3x2", + "m2t3xB", + "matxBBf", + "maRR3xf", + "maVV0Lf", "mat3x2f", - "ma7ppl2f", - "mNNt3xg", - "uub3XX2f", - "matx2h", - "Qt882h", - "mt9q2h", + "a3OOK2f", + "magw3xf", + "hht3L2f", + "aKii3xh", + "ma3x2h", + "UUa3882", "mat3x2h", - "m11t3x2h", - "22at3iih", - "at3x277", - "NNa323", - "VVat3x3", - "ma11F3w3", + "rrvvt3x2h", + "m3xwmm", + "j443x2h", + "matXx3", + "m8t3x3", + "mat3vEE", "mat3x3", - "matww3", - "mat3D3", - "maKx3", - "mat31PPhf", - "mat33f", - "mYYt3x3f", - "mat3x3f", - "mttHH3kk", - "mat3rr3f", - "WWas3x3f", - "Yt3x3h", - "mt3qfh", - "mav223xuh", - "mat3x3h", - "t3x3h", - "YYat3h", - "may3x3EYY", - "Moatd4", - "mt3xMM", - "m55t3x4", - "mat3x4", - "maN4", - "ma33x4", + "mzzi3x", + "maGGQ3JJ3", + "mat3ss3", + "matKxPf", + "mat3ttf", "mt3x3", - "mm66Issf", - "mat3x1f", - "Xt3x4", + "mat3x3f", + "mMMt3x3f", + "maJ03x3f", + "V8x3", + "maKggx3hh", + "maf3x3h", + "matQ7x3h", + "mat3x3h", + "mat3YYh", + "mak3x3", + "man3x2", + "mFFx4", + "GGatPPuUU", + "mEEFa4", + "mat3x4", + "mBet3dd4", + "55atExcc", + "txKK", + "mat3x4R", + "maDx49", + "mt3x4f", "mat3x4f", - "LatIx4f", - "at3ff", - "mYtURD4", - "mah3x4h", - "uuIqt3x", - "maH3x4h", + "aaat3I", + "m77t3x4f", + "matIx4f", + "md3x4h", + "mat34h", + "mtt4h", "mat3x4h", - "at3QQvv", - "at66eh", - "ma7O4h", - "m0t55DD2", - "IIaH4x2", - "mat4x", + "ma3XX3x4h", + "Eat34h", + "maXX3x4", + "mxBt4x2", + "Wt4x", + "mat66x2", "mat4x2", - "mt4r2", - "mat4xl", - "mGttx2", - "mat4y2", - "mt4x2f", - "IIaBB4x2f", + "atTv0", + "kt", + "mpt4x", + "at112f", + "EaJ4yBBf", + "mqIm4x2f", "mat4x2f", - "TTat4x833", - "ddUUnntYYx2f", - "m5CCxxdZ", - "matkkq2h", - "005itpxh", - "maIInnx2h", + "ma4xFf", + "Yt4x2f", + "mHHtDh2f", + "Ht22h", + "matx2", + "matx2h", "mat4x2h", - "Ka4Wcc", - "m42KK", - "mat66x2h", - "mKKtPx", - "maxx43", - "matqx3", + "matx2h", + "matddx2h", + "Oat4x2h", + "bbtB3", + "m00tx3", + "hat4x3", "mat4x3", - "rMtyyxSS", - "uat4", - "tx3", - "ma5F4x3f", - "rra444z3f", - "matWW", + "matgYx", + "Oat4x3", + "mhx3", + "fpaEEx3f", + "mavx3f", + "mzztx3f", "mat4x3f", - "CatZJXx3f", - "maPPx3f", - "mat4c3f", - "matPPll6h", - "mat993yy", - "mat4JKKh", + "ma4x3f", + "OOaJxii", + "mft4G3f", + "mat4x322T", + "datlx3h", + "bat4x3h", "mat4x3h", - "mat4_h", - "ayK3h", - "mzt4V3k", - "qaSKx4", - "mat44", - "ma4xVV", + "BBatx3h", + "PPIXt4S3h", + "mjjt4x3h", + "macc4_4", + "SS6zz4xx", + "mtx", "mat4x4", - "mAAt4xI", - "jb44", - "t4YYx", - "mao4x4", - "mtx114f", - "mtmxccf", + "mxxtvN", + "AA00t44", + "tyexy", + "mabWWx0f", + "ttatMMxmf", + "madf", "mat4x4f", - "aJJ4x4f", - "fCCDD4x4U", - "mgt4x4f", - "CCx4h", - "mat4x66", - "maN4M4h", + "mat_4f", + "Vat4EE4f", + "mat44f", + "mRIxah", + "ma4mmh", + "at4x4p", "mat4x4h", - "mattth", - "maKWxh", - "mateezx4h", - "", - "w9", - "4tnn", + "mat4xh", + "aaxh", + "mad4x4h", + "pCPtd", + "p", + "5tr", "ptr", - "tll", - "4to", - "wEgg", - "gamler", - "spleS", - "aampl", + "ff99j", + "YYvXR", + "r", + "XX8m5le", + "mpler", + "sccmlppr", "sampler", - "TamplZRRr", - "sa8TplOr", - "m0ampler", - "sampler_Bmomparison", - "Mamper_ppomarison", - "samper_compOOrison", + "sampver", + "EESSmplr", + "smplr", + "samplecomp_risa", + "sampler_co_prwwson", + "samplerdd99omparison", "sampler_comparison", - "sampler_compGrGGon", - "sHHm11ler_comparison", - "sa6ler_FFeemparison", - "texure_1", - "tKiilure_1d", - "exture_1d", + "ampler_o99paPPison", + "saplerKKcomparison", + "saMpler_oomDDarison", + "teiie_1B", + "txureq1d", + "txt00rLL_d", "texture_1d", - "99etvIIre_1d", - "texture_d", - "texture_hd", - "llxzzure_PPd", - "exue2d", - "tffqqtre_2d", + "tnxture_16vv", + "trrxtur_nd", + "xxture_eed", + "CCNOxture_2d", + "txture_2d", + "tex4uae_2d", "texture_2d", - "texturJdd_d", - "trXXtu_2zz", - "textu2e2d", - "tNyyture_2d_array", - "txture_2d_rOOa", - "textureErduaZPay", + "extuNNe_2NN", + "texture2d", + "tuxtre2d", + "teYYtuAe_2d_arESy", + "texture_2d_0rray", + "texture_2d_aarray", "texture_2d_array", - "exl22re_2dd_areeay", - "mextureVV_ar9ay", - "teIItu1_2d_array", - "tebture_3d", - "ie7ure3d", - "teotiire_3d", + "texturmm_2d_arra", + "texture_2d_aray", + "teEuUUe_2darray", + "tKKture_Dd", + "text__r0_3d", + "tAtuel3p", "texture_3d", - "extre_35", - "textre_iS", - "t22xtur_3", - "teC711recuGe", - "texture8cffbe", - "textue_cue", + "textue_3d", + "texturBB_3d", + "nnbb99re_3d", + "AAEExture_cub", + "t66Ttu5e_cube", + "textuHe_cube", "texture_cube", - "texture_SSJJbe", - "textrecu9e", - "TTeJJbture_cube", - "t66ture_cube_aray", - "textur66_cubu_arra", - "textureWubeyarray", + "textrexxHcub", + "tzx0uryy_cnbe", + "texture_cue", + "texurH_kube_array", + "exture_cube_array", + "ooexrrre_cbe_array", "texture_cube_array", - "texture_cube_ara", - "texture_ube_array", - "rexture_cube_array", - "tex2ure_depth_2B", - "texture_dpBBh_2d", - "texture_dpth_RRd", + "textre_cubJJarray", + "tCCxtu0e_cube_arry", + "texturAAcxbe99aFray", + "textcre_depth_2d", + "texture_Septh_2d", + "textureodpthBB2d", "texture_depth_2d", - "tLL0Vure_deth_2d", - "tetKKredOOpth_2d", - "textgwre_dpth_2d", - "textue_depthLh2d_arpay", - "texture_depEh2diiKrray", - "texture_dept_2d_array", + "texture_dept_2d", + "textummedepth_2d", + "toxture_ggeQQtPP2d", + "tetur_dptB_2d_rray", + "texNure_deKKh2d_arrlly", + "teture_dpth_2d_arrray", "texture_depth_2d_array", - "t88xtuUUe_deph_2d_rray", - "texrruvve_depth_2d_array", - "texture_depmm_2d_wray", - "tjture_d44pth_cube", - "texture_depth_cXbe", - "t8xture_depth_cube", + "texture_epth_ppd_array", + "teyturPP_depth_2d_array", + "texture_ZZptcc_2d_arry", + "txtue_depth_cube", + "texture00depth_cube", + "texPPuBB_deJth_cusse", "texture_depth_cube", - "textur_devvth_cubEE", - "tzxturi99epth_cube", - "teQQtuJJGe_nepth_cuAe", - "texture_depth_cusse_array", - "texture_Pepth_cKbe_array", - "texture_dppp_cube_attray", + "texJJre_dffpwwh_fube", + "textIre_depXXh_cub", + "textur_ph_cue", + "textue_depth_cube_array", + "tKKxtue_depth_cube_array", + "teture_d44ptm_cube_adray", "texture_depth_cube_array", - "exture_deth_cube_array", - "texture_depth_MMube_array", - "tJJxture_depth_cube_a0ray", - "textu8_dpth_mulisampled_2V", - "texture_dhhpth_mKltisggmpled_2d", - "texture_depth_multisampledf2d", + "pexture_deoth_cube_array", + "thhHxtureNdepth_cubejarray", + "texwwuUUe_depthEc33be_array", + "texture_dept_multiuuampled_2", + "ddextKre_depth_ultisampcerr_2d", + "textuPPe_depr_multttsample2_2d", "texture_depth_multisampled_2d", - "te77ture_depth_multisamQled_2d", - "teture_depthYYmultisampled_2d", - "texture_deptk_multiampled_Sd", - "txturn_ext2rnal", - "txture_FFternal", - "texUPPIre_GGxuernal", + "1exture_depth_wwsltisampled_2d", + "textuce_depth_mnnltisamp11ed_2d", + "texture_depth_multisapled_2d", + "texture_externl", + "teSS66ue_exaaeInal", + "textuEEe_extenal", "texture_external", - "taxtuvEE_externl", - "textureexBddernDDl", - "tEEMtur_e55tccrnal", - "texturemuKKtisample_d", - "texture_multisRmpled_2d", - "texturemulDisampl9d_2d", + "ccexture_exVerIRl", + "te9tue_extrnal", + "taaxture_exterha", + "texture_multisamLLeS_2d", + "tefurmm_mutisampled_2d", + "texture_mul4isampld_qV", "texture_multisampled_2d", - "teture_multisampled_2d", - "textuIa_multisampld_2d", - "texture_multisamp77ed_2d", - "texIure_storage_1d", - "texture_storagedd", - "texture_storae_1d", + "texture_multisa_pled_d", + "texure_multisampledQd", + "texRRuremultisampledEd2d", + "textur_st9rage_1d", + "tCCx0ure_strag_1", + "textuezstorae_1d", "texture_storage_1d", - "texture_strate_d", - "texture33stoXXcge_1d", - "texturestorage_1E", - "textuXXestorage_2d", - "texture_stoBaxxe_2d", - "texte_storWge_2G", + "texccure_storage_1d", + "textureOQQ_orge_1d", + "teturettstrage_1d", + "textCCrepzzstEr33ge_2d", + "textudde_storaghh_2d", + "_etur77_66torage_2d", "texture_storage_2d", - "texture_s66orage_2d", - "textvTr_so0age_2d", - "textureorgek2d", - "textppre_stoae_2d_array", - "textre_stora11e_d_array", - "textureystorBEgeJ2d_array", + "texture_storaPe_2d", + "twxture_storage_2d", + "textur_straguu_2", + "textureXXstorage_6d_array", + "extuRRestorage_2d_aray", + "txtre_storage_2dVVarr1", "texture_storage_2d_array", - "texture_mtorage_2dxIqrray", - "teture_storageF2d_array", - "textur_Ytorage_2d_array", - "heDture_sHHorage_3d", - "texturstorage23H", - "teture_strage_3d", + "GGexture_storHHge_2d_array", + "MFFxt7re_storage_2d_array", + "texture_storage_d_array", + "3xTugge_stoage_3d", + "texturP_QtKKrag1__d", + "textre_storageE3d", "texture_storage_3d", - "texture_storage_d", - "texturestorage_3d", - "ddexture_storage_3d", - "uPO", - "ba", - "u02", + "tMture_storage_d", + "t77xture_sGGorSSe_3d", + "txtttre_storage_3FF", + "uZZss2", + "u2", + "u3l", "u32", - "h32", - "gYY", - "O32", - "eh", - "ppfe2", - "vev", + "u3h", + "uTT", + "ww2", + "vKvjj", + "vYY", + "EcI2", "vec2", - "vzz2", - "vc2", - "OOii", - "vGc2f", - "22ecTTf", - "dlc2f", + "vecQQ", + "Pc", + "veffH", + "vec2n", + "g6F2f", + "vssh8f", "vec2f", - "vecbf", - "ec2BB", - "IIScXPP", - "jjec2h", - "cc_c2h", - "zz6xx2h", + "veFllf", + "00e2j", + "gec2f", + "vece", + "ffc2h", + "ve", "vec2h", - "c2", - "4xx2N", - "p0AAeh", - "vey2", - "vbWW0i", - "meMMtti", + "ve2h", + "vqc2h", + "AAe", + "ec2i", + "vec2j", + "ZZec2i", "vec2i", - "di", - "vvc_", - "VEEc2i", - "vec24", - "VVeX2u", - "veVou", + "PPecII2", + "ZZec2i", + "vnnc2i", + "HekkZ222", + "ec2", + "RcNQQ", "vec2u", - "ve2u", - "ecKKt", - "eG", - "ea3", - "OOc", - "G", + "eDu", + "s3c2cu", + "vRR2u", + "vlJJ3", + "MM", + "vT63", "vec3", - "ve53", - "9fjec3", - "vvXcRY", - "ccf", - "v8XX5", - "ec3", + "QQec3", + "vuA", + "e3", + "yeq3", + "vec3xx", + "crr", "vec3f", - "ppc3cc", - "vecvf", - "eEE3SS", - "vec", - "eh", - "ec3ww", + "v99cf", + "vecf", + "ecHl", + "e_h", + "uec3", + "vc3h", "vec3h", - "vecd99h", - "ve99P", - "KKec3", - "ooMcDD", - "vei", - "vqi", + "EEtmec3h", + "vec", + "ec3rr", + "xc3i", + "vezz", + "vec3e", "vec3i", - "veL30", - "vncvv66", - "vrrn3", - "vxxce", - "NCCOc3u", - "vc3u", + "uc3Zp", + "00uc7TT", + "vvJJ", + "vecQu", + "ve3R", + "e", "vec3u", - "aec4u", - "NNc3NN", - "ve3u", - "vc", - "vAYS4", - "vec0", + "veprPP", + "xxeDD88u", + "lldmYYqqu", + "FFec4", + "rGecNN", + "Mecl", "vec4", - "vecaa", - "mmcq", - "vc4", - "vE4U", - "veKD4", - "v0t4__", + "c", + "qxl4", + "ve4", + "ae44f", + "vec4WW", + "vecjj", "vec4f", - "cpA", - "ec4f", - "vBBc4f", - "vbnn99", - "EEcAAh", - "v5c66h", + "vjjc4f", + "vj1f", + "vc4f", + "vec499", + "vyVV4h", + "ec4xZ", "vec4h", - "vHc4h", - "vecxh", - "vzyn40", - "ve4i", - "kH4i", - "veci", + "v33vvh", + "vecs9", + "veF4", + "uec4i", + "eIKK", + "ve4J", "vec4i", - "oo4rr", - "JJc4", - "vcCC0", - "xAA99F", - "veccu", - "vec4S", + "vSSCCXXi", + "JecWW6ZZ", + "ecd5", + "vBBcBU", + "JJ0c411", + "vecttu", "vec4u", - "vocBB", - "ec4u", - "veemm", + "vttc", + "veL4u", + "v1c4u", }; for (auto _ : state) { for (auto* str : kStrings) { @@ -528,7 +710,7 @@ void BuiltinParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(BuiltinParser); diff --git a/src/tint/builtin/builtin_test.cc b/src/tint/builtin/builtin_test.cc index 9ff8d99228..ab8c4c6fe0 100644 --- a/src/tint/builtin/builtin_test.cc +++ b/src/tint/builtin/builtin_test.cc @@ -43,6 +43,32 @@ inline std::ostream& operator<<(std::ostream& out, Case c) { } 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}, {"array", Builtin::kArray}, {"atomic", Builtin::kAtomic}, @@ -116,216 +142,294 @@ static constexpr Case kValidCases[] = { }; static constexpr Case kInvalidCases[] = { - {"__acked_veccc", Builtin::kUndefined}, - {"_pac3ed_v3", Builtin::kUndefined}, - {"__packeV_vec3", Builtin::kUndefined}, - {"arra1", Builtin::kUndefined}, - {"qqrJy", Builtin::kUndefined}, - {"arrll7y", Builtin::kUndefined}, - {"atppmHHc", Builtin::kUndefined}, - {"cto", Builtin::kUndefined}, - {"abGmi", Builtin::kUndefined}, - {"bovii", Builtin::kUndefined}, - {"boWWl", Builtin::kUndefined}, - {"Mxxl", Builtin::kUndefined}, - {"fgg", Builtin::kUndefined}, - {"X", Builtin::kUndefined}, - {"316", Builtin::kUndefined}, - {"fE2", Builtin::kUndefined}, - {"fPTT", Builtin::kUndefined}, - {"dxx2", Builtin::kUndefined}, - {"4432", Builtin::kUndefined}, - {"iSVV2", Builtin::kUndefined}, - {"RR2", Builtin::kUndefined}, - {"at292", Builtin::kUndefined}, - {"mat2x", Builtin::kUndefined}, - {"Vat2OR2", Builtin::kUndefined}, - {"ma2xyf", Builtin::kUndefined}, - {"llnarr2772f", Builtin::kUndefined}, - {"mat24200", Builtin::kUndefined}, - {"a2ooh", Builtin::kUndefined}, - {"zz2x2h", Builtin::kUndefined}, - {"miitppx1", Builtin::kUndefined}, - {"maXX2x3", Builtin::kUndefined}, - {"55IIt2nn99", Builtin::kUndefined}, - {"aHHrrt2xSS", Builtin::kUndefined}, - {"makkf", Builtin::kUndefined}, - {"jatgRx", Builtin::kUndefined}, - {"mb2x3", Builtin::kUndefined}, - {"mat2xjh", Builtin::kUndefined}, - {"at2x3h", Builtin::kUndefined}, - {"q2x3h", Builtin::kUndefined}, - {"mNN2x4", Builtin::kUndefined}, - {"mavv4", Builtin::kUndefined}, - {"maQQx4", Builtin::kUndefined}, - {"maffxr", Builtin::kUndefined}, - {"mat2xjf", Builtin::kUndefined}, - {"mNNw2x48", Builtin::kUndefined}, - {"matx4h", Builtin::kUndefined}, - {"mrrt2x4h", Builtin::kUndefined}, - {"Gat2x4h", Builtin::kUndefined}, - {"matFFx2", Builtin::kUndefined}, - {"mtx", Builtin::kUndefined}, - {"mrrt3x", Builtin::kUndefined}, + {"__atomic_compareexchangeccresult_i32", Builtin::kUndefined}, + {"__atoml3_compare_exchane_resulti2", Builtin::kUndefined}, + {"__atomic_compare_Vxchange_result_i32", Builtin::kUndefined}, + {"__atomic_com1are_exchange_result_u32", Builtin::kUndefined}, + {"__atomic_qqompare_exchage_resulJ_u32", Builtin::kUndefined}, + {"__atllmic_compare_exchange_result_u377", Builtin::kUndefined}, + {"qpp_frexp_resultHHbstract", Builtin::kUndefined}, + {"__fep_esulv_abstract", Builtin::kUndefined}, + {"__Gbexp_resul_abstract", Builtin::kUndefined}, + {"_vfrexp_resiilt_f16", Builtin::kUndefined}, + {"__fr8xp_resultWWf16", Builtin::kUndefined}, + {"__frxp_result_fMxx", Builtin::kUndefined}, + {"gg_fXexp_reslt_f32", Builtin::kUndefined}, + {"__frXxpresul_V32", Builtin::kUndefined}, + {"__frexp_r3sult_f32", Builtin::kUndefined}, + {"__frexpEresult_vec2_abstract", Builtin::kUndefined}, + {"__frex_rPPsult_vTTc2_abstract", Builtin::kUndefined}, + {"__frexp_resuddt_ec2_xxbstract", Builtin::kUndefined}, + {"__frexp_result_ve442_f16", Builtin::kUndefined}, + {"_SSfrexp_resulVV_vec2_f16", Builtin::kUndefined}, + {"__fRxpRr22sult_vec2_f16", Builtin::kUndefined}, + {"__frexp_res9lt_vec_fF2", Builtin::kUndefined}, + {"__frexp_result_ve2_f32", Builtin::kUndefined}, + {"_OOfrexp_result_VeHRRf32", Builtin::kUndefined}, + {"__frexp_reyult_vec3_absract", Builtin::kUndefined}, + {"__frexp_re77ulll_vecG_arrnstract", Builtin::kUndefined}, + {"__4rexp_result_vec3_00bstract", Builtin::kUndefined}, + {"__oorxp_result_vec316", Builtin::kUndefined}, + {"zz_frexp_esult_ec3_f16", Builtin::kUndefined}, + {"__iirex11_result_vp3_f16", Builtin::kUndefined}, + {"__frXXxp_result_vec3_f32", Builtin::kUndefined}, + {"__fnnexp99resIIlt_vec3_f355", Builtin::kUndefined}, + {"__faSSerrp_result_vHHc3_fY2", Builtin::kUndefined}, + {"__freHp_resutve4_abstkkact", Builtin::kUndefined}, + {"jfrexpgresult_veRR4_abstrac", Builtin::kUndefined}, + {"__frexp_resul_vec4_absbrac", Builtin::kUndefined}, + {"_jfrexp_result_vec4_f16", Builtin::kUndefined}, + {"__frexp_resultvec4_f16", Builtin::kUndefined}, + {"__freqpresultvec4_f16", Builtin::kUndefined}, + {"__frexNN_result_vec_f32", Builtin::kUndefined}, + {"__frexp_resvvlt_vc4_f3", Builtin::kUndefined}, + {"__frexp_esult_vec4_f3QQ", Builtin::kUndefined}, + {"rmodf_reffultabstract", Builtin::kUndefined}, + {"__jodf_result_abstract", Builtin::kUndefined}, + {"_mNNwdf_r2sult8abstract", Builtin::kUndefined}, + {"__mdf_result_f16", Builtin::kUndefined}, + {"__modrr_result_f16", Builtin::kUndefined}, + {"__mGdf_result_f16", Builtin::kUndefined}, + {"__modf_resulFF_f32", Builtin::kUndefined}, + {"__modf_eult_E3", 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}, - {"Da3xJJf", Builtin::kUndefined}, - {"ma82", Builtin::kUndefined}, - {"m11k2", Builtin::kUndefined}, - {"matx2h", Builtin::kUndefined}, - {"maJx2h", Builtin::kUndefined}, - {"cat3x3", Builtin::kUndefined}, - {"mat3O3", Builtin::kUndefined}, - {"ttKavv3x__", Builtin::kUndefined}, - {"xx83x3f", Builtin::kUndefined}, - {"__qatF3", Builtin::kUndefined}, - {"matqx3f", Builtin::kUndefined}, - {"33atOx3h", Builtin::kUndefined}, - {"mtt63x9oQQ", Builtin::kUndefined}, - {"ma3x66h", Builtin::kUndefined}, - {"66aOzx4", Builtin::kUndefined}, - {"myyt3x4", Builtin::kUndefined}, - {"HHZx4", Builtin::kUndefined}, - {"4WWt3q4f", Builtin::kUndefined}, - {"mOO3x4f", Builtin::kUndefined}, - {"oatY4f", Builtin::kUndefined}, - {"matx", Builtin::kUndefined}, - {"ma3xFh", Builtin::kUndefined}, - {"at3x4w", Builtin::kUndefined}, - {"ma4Gf", Builtin::kUndefined}, - {"qatKKx2", Builtin::kUndefined}, - {"mmmt4x2", Builtin::kUndefined}, - {"at4x2f", Builtin::kUndefined}, - {"mt4x2q", Builtin::kUndefined}, - {"mat4xbb", Builtin::kUndefined}, - {"mi4x2h", Builtin::kUndefined}, - {"maOO4xq", Builtin::kUndefined}, - {"matTvvx2h", Builtin::kUndefined}, - {"mat4FF3", Builtin::kUndefined}, - {"mtQ00P", Builtin::kUndefined}, - {"maP4x3", Builtin::kUndefined}, - {"ma774xss", Builtin::kUndefined}, - {"RRCbb4x3f", Builtin::kUndefined}, - {"mXXt4x3f", Builtin::kUndefined}, - {"qaCC4xOOh", Builtin::kUndefined}, - {"ma4s3L", Builtin::kUndefined}, - {"mXt4x3h", Builtin::kUndefined}, - {"mat4x", Builtin::kUndefined}, - {"qqt4", Builtin::kUndefined}, - {"mat4x22", Builtin::kUndefined}, - {"myzz40XX", Builtin::kUndefined}, - {"matVViP", Builtin::kUndefined}, - {"mannC4f", Builtin::kUndefined}, - {"atx4AHHq", Builtin::kUndefined}, - {"may4x4", Builtin::kUndefined}, - {"aOOOZZh", Builtin::kUndefined}, - {"V", Builtin::kUndefined}, - {"ptf__", Builtin::kUndefined}, - {"4lMT", Builtin::kUndefined}, - {"sNNmplg", Builtin::kUndefined}, - {"uubpXXer", Builtin::kUndefined}, - {"samler", Builtin::kUndefined}, - {"m88ler_cQmparisoK", Builtin::kUndefined}, - {"qa9ler_comparison", Builtin::kUndefined}, - {"sampler_comparis11n", Builtin::kUndefined}, - {"teiiu22eF1d", Builtin::kUndefined}, - {"tex77ur_1d", Builtin::kUndefined}, - {"te2urNN_1d", Builtin::kUndefined}, - {"texturVV_2d", Builtin::kUndefined}, - {"WWFxtu11e_wd", Builtin::kUndefined}, - {"txture_2ww", Builtin::kUndefined}, - {"texture_2d_arrDy", Builtin::kUndefined}, - {"teKtre_2d_array", Builtin::kUndefined}, - {"texhure_2fra11raPP", Builtin::kUndefined}, - {"texture3d", Builtin::kUndefined}, - {"texture_3YY", Builtin::kUndefined}, - {"HHtxtrkk_3d", Builtin::kUndefined}, - {"texrrure_cube", Builtin::kUndefined}, - {"tssxturWW_cue", Builtin::kUndefined}, - {"teYure_cube", Builtin::kUndefined}, - {"txture_Lufe_arraq", Builtin::kUndefined}, - {"te22ture_uuubevvfray", Builtin::kUndefined}, - {"texturecube_aray", Builtin::kUndefined}, - {"texture_Yepth_2", Builtin::kUndefined}, - {"teytYYEe_77epth_2d", Builtin::kUndefined}, - {"teMture_deootd2d", Builtin::kUndefined}, - {"texMMre_depth_2d_array", Builtin::kUndefined}, - {"texture_depth_2d_arra55", Builtin::kUndefined}, - {"texture_deh_2d_aNray", Builtin::kUndefined}, - {"te3ture_dpth_cO3be", Builtin::kUndefined}, - {"texture_depth_cub3", Builtin::kUndefined}, - {"Iexturedepth_cume", Builtin::kUndefined}, - {"texture_depthnncube_Krrry", Builtin::kUndefined}, - {"texture_dth_XXube_rra", Builtin::kUndefined}, - {"textIre_depph_ubeLLarray", Builtin::kUndefined}, - {"txtfre_depthmultisampled_2d", Builtin::kUndefined}, - {"texURuYe_Depthmultisampled_2d", Builtin::kUndefined}, - {"texture_depth_multisamphed_2d", Builtin::kUndefined}, - {"teqtureuIIextnal", Builtin::kUndefined}, - {"texture_externaH", Builtin::kUndefined}, - {"texre_externaQvv", Builtin::kUndefined}, - {"textureemultismp66ed_d", Builtin::kUndefined}, - {"tW7trO_multisampled_2d", Builtin::kUndefined}, - {"texture_mult550ampled_2DD", Builtin::kUndefined}, - {"teHture_storIIge_1d", Builtin::kUndefined}, - {"textue_storage_1d", Builtin::kUndefined}, - {"rexture_storae_1d", Builtin::kUndefined}, - {"texture_stolage_2d", Builtin::kUndefined}, - {"txture_JJtGrgtt_2d", Builtin::kUndefined}, - {"yexture_storage2d", Builtin::kUndefined}, - {"texture_storage_2d_rray", Builtin::kUndefined}, - {"texture_IItorage_2d_BBrray", Builtin::kUndefined}, - {"33exture_TTtorge_Kd_ar88ay", Builtin::kUndefined}, - {"texSnnYUUure_storage_3d", Builtin::kUndefined}, - {"textuxe_5torCCdZ_3d", Builtin::kUndefined}, - {"tkkxture_storaqe_3d", Builtin::kUndefined}, - {"5i00", Builtin::kUndefined}, - {"unII2", Builtin::kUndefined}, - {"cc", Builtin::kUndefined}, - {"KK", Builtin::kUndefined}, - {"66ec2", Builtin::kUndefined}, - {"PPEK", Builtin::kUndefined}, - {"vexxf", Builtin::kUndefined}, - {"qec2f", Builtin::kUndefined}, - {"veSyMMr", Builtin::kUndefined}, - {"v2u", Builtin::kUndefined}, - {"ec", Builtin::kUndefined}, - {"5eFF2h", Builtin::kUndefined}, - {"rrecz44", Builtin::kUndefined}, - {"vWW", Builtin::kUndefined}, - {"XJecCZZ", Builtin::kUndefined}, - {"vePP2", Builtin::kUndefined}, - {"vec2c", Builtin::kUndefined}, - {"ve6ll2u", Builtin::kUndefined}, - {"vey99", Builtin::kUndefined}, - {"vKKc3", Builtin::kUndefined}, - {"x_3", Builtin::kUndefined}, - {"Ky3", Builtin::kUndefined}, - {"zek3f", Builtin::kUndefined}, - {"veKSf", Builtin::kUndefined}, - {"vc3h", Builtin::kUndefined}, - {"ec3VV", Builtin::kUndefined}, - {"IAAc3h", Builtin::kUndefined}, - {"jbR", Builtin::kUndefined}, - {"veY4", Builtin::kUndefined}, - {"ec3i", Builtin::kUndefined}, - {"vc911", Builtin::kUndefined}, - {"mmccu", Builtin::kUndefined}, - {"vJJcu", Builtin::kUndefined}, - {"lDCfcU", Builtin::kUndefined}, - {"veg4", Builtin::kUndefined}, - {"CC", Builtin::kUndefined}, - {"ec4f", Builtin::kUndefined}, - {"vIc__f", Builtin::kUndefined}, - {"ePPtt", Builtin::kUndefined}, - {"v3dc4h", Builtin::kUndefined}, - {"vcyyh", Builtin::kUndefined}, - {"u4", Builtin::kUndefined}, - {"v03nni", Builtin::kUndefined}, - {"Cuuecnv", Builtin::kUndefined}, - {"vX4ll", Builtin::kUndefined}, - {"vocppu", Builtin::kUndefined}, - {"vwwc4", Builtin::kUndefined}, - {"veuug", Builtin::kUndefined}, + {"YYat3f", Builtin::kUndefined}, + {"may3x2EYY", Builtin::kUndefined}, + {"da3xMoh", Builtin::kUndefined}, + {"matMMx2", Builtin::kUndefined}, + {"mat3x55h", Builtin::kUndefined}, + {"maN3", Builtin::kUndefined}, + {"ma33x3", Builtin::kUndefined}, + {"mt3x3", Builtin::kUndefined}, + {"mm66Issf", Builtin::kUndefined}, + {"mat3x1f", Builtin::kUndefined}, + {"Xt3x3", Builtin::kUndefined}, + {"LatIx3h", Builtin::kUndefined}, + {"at3fh", Builtin::kUndefined}, + {"mYtURD3", Builtin::kUndefined}, + {"mah3x4", Builtin::kUndefined}, + {"muqII4", Builtin::kUndefined}, + {"mat3xH", Builtin::kUndefined}, + {"at3QQvv", Builtin::kUndefined}, + {"at66ef", Builtin::kUndefined}, + {"ma7O4f", Builtin::kUndefined}, + {"m55t3x0DD", Builtin::kUndefined}, + {"maH3x4II", Builtin::kUndefined}, + {"at3x4", Builtin::kUndefined}, + {"ma994x2", Builtin::kUndefined}, + {"mWWt4Gt2", Builtin::kUndefined}, + {"ay42", Builtin::kUndefined}, + {"mt4x2f", Builtin::kUndefined}, + {"IIaBB4x2f", Builtin::kUndefined}, + {"TTat4x833", Builtin::kUndefined}, + {"ddUUnntYYx2h", Builtin::kUndefined}, + {"m5CCxxdZ", Builtin::kUndefined}, + {"matkkq2h", Builtin::kUndefined}, + {"5iitp00", Builtin::kUndefined}, + {"mnntIIx3", Builtin::kUndefined}, + {"ccaKx", Builtin::kUndefined}, + {"m43KK", Builtin::kUndefined}, + {"mat66x3f", Builtin::kUndefined}, + {"Et4PP3K", Builtin::kUndefined}, + {"xxatx3h", Builtin::kUndefined}, + {"qat4x3h", Builtin::kUndefined}, + {"MMayySrxh", Builtin::kUndefined}, + {"uat4", Builtin::kUndefined}, + {"tx4", Builtin::kUndefined}, + {"ma54FF4", Builtin::kUndefined}, + {"rra444z4f", Builtin::kUndefined}, + {"matWW", Builtin::kUndefined}, + {"CatZJXx4f", Builtin::kUndefined}, + {"maPPx4h", Builtin::kUndefined}, + {"mat4c4h", Builtin::kUndefined}, + {"matPPll6h", Builtin::kUndefined}, + {"9tyy", Builtin::kUndefined}, + {"ptKK", Builtin::kUndefined}, + {"x_", Builtin::kUndefined}, + {"ayKer", Builtin::kUndefined}, + {"szmpVek", Builtin::kUndefined}, + {"sampqeK", Builtin::kUndefined}, + {"sampler_comparisn", Builtin::kUndefined}, + {"sapler_comparisVVn", Builtin::kUndefined}, + {"samplerIcompaAUison", Builtin::kUndefined}, + {"jexurbRd", Builtin::kUndefined}, + {"exure_YYd", Builtin::kUndefined}, + {"exture_1d", Builtin::kUndefined}, + {"texxxur_1d", Builtin::kUndefined}, + {"tJxucce_2d", Builtin::kUndefined}, + {"texure_JJd", Builtin::kUndefined}, + {"lDexture_fCC_arraU", Builtin::kUndefined}, + {"tegture_2d_array", Builtin::kUndefined}, + {"teCCure2d_arra", Builtin::kUndefined}, + {"textue_3d", Builtin::kUndefined}, + {"tIx__ure_3d", Builtin::kUndefined}, + {"texurettPP", Builtin::kUndefined}, + {"tddx3ure_cube", Builtin::kUndefined}, + {"teKyyur_cube", Builtin::kUndefined}, + {"teturecub", Builtin::kUndefined}, + {"textinne_c03e_array", Builtin::kUndefined}, + {"nextCCruuvcubK_array", Builtin::kUndefined}, + {"tXxturellcbe_array", Builtin::kUndefined}, + {"tppxture_depth_2d", Builtin::kUndefined}, + {"txture_deptww_2d", Builtin::kUndefined}, + {"gexturedemmthuu2", Builtin::kUndefined}, + {"texmmre_depthaa2daray", Builtin::kUndefined}, + {"texture_RRepth_Td_ccZray", Builtin::kUndefined}, + {"text88re_depthTOd_array", Builtin::kUndefined}, + {"texture_depth_cm00e", Builtin::kUndefined}, + {"texture_Bmepth_cube", Builtin::kUndefined}, + {"Mextre_ppeph_cube", Builtin::kUndefined}, + {"texturOO_depth_cub_array", Builtin::kUndefined}, + {"GeGGture_depthcube_array", Builtin::kUndefined}, + {"texture11Hdepth_cube_array", Builtin::kUndefined}, + {"textu6e_FFepth_multiameeled_2d", Builtin::kUndefined}, + {"texture_epth_mltisampled_2d", Builtin::kUndefined}, + {"texture_depth_mullKsaiipled_2d", Builtin::kUndefined}, + {"texture_extenal", Builtin::kUndefined}, + {"IIext99reexvvernal", Builtin::kUndefined}, + {"texture_externl", Builtin::kUndefined}, + {"texture_mhltisampled_2d", Builtin::kUndefined}, + {"texturemPllltisampzzed_2d", Builtin::kUndefined}, + {"exture_mltisamed_2d", Builtin::kUndefined}, + {"texture_qqtoragff_1", Builtin::kUndefined}, + {"textre_JJddorage_1W", Builtin::kUndefined}, + {"XXrxture_storae1zz", Builtin::kUndefined}, + {"texturestorag2_2d", Builtin::kUndefined}, + {"yyNxture_storage_2d", Builtin::kUndefined}, + {"etue_storage_2OO", Builtin::kUndefined}, + {"reutuPe_storZgeE2d_array", Builtin::kUndefined}, + {"texlure_storddeee_d_22rray", Builtin::kUndefined}, + {"texture_mtorage_2V_a9ra", Builtin::kUndefined}, + {"teII1re_storage_3d", Builtin::kUndefined}, + {"texture_storagb_3d", Builtin::kUndefined}, + {"texizrestorge73d", Builtin::kUndefined}, + {"u3oi", Builtin::kUndefined}, + {"3", Builtin::kUndefined}, + {"S2", Builtin::kUndefined}, + {"e22", Builtin::kUndefined}, + {"1eC2", Builtin::kUndefined}, + {"vf8c2", Builtin::kUndefined}, + {"c2f", Builtin::kUndefined}, + {"JJecSSf", Builtin::kUndefined}, + {"92f", Builtin::kUndefined}, + {"vbbJJ2TT", Builtin::kUndefined}, + {"e66h", Builtin::kUndefined}, + {"u662h", Builtin::kUndefined}, + {"vW2i", Builtin::kUndefined}, + {"v2i", Builtin::kUndefined}, + {"veci", Builtin::kUndefined}, + {"rec2u", Builtin::kUndefined}, + {"2ec2B", Builtin::kUndefined}, + {"vcBBu", Builtin::kUndefined}, + {"veRR", Builtin::kUndefined}, + {"VLL0", Builtin::kUndefined}, + {"KOe3", Builtin::kUndefined}, + {"vgwcf", Builtin::kUndefined}, + {"vLphf", Builtin::kUndefined}, + {"eiiEf", Builtin::kUndefined}, + {"ec3h", Builtin::kUndefined}, + {"UU883", Builtin::kUndefined}, + {"rrecvvh", Builtin::kUndefined}, + {"ecmm", Builtin::kUndefined}, + {"vec4j", Builtin::kUndefined}, + {"vec3X", Builtin::kUndefined}, + {"vec38", Builtin::kUndefined}, + {"vecvEE", Builtin::kUndefined}, + {"z99ci", Builtin::kUndefined}, + {"JJGeQQ4", Builtin::kUndefined}, + {"ssec4", Builtin::kUndefined}, + {"PecK", Builtin::kUndefined}, + {"tpc4f", Builtin::kUndefined}, + {"vec", Builtin::kUndefined}, + {"MMec4f", Builtin::kUndefined}, + {"vJJc40", Builtin::kUndefined}, + {"8c", Builtin::kUndefined}, + {"vecggKh", Builtin::kUndefined}, + {"vecfi", Builtin::kUndefined}, + {"vec47Q", Builtin::kUndefined}, + {"veY4i", Builtin::kUndefined}, + {"keSu", Builtin::kUndefined}, + {"n422", Builtin::kUndefined}, + {"vFFu", Builtin::kUndefined}, }; using BuiltinParseTest = testing::TestWithParam; diff --git a/src/tint/builtin/builtin_value_bench.cc b/src/tint/builtin/builtin_value_bench.cc index 4d2e562308..28c01ac50d 100644 --- a/src/tint/builtin/builtin_value_bench.cc +++ b/src/tint/builtin/builtin_value_bench.cc @@ -129,7 +129,7 @@ void BuiltinValueParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(BuiltinValueParser); diff --git a/src/tint/builtin/diagnostic_rule_bench.cc b/src/tint/builtin/diagnostic_rule_bench.cc index fb6058098a..efa6977a1e 100644 --- a/src/tint/builtin/diagnostic_rule_bench.cc +++ b/src/tint/builtin/diagnostic_rule_bench.cc @@ -41,7 +41,7 @@ void CoreDiagnosticRuleParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(CoreDiagnosticRuleParser); @@ -56,7 +56,7 @@ void ChromiumDiagnosticRuleParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(ChromiumDiagnosticRuleParser); diff --git a/src/tint/builtin/diagnostic_severity_bench.cc b/src/tint/builtin/diagnostic_severity_bench.cc index 13c3ff1a74..35fe04646d 100644 --- a/src/tint/builtin/diagnostic_severity_bench.cc +++ b/src/tint/builtin/diagnostic_severity_bench.cc @@ -42,7 +42,7 @@ void DiagnosticSeverityParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(DiagnosticSeverityParser); diff --git a/src/tint/builtin/extension_bench.cc b/src/tint/builtin/extension_bench.cc index b3e410e591..7a50281c7f 100644 --- a/src/tint/builtin/extension_bench.cc +++ b/src/tint/builtin/extension_bench.cc @@ -80,7 +80,7 @@ void ExtensionParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(ExtensionParser); diff --git a/src/tint/builtin/interpolation_sampling_bench.cc b/src/tint/builtin/interpolation_sampling_bench.cc index 4fb7e09869..deb4de0b84 100644 --- a/src/tint/builtin/interpolation_sampling_bench.cc +++ b/src/tint/builtin/interpolation_sampling_bench.cc @@ -40,7 +40,7 @@ void InterpolationSamplingParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(InterpolationSamplingParser); diff --git a/src/tint/builtin/interpolation_type_bench.cc b/src/tint/builtin/interpolation_type_bench.cc index 5098f26973..70b5961a2c 100644 --- a/src/tint/builtin/interpolation_type_bench.cc +++ b/src/tint/builtin/interpolation_type_bench.cc @@ -43,7 +43,7 @@ void InterpolationTypeParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(InterpolationTypeParser); diff --git a/src/tint/builtin/texel_format_bench.cc b/src/tint/builtin/texel_format_bench.cc index 5d202eacd2..0773b9a813 100644 --- a/src/tint/builtin/texel_format_bench.cc +++ b/src/tint/builtin/texel_format_bench.cc @@ -62,7 +62,7 @@ void TexelFormatParser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK(TexelFormatParser); diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def index ee4b517b9b..0d667ab34e 100644 --- a/src/tint/intrinsics.def +++ b/src/tint/intrinsics.def @@ -218,6 +218,32 @@ enum builtin_type { // Internal types. __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 diff --git a/src/tint/program_builder.h b/src/tint/program_builder.h index b4f5eea1b0..49b14fc2a5 100644 --- a/src/tint/program_builder.h +++ b/src/tint/program_builder.h @@ -1470,7 +1470,7 @@ class ProgramBuilder { /// @param name the symbol string /// @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 /// @return a Symbol with the given enum value diff --git a/src/tint/resolver/builtin_structs.cc b/src/tint/resolver/builtin_structs.cc index 02f170897b..ce4f38c7d9 100644 --- a/src/tint/resolver/builtin_structs.cc +++ b/src/tint/resolver/builtin_structs.cc @@ -29,12 +29,12 @@ namespace tint::resolver { namespace { struct NameAndType { - std::string name; + std::string_view name; const type::Type* type; }; sem::Struct* BuildStruct(ProgramBuilder& b, - std::string name, + builtin::Builtin name, std::initializer_list member_names_and_types) { uint32_t offset = 0; uint32_t max_align = 0; @@ -69,45 +69,64 @@ sem::Struct* BuildStruct(ProgramBuilder& b, } } // 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) { return Switch( ty, [&](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*) { - return BuildStruct(b, "__modf_result_f16", {{"fract", ty}, {"whole", ty}}); + return BuildStruct(b, builtin::Builtin::kModfResultF16, {{"fract", ty}, {"whole", ty}}); }, [&](const type::AbstractFloat*) { - auto* abstract = - BuildStruct(b, "__modf_result_abstract", {{"fract", ty}, {"whole", ty}}); + auto* abstract = BuildStruct(b, builtin::Builtin::kModfResultAbstract, + {{"fract", ty}, {"whole", ty}}); auto* f32 = b.create(); auto* f16 = b.create(); abstract->SetConcreteTypes(utils::Vector{ - BuildStruct(b, "__modf_result_f32", {{"fract", f32}, {"whole", f32}}), - BuildStruct(b, "__modf_result_f16", {{"fract", f16}, {"whole", f16}}), + BuildStruct(b, builtin::Builtin::kModfResultF32, {{"fract", f32}, {"whole", f32}}), + BuildStruct(b, builtin::Builtin::kModfResultF16, {{"fract", f16}, {"whole", f16}}), }); return abstract; }, [&](const type::Vector* vec) { auto width = vec->Width(); - auto prefix = "__modf_result_vec" + std::to_string(width); return Switch( vec->type(), // [&](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*) { - return BuildStruct(b, prefix + "_f16", {{"fract", vec}, {"whole", vec}}); + return BuildStruct(b, kModfVecF16Names[width - 2], + {{"fract", vec}, {"whole", vec}}); }, [&](const type::AbstractFloat*) { auto* vec_f32 = b.create(b.create(), width); auto* vec_f16 = b.create(b.create(), width); - auto* abstract = - BuildStruct(b, prefix + "_abstract", {{"fract", vec}, {"whole", vec}}); + auto* abstract = BuildStruct(b, kModfVecAbstractNames[width - 2], + {{"fract", vec}, {"whole", vec}}); abstract->SetConcreteTypes(utils::Vector{ - BuildStruct(b, prefix + "_f32", {{"fract", vec_f32}, {"whole", vec_f32}}), - BuildStruct(b, prefix + "_f16", {{"fract", vec_f16}, {"whole", vec_f16}}), + BuildStruct(b, kModfVecF32Names[width - 2], + {{"fract", vec_f32}, {"whole", vec_f32}}), + BuildStruct(b, kModfVecF16Names[width - 2], + {{"fract", vec_f16}, {"whole", vec_f16}}), }); 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) { return Switch( ty, // [&](const type::F32*) { auto* i32 = b.create(); - return BuildStruct(b, "__frexp_result_f32", {{"fract", ty}, {"exp", i32}}); + return BuildStruct(b, builtin::Builtin::kFrexpResultF32, {{"fract", ty}, {"exp", i32}}); }, [&](const type::F16*) { auto* i32 = b.create(); - return BuildStruct(b, "__frexp_result_f16", {{"fract", ty}, {"exp", i32}}); + return BuildStruct(b, builtin::Builtin::kFrexpResultF16, {{"fract", ty}, {"exp", i32}}); }, [&](const type::AbstractFloat*) { auto* f32 = b.create(); auto* f16 = b.create(); auto* i32 = b.create(); auto* ai = b.create(); - auto* abstract = - BuildStruct(b, "__frexp_result_abstract", {{"fract", ty}, {"exp", ai}}); + auto* abstract = BuildStruct(b, builtin::Builtin::kFrexpResultAbstract, + {{"fract", ty}, {"exp", ai}}); abstract->SetConcreteTypes(utils::Vector{ - BuildStruct(b, "__frexp_result_f32", {{"fract", f32}, {"exp", i32}}), - BuildStruct(b, "__frexp_result_f16", {{"fract", f16}, {"exp", i32}}), + BuildStruct(b, builtin::Builtin::kFrexpResultF32, {{"fract", f32}, {"exp", i32}}), + BuildStruct(b, builtin::Builtin::kFrexpResultF16, {{"fract", f16}, {"exp", i32}}), }); return abstract; }, [&](const type::Vector* vec) { auto width = vec->Width(); - auto prefix = "__frexp_result_vec" + std::to_string(width); return Switch( vec->type(), // [&](const type::F32*) { auto* vec_i32 = b.create(b.create(), 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*) { auto* vec_i32 = b.create(b.create(), 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*) { auto* vec_f32 = b.create(b.create(), width); auto* vec_f16 = b.create(b.create(), width); auto* vec_i32 = b.create(b.create(), width); auto* vec_ai = b.create(b.create(), width); - auto* abstract = - BuildStruct(b, prefix + "_abstract", {{"fract", ty}, {"exp", vec_ai}}); + auto* abstract = BuildStruct(b, kFrexpVecAbstractNames[width - 2], + {{"fract", ty}, {"exp", vec_ai}}); abstract->SetConcreteTypes(utils::Vector{ - BuildStruct(b, prefix + "_f32", {{"fract", vec_f32}, {"exp", vec_i32}}), - BuildStruct(b, prefix + "_f16", {{"fract", vec_f16}, {"exp", vec_i32}}), + BuildStruct(b, kFrexpVecF32Names[width - 2], + {{"fract", vec_f32}, {"exp", vec_i32}}), + BuildStruct(b, kFrexpVecF16Names[width - 2], + {{"fract", vec_f16}, {"exp", vec_i32}}), }); return abstract; }, @@ -186,9 +223,21 @@ sem::Struct* CreateFrexpResult(ProgramBuilder& b, const type::Type* ty) { } sem::Struct* CreateAtomicCompareExchangeResult(ProgramBuilder& b, const type::Type* ty) { - return BuildStruct( - b, "__atomic_compare_exchange_result" + ty->FriendlyName(), - {{"old_value", const_cast(ty)}, {"exchanged", b.create()}}); + return Switch( + ty, // + [&](const type::I32*) { + return BuildStruct(b, builtin::Builtin::kAtomicCompareExchangeResultI32, + {{"old_value", ty}, {"exchanged", b.create()}}); + }, + [&](const type::U32*) { + return BuildStruct(b, builtin::Builtin::kAtomicCompareExchangeResultU32, + {{"old_value", ty}, {"exchanged", b.create()}}); + }, + [&](Default) { + TINT_ICE(Resolver, b.Diagnostics()) + << "unhandled atomic_compare_exchange type: " << b.FriendlyName(ty); + return nullptr; + }); } } // namespace tint::resolver diff --git a/src/tint/resolver/builtin_structs_test.cc b/src/tint/resolver/builtin_structs_test.cc new file mode 100644 index 0000000000..613bcb6ba8 --- /dev/null +++ b/src/tint/resolver/builtin_structs_test.cc @@ -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; + +TEST_P(ResolverBuiltinStructs, Resolve) { + Enable(builtin::Extension::kF16); + + // var p : NAME; + auto* var = GlobalVar("p", ty(GetParam()), builtin::AddressSpace::kPrivate); + + ASSERT_TRUE(r()->Resolve()) << r()->error(); + auto* str = As(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 diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc index 616acc6614..99e68d0cbf 100644 --- a/src/tint/resolver/resolver.cc +++ b/src/tint/resolver/resolver.cc @@ -43,6 +43,7 @@ #include "src/tint/ast/while_statement.h" #include "src/tint/ast/workgroup_attribute.h" #include "src/tint/builtin/builtin.h" +#include "src/tint/resolver/builtin_structs.h" #include "src/tint/resolver/uniformity.h" #include "src/tint/sem/break_if_statement.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* { return TINT_LIKELY(CheckNotTemplated("type", ident)) ? ty : nullptr; }; + auto af = [&] { return b.create(); }; auto f32 = [&] { return b.create(); }; auto i32 = [&] { return b.create(); }; auto u32 = [&] { return b.create(); }; @@ -2737,9 +2739,60 @@ type::Type* Resolver::BuiltinType(builtin::Builtin builtin_ty, const ast::Identi return storage_texture(type::TextureDimension::k2dArray); case builtin::Builtin::kTextureStorage3D: return storage_texture(type::TextureDimension::k3d); - case builtin::Builtin::kPackedVec3: { + case builtin::Builtin::kPackedVec3: 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: break; } diff --git a/src/tint/templates/enums.tmpl.inc b/src/tint/templates/enums.tmpl.inc index ad4694220e..adc06e05f0 100644 --- a/src/tint/templates/enums.tmpl.inc +++ b/src/tint/templates/enums.tmpl.inc @@ -191,7 +191,7 @@ void {{$enum}}Parser(::benchmark::State& state) { benchmark::DoNotOptimize(result); } } -} +} // NOLINT(readability/fn_size) BENCHMARK({{$enum}}Parser); {{- end -}} diff --git a/test/tint/bug/chromium/1430309.wgsl b/test/tint/bug/chromium/1430309.wgsl new file mode 100644 index 0000000000..fee1a69f80 --- /dev/null +++ b/test/tint/bug/chromium/1430309.wgsl @@ -0,0 +1,11 @@ +struct frexp_result_f32 { + f : f32, +} + +var a : frexp_result_f32; +var b = frexp(1); + +@fragment +fn main() -> @location(0) vec4f { + return vec4f(a.f, b.fract, 0, 0); +} diff --git a/test/tint/bug/chromium/1430309.wgsl.expected.dxc.hlsl b/test/tint/bug/chromium/1430309.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..9f62ac4b7d --- /dev/null +++ b/test/tint/bug/chromium/1430309.wgsl.expected.dxc.hlsl @@ -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; +} diff --git a/test/tint/bug/chromium/1430309.wgsl.expected.fxc.hlsl b/test/tint/bug/chromium/1430309.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..9f62ac4b7d --- /dev/null +++ b/test/tint/bug/chromium/1430309.wgsl.expected.fxc.hlsl @@ -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; +} diff --git a/test/tint/bug/chromium/1430309.wgsl.expected.glsl b/test/tint/bug/chromium/1430309.wgsl.expected.glsl new file mode 100644 index 0000000000..e06c7d277b --- /dev/null +++ b/test/tint/bug/chromium/1430309.wgsl.expected.glsl @@ -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; +} diff --git a/test/tint/bug/chromium/1430309.wgsl.expected.msl b/test/tint/bug/chromium/1430309.wgsl.expected.msl new file mode 100644 index 0000000000..5948214511 --- /dev/null +++ b/test/tint/bug/chromium/1430309.wgsl.expected.msl @@ -0,0 +1,34 @@ +#include + +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; +} + diff --git a/test/tint/bug/chromium/1430309.wgsl.expected.spvasm b/test/tint/bug/chromium/1430309.wgsl.expected.spvasm new file mode 100644 index 0000000000..8357eb20d3 --- /dev/null +++ b/test/tint/bug/chromium/1430309.wgsl.expected.spvasm @@ -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 diff --git a/test/tint/bug/chromium/1430309.wgsl.expected.wgsl b/test/tint/bug/chromium/1430309.wgsl.expected.wgsl new file mode 100644 index 0000000000..a4ce7b9ae7 --- /dev/null +++ b/test/tint/bug/chromium/1430309.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +struct frexp_result_f32 { + f : f32, +} + +var a : frexp_result_f32; + +var b = frexp(1); + +@fragment +fn main() -> @location(0) vec4f { + return vec4f(a.f, b.fract, 0, 0); +} diff --git a/test/tint/bug/tint/1573.wgsl.expected.glsl b/test/tint/bug/tint/1573.wgsl.expected.glsl index 55803e2858..55ff9c9058 100644 --- a/test/tint/bug/tint/1573.wgsl.expected.glsl +++ b/test/tint/bug/tint/1573.wgsl.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -12,10 +12,10 @@ layout(binding = 0, std430) buffer a_block_ssbo { void tint_symbol() { 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.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; diff --git a/test/tint/bug/tint/1573.wgsl.expected.msl b/test/tint/bug/tint/1573.wgsl.expected.msl index 0aef1728fa..ef0c0b8964 100644 --- a/test/tint/bug/tint/1573.wgsl.expected.msl +++ b/test/tint/bug/tint/1573.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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)]]) { 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; } diff --git a/test/tint/bug/tint/1573.wgsl.expected.spvasm b/test/tint/bug/tint/1573.wgsl.expected.spvasm index c24a87f84a..b81df68e2e 100644 --- a/test/tint/bug/tint/1573.wgsl.expected.spvasm +++ b/test/tint/bug/tint/1573.wgsl.expected.spvasm @@ -12,15 +12,15 @@ OpName %a "a" OpName %main "main" OpName %value "value" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpDecorate %a_block Block OpMemberDecorate %a_block 0 Offset 0 OpDecorate %a DescriptorSet 0 OpDecorate %a Binding 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %a_block = OpTypeStruct %uint %_ptr_StorageBuffer_a_block = OpTypePointer StorageBuffer %a_block @@ -31,7 +31,7 @@ %_ptr_Function_uint = OpTypePointer Function %uint %12 = OpConstantNull %uint %bool = OpTypeBool -%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool +%__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint @@ -43,6 +43,6 @@ %21 = OpLoad %uint %value %22 = OpAtomicCompareExchange %uint %20 %uint_1 %uint_0 %uint_0 %21 %12 %23 = OpIEqual %bool %22 %12 - %13 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %22 %23 + %13 = OpCompositeConstruct %__atomic_compare_exchange_result_u32 %22 %23 OpReturn OpFunctionEnd diff --git a/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl index 1ca845885f..8a4a1bac1a 100644 --- a/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl +++ b/test/tint/bug/tint/1574.wgsl.expected.dxc.hlsl @@ -1,8 +1,8 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -62,39 +62,39 @@ void main_inner(uint local_invocation_index) { } { 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; InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value); atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value; - const atomic_compare_exchange_resultu32 r1 = atomic_result_2; - atomic_compare_exchange_resultu32 atomic_result_3 = (atomic_compare_exchange_resultu32)0; + const atomic_compare_exchange_result_u32 r1 = atomic_result_2; + atomic_compare_exchange_result_u32 atomic_result_3 = (atomic_compare_exchange_result_u32)0; uint atomic_compare_value_1 = 0u; 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; - const atomic_compare_exchange_resultu32 r2 = atomic_result_3; - atomic_compare_exchange_resultu32 atomic_result_4 = (atomic_compare_exchange_resultu32)0; + const atomic_compare_exchange_result_u32 r2 = atomic_result_3; + atomic_compare_exchange_result_u32 atomic_result_4 = (atomic_compare_exchange_result_u32)0; uint atomic_compare_value_2 = 0u; 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; - const atomic_compare_exchange_resultu32 r3 = atomic_result_4; + const atomic_compare_exchange_result_u32 r3 = atomic_result_4; } { 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; 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; - const atomic_compare_exchange_resulti32 r1 = atomic_result_5; - atomic_compare_exchange_resulti32 atomic_result_6 = (atomic_compare_exchange_resulti32)0; + const atomic_compare_exchange_result_i32 r1 = atomic_result_5; + atomic_compare_exchange_result_i32 atomic_result_6 = (atomic_compare_exchange_result_i32)0; int atomic_compare_value_4 = 0; 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; - const atomic_compare_exchange_resulti32 r2 = atomic_result_6; - atomic_compare_exchange_resulti32 atomic_result_7 = (atomic_compare_exchange_resulti32)0; + const atomic_compare_exchange_result_i32 r2 = atomic_result_6; + atomic_compare_exchange_result_i32 atomic_result_7 = (atomic_compare_exchange_result_i32)0; int atomic_compare_value_5 = 0; 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; - const atomic_compare_exchange_resulti32 r3 = atomic_result_7; + const atomic_compare_exchange_result_i32 r3 = atomic_result_7; } } diff --git a/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl index 1ca845885f..8a4a1bac1a 100644 --- a/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl +++ b/test/tint/bug/tint/1574.wgsl.expected.fxc.hlsl @@ -1,8 +1,8 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -62,39 +62,39 @@ void main_inner(uint local_invocation_index) { } { 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; InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value); atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value; - const atomic_compare_exchange_resultu32 r1 = atomic_result_2; - atomic_compare_exchange_resultu32 atomic_result_3 = (atomic_compare_exchange_resultu32)0; + const atomic_compare_exchange_result_u32 r1 = atomic_result_2; + atomic_compare_exchange_result_u32 atomic_result_3 = (atomic_compare_exchange_result_u32)0; uint atomic_compare_value_1 = 0u; 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; - const atomic_compare_exchange_resultu32 r2 = atomic_result_3; - atomic_compare_exchange_resultu32 atomic_result_4 = (atomic_compare_exchange_resultu32)0; + const atomic_compare_exchange_result_u32 r2 = atomic_result_3; + atomic_compare_exchange_result_u32 atomic_result_4 = (atomic_compare_exchange_result_u32)0; uint atomic_compare_value_2 = 0u; 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; - const atomic_compare_exchange_resultu32 r3 = atomic_result_4; + const atomic_compare_exchange_result_u32 r3 = atomic_result_4; } { 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; 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; - const atomic_compare_exchange_resulti32 r1 = atomic_result_5; - atomic_compare_exchange_resulti32 atomic_result_6 = (atomic_compare_exchange_resulti32)0; + const atomic_compare_exchange_result_i32 r1 = atomic_result_5; + atomic_compare_exchange_result_i32 atomic_result_6 = (atomic_compare_exchange_result_i32)0; int atomic_compare_value_4 = 0; 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; - const atomic_compare_exchange_resulti32 r2 = atomic_result_6; - atomic_compare_exchange_resulti32 atomic_result_7 = (atomic_compare_exchange_resulti32)0; + const atomic_compare_exchange_result_i32 r2 = atomic_result_6; + atomic_compare_exchange_result_i32 atomic_result_7 = (atomic_compare_exchange_result_i32)0; int atomic_compare_value_5 = 0; 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; - const atomic_compare_exchange_resulti32 r3 = atomic_result_7; + const atomic_compare_exchange_result_i32 r3 = atomic_result_7; } } diff --git a/test/tint/bug/tint/1574.wgsl.expected.glsl b/test/tint/bug/tint/1574.wgsl.expected.glsl index 3c839bf32a..b7d6dc95a9 100644 --- a/test/tint/bug/tint/1574.wgsl.expected.glsl +++ b/test/tint/bug/tint/1574.wgsl.expected.glsl @@ -1,11 +1,11 @@ #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -29,63 +29,63 @@ void tint_symbol(uint local_invocation_index) { barrier(); { 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.exchanged = atomic_compare_result.old_value == 0u; - atomic_compare_exchange_resultu32 r1 = atomic_compare_result; - atomic_compare_exchange_resultu32 atomic_compare_result_1; + atomic_compare_exchange_result_u32 r1 = atomic_compare_result; + 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.exchanged = atomic_compare_result_1.old_value == 0u; - atomic_compare_exchange_resultu32 r2 = atomic_compare_result_1; - atomic_compare_exchange_resultu32 atomic_compare_result_2; + atomic_compare_exchange_result_u32 r2 = atomic_compare_result_1; + 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.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; - 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.exchanged = atomic_compare_result_3.old_value == 0; - atomic_compare_exchange_resulti32 r1 = atomic_compare_result_3; - atomic_compare_exchange_resulti32 atomic_compare_result_4; + atomic_compare_exchange_result_i32 r1 = atomic_compare_result_3; + 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.exchanged = atomic_compare_result_4.old_value == 0; - atomic_compare_exchange_resulti32 r2 = atomic_compare_result_4; - atomic_compare_exchange_resulti32 atomic_compare_result_5; + atomic_compare_exchange_result_i32 r2 = atomic_compare_result_4; + 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.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; - 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.exchanged = atomic_compare_result_6.old_value == 0u; - atomic_compare_exchange_resultu32 r1 = atomic_compare_result_6; - atomic_compare_exchange_resultu32 atomic_compare_result_7; + atomic_compare_exchange_result_u32 r1 = atomic_compare_result_6; + atomic_compare_exchange_result_u32 atomic_compare_result_7; 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_exchange_resultu32 r2 = atomic_compare_result_7; - atomic_compare_exchange_resultu32 atomic_compare_result_8; + atomic_compare_exchange_result_u32 r2 = atomic_compare_result_7; + atomic_compare_exchange_result_u32 atomic_compare_result_8; 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_exchange_resultu32 r3 = atomic_compare_result_8; + atomic_compare_exchange_result_u32 r3 = atomic_compare_result_8; } { 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.exchanged = atomic_compare_result_9.old_value == 0; - atomic_compare_exchange_resulti32 r1 = atomic_compare_result_9; - atomic_compare_exchange_resulti32 atomic_compare_result_10; + atomic_compare_exchange_result_i32 r1 = atomic_compare_result_9; + atomic_compare_exchange_result_i32 atomic_compare_result_10; 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_exchange_resulti32 r2 = atomic_compare_result_10; - atomic_compare_exchange_resulti32 atomic_compare_result_11; + atomic_compare_exchange_result_i32 r2 = atomic_compare_result_10; + atomic_compare_exchange_result_i32 atomic_compare_result_11; 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_exchange_resulti32 r3 = atomic_compare_result_11; + atomic_compare_exchange_result_i32 r3 = atomic_compare_result_11; } } diff --git a/test/tint/bug/tint/1574.wgsl.expected.msl b/test/tint/bug/tint/1574.wgsl.expected.msl index 80ab0b68be..78b34aa6e7 100644 --- a/test/tint/bug/tint/1574.wgsl.expected.msl +++ b/test/tint/bug/tint/1574.wgsl.expected.msl @@ -2,33 +2,33 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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); { uint value = 42u; - atomic_compare_exchange_resultu32 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_resultu32 const r3 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); + atomic_compare_exchange_result_u32 const r1 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); + atomic_compare_exchange_result_u32 const r2 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); + atomic_compare_exchange_result_u32 const r3 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value); } { int value = 42; - atomic_compare_exchange_resulti32 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_resulti32 const r3 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); + atomic_compare_exchange_result_i32 const r1 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); + atomic_compare_exchange_result_i32 const r2 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); + atomic_compare_exchange_result_i32 const r3 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value); } { uint value = 42u; - atomic_compare_exchange_resultu32 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_resultu32 const r3 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); + atomic_compare_exchange_result_u32 const r1 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); + atomic_compare_exchange_result_u32 const r2 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); + atomic_compare_exchange_result_u32 const r3 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value); } { int value = 42; - atomic_compare_exchange_resulti32 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_resulti32 const r3 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); + atomic_compare_exchange_result_i32 const r1 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); + atomic_compare_exchange_result_i32 const r2 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); + atomic_compare_exchange_result_i32 const r3 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value); } } diff --git a/test/tint/bug/tint/1574.wgsl.expected.spvasm b/test/tint/bug/tint/1574.wgsl.expected.spvasm index 79f74c63d7..b44399fb4f 100644 --- a/test/tint/bug/tint/1574.wgsl.expected.spvasm +++ b/test/tint/bug/tint/1574.wgsl.expected.spvasm @@ -19,13 +19,13 @@ OpName %main_inner "main_inner" OpName %local_invocation_index "local_invocation_index" OpName %value "value" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %value_0 "value" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %value_1 "value" OpName %value_2 "value" OpName %main "main" @@ -38,10 +38,10 @@ OpMemberDecorate %a_i32_block 0 Offset 0 OpDecorate %a_i32 DescriptorSet 0 OpDecorate %a_i32 Binding 1 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input @@ -67,11 +67,11 @@ %uint_264 = OpConstant %uint 264 %uint_42 = OpConstant %uint 42 %_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 %int_42 = OpConstant %int 42 %_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 %113 = OpTypeFunction %void %main_inner = OpFunction %void None %15 @@ -95,59 +95,59 @@ %43 = OpLoad %uint %value %44 = OpAtomicCompareExchange %uint %42 %uint_1 %uint_0 %uint_0 %43 %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 %49 = OpLoad %uint %value %50 = OpAtomicCompareExchange %uint %48 %uint_1 %uint_0 %uint_0 %49 %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 %55 = OpLoad %uint %value %56 = OpAtomicCompareExchange %uint %54 %uint_1 %uint_0 %uint_0 %55 %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 %65 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0 %66 = OpLoad %int %value_0 %67 = OpAtomicCompareExchange %int %65 %uint_1 %uint_0 %uint_0 %66 %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 %72 = OpLoad %int %value_0 %73 = OpAtomicCompareExchange %int %71 %uint_1 %uint_0 %uint_0 %72 %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 %78 = OpLoad %int %value_0 %79 = OpAtomicCompareExchange %int %77 %uint_1 %uint_0 %uint_0 %78 %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 %84 = OpLoad %uint %value_1 %85 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %84 %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 %90 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %89 %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 %95 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %94 %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 %100 = OpLoad %int %value_2 %101 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %100 %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 %106 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %105 %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 %111 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %110 %32 %112 = OpIEqual %bool %111 %32 - %108 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %111 %112 + %108 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %111 %112 OpReturn OpFunctionEnd %main = OpFunction %void None %113 diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl index f1b92401fd..9e34babfbc 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -26,10 +26,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_1bd40a() { 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.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 x_19 = old_value_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 -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -77,10 +77,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_1bd40a() { 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.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 x_19 = old_value_1; x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_19, (x_19 == 1)); diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl index c4f2e90ddf..1a36daec3c 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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) { 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 x_19 = old_value_1; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_19, .exchanged=(x_19 == 1)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm index 6b85b45492..a873292e70 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm @@ -19,9 +19,9 @@ OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main "fragment_main" OpName %compute_main_1 "compute_main_1" @@ -33,8 +33,8 @@ OpDecorate %sb_rw Binding 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %int = OpTypeInt 32 1 %SB_RW_atomic = OpTypeStruct %int %sb_rw_block = OpTypeStruct %SB_RW_atomic @@ -46,7 +46,7 @@ %x__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %12 = OpConstantNull %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_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 @@ -59,7 +59,7 @@ %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 %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 %27 = OpIEqual %bool %26 %int_1 %28 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %26 %27 diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl index aa3ae45fa9..cb7126a472 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -26,10 +26,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_63d8e6() { 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.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 x_17 = old_value_1; 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 -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -77,10 +77,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_63d8e6() { 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.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 x_17 = old_value_1; x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u)); diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl index dae3d285ea..4bc8a6be3d 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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) { 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 x_17 = old_value_1; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_17, .exchanged=(x_17 == 1u)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm index b133b0adc1..319fbeff39 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm @@ -19,9 +19,9 @@ OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main "fragment_main" OpName %compute_main_1 "compute_main_1" @@ -33,8 +33,8 @@ OpDecorate %sb_rw Binding 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %SB_RW_atomic = OpTypeStruct %uint %sb_rw_block = OpTypeStruct %SB_RW_atomic @@ -46,7 +46,7 @@ %x__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %12 = OpConstantNull %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_0 = OpConstant %uint 0 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint @@ -57,7 +57,7 @@ %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 %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 %25 = OpIEqual %bool %24 %uint_1 %26 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %24 %25 diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl index 228a3a78c6..b86a5ff3c7 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -12,11 +12,11 @@ groupshared int arg_0; void atomicCompareExchangeWeak_e88938() { 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_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 x_18 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl index 228a3a78c6..b86a5ff3c7 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -12,11 +12,11 @@ groupshared int arg_0; void atomicCompareExchangeWeak_e88938() { 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_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 x_18 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_18, (x_18 == 1)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl index 11092c71de..e84462ab0f 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -15,10 +15,10 @@ uint local_invocation_index_1 = 0u; shared int arg_0; void atomicCompareExchangeWeak_e88938() { 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.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 x_18 = old_value_1; x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_18, (x_18 == 1)); diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl index 136bec8082..c9447ee7cd 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -23,7 +23,7 @@ struct x__atomic_compare_exchange_resulti32 { void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol_2) { 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 x_18 = old_value_1; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_18, .exchanged=(x_18 == 1)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm index e4be06c9df..e143e3cdcd 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm @@ -15,9 +15,9 @@ OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index_2 "local_invocation_index_2" OpName %compute_main_1 "compute_main_1" @@ -27,8 +27,8 @@ OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input @@ -44,7 +44,7 @@ %x__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool %16 = OpConstantNull %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_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 @@ -57,7 +57,7 @@ OpStore %res %16 %25 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %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 %28 = OpIEqual %bool %27 %int_1 %29 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %27 %28 diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl index 3ba0e3f1c3..b5a010d78c 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -12,11 +12,11 @@ groupshared uint arg_0; void atomicCompareExchangeWeak_83580d() { 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_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 x_17 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl index 3ba0e3f1c3..b5a010d78c 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -12,11 +12,11 @@ groupshared uint arg_0; void atomicCompareExchangeWeak_83580d() { 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_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 x_17 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_17, (x_17 == 1u)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl index 59a68d220d..b8db0ed5ed 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -15,10 +15,10 @@ uint local_invocation_index_1 = 0u; shared uint arg_0; void atomicCompareExchangeWeak_83580d() { 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.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 x_17 = old_value_1; x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_17, (x_17 == 1u)); diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl index a704ffaf30..d6e1a76dbb 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -23,7 +23,7 @@ struct x__atomic_compare_exchange_resultu32 { void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol_2) { 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 x_17 = old_value_1; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_17, .exchanged=(x_17 == 1u)}; diff --git a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm index 6200f8954b..44e61e0a0d 100644 --- a/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/literal/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm @@ -15,9 +15,9 @@ OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index_2 "local_invocation_index_2" OpName %compute_main_1 "compute_main_1" @@ -27,8 +27,8 @@ OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input @@ -43,7 +43,7 @@ %x__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool %15 = OpConstantNull %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_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 @@ -55,7 +55,7 @@ OpStore %res %15 %24 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %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 %27 = OpIEqual %bool %26 %uint_1 %28 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %26 %27 diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl index af721737eb..dffc4024d4 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -32,10 +32,10 @@ void atomicCompareExchangeWeak_1bd40a() { arg_2 = 1; int x_23 = arg_2; 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.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 x_25 = old_value_1; 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 -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -89,10 +89,10 @@ void atomicCompareExchangeWeak_1bd40a() { arg_2 = 1; int x_23 = arg_2; 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.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 x_25 = old_value_1; x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_25, (x_25 == x_23)); diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl index f629b29269..e281e03bbe 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -33,7 +33,7 @@ void atomicCompareExchangeWeak_1bd40a(device SB_RW_atomic* const tint_symbol_2) arg_2 = 1; int const x_23 = arg_2; 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 x_25 = old_value_1; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_25, .exchanged=(x_25 == x_23)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm index c90584f806..24b0e9b535 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_i32.spvasm.expected.spvasm @@ -21,9 +21,9 @@ OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main "fragment_main" OpName %compute_main_1 "compute_main_1" @@ -35,8 +35,8 @@ OpDecorate %sb_rw Binding 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %int = OpTypeInt 32 1 %SB_RW_atomic = OpTypeStruct %int %sb_rw_block = OpTypeStruct %SB_RW_atomic @@ -51,7 +51,7 @@ %16 = OpConstantNull %x__atomic_compare_exchange_resulti32 %_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32 %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_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 @@ -71,7 +71,7 @@ %29 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0 %30 = OpAtomicCompareExchange %int %29 %uint_1 %uint_0 %uint_0 %20 %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 %33 = OpIEqual %bool %32 %20 %34 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %32 %33 diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl index e72b539e0c..7c0b63fc94 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -32,10 +32,10 @@ void atomicCompareExchangeWeak_63d8e6() { arg_2 = 1u; uint x_21 = arg_2; 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.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 x_23 = old_value_1; 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 -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -89,10 +89,10 @@ void atomicCompareExchangeWeak_63d8e6() { arg_2 = 1u; uint x_21 = arg_2; 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.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 x_23 = old_value_1; x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21)); diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl index 7aca7ef3fb..34931277e6 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -33,7 +33,7 @@ void atomicCompareExchangeWeak_63d8e6(device SB_RW_atomic* const tint_symbol_2) arg_2 = 1u; uint const x_21 = arg_2; 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 x_23 = old_value_1; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_23, .exchanged=(x_23 == x_21)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm index ee2758aade..1321c54a85 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/storage_u32.spvasm.expected.spvasm @@ -21,9 +21,9 @@ OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %fragment_main_1 "fragment_main_1" OpName %fragment_main "fragment_main" OpName %compute_main_1 "compute_main_1" @@ -35,8 +35,8 @@ OpDecorate %sb_rw Binding 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %SB_RW_atomic = OpTypeStruct %uint %sb_rw_block = OpTypeStruct %SB_RW_atomic @@ -51,7 +51,7 @@ %16 = OpConstantNull %x__atomic_compare_exchange_resultu32 %_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32 %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 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint %atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6 @@ -69,7 +69,7 @@ %27 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0 %28 = OpAtomicCompareExchange %uint %27 %uint_1 %uint_0 %uint_0 %20 %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 %31 = OpIEqual %bool %30 %20 %32 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %30 %31 diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl index 0f21d8d213..3326ee4634 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.dxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_e88938() { arg_2 = 1; const int x_22 = arg_2; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_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 x_24 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl index 0f21d8d213..3326ee4634 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.fxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_e88938() { arg_2 = 1; const int x_22 = arg_2; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, x_22, atomic_result.old_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 x_24 = old_value_1; const x__atomic_compare_exchange_resulti32 tint_symbol_3 = {x_24, (x_24 == x_22)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl index bfa323328d..c93cea245d 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -21,10 +21,10 @@ void atomicCompareExchangeWeak_e88938() { arg_2 = 1; int x_22 = arg_2; 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.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 x_24 = old_value_1; x__atomic_compare_exchange_resulti32 tint_symbol_1 = x__atomic_compare_exchange_resulti32(x_24, (x_24 == x_22)); diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl index 412c26a525..42008ee8b3 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -29,7 +29,7 @@ void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol_ arg_2 = 1; int const x_22 = arg_2; 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 x_24 = old_value_1; x__atomic_compare_exchange_resulti32 const tint_symbol_1 = {.old_value=x_24, .exchanged=(x_24 == x_22)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm index ab701933da..08d81f19b2 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_i32.spvasm.expected.spvasm @@ -17,9 +17,9 @@ OpMemberName %x__atomic_compare_exchange_resulti32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resulti32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index_2 "local_invocation_index_2" OpName %compute_main_1 "compute_main_1" @@ -29,8 +29,8 @@ OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpMemberDecorate %x__atomic_compare_exchange_resulti32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resulti32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input @@ -49,7 +49,7 @@ %20 = OpConstantNull %x__atomic_compare_exchange_resulti32 %_ptr_Function_x__atomic_compare_exchange_resulti32 = OpTypePointer Function %x__atomic_compare_exchange_resulti32 %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_0 = OpConstant %uint 0 %36 = OpTypeFunction %void %uint @@ -68,7 +68,7 @@ %25 = OpLoad %int %arg_1 %31 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %24 %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 %34 = OpIEqual %bool %33 %24 %35 = OpCompositeConstruct %x__atomic_compare_exchange_resulti32 %33 %34 diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl index 46c7526d3c..40710ff7fd 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.dxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_83580d() { arg_2 = 1u; const uint x_21 = arg_2; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_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 x_23 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl index 46c7526d3c..40710ff7fd 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.fxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -18,11 +18,11 @@ void atomicCompareExchangeWeak_83580d() { arg_2 = 1u; const uint x_21 = arg_2; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, x_21, atomic_result.old_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 x_23 = old_value_1; const x__atomic_compare_exchange_resultu32 tint_symbol_3 = {x_23, (x_23 == x_21)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl index aeaaf5f701..920a16edf0 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -21,10 +21,10 @@ void atomicCompareExchangeWeak_83580d() { arg_2 = 1u; uint x_21 = arg_2; 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.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 x_23 = old_value_1; x__atomic_compare_exchange_resultu32 tint_symbol_1 = x__atomic_compare_exchange_resultu32(x_23, (x_23 == x_21)); diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl index 7984286995..82ddffc5eb 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -29,7 +29,7 @@ void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol arg_2 = 1u; uint const x_21 = arg_2; 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 x_23 = old_value_1; x__atomic_compare_exchange_resultu32 const tint_symbol_1 = {.old_value=x_23, .exchanged=(x_23 == x_21)}; diff --git a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm index aae3c6ca81..9091607e0a 100644 --- a/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm +++ b/test/tint/builtins/atomics/from_gen/var/atomicCompareExchangeWeak/workgroup_u32.spvasm.expected.spvasm @@ -17,9 +17,9 @@ OpMemberName %x__atomic_compare_exchange_resultu32 0 "old_value" OpMemberName %x__atomic_compare_exchange_resultu32 1 "exchanged" OpName %res "res" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index_2 "local_invocation_index_2" OpName %compute_main_1 "compute_main_1" @@ -29,8 +29,8 @@ OpDecorate %local_invocation_index_1_param_1 BuiltIn LocalInvocationIndex OpMemberDecorate %x__atomic_compare_exchange_resultu32 0 Offset 0 OpMemberDecorate %x__atomic_compare_exchange_resultu32 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1_param_1 = OpVariable %_ptr_Input_uint Input @@ -47,7 +47,7 @@ %18 = OpConstantNull %x__atomic_compare_exchange_resultu32 %_ptr_Function_x__atomic_compare_exchange_resultu32 = OpTypePointer Function %x__atomic_compare_exchange_resultu32 %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_0 = OpConstant %uint 0 %34 = OpTypeFunction %void %uint @@ -66,7 +66,7 @@ %23 = OpLoad %uint %arg_1 %29 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %22 %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 %32 = OpIEqual %bool %31 %22 %33 = OpCompositeConstruct %x__atomic_compare_exchange_resultu32 %31 %32 diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl index e3d50476a7..d7e3d34ccd 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -16,10 +16,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { } sb_rw; 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.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() { @@ -32,7 +32,7 @@ void main() { } #version 310 es -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -47,10 +47,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { } sb_rw; 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.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() { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl index 9f0ca2d9ad..5a81ab9fa8 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -17,7 +17,7 @@ struct SB_RW { }; 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)]]) { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm index 597b8974ea..08cb660f9b 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm @@ -15,9 +15,9 @@ OpMemberName %SB_RW 0 "arg_0" OpName %sb_rw "sb_rw" OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %res "res" OpName %fragment_main "fragment_main" OpName %compute_main "compute_main" @@ -26,8 +26,8 @@ OpMemberDecorate %SB_RW 0 Offset 0 OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw Binding 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %int = OpTypeInt 32 1 %SB_RW = OpTypeStruct %int %sb_rw_block = OpTypeStruct %SB_RW @@ -36,21 +36,21 @@ %void = OpTypeVoid %6 = OpTypeFunction %void %bool = OpTypeBool -%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool +%__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool %uint = OpTypeInt 32 0 %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int %int_1 = OpConstant %int 1 -%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 - %24 = OpConstantNull %__atomic_compare_exchange_resulti32 +%_ptr_Function___atomic_compare_exchange_result_i32 = OpTypePointer Function %__atomic_compare_exchange_result_i32 + %24 = OpConstantNull %__atomic_compare_exchange_result_i32 %atomicCompareExchangeWeak_1bd40a = OpFunction %void None %6 %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 %20 = OpAtomicCompareExchange %int %18 %uint_1 %uint_0 %uint_0 %int_1 %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 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl index 5498ecb02f..9a64da2844 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -16,10 +16,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { } sb_rw; 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.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() { @@ -32,7 +32,7 @@ void main() { } #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -47,10 +47,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { } sb_rw; 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.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() { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl index 9437613622..ef6ed6dad6 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -17,7 +17,7 @@ struct SB_RW { }; 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)]]) { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm index 6d95811d6c..3ed657e8ff 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm @@ -15,9 +15,9 @@ OpMemberName %SB_RW 0 "arg_0" OpName %sb_rw "sb_rw" OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %res "res" OpName %fragment_main "fragment_main" OpName %compute_main "compute_main" @@ -26,8 +26,8 @@ OpMemberDecorate %SB_RW 0 Offset 0 OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw Binding 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %SB_RW = OpTypeStruct %uint %sb_rw_block = OpTypeStruct %SB_RW @@ -36,19 +36,19 @@ %void = OpTypeVoid %6 = OpTypeFunction %void %bool = OpTypeBool -%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool +%__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint -%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 - %22 = OpConstantNull %__atomic_compare_exchange_resultu32 +%_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32 + %22 = OpConstantNull %__atomic_compare_exchange_result_u32 %atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6 %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 %18 = OpAtomicCompareExchange %uint %17 %uint_1 %uint_0 %uint_0 %uint_1 %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 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl index 3f46597640..4ba5820574 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl @@ -1,15 +1,15 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; groupshared uint arg_0; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value); atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; - atomic_compare_exchange_resultu32 res = atomic_result; + atomic_compare_exchange_result_u32 res = atomic_result; } struct tint_symbol_1 { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl index 3f46597640..4ba5820574 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl @@ -1,15 +1,15 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; groupshared uint arg_0; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value); atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; - atomic_compare_exchange_resultu32 res = atomic_result; + atomic_compare_exchange_result_u32 res = atomic_result; } struct tint_symbol_1 { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl index 589a98f416..27c55a218f 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -8,10 +8,10 @@ struct atomic_compare_exchange_resultu32 { shared uint arg_0; 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.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) { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl index effc1417eb..bd547cd67d 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl @@ -2,18 +2,18 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) { - atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u); + 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) { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm index 826498596a..98b61402b3 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm @@ -10,16 +10,16 @@ OpName %local_invocation_index_1 "local_invocation_index_1" OpName %arg_0 "arg_0" OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %res "res" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index "local_invocation_index" OpName %compute_main "compute_main" OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input @@ -28,21 +28,21 @@ %void = OpTypeVoid %6 = OpTypeFunction %void %bool = OpTypeBool -%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool +%__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 -%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 - %21 = OpConstantNull %__atomic_compare_exchange_resultu32 +%_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32 + %21 = OpConstantNull %__atomic_compare_exchange_result_u32 %22 = OpTypeFunction %void %uint %28 = OpConstantNull %uint %uint_264 = OpConstant %uint 264 %atomicCompareExchangeWeak_83580d = OpFunction %void None %6 %9 = OpLabel - %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %21 + %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 %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 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl index 4d201c49d5..26abb60402 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl @@ -1,15 +1,15 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; groupshared int arg_0; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value); atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; - atomic_compare_exchange_resulti32 res = atomic_result; + atomic_compare_exchange_result_i32 res = atomic_result; } struct tint_symbol_1 { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl index 4d201c49d5..26abb60402 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl @@ -1,15 +1,15 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; groupshared int arg_0; 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; InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value); atomic_result.exchanged = atomic_result.old_value == atomic_compare_value; - atomic_compare_exchange_resulti32 res = atomic_result; + atomic_compare_exchange_result_i32 res = atomic_result; } struct tint_symbol_1 { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl index ff5e7a1b5e..485e8b4a89 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -8,10 +8,10 @@ struct atomic_compare_exchange_resulti32 { shared int arg_0; 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.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) { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl index 2843d71cf2..97f5501096 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl @@ -2,18 +2,18 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; } void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol) { - atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1); + 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) { diff --git a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm index a0f338c664..e1a4f731ca 100644 --- a/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm @@ -10,16 +10,16 @@ OpName %local_invocation_index_1 "local_invocation_index_1" OpName %arg_0 "arg_0" OpName %atomicCompareExchangeWeak_e88938 "atomicCompareExchangeWeak_e88938" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %res "res" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index "local_invocation_index" OpName %compute_main "compute_main" OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input @@ -29,21 +29,21 @@ %void = OpTypeVoid %7 = OpTypeFunction %void %bool = OpTypeBool -%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool +%__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 %int_1 = OpConstant %int 1 -%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 - %22 = OpConstantNull %__atomic_compare_exchange_resulti32 +%_ptr_Function___atomic_compare_exchange_result_i32 = OpTypePointer Function %__atomic_compare_exchange_result_i32 + %22 = OpConstantNull %__atomic_compare_exchange_result_i32 %23 = OpTypeFunction %void %uint %29 = OpConstantNull %int %uint_264 = OpConstant %uint 264 %atomicCompareExchangeWeak_e88938 = OpFunction %void None %7 %10 = OpLabel - %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %22 + %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 %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 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl index 481d83a3ed..937662c85f 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -18,10 +18,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_1bd40a() { int arg_1 = 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.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() { @@ -34,7 +34,7 @@ void main() { } #version 310 es -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -51,10 +51,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_1bd40a() { int arg_1 = 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.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() { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl index d2b66917f6..dda297221d 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -19,7 +19,7 @@ struct SB_RW { void atomicCompareExchangeWeak_1bd40a(device SB_RW* const tint_symbol) { int arg_1 = 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)]]) { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm index b4bd3edf42..8bedd3ee89 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm @@ -17,9 +17,9 @@ OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a" OpName %arg_1 "arg_1" OpName %arg_2 "arg_2" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %res "res" OpName %fragment_main "fragment_main" OpName %compute_main "compute_main" @@ -28,8 +28,8 @@ OpMemberDecorate %SB_RW 0 Offset 0 OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw Binding 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %int = OpTypeInt 32 1 %SB_RW = OpTypeStruct %int %sb_rw_block = OpTypeStruct %SB_RW @@ -41,18 +41,18 @@ %_ptr_Function_int = OpTypePointer Function %int %13 = OpConstantNull %int %bool = OpTypeBool -%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool +%__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool %uint = OpTypeInt 32 0 %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int -%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 - %30 = OpConstantNull %__atomic_compare_exchange_resulti32 +%_ptr_Function___atomic_compare_exchange_result_i32 = OpTypePointer Function %__atomic_compare_exchange_result_i32 + %30 = OpConstantNull %__atomic_compare_exchange_result_i32 %atomicCompareExchangeWeak_1bd40a = OpFunction %void None %6 %9 = OpLabel %arg_1 = 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_2 %int_1 %23 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0 %uint_0 @@ -60,7 +60,7 @@ %25 = OpLoad %int %arg_1 %26 = OpAtomicCompareExchange %int %23 %uint_1 %uint_0 %uint_0 %24 %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 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl index 835d61c777..6edcf7ef97 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -18,10 +18,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_63d8e6() { uint arg_1 = 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.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() { @@ -34,7 +34,7 @@ void main() { } #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -51,10 +51,10 @@ layout(binding = 0, std430) buffer sb_rw_block_ssbo { void atomicCompareExchangeWeak_63d8e6() { uint arg_1 = 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.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() { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl index f4fda565e0..85149d6828 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -19,7 +19,7 @@ struct SB_RW { void atomicCompareExchangeWeak_63d8e6(device SB_RW* const tint_symbol) { uint arg_1 = 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)]]) { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm index 1bd3a82469..96676e462f 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm @@ -17,9 +17,9 @@ OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6" OpName %arg_1 "arg_1" OpName %arg_2 "arg_2" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %res "res" OpName %fragment_main "fragment_main" OpName %compute_main "compute_main" @@ -28,8 +28,8 @@ OpMemberDecorate %SB_RW 0 Offset 0 OpDecorate %sb_rw DescriptorSet 0 OpDecorate %sb_rw Binding 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %SB_RW = OpTypeStruct %uint %sb_rw_block = OpTypeStruct %SB_RW @@ -41,16 +41,16 @@ %_ptr_Function_uint = OpTypePointer Function %uint %13 = OpConstantNull %uint %bool = OpTypeBool -%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool +%__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool %uint_0 = OpConstant %uint 0 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint -%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 - %28 = OpConstantNull %__atomic_compare_exchange_resultu32 +%_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32 + %28 = OpConstantNull %__atomic_compare_exchange_result_u32 %atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %6 %9 = OpLabel %arg_1 = 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_2 %uint_1 %21 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0 %uint_0 @@ -58,7 +58,7 @@ %23 = OpLoad %uint %arg_1 %24 = OpAtomicCompareExchange %uint %21 %uint_1 %uint_0 %uint_0 %22 %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 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl index a445a1d0ab..e9a52b41d8 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.dxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -7,11 +7,11 @@ groupshared uint arg_0; void atomicCompareExchangeWeak_83580d() { uint arg_1 = 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; InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_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 { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl index a445a1d0ab..e9a52b41d8 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.fxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -7,11 +7,11 @@ groupshared uint arg_0; void atomicCompareExchangeWeak_83580d() { uint arg_1 = 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; InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_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 { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl index a9569da165..49b2a93963 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; bool exchanged; }; @@ -10,10 +10,10 @@ shared uint arg_0; void atomicCompareExchangeWeak_83580d() { uint arg_1 = 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.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) { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl index 811dd26244..750b1f979e 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resultu32 { +struct atomic_compare_exchange_result_u32 { uint old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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) { uint arg_1 = 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) { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm index 05cd15a8ea..521c9c4949 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm @@ -12,16 +12,16 @@ OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d" OpName %arg_1 "arg_1" OpName %arg_2 "arg_2" - OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32" - OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged" + OpName %__atomic_compare_exchange_result_u32 "__atomic_compare_exchange_result_u32" + OpMemberName %__atomic_compare_exchange_result_u32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_u32 1 "exchanged" OpName %res "res" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index "local_invocation_index" OpName %compute_main "compute_main" OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex - OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_u32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_u32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input @@ -33,25 +33,25 @@ %_ptr_Function_uint = OpTypePointer Function %uint %13 = OpConstantNull %uint %bool = OpTypeBool -%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool +%__atomic_compare_exchange_result_u32 = OpTypeStruct %uint %bool %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 -%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32 - %27 = OpConstantNull %__atomic_compare_exchange_resultu32 +%_ptr_Function___atomic_compare_exchange_result_u32 = OpTypePointer Function %__atomic_compare_exchange_result_u32 + %27 = OpConstantNull %__atomic_compare_exchange_result_u32 %28 = OpTypeFunction %void %uint %uint_264 = OpConstant %uint 264 %atomicCompareExchangeWeak_83580d = OpFunction %void None %6 %9 = OpLabel %arg_1 = 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_2 %uint_1 %21 = OpLoad %uint %arg_2 %22 = OpLoad %uint %arg_1 %23 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %21 %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 OpReturn OpFunctionEnd diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl index c14dade47d..d38a5208e5 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.dxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -7,11 +7,11 @@ groupshared int arg_0; void atomicCompareExchangeWeak_e88938() { int arg_1 = 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; InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_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 { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl index c14dade47d..d38a5208e5 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.fxc.hlsl @@ -1,4 +1,4 @@ -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -7,11 +7,11 @@ groupshared int arg_0; void atomicCompareExchangeWeak_e88938() { int arg_1 = 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; InterlockedCompareExchange(arg_0, atomic_compare_value, arg_2, atomic_result.old_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 { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl index 2c5d1a394a..281efee839 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl @@ -1,6 +1,6 @@ #version 310 es -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -10,10 +10,10 @@ shared int arg_0; void atomicCompareExchangeWeak_e88938() { int arg_1 = 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.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) { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl index de5d73f584..e23cec2c36 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); 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) { int arg_1 = 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) { diff --git a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm index 531d46877e..b0cb480037 100644 --- a/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm +++ b/test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm @@ -12,16 +12,16 @@ OpName %atomicCompareExchangeWeak_e88938 "atomicCompareExchangeWeak_e88938" OpName %arg_1 "arg_1" OpName %arg_2 "arg_2" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %res "res" OpName %compute_main_inner "compute_main_inner" OpName %local_invocation_index "local_invocation_index" OpName %compute_main "compute_main" OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %uint = OpTypeInt 32 0 %_ptr_Input_uint = OpTypePointer Input %uint %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input @@ -34,25 +34,25 @@ %_ptr_Function_int = OpTypePointer Function %int %14 = OpConstantNull %int %bool = OpTypeBool -%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool +%__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool %uint_2 = OpConstant %uint 2 %uint_0 = OpConstant %uint 0 -%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32 - %28 = OpConstantNull %__atomic_compare_exchange_resulti32 +%_ptr_Function___atomic_compare_exchange_result_i32 = OpTypePointer Function %__atomic_compare_exchange_result_i32 + %28 = OpConstantNull %__atomic_compare_exchange_result_i32 %29 = OpTypeFunction %void %uint %uint_264 = OpConstant %uint 264 %atomicCompareExchangeWeak_e88938 = OpFunction %void None %7 %10 = OpLabel %arg_1 = OpVariable %_ptr_Function_int Function %14 %arg_2 = OpVariable %_ptr_Function_int Function %14 - %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %28 + %res = OpVariable %_ptr_Function___atomic_compare_exchange_result_i32 Function %28 OpStore %arg_1 %int_1 OpStore %arg_2 %int_1 %22 = OpLoad %int %arg_2 %23 = OpLoad %int %arg_1 %24 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %22 %23 %25 = OpIEqual %bool %24 %23 - %16 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %24 %25 + %16 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %24 %25 OpStore %res %16 OpReturn OpFunctionEnd diff --git a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.glsl b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.glsl index 6a768d3cce..6e53b760ff 100644 --- a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.glsl +++ b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.glsl @@ -1,7 +1,7 @@ #version 310 es precision highp float; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; bool exchanged; }; @@ -23,10 +23,10 @@ int foo() { int x = 0; tint_symbol_1 tint_symbol = tint_symbol_1(0, false); if (!(tint_discarded)) { - atomic_compare_exchange_resulti32 atomic_compare_result; + atomic_compare_exchange_result_i32 atomic_compare_result; atomic_compare_result.old_value = atomicCompSwap(a.inner, 0, 1); atomic_compare_result.exchanged = atomic_compare_result.old_value == 0; - atomic_compare_exchange_resulti32 tint_symbol_2 = atomic_compare_result; + atomic_compare_exchange_result_i32 tint_symbol_2 = atomic_compare_result; tint_symbol.old_value = tint_symbol_2.old_value; tint_symbol.exchanged = tint_symbol_2.exchanged; } diff --git a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.msl b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.msl index db17a1dfc7..51643495fb 100644 --- a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.msl +++ b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.msl @@ -2,11 +2,11 @@ using namespace metal; -struct atomic_compare_exchange_resulti32 { +struct atomic_compare_exchange_result_i32 { int old_value; 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; bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed); return {old_value, exchanged}; @@ -30,7 +30,7 @@ int foo_inner(thread tint_private_vars_struct* const tint_private_vars, device a int x = 0; tint_symbol_2 tint_symbol_1 = {}; if (!((*(tint_private_vars)).tint_discarded)) { - atomic_compare_exchange_resulti32 const tint_symbol_3 = atomicCompareExchangeWeak_1(tint_symbol_4, 0, 1); + atomic_compare_exchange_result_i32 const tint_symbol_3 = atomicCompareExchangeWeak_1(tint_symbol_4, 0, 1); tint_symbol_1.old_value = tint_symbol_3.old_value; tint_symbol_1.exchanged = tint_symbol_3.exchanged; } diff --git a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.spvasm b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.spvasm index b210415256..fd658dcc16 100644 --- a/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.spvasm +++ b/test/tint/statements/discard/atomic_cmpxchg.wgsl.expected.spvasm @@ -18,9 +18,9 @@ OpMemberName %tint_symbol_1 0 "old_value" OpMemberName %tint_symbol_1 1 "exchanged" OpName %tint_symbol "tint_symbol" - OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32" - OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value" - OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged" + OpName %__atomic_compare_exchange_result_i32 "__atomic_compare_exchange_result_i32" + OpMemberName %__atomic_compare_exchange_result_i32 0 "old_value" + OpMemberName %__atomic_compare_exchange_result_i32 1 "exchanged" OpName %foo "foo" OpDecorate %value Location 0 OpDecorate %a_block Block @@ -29,8 +29,8 @@ OpDecorate %a Binding 0 OpMemberDecorate %tint_symbol_1 0 Offset 0 OpMemberDecorate %tint_symbol_1 1 Offset 4 - OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0 - OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4 + OpMemberDecorate %__atomic_compare_exchange_result_i32 0 Offset 0 + OpMemberDecorate %__atomic_compare_exchange_result_i32 1 Offset 4 %bool = OpTypeBool %2 = OpConstantNull %bool %_ptr_Private_bool = OpTypePointer Private %bool @@ -48,7 +48,7 @@ %tint_symbol_1 = OpTypeStruct %int %bool %_ptr_Function_tint_symbol_1 = OpTypePointer Function %tint_symbol_1 %21 = OpConstantNull %tint_symbol_1 -%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool +%__atomic_compare_exchange_result_i32 = OpTypeStruct %int %bool %uint = OpTypeInt 32 0 %uint_1 = OpConstant %uint 1 %uint_0 = OpConstant %uint 0 @@ -71,7 +71,7 @@ %33 = OpAccessChain %_ptr_StorageBuffer_int %a %uint_0 %35 = OpAtomicCompareExchange %int %33 %uint_1 %uint_0 %uint_0 %int_1 %8 %36 = OpIEqual %bool %35 %8 - %26 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %35 %36 + %26 = OpCompositeConstruct %__atomic_compare_exchange_result_i32 %35 %36 %37 = OpAccessChain %_ptr_Function_int %tint_symbol %uint_0 %38 = OpCompositeExtract %int %26 0 OpStore %37 %38