@@ -1439,15 +1439,15 @@ static bool generateKernelClockInst(const SPIRV::IncomingCall *Call,
1439
1439
return true ;
1440
1440
}
1441
1441
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.
1448
1448
//
1449
1449
// For a constant index >= 3 we generate:
1450
- // %res = OpConstant %SizeT 0
1450
+ // %res = OpConstant %SizeT DefaultValue
1451
1451
//
1452
1452
// For other indices we generate:
1453
1453
// %g = OpVariable %ptr_V3_SizeT Input
@@ -1461,7 +1461,7 @@ static bool generateKernelClockInst(const SPIRV::IncomingCall *Call,
1461
1461
// If the index is dynamic, we generate:
1462
1462
// %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
1463
1463
// %cmp = OpULessThan %bool %idx %const_3
1464
- // %res = OpSelect %SizeT %cmp %tmp %const_0
1464
+ // %res = OpSelect %SizeT %cmp %tmp %const_<DefaultValue>
1465
1465
//
1466
1466
// If the bitwidth of %res does not match the expected return type, we add an
1467
1467
// extend or truncate.
@@ -1840,11 +1840,11 @@ static bool generateGetQueryInst(const SPIRV::IncomingCall *Call,
1840
1840
// Lookup the builtin record.
1841
1841
SPIRV::BuiltIn::BuiltIn Value =
1842
1842
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 );
1848
1848
}
1849
1849
1850
1850
static bool generateImageSizeQueryInst (const SPIRV::IncomingCall *Call,
0 commit comments