spirv_common.hpp 51 KB

12345678910111213141516171819202122232425262728293031323334353637383940414243444546474849505152535455565758596061626364656667686970717273747576777879808182838485868788899091929394959697989910010110210310410510610710810911011111211311411511611711811912012112212312412512612712812913013113213313413513613713813914014114214314414514614714814915015115215315415515615715815916016116216316416516616716816917017117217317417517617717817918018118218318418518618718818919019119219319419519619719819920020120220320420520620720820921021121221321421521621721821922022122222322422522622722822923023123223323423523623723823924024124224324424524624724824925025125225325425525625725825926026126226326426526626726826927027127227327427527627727827928028128228328428528628728828929029129229329429529629729829930030130230330430530630730830931031131231331431531631731831932032132232332432532632732832933033133233333433533633733833934034134234334434534634734834935035135235335435535635735835936036136236336436536636736836937037137237337437537637737837938038138238338438538638738838939039139239339439539639739839940040140240340440540640740840941041141241341441541641741841942042142242342442542642742842943043143243343443543643743843944044144244344444544644744844945045145245345445545645745845946046146246346446546646746846947047147247347447547647747847948048148248348448548648748848949049149249349449549649749849950050150250350450550650750850951051151251351451551651751851952052152252352452552652752852953053153253353453553653753853954054154254354454554654754854955055155255355455555655755855956056156256356456556656756856957057157257357457557657757857958058158258358458558658758858959059159259359459559659759859960060160260360460560660760860961061161261361461561661761861962062162262362462562662762862963063163263363463563663763863964064164264364464564664764864965065165265365465565665765865966066166266366466566666766866967067167267367467567667767867968068168268368468568668768868969069169269369469569669769869970070170270370470570670770870971071171271371471571671771871972072172272372472572672772872973073173273373473573673773873974074174274374474574674774874975075175275375475575675775875976076176276376476576676776876977077177277377477577677777877978078178278378478578678778878979079179279379479579679779879980080180280380480580680780880981081181281381481581681781881982082182282382482582682782882983083183283383483583683783883984084184284384484584684784884985085185285385485585685785885986086186286386486586686786886987087187287387487587687787887988088188288388488588688788888989089189289389489589689789889990090190290390490590690790890991091191291391491591691791891992092192292392492592692792892993093193293393493593693793893994094194294394494594694794894995095195295395495595695795895996096196296396496596696796896997097197297397497597697797897998098198298398498598698798898999099199299399499599699799899910001001100210031004100510061007100810091010101110121013101410151016101710181019102010211022102310241025102610271028102910301031103210331034103510361037103810391040104110421043104410451046104710481049105010511052105310541055105610571058105910601061106210631064106510661067106810691070107110721073107410751076107710781079108010811082108310841085108610871088108910901091109210931094109510961097109810991100110111021103110411051106110711081109111011111112111311141115111611171118111911201121112211231124112511261127112811291130113111321133113411351136113711381139114011411142114311441145114611471148114911501151115211531154115511561157115811591160116111621163116411651166116711681169117011711172117311741175117611771178117911801181118211831184118511861187118811891190119111921193119411951196119711981199120012011202120312041205120612071208120912101211121212131214121512161217121812191220122112221223122412251226122712281229123012311232123312341235123612371238123912401241124212431244124512461247124812491250125112521253125412551256125712581259126012611262126312641265126612671268126912701271127212731274127512761277127812791280128112821283128412851286128712881289129012911292129312941295129612971298129913001301130213031304130513061307130813091310131113121313131413151316131713181319132013211322132313241325132613271328132913301331133213331334133513361337133813391340134113421343134413451346134713481349135013511352135313541355135613571358135913601361136213631364136513661367136813691370137113721373137413751376137713781379138013811382138313841385138613871388138913901391139213931394139513961397139813991400140114021403140414051406140714081409141014111412141314141415141614171418141914201421142214231424142514261427142814291430143114321433143414351436143714381439144014411442144314441445144614471448144914501451145214531454145514561457145814591460146114621463146414651466146714681469147014711472147314741475147614771478147914801481148214831484148514861487148814891490149114921493149414951496149714981499150015011502150315041505150615071508150915101511151215131514151515161517151815191520152115221523152415251526152715281529153015311532153315341535153615371538153915401541154215431544154515461547154815491550155115521553155415551556155715581559156015611562156315641565156615671568156915701571157215731574157515761577157815791580158115821583158415851586158715881589159015911592159315941595159615971598159916001601160216031604160516061607160816091610161116121613161416151616161716181619162016211622162316241625162616271628162916301631163216331634163516361637163816391640164116421643164416451646164716481649165016511652165316541655165616571658165916601661166216631664166516661667166816691670167116721673167416751676167716781679168016811682168316841685168616871688168916901691169216931694169516961697169816991700170117021703170417051706170717081709171017111712171317141715171617171718171917201721172217231724172517261727172817291730173117321733173417351736173717381739174017411742174317441745174617471748174917501751175217531754175517561757175817591760176117621763176417651766176717681769177017711772177317741775177617771778177917801781178217831784178517861787178817891790179117921793179417951796179717981799180018011802180318041805180618071808180918101811181218131814181518161817181818191820182118221823182418251826182718281829183018311832183318341835183618371838183918401841184218431844184518461847184818491850185118521853185418551856185718581859186018611862186318641865186618671868186918701871187218731874187518761877187818791880188118821883188418851886188718881889189018911892189318941895189618971898189919001901190219031904190519061907190819091910191119121913191419151916191719181919192019211922192319241925192619271928192919301931193219331934193519361937193819391940194119421943194419451946194719481949195019511952195319541955195619571958195919601961196219631964196519661967196819691970197119721973197419751976197719781979198019811982198319841985198619871988198919901991199219931994199519961997199819992000200120022003200420052006200720082009201020112012201320142015201620172018201920202021202220232024202520262027202820292030203120322033203420352036203720382039204020412042204320442045204620472048204920502051205220532054205520562057205820592060206120622063206420652066206720682069207020712072
  1. /*
  2. * Copyright 2015-2021 Arm Limited
  3. * SPDX-License-Identifier: Apache-2.0 OR MIT
  4. *
  5. * Licensed under the Apache License, Version 2.0 (the "License");
  6. * you may not use this file except in compliance with the License.
  7. * You may obtain a copy of the License at
  8. *
  9. * http://www.apache.org/licenses/LICENSE-2.0
  10. *
  11. * Unless required by applicable law or agreed to in writing, software
  12. * distributed under the License is distributed on an "AS IS" BASIS,
  13. * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  14. * See the License for the specific language governing permissions and
  15. * limitations under the License.
  16. */
  17. /*
  18. * At your option, you may choose to accept this material under either:
  19. * 1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or
  20. * 2. The MIT License, found at <http://opensource.org/licenses/MIT>.
  21. */
  22. #ifndef SPIRV_CROSS_COMMON_HPP
  23. #define SPIRV_CROSS_COMMON_HPP
  24. #ifndef SPV_ENABLE_UTILITY_CODE
  25. #define SPV_ENABLE_UTILITY_CODE
  26. #endif
  27. // Pragmatic hack to avoid symbol conflicts when including both hpp11 and hpp headers in same translation unit.
  28. // This is an unfortunate SPIRV-Headers issue that we cannot easily deal with ourselves.
  29. #ifdef SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
  30. #define spv SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
  31. #define SPIRV_CROSS_SPV_HEADER_NAMESPACE SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
  32. #else
  33. #define SPIRV_CROSS_SPV_HEADER_NAMESPACE spv
  34. #endif
  35. #include "spirv.hpp"
  36. #include "spirv_cross_containers.hpp"
  37. #include "spirv_cross_error_handling.hpp"
  38. #include <functional>
  39. // A bit crude, but allows projects which embed SPIRV-Cross statically to
  40. // effectively hide all the symbols from other projects.
  41. // There is a case where we have:
  42. // - Project A links against SPIRV-Cross statically.
  43. // - Project A links against Project B statically.
  44. // - Project B links against SPIRV-Cross statically (might be a different version).
  45. // This leads to a conflict with extremely bizarre results.
  46. // By overriding the namespace in one of the project builds, we can work around this.
  47. // If SPIRV-Cross is embedded in dynamic libraries,
  48. // prefer using -fvisibility=hidden on GCC/Clang instead.
  49. #ifdef SPIRV_CROSS_NAMESPACE_OVERRIDE
  50. #define SPIRV_CROSS_NAMESPACE SPIRV_CROSS_NAMESPACE_OVERRIDE
  51. #else
  52. #define SPIRV_CROSS_NAMESPACE spirv_cross
  53. #endif
  54. namespace SPIRV_CROSS_NAMESPACE
  55. {
  56. namespace inner
  57. {
  58. template <typename T>
  59. void join_helper(StringStream<> &stream, T &&t)
  60. {
  61. stream << std::forward<T>(t);
  62. }
  63. template <typename T, typename... Ts>
  64. void join_helper(StringStream<> &stream, T &&t, Ts &&... ts)
  65. {
  66. stream << std::forward<T>(t);
  67. join_helper(stream, std::forward<Ts>(ts)...);
  68. }
  69. } // namespace inner
  70. class Bitset
  71. {
  72. public:
  73. Bitset() = default;
  74. explicit inline Bitset(uint64_t lower_)
  75. : lower(lower_)
  76. {
  77. }
  78. inline bool get(uint32_t bit) const
  79. {
  80. if (bit < 64)
  81. return (lower & (1ull << bit)) != 0;
  82. else
  83. return higher.count(bit) != 0;
  84. }
  85. inline void set(uint32_t bit)
  86. {
  87. if (bit < 64)
  88. lower |= 1ull << bit;
  89. else
  90. higher.insert(bit);
  91. }
  92. inline void clear(uint32_t bit)
  93. {
  94. if (bit < 64)
  95. lower &= ~(1ull << bit);
  96. else
  97. higher.erase(bit);
  98. }
  99. inline uint64_t get_lower() const
  100. {
  101. return lower;
  102. }
  103. inline void reset()
  104. {
  105. lower = 0;
  106. higher.clear();
  107. }
  108. inline void merge_and(const Bitset &other)
  109. {
  110. lower &= other.lower;
  111. std::unordered_set<uint32_t> tmp_set;
  112. for (auto &v : higher)
  113. if (other.higher.count(v) != 0)
  114. tmp_set.insert(v);
  115. higher = std::move(tmp_set);
  116. }
  117. inline void merge_or(const Bitset &other)
  118. {
  119. lower |= other.lower;
  120. for (auto &v : other.higher)
  121. higher.insert(v);
  122. }
  123. inline bool operator==(const Bitset &other) const
  124. {
  125. if (lower != other.lower)
  126. return false;
  127. if (higher.size() != other.higher.size())
  128. return false;
  129. for (auto &v : higher)
  130. if (other.higher.count(v) == 0)
  131. return false;
  132. return true;
  133. }
  134. inline bool operator!=(const Bitset &other) const
  135. {
  136. return !(*this == other);
  137. }
  138. template <typename Op>
  139. void for_each_bit(const Op &op) const
  140. {
  141. // TODO: Add ctz-based iteration.
  142. for (uint32_t i = 0; i < 64; i++)
  143. {
  144. if (lower & (1ull << i))
  145. op(i);
  146. }
  147. if (higher.empty())
  148. return;
  149. // Need to enforce an order here for reproducible results,
  150. // but hitting this path should happen extremely rarely, so having this slow path is fine.
  151. SmallVector<uint32_t> bits;
  152. bits.reserve(higher.size());
  153. for (auto &v : higher)
  154. bits.push_back(v);
  155. std::sort(std::begin(bits), std::end(bits));
  156. for (auto &v : bits)
  157. op(v);
  158. }
  159. inline bool empty() const
  160. {
  161. return lower == 0 && higher.empty();
  162. }
  163. private:
  164. // The most common bits to set are all lower than 64,
  165. // so optimize for this case. Bits spilling outside 64 go into a slower data structure.
  166. // In almost all cases, higher data structure will not be used.
  167. uint64_t lower = 0;
  168. std::unordered_set<uint32_t> higher;
  169. };
  170. // Helper template to avoid lots of nasty string temporary munging.
  171. template <typename... Ts>
  172. std::string join(Ts &&... ts)
  173. {
  174. StringStream<> stream;
  175. inner::join_helper(stream, std::forward<Ts>(ts)...);
  176. return stream.str();
  177. }
  178. inline std::string merge(const SmallVector<std::string> &list, const char *between = ", ")
  179. {
  180. StringStream<> stream;
  181. for (auto &elem : list)
  182. {
  183. stream << elem;
  184. if (&elem != &list.back())
  185. stream << between;
  186. }
  187. return stream.str();
  188. }
  189. // Make sure we don't accidentally call this with float or doubles with SFINAE.
  190. // Have to use the radix-aware overload.
  191. template <typename T, typename std::enable_if<!std::is_floating_point<T>::value, int>::type = 0>
  192. inline std::string convert_to_string(const T &t)
  193. {
  194. return std::to_string(t);
  195. }
  196. static inline std::string convert_to_string(int32_t value)
  197. {
  198. // INT_MIN is ... special on some backends. If we use a decimal literal, and negate it, we
  199. // could accidentally promote the literal to long first, then negate.
  200. // To workaround it, emit int(0x80000000) instead.
  201. if (value == (std::numeric_limits<int32_t>::min)())
  202. return "int(0x80000000)";
  203. else
  204. return std::to_string(value);
  205. }
  206. static inline std::string convert_to_string(int64_t value, const std::string &int64_type, bool long_long_literal_suffix)
  207. {
  208. // INT64_MIN is ... special on some backends.
  209. // If we use a decimal literal, and negate it, we might overflow the representable numbers.
  210. // To workaround it, emit int(0x80000000) instead.
  211. if (value == (std::numeric_limits<int64_t>::min)())
  212. return join(int64_type, "(0x8000000000000000u", (long_long_literal_suffix ? "ll" : "l"), ")");
  213. else
  214. return std::to_string(value) + (long_long_literal_suffix ? "ll" : "l");
  215. }
  216. // Allow implementations to set a convenient standard precision
  217. #ifndef SPIRV_CROSS_FLT_FMT
  218. #define SPIRV_CROSS_FLT_FMT "%.32g"
  219. #endif
  220. // Disable sprintf and strcat warnings.
  221. // We cannot rely on snprintf and family existing because, ..., MSVC.
  222. #if defined(__clang__) || defined(__GNUC__)
  223. #pragma GCC diagnostic push
  224. #pragma GCC diagnostic ignored "-Wdeprecated-declarations"
  225. #elif defined(_MSC_VER)
  226. #pragma warning(push)
  227. #pragma warning(disable : 4996)
  228. #endif
  229. static inline void fixup_radix_point(char *str, char radix_point)
  230. {
  231. // Setting locales is a very risky business in multi-threaded program,
  232. // so just fixup locales instead. We only need to care about the radix point.
  233. if (radix_point != '.')
  234. {
  235. while (*str != '\0')
  236. {
  237. if (*str == radix_point)
  238. *str = '.';
  239. str++;
  240. }
  241. }
  242. }
  243. inline std::string convert_to_string(float t, char locale_radix_point)
  244. {
  245. // std::to_string for floating point values is broken.
  246. // Fallback to something more sane.
  247. char buf[64];
  248. sprintf(buf, SPIRV_CROSS_FLT_FMT, t);
  249. fixup_radix_point(buf, locale_radix_point);
  250. // Ensure that the literal is float.
  251. if (!strchr(buf, '.') && !strchr(buf, 'e'))
  252. strcat(buf, ".0");
  253. return buf;
  254. }
  255. inline std::string convert_to_string(double t, char locale_radix_point)
  256. {
  257. // std::to_string for floating point values is broken.
  258. // Fallback to something more sane.
  259. char buf[64];
  260. sprintf(buf, SPIRV_CROSS_FLT_FMT, t);
  261. fixup_radix_point(buf, locale_radix_point);
  262. // Ensure that the literal is float.
  263. if (!strchr(buf, '.') && !strchr(buf, 'e'))
  264. strcat(buf, ".0");
  265. return buf;
  266. }
  267. #if defined(__clang__) || defined(__GNUC__)
  268. #pragma GCC diagnostic pop
  269. #elif defined(_MSC_VER)
  270. #pragma warning(pop)
  271. #endif
  272. class FloatFormatter
  273. {
  274. public:
  275. virtual ~FloatFormatter() = default;
  276. virtual std::string format_float(float value) = 0;
  277. virtual std::string format_double(double value) = 0;
  278. };
  279. template <typename T>
  280. struct ValueSaver
  281. {
  282. explicit ValueSaver(T &current_)
  283. : current(current_)
  284. , saved(current_)
  285. {
  286. }
  287. void release()
  288. {
  289. current = saved;
  290. }
  291. ~ValueSaver()
  292. {
  293. release();
  294. }
  295. T &current;
  296. T saved;
  297. };
  298. struct Instruction
  299. {
  300. uint16_t op = 0;
  301. uint16_t count = 0;
  302. // If offset is 0 (not a valid offset into the instruction stream),
  303. // we have an instruction stream which is embedded in the object.
  304. uint32_t offset = 0;
  305. uint32_t length = 0;
  306. inline bool is_embedded() const
  307. {
  308. return offset == 0;
  309. }
  310. };
  311. struct EmbeddedInstruction : Instruction
  312. {
  313. SmallVector<uint32_t> ops;
  314. };
  315. enum Types
  316. {
  317. TypeNone,
  318. TypeType,
  319. TypeVariable,
  320. TypeConstant,
  321. TypeFunction,
  322. TypeFunctionPrototype,
  323. TypeBlock,
  324. TypeExtension,
  325. TypeExpression,
  326. TypeConstantOp,
  327. TypeCombinedImageSampler,
  328. TypeAccessChain,
  329. TypeUndef,
  330. TypeString,
  331. TypeDebugLocalVariable,
  332. TypeCount
  333. };
  334. template <Types type>
  335. class TypedID;
  336. template <>
  337. class TypedID<TypeNone>
  338. {
  339. public:
  340. TypedID() = default;
  341. TypedID(uint32_t id_)
  342. : id(id_)
  343. {
  344. }
  345. template <Types U>
  346. TypedID(const TypedID<U> &other)
  347. {
  348. *this = other;
  349. }
  350. template <Types U>
  351. TypedID &operator=(const TypedID<U> &other)
  352. {
  353. id = uint32_t(other);
  354. return *this;
  355. }
  356. // Implicit conversion to u32 is desired here.
  357. // As long as we block implicit conversion between TypedID<A> and TypedID<B> we're good.
  358. operator uint32_t() const
  359. {
  360. return id;
  361. }
  362. template <Types U>
  363. operator TypedID<U>() const
  364. {
  365. return TypedID<U>(*this);
  366. }
  367. private:
  368. uint32_t id = 0;
  369. };
  370. template <Types type>
  371. class TypedID
  372. {
  373. public:
  374. TypedID() = default;
  375. TypedID(uint32_t id_)
  376. : id(id_)
  377. {
  378. }
  379. explicit TypedID(const TypedID<TypeNone> &other)
  380. : id(uint32_t(other))
  381. {
  382. }
  383. operator uint32_t() const
  384. {
  385. return id;
  386. }
  387. private:
  388. uint32_t id = 0;
  389. };
  390. using VariableID = TypedID<TypeVariable>;
  391. using TypeID = TypedID<TypeType>;
  392. using ConstantID = TypedID<TypeConstant>;
  393. using FunctionID = TypedID<TypeFunction>;
  394. using BlockID = TypedID<TypeBlock>;
  395. using ID = TypedID<TypeNone>;
  396. // Helper for Variant interface.
  397. struct IVariant
  398. {
  399. virtual ~IVariant() = default;
  400. virtual IVariant *clone(ObjectPoolBase *pool) = 0;
  401. ID self = 0;
  402. protected:
  403. IVariant() = default;
  404. IVariant(const IVariant&) = default;
  405. IVariant &operator=(const IVariant&) = default;
  406. };
  407. #define SPIRV_CROSS_DECLARE_CLONE(T) \
  408. IVariant *clone(ObjectPoolBase *pool) override \
  409. { \
  410. return static_cast<ObjectPool<T> *>(pool)->allocate(*this); \
  411. }
  412. struct SPIRUndef : IVariant
  413. {
  414. enum
  415. {
  416. type = TypeUndef
  417. };
  418. explicit SPIRUndef(TypeID basetype_)
  419. : basetype(basetype_)
  420. {
  421. }
  422. TypeID basetype;
  423. SPIRV_CROSS_DECLARE_CLONE(SPIRUndef)
  424. };
  425. struct SPIRString : IVariant
  426. {
  427. enum
  428. {
  429. type = TypeString
  430. };
  431. explicit SPIRString(std::string str_)
  432. : str(std::move(str_))
  433. {
  434. }
  435. std::string str;
  436. SPIRV_CROSS_DECLARE_CLONE(SPIRString)
  437. };
  438. struct SPIRDebugLocalVariable : IVariant
  439. {
  440. enum
  441. {
  442. type = TypeDebugLocalVariable
  443. };
  444. uint32_t name_id;
  445. SPIRV_CROSS_DECLARE_CLONE(SPIRDebugLocalVariable)
  446. };
  447. // This type is only used by backends which need to access the combined image and sampler IDs separately after
  448. // the OpSampledImage opcode.
  449. struct SPIRCombinedImageSampler : IVariant
  450. {
  451. enum
  452. {
  453. type = TypeCombinedImageSampler
  454. };
  455. SPIRCombinedImageSampler(TypeID type_, VariableID image_, VariableID sampler_)
  456. : combined_type(type_)
  457. , image(image_)
  458. , sampler(sampler_)
  459. {
  460. }
  461. TypeID combined_type;
  462. VariableID image;
  463. VariableID sampler;
  464. SPIRV_CROSS_DECLARE_CLONE(SPIRCombinedImageSampler)
  465. };
  466. struct SPIRConstantOp : IVariant
  467. {
  468. enum
  469. {
  470. type = TypeConstantOp
  471. };
  472. SPIRConstantOp(TypeID result_type, spv::Op op, const uint32_t *args, uint32_t length)
  473. : opcode(op)
  474. , basetype(result_type)
  475. {
  476. arguments.reserve(length);
  477. for (uint32_t i = 0; i < length; i++)
  478. arguments.push_back(args[i]);
  479. }
  480. spv::Op opcode;
  481. SmallVector<uint32_t> arguments;
  482. TypeID basetype;
  483. SPIRV_CROSS_DECLARE_CLONE(SPIRConstantOp)
  484. };
  485. struct SPIRType : IVariant
  486. {
  487. enum
  488. {
  489. type = TypeType
  490. };
  491. spv::Op op = spv::Op::OpNop;
  492. explicit SPIRType(spv::Op op_) : op(op_) {}
  493. enum BaseType
  494. {
  495. Unknown,
  496. Void,
  497. Boolean,
  498. SByte,
  499. UByte,
  500. Short,
  501. UShort,
  502. Int,
  503. UInt,
  504. Int64,
  505. UInt64,
  506. AtomicCounter,
  507. Half,
  508. Float,
  509. Double,
  510. Struct,
  511. Image,
  512. SampledImage,
  513. Sampler,
  514. AccelerationStructure,
  515. RayQuery,
  516. CoopVecNV,
  517. // Keep internal types at the end.
  518. ControlPointArray,
  519. Interpolant,
  520. Char,
  521. // MSL specific type, that is used by 'object'(analog of 'task' from glsl) shader.
  522. MeshGridProperties,
  523. BFloat16,
  524. FloatE4M3,
  525. FloatE5M2,
  526. Tensor
  527. };
  528. // Scalar/vector/matrix support.
  529. BaseType basetype = Unknown;
  530. uint32_t width = 0;
  531. uint32_t vecsize = 1;
  532. uint32_t columns = 1;
  533. // Arrays, support array of arrays by having a vector of array sizes.
  534. SmallVector<uint32_t> array;
  535. // Array elements can be either specialization constants or specialization ops.
  536. // This array determines how to interpret the array size.
  537. // If an element is true, the element is a literal,
  538. // otherwise, it's an expression, which must be resolved on demand.
  539. // The actual size is not really known until runtime.
  540. SmallVector<bool> array_size_literal;
  541. // Pointers
  542. // Keep track of how many pointer layers we have.
  543. uint32_t pointer_depth = 0;
  544. bool pointer = false;
  545. bool forward_pointer = false;
  546. union
  547. {
  548. struct
  549. {
  550. uint32_t use_id;
  551. uint32_t rows_id;
  552. uint32_t columns_id;
  553. uint32_t scope_id;
  554. } cooperative;
  555. struct
  556. {
  557. uint32_t component_type_id;
  558. uint32_t component_count_id;
  559. } coopVecNV;
  560. struct
  561. {
  562. uint32_t type;
  563. uint32_t rank;
  564. uint32_t shape;
  565. } tensor;
  566. } ext;
  567. spv::StorageClass storage = spv::StorageClassGeneric;
  568. SmallVector<TypeID> member_types;
  569. // If member order has been rewritten to handle certain scenarios with Offset,
  570. // allow codegen to rewrite the index.
  571. SmallVector<uint32_t> member_type_index_redirection;
  572. struct ImageType
  573. {
  574. TypeID type;
  575. spv::Dim dim;
  576. bool depth;
  577. bool arrayed;
  578. bool ms;
  579. uint32_t sampled;
  580. spv::ImageFormat format;
  581. spv::AccessQualifier access;
  582. } image = {};
  583. // Structs can be declared multiple times if they are used as part of interface blocks.
  584. // We want to detect this so that we only emit the struct definition once.
  585. // Since we cannot rely on OpName to be equal, we need to figure out aliases.
  586. TypeID type_alias = 0;
  587. // Denotes the type which this type is based on.
  588. // Allows the backend to traverse how a complex type is built up during access chains.
  589. TypeID parent_type = 0;
  590. // Used in backends to avoid emitting members with conflicting names.
  591. std::unordered_set<std::string> member_name_cache;
  592. SPIRV_CROSS_DECLARE_CLONE(SPIRType)
  593. };
  594. struct SPIRExtension : IVariant
  595. {
  596. enum
  597. {
  598. type = TypeExtension
  599. };
  600. enum Extension
  601. {
  602. Unsupported,
  603. GLSL,
  604. SPV_debug_info,
  605. SPV_AMD_shader_ballot,
  606. SPV_AMD_shader_explicit_vertex_parameter,
  607. SPV_AMD_shader_trinary_minmax,
  608. SPV_AMD_gcn_shader,
  609. NonSemanticDebugPrintf,
  610. NonSemanticShaderDebugInfo,
  611. NonSemanticGeneric
  612. };
  613. enum ShaderDebugInfoOps
  614. {
  615. DebugLine = 103,
  616. DebugSource = 35
  617. };
  618. explicit SPIRExtension(Extension ext_)
  619. : ext(ext_)
  620. {
  621. }
  622. Extension ext;
  623. SPIRV_CROSS_DECLARE_CLONE(SPIRExtension)
  624. };
  625. // SPIREntryPoint is not a variant since its IDs are used to decorate OpFunction,
  626. // so in order to avoid conflicts, we can't stick them in the ids array.
  627. struct SPIREntryPoint
  628. {
  629. SPIREntryPoint(FunctionID self_, spv::ExecutionModel execution_model, const std::string &entry_name)
  630. : self(self_)
  631. , name(entry_name)
  632. , orig_name(entry_name)
  633. , model(execution_model)
  634. {
  635. }
  636. SPIREntryPoint() = default;
  637. FunctionID self = 0;
  638. std::string name;
  639. std::string orig_name;
  640. std::unordered_map<uint32_t, uint32_t> fp_fast_math_defaults;
  641. bool signed_zero_inf_nan_preserve_8 = false;
  642. bool signed_zero_inf_nan_preserve_16 = false;
  643. bool signed_zero_inf_nan_preserve_32 = false;
  644. bool signed_zero_inf_nan_preserve_64 = false;
  645. SmallVector<VariableID> interface_variables;
  646. Bitset flags;
  647. struct WorkgroupSize
  648. {
  649. uint32_t x = 0, y = 0, z = 0;
  650. uint32_t id_x = 0, id_y = 0, id_z = 0;
  651. uint32_t constant = 0; // Workgroup size can be expressed as a constant/spec-constant instead.
  652. } workgroup_size;
  653. uint32_t invocations = 0;
  654. uint32_t output_vertices = 0;
  655. uint32_t output_primitives = 0;
  656. spv::ExecutionModel model = spv::ExecutionModelMax;
  657. bool geometry_passthrough = false;
  658. };
  659. struct SPIRExpression : IVariant
  660. {
  661. enum
  662. {
  663. type = TypeExpression
  664. };
  665. // Only created by the backend target to avoid creating tons of temporaries.
  666. SPIRExpression(std::string expr, TypeID expression_type_, bool immutable_)
  667. : expression(std::move(expr))
  668. , expression_type(expression_type_)
  669. , immutable(immutable_)
  670. {
  671. }
  672. // If non-zero, prepend expression with to_expression(base_expression).
  673. // Used in amortizing multiple calls to to_expression()
  674. // where in certain cases that would quickly force a temporary when not needed.
  675. ID base_expression = 0;
  676. std::string expression;
  677. TypeID expression_type = 0;
  678. // If this expression is a forwarded load,
  679. // allow us to reference the original variable.
  680. ID loaded_from = 0;
  681. // If this expression will never change, we can avoid lots of temporaries
  682. // in high level source.
  683. // An expression being immutable can be speculative,
  684. // it is assumed that this is true almost always.
  685. bool immutable = false;
  686. // Before use, this expression must be transposed.
  687. // This is needed for targets which don't support row_major layouts.
  688. bool need_transpose = false;
  689. // Whether or not this is an access chain expression.
  690. bool access_chain = false;
  691. // Whether or not gl_MeshVerticesEXT[].gl_Position (as a whole or .y) is referenced
  692. bool access_meshlet_position_y = false;
  693. // A list of expressions which this expression depends on.
  694. SmallVector<ID> expression_dependencies;
  695. // Similar as expression dependencies, but does not stop the tracking for force-temporary variables.
  696. // We need to know the full chain from store back to any SSA variable.
  697. SmallVector<ID> invariance_dependencies;
  698. // By reading this expression, we implicitly read these expressions as well.
  699. // Used by access chain Store and Load since we read multiple expressions in this case.
  700. SmallVector<ID> implied_read_expressions;
  701. // The expression was emitted at a certain scope. Lets us track when an expression read means multiple reads.
  702. uint32_t emitted_loop_level = 0;
  703. SPIRV_CROSS_DECLARE_CLONE(SPIRExpression)
  704. };
  705. struct SPIRFunctionPrototype : IVariant
  706. {
  707. enum
  708. {
  709. type = TypeFunctionPrototype
  710. };
  711. explicit SPIRFunctionPrototype(TypeID return_type_)
  712. : return_type(return_type_)
  713. {
  714. }
  715. TypeID return_type;
  716. SmallVector<uint32_t> parameter_types;
  717. SPIRV_CROSS_DECLARE_CLONE(SPIRFunctionPrototype)
  718. };
  719. struct SPIRBlock : IVariant
  720. {
  721. enum
  722. {
  723. type = TypeBlock
  724. };
  725. enum Terminator
  726. {
  727. Unknown,
  728. Direct, // Emit next block directly without a particular condition.
  729. Select, // Block ends with an if/else block.
  730. MultiSelect, // Block ends with switch statement.
  731. Return, // Block ends with return.
  732. Unreachable, // Noop
  733. Kill, // Discard
  734. IgnoreIntersection, // Ray Tracing
  735. TerminateRay, // Ray Tracing
  736. EmitMeshTasks // Mesh shaders
  737. };
  738. enum Merge
  739. {
  740. MergeNone,
  741. MergeLoop,
  742. MergeSelection
  743. };
  744. enum Hints
  745. {
  746. HintNone,
  747. HintUnroll,
  748. HintDontUnroll,
  749. HintFlatten,
  750. HintDontFlatten
  751. };
  752. enum Method
  753. {
  754. MergeToSelectForLoop,
  755. MergeToDirectForLoop,
  756. MergeToSelectContinueForLoop
  757. };
  758. enum ContinueBlockType
  759. {
  760. ContinueNone,
  761. // Continue block is branchless and has at least one instruction.
  762. ForLoop,
  763. // Noop continue block.
  764. WhileLoop,
  765. // Continue block is conditional.
  766. DoWhileLoop,
  767. // Highly unlikely that anything will use this,
  768. // since it is really awkward/impossible to express in GLSL.
  769. ComplexLoop
  770. };
  771. enum : uint32_t
  772. {
  773. NoDominator = 0xffffffffu
  774. };
  775. Terminator terminator = Unknown;
  776. Merge merge = MergeNone;
  777. Hints hint = HintNone;
  778. BlockID next_block = 0;
  779. BlockID merge_block = 0;
  780. BlockID continue_block = 0;
  781. ID return_value = 0; // If 0, return nothing (void).
  782. ID condition = 0;
  783. BlockID true_block = 0;
  784. BlockID false_block = 0;
  785. BlockID default_block = 0;
  786. // If terminator is EmitMeshTasksEXT.
  787. struct
  788. {
  789. ID groups[3];
  790. ID payload;
  791. } mesh = {};
  792. SmallVector<Instruction> ops;
  793. struct Phi
  794. {
  795. ID local_variable; // flush local variable ...
  796. BlockID parent; // If we're in from_block and want to branch into this block ...
  797. VariableID function_variable; // to this function-global "phi" variable first.
  798. };
  799. // Before entering this block flush out local variables to magical "phi" variables.
  800. SmallVector<Phi> phi_variables;
  801. // Declare these temporaries before beginning the block.
  802. // Used for handling complex continue blocks which have side effects.
  803. SmallVector<std::pair<TypeID, ID>> declare_temporary;
  804. // Declare these temporaries, but only conditionally if this block turns out to be
  805. // a complex loop header.
  806. SmallVector<std::pair<TypeID, ID>> potential_declare_temporary;
  807. struct Case
  808. {
  809. uint64_t value;
  810. BlockID block;
  811. };
  812. SmallVector<Case> cases_32bit;
  813. SmallVector<Case> cases_64bit;
  814. // If we have tried to optimize code for this block but failed,
  815. // keep track of this.
  816. bool disable_block_optimization = false;
  817. // If the continue block is complex, fallback to "dumb" for loops.
  818. bool complex_continue = false;
  819. // Do we need a ladder variable to defer breaking out of a loop construct after a switch block?
  820. bool need_ladder_break = false;
  821. // If marked, we have explicitly handled Phi from this block, so skip any flushes related to that on a branch.
  822. // Used to handle an edge case with switch and case-label fallthrough where fall-through writes to Phi.
  823. BlockID ignore_phi_from_block = 0;
  824. // The dominating block which this block might be within.
  825. // Used in continue; blocks to determine if we really need to write continue.
  826. BlockID loop_dominator = 0;
  827. // All access to these variables are dominated by this block,
  828. // so before branching anywhere we need to make sure that we declare these variables.
  829. SmallVector<VariableID> dominated_variables;
  830. SmallVector<bool> rearm_dominated_variables;
  831. // These are variables which should be declared in a for loop header, if we
  832. // fail to use a classic for-loop,
  833. // we remove these variables, and fall back to regular variables outside the loop.
  834. SmallVector<VariableID> loop_variables;
  835. // Some expressions are control-flow dependent, i.e. any instruction which relies on derivatives or
  836. // sub-group-like operations.
  837. // Make sure that we only use these expressions in the original block.
  838. SmallVector<ID> invalidate_expressions;
  839. SPIRV_CROSS_DECLARE_CLONE(SPIRBlock)
  840. };
  841. struct SPIRFunction : IVariant
  842. {
  843. enum
  844. {
  845. type = TypeFunction
  846. };
  847. SPIRFunction(TypeID return_type_, TypeID function_type_)
  848. : return_type(return_type_)
  849. , function_type(function_type_)
  850. {
  851. }
  852. struct Parameter
  853. {
  854. TypeID type;
  855. ID id;
  856. uint32_t read_count;
  857. uint32_t write_count;
  858. // Set to true if this parameter aliases a global variable,
  859. // used mostly in Metal where global variables
  860. // have to be passed down to functions as regular arguments.
  861. // However, for this kind of variable, we should not care about
  862. // read and write counts as access to the function arguments
  863. // is not local to the function in question.
  864. bool alias_global_variable;
  865. };
  866. // When calling a function, and we're remapping separate image samplers,
  867. // resolve these arguments into combined image samplers and pass them
  868. // as additional arguments in this order.
  869. // It gets more complicated as functions can pull in their own globals
  870. // and combine them with parameters,
  871. // so we need to distinguish if something is local parameter index
  872. // or a global ID.
  873. struct CombinedImageSamplerParameter
  874. {
  875. VariableID id;
  876. VariableID image_id;
  877. VariableID sampler_id;
  878. bool global_image;
  879. bool global_sampler;
  880. bool depth;
  881. };
  882. TypeID return_type;
  883. TypeID function_type;
  884. SmallVector<Parameter> arguments;
  885. // Can be used by backends to add magic arguments.
  886. // Currently used by combined image/sampler implementation.
  887. SmallVector<Parameter> shadow_arguments;
  888. SmallVector<VariableID> local_variables;
  889. BlockID entry_block = 0;
  890. SmallVector<BlockID> blocks;
  891. SmallVector<CombinedImageSamplerParameter> combined_parameters;
  892. struct EntryLine
  893. {
  894. uint32_t file_id = 0;
  895. uint32_t line_literal = 0;
  896. };
  897. EntryLine entry_line;
  898. void add_local_variable(VariableID id)
  899. {
  900. local_variables.push_back(id);
  901. }
  902. void add_parameter(TypeID parameter_type, ID id, bool alias_global_variable = false)
  903. {
  904. // Arguments are read-only until proven otherwise.
  905. arguments.push_back({ parameter_type, id, 0u, 0u, alias_global_variable });
  906. }
  907. // Hooks to be run when the function returns.
  908. // Mostly used for lowering internal data structures onto flattened structures.
  909. // Need to defer this, because they might rely on things which change during compilation.
  910. // Intentionally not a small vector, this one is rare, and std::function can be large.
  911. Vector<std::function<void()>> fixup_hooks_out;
  912. // Hooks to be run when the function begins.
  913. // Mostly used for populating internal data structures from flattened structures.
  914. // Need to defer this, because they might rely on things which change during compilation.
  915. // Intentionally not a small vector, this one is rare, and std::function can be large.
  916. Vector<std::function<void()>> fixup_hooks_in;
  917. // On function entry, make sure to copy a constant array into thread addr space to work around
  918. // the case where we are passing a constant array by value to a function on backends which do not
  919. // consider arrays value types.
  920. SmallVector<ID> constant_arrays_needed_on_stack;
  921. // Does this function (or any function called by it), emit geometry?
  922. bool emits_geometry = false;
  923. bool active = false;
  924. bool flush_undeclared = true;
  925. bool do_combined_parameters = true;
  926. SPIRV_CROSS_DECLARE_CLONE(SPIRFunction)
  927. };
  928. struct SPIRAccessChain : IVariant
  929. {
  930. enum
  931. {
  932. type = TypeAccessChain
  933. };
  934. SPIRAccessChain(TypeID basetype_, spv::StorageClass storage_, std::string base_, std::string dynamic_index_,
  935. int32_t static_index_)
  936. : basetype(basetype_)
  937. , storage(storage_)
  938. , base(std::move(base_))
  939. , dynamic_index(std::move(dynamic_index_))
  940. , static_index(static_index_)
  941. {
  942. }
  943. // The access chain represents an offset into a buffer.
  944. // Some backends need more complicated handling of access chains to be able to use buffers, like HLSL
  945. // which has no usable buffer type ala GLSL SSBOs.
  946. // StructuredBuffer is too limited, so our only option is to deal with ByteAddressBuffer which works with raw addresses.
  947. TypeID basetype;
  948. spv::StorageClass storage;
  949. std::string base;
  950. std::string dynamic_index;
  951. int32_t static_index;
  952. VariableID loaded_from = 0;
  953. uint32_t matrix_stride = 0;
  954. uint32_t array_stride = 0;
  955. bool row_major_matrix = false;
  956. bool immutable = false;
  957. // By reading this expression, we implicitly read these expressions as well.
  958. // Used by access chain Store and Load since we read multiple expressions in this case.
  959. SmallVector<ID> implied_read_expressions;
  960. SPIRV_CROSS_DECLARE_CLONE(SPIRAccessChain)
  961. };
  962. struct SPIRVariable : IVariant
  963. {
  964. enum
  965. {
  966. type = TypeVariable
  967. };
  968. SPIRVariable() = default;
  969. SPIRVariable(TypeID basetype_, spv::StorageClass storage_, ID initializer_ = 0, VariableID basevariable_ = 0)
  970. : basetype(basetype_)
  971. , storage(storage_)
  972. , initializer(initializer_)
  973. , basevariable(basevariable_)
  974. {
  975. }
  976. TypeID basetype = 0;
  977. spv::StorageClass storage = spv::StorageClassGeneric;
  978. uint32_t decoration = 0;
  979. ID initializer = 0;
  980. VariableID basevariable = 0;
  981. SmallVector<uint32_t> dereference_chain;
  982. bool compat_builtin = false;
  983. // If a variable is shadowed, we only statically assign to it
  984. // and never actually emit a statement for it.
  985. // When we read the variable as an expression, just forward
  986. // shadowed_id as the expression.
  987. bool statically_assigned = false;
  988. ID static_expression = 0;
  989. // Temporaries which can remain forwarded as long as this variable is not modified.
  990. SmallVector<ID> dependees;
  991. // ShaderDebugInfo local variables attached to this variable via DebugDeclare
  992. SmallVector<ID> debug_local_variables;
  993. bool deferred_declaration = false;
  994. bool phi_variable = false;
  995. // Used to deal with Phi variable flushes. See flush_phi().
  996. bool allocate_temporary_copy = false;
  997. bool remapped_variable = false;
  998. uint32_t remapped_components = 0;
  999. // The block which dominates all access to this variable.
  1000. BlockID dominator = 0;
  1001. // If true, this variable is a loop variable, when accessing the variable
  1002. // outside a loop,
  1003. // we should statically forward it.
  1004. bool loop_variable = false;
  1005. // Set to true while we're inside the for loop.
  1006. bool loop_variable_enable = false;
  1007. // Used to find global LUTs
  1008. bool is_written_to = false;
  1009. SPIRFunction::Parameter *parameter = nullptr;
  1010. SPIRV_CROSS_DECLARE_CLONE(SPIRVariable)
  1011. };
  1012. struct SPIRConstant : IVariant
  1013. {
  1014. enum
  1015. {
  1016. type = TypeConstant
  1017. };
  1018. union Constant
  1019. {
  1020. uint32_t u32;
  1021. int32_t i32;
  1022. float f32;
  1023. uint64_t u64;
  1024. int64_t i64;
  1025. double f64;
  1026. };
  1027. struct ConstantVector
  1028. {
  1029. Constant r[4];
  1030. // If != 0, this element is a specialization constant, and we should keep track of it as such.
  1031. ID id[4];
  1032. uint32_t vecsize = 1;
  1033. ConstantVector()
  1034. {
  1035. memset(r, 0, sizeof(r));
  1036. }
  1037. };
  1038. struct ConstantMatrix
  1039. {
  1040. ConstantVector c[4];
  1041. // If != 0, this column is a specialization constant, and we should keep track of it as such.
  1042. ID id[4];
  1043. uint32_t columns = 1;
  1044. };
  1045. static inline float f16_to_f32(uint16_t u16_value)
  1046. {
  1047. // Based on the GLM implementation.
  1048. int s = (u16_value >> 15) & 0x1;
  1049. int e = (u16_value >> 10) & 0x1f;
  1050. int m = (u16_value >> 0) & 0x3ff;
  1051. union
  1052. {
  1053. float f32;
  1054. uint32_t u32;
  1055. } u;
  1056. if (e == 0)
  1057. {
  1058. if (m == 0)
  1059. {
  1060. u.u32 = uint32_t(s) << 31;
  1061. return u.f32;
  1062. }
  1063. else
  1064. {
  1065. while ((m & 0x400) == 0)
  1066. {
  1067. m <<= 1;
  1068. e--;
  1069. }
  1070. e++;
  1071. m &= ~0x400;
  1072. }
  1073. }
  1074. else if (e == 31)
  1075. {
  1076. if (m == 0)
  1077. {
  1078. u.u32 = (uint32_t(s) << 31) | 0x7f800000u;
  1079. return u.f32;
  1080. }
  1081. else
  1082. {
  1083. u.u32 = (uint32_t(s) << 31) | 0x7f800000u | (m << 13);
  1084. return u.f32;
  1085. }
  1086. }
  1087. e += 127 - 15;
  1088. m <<= 13;
  1089. u.u32 = (uint32_t(s) << 31) | (e << 23) | m;
  1090. return u.f32;
  1091. }
  1092. static inline float fe4m3_to_f32(uint8_t v)
  1093. {
  1094. if ((v & 0x7f) == 0x7f)
  1095. {
  1096. union
  1097. {
  1098. float f32;
  1099. uint32_t u32;
  1100. } u;
  1101. u.u32 = (v & 0x80) ? 0xffffffffu : 0x7fffffffu;
  1102. return u.f32;
  1103. }
  1104. else
  1105. {
  1106. // Reuse the FP16 to FP32 code. Cute bit-hackery.
  1107. return f16_to_f32((int16_t(int8_t(v)) << 7) & (0xffff ^ 0x4000)) * 256.0f;
  1108. }
  1109. }
  1110. inline uint32_t specialization_constant_id(uint32_t col, uint32_t row) const
  1111. {
  1112. return m.c[col].id[row];
  1113. }
  1114. inline uint32_t specialization_constant_id(uint32_t col) const
  1115. {
  1116. return m.id[col];
  1117. }
  1118. inline uint32_t scalar(uint32_t col = 0, uint32_t row = 0) const
  1119. {
  1120. return m.c[col].r[row].u32;
  1121. }
  1122. inline int16_t scalar_i16(uint32_t col = 0, uint32_t row = 0) const
  1123. {
  1124. return int16_t(m.c[col].r[row].u32 & 0xffffu);
  1125. }
  1126. inline uint16_t scalar_u16(uint32_t col = 0, uint32_t row = 0) const
  1127. {
  1128. return uint16_t(m.c[col].r[row].u32 & 0xffffu);
  1129. }
  1130. inline int8_t scalar_i8(uint32_t col = 0, uint32_t row = 0) const
  1131. {
  1132. return int8_t(m.c[col].r[row].u32 & 0xffu);
  1133. }
  1134. inline uint8_t scalar_u8(uint32_t col = 0, uint32_t row = 0) const
  1135. {
  1136. return uint8_t(m.c[col].r[row].u32 & 0xffu);
  1137. }
  1138. inline float scalar_f16(uint32_t col = 0, uint32_t row = 0) const
  1139. {
  1140. return f16_to_f32(scalar_u16(col, row));
  1141. }
  1142. inline float scalar_bf16(uint32_t col = 0, uint32_t row = 0) const
  1143. {
  1144. uint32_t v = scalar_u16(col, row) << 16;
  1145. float fp32;
  1146. memcpy(&fp32, &v, sizeof(float));
  1147. return fp32;
  1148. }
  1149. inline float scalar_floate4m3(uint32_t col = 0, uint32_t row = 0) const
  1150. {
  1151. return fe4m3_to_f32(scalar_u8(col, row));
  1152. }
  1153. inline float scalar_bf8(uint32_t col = 0, uint32_t row = 0) const
  1154. {
  1155. return f16_to_f32(uint16_t(scalar_u8(col, row) << 8));
  1156. }
  1157. inline float scalar_f32(uint32_t col = 0, uint32_t row = 0) const
  1158. {
  1159. return m.c[col].r[row].f32;
  1160. }
  1161. inline int32_t scalar_i32(uint32_t col = 0, uint32_t row = 0) const
  1162. {
  1163. return m.c[col].r[row].i32;
  1164. }
  1165. inline double scalar_f64(uint32_t col = 0, uint32_t row = 0) const
  1166. {
  1167. return m.c[col].r[row].f64;
  1168. }
  1169. inline int64_t scalar_i64(uint32_t col = 0, uint32_t row = 0) const
  1170. {
  1171. return m.c[col].r[row].i64;
  1172. }
  1173. inline uint64_t scalar_u64(uint32_t col = 0, uint32_t row = 0) const
  1174. {
  1175. return m.c[col].r[row].u64;
  1176. }
  1177. inline const ConstantVector &vector() const
  1178. {
  1179. return m.c[0];
  1180. }
  1181. inline uint32_t vector_size() const
  1182. {
  1183. return m.c[0].vecsize;
  1184. }
  1185. inline uint32_t columns() const
  1186. {
  1187. return m.columns;
  1188. }
  1189. inline void make_null(const SPIRType &constant_type_)
  1190. {
  1191. m = {};
  1192. m.columns = constant_type_.columns;
  1193. for (auto &c : m.c)
  1194. c.vecsize = constant_type_.vecsize;
  1195. }
  1196. inline bool constant_is_null() const
  1197. {
  1198. if (specialization)
  1199. return false;
  1200. if (!subconstants.empty())
  1201. return false;
  1202. for (uint32_t col = 0; col < columns(); col++)
  1203. for (uint32_t row = 0; row < vector_size(); row++)
  1204. if (scalar_u64(col, row) != 0)
  1205. return false;
  1206. return true;
  1207. }
  1208. explicit SPIRConstant(uint32_t constant_type_)
  1209. : constant_type(constant_type_)
  1210. {
  1211. }
  1212. SPIRConstant() = default;
  1213. SPIRConstant(TypeID constant_type_, const uint32_t *elements, uint32_t num_elements, bool specialized, bool replicated_ = false)
  1214. : constant_type(constant_type_)
  1215. , specialization(specialized)
  1216. , replicated(replicated_)
  1217. {
  1218. subconstants.reserve(num_elements);
  1219. for (uint32_t i = 0; i < num_elements; i++)
  1220. subconstants.push_back(elements[i]);
  1221. specialization = specialized;
  1222. }
  1223. // Construct scalar (32-bit).
  1224. SPIRConstant(TypeID constant_type_, uint32_t v0, bool specialized)
  1225. : constant_type(constant_type_)
  1226. , specialization(specialized)
  1227. {
  1228. m.c[0].r[0].u32 = v0;
  1229. m.c[0].vecsize = 1;
  1230. m.columns = 1;
  1231. }
  1232. // Construct scalar (64-bit).
  1233. SPIRConstant(TypeID constant_type_, uint64_t v0, bool specialized)
  1234. : constant_type(constant_type_)
  1235. , specialization(specialized)
  1236. {
  1237. m.c[0].r[0].u64 = v0;
  1238. m.c[0].vecsize = 1;
  1239. m.columns = 1;
  1240. }
  1241. // Construct vectors and matrices.
  1242. SPIRConstant(TypeID constant_type_, const SPIRConstant *const *vector_elements, uint32_t num_elements,
  1243. bool specialized)
  1244. : constant_type(constant_type_)
  1245. , specialization(specialized)
  1246. {
  1247. bool matrix = vector_elements[0]->m.c[0].vecsize > 1;
  1248. if (matrix)
  1249. {
  1250. m.columns = num_elements;
  1251. for (uint32_t i = 0; i < num_elements; i++)
  1252. {
  1253. m.c[i] = vector_elements[i]->m.c[0];
  1254. if (vector_elements[i]->specialization)
  1255. m.id[i] = vector_elements[i]->self;
  1256. }
  1257. }
  1258. else
  1259. {
  1260. m.c[0].vecsize = num_elements;
  1261. m.columns = 1;
  1262. for (uint32_t i = 0; i < num_elements; i++)
  1263. {
  1264. m.c[0].r[i] = vector_elements[i]->m.c[0].r[0];
  1265. if (vector_elements[i]->specialization)
  1266. m.c[0].id[i] = vector_elements[i]->self;
  1267. }
  1268. }
  1269. }
  1270. TypeID constant_type = 0;
  1271. ConstantMatrix m;
  1272. // If this constant is a specialization constant (i.e. created with OpSpecConstant*).
  1273. bool specialization = false;
  1274. // If this constant is used as an array length which creates specialization restrictions on some backends.
  1275. bool is_used_as_array_length = false;
  1276. // If true, this is a LUT, and should always be declared in the outer scope.
  1277. bool is_used_as_lut = false;
  1278. // If this is a null constant of array type with specialized length.
  1279. // May require special handling in initializer
  1280. bool is_null_array_specialized_length = false;
  1281. // For composites which are constant arrays, etc.
  1282. SmallVector<ConstantID> subconstants;
  1283. // Whether the subconstants are intended to be replicated (e.g. OpConstantCompositeReplicateEXT)
  1284. bool replicated = false;
  1285. // Non-Vulkan GLSL, HLSL and sometimes MSL emits defines for each specialization constant,
  1286. // and uses them to initialize the constant. This allows the user
  1287. // to still be able to specialize the value by supplying corresponding
  1288. // preprocessor directives before compiling the shader.
  1289. std::string specialization_constant_macro_name;
  1290. SPIRV_CROSS_DECLARE_CLONE(SPIRConstant)
  1291. };
  1292. // Variants have a very specific allocation scheme.
  1293. struct ObjectPoolGroup
  1294. {
  1295. std::unique_ptr<ObjectPoolBase> pools[TypeCount];
  1296. };
  1297. class Variant
  1298. {
  1299. public:
  1300. explicit Variant(ObjectPoolGroup *group_)
  1301. : group(group_)
  1302. {
  1303. }
  1304. ~Variant()
  1305. {
  1306. if (holder)
  1307. group->pools[type]->deallocate_opaque(holder);
  1308. }
  1309. // Marking custom move constructor as noexcept is important.
  1310. Variant(Variant &&other) SPIRV_CROSS_NOEXCEPT
  1311. {
  1312. *this = std::move(other);
  1313. }
  1314. // We cannot copy from other variant without our own pool group.
  1315. // Have to explicitly copy.
  1316. Variant(const Variant &variant) = delete;
  1317. // Marking custom move constructor as noexcept is important.
  1318. Variant &operator=(Variant &&other) SPIRV_CROSS_NOEXCEPT
  1319. {
  1320. if (this != &other)
  1321. {
  1322. if (holder)
  1323. group->pools[type]->deallocate_opaque(holder);
  1324. holder = other.holder;
  1325. group = other.group;
  1326. type = other.type;
  1327. allow_type_rewrite = other.allow_type_rewrite;
  1328. other.holder = nullptr;
  1329. other.type = TypeNone;
  1330. }
  1331. return *this;
  1332. }
  1333. // This copy/clone should only be called in the Compiler constructor.
  1334. // If this is called inside ::compile(), we invalidate any references we took higher in the stack.
  1335. // This should never happen.
  1336. Variant &operator=(const Variant &other)
  1337. {
  1338. //#define SPIRV_CROSS_COPY_CONSTRUCTOR_SANITIZE
  1339. #ifdef SPIRV_CROSS_COPY_CONSTRUCTOR_SANITIZE
  1340. abort();
  1341. #endif
  1342. if (this != &other)
  1343. {
  1344. if (holder)
  1345. group->pools[type]->deallocate_opaque(holder);
  1346. if (other.holder)
  1347. holder = other.holder->clone(group->pools[other.type].get());
  1348. else
  1349. holder = nullptr;
  1350. type = other.type;
  1351. allow_type_rewrite = other.allow_type_rewrite;
  1352. }
  1353. return *this;
  1354. }
  1355. void set(IVariant *val, Types new_type)
  1356. {
  1357. if (holder)
  1358. group->pools[type]->deallocate_opaque(holder);
  1359. holder = nullptr;
  1360. if (!allow_type_rewrite && type != TypeNone && type != new_type)
  1361. {
  1362. if (val)
  1363. group->pools[new_type]->deallocate_opaque(val);
  1364. SPIRV_CROSS_THROW("Overwriting a variant with new type.");
  1365. }
  1366. holder = val;
  1367. type = new_type;
  1368. allow_type_rewrite = false;
  1369. }
  1370. template <typename T, typename... Ts>
  1371. T *allocate_and_set(Types new_type, Ts &&... ts)
  1372. {
  1373. T *val = static_cast<ObjectPool<T> &>(*group->pools[new_type]).allocate(std::forward<Ts>(ts)...);
  1374. set(val, new_type);
  1375. return val;
  1376. }
  1377. template <typename T>
  1378. T &get()
  1379. {
  1380. if (!holder)
  1381. SPIRV_CROSS_THROW("nullptr");
  1382. if (static_cast<Types>(T::type) != type)
  1383. SPIRV_CROSS_THROW("Bad cast");
  1384. return *static_cast<T *>(holder);
  1385. }
  1386. template <typename T>
  1387. const T &get() const
  1388. {
  1389. if (!holder)
  1390. SPIRV_CROSS_THROW("nullptr");
  1391. if (static_cast<Types>(T::type) != type)
  1392. SPIRV_CROSS_THROW("Bad cast");
  1393. return *static_cast<const T *>(holder);
  1394. }
  1395. Types get_type() const
  1396. {
  1397. return type;
  1398. }
  1399. ID get_id() const
  1400. {
  1401. return holder ? holder->self : ID(0);
  1402. }
  1403. bool empty() const
  1404. {
  1405. return !holder;
  1406. }
  1407. void reset()
  1408. {
  1409. if (holder)
  1410. group->pools[type]->deallocate_opaque(holder);
  1411. holder = nullptr;
  1412. type = TypeNone;
  1413. }
  1414. void set_allow_type_rewrite()
  1415. {
  1416. allow_type_rewrite = true;
  1417. }
  1418. private:
  1419. ObjectPoolGroup *group = nullptr;
  1420. IVariant *holder = nullptr;
  1421. Types type = TypeNone;
  1422. bool allow_type_rewrite = false;
  1423. };
  1424. template <typename T>
  1425. T &variant_get(Variant &var)
  1426. {
  1427. return var.get<T>();
  1428. }
  1429. template <typename T>
  1430. const T &variant_get(const Variant &var)
  1431. {
  1432. return var.get<T>();
  1433. }
  1434. template <typename T, typename... P>
  1435. T &variant_set(Variant &var, P &&... args)
  1436. {
  1437. auto *ptr = var.allocate_and_set<T>(static_cast<Types>(T::type), std::forward<P>(args)...);
  1438. return *ptr;
  1439. }
  1440. struct AccessChainMeta
  1441. {
  1442. uint32_t storage_physical_type = 0;
  1443. bool need_transpose = false;
  1444. bool storage_is_packed = false;
  1445. bool storage_is_invariant = false;
  1446. bool flattened_struct = false;
  1447. bool relaxed_precision = false;
  1448. bool access_meshlet_position_y = false;
  1449. bool chain_is_builtin = false;
  1450. spv::BuiltIn builtin = {};
  1451. };
  1452. enum ExtendedDecorations
  1453. {
  1454. // Marks if a buffer block is re-packed, i.e. member declaration might be subject to PhysicalTypeID remapping and padding.
  1455. SPIRVCrossDecorationBufferBlockRepacked = 0,
  1456. // A type in a buffer block might be declared with a different physical type than the logical type.
  1457. // If this is not set, PhysicalTypeID == the SPIR-V type as declared.
  1458. SPIRVCrossDecorationPhysicalTypeID,
  1459. // Marks if the physical type is to be declared with tight packing rules, i.e. packed_floatN on MSL and friends.
  1460. // If this is set, PhysicalTypeID might also be set. It can be set to same as logical type if all we're doing
  1461. // is converting float3 to packed_float3 for example.
  1462. // If this is marked on a struct, it means the struct itself must use only Packed types for all its members.
  1463. SPIRVCrossDecorationPhysicalTypePacked,
  1464. // The padding in bytes before declaring this struct member.
  1465. // If used on a struct type, marks the target size of a struct.
  1466. SPIRVCrossDecorationPaddingTarget,
  1467. SPIRVCrossDecorationInterfaceMemberIndex,
  1468. SPIRVCrossDecorationInterfaceOrigID,
  1469. SPIRVCrossDecorationResourceIndexPrimary,
  1470. // Used for decorations like resource indices for samplers when part of combined image samplers.
  1471. // A variable might need to hold two resource indices in this case.
  1472. SPIRVCrossDecorationResourceIndexSecondary,
  1473. // Used for resource indices for multiplanar images when part of combined image samplers.
  1474. SPIRVCrossDecorationResourceIndexTertiary,
  1475. SPIRVCrossDecorationResourceIndexQuaternary,
  1476. // Marks a buffer block for using explicit offsets (GLSL/HLSL).
  1477. SPIRVCrossDecorationExplicitOffset,
  1478. // Apply to a variable in the Input storage class; marks it as holding the base group passed to vkCmdDispatchBase(),
  1479. // or the base vertex and instance indices passed to vkCmdDrawIndexed().
  1480. // In MSL, this is used to adjust the WorkgroupId and GlobalInvocationId variables in compute shaders,
  1481. // and to hold the BaseVertex and BaseInstance variables in vertex shaders.
  1482. SPIRVCrossDecorationBuiltInDispatchBase,
  1483. // Apply to a variable that is a function parameter; marks it as being a "dynamic"
  1484. // combined image-sampler. In MSL, this is used when a function parameter might hold
  1485. // either a regular combined image-sampler or one that has an attached sampler
  1486. // Y'CbCr conversion.
  1487. SPIRVCrossDecorationDynamicImageSampler,
  1488. // Apply to a variable in the Input storage class; marks it as holding the size of the stage
  1489. // input grid.
  1490. // In MSL, this is used to hold the vertex and instance counts in a tessellation pipeline
  1491. // vertex shader.
  1492. SPIRVCrossDecorationBuiltInStageInputSize,
  1493. // Apply to any access chain of a tessellation I/O variable; stores the type of the sub-object
  1494. // that was chained to, as recorded in the input variable itself. This is used in case the pointer
  1495. // is itself used as the base of an access chain, to calculate the original type of the sub-object
  1496. // chained to, in case a swizzle needs to be applied. This should not happen normally with valid
  1497. // SPIR-V, but the MSL backend can change the type of input variables, necessitating the
  1498. // addition of swizzles to keep the generated code compiling.
  1499. SPIRVCrossDecorationTessIOOriginalInputTypeID,
  1500. // Apply to any access chain of an interface variable used with pull-model interpolation, where the variable is a
  1501. // vector but the resulting pointer is a scalar; stores the component index that is to be accessed by the chain.
  1502. // This is used when emitting calls to interpolation functions on the chain in MSL: in this case, the component
  1503. // must be applied to the result, since pull-model interpolants in MSL cannot be swizzled directly, but the
  1504. // results of interpolation can.
  1505. SPIRVCrossDecorationInterpolantComponentExpr,
  1506. // Apply to any struct type that is used in the Workgroup storage class.
  1507. // This causes matrices in MSL prior to Metal 3.0 to be emitted using a special
  1508. // class that is convertible to the standard matrix type, to work around the
  1509. // lack of constructors in the 'threadgroup' address space.
  1510. SPIRVCrossDecorationWorkgroupStruct,
  1511. SPIRVCrossDecorationOverlappingBinding,
  1512. SPIRVCrossDecorationCount
  1513. };
  1514. struct Meta
  1515. {
  1516. struct Decoration
  1517. {
  1518. std::string alias;
  1519. std::string qualified_alias;
  1520. std::string user_semantic;
  1521. std::string user_type;
  1522. Bitset decoration_flags;
  1523. spv::BuiltIn builtin_type = spv::BuiltInMax;
  1524. uint32_t location = 0;
  1525. uint32_t component = 0;
  1526. uint32_t set = 0;
  1527. uint32_t binding = 0;
  1528. uint32_t offset = 0;
  1529. uint32_t xfb_buffer = 0;
  1530. uint32_t xfb_stride = 0;
  1531. uint32_t stream = 0;
  1532. uint32_t array_stride = 0;
  1533. uint32_t matrix_stride = 0;
  1534. uint32_t input_attachment = 0;
  1535. uint32_t spec_id = 0;
  1536. uint32_t index = 0;
  1537. spv::FPRoundingMode fp_rounding_mode = spv::FPRoundingModeMax;
  1538. spv::FPFastMathModeMask fp_fast_math_mode = spv::FPFastMathModeMaskNone;
  1539. bool builtin = false;
  1540. bool qualified_alias_explicit_override = false;
  1541. struct Extended
  1542. {
  1543. Extended()
  1544. {
  1545. // MSVC 2013 workaround to init like this.
  1546. for (auto &v : values)
  1547. v = 0;
  1548. }
  1549. Bitset flags;
  1550. uint32_t values[SPIRVCrossDecorationCount];
  1551. } extended;
  1552. };
  1553. Decoration decoration;
  1554. // Intentionally not a SmallVector. Decoration is large and somewhat rare.
  1555. Vector<Decoration> members;
  1556. std::unordered_map<uint32_t, uint32_t> decoration_word_offset;
  1557. // For SPV_GOOGLE_hlsl_functionality1.
  1558. bool hlsl_is_magic_counter_buffer = false;
  1559. // ID for the sibling counter buffer.
  1560. uint32_t hlsl_magic_counter_buffer = 0;
  1561. };
  1562. // A user callback that remaps the type of any variable.
  1563. // var_name is the declared name of the variable.
  1564. // name_of_type is the textual name of the type which will be used in the code unless written to by the callback.
  1565. using VariableTypeRemapCallback =
  1566. std::function<void(const SPIRType &type, const std::string &var_name, std::string &name_of_type)>;
  1567. class Hasher
  1568. {
  1569. public:
  1570. inline void u32(uint32_t value)
  1571. {
  1572. h = (h * 0x100000001b3ull) ^ value;
  1573. }
  1574. inline uint64_t get() const
  1575. {
  1576. return h;
  1577. }
  1578. private:
  1579. uint64_t h = 0xcbf29ce484222325ull;
  1580. };
  1581. static inline bool type_is_floating_point(const SPIRType &type)
  1582. {
  1583. return type.basetype == SPIRType::Half || type.basetype == SPIRType::Float || type.basetype == SPIRType::Double ||
  1584. type.basetype == SPIRType::BFloat16 || type.basetype == SPIRType::FloatE5M2 || type.basetype == SPIRType::FloatE4M3;
  1585. }
  1586. static inline bool type_is_integral(const SPIRType &type)
  1587. {
  1588. return type.basetype == SPIRType::SByte || type.basetype == SPIRType::UByte || type.basetype == SPIRType::Short ||
  1589. type.basetype == SPIRType::UShort || type.basetype == SPIRType::Int || type.basetype == SPIRType::UInt ||
  1590. type.basetype == SPIRType::Int64 || type.basetype == SPIRType::UInt64;
  1591. }
  1592. static inline SPIRType::BaseType to_signed_basetype(uint32_t width)
  1593. {
  1594. switch (width)
  1595. {
  1596. case 8:
  1597. return SPIRType::SByte;
  1598. case 16:
  1599. return SPIRType::Short;
  1600. case 32:
  1601. return SPIRType::Int;
  1602. case 64:
  1603. return SPIRType::Int64;
  1604. default:
  1605. SPIRV_CROSS_THROW("Invalid bit width.");
  1606. }
  1607. }
  1608. static inline SPIRType::BaseType to_unsigned_basetype(uint32_t width)
  1609. {
  1610. switch (width)
  1611. {
  1612. case 8:
  1613. return SPIRType::UByte;
  1614. case 16:
  1615. return SPIRType::UShort;
  1616. case 32:
  1617. return SPIRType::UInt;
  1618. case 64:
  1619. return SPIRType::UInt64;
  1620. default:
  1621. SPIRV_CROSS_THROW("Invalid bit width.");
  1622. }
  1623. }
  1624. // Returns true if an arithmetic operation does not change behavior depending on signedness.
  1625. static inline bool opcode_is_sign_invariant(spv::Op opcode)
  1626. {
  1627. switch (opcode)
  1628. {
  1629. case spv::OpIEqual:
  1630. case spv::OpINotEqual:
  1631. case spv::OpISub:
  1632. case spv::OpIAdd:
  1633. case spv::OpIMul:
  1634. case spv::OpShiftLeftLogical:
  1635. case spv::OpBitwiseOr:
  1636. case spv::OpBitwiseXor:
  1637. case spv::OpBitwiseAnd:
  1638. return true;
  1639. default:
  1640. return false;
  1641. }
  1642. }
  1643. static inline bool opcode_can_promote_integer_implicitly(spv::Op opcode)
  1644. {
  1645. switch (opcode)
  1646. {
  1647. case spv::OpSNegate:
  1648. case spv::OpNot:
  1649. case spv::OpBitwiseAnd:
  1650. case spv::OpBitwiseOr:
  1651. case spv::OpBitwiseXor:
  1652. case spv::OpShiftLeftLogical:
  1653. case spv::OpShiftRightLogical:
  1654. case spv::OpShiftRightArithmetic:
  1655. case spv::OpIAdd:
  1656. case spv::OpISub:
  1657. case spv::OpIMul:
  1658. case spv::OpSDiv:
  1659. case spv::OpUDiv:
  1660. case spv::OpSRem:
  1661. case spv::OpUMod:
  1662. case spv::OpSMod:
  1663. return true;
  1664. default:
  1665. return false;
  1666. }
  1667. }
  1668. struct SetBindingPair
  1669. {
  1670. uint32_t desc_set;
  1671. uint32_t binding;
  1672. inline bool operator==(const SetBindingPair &other) const
  1673. {
  1674. return desc_set == other.desc_set && binding == other.binding;
  1675. }
  1676. inline bool operator<(const SetBindingPair &other) const
  1677. {
  1678. return desc_set < other.desc_set || (desc_set == other.desc_set && binding < other.binding);
  1679. }
  1680. };
  1681. struct LocationComponentPair
  1682. {
  1683. uint32_t location;
  1684. uint32_t component;
  1685. inline bool operator==(const LocationComponentPair &other) const
  1686. {
  1687. return location == other.location && component == other.component;
  1688. }
  1689. inline bool operator<(const LocationComponentPair &other) const
  1690. {
  1691. return location < other.location || (location == other.location && component < other.component);
  1692. }
  1693. };
  1694. struct StageSetBinding
  1695. {
  1696. spv::ExecutionModel model;
  1697. uint32_t desc_set;
  1698. uint32_t binding;
  1699. inline bool operator==(const StageSetBinding &other) const
  1700. {
  1701. return model == other.model && desc_set == other.desc_set && binding == other.binding;
  1702. }
  1703. };
  1704. struct InternalHasher
  1705. {
  1706. inline size_t operator()(const SetBindingPair &value) const
  1707. {
  1708. // Quality of hash doesn't really matter here.
  1709. auto hash_set = std::hash<uint32_t>()(value.desc_set);
  1710. auto hash_binding = std::hash<uint32_t>()(value.binding);
  1711. return (hash_set * 0x10001b31) ^ hash_binding;
  1712. }
  1713. inline size_t operator()(const LocationComponentPair &value) const
  1714. {
  1715. // Quality of hash doesn't really matter here.
  1716. auto hash_set = std::hash<uint32_t>()(value.location);
  1717. auto hash_binding = std::hash<uint32_t>()(value.component);
  1718. return (hash_set * 0x10001b31) ^ hash_binding;
  1719. }
  1720. inline size_t operator()(const StageSetBinding &value) const
  1721. {
  1722. // Quality of hash doesn't really matter here.
  1723. auto hash_model = std::hash<uint32_t>()(value.model);
  1724. auto hash_set = std::hash<uint32_t>()(value.desc_set);
  1725. auto tmp_hash = (hash_model * 0x10001b31) ^ hash_set;
  1726. return (tmp_hash * 0x10001b31) ^ value.binding;
  1727. }
  1728. };
  1729. // Special constant used in a {MSL,HLSL}ResourceBinding desc_set
  1730. // element to indicate the bindings for the push constants.
  1731. static const uint32_t ResourceBindingPushConstantDescriptorSet = ~(0u);
  1732. // Special constant used in a {MSL,HLSL}ResourceBinding binding
  1733. // element to indicate the bindings for the push constants.
  1734. static const uint32_t ResourceBindingPushConstantBinding = 0;
  1735. } // namespace SPIRV_CROSS_NAMESPACE
  1736. namespace std
  1737. {
  1738. template <SPIRV_CROSS_NAMESPACE::Types type>
  1739. struct hash<SPIRV_CROSS_NAMESPACE::TypedID<type>>
  1740. {
  1741. size_t operator()(const SPIRV_CROSS_NAMESPACE::TypedID<type> &value) const
  1742. {
  1743. return std::hash<uint32_t>()(value);
  1744. }
  1745. };
  1746. } // namespace std
  1747. #ifdef SPIRV_CROSS_SPV_HEADER_NAMESPACE_OVERRIDE
  1748. #undef spv
  1749. #endif
  1750. #endif