encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity
dawn-cmake/test/tint/var/uses/many_workgroup_vars.wgsl.ex...

418 lines
22 KiB
Plaintext
Raw Normal View History

msl: Use a struct for threadgroup memory arguments MSL has a limit on the number of threadgroup memory arguments, so use a struct to support an arbitrary number of workgroup variables. This commit introduces a `State` object to this transform, which is used to track which structs have been cloned eagerly, in order to avoid duplicating them. Bug: tint:938 Change-Id: Ia467db186e176a08f160455eab5fd3b3662f56b8 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65360 Auto-Submit: James Price <jrprice@google.com> Kokoro: James Price <jrprice@google.com> Commit-Queue: James Price <jrprice@google.com> Reviewed-by: Ben Clayton <bclayton@google.com>
2021-09-29 18:56:17 +00:00
#include <metal_stdlib>
using namespace metal;
struct tint_symbol_202 {
float2x2 m00;
float2x2 m01;
float2x2 m02;
float2x2 m03;
float2x2 m04;
float2x2 m05;
float2x2 m06;
float2x2 m07;
float2x2 m08;
float2x2 m09;
float2x2 m10;
float2x2 m11;
float2x2 m12;
float2x2 m13;
float2x2 m14;
float2x2 m15;
float2x2 m16;
float2x2 m17;
float2x2 m18;
float2x2 m19;
float2x2 m20;
float2x2 m21;
float2x2 m22;
float2x2 m23;
float2x2 m24;
float2x2 m25;
float2x2 m26;
float2x2 m27;
float2x2 m28;
float2x2 m29;
float2x2 m30;
float2x2 m31;
float2x2 m32;
float2x2 m33;
float2x2 m34;
float2x2 m35;
float2x2 m36;
float2x2 m37;
float2x2 m38;
float2x2 m39;
float2x2 m40;
float2x2 m41;
float2x2 m42;
float2x2 m43;
float2x2 m44;
float2x2 m45;
float2x2 m46;
float2x2 m47;
float2x2 m48;
float2x2 m49;
float2x2 m50;
float2x2 m51;
float2x2 m52;
float2x2 m53;
float2x2 m54;
float2x2 m55;
float2x2 m56;
float2x2 m57;
float2x2 m58;
float2x2 m59;
float2x2 m60;
float2x2 m61;
float2x2 m62;
float2x2 m63;
float2x2 m64;
float2x2 m65;
float2x2 m66;
float2x2 m67;
float2x2 m68;
float2x2 m69;
float2x2 m70;
float2x2 m71;
float2x2 m72;
float2x2 m73;
float2x2 m74;
float2x2 m75;
float2x2 m76;
float2x2 m77;
float2x2 m78;
float2x2 m79;
float2x2 m80;
float2x2 m81;
float2x2 m82;
float2x2 m83;
float2x2 m84;
float2x2 m85;
float2x2 m86;
float2x2 m87;
float2x2 m88;
float2x2 m89;
float2x2 m90;
float2x2 m91;
float2x2 m92;
float2x2 m93;
float2x2 m94;
float2x2 m95;
float2x2 m96;
float2x2 m97;
float2x2 m98;
float2x2 m99;
};
void tint_symbol_inner(uint idx, threadgroup float2x2* const tint_symbol_1, threadgroup float2x2* const tint_symbol_2, threadgroup float2x2* const tint_symbol_3, threadgroup float2x2* const tint_symbol_4, threadgroup float2x2* const tint_symbol_5, threadgroup float2x2* const tint_symbol_6, threadgroup float2x2* const tint_symbol_7, threadgroup float2x2* const tint_symbol_8, threadgroup float2x2* const tint_symbol_9, threadgroup float2x2* const tint_symbol_10, threadgroup float2x2* const tint_symbol_11, threadgroup float2x2* const tint_symbol_12, threadgroup float2x2* const tint_symbol_13, threadgroup float2x2* const tint_symbol_14, threadgroup float2x2* const tint_symbol_15, threadgroup float2x2* const tint_symbol_16, threadgroup float2x2* const tint_symbol_17, threadgroup float2x2* const tint_symbol_18, threadgroup float2x2* const tint_symbol_19, threadgroup float2x2* const tint_symbol_20, threadgroup float2x2* const tint_symbol_21, threadgroup float2x2* const tint_symbol_22, threadgroup float2x2* const tint_symbol_23, threadgroup float2x2* const tint_symbol_24, threadgroup float2x2* const tint_symbol_25, threadgroup float2x2* const tint_symbol_26, threadgroup float2x2* const tint_symbol_27, threadgroup float2x2* const tint_symbol_28, threadgroup float2x2* const tint_symbol_29, threadgroup float2x2* const tint_symbol_30, threadgroup float2x2* const tint_symbol_31, threadgroup float2x2* const tint_symbol_32, threadgroup float2x2* const tint_symbol_33, threadgroup float2x2* const tint_symbol_34, threadgroup float2x2* const tint_symbol_35, threadgroup float2x2* const tint_symbol_36, threadgroup float2x2* const tint_symbol_37, threadgroup float2x2* const tint_symbol_38, threadgroup float2x2* const tint_symbol_39, threadgroup float2x2* const tint_symbol_40, threadgroup float2x2* const tint_symbol_41, threadgroup float2x2* const tint_symbol_42, threadgroup float2x2* const tint_symbol_43, threadgroup float2x2* const tint_symbol_44, threadgroup float2x2* const tint_symbol_45, threadgroup float2x2* const tint_symbol_46, threadgroup float2x2* const tint_symbol_47, threadgroup float2x2* const tint_symbol_48, threadgroup float2x2* const tint_symbol_49, threadgroup float2x2* const tint_symbol_50, threadgroup float2x2* const tint_symbol_51, threadgroup float2x2* const tint_symbol_52, threadgroup float2x2* const tint_symbol_53, threadgroup float2x2* const tint_symbol_54, threadgroup float2x2* const tint_symbol_55, threadgroup float2x2* const tint_symbol_56, threadgroup float2x2* const tint_symbol_57, threadgroup float2x2* const tint_symbol_58, threadgroup float2x2* const tint_symbol_59, threadgroup float2x2* const tint_symbol_60, threadgroup float2x2* const tint_symbol_61, threadgroup float2x2* const tint_symbol_62, threadgroup float2x2* const tint_symbol_63, threadgroup float2x2* const tint_symbol_64, threadgroup float2x2* const tint_symbol_65, threadgroup float2x2* const tint_symbol_66, threadgroup float2x2* const tint_symbol_67, threadgroup float2x2* const tint_symbol_68, threadgroup float2x2* const tint_symbol_69, threadgroup float2x2* const tint_symbol_70, threadgroup float2x2* const tint_symbol_71, threadgroup float2x2* const tint_symbol_72, threadgroup float2x2* const tint_symbol_73, threadgroup float2x2* const tint_symbol_74, threadgroup float2x2* const tint_symbol_75, threadgroup float2x2* const tint_symbol_76, threadgroup float2x2* const tint_symbol_77, threadgroup float2x2* const tint_symbol_78, threadgroup float2x2* const tint_symbol_79, threadgroup float2x2* const tint_symbol_80, threadgroup float2x2* const tint_symbol_81, threadgroup float2x2* const tint_symbol_82, threadgroup float2x2* const tint_symbol_83, threadgroup float2x2* const tint_symbol_84, threadgroup float2x2* const tint_symbol_85, threadgroup float2x2* const tint_symbol_86, threadgroup float2x2* const tint_symbol_87, threadgroup float2x2* const tint_symbol_88, threadgroup float2x2* const tint_symbol_89, threadgroup float2x2* const tint_symbol_90, threadgroup float2x2* const tint_symbol_91, threadgroup float2x2* const tint_symbol_92, threadgroup float2x2* c
{
*(tint_symbol_1) = float2x2();
*(tint_symbol_2) = float2x2();
*(tint_symbol_3) = float2x2();
*(tint_symbol_4) = float2x2();
*(tint_symbol_5) = float2x2();
*(tint_symbol_6) = float2x2();
*(tint_symbol_7) = float2x2();
*(tint_symbol_8) = float2x2();
*(tint_symbol_9) = float2x2();
*(tint_symbol_10) = float2x2();
*(tint_symbol_11) = float2x2();
*(tint_symbol_12) = float2x2();
*(tint_symbol_13) = float2x2();
*(tint_symbol_14) = float2x2();
*(tint_symbol_15) = float2x2();
*(tint_symbol_16) = float2x2();
*(tint_symbol_17) = float2x2();
*(tint_symbol_18) = float2x2();
*(tint_symbol_19) = float2x2();
*(tint_symbol_20) = float2x2();
*(tint_symbol_21) = float2x2();
*(tint_symbol_22) = float2x2();
*(tint_symbol_23) = float2x2();
*(tint_symbol_24) = float2x2();
*(tint_symbol_25) = float2x2();
*(tint_symbol_26) = float2x2();
*(tint_symbol_27) = float2x2();
*(tint_symbol_28) = float2x2();
*(tint_symbol_29) = float2x2();
*(tint_symbol_30) = float2x2();
*(tint_symbol_31) = float2x2();
*(tint_symbol_32) = float2x2();
*(tint_symbol_33) = float2x2();
*(tint_symbol_34) = float2x2();
*(tint_symbol_35) = float2x2();
*(tint_symbol_36) = float2x2();
*(tint_symbol_37) = float2x2();
*(tint_symbol_38) = float2x2();
*(tint_symbol_39) = float2x2();
*(tint_symbol_40) = float2x2();
*(tint_symbol_41) = float2x2();
*(tint_symbol_42) = float2x2();
*(tint_symbol_43) = float2x2();
*(tint_symbol_44) = float2x2();
*(tint_symbol_45) = float2x2();
*(tint_symbol_46) = float2x2();
*(tint_symbol_47) = float2x2();
*(tint_symbol_48) = float2x2();
*(tint_symbol_49) = float2x2();
*(tint_symbol_50) = float2x2();
*(tint_symbol_51) = float2x2();
*(tint_symbol_52) = float2x2();
*(tint_symbol_53) = float2x2();
*(tint_symbol_54) = float2x2();
*(tint_symbol_55) = float2x2();
*(tint_symbol_56) = float2x2();
*(tint_symbol_57) = float2x2();
*(tint_symbol_58) = float2x2();
*(tint_symbol_59) = float2x2();
*(tint_symbol_60) = float2x2();
*(tint_symbol_61) = float2x2();
*(tint_symbol_62) = float2x2();
*(tint_symbol_63) = float2x2();
*(tint_symbol_64) = float2x2();
*(tint_symbol_65) = float2x2();
*(tint_symbol_66) = float2x2();
*(tint_symbol_67) = float2x2();
*(tint_symbol_68) = float2x2();
*(tint_symbol_69) = float2x2();
*(tint_symbol_70) = float2x2();
*(tint_symbol_71) = float2x2();
*(tint_symbol_72) = float2x2();
*(tint_symbol_73) = float2x2();
*(tint_symbol_74) = float2x2();
*(tint_symbol_75) = float2x2();
*(tint_symbol_76) = float2x2();
*(tint_symbol_77) = float2x2();
*(tint_symbol_78) = float2x2();
*(tint_symbol_79) = float2x2();
*(tint_symbol_80) = float2x2();
*(tint_symbol_81) = float2x2();
*(tint_symbol_82) = float2x2();
*(tint_symbol_83) = float2x2();
*(tint_symbol_84) = float2x2();
*(tint_symbol_85) = float2x2();
*(tint_symbol_86) = float2x2();
*(tint_symbol_87) = float2x2();
*(tint_symbol_88) = float2x2();
*(tint_symbol_89) = float2x2();
*(tint_symbol_90) = float2x2();
*(tint_symbol_91) = float2x2();
*(tint_symbol_92) = float2x2();
*(tint_symbol_93) = float2x2();
*(tint_symbol_94) = float2x2();
*(tint_symbol_95) = float2x2();
*(tint_symbol_96) = float2x2();
*(tint_symbol_97) = float2x2();
*(tint_symbol_98) = float2x2();
*(tint_symbol_99) = float2x2();
*(tint_symbol_100) = float2x2();
}
threadgroup_barrier(mem_flags::mem_threadgroup);
(*(tint_symbol_1))[0][0] = 1.0f;
(*(tint_symbol_2))[0][0] = 1.0f;
(*(tint_symbol_3))[0][0] = 1.0f;
(*(tint_symbol_4))[0][0] = 1.0f;
(*(tint_symbol_5))[0][0] = 1.0f;
(*(tint_symbol_6))[0][0] = 1.0f;
(*(tint_symbol_7))[0][0] = 1.0f;
(*(tint_symbol_8))[0][0] = 1.0f;
(*(tint_symbol_9))[0][0] = 1.0f;
(*(tint_symbol_10))[0][0] = 1.0f;
(*(tint_symbol_11))[0][0] = 1.0f;
(*(tint_symbol_12))[0][0] = 1.0f;
(*(tint_symbol_13))[0][0] = 1.0f;
(*(tint_symbol_14))[0][0] = 1.0f;
(*(tint_symbol_15))[0][0] = 1.0f;
(*(tint_symbol_16))[0][0] = 1.0f;
(*(tint_symbol_17))[0][0] = 1.0f;
(*(tint_symbol_18))[0][0] = 1.0f;
(*(tint_symbol_19))[0][0] = 1.0f;
(*(tint_symbol_20))[0][0] = 1.0f;
(*(tint_symbol_21))[0][0] = 1.0f;
(*(tint_symbol_22))[0][0] = 1.0f;
(*(tint_symbol_23))[0][0] = 1.0f;
(*(tint_symbol_24))[0][0] = 1.0f;
(*(tint_symbol_25))[0][0] = 1.0f;
(*(tint_symbol_26))[0][0] = 1.0f;
(*(tint_symbol_27))[0][0] = 1.0f;
(*(tint_symbol_28))[0][0] = 1.0f;
(*(tint_symbol_29))[0][0] = 1.0f;
(*(tint_symbol_30))[0][0] = 1.0f;
(*(tint_symbol_31))[0][0] = 1.0f;
(*(tint_symbol_32))[0][0] = 1.0f;
(*(tint_symbol_33))[0][0] = 1.0f;
(*(tint_symbol_34))[0][0] = 1.0f;
(*(tint_symbol_35))[0][0] = 1.0f;
(*(tint_symbol_36))[0][0] = 1.0f;
(*(tint_symbol_37))[0][0] = 1.0f;
(*(tint_symbol_38))[0][0] = 1.0f;
(*(tint_symbol_39))[0][0] = 1.0f;
(*(tint_symbol_40))[0][0] = 1.0f;
(*(tint_symbol_41))[0][0] = 1.0f;
(*(tint_symbol_42))[0][0] = 1.0f;
(*(tint_symbol_43))[0][0] = 1.0f;
(*(tint_symbol_44))[0][0] = 1.0f;
(*(tint_symbol_45))[0][0] = 1.0f;
(*(tint_symbol_46))[0][0] = 1.0f;
(*(tint_symbol_47))[0][0] = 1.0f;
(*(tint_symbol_48))[0][0] = 1.0f;
(*(tint_symbol_49))[0][0] = 1.0f;
(*(tint_symbol_50))[0][0] = 1.0f;
(*(tint_symbol_51))[0][0] = 1.0f;
(*(tint_symbol_52))[0][0] = 1.0f;
(*(tint_symbol_53))[0][0] = 1.0f;
(*(tint_symbol_54))[0][0] = 1.0f;
(*(tint_symbol_55))[0][0] = 1.0f;
(*(tint_symbol_56))[0][0] = 1.0f;
(*(tint_symbol_57))[0][0] = 1.0f;
(*(tint_symbol_58))[0][0] = 1.0f;
(*(tint_symbol_59))[0][0] = 1.0f;
(*(tint_symbol_60))[0][0] = 1.0f;
(*(tint_symbol_61))[0][0] = 1.0f;
(*(tint_symbol_62))[0][0] = 1.0f;
(*(tint_symbol_63))[0][0] = 1.0f;
(*(tint_symbol_64))[0][0] = 1.0f;
(*(tint_symbol_65))[0][0] = 1.0f;
(*(tint_symbol_66))[0][0] = 1.0f;
(*(tint_symbol_67))[0][0] = 1.0f;
(*(tint_symbol_68))[0][0] = 1.0f;
(*(tint_symbol_69))[0][0] = 1.0f;
(*(tint_symbol_70))[0][0] = 1.0f;
(*(tint_symbol_71))[0][0] = 1.0f;
(*(tint_symbol_72))[0][0] = 1.0f;
(*(tint_symbol_73))[0][0] = 1.0f;
(*(tint_symbol_74))[0][0] = 1.0f;
(*(tint_symbol_75))[0][0] = 1.0f;
(*(tint_symbol_76))[0][0] = 1.0f;
(*(tint_symbol_77))[0][0] = 1.0f;
(*(tint_symbol_78))[0][0] = 1.0f;
(*(tint_symbol_79))[0][0] = 1.0f;
(*(tint_symbol_80))[0][0] = 1.0f;
(*(tint_symbol_81))[0][0] = 1.0f;
(*(tint_symbol_82))[0][0] = 1.0f;
(*(tint_symbol_83))[0][0] = 1.0f;
(*(tint_symbol_84))[0][0] = 1.0f;
(*(tint_symbol_85))[0][0] = 1.0f;
(*(tint_symbol_86))[0][0] = 1.0f;
(*(tint_symbol_87))[0][0] = 1.0f;
(*(tint_symbol_88))[0][0] = 1.0f;
(*(tint_symbol_89))[0][0] = 1.0f;
(*(tint_symbol_90))[0][0] = 1.0f;
(*(tint_symbol_91))[0][0] = 1.0f;
(*(tint_symbol_92))[0][0] = 1.0f;
(*(tint_symbol_93))[0][0] = 1.0f;
(*(tint_symbol_94))[0][0] = 1.0f;
(*(tint_symbol_95))[0][0] = 1.0f;
(*(tint_symbol_96))[0][0] = 1.0f;
(*(tint_symbol_97))[0][0] = 1.0f;
(*(tint_symbol_98))[0][0] = 1.0f;
(*(tint_symbol_99))[0][0] = 1.0f;
(*(tint_symbol_100))[0][0] = 1.0f;
}
kernel void tint_symbol(threadgroup tint_symbol_202* tint_symbol_102 [[threadgroup(0)]], uint idx [[thread_index_in_threadgroup]]) {
msl: Handle buffer variables in transform This removes a lot of awkward logic from the MSL writer, and means that we now handle all module-scope variables with the same transform. Change-Id: I782e36a4b88dafbc3f8364f7caa7f95c6ae3f5f1 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/67643 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Ben Clayton <bclayton@google.com>
2021-10-28 15:00:39 +00:00
threadgroup float2x2* const tint_symbol_101 = &((*(tint_symbol_102)).m00);
threadgroup float2x2* const tint_symbol_103 = &((*(tint_symbol_102)).m01);
threadgroup float2x2* const tint_symbol_104 = &((*(tint_symbol_102)).m02);
threadgroup float2x2* const tint_symbol_105 = &((*(tint_symbol_102)).m03);
threadgroup float2x2* const tint_symbol_106 = &((*(tint_symbol_102)).m04);
threadgroup float2x2* const tint_symbol_107 = &((*(tint_symbol_102)).m05);
threadgroup float2x2* const tint_symbol_108 = &((*(tint_symbol_102)).m06);
threadgroup float2x2* const tint_symbol_109 = &((*(tint_symbol_102)).m07);
threadgroup float2x2* const tint_symbol_110 = &((*(tint_symbol_102)).m08);
threadgroup float2x2* const tint_symbol_111 = &((*(tint_symbol_102)).m09);
threadgroup float2x2* const tint_symbol_112 = &((*(tint_symbol_102)).m10);
threadgroup float2x2* const tint_symbol_113 = &((*(tint_symbol_102)).m11);
threadgroup float2x2* const tint_symbol_114 = &((*(tint_symbol_102)).m12);
threadgroup float2x2* const tint_symbol_115 = &((*(tint_symbol_102)).m13);
threadgroup float2x2* const tint_symbol_116 = &((*(tint_symbol_102)).m14);
threadgroup float2x2* const tint_symbol_117 = &((*(tint_symbol_102)).m15);
threadgroup float2x2* const tint_symbol_118 = &((*(tint_symbol_102)).m16);
threadgroup float2x2* const tint_symbol_119 = &((*(tint_symbol_102)).m17);
threadgroup float2x2* const tint_symbol_120 = &((*(tint_symbol_102)).m18);
threadgroup float2x2* const tint_symbol_121 = &((*(tint_symbol_102)).m19);
threadgroup float2x2* const tint_symbol_122 = &((*(tint_symbol_102)).m20);
threadgroup float2x2* const tint_symbol_123 = &((*(tint_symbol_102)).m21);
threadgroup float2x2* const tint_symbol_124 = &((*(tint_symbol_102)).m22);
threadgroup float2x2* const tint_symbol_125 = &((*(tint_symbol_102)).m23);
threadgroup float2x2* const tint_symbol_126 = &((*(tint_symbol_102)).m24);
threadgroup float2x2* const tint_symbol_127 = &((*(tint_symbol_102)).m25);
threadgroup float2x2* const tint_symbol_128 = &((*(tint_symbol_102)).m26);
threadgroup float2x2* const tint_symbol_129 = &((*(tint_symbol_102)).m27);
threadgroup float2x2* const tint_symbol_130 = &((*(tint_symbol_102)).m28);
threadgroup float2x2* const tint_symbol_131 = &((*(tint_symbol_102)).m29);
threadgroup float2x2* const tint_symbol_132 = &((*(tint_symbol_102)).m30);
threadgroup float2x2* const tint_symbol_133 = &((*(tint_symbol_102)).m31);
threadgroup float2x2* const tint_symbol_134 = &((*(tint_symbol_102)).m32);
threadgroup float2x2* const tint_symbol_135 = &((*(tint_symbol_102)).m33);
threadgroup float2x2* const tint_symbol_136 = &((*(tint_symbol_102)).m34);
threadgroup float2x2* const tint_symbol_137 = &((*(tint_symbol_102)).m35);
threadgroup float2x2* const tint_symbol_138 = &((*(tint_symbol_102)).m36);
threadgroup float2x2* const tint_symbol_139 = &((*(tint_symbol_102)).m37);
threadgroup float2x2* const tint_symbol_140 = &((*(tint_symbol_102)).m38);
threadgroup float2x2* const tint_symbol_141 = &((*(tint_symbol_102)).m39);
threadgroup float2x2* const tint_symbol_142 = &((*(tint_symbol_102)).m40);
threadgroup float2x2* const tint_symbol_143 = &((*(tint_symbol_102)).m41);
threadgroup float2x2* const tint_symbol_144 = &((*(tint_symbol_102)).m42);
threadgroup float2x2* const tint_symbol_145 = &((*(tint_symbol_102)).m43);
threadgroup float2x2* const tint_symbol_146 = &((*(tint_symbol_102)).m44);
threadgroup float2x2* const tint_symbol_147 = &((*(tint_symbol_102)).m45);
threadgroup float2x2* const tint_symbol_148 = &((*(tint_symbol_102)).m46);
threadgroup float2x2* const tint_symbol_149 = &((*(tint_symbol_102)).m47);
threadgroup float2x2* const tint_symbol_150 = &((*(tint_symbol_102)).m48);
threadgroup float2x2* const tint_symbol_151 = &((*(tint_symbol_102)).m49);
threadgroup float2x2* const tint_symbol_152 = &((*(tint_symbol_102)).m50);
threadgroup float2x2* const tint_symbol_153 = &((*(tint_symbol_102)).m51);
threadgroup float2x2* const tint_symbol_154 = &((*(tint_symbol_102)).m52);
threadgroup float2x2* const tint_symbol_155 = &((*(tint_symbol_102)).m53);
threadgroup float2x2* const tint_symbol_156 = &((*(tint_symbol_102)).m54);
threadgroup float2x2* const tint_symbol_157 = &((*(tint_symbol_102)).m55);
threadgroup float2x2* const tint_symbol_158 = &((*(tint_symbol_102)).m56);
threadgroup float2x2* const tint_symbol_159 = &((*(tint_symbol_102)).m57);
threadgroup float2x2* const tint_symbol_160 = &((*(tint_symbol_102)).m58);
threadgroup float2x2* const tint_symbol_161 = &((*(tint_symbol_102)).m59);
threadgroup float2x2* const tint_symbol_162 = &((*(tint_symbol_102)).m60);
threadgroup float2x2* const tint_symbol_163 = &((*(tint_symbol_102)).m61);
threadgroup float2x2* const tint_symbol_164 = &((*(tint_symbol_102)).m62);
threadgroup float2x2* const tint_symbol_165 = &((*(tint_symbol_102)).m63);
threadgroup float2x2* const tint_symbol_166 = &((*(tint_symbol_102)).m64);
threadgroup float2x2* const tint_symbol_167 = &((*(tint_symbol_102)).m65);
threadgroup float2x2* const tint_symbol_168 = &((*(tint_symbol_102)).m66);
threadgroup float2x2* const tint_symbol_169 = &((*(tint_symbol_102)).m67);
threadgroup float2x2* const tint_symbol_170 = &((*(tint_symbol_102)).m68);
threadgroup float2x2* const tint_symbol_171 = &((*(tint_symbol_102)).m69);
threadgroup float2x2* const tint_symbol_172 = &((*(tint_symbol_102)).m70);
threadgroup float2x2* const tint_symbol_173 = &((*(tint_symbol_102)).m71);
threadgroup float2x2* const tint_symbol_174 = &((*(tint_symbol_102)).m72);
threadgroup float2x2* const tint_symbol_175 = &((*(tint_symbol_102)).m73);
threadgroup float2x2* const tint_symbol_176 = &((*(tint_symbol_102)).m74);
threadgroup float2x2* const tint_symbol_177 = &((*(tint_symbol_102)).m75);
threadgroup float2x2* const tint_symbol_178 = &((*(tint_symbol_102)).m76);
threadgroup float2x2* const tint_symbol_179 = &((*(tint_symbol_102)).m77);
threadgroup float2x2* const tint_symbol_180 = &((*(tint_symbol_102)).m78);
threadgroup float2x2* const tint_symbol_181 = &((*(tint_symbol_102)).m79);
threadgroup float2x2* const tint_symbol_182 = &((*(tint_symbol_102)).m80);
threadgroup float2x2* const tint_symbol_183 = &((*(tint_symbol_102)).m81);
threadgroup float2x2* const tint_symbol_184 = &((*(tint_symbol_102)).m82);
threadgroup float2x2* const tint_symbol_185 = &((*(tint_symbol_102)).m83);
threadgroup float2x2* const tint_symbol_186 = &((*(tint_symbol_102)).m84);
threadgroup float2x2* const tint_symbol_187 = &((*(tint_symbol_102)).m85);
threadgroup float2x2* const tint_symbol_188 = &((*(tint_symbol_102)).m86);
threadgroup float2x2* const tint_symbol_189 = &((*(tint_symbol_102)).m87);
threadgroup float2x2* const tint_symbol_190 = &((*(tint_symbol_102)).m88);
threadgroup float2x2* const tint_symbol_191 = &((*(tint_symbol_102)).m89);
threadgroup float2x2* const tint_symbol_192 = &((*(tint_symbol_102)).m90);
threadgroup float2x2* const tint_symbol_193 = &((*(tint_symbol_102)).m91);
threadgroup float2x2* const tint_symbol_194 = &((*(tint_symbol_102)).m92);
threadgroup float2x2* const tint_symbol_195 = &((*(tint_symbol_102)).m93);
threadgroup float2x2* const tint_symbol_196 = &((*(tint_symbol_102)).m94);
threadgroup float2x2* const tint_symbol_197 = &((*(tint_symbol_102)).m95);
threadgroup float2x2* const tint_symbol_198 = &((*(tint_symbol_102)).m96);
threadgroup float2x2* const tint_symbol_199 = &((*(tint_symbol_102)).m97);
threadgroup float2x2* const tint_symbol_200 = &((*(tint_symbol_102)).m98);
threadgroup float2x2* const tint_symbol_201 = &((*(tint_symbol_102)).m99);
tint_symbol_inner(idx, tint_symbol_101, tint_symbol_103, tint_symbol_104, tint_symbol_105, tint_symbol_106, tint_symbol_107, tint_symbol_108, tint_symbol_109, tint_symbol_110, tint_symbol_111, tint_symbol_112, tint_symbol_113, tint_symbol_114, tint_symbol_115, tint_symbol_116, tint_symbol_117, tint_symbol_118, tint_symbol_119, tint_symbol_120, tint_symbol_121, tint_symbol_122, tint_symbol_123, tint_symbol_124, tint_symbol_125, tint_symbol_126, tint_symbol_127, tint_symbol_128, tint_symbol_129, tint_symbol_130, tint_symbol_131, tint_symbol_132, tint_symbol_133, tint_symbol_134, tint_symbol_135, tint_symbol_136, tint_symbol_137, tint_symbol_138, tint_symbol_139, tint_symbol_140, tint_symbol_141, tint_symbol_142, tint_symbol_143, tint_symbol_144, tint_symbol_145, tint_symbol_146, tint_symbol_147, tint_symbol_148, tint_symbol_149, tint_symbol_150, tint_symbol_151, tint_symbol_152, tint_symbol_153, tint_symbol_154, tint_symbol_155, tint_symbol_156, tint_symbol_157, tint_symbol_158, tint_symbol_159, tint_symbol_160, tint_symbol_161, tint_symbol_162, tint_symbol_163, tint_symbol_164, tint_symbol_165, tint_symbol_166, tint_symbol_167, tint_symbol_168, tint_symbol_169, tint_symbol_170, tint_symbol_171, tint_symbol_172, tint_symbol_173, tint_symbol_174, tint_symbol_175, tint_symbol_176, tint_symbol_177, tint_symbol_178, tint_symbol_179, tint_symbol_180, tint_symbol_181, tint_symbol_182, tint_symbol_183, tint_symbol_184, tint_symbol_185, tint_symbol_186, tint_symbol_187, tint_symbol_188, tint_symbol_189, tint_symbol_190, tint_symbol_191, tint_symbol_192, tint_symbol_193, tint_symbol_194, tint_symbol_195, tint_symbol_196, tint_symbol_197, tint_symbol_198, tint_symbol_199, tint_symbol_200, tint_symbol_201);
msl: Use a struct for threadgroup memory arguments MSL has a limit on the number of threadgroup memory arguments, so use a struct to support an arbitrary number of workgroup variables. This commit introduces a `State` object to this transform, which is used to track which structs have been cloned eagerly, in order to avoid duplicating them. Bug: tint:938 Change-Id: Ia467db186e176a08f160455eab5fd3b3662f56b8 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65360 Auto-Submit: James Price <jrprice@google.com> Kokoro: James Price <jrprice@google.com> Commit-Queue: James Price <jrprice@google.com> Reviewed-by: Ben Clayton <bclayton@google.com>
2021-09-29 18:56:17 +00:00
return;
}