ShaderProgramReflection.cpp 26 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622623624625626627628629630631632633634635636637638639640641642643644645646647648649650651652653654655656657658659660661662663664665666667668669670671672673674675676677678679680681682683684685686687688689690691692693694695696697698699700701702703704705706707708709710711712713714715716717718719720721722723724725726727728729730731732733734735736737738739740741742743744745746747748749750751752753754755756757758759760761762763764765766767768769770771772773774775776777778779780781782783784785786787788789790791792793794795796797798799800801802803804805806807808809810811812813814815816817818819820821822823824825826827828829830831832833834835836837838839840841842843844845846847848849850851852853854855856857858859860861862863864865866867868869870871872873874875876877878879880881882883884885886887888889890891892893894895896897898899900901902903904905906907908909910911912913914915916917918919920921922923924925926927928929930931932933934935936937938939940941942943944945946947948949950951952953954955956957958959960961962963964965966967968969970971972973974975
  1. // Copyright (C) 2009-2023, Panagiotis Christopoulos Charitos and contributors.
  2. // All rights reserved.
  3. // Code licensed under the BSD License.
  4. // http://www.anki3d.org/LICENSE
  5. #include <AnKi/ShaderCompiler/ShaderProgramReflection.h>
  6. #include <AnKi/Gr/Utils/Functions.h>
  7. #include <SpirvCross/spirv_glsl.hpp>
  8. namespace anki {
  9. static ShaderVariableDataType spirvcrossBaseTypeToAnki(spirv_cross::SPIRType::BaseType cross)
  10. {
  11. ShaderVariableDataType out = ShaderVariableDataType::kNone;
  12. switch(cross)
  13. {
  14. case spirv_cross::SPIRType::SByte:
  15. out = ShaderVariableDataType::kI8;
  16. break;
  17. case spirv_cross::SPIRType::UByte:
  18. out = ShaderVariableDataType::kU8;
  19. break;
  20. case spirv_cross::SPIRType::Short:
  21. out = ShaderVariableDataType::kI16;
  22. break;
  23. case spirv_cross::SPIRType::UShort:
  24. out = ShaderVariableDataType::kU16;
  25. break;
  26. case spirv_cross::SPIRType::Int:
  27. out = ShaderVariableDataType::kI32;
  28. break;
  29. case spirv_cross::SPIRType::UInt:
  30. out = ShaderVariableDataType::kU32;
  31. break;
  32. case spirv_cross::SPIRType::Int64:
  33. out = ShaderVariableDataType::kI64;
  34. break;
  35. case spirv_cross::SPIRType::UInt64:
  36. out = ShaderVariableDataType::kU64;
  37. break;
  38. case spirv_cross::SPIRType::Half:
  39. out = ShaderVariableDataType::kF16;
  40. break;
  41. case spirv_cross::SPIRType::Float:
  42. out = ShaderVariableDataType::kF32;
  43. break;
  44. default:
  45. break;
  46. }
  47. return out;
  48. }
  49. /// Populates the reflection info.
  50. class SpirvReflector : public spirv_cross::Compiler
  51. {
  52. public:
  53. SpirvReflector(const U32* ir, PtrSize wordCount, ShaderReflectionVisitorInterface* interface)
  54. : spirv_cross::Compiler(ir, wordCount)
  55. , m_interface(interface)
  56. {
  57. }
  58. [[nodiscard]] static Error performSpirvReflection(Array<ConstWeakArray<U8>, U32(ShaderType::kCount)> spirv,
  59. ShaderReflectionVisitorInterface& interface);
  60. private:
  61. class Var
  62. {
  63. public:
  64. String m_name;
  65. ShaderVariableBlockInfo m_blockInfo;
  66. ShaderVariableDataType m_type = ShaderVariableDataType::kNone;
  67. };
  68. class Block
  69. {
  70. public:
  71. String m_name;
  72. DynamicArray<Var> m_vars;
  73. U32 m_binding = kMaxU32;
  74. U32 m_set = kMaxU32;
  75. U32 m_size = kMaxU32;
  76. };
  77. class Opaque
  78. {
  79. public:
  80. String m_name;
  81. ShaderVariableDataType m_type = ShaderVariableDataType::kNone;
  82. U32 m_binding = kMaxU32;
  83. U32 m_set = kMaxU32;
  84. U32 m_arraySize = kMaxU32;
  85. };
  86. class Const
  87. {
  88. public:
  89. String m_name;
  90. ShaderVariableDataType m_type = ShaderVariableDataType::kNone;
  91. U32 m_constantId = kMaxU32;
  92. };
  93. class StructMember
  94. {
  95. public:
  96. String m_name;
  97. ShaderVariableDataType m_type = ShaderVariableDataType::kNone;
  98. U32 m_structIndex = kMaxU32; ///< The member is actually a struct.
  99. U32 m_offset = kMaxU32;
  100. U32 m_arraySize = kMaxU32;
  101. };
  102. class Struct
  103. {
  104. public:
  105. String m_name;
  106. DynamicArray<StructMember> m_members;
  107. U32 m_size = 0;
  108. U32 m_alignment = 0;
  109. };
  110. ShaderReflectionVisitorInterface* m_interface = nullptr;
  111. Error spirvTypeToAnki(const spirv_cross::SPIRType& type, ShaderVariableDataType& out) const;
  112. Error blockReflection(const spirv_cross::Resource& res, Bool isStorage, DynamicArray<Block>& blocks) const;
  113. Error opaqueReflection(const spirv_cross::Resource& res, DynamicArray<Opaque>& opaques) const;
  114. Error constsReflection(DynamicArray<Const>& consts) const;
  115. Error blockVariablesReflection(spirv_cross::TypeID resourceId, DynamicArray<Var>& vars) const;
  116. Error blockVariableReflection(const spirv_cross::SPIRType& type, CString parentVariable, U32 baseOffset, DynamicArray<Var>& vars) const;
  117. Error workgroupSizes(U32& sizex, U32& sizey, U32& sizez, U32& specConstMask);
  118. Error structsReflection(DynamicArray<Struct>& structs) const;
  119. Error structReflection(uint32_t id, const spirv_cross::SPIRType& type, U32 depth, Bool& skipped, DynamicArray<Struct>& structs,
  120. U32& structIndexInStructsArr) const;
  121. };
  122. Error SpirvReflector::structsReflection(DynamicArray<Struct>& structs) const
  123. {
  124. Error err = Error::kNone;
  125. ir.for_each_typed_id<spirv_cross::SPIRType>([&err, &structs, this](uint32_t id, const spirv_cross::SPIRType& type) {
  126. if(err)
  127. {
  128. return;
  129. }
  130. if(type.basetype != spirv_cross::SPIRType::Struct || type.pointer || !type.array.empty() || has_decoration(type.self, spv::DecorationBlock))
  131. {
  132. return;
  133. }
  134. U32 idx;
  135. Bool skipped;
  136. err = structReflection(id, type, 0, skipped, structs, idx);
  137. });
  138. return err;
  139. }
  140. Error SpirvReflector::structReflection(uint32_t id, const spirv_cross::SPIRType& type, U32 depth, Bool& skipped, DynamicArray<Struct>& structs,
  141. U32& structIndexInStructsArr) const
  142. {
  143. skipped = false;
  144. // Name
  145. std::string name = to_name(id);
  146. // Skip GL builtins, SPIRV-Cross things and symbols that should be skipped
  147. if(CString(name.c_str()).find("gl_") == 0 || CString(name.c_str()).find("_") == 0 || (depth == 0 && m_interface->skipSymbol(name.c_str())))
  148. {
  149. skipped = true;
  150. return Error::kNone;
  151. }
  152. // Check if the struct is already there
  153. structIndexInStructsArr = 0;
  154. for(const Struct& s : structs)
  155. {
  156. if(s.m_name == name.c_str())
  157. {
  158. return Error::kNone;
  159. }
  160. ++structIndexInStructsArr;
  161. }
  162. // Create new struct
  163. Struct cstruct;
  164. cstruct.m_name = name.c_str();
  165. U32 membersOffset = 0;
  166. Bool aMemberWasSkipped = false;
  167. // Members
  168. for(U32 i = 0; i < type.member_types.size(); ++i)
  169. {
  170. StructMember& member = *cstruct.m_members.emplaceBack();
  171. const spirv_cross::SPIRType& memberType = get<spirv_cross::SPIRType>(type.member_types[i]);
  172. // Get name
  173. const spirv_cross::Meta* meta = ir.find_meta(type.self);
  174. ANKI_ASSERT(meta);
  175. ANKI_ASSERT(i < meta->members.size());
  176. ANKI_ASSERT(!meta->members[i].alias.empty());
  177. member.m_name = meta->members[i].alias.c_str();
  178. // Array size
  179. if(!memberType.array.empty())
  180. {
  181. if(memberType.array.size() > 1)
  182. {
  183. ANKI_SHADER_COMPILER_LOGE("Can't support multi-dimentional arrays at the moment");
  184. return Error::kUserData;
  185. }
  186. const Bool notSpecConstantArraySize = memberType.array_size_literal[0];
  187. if(notSpecConstantArraySize)
  188. {
  189. // Have a min to acount for unsized arrays of SSBOs
  190. member.m_arraySize = max(memberType.array[0], 1u);
  191. }
  192. else
  193. {
  194. ANKI_SHADER_COMPILER_LOGE("Arrays with spec constant size are not allowed: %s", member.m_name.cstr());
  195. return Error::kFunctionFailed;
  196. }
  197. }
  198. else
  199. {
  200. member.m_arraySize = 1;
  201. }
  202. // Type
  203. const ShaderVariableDataType baseType = spirvcrossBaseTypeToAnki(memberType.basetype);
  204. const Bool isNumeric = baseType != ShaderVariableDataType::kNone;
  205. ShaderVariableDataType actualType = ShaderVariableDataType::kNone;
  206. U32 memberSize = 0;
  207. U32 memberAlignment = 0;
  208. if(isNumeric)
  209. {
  210. const Bool isMatrix = memberType.columns > 1;
  211. if(0)
  212. {
  213. }
  214. #define ANKI_SVDT_MACRO(type, baseType_, rowCount, columnCount, isIntagralType) \
  215. else if(ShaderVariableDataType::k##baseType_ == baseType && isMatrix && memberType.vecsize == rowCount && memberType.columns == columnCount) \
  216. { \
  217. actualType = ShaderVariableDataType::k##type; \
  218. memberSize = sizeof(type); \
  219. memberAlignment = alignof(baseType_); \
  220. } \
  221. else if(ShaderVariableDataType::k##baseType_ == baseType && !isMatrix && memberType.vecsize == rowCount) \
  222. { \
  223. actualType = ShaderVariableDataType::k##type; \
  224. memberSize = sizeof(type); \
  225. memberAlignment = alignof(baseType_); \
  226. }
  227. #include <AnKi/Gr/ShaderVariableDataType.defs.h>
  228. #undef ANKI_SVDT_MACRO
  229. member.m_type = actualType;
  230. }
  231. else if(memberType.basetype == spirv_cross::SPIRType::Struct)
  232. {
  233. U32 idx = kMaxU32;
  234. Bool memberSkipped = false;
  235. ANKI_CHECK(structReflection(type.member_types[i], memberType, depth + 1, memberSkipped, structs, idx));
  236. if(memberSkipped)
  237. {
  238. aMemberWasSkipped = true;
  239. break;
  240. }
  241. else
  242. {
  243. ANKI_ASSERT(idx < structs.getSize());
  244. member.m_structIndex = idx;
  245. memberSize = structs[idx].m_size;
  246. memberAlignment = structs[idx].m_alignment;
  247. }
  248. }
  249. else
  250. {
  251. ANKI_SHADER_COMPILER_LOGE("Unhandled base type for member: %s", name.c_str());
  252. return Error::kFunctionFailed;
  253. }
  254. // Update offsets and alignments
  255. memberSize *= member.m_arraySize;
  256. member.m_offset = getAlignedRoundUp(memberAlignment, membersOffset);
  257. cstruct.m_alignment = max(cstruct.m_alignment, memberAlignment);
  258. cstruct.m_size = member.m_offset + memberSize;
  259. membersOffset = member.m_offset + memberSize;
  260. }
  261. if(!aMemberWasSkipped)
  262. {
  263. // Now you can create the struct
  264. alignRoundUp(cstruct.m_alignment, cstruct.m_size);
  265. Struct& newStruct = *structs.emplaceBack();
  266. newStruct = std::move(cstruct);
  267. }
  268. else
  269. {
  270. skipped = true;
  271. }
  272. return Error::kNone;
  273. }
  274. Error SpirvReflector::blockVariablesReflection(spirv_cross::TypeID resourceId, DynamicArray<Var>& vars) const
  275. {
  276. Bool found = false;
  277. Error err = Error::kNone;
  278. ir.for_each_typed_id<spirv_cross::SPIRType>([&](uint32_t, const spirv_cross::SPIRType& type) {
  279. if(err)
  280. {
  281. return;
  282. }
  283. if(type.basetype == spirv_cross::SPIRType::Struct && !type.pointer && type.array.empty())
  284. {
  285. if(type.self == resourceId)
  286. {
  287. found = true;
  288. err = blockVariableReflection(type, CString(), 0, vars);
  289. }
  290. }
  291. });
  292. ANKI_CHECK(err);
  293. if(!found)
  294. {
  295. ANKI_SHADER_COMPILER_LOGE("Can't determine the type of a block");
  296. return Error::kUserData;
  297. }
  298. return Error::kNone;
  299. }
  300. Error SpirvReflector::blockVariableReflection(const spirv_cross::SPIRType& type, CString parentVariable, U32 baseOffset,
  301. DynamicArray<Var>& vars) const
  302. {
  303. ANKI_ASSERT(type.basetype == spirv_cross::SPIRType::Struct);
  304. for(U32 i = 0; i < type.member_types.size(); ++i)
  305. {
  306. Var var;
  307. const spirv_cross::SPIRType& memberType = get<spirv_cross::SPIRType>(type.member_types[i]);
  308. // Name
  309. {
  310. const spirv_cross::Meta* meta = ir.find_meta(type.self);
  311. ANKI_ASSERT(meta);
  312. ANKI_ASSERT(i < meta->members.size());
  313. ANKI_ASSERT(!meta->members[i].alias.empty());
  314. const std::string& name = meta->members[i].alias;
  315. if(parentVariable.isEmpty())
  316. {
  317. var.m_name = name.c_str();
  318. }
  319. else
  320. {
  321. var.m_name.sprintf("%s.%s", parentVariable.cstr(), name.c_str());
  322. }
  323. }
  324. // Offset
  325. {
  326. auto it = ir.meta.find(type.self);
  327. ANKI_ASSERT(it != ir.meta.end());
  328. const spirv_cross::Vector<spirv_cross::Meta::Decoration>& memb = it->second.members;
  329. ANKI_ASSERT(i < memb.size());
  330. const spirv_cross::Meta::Decoration& dec = memb[i];
  331. ANKI_ASSERT(dec.decoration_flags.get(spv::DecorationOffset));
  332. var.m_blockInfo.m_offset = I16(dec.offset + baseOffset);
  333. }
  334. // Array size
  335. Bool isArray = false;
  336. {
  337. if(!memberType.array.empty())
  338. {
  339. if(memberType.array.size() > 1)
  340. {
  341. ANKI_SHADER_COMPILER_LOGE("Can't support multi-dimentional arrays at the moment");
  342. return Error::kUserData;
  343. }
  344. const Bool notSpecConstantArraySize = memberType.array_size_literal[0];
  345. if(notSpecConstantArraySize)
  346. {
  347. // Have a min to acount for unsized arrays of SSBOs
  348. var.m_blockInfo.m_arraySize = max<I16>(I16(memberType.array[0]), 1);
  349. isArray = true;
  350. }
  351. else
  352. {
  353. var.m_blockInfo.m_arraySize = 1;
  354. isArray = true;
  355. }
  356. }
  357. else
  358. {
  359. var.m_blockInfo.m_arraySize = 1;
  360. }
  361. }
  362. // Array stride
  363. if(has_decoration(type.member_types[i], spv::DecorationArrayStride))
  364. {
  365. var.m_blockInfo.m_arrayStride = I16(get_decoration(type.member_types[i], spv::DecorationArrayStride));
  366. }
  367. const ShaderVariableDataType baseType = spirvcrossBaseTypeToAnki(memberType.basetype);
  368. const Bool isNumeric = baseType != ShaderVariableDataType::kNone;
  369. if(memberType.basetype == spirv_cross::SPIRType::Struct)
  370. {
  371. if(var.m_blockInfo.m_arraySize == 1 && !isArray)
  372. {
  373. ANKI_CHECK(blockVariableReflection(memberType, var.m_name, var.m_blockInfo.m_offset, vars));
  374. }
  375. else
  376. {
  377. for(U32 i = 0; i < U32(var.m_blockInfo.m_arraySize); ++i)
  378. {
  379. String newName;
  380. newName.sprintf("%s[%u]", var.m_name.getBegin(), i);
  381. ANKI_CHECK(blockVariableReflection(memberType, newName, var.m_blockInfo.m_offset + var.m_blockInfo.m_arrayStride * i, vars));
  382. }
  383. }
  384. }
  385. else if(isNumeric)
  386. {
  387. const Bool isMatrix = memberType.columns > 1;
  388. if(0)
  389. {
  390. }
  391. #define ANKI_SVDT_MACRO(type_, baseType_, rowCount, columnCount, isIntagralType) \
  392. else if(ShaderVariableDataType::k##baseType_ == baseType && isMatrix && memberType.vecsize == rowCount && memberType.columns == columnCount) \
  393. { \
  394. var.m_type = ShaderVariableDataType::k##type_; \
  395. auto it = ir.meta.find(type.self); \
  396. ANKI_ASSERT(it != ir.meta.end()); \
  397. const spirv_cross::Vector<spirv_cross::Meta::Decoration>& memberDecorations = it->second.members; \
  398. ANKI_ASSERT(i < memberDecorations.size()); \
  399. var.m_blockInfo.m_matrixStride = I16(memberDecorations[i].matrix_stride); \
  400. } \
  401. else if(ShaderVariableDataType::k##baseType_ == baseType && !isMatrix && memberType.vecsize == rowCount) \
  402. { \
  403. var.m_type = ShaderVariableDataType::k##type_; \
  404. }
  405. #include <AnKi/Gr/ShaderVariableDataType.defs.h>
  406. #undef ANKI_SVDT_MACRO
  407. if(var.m_type == ShaderVariableDataType::kNone)
  408. {
  409. ANKI_SHADER_COMPILER_LOGE("Unhandled numeric member: %s", var.m_name.cstr());
  410. return Error::kFunctionFailed;
  411. }
  412. }
  413. else
  414. {
  415. ANKI_SHADER_COMPILER_LOGE("Unhandled base type for member: %s", var.m_name.cstr());
  416. return Error::kFunctionFailed;
  417. }
  418. // Store the member if it's no struct
  419. if(var.m_type != ShaderVariableDataType::kNone)
  420. {
  421. vars.emplaceBack(std::move(var));
  422. }
  423. }
  424. return Error::kNone;
  425. }
  426. Error SpirvReflector::blockReflection(const spirv_cross::Resource& res, [[maybe_unused]] Bool isStorage, DynamicArray<Block>& blocks) const
  427. {
  428. Block newBlock;
  429. const spirv_cross::SPIRType type = get_type(res.type_id);
  430. const spirv_cross::Bitset decorationMask = get_decoration_bitset(res.id);
  431. const Bool isPushConstant = get_storage_class(res.id) == spv::StorageClassPushConstant;
  432. // Name
  433. {
  434. const std::string name = (!res.name.empty()) ? res.name : to_name(res.base_type_id);
  435. if(name.length() == 0)
  436. {
  437. ANKI_SHADER_COMPILER_LOGE("Can't accept zero name length");
  438. return Error::kUserData;
  439. }
  440. if(m_interface->skipSymbol(name.c_str()))
  441. {
  442. return Error::kNone;
  443. }
  444. newBlock.m_name = name.c_str();
  445. }
  446. // Set
  447. if(!isPushConstant)
  448. {
  449. newBlock.m_set = get_decoration(res.id, spv::DecorationDescriptorSet);
  450. if(newBlock.m_set >= kMaxDescriptorSets)
  451. {
  452. ANKI_SHADER_COMPILER_LOGE("Too high descriptor set: %u", newBlock.m_set);
  453. return Error::kUserData;
  454. }
  455. }
  456. // Binding
  457. if(!isPushConstant)
  458. {
  459. newBlock.m_binding = get_decoration(res.id, spv::DecorationBinding);
  460. }
  461. // Size
  462. newBlock.m_size = U32(get_declared_struct_size(get_type(res.base_type_id)));
  463. ANKI_ASSERT(isStorage || newBlock.m_size > 0);
  464. // Add it
  465. const Block* otherFound = nullptr;
  466. for(const Block& other : blocks)
  467. {
  468. const Bool bindingSame = other.m_set == newBlock.m_set && other.m_binding == newBlock.m_binding;
  469. const Bool nameSame = strcmp(other.m_name.getBegin(), newBlock.m_name.getBegin()) == 0;
  470. const Bool sizeSame = other.m_size == newBlock.m_size;
  471. const Bool err0 = bindingSame && (!nameSame || !sizeSame);
  472. const Bool err1 = nameSame && (!bindingSame || !sizeSame);
  473. if(err0 || err1)
  474. {
  475. ANKI_SHADER_COMPILER_LOGE("Linking error. Blocks %s and %s", other.m_name.cstr(), newBlock.m_name.cstr());
  476. return Error::kUserData;
  477. }
  478. if(bindingSame)
  479. {
  480. otherFound = &other;
  481. break;
  482. }
  483. }
  484. if(!otherFound)
  485. {
  486. // Get the variables
  487. ANKI_CHECK(blockVariablesReflection(res.base_type_id, newBlock.m_vars));
  488. // Store the block
  489. blocks.emplaceBack(std::move(newBlock));
  490. }
  491. #if ANKI_ASSERTIONS_ENABLED
  492. else
  493. {
  494. DynamicArray<Var> vars;
  495. ANKI_CHECK(blockVariablesReflection(res.base_type_id, vars));
  496. ANKI_ASSERT(vars.getSize() == otherFound->m_vars.getSize() && "Expecting same vars");
  497. }
  498. #endif
  499. return Error::kNone;
  500. }
  501. Error SpirvReflector::spirvTypeToAnki(const spirv_cross::SPIRType& type, ShaderVariableDataType& out) const
  502. {
  503. switch(type.basetype)
  504. {
  505. case spirv_cross::SPIRType::Image:
  506. case spirv_cross::SPIRType::SampledImage:
  507. {
  508. switch(type.image.dim)
  509. {
  510. case spv::Dim1D:
  511. out = (type.image.arrayed) ? ShaderVariableDataType::kTexture1DArray : ShaderVariableDataType::kTexture1D;
  512. break;
  513. case spv::Dim2D:
  514. out = (type.image.arrayed) ? ShaderVariableDataType::kTexture2DArray : ShaderVariableDataType::kTexture2D;
  515. break;
  516. case spv::Dim3D:
  517. out = ShaderVariableDataType::kTexture3D;
  518. break;
  519. case spv::DimCube:
  520. out = (type.image.arrayed) ? ShaderVariableDataType::kTextureCubeArray : ShaderVariableDataType::kTextureCube;
  521. break;
  522. default:
  523. ANKI_ASSERT(0);
  524. }
  525. break;
  526. }
  527. case spirv_cross::SPIRType::Sampler:
  528. out = ShaderVariableDataType::kSampler;
  529. break;
  530. default:
  531. ANKI_SHADER_COMPILER_LOGE("Can't determine the type");
  532. return Error::kUserData;
  533. }
  534. return Error::kNone;
  535. }
  536. Error SpirvReflector::opaqueReflection(const spirv_cross::Resource& res, DynamicArray<Opaque>& opaques) const
  537. {
  538. Opaque newOpaque;
  539. const spirv_cross::SPIRType type = get_type(res.type_id);
  540. const spirv_cross::Bitset decorationMask = get_decoration_bitset(res.id);
  541. const spirv_cross::ID fallbackId = spirv_cross::ID(res.id);
  542. // Name
  543. const std::string name = (!res.name.empty()) ? res.name : get_fallback_name(fallbackId);
  544. if(name.length() == 0)
  545. {
  546. ANKI_SHADER_COMPILER_LOGE("Can't accept zero length name");
  547. return Error::kUserData;
  548. }
  549. if(m_interface->skipSymbol(name.c_str()))
  550. {
  551. return Error::kNone;
  552. }
  553. newOpaque.m_name = name.c_str();
  554. // Type
  555. ANKI_CHECK(spirvTypeToAnki(type, newOpaque.m_type));
  556. // Set
  557. newOpaque.m_set = get_decoration(res.id, spv::DecorationDescriptorSet);
  558. if(newOpaque.m_set >= kMaxDescriptorSets)
  559. {
  560. ANKI_SHADER_COMPILER_LOGE("Too high descriptor set: %u", newOpaque.m_set);
  561. return Error::kUserData;
  562. }
  563. // Binding
  564. newOpaque.m_binding = get_decoration(res.id, spv::DecorationBinding);
  565. // Size
  566. if(type.array.size() == 0)
  567. {
  568. newOpaque.m_arraySize = 1;
  569. }
  570. else if(type.array.size() == 1)
  571. {
  572. newOpaque.m_arraySize = type.array[0];
  573. }
  574. else
  575. {
  576. ANKI_SHADER_COMPILER_LOGE("Can't support multi-dimensional arrays: %s", newOpaque.m_name.cstr());
  577. return Error::kUserData;
  578. }
  579. // Add it
  580. Bool found = false;
  581. for(const Opaque& other : opaques)
  582. {
  583. const Bool bindingSame = other.m_set == newOpaque.m_set && other.m_binding == newOpaque.m_binding;
  584. const Bool nameSame = other.m_name == newOpaque.m_name;
  585. const Bool sizeSame = other.m_arraySize == newOpaque.m_arraySize;
  586. const Bool typeSame = other.m_type == newOpaque.m_type;
  587. const Bool err = nameSame && (!bindingSame || !sizeSame || !typeSame);
  588. if(err)
  589. {
  590. ANKI_SHADER_COMPILER_LOGE("Linking error");
  591. return Error::kUserData;
  592. }
  593. if(nameSame)
  594. {
  595. found = true;
  596. break;
  597. }
  598. }
  599. if(!found)
  600. {
  601. opaques.emplaceBack(std::move(newOpaque));
  602. }
  603. return Error::kNone;
  604. }
  605. Error SpirvReflector::constsReflection(DynamicArray<Const>& consts) const
  606. {
  607. spirv_cross::SmallVector<spirv_cross::SpecializationConstant> specConsts = get_specialization_constants();
  608. for(const spirv_cross::SpecializationConstant& c : specConsts)
  609. {
  610. Const newConst;
  611. const spirv_cross::SPIRConstant cc = get<spirv_cross::SPIRConstant>(c.id);
  612. const spirv_cross::SPIRType type = get<spirv_cross::SPIRType>(cc.constant_type);
  613. const std::string name = get_name(c.id);
  614. if(name.length() == 0)
  615. {
  616. ANKI_SHADER_COMPILER_LOGE("Can't accept zero legth name");
  617. return Error::kUserData;
  618. }
  619. newConst.m_name = name.c_str();
  620. newConst.m_constantId = c.constant_id;
  621. switch(type.basetype)
  622. {
  623. case spirv_cross::SPIRType::UInt:
  624. newConst.m_type = ShaderVariableDataType::kU32;
  625. break;
  626. case spirv_cross::SPIRType::Int:
  627. newConst.m_type = ShaderVariableDataType::kI32;
  628. break;
  629. case spirv_cross::SPIRType::Float:
  630. newConst.m_type = ShaderVariableDataType::kF32;
  631. break;
  632. default:
  633. ANKI_SHADER_COMPILER_LOGE("Can't determine the type of the spec constant: %s", name.c_str());
  634. return Error::kUserData;
  635. }
  636. // Search for it
  637. Const* foundConst = nullptr;
  638. for(Const& other : consts)
  639. {
  640. const Bool nameSame = other.m_name == newConst.m_name;
  641. const Bool typeSame = other.m_type == newConst.m_type;
  642. const Bool idSame = other.m_constantId == newConst.m_constantId;
  643. const Bool err0 = nameSame && (!typeSame || !idSame);
  644. const Bool err1 = idSame && (!nameSame || !typeSame);
  645. if(err0 || err1)
  646. {
  647. ANKI_SHADER_COMPILER_LOGE("Linking error: %s", newConst.m_name.cstr());
  648. return Error::kUserData;
  649. }
  650. if(idSame)
  651. {
  652. foundConst = &other;
  653. break;
  654. }
  655. }
  656. // Add it or update it
  657. if(foundConst == nullptr)
  658. {
  659. consts.emplaceBack(std::move(newConst));
  660. }
  661. }
  662. return Error::kNone;
  663. }
  664. Error SpirvReflector::workgroupSizes(U32& sizex, U32& sizey, U32& sizez, U32& specConstMask)
  665. {
  666. sizex = sizey = sizez = specConstMask = 0;
  667. auto entries = get_entry_points_and_stages();
  668. for(const auto& e : entries)
  669. {
  670. if(e.execution_model == spv::ExecutionModelGLCompute)
  671. {
  672. const auto& spvEntry = get_entry_point(e.name, e.execution_model);
  673. spirv_cross::SpecializationConstant specx, specy, specz;
  674. get_work_group_size_specialization_constants(specx, specy, specz);
  675. if(specx.id != spirv_cross::ID(0))
  676. {
  677. specConstMask |= 1;
  678. sizex = specx.constant_id;
  679. }
  680. else
  681. {
  682. sizex = spvEntry.workgroup_size.x;
  683. }
  684. if(specy.id != spirv_cross::ID(0))
  685. {
  686. specConstMask |= 2;
  687. sizey = specy.constant_id;
  688. }
  689. else
  690. {
  691. sizey = spvEntry.workgroup_size.y;
  692. }
  693. if(specz.id != spirv_cross::ID(0))
  694. {
  695. specConstMask |= 4;
  696. sizez = specz.constant_id;
  697. }
  698. else
  699. {
  700. sizez = spvEntry.workgroup_size.z;
  701. }
  702. }
  703. }
  704. return Error::kNone;
  705. }
  706. Error SpirvReflector::performSpirvReflection(Array<ConstWeakArray<U8>, U32(ShaderType::kCount)> spirv, ShaderReflectionVisitorInterface& interface)
  707. {
  708. DynamicArray<Block> uniformBlocks;
  709. DynamicArray<Block> storageBlocks;
  710. DynamicArray<Block> pushConstantBlock;
  711. DynamicArray<Opaque> opaques;
  712. DynamicArray<Const> specializationConstants;
  713. Array<U32, 3> workgroupSizes = {};
  714. U32 workgroupSizeSpecConstMask = 0;
  715. DynamicArray<Struct> structs;
  716. // Perform reflection for each stage
  717. for(const ShaderType type : EnumIterable<ShaderType>())
  718. {
  719. if(spirv[type].getSize() == 0)
  720. {
  721. continue;
  722. }
  723. // Parse SPIR-V
  724. const unsigned int* spvb = reinterpret_cast<const unsigned int*>(spirv[type].getBegin());
  725. SpirvReflector compiler(spvb, spirv[type].getSizeInBytes() / sizeof(unsigned int), &interface);
  726. // Uniform blocks
  727. for(const spirv_cross::Resource& res : compiler.get_shader_resources().uniform_buffers)
  728. {
  729. ANKI_CHECK(compiler.blockReflection(res, false, uniformBlocks));
  730. }
  731. // Sorage blocks
  732. for(const spirv_cross::Resource& res : compiler.get_shader_resources().storage_buffers)
  733. {
  734. ANKI_CHECK(compiler.blockReflection(res, true, storageBlocks));
  735. }
  736. // Push constants
  737. if(compiler.get_shader_resources().push_constant_buffers.size() == 1)
  738. {
  739. ANKI_CHECK(compiler.blockReflection(compiler.get_shader_resources().push_constant_buffers[0], false, pushConstantBlock));
  740. }
  741. else if(compiler.get_shader_resources().push_constant_buffers.size() > 1)
  742. {
  743. ANKI_SHADER_COMPILER_LOGE("Expecting only a single push constants block");
  744. return Error::kUserData;
  745. }
  746. // Opaque
  747. for(const spirv_cross::Resource& res : compiler.get_shader_resources().separate_images)
  748. {
  749. ANKI_CHECK(compiler.opaqueReflection(res, opaques));
  750. }
  751. for(const spirv_cross::Resource& res : compiler.get_shader_resources().storage_images)
  752. {
  753. ANKI_CHECK(compiler.opaqueReflection(res, opaques));
  754. }
  755. for(const spirv_cross::Resource& res : compiler.get_shader_resources().separate_samplers)
  756. {
  757. ANKI_CHECK(compiler.opaqueReflection(res, opaques));
  758. }
  759. // Spec consts
  760. ANKI_CHECK(compiler.constsReflection(specializationConstants));
  761. // Workgroup sizes
  762. if(type == ShaderType::kCompute)
  763. {
  764. ANKI_CHECK(compiler.workgroupSizes(workgroupSizes[0], workgroupSizes[1], workgroupSizes[2], workgroupSizeSpecConstMask));
  765. }
  766. // Structs
  767. ANKI_CHECK(compiler.structsReflection(structs));
  768. }
  769. // Inform through the interface
  770. ANKI_CHECK(interface.setCounts(uniformBlocks.getSize(), storageBlocks.getSize(), opaques.getSize(), pushConstantBlock.getSize() == 1,
  771. specializationConstants.getSize(), structs.getSize()));
  772. for(U32 i = 0; i < uniformBlocks.getSize(); ++i)
  773. {
  774. const Block& block = uniformBlocks[i];
  775. ANKI_CHECK(interface.visitUniformBlock(i, block.m_name, block.m_set, block.m_binding, block.m_size, block.m_vars.getSize()));
  776. for(U32 j = 0; j < block.m_vars.getSize(); ++j)
  777. {
  778. const Var& var = block.m_vars[j];
  779. ANKI_CHECK(interface.visitUniformVariable(i, j, var.m_name, var.m_type, var.m_blockInfo));
  780. }
  781. }
  782. for(U32 i = 0; i < storageBlocks.getSize(); ++i)
  783. {
  784. const Block& block = storageBlocks[i];
  785. ANKI_CHECK(interface.visitStorageBlock(i, block.m_name, block.m_set, block.m_binding, block.m_size, block.m_vars.getSize()));
  786. for(U32 j = 0; j < block.m_vars.getSize(); ++j)
  787. {
  788. const Var& var = block.m_vars[j];
  789. ANKI_CHECK(interface.visitStorageVariable(i, j, var.m_name, var.m_type, var.m_blockInfo));
  790. }
  791. }
  792. if(pushConstantBlock.getSize() == 1)
  793. {
  794. ANKI_CHECK(
  795. interface.visitPushConstantsBlock(pushConstantBlock[0].m_name, pushConstantBlock[0].m_size, pushConstantBlock[0].m_vars.getSize()));
  796. for(U32 j = 0; j < pushConstantBlock[0].m_vars.getSize(); ++j)
  797. {
  798. const Var& var = pushConstantBlock[0].m_vars[j];
  799. ANKI_CHECK(interface.visitPushConstant(j, var.m_name, var.m_type, var.m_blockInfo));
  800. }
  801. }
  802. for(U32 i = 0; i < opaques.getSize(); ++i)
  803. {
  804. const Opaque& o = opaques[i];
  805. ANKI_CHECK(interface.visitOpaque(i, o.m_name, o.m_type, o.m_set, o.m_binding, o.m_arraySize));
  806. }
  807. for(U32 i = 0; i < specializationConstants.getSize(); ++i)
  808. {
  809. const Const& c = specializationConstants[i];
  810. ANKI_CHECK(interface.visitConstant(i, c.m_name, c.m_type, c.m_constantId));
  811. }
  812. if(spirv[ShaderType::kCompute].getSize())
  813. {
  814. ANKI_CHECK(interface.setWorkgroupSizes(workgroupSizes[0], workgroupSizes[1], workgroupSizes[2], workgroupSizeSpecConstMask));
  815. }
  816. for(U32 i = 0; i < structs.getSize(); ++i)
  817. {
  818. const Struct& s = structs[i];
  819. ANKI_CHECK(interface.visitStruct(i, s.m_name, s.m_members.getSize(), s.m_size));
  820. for(U32 j = 0; j < s.m_members.getSize(); ++j)
  821. {
  822. const StructMember& sm = s.m_members[j];
  823. ANKI_CHECK(interface.visitStructMember(i, s.m_name, j, sm.m_name, sm.m_type,
  824. (sm.m_structIndex != kMaxU32) ? structs[sm.m_structIndex].m_name.toCString() : CString(),
  825. sm.m_offset, sm.m_arraySize));
  826. }
  827. }
  828. return Error::kNone;
  829. }
  830. Error performSpirvReflection(Array<ConstWeakArray<U8>, U32(ShaderType::kCount)> spirv, ShaderReflectionVisitorInterface& interface)
  831. {
  832. return SpirvReflector::performSpirvReflection(spirv, interface);
  833. }
  834. } // end namespace anki