@@ -1439,15 +1439,15 @@ static bool generateKernelClockInst(const SPIRV::IncomingCall *Call,
14391439 return true ;
14401440}
14411441
1442- // These queries ask for a single size_t result for a given dimension index, e.g
1443- // size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
1444- // these values are all vec3 types, so we need to extract the correct index or
1445- // return defaultVal (0 or 1 depending on the query). We also handle extending
1446- // or tuncating in case size_t does not match the expected result type's
1447- // bitwidth.
1442+ // These queries ask for a single size_t result for a given dimension index,
1443+ // e.g. size_t get_global_id(uint dimindex). In SPIR-V, the builtins
1444+ // corresponding to these values are all vec3 types, so we need to extract the
1445+ // correct index or return DefaultValue (0 or 1 depending on the query). We also
1446+ // handle extending or truncating in case size_t does not match the expected
1447+ // result type's bitwidth.
14481448//
14491449// For a constant index >= 3 we generate:
1450- // %res = OpConstant %SizeT 0
1450+ // %res = OpConstant %SizeT DefaultValue
14511451//
14521452// For other indices we generate:
14531453// %g = OpVariable %ptr_V3_SizeT Input
@@ -1461,7 +1461,7 @@ static bool generateKernelClockInst(const SPIRV::IncomingCall *Call,
14611461// If the index is dynamic, we generate:
14621462// %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
14631463// %cmp = OpULessThan %bool %idx %const_3
1464- // %res = OpSelect %SizeT %cmp %tmp %const_0
1464+ // %res = OpSelect %SizeT %cmp %tmp %const_<DefaultValue>
14651465//
14661466// If the bitwidth of %res does not match the expected return type, we add an
14671467// extend or truncate.
@@ -1840,11 +1840,11 @@ static bool generateGetQueryInst(const SPIRV::IncomingCall *Call,
18401840 // Lookup the builtin record.
18411841 SPIRV::BuiltIn::BuiltIn Value =
18421842 SPIRV::lookupGetBuiltin (Call->Builtin ->Name , Call->Builtin ->Set )->Value ;
1843- uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||
1844- Value == SPIRV::BuiltIn::NumWorkgroups ||
1845- Value == SPIRV::BuiltIn::WorkgroupSize ||
1846- Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
1847- return genWorkgroupQuery (Call, MIRBuilder, GR, Value, IsDefault ? 1 : 0 );
1843+ const bool IsDefaultOne = (Value == SPIRV::BuiltIn::GlobalSize ||
1844+ Value == SPIRV::BuiltIn::NumWorkgroups ||
1845+ Value == SPIRV::BuiltIn::WorkgroupSize ||
1846+ Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
1847+ return genWorkgroupQuery (Call, MIRBuilder, GR, Value, IsDefaultOne ? 1 : 0 );
18481848}
18491849
18501850static bool generateImageSizeQueryInst (const SPIRV::IncomingCall *Call,
0 commit comments