diff --git a/README.md b/README.md index 20fdda7..c1c963d 100644 --- a/README.md +++ b/README.md @@ -41,9 +41,9 @@ Minimal is `0.15.1`. But you know try your version and believe. ## Bgfx version -- [BX](https://github.com/bkaradzic/bx//compare/5dc415ee2e9935089b21186518436681c2d03b47...master) +- [BX](https://github.com/bkaradzic/bx//compare/ce31b1445475ecd4b090471144c4c30a1cbdd871...master) - [BImg](https://github.com/bkaradzic/bimg/compare/bf10ffbb3df1f9f12ad7a9105e5e96e11a9c5a0c...master) -- [BGFX](https://github.com/bkaradzic/bgfx/compare/8a60697cfdfe6181b87ea0c49dff58e43448f712...master) +- [BGFX](https://github.com/bkaradzic/bgfx/compare/56eb016280731451c3b7f18433dc114df035d52a...master) ## Getting started diff --git a/build.zig.zon b/build.zig.zon index 9089aee..4f2fcb9 100644 --- a/build.zig.zon +++ b/build.zig.zon @@ -1,7 +1,7 @@ .{ .name = .zbgfx, .fingerprint = 0xc48ed871c4086e4a, - .version = "0.6.0", + .version = "0.7.0", .minimum_zig_version = "0.15.2", .paths = .{ "includes", diff --git a/libs/bgfx/3rdparty/glslang/SPIRV/GlslangToSpv.cpp b/libs/bgfx/3rdparty/glslang/SPIRV/GlslangToSpv.cpp index c42c139..1a5c40c 100644 --- a/libs/bgfx/3rdparty/glslang/SPIRV/GlslangToSpv.cpp +++ b/libs/bgfx/3rdparty/glslang/SPIRV/GlslangToSpv.cpp @@ -5167,6 +5167,16 @@ spv::Id TGlslangToSpvTraverser::createSpvVariable(const glslang::TIntermSymbol* } spv::Id var = builder.createVariable(spv::NoPrecision, storageClass, spvType, name, initializer, false); + + if (options.emitNonSemanticShaderDebugInfo && storageClass != spv::StorageClass::Function) { + // Create variable alias for retargeted symbols if any. + // Notably, this is only applicable to built-in variables so that it is okay to only use name as the key. + auto [itBegin, itEnd] = glslangIntermediate->getBuiltinAliasLookup().equal_range(name); + for (auto it = itBegin; it != itEnd; ++it) { + builder.createDebugGlobalVariable(builder.getDebugType(spvType), it->second.c_str(), var); + } + } + std::vector topLevelDecorations; glslang::TQualifier typeQualifier = node->getType().getQualifier(); TranslateMemoryDecoration(typeQualifier, topLevelDecorations, glslangIntermediate->usingVulkanMemoryModel()); @@ -5675,7 +5685,8 @@ spv::Id TGlslangToSpvTraverser::convertGlslangToSpvType(const glslang::TType& ty if (type.isSizedArray()) spvType = builder.makeArrayType(spvType, makeArraySizeId(*type.getArraySizes(), 0), stride); else { - if (!lastBufferBlockMember) { + // If we see an runtime array in a buffer_reference, it is not a descriptor + if (!lastBufferBlockMember && type.getBasicType() != glslang::EbtReference) { builder.addIncorporatedExtension("SPV_EXT_descriptor_indexing", spv::Spv_1_5); builder.addCapability(spv::Capability::RuntimeDescriptorArrayEXT); } diff --git a/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.cpp b/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.cpp index 6e7f3cc..0e4237f 100644 --- a/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.cpp +++ b/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.cpp @@ -97,7 +97,7 @@ Id Builder::makeVoidType() module.mapInstruction(type); // Core OpTypeVoid used for debug void type if (emitNonSemanticShaderDebugInfo) - debugId[typeId] = typeId; + debugTypeIdLookup[typeId] = typeId; } else type = groupedTypes[enumCast(Op::OpTypeVoid)].back(); @@ -115,7 +115,7 @@ Id Builder::makeBoolType() if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeBoolDebugType(32); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } } else @@ -139,7 +139,7 @@ Id Builder::makeSamplerType(const char* debugName) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeOpaqueDebugType(debugName); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -167,7 +167,7 @@ Id Builder::makePointer(StorageClass storageClass, Id pointee) if (emitNonSemanticShaderDebugInfo) { const Id debugResultId = makePointerDebugType(storageClass, pointee); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -185,7 +185,7 @@ Id Builder::makeForwardPointer(StorageClass storageClass) if (emitNonSemanticShaderDebugInfo) { const Id debugResultId = makeForwardPointerDebugType(storageClass); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); } @@ -213,9 +213,9 @@ Id Builder::makePointerFromForwardPointer(StorageClass storageClass, Id forwardP // that was emitted alongside the forward pointer, now that we have a pointee debug // type for it to point to. if (emitNonSemanticShaderDebugInfo) { - Instruction *debugForwardPointer = module.getInstruction(debugId[forwardPointerType]); - assert(debugId[pointee]); - debugForwardPointer->setIdOperand(2, debugId[pointee]); + Instruction *debugForwardPointer = module.getInstruction(getDebugType(forwardPointerType)); + assert(getDebugType(pointee)); + debugForwardPointer->setIdOperand(2, getDebugType(pointee)); } return type->getResultId(); @@ -257,7 +257,7 @@ Id Builder::makeIntegerType(int width, bool hasSign) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeIntegerDebugType(width, hasSign); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -298,7 +298,7 @@ Id Builder::makeFloatType(int width) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeFloatDebugType(width); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -334,7 +334,7 @@ Id Builder::makeBFloat16Type() if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeFloatDebugType(width); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } #endif @@ -371,7 +371,7 @@ Id Builder::makeFloatE5M2Type() if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeFloatDebugType(width); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } #endif @@ -408,7 +408,7 @@ Id Builder::makeFloatE4M3Type() if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeFloatDebugType(width); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } #endif @@ -439,7 +439,7 @@ Id Builder::makeStructType(const std::vector& members, const std::vectorgetResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -492,7 +492,7 @@ Id Builder::makeVectorType(Id component, int size) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeVectorDebugType(component, size); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -525,7 +525,7 @@ Id Builder::makeMatrixType(Id component, int cols, int rows) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeMatrixDebugType(column, cols); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -562,7 +562,7 @@ Id Builder::makeCooperativeMatrixTypeKHR(Id component, Id scope, Id rows, Id col // Find a name for one of the parameters. It can either come from debuginfo for another // type, or an OpName from a constant. auto const findName = [&](Id id) { - Id id2 = debugId[id]; + Id id2 = getDebugType(id); for (auto &t : groupedDebugTypes[NonSemanticShaderDebugInfo100DebugTypeBasic]) { if (t->getResultId() == id2) { for (auto &s : strings) { @@ -591,7 +591,7 @@ Id Builder::makeCooperativeMatrixTypeKHR(Id component, Id scope, Id rows, Id col // There's no nonsemantic debug info instruction for cooperative matrix types, // use opaque composite instead. auto const debugResultId = makeOpaqueDebugType(debugName.c_str()); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -742,7 +742,7 @@ Id Builder::makeArrayType(Id element, Id sizeId, int stride) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeArrayDebugType(element, sizeId); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -758,7 +758,7 @@ Id Builder::makeRuntimeArray(Id element) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeArrayDebugType(element, makeUintConstant(0)); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -785,12 +785,12 @@ Id Builder::makeFunctionType(Id returnType, const std::vector& paramTypes) // function type is created for the wrapper function. However, nonsemantic shader debug information is disabled // while creating the HLSL wrapper. Consequently, if we encounter another void(void) function, we need to create // the associated debug function type if it hasn't been created yet. - if(emitNonSemanticShaderDebugInfo && debugId[type->getResultId()] == 0) { + if(emitNonSemanticShaderDebugInfo && getDebugType(type->getResultId()) == NoType) { assert(sourceLang == spv::SourceLanguage::HLSL); assert(getTypeClass(returnType) == Op::OpTypeVoid && paramTypes.size() == 0); - Id debugTypeId = makeDebugFunctionType(returnType, {}); - debugId[type->getResultId()] = debugTypeId; + Id id = makeDebugFunctionType(returnType, {}); + debugTypeIdLookup[type->getResultId()] = id; } return type->getResultId(); } @@ -810,7 +810,7 @@ Id Builder::makeFunctionType(Id returnType, const std::vector& paramTypes) // make debug type and map it if (emitNonSemanticShaderDebugInfo) { Id debugTypeId = makeDebugFunctionType(returnType, paramTypes); - debugId[typeId] = debugTypeId; + debugTypeIdLookup[typeId] = debugTypeId; } return type->getResultId(); @@ -818,7 +818,7 @@ Id Builder::makeFunctionType(Id returnType, const std::vector& paramTypes) Id Builder::makeDebugFunctionType(Id returnType, const std::vector& paramTypes) { - assert(debugId[returnType] != 0); + assert(getDebugType(returnType) != NoType); Id typeId = getUniqueId(); auto type = new Instruction(typeId, makeVoidType(), Op::OpExtInst); @@ -826,13 +826,13 @@ Id Builder::makeDebugFunctionType(Id returnType, const std::vector& paramTyp type->addIdOperand(nonSemanticShaderDebugInfo); type->addImmediateOperand(NonSemanticShaderDebugInfo100DebugTypeFunction); type->addIdOperand(makeUintConstant(NonSemanticShaderDebugInfo100FlagIsPublic)); - type->addIdOperand(debugId[returnType]); + type->addIdOperand(getDebugType(returnType)); for (auto const paramType : paramTypes) { if (isPointerType(paramType) || isArrayType(paramType)) { - type->addIdOperand(debugId[getContainedTypeId(paramType)]); + type->addIdOperand(getDebugType(getContainedTypeId(paramType))); } else { - type->addIdOperand(debugId[paramType]); + type->addIdOperand(getDebugType(paramType)); } } constantsTypesGlobals.push_back(std::unique_ptr(type)); @@ -923,7 +923,7 @@ Id Builder::makeImageType(Id sampledType, Dim dim, bool depth, bool arrayed, boo if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeOpaqueDebugType(debugName); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -950,7 +950,7 @@ Id Builder::makeSampledImageType(Id imageType, const char* debugName) if (emitNonSemanticShaderDebugInfo) { auto const debugResultId = makeOpaqueDebugType(debugName); - debugId[type->getResultId()] = debugResultId; + debugTypeIdLookup[type->getResultId()] = debugResultId; } return type->getResultId(); @@ -1099,7 +1099,7 @@ Id Builder::makeSequentialDebugType(Id const baseType, Id const componentCount, type->reserveOperands(4); type->addIdOperand(nonSemanticShaderDebugInfo); type->addImmediateOperand(sequenceType); - type->addIdOperand(debugId[baseType]); // base type + type->addIdOperand(getDebugType(baseType)); // base type type->addIdOperand(componentCount); // component count groupedDebugTypes[sequenceType].push_back(type); @@ -1135,7 +1135,7 @@ Id Builder::makeMatrixDebugType(Id const vectorType, int const vectorCount, bool type->reserveOperands(5); type->addIdOperand(nonSemanticShaderDebugInfo); type->addImmediateOperand(NonSemanticShaderDebugInfo100DebugTypeMatrix); - type->addIdOperand(debugId[vectorType]); // vector type id + type->addIdOperand(getDebugType(vectorType)); // vector type id type->addIdOperand(makeUintConstant(vectorCount)); // component count id type->addIdOperand(makeBoolConstant(columnMajor)); // column-major id @@ -1148,7 +1148,7 @@ Id Builder::makeMatrixDebugType(Id const vectorType, int const vectorCount, bool Id Builder::makeMemberDebugType(Id const memberType, StructMemberDebugInfo const& debugTypeLoc) { - assert(debugId[memberType] != 0); + assert(getDebugType(memberType) != NoType); Instruction* type = new Instruction(getUniqueId(), makeVoidType(), Op::OpExtInst); type->reserveOperands(10); @@ -1156,7 +1156,7 @@ Id Builder::makeMemberDebugType(Id const memberType, StructMemberDebugInfo const type->addImmediateOperand(NonSemanticShaderDebugInfo100DebugTypeMember); type->addIdOperand(getStringId(debugTypeLoc.name)); // name id type->addIdOperand(debugTypeLoc.debugTypeOverride != 0 ? debugTypeLoc.debugTypeOverride - : debugId[memberType]); // type id + : getDebugType(memberType)); // type id type->addIdOperand(makeDebugSource(currentFileId)); // source id type->addIdOperand(makeUintConstant(debugTypeLoc.line)); // line id TODO: currentLine is always zero type->addIdOperand(makeUintConstant(debugTypeLoc.column)); // TODO: column id @@ -1178,7 +1178,7 @@ Id Builder::makeCompositeDebugType(std::vector const& memberTypes, std::vect std::vector memberDebugTypes; assert(memberTypes.size() == memberDebugInfo.size()); for (size_t i = 0; i < memberTypes.size(); i++) { - if (debugId[memberTypes[i]]) { + if (getDebugType(memberTypes[i]) != NoType) { memberDebugTypes.emplace_back(makeMemberDebugType(memberTypes[i], memberDebugInfo[i])); } } @@ -1238,7 +1238,7 @@ Id Builder::makeOpaqueDebugType(char const* const name) Id Builder::makePointerDebugType(StorageClass storageClass, Id const baseType) { - const Id debugBaseType = debugId[baseType]; + const Id debugBaseType = getDebugType(baseType); if (!debugBaseType) { return makeDebugInfoNone(); } @@ -1474,7 +1474,7 @@ Id Builder::makeAccelerationStructureType() module.mapInstruction(type); if (emitNonSemanticShaderDebugInfo) { spv::Id debugType = makeOpaqueDebugType("accelerationStructure"); - debugId[type->getResultId()] = debugType; + debugTypeIdLookup[type->getResultId()] = debugType; } } else { type = groupedTypes[enumCast(Op::OpTypeAccelerationStructureKHR)].back(); @@ -1493,7 +1493,7 @@ Id Builder::makeRayQueryType() module.mapInstruction(type); if (emitNonSemanticShaderDebugInfo) { spv::Id debugType = makeOpaqueDebugType("rayQuery"); - debugId[type->getResultId()] = debugType; + debugTypeIdLookup[type->getResultId()] = debugType; } } else { type = groupedTypes[enumCast(Op::OpTypeRayQueryKHR)].back(); @@ -1526,7 +1526,7 @@ Id Builder::makeHitObjectNVType() module.mapInstruction(type); if (emitNonSemanticShaderDebugInfo) { spv::Id debugType = makeOpaqueDebugType("hitObjectNV"); - debugId[type->getResultId()] = debugType; + debugTypeIdLookup[type->getResultId()] = debugType; } } else { type = groupedTypes[enumCast(Op::OpTypeHitObjectNV)].back(); @@ -1734,32 +1734,17 @@ bool Builder::containsPhysicalStorageBufferOrArray(Id typeId) const // can be reused rather than duplicated. (Required by the specification). Id Builder::findScalarConstant(Op typeClass, Op opcode, Id typeId, unsigned value) { - Instruction* constant; - for (int i = 0; i < (int)groupedConstants[enumCast(typeClass)].size(); ++i) { - constant = groupedConstants[enumCast(typeClass)][i]; - if (constant->getOpCode() == opcode && - constant->getTypeId() == typeId && - constant->getImmediateOperand(0) == value) - return constant->getResultId(); - } - - return 0; + ScalarConstantKey key{ enumCast(typeClass), enumCast(opcode), typeId, value, 0 }; + auto it = groupedScalarConstantResultIDs.find(key); + return (it != groupedScalarConstantResultIDs.end()) ? it->second : 0; } // Version of findScalarConstant (see above) for scalars that take two operands (e.g. a 'double' or 'int64'). Id Builder::findScalarConstant(Op typeClass, Op opcode, Id typeId, unsigned v1, unsigned v2) { - Instruction* constant; - for (int i = 0; i < (int)groupedConstants[enumCast(typeClass)].size(); ++i) { - constant = groupedConstants[enumCast(typeClass)][i]; - if (constant->getOpCode() == opcode && - constant->getTypeId() == typeId && - constant->getImmediateOperand(0) == v1 && - constant->getImmediateOperand(1) == v2) - return constant->getResultId(); - } - - return 0; + ScalarConstantKey key{ enumCast(typeClass), enumCast(opcode), typeId, v1, v2 }; + auto it = groupedScalarConstantResultIDs.find(key); + return (it != groupedScalarConstantResultIDs.end()) ? it->second : 0; } // Return true if consuming 'opcode' means consuming a constant. @@ -1831,19 +1816,12 @@ Id Builder::makeNullConstant(Id typeId) Id Builder::makeBoolConstant(bool b, bool specConstant) { Id typeId = makeBoolType(); - Instruction* constant; Op opcode = specConstant ? (b ? Op::OpSpecConstantTrue : Op::OpSpecConstantFalse) : (b ? Op::OpConstantTrue : Op::OpConstantFalse); // See if we already made it. Applies only to regular constants, because specialization constants // must remain distinct for the purpose of applying a SpecId decoration. - if (! specConstant) { - Id existing = 0; - for (int i = 0; i < (int)groupedConstants[enumCast(Op::OpTypeBool)].size(); ++i) { - constant = groupedConstants[enumCast(Op::OpTypeBool)][i]; - if (constant->getTypeId() == typeId && constant->getOpCode() == opcode) - existing = constant->getResultId(); - } - + if (!specConstant) { + Id existing = findScalarConstant(Op::OpTypeBool, opcode, typeId, 0); if (existing) return existing; } @@ -1851,10 +1829,14 @@ Id Builder::makeBoolConstant(bool b, bool specConstant) // Make it Instruction* c = new Instruction(getUniqueId(), typeId, opcode); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeBool)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{enumCast(Op::OpTypeBool), enumCast(opcode), typeId, 0, 0}; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeIntConstant(Id typeId, unsigned value, bool specConstant) @@ -1872,10 +1854,14 @@ Id Builder::makeIntConstant(Id typeId, unsigned value, bool specConstant) Instruction* c = new Instruction(getUniqueId(), typeId, opcode); c->addImmediateOperand(value); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeInt)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{ enumCast(Op::OpTypeInt), enumCast(opcode), typeId, value, 0 }; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeInt64Constant(Id typeId, unsigned long long value, bool specConstant) @@ -1898,10 +1884,14 @@ Id Builder::makeInt64Constant(Id typeId, unsigned long long value, bool specCons c->addImmediateOperand(op1); c->addImmediateOperand(op2); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeInt)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{ enumCast(Op::OpTypeInt), enumCast(opcode), typeId, op1, op2 }; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeFloatConstant(float f, bool specConstant) @@ -1923,10 +1913,14 @@ Id Builder::makeFloatConstant(float f, bool specConstant) Instruction* c = new Instruction(getUniqueId(), typeId, opcode); c->addImmediateOperand(value); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeFloat)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{ enumCast(Op::OpTypeFloat), enumCast(opcode), typeId, value, 0 }; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeDoubleConstant(double d, bool specConstant) @@ -1952,10 +1946,14 @@ Id Builder::makeDoubleConstant(double d, bool specConstant) c->addImmediateOperand(op1); c->addImmediateOperand(op2); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeFloat)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{ enumCast(Op::OpTypeFloat), enumCast(opcode), typeId, op1, op2 }; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeFloat16Constant(float f16, bool specConstant) @@ -1980,10 +1978,14 @@ Id Builder::makeFloat16Constant(float f16, bool specConstant) Instruction* c = new Instruction(getUniqueId(), typeId, opcode); c->addImmediateOperand(value); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeFloat)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{ enumCast(Op::OpTypeFloat), enumCast(opcode), typeId, value, 0 }; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeBFloat16Constant(float bf16, bool specConstant) @@ -2011,10 +2013,14 @@ Id Builder::makeBFloat16Constant(float bf16, bool specConstant) Instruction* c = new Instruction(getUniqueId(), typeId, opcode); c->addImmediateOperand(value); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeFloat)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{ enumCast(Op::OpTypeFloat), enumCast(opcode), typeId, value, 0 }; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeFloatE5M2Constant(float fe5m2, bool specConstant) @@ -2039,10 +2045,14 @@ Id Builder::makeFloatE5M2Constant(float fe5m2, bool specConstant) Instruction* c = new Instruction(getUniqueId(), typeId, opcode); c->addImmediateOperand(value); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeFloat)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{enumCast(Op::OpTypeFloat), enumCast(opcode), typeId, value, 0}; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeFloatE4M3Constant(float fe4m3, bool specConstant) @@ -2067,10 +2077,14 @@ Id Builder::makeFloatE4M3Constant(float fe4m3, bool specConstant) Instruction* c = new Instruction(getUniqueId(), typeId, opcode); c->addImmediateOperand(value); constantsTypesGlobals.push_back(std::unique_ptr(c)); - groupedConstants[enumCast(Op::OpTypeFloat)].push_back(c); module.mapInstruction(c); - return c->getResultId(); + Id resultId = c->getResultId(); + if (!specConstant) { + ScalarConstantKey key{enumCast(Op::OpTypeFloat), enumCast(opcode), typeId, value, 0}; + groupedScalarConstantResultIDs[key] = resultId; + } + return resultId; } Id Builder::makeFpConstant(Id type, double d, bool specConstant) @@ -2111,8 +2125,8 @@ Id Builder::findCompositeConstant(Op typeClass, Op opcode, Id typeId, const std: { Instruction* constant = nullptr; bool found = false; - for (int i = 0; i < (int)groupedConstants[enumCast(typeClass)].size(); ++i) { - constant = groupedConstants[enumCast(typeClass)][i]; + for (int i = 0; i < (int)groupedCompositeConstants[enumCast(typeClass)].size(); ++i) { + constant = groupedCompositeConstants[enumCast(typeClass)][i]; if (constant->getTypeId() != typeId) continue; @@ -2222,7 +2236,7 @@ Id Builder::makeCompositeConstant(Id typeId, const std::vector& members, boo if (typeClass == Op::OpTypeStruct) groupedStructConstants[typeId].push_back(c); else - groupedConstants[enumCast(typeClass)].push_back(c); + groupedCompositeConstants[enumCast(typeClass)].push_back(c); module.mapInstruction(c); return c->getResultId(); @@ -2603,14 +2617,14 @@ void Builder::setupFunctionDebugInfo(Function* function, const char* name, const Id nameId = getStringId(unmangleFunctionName(name)); Id funcTypeId = function->getFuncTypeId(); - assert(debugId[funcTypeId] != 0); + assert(getDebugType(funcTypeId) != NoType); Id funcId = function->getId(); assert(funcId != 0); // Make the debug function instruction Id debugFuncId = makeDebugFunction(function, nameId, funcTypeId); - debugId[funcId] = debugFuncId; + debugFuncIdLookup[funcId] = debugFuncId; currentDebugScopeId.push(debugFuncId); // DebugScope and DebugLine for parameter DebugDeclares @@ -2629,9 +2643,8 @@ void Builder::setupFunctionDebugInfo(Function* function, const char* name, const } auto const& paramName = paramNames[p]; - auto const debugLocalVariableId = createDebugLocalVariable(debugId[paramTypeId], paramName, p + 1); + auto const debugLocalVariableId = createDebugLocalVariable(getDebugType(paramTypeId), paramName, p + 1); auto const paramId = static_cast(firstParamId + p); - debugId[paramId] = debugLocalVariableId; if (passByRef) { makeDebugDeclare(debugLocalVariableId, paramId); @@ -2651,7 +2664,7 @@ Id Builder::makeDebugFunction([[maybe_unused]] Function* function, Id nameId, Id assert(function != nullptr); assert(nameId != 0); assert(funcTypeId != 0); - assert(debugId[funcTypeId] != 0); + assert(getDebugType(funcTypeId) != NoType); Id funcId = getUniqueId(); auto type = new Instruction(funcId, makeVoidType(), Op::OpExtInst); @@ -2659,7 +2672,7 @@ Id Builder::makeDebugFunction([[maybe_unused]] Function* function, Id nameId, Id type->addIdOperand(nonSemanticShaderDebugInfo); type->addImmediateOperand(NonSemanticShaderDebugInfo100DebugFunction); type->addIdOperand(nameId); - type->addIdOperand(debugId[funcTypeId]); + type->addIdOperand(getDebugType(funcTypeId)); type->addIdOperand(makeDebugSource(currentFileId)); // TODO: This points to file of definition instead of declaration type->addIdOperand(makeUintConstant(currentLine)); // TODO: This points to line of definition instead of declaration type->addIdOperand(makeUintConstant(0)); // column @@ -2754,14 +2767,15 @@ void Builder::enterFunction(Function const* function) if (emitNonSemanticShaderDebugInfo) { // Initialize scope state Id funcId = function->getFuncId(); - currentDebugScopeId.push(debugId[funcId]); + Id debugFuncId = getDebugFunction(funcId); + currentDebugScopeId.push(debugFuncId); // Create DebugFunctionDefinition spv::Id resultId = getUniqueId(); Instruction* defInst = new Instruction(resultId, makeVoidType(), Op::OpExtInst); defInst->reserveOperands(4); defInst->addIdOperand(nonSemanticShaderDebugInfo); defInst->addImmediateOperand(NonSemanticShaderDebugInfo100DebugFunctionDefinition); - defInst->addIdOperand(debugId[funcId]); + defInst->addIdOperand(debugFuncId); defInst->addIdOperand(funcId); addInstruction(std::unique_ptr(defInst)); } @@ -2855,14 +2869,11 @@ Id Builder::createVariable(Decoration precision, StorageClass storageClass, Id t // We may emulate some local variables as global variable with private storage in SPIR-V, but we still want to // treat them as local variables in debug info. if (storageClass == StorageClass::Function || (currentFunction && storageClass == StorageClass::Private)) { - auto const debugLocalVariableId = createDebugLocalVariable(debugId[type], name); - debugId[inst->getResultId()] = debugLocalVariableId; - + auto const debugLocalVariableId = createDebugLocalVariable(getDebugType(type), name); makeDebugDeclare(debugLocalVariableId, inst->getResultId()); } else { - auto const debugResultId = createDebugGlobalVariable(debugId[type], name, inst->getResultId()); - debugId[inst->getResultId()] = debugResultId; + createDebugGlobalVariable(getDebugType(type), name, inst->getResultId()); } } diff --git a/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.h b/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.h index 68b4764..70c1def 100644 --- a/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.h +++ b/libs/bgfx/3rdparty/glslang/SPIRV/SpvBuilder.h @@ -208,10 +208,20 @@ class Builder { // Maps the given OpType Id to a Non-Semantic DebugType Id. Id getDebugType(Id type) { - if (emitNonSemanticShaderDebugInfo) { - return debugId[type]; + if (auto it = debugTypeIdLookup.find(type); it != debugTypeIdLookup.end()) { + return it->second; } - return 0; + + return NoType; + } + + // Maps the given OpFunction Id to a Non-Semantic DebugFunction Id. + Id getDebugFunction(Id func) { + if (auto it = debugFuncIdLookup.find(func); it != debugFuncIdLookup.end()) { + return it->second; + } + + return NoResult; } // For creating new types (will return old type if the requested one was already made). @@ -1031,8 +1041,58 @@ class Builder { // not output, internally used for quick & dirty canonical (unique) creation + // Key for scalar constants (handles both 32-bit and 64-bit) + struct ScalarConstantKey { + unsigned int typeClass; // OpTypeInt, OpTypeFloat, OpTypeBool + unsigned int opcode; // OpConstant, OpSpecConstant, OpConstantTrue, etc. + Id typeId; // The specific type + unsigned value1; // First operand (or only operand) + unsigned value2; // Second operand (0 for single-operand constants) + + bool operator==(const ScalarConstantKey& other) const { + return typeClass == other.typeClass && + opcode == other.opcode && + typeId == other.typeId && + value1 == other.value1 && + value2 == other.value2; + } + }; + + struct ScalarConstantKeyHash { + // 64/32 bit mix function from MurmurHash3 + inline std::size_t hash_mix(std::size_t h) const { + if constexpr (sizeof(std::size_t) == 8) { + h ^= h >> 33; + h *= UINT64_C(0xff51afd7ed558ccd); + h ^= h >> 33; + h *= UINT64_C(0xc4ceb9fe1a85ec53); + h ^= h >> 33; + return h; + } else { + h ^= h >> 16; + h *= UINT32_C(0x85ebca6b); + h ^= h >> 13; + h *= UINT32_C(0xc2b2ae35); + h ^= h >> 16; + return h; + } + } + + // Hash combine from boost + inline std::size_t hash_combine(std::size_t seed, std::size_t v) const { + return hash_mix(seed + 0x9e3779b9 + v); + } + + std::size_t operator()(const ScalarConstantKey& k) const { + size_t hash1 = hash_combine(std::hash{}(k.typeClass), std::hash{}(k.opcode)); + size_t hash2 = hash_combine(std::hash{}(k.value1), std::hash{}(k.value2)); + size_t hash3 = hash_combine(hash1, hash2); + return hash_combine(hash3, std::hash{}(k.typeId)); + } + }; + // map type opcodes to constant inst. - std::unordered_map> groupedConstants; + std::unordered_map> groupedCompositeConstants; // map struct-id to constant instructions std::unordered_map> groupedStructConstants; // map type opcodes to type instructions @@ -1041,6 +1101,8 @@ class Builder { std::unordered_map> groupedDebugTypes; // list of OpConstantNull instructions std::vector nullConstants; + // map scalar constants to result IDs + std::unordered_map groupedScalarConstantResultIDs; // Track which types have explicit layouts, to avoid reusing in storage classes without layout. // Currently only tracks array types. @@ -1058,8 +1120,11 @@ class Builder { // map from include file name ids to their contents std::map includeFiles; - // map from core id to debug id - std::map debugId; + // maps from OpTypeXXX id to DebugTypeXXX id + std::unordered_map debugTypeIdLookup; + + // maps from OpFunction id to DebugFunction id + std::unordered_map debugFuncIdLookup; // map from file name string id to DebugSource id std::unordered_map debugSourceId; diff --git a/libs/bgfx/3rdparty/glslang/build_info.h b/libs/bgfx/3rdparty/glslang/build_info.h index d8f2a74..7816b66 100644 --- a/libs/bgfx/3rdparty/glslang/build_info.h +++ b/libs/bgfx/3rdparty/glslang/build_info.h @@ -35,7 +35,7 @@ #define GLSLANG_BUILD_INFO #define GLSLANG_VERSION_MAJOR 16 -#define GLSLANG_VERSION_MINOR 0 +#define GLSLANG_VERSION_MINOR 1 #define GLSLANG_VERSION_PATCH 0 #define GLSLANG_VERSION_FLAVOR "" diff --git a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/ParseHelper.cpp b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/ParseHelper.cpp index 667ed67..38ce7bd 100644 --- a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/ParseHelper.cpp +++ b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/ParseHelper.cpp @@ -4188,6 +4188,10 @@ void TParseContext::reservedErrorCheck(const TSourceLoc& loc, const TString& ide // "Identifiers starting with "gl_" are reserved for use by OpenGL, and may not be // declared in a shader; this results in a compile-time error." if (! symbolTable.atBuiltInLevel()) { + // The extension GL_EXT_conservative_depth allows us to declare "gl_FragDepth". + if (identifier == "gl_FragDepth" && extensionTurnedOn(E_GL_EXT_conservative_depth)) + return; + if (builtInName(identifier) && !extensionTurnedOn(E_GL_EXT_spirv_intrinsics)) // The extension GL_EXT_spirv_intrinsics allows us to declare identifiers starting with "gl_". error(loc, "identifiers starting with \"gl_\" are reserved", identifier.c_str(), ""); @@ -5830,7 +5834,8 @@ TSymbol* TParseContext::redeclareBuiltinVariable(const TSourceLoc& loc, const TS bool nonEsRedecls = (!isEsProfile() && (version >= 130 || identifier == "gl_TexCoord")); bool esRedecls = (isEsProfile() && - (version >= 320 || extensionsTurnedOn(Num_AEP_shader_io_blocks, AEP_shader_io_blocks))); + (version >= 320 || extensionsTurnedOn(Num_AEP_shader_io_blocks, AEP_shader_io_blocks) || + (identifier == "gl_FragDepth" && extensionTurnedOn(E_GL_EXT_conservative_depth)))); if (! esRedecls && ! nonEsRedecls) return nullptr; @@ -6535,6 +6540,9 @@ void TParseContext::finish() if (parsingBuiltins) return; + // Forward builtin alias to AST for later use + intermediate.setBuiltinAliasLookup(symbolTable.collectBuiltinAlias()); + // Check on array indexes for ES 2.0 (version 100) limitations. for (size_t i = 0; i < needsIndexLimitationChecking.size(); ++i) constantIndexExpressionCheck(needsIndexLimitationChecking[i]); diff --git a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/SymbolTable.h b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/SymbolTable.h index edeb8d5..9161927 100644 --- a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/SymbolTable.h +++ b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/SymbolTable.h @@ -70,6 +70,7 @@ #include "../Include/InfoSink.h" #include +#include namespace glslang { @@ -505,6 +506,11 @@ class TSymbolTableLevel { retargetedSymbols.push_back({from, to}); } + void collectRetargetedSymbols(std::unordered_multimap &out) const { + for (const auto &[fromName, toName] : retargetedSymbols) + out.insert({std::string{toName}, std::string{fromName}}); + } + TSymbol* find(const TString& name) const { tLevel::const_iterator it = level.find(name); @@ -662,9 +668,10 @@ class TSymbolTable { // protected: static const uint32_t LevelFlagBitOffset = 56; - static const int globalLevel = 3; + static constexpr int builtinLevel = 2; + static constexpr int globalLevel = 3; static bool isSharedLevel(int level) { return level <= 1; } // exclude all per-compile levels - static bool isBuiltInLevel(int level) { return level <= 2; } // exclude user globals + static bool isBuiltInLevel(int level) { return level <= builtinLevel; } // exclude user globals static bool isGlobalLevel(int level) { return level <= globalLevel; } // include user globals public: bool isEmpty() { return table.size() == 0; } @@ -829,6 +836,13 @@ class TSymbolTable { table[level]->retargetSymbol(from, to); } + std::unordered_multimap collectBuiltinAlias() { + std::unordered_multimap allRetargets; + for (int level = 0; level <= std::min(currentLevel(), builtinLevel); ++level) + table[level]->collectRetargetedSymbols(allRetargets); + + return allRetargets; + } // Find of a symbol that returns how many layers deep of nested // structures-with-member-functions ('this' scopes) deep the symbol was diff --git a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/Versions.cpp b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/Versions.cpp index e8ce150..e01f96b 100644 --- a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/Versions.cpp +++ b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/Versions.cpp @@ -728,8 +728,8 @@ void TParseVersions::getPreamble(std::string& preamble) case EShLangClosestHit: preamble += "#define GL_CLOSEST_HIT_SHADER_EXT 1 \n"; break; case EShLangMiss: preamble += "#define GL_MISS_SHADER_EXT 1 \n"; break; case EShLangCallable: preamble += "#define GL_CALLABLE_SHADER_EXT 1 \n"; break; - case EShLangTask: preamble += "#define GL_TASK_SHADER_NV 1 \n"; break; - case EShLangMesh: preamble += "#define GL_MESH_SHADER_NV 1 \n"; break; + case EShLangTask: preamble += "#define GL_TASK_SHADER_EXT 1 \n"; break; + case EShLangMesh: preamble += "#define GL_MESH_SHADER_EXT 1 \n"; break; default: break; } } diff --git a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/localintermediate.h b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/localintermediate.h index 2e10744..ce57ce4 100644 --- a/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/localintermediate.h +++ b/libs/bgfx/3rdparty/glslang/glslang/MachineIndependent/localintermediate.h @@ -48,6 +48,7 @@ #include #include #include +#include #include class TInfoSink; @@ -1161,6 +1162,13 @@ class TIntermediate { void updateNumericFeature(TNumericFeatures::feature f, bool on) { on ? numericFeatures.insert(f) : numericFeatures.erase(f); } + void setBuiltinAliasLookup(std::unordered_multimap symbolMap) { + builtinAliasLookup = std::move(symbolMap); + } + const std::unordered_multimap& getBuiltinAliasLookup() const { + return builtinAliasLookup; + } + protected: TIntermSymbol* addSymbol(long long Id, const TString&, const TString&, const TType&, const TConstUnionArray&, TIntermTyped* subtree, const TSourceLoc&); void error(TInfoSink& infoSink, const TSourceLoc* loc, EShMessages messages, const char*, EShLanguage unitStage = EShLangCount); @@ -1335,6 +1343,9 @@ class TIntermediate { // Included text. First string is a name, second is the included text std::map includeText; + // Maps from canonical symbol name to alias symbol names + std::unordered_multimap builtinAliasLookup; + // for OpModuleProcessed, or equivalent TProcesses processes; diff --git a/libs/bgfx/3rdparty/khronos/vulkan-local/vulkan_core.h b/libs/bgfx/3rdparty/khronos/vulkan-local/vulkan_core.h index 51a6189..922db39 100644 --- a/libs/bgfx/3rdparty/khronos/vulkan-local/vulkan_core.h +++ b/libs/bgfx/3rdparty/khronos/vulkan-local/vulkan_core.h @@ -66,7 +66,7 @@ extern "C" { //#define VK_API_VERSION VK_MAKE_API_VERSION(0, 1, 0, 0) // Patch version should always be set to 0 // Version of this file -#define VK_HEADER_VERSION 333 +#define VK_HEADER_VERSION 336 // Complete version of this file #define VK_HEADER_VERSION_COMPLETE VK_MAKE_API_VERSION(0, 1, 4, VK_HEADER_VERSION) @@ -177,6 +177,7 @@ typedef enum VkResult { VK_ERROR_VIDEO_PROFILE_CODEC_NOT_SUPPORTED_KHR = -1000023004, VK_ERROR_VIDEO_STD_VERSION_NOT_SUPPORTED_KHR = -1000023005, VK_ERROR_INVALID_DRM_FORMAT_MODIFIER_PLANE_LAYOUT_EXT = -1000158000, + VK_ERROR_PRESENT_TIMING_QUEUE_FULL_EXT = -1000208000, VK_ERROR_FULL_SCREEN_EXCLUSIVE_MODE_LOST_EXT = -1000255000, VK_THREAD_IDLE_KHR = 1000268000, VK_THREAD_DONE_KHR = 1000268001, @@ -746,6 +747,16 @@ typedef enum VkStructureType { VK_STRUCTURE_TYPE_QUEUE_FAMILY_CHECKPOINT_PROPERTIES_NV = 1000206001, VK_STRUCTURE_TYPE_QUEUE_FAMILY_CHECKPOINT_PROPERTIES_2_NV = 1000314008, VK_STRUCTURE_TYPE_CHECKPOINT_DATA_2_NV = 1000314009, + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PRESENT_TIMING_FEATURES_EXT = 1000208000, + VK_STRUCTURE_TYPE_SWAPCHAIN_TIMING_PROPERTIES_EXT = 1000208001, + VK_STRUCTURE_TYPE_SWAPCHAIN_TIME_DOMAIN_PROPERTIES_EXT = 1000208002, + VK_STRUCTURE_TYPE_PRESENT_TIMINGS_INFO_EXT = 1000208003, + VK_STRUCTURE_TYPE_PRESENT_TIMING_INFO_EXT = 1000208004, + VK_STRUCTURE_TYPE_PAST_PRESENTATION_TIMING_INFO_EXT = 1000208005, + VK_STRUCTURE_TYPE_PAST_PRESENTATION_TIMING_PROPERTIES_EXT = 1000208006, + VK_STRUCTURE_TYPE_PAST_PRESENTATION_TIMING_EXT = 1000208007, + VK_STRUCTURE_TYPE_PRESENT_TIMING_SURFACE_CAPABILITIES_EXT = 1000208008, + VK_STRUCTURE_TYPE_SWAPCHAIN_CALIBRATED_TIMESTAMP_INFO_EXT = 1000208009, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_INTEGER_FUNCTIONS_2_FEATURES_INTEL = 1000209000, VK_STRUCTURE_TYPE_QUERY_POOL_PERFORMANCE_QUERY_CREATE_INFO_INTEL = 1000210000, VK_STRUCTURE_TYPE_INITIALIZE_PERFORMANCE_API_INFO_INTEL = 1000210001, @@ -1134,6 +1145,7 @@ typedef enum VkStructureType { VK_STRUCTURE_TYPE_SWAPCHAIN_PRESENT_SCALING_CREATE_INFO_KHR = 1000275004, VK_STRUCTURE_TYPE_RELEASE_SWAPCHAIN_IMAGES_INFO_KHR = 1000275005, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_MULTIVIEW_PER_VIEW_VIEWPORTS_FEATURES_QCOM = 1000488000, + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_RAY_TRACING_INVOCATION_REORDER_FEATURES_NV = 1000490000, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_RAY_TRACING_INVOCATION_REORDER_PROPERTIES_NV = 1000490001, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_VECTOR_FEATURES_NV = 1000491000, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COOPERATIVE_VECTOR_PROPERTIES_NV = 1000491001, @@ -1371,6 +1383,8 @@ typedef enum VkStructureType { VK_STRUCTURE_TYPE_RESOLVE_IMAGE_MODE_INFO_KHR = 1000630004, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PIPELINE_CACHE_INCREMENTAL_MODE_FEATURES_SEC = 1000637000, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_UNIFORM_BUFFER_UNSIZED_ARRAY_FEATURES_EXT = 1000642000, + VK_STRUCTURE_TYPE_COMPUTE_OCCUPANCY_PRIORITY_PARAMETERS_NV = 1000645000, + VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_COMPUTE_OCCUPANCY_PRIORITY_FEATURES_NV = 1000645001, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VARIABLE_POINTER_FEATURES = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VARIABLE_POINTERS_FEATURES, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_DRAW_PARAMETER_FEATURES = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_DRAW_PARAMETERS_FEATURES, // VK_STRUCTURE_TYPE_DEBUG_REPORT_CREATE_INFO_EXT is a legacy alias @@ -1609,7 +1623,6 @@ typedef enum VkStructureType { VK_STRUCTURE_TYPE_PIPELINE_CREATE_FLAGS_2_CREATE_INFO_KHR = VK_STRUCTURE_TYPE_PIPELINE_CREATE_FLAGS_2_CREATE_INFO, VK_STRUCTURE_TYPE_BUFFER_USAGE_FLAGS_2_CREATE_INFO_KHR = VK_STRUCTURE_TYPE_BUFFER_USAGE_FLAGS_2_CREATE_INFO, VK_STRUCTURE_TYPE_SHADER_REQUIRED_SUBGROUP_SIZE_CREATE_INFO_EXT = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_REQUIRED_SUBGROUP_SIZE_CREATE_INFO, - VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_RAY_TRACING_INVOCATION_REORDER_FEATURES_NV = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_RAY_TRACING_INVOCATION_REORDER_FEATURES_EXT, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VERTEX_ATTRIBUTE_DIVISOR_PROPERTIES_KHR = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VERTEX_ATTRIBUTE_DIVISOR_PROPERTIES, VK_STRUCTURE_TYPE_PIPELINE_VERTEX_INPUT_DIVISOR_STATE_CREATE_INFO_KHR = VK_STRUCTURE_TYPE_PIPELINE_VERTEX_INPUT_DIVISOR_STATE_CREATE_INFO, VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VERTEX_ATTRIBUTE_DIVISOR_FEATURES_KHR = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_VERTEX_ATTRIBUTE_DIVISOR_FEATURES, @@ -8881,6 +8894,7 @@ typedef enum VkSwapchainCreateFlagBitsKHR { VK_SWAPCHAIN_CREATE_SPLIT_INSTANCE_BIND_REGIONS_BIT_KHR = 0x00000001, VK_SWAPCHAIN_CREATE_PROTECTED_BIT_KHR = 0x00000002, VK_SWAPCHAIN_CREATE_MUTABLE_FORMAT_BIT_KHR = 0x00000004, + VK_SWAPCHAIN_CREATE_PRESENT_TIMING_BIT_EXT = 0x00000200, VK_SWAPCHAIN_CREATE_PRESENT_ID_2_BIT_KHR = 0x00000040, VK_SWAPCHAIN_CREATE_PRESENT_WAIT_2_BIT_KHR = 0x00000080, VK_SWAPCHAIN_CREATE_DEFERRED_MEMORY_ALLOCATION_BIT_KHR = 0x00000008, @@ -13376,6 +13390,8 @@ typedef enum VkTimeDomainKHR { VK_TIME_DOMAIN_CLOCK_MONOTONIC_KHR = 1, VK_TIME_DOMAIN_CLOCK_MONOTONIC_RAW_KHR = 2, VK_TIME_DOMAIN_QUERY_PERFORMANCE_COUNTER_KHR = 3, + VK_TIME_DOMAIN_PRESENT_STAGE_LOCAL_EXT = 1000208000, + VK_TIME_DOMAIN_SWAPCHAIN_LOCAL_EXT = 1000208001, VK_TIME_DOMAIN_DEVICE_EXT = VK_TIME_DOMAIN_DEVICE_KHR, VK_TIME_DOMAIN_CLOCK_MONOTONIC_EXT = VK_TIME_DOMAIN_CLOCK_MONOTONIC_KHR, VK_TIME_DOMAIN_CLOCK_MONOTONIC_RAW_EXT = VK_TIME_DOMAIN_CLOCK_MONOTONIC_RAW_KHR, @@ -16812,6 +16828,161 @@ VKAPI_ATTR void VKAPI_CALL vkGetQueueCheckpointData2NV( #endif +// VK_EXT_present_timing is a preprocessor guard. Do not pass it to API calls. +#define VK_EXT_present_timing 1 +#define VK_EXT_PRESENT_TIMING_SPEC_VERSION 3 +#define VK_EXT_PRESENT_TIMING_EXTENSION_NAME "VK_EXT_present_timing" + +typedef enum VkPresentStageFlagBitsEXT { + VK_PRESENT_STAGE_QUEUE_OPERATIONS_END_BIT_EXT = 0x00000001, + VK_PRESENT_STAGE_REQUEST_DEQUEUED_BIT_EXT = 0x00000002, + VK_PRESENT_STAGE_IMAGE_FIRST_PIXEL_OUT_BIT_EXT = 0x00000004, + VK_PRESENT_STAGE_IMAGE_FIRST_PIXEL_VISIBLE_BIT_EXT = 0x00000008, + VK_PRESENT_STAGE_FLAG_BITS_MAX_ENUM_EXT = 0x7FFFFFFF +} VkPresentStageFlagBitsEXT; +typedef VkFlags VkPresentStageFlagsEXT; + +typedef enum VkPastPresentationTimingFlagBitsEXT { + VK_PAST_PRESENTATION_TIMING_ALLOW_PARTIAL_RESULTS_BIT_EXT = 0x00000001, + VK_PAST_PRESENTATION_TIMING_ALLOW_OUT_OF_ORDER_RESULTS_BIT_EXT = 0x00000002, + VK_PAST_PRESENTATION_TIMING_FLAG_BITS_MAX_ENUM_EXT = 0x7FFFFFFF +} VkPastPresentationTimingFlagBitsEXT; +typedef VkFlags VkPastPresentationTimingFlagsEXT; + +typedef enum VkPresentTimingInfoFlagBitsEXT { + VK_PRESENT_TIMING_INFO_PRESENT_AT_RELATIVE_TIME_BIT_EXT = 0x00000001, + VK_PRESENT_TIMING_INFO_PRESENT_AT_NEAREST_REFRESH_CYCLE_BIT_EXT = 0x00000002, + VK_PRESENT_TIMING_INFO_FLAG_BITS_MAX_ENUM_EXT = 0x7FFFFFFF +} VkPresentTimingInfoFlagBitsEXT; +typedef VkFlags VkPresentTimingInfoFlagsEXT; +typedef struct VkPhysicalDevicePresentTimingFeaturesEXT { + VkStructureType sType; + void* pNext; + VkBool32 presentTiming; + VkBool32 presentAtAbsoluteTime; + VkBool32 presentAtRelativeTime; +} VkPhysicalDevicePresentTimingFeaturesEXT; + +typedef struct VkPresentTimingSurfaceCapabilitiesEXT { + VkStructureType sType; + void* pNext; + VkBool32 presentTimingSupported; + VkBool32 presentAtAbsoluteTimeSupported; + VkBool32 presentAtRelativeTimeSupported; + VkPresentStageFlagsEXT presentStageQueries; +} VkPresentTimingSurfaceCapabilitiesEXT; + +typedef struct VkSwapchainCalibratedTimestampInfoEXT { + VkStructureType sType; + const void* pNext; + VkSwapchainKHR swapchain; + VkPresentStageFlagsEXT presentStage; + uint64_t timeDomainId; +} VkSwapchainCalibratedTimestampInfoEXT; + +typedef struct VkSwapchainTimingPropertiesEXT { + VkStructureType sType; + void* pNext; + uint64_t refreshDuration; + uint64_t refreshInterval; +} VkSwapchainTimingPropertiesEXT; + +typedef struct VkSwapchainTimeDomainPropertiesEXT { + VkStructureType sType; + void* pNext; + uint32_t timeDomainCount; + VkTimeDomainKHR* pTimeDomains; + uint64_t* pTimeDomainIds; +} VkSwapchainTimeDomainPropertiesEXT; + +typedef struct VkPastPresentationTimingInfoEXT { + VkStructureType sType; + const void* pNext; + VkPastPresentationTimingFlagsEXT flags; + VkSwapchainKHR swapchain; +} VkPastPresentationTimingInfoEXT; + +typedef struct VkPresentStageTimeEXT { + VkPresentStageFlagsEXT stage; + uint64_t time; +} VkPresentStageTimeEXT; + +typedef struct VkPastPresentationTimingEXT { + VkStructureType sType; + void* pNext; + uint64_t presentId; + uint64_t targetTime; + uint32_t presentStageCount; + VkPresentStageTimeEXT* pPresentStages; + VkTimeDomainKHR timeDomain; + uint64_t timeDomainId; + VkBool32 reportComplete; +} VkPastPresentationTimingEXT; + +typedef struct VkPastPresentationTimingPropertiesEXT { + VkStructureType sType; + void* pNext; + uint64_t timingPropertiesCounter; + uint64_t timeDomainsCounter; + uint32_t presentationTimingCount; + VkPastPresentationTimingEXT* pPresentationTimings; +} VkPastPresentationTimingPropertiesEXT; + +typedef struct VkPresentTimingInfoEXT { + VkStructureType sType; + const void* pNext; + VkPresentTimingInfoFlagsEXT flags; + uint64_t targetTime; + uint64_t timeDomainId; + VkPresentStageFlagsEXT presentStageQueries; + VkPresentStageFlagsEXT targetTimeDomainPresentStage; +} VkPresentTimingInfoEXT; + +typedef struct VkPresentTimingsInfoEXT { + VkStructureType sType; + const void* pNext; + uint32_t swapchainCount; + const VkPresentTimingInfoEXT* pTimingInfos; +} VkPresentTimingsInfoEXT; + +typedef VkResult (VKAPI_PTR *PFN_vkSetSwapchainPresentTimingQueueSizeEXT)(VkDevice device, VkSwapchainKHR swapchain, uint32_t size); +typedef VkResult (VKAPI_PTR *PFN_vkGetSwapchainTimingPropertiesEXT)(VkDevice device, VkSwapchainKHR swapchain, VkSwapchainTimingPropertiesEXT* pSwapchainTimingProperties, uint64_t* pSwapchainTimingPropertiesCounter); +typedef VkResult (VKAPI_PTR *PFN_vkGetSwapchainTimeDomainPropertiesEXT)(VkDevice device, VkSwapchainKHR swapchain, VkSwapchainTimeDomainPropertiesEXT* pSwapchainTimeDomainProperties, uint64_t* pTimeDomainsCounter); +typedef VkResult (VKAPI_PTR *PFN_vkGetPastPresentationTimingEXT)(VkDevice device, const VkPastPresentationTimingInfoEXT* pPastPresentationTimingInfo, VkPastPresentationTimingPropertiesEXT* pPastPresentationTimingProperties); + +#ifndef VK_NO_PROTOTYPES +#ifndef VK_ONLY_EXPORTED_PROTOTYPES +VKAPI_ATTR VkResult VKAPI_CALL vkSetSwapchainPresentTimingQueueSizeEXT( + VkDevice device, + VkSwapchainKHR swapchain, + uint32_t size); +#endif + +#ifndef VK_ONLY_EXPORTED_PROTOTYPES +VKAPI_ATTR VkResult VKAPI_CALL vkGetSwapchainTimingPropertiesEXT( + VkDevice device, + VkSwapchainKHR swapchain, + VkSwapchainTimingPropertiesEXT* pSwapchainTimingProperties, + uint64_t* pSwapchainTimingPropertiesCounter); +#endif + +#ifndef VK_ONLY_EXPORTED_PROTOTYPES +VKAPI_ATTR VkResult VKAPI_CALL vkGetSwapchainTimeDomainPropertiesEXT( + VkDevice device, + VkSwapchainKHR swapchain, + VkSwapchainTimeDomainPropertiesEXT* pSwapchainTimeDomainProperties, + uint64_t* pTimeDomainsCounter); +#endif + +#ifndef VK_ONLY_EXPORTED_PROTOTYPES +VKAPI_ATTR VkResult VKAPI_CALL vkGetPastPresentationTimingEXT( + VkDevice device, + const VkPastPresentationTimingInfoEXT* pPastPresentationTimingInfo, + VkPastPresentationTimingPropertiesEXT* pPastPresentationTimingProperties); +#endif +#endif + + // VK_INTEL_shader_integer_functions2 is a preprocessor guard. Do not pass it to API calls. #define VK_INTEL_shader_integer_functions2 1 #define VK_INTEL_SHADER_INTEGER_FUNCTIONS_2_SPEC_VERSION 1 @@ -21866,13 +22037,11 @@ typedef struct VkPhysicalDeviceRayTracingInvocationReorderPropertiesNV { VkRayTracingInvocationReorderModeEXT rayTracingInvocationReorderReorderingHint; } VkPhysicalDeviceRayTracingInvocationReorderPropertiesNV; -typedef struct VkPhysicalDeviceRayTracingInvocationReorderFeaturesEXT { +typedef struct VkPhysicalDeviceRayTracingInvocationReorderFeaturesNV { VkStructureType sType; void* pNext; VkBool32 rayTracingInvocationReorder; -} VkPhysicalDeviceRayTracingInvocationReorderFeaturesEXT; - -typedef VkPhysicalDeviceRayTracingInvocationReorderFeaturesEXT VkPhysicalDeviceRayTracingInvocationReorderFeaturesNV; +} VkPhysicalDeviceRayTracingInvocationReorderFeaturesNV; @@ -23657,6 +23826,12 @@ typedef struct VkPhysicalDeviceRayTracingInvocationReorderPropertiesEXT { uint32_t maxShaderBindingTableRecordIndex; } VkPhysicalDeviceRayTracingInvocationReorderPropertiesEXT; +typedef struct VkPhysicalDeviceRayTracingInvocationReorderFeaturesEXT { + VkStructureType sType; + void* pNext; + VkBool32 rayTracingInvocationReorder; +} VkPhysicalDeviceRayTracingInvocationReorderFeaturesEXT; + // VK_EXT_depth_clamp_control is a preprocessor guard. Do not pass it to API calls. @@ -24018,6 +24193,37 @@ typedef struct VkPhysicalDeviceShaderUniformBufferUnsizedArrayFeaturesEXT { +// VK_NV_compute_occupancy_priority is a preprocessor guard. Do not pass it to API calls. +#define VK_NV_compute_occupancy_priority 1 +#define VK_NV_COMPUTE_OCCUPANCY_PRIORITY_SPEC_VERSION 1 +#define VK_NV_COMPUTE_OCCUPANCY_PRIORITY_EXTENSION_NAME "VK_NV_compute_occupancy_priority" +#define VK_COMPUTE_OCCUPANCY_PRIORITY_LOW_NV 0.25f +#define VK_COMPUTE_OCCUPANCY_PRIORITY_NORMAL_NV 0.50f +#define VK_COMPUTE_OCCUPANCY_PRIORITY_HIGH_NV 0.75f +typedef struct VkComputeOccupancyPriorityParametersNV { + VkStructureType sType; + const void* pNext; + float occupancyPriority; + float occupancyThrottling; +} VkComputeOccupancyPriorityParametersNV; + +typedef struct VkPhysicalDeviceComputeOccupancyPriorityFeaturesNV { + VkStructureType sType; + void* pNext; + VkBool32 computeOccupancyPriority; +} VkPhysicalDeviceComputeOccupancyPriorityFeaturesNV; + +typedef void (VKAPI_PTR *PFN_vkCmdSetComputeOccupancyPriorityNV)(VkCommandBuffer commandBuffer, const VkComputeOccupancyPriorityParametersNV* pParameters); + +#ifndef VK_NO_PROTOTYPES +#ifndef VK_ONLY_EXPORTED_PROTOTYPES +VKAPI_ATTR void VKAPI_CALL vkCmdSetComputeOccupancyPriorityNV( + VkCommandBuffer commandBuffer, + const VkComputeOccupancyPriorityParametersNV* pParameters); +#endif +#endif + + // VK_KHR_acceleration_structure is a preprocessor guard. Do not pass it to API calls. #define VK_KHR_acceleration_structure 1 #define VK_KHR_ACCELERATION_STRUCTURE_SPEC_VERSION 13 diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_cross.cpp b/libs/bgfx/3rdparty/spirv-cross/spirv_cross.cpp index c99febe..2204e39 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_cross.cpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_cross.cpp @@ -742,6 +742,14 @@ bool Compiler::is_physical_pointer(const SPIRType &type) const return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer; } +bool Compiler::is_physical_or_buffer_pointer(const SPIRType &type) const +{ + return type.op == OpTypePointer && + (type.storage == StorageClassPhysicalStorageBuffer || type.storage == StorageClassUniform || + type.storage == StorageClassStorageBuffer || type.storage == StorageClassWorkgroup || + type.storage == StorageClassPushConstant); +} + bool Compiler::is_physical_pointer_to_buffer_block(const SPIRType &type) const { return is_physical_pointer(type) && get_pointee_type(type).self == type.parent_type && diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_cross.hpp b/libs/bgfx/3rdparty/spirv-cross/spirv_cross.hpp index 601ca26..06f2b73 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_cross.hpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_cross.hpp @@ -694,6 +694,7 @@ class Compiler bool is_array(const SPIRType &type) const; bool is_pointer(const SPIRType &type) const; bool is_physical_pointer(const SPIRType &type) const; + bool is_physical_or_buffer_pointer(const SPIRType &type) const; bool is_physical_pointer_to_buffer_block(const SPIRType &type) const; static bool is_runtime_size_array(const SPIRType &type); uint32_t expression_type_id(uint32_t id) const; diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.cpp b/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.cpp index 3cb08f2..f9ce6cb 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.cpp @@ -5051,10 +5051,10 @@ void CompilerGLSL::emit_polyfills(uint32_t polyfills, bool relaxed) // Returns a string representation of the ID, usable as a function arg. // Default is to simply return the expression representation fo the arg ID. // Subclasses may override to modify the return value. -string CompilerGLSL::to_func_call_arg(const SPIRFunction::Parameter &, uint32_t id) +string CompilerGLSL::to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id) { // BDA expects pointers through function interface. - if (is_physical_pointer(expression_type(id))) + if (!arg.alias_global_variable && is_physical_or_buffer_pointer(expression_type(id))) return to_pointer_expression(id); // Make sure that we use the name of the original variable, and not the parameter alias. @@ -6896,6 +6896,16 @@ void CompilerGLSL::emit_uninitialized_temporary(uint32_t result_type, uint32_t r } } +bool CompilerGLSL::can_declare_inline_temporary(uint32_t id) const +{ + if (!block_temporary_hoisting && current_continue_block && !hoisted_temporaries.count(id)) + return false; + if (hoisted_temporaries.count(id)) + return false; + + return true; +} + string CompilerGLSL::declare_temporary(uint32_t result_type, uint32_t result_id) { auto &type = get(result_type); @@ -6973,6 +6983,42 @@ SPIRExpression &CompilerGLSL::emit_op(uint32_t result_type, uint32_t result_id, } } +void CompilerGLSL::emit_transposed_op(uint32_t result_type, uint32_t result_id, const string &rhs, bool forwarding) +{ + if (forwarding && (forced_temporaries.find(result_id) == end(forced_temporaries))) + { + // Just forward it without temporary. + // If the forward is trivial, we do not force flushing to temporary for this expression. + forwarded_temporaries.insert(result_id); + auto &e = set(result_id, rhs, result_type, true); + e.need_transpose = true; + } + else if (can_declare_inline_temporary(result_id)) + { + // If expression isn't immutable, bind it to a temporary and make the new temporary immutable (they always are). + // Since the expression is transposed, we have to ensure the temporary is the transposed type. + + auto &transposed_type_id = extra_sub_expressions[result_id]; + if (!transposed_type_id) + { + auto dummy_type = get(result_type); + std::swap(dummy_type.columns, dummy_type.vecsize); + transposed_type_id = ir.increase_bound_by(1); + set(transposed_type_id, dummy_type); + } + + statement(declare_temporary(transposed_type_id, result_id), rhs, ";"); + auto &e = set(result_id, to_name(result_id), result_type, true); + e.need_transpose = true; + } + else + { + // If we cannot declare the temporary because it's already been hoisted, we don't have the + // chance to override the temporary type ourselves. Just transpose() the expression. + emit_op(result_type, result_id, join("transpose(", rhs, ")"), forwarding); + } +} + void CompilerGLSL::emit_unary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op) { bool forward = should_forward(op0); @@ -11581,7 +11627,7 @@ bool CompilerGLSL::should_dereference(uint32_t id) // If id is a variable but not a phi variable, we should not dereference it. // BDA passed around as parameters are always pointers. if (auto *var = maybe_get(id)) - return (var->parameter && is_physical_pointer(type)) || var->phi_variable; + return (var->parameter && is_physical_or_buffer_pointer(type)) || var->phi_variable; if (auto *expr = maybe_get(id)) { @@ -11617,8 +11663,8 @@ bool CompilerGLSL::should_dereference(uint32_t id) bool CompilerGLSL::should_dereference_caller_param(uint32_t id) { const auto &type = expression_type(id); - // BDA is always passed around as pointers. - if (is_physical_pointer(type)) + // BDA is always passed around as pointers. Similarly, we need to pass variable buffer pointers as pointers. + if (is_physical_or_buffer_pointer(type)) return false; return should_dereference(id); @@ -13507,8 +13553,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) auto expr = join(enclose_expression(to_unpacked_row_major_matrix_expression(ops[3])), " * ", enclose_expression(to_unpacked_row_major_matrix_expression(ops[2]))); bool forward = should_forward(ops[2]) && should_forward(ops[3]); - auto &e = emit_op(ops[0], ops[1], expr, forward); - e.need_transpose = true; + emit_transposed_op(ops[0], ops[1], expr, forward); a->need_transpose = true; b->need_transpose = true; inherit_expression_dependencies(ops[1], ops[2]); @@ -13531,8 +13576,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) auto expr = join(enclose_expression(to_unpacked_row_major_matrix_expression(ops[2])), " * ", to_enclosed_unpacked_expression(ops[3])); bool forward = should_forward(ops[2]) && should_forward(ops[3]); - auto &e = emit_op(ops[0], ops[1], expr, forward); - e.need_transpose = true; + emit_transposed_op(ops[0], ops[1], expr, forward); a->need_transpose = true; inherit_expression_dependencies(ops[1], ops[2]); inherit_expression_dependencies(ops[1], ops[3]); diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.hpp b/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.hpp index 78bff2d..98e93ee 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_glsl.hpp @@ -763,6 +763,7 @@ class CompilerGLSL : public Compiler bool expression_read_implies_multiple_reads(uint32_t id) const; SPIRExpression &emit_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs, bool suppress_usage_tracking = false); + void emit_transposed_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs); void access_chain_internal_append_index(std::string &expr, uint32_t base, const SPIRType *type, AccessChainFlags flags, bool &access_chain_is_arrayed, uint32_t index); @@ -805,6 +806,7 @@ class CompilerGLSL : public Compiler const char *index_to_swizzle(uint32_t index); std::string remap_swizzle(const SPIRType &result_type, uint32_t input_components, const std::string &expr); std::string declare_temporary(uint32_t type, uint32_t id); + bool can_declare_inline_temporary(uint32_t id) const; void emit_uninitialized_temporary(uint32_t type, uint32_t id); SPIRExpression &emit_uninitialized_temporary_expression(uint32_t type, uint32_t id); virtual void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector &arglist); diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.cpp b/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.cpp index 7394a09..098beca 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.cpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.cpp @@ -491,9 +491,9 @@ string CompilerHLSL::type_to_glsl(const SPIRType &type, uint32_t id) case SPIRType::Double: return join("double", type.vecsize); case SPIRType::Int64: - return join("i64vec", type.vecsize); + return join("int64_t", type.vecsize); case SPIRType::UInt64: - return join("u64vec", type.vecsize); + return join("uint64_t", type.vecsize); default: return "???"; } @@ -1595,6 +1595,7 @@ void CompilerHLSL::replace_illegal_names() "Texture3D", "TextureCube", "TextureCubeArray", "true", "typedef", "triangle", "triangleadj", "TriangleStream", "uint", "uniform", "unorm", "unsigned", "vector", "vertexfragment", "VertexShader", "vertices", "void", "volatile", "while", + "signed", }; CompilerGLSL::replace_illegal_names(keywords); @@ -1709,9 +1710,11 @@ void CompilerHLSL::emit_resources() ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { auto &type = this->get(var.basetype); + bool is_hidden = is_hidden_io_variable(var); + if (var.storage != StorageClassFunction && !var.remapped_variable && type.pointer && (var.storage == StorageClassInput || var.storage == StorageClassOutput) && !is_builtin_variable(var) && - interface_variable_exists_in_entry_point(var.self)) + interface_variable_exists_in_entry_point(var.self) && !is_hidden) { // Builtin variables are handled separately. emit_interface_block_globally(var); @@ -1747,8 +1750,10 @@ void CompilerHLSL::emit_resources() if (var.storage != StorageClassInput && var.storage != StorageClassOutput) return; + bool is_hidden = is_hidden_io_variable(var); + if (!var.remapped_variable && type.pointer && !is_builtin_variable(var) && - interface_variable_exists_in_entry_point(var.self)) + interface_variable_exists_in_entry_point(var.self) && !is_hidden) { if (block) { @@ -3482,10 +3487,12 @@ void CompilerHLSL::emit_hlsl_entry_point() if (var.storage != StorageClassInput) return; + bool is_hidden = is_hidden_io_variable(var); + bool need_matrix_unroll = var.storage == StorageClassInput && execution.model == ExecutionModelVertex; if (!var.remapped_variable && type.pointer && !is_builtin_variable(var) && - interface_variable_exists_in_entry_point(var.self)) + interface_variable_exists_in_entry_point(var.self) && !is_hidden) { if (block) { @@ -7119,6 +7126,30 @@ bool CompilerHLSL::is_hlsl_force_storage_buffer_as_uav(ID id) const return (force_uav_buffer_bindings.find({ desc_set, binding }) != force_uav_buffer_bindings.end()); } +bool CompilerHLSL::is_hidden_io_variable(const SPIRVariable &var) const +{ + if (!is_hidden_variable(var)) + return false; + + // It is too risky to remove stage IO variables that are linkable since it affects link compatibility. + // For vertex inputs and fragment outputs, it's less of a concern and we want reflection data + // to match reality. + + bool is_external_linkage = + (get_execution_model() == ExecutionModelVertex && var.storage == StorageClassInput) || + (get_execution_model() == ExecutionModelFragment && var.storage == StorageClassOutput); + + if (!is_external_linkage) + return false; + + // Unused output I/O variables might still be required to implement framebuffer fetch. + if (var.storage == StorageClassOutput && !is_legacy() && + location_is_framebuffer_fetch(get_decoration(var.self, DecorationLocation)) != 0) + return false; + + return true; +} + void CompilerHLSL::set_hlsl_force_storage_buffer_as_uav(uint32_t desc_set, uint32_t binding) { SetBindingPair pair = { desc_set, binding }; diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.hpp b/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.hpp index e4979db..5f9c313 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.hpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_hlsl.hpp @@ -298,6 +298,7 @@ class CompilerHLSL : public CompilerGLSL SPIRType::BaseType get_builtin_basetype(BuiltIn builtin, SPIRType::BaseType default_type) override; bool is_hlsl_force_storage_buffer_as_uav(ID id) const; + bool is_hidden_io_variable(const SPIRVariable &var) const; Options hlsl_options; diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_msl.cpp b/libs/bgfx/3rdparty/spirv-cross/spirv_msl.cpp index 1c07b52..2dc1487 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_msl.cpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_msl.cpp @@ -1514,7 +1514,7 @@ void CompilerMSL::emit_entry_point_declarations() if (is_array(type)) { is_using_builtin_array = true; - statement(get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, true), name, + statement(get_variable_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, true), name, type_to_array_glsl(type, var_id), " ="); uint32_t array_size = get_resource_array_size(type, var_id); @@ -1525,8 +1525,8 @@ void CompilerMSL::emit_entry_point_declarations() for (uint32_t i = 0; i < array_size; i++) { - statement("(", get_argument_address_space(var), " ", type_to_glsl(type), "* ", - to_restrict(var_id, false), ")((", get_argument_address_space(var), " char* ", + statement("(", get_variable_address_space(var), " ", type_to_glsl(type), "* ", + to_restrict(var_id, false), ")((", get_variable_address_space(var), " char* ", to_restrict(var_id, false), ")", to_name(arg_id), ".", dynamic_buffer.second.mbr_name, "[", i, "]", " + ", to_name(dynamic_offsets_buffer_id), "[", base_index + i, "]),"); } @@ -1537,9 +1537,9 @@ void CompilerMSL::emit_entry_point_declarations() } else { - statement(get_argument_address_space(var), " auto& ", to_restrict(var_id, true), name, " = *(", - get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, false), ")((", - get_argument_address_space(var), " char* ", to_restrict(var_id, false), ")", to_name(arg_id), ".", + statement(get_variable_address_space(var), " auto& ", to_restrict(var_id, true), name, " = *(", + get_variable_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, false), ")((", + get_variable_address_space(var), " char* ", to_restrict(var_id, false), ")", to_name(arg_id), ".", dynamic_buffer.second.mbr_name, " + ", to_name(dynamic_offsets_buffer_id), "[", base_index, "]);"); } } @@ -1594,7 +1594,7 @@ void CompilerMSL::emit_entry_point_declarations() statement("spvDescriptorArray ", name, "Smplr {", resource_name, "Smplr};"); break; case SPIRType::Struct: - statement("spvDescriptorArray<", get_argument_address_space(var), " ", type_to_glsl(buffer_type), "*> ", + statement("spvDescriptorArray<", get_variable_address_space(var), " ", type_to_glsl(buffer_type), "*> ", name, " {", resource_name, "};"); break; default: @@ -1605,7 +1605,7 @@ void CompilerMSL::emit_entry_point_declarations() else if (!type.array.empty() && type.basetype == SPIRType::Struct) { // Emit only buffer arrays here. - statement(get_argument_address_space(var), " ", type_to_glsl(buffer_type), "* ", + statement(get_variable_address_space(var), " ", type_to_glsl(buffer_type), "* ", to_restrict(var.self, true), name, "[] ="); begin_scope(); uint32_t array_size = get_resource_array_size(type, var.self); @@ -1629,7 +1629,7 @@ void CompilerMSL::emit_entry_point_declarations() continue; const auto &type = get_variable_data_type(var); - auto addr_space = get_argument_address_space(var); + auto addr_space = get_variable_address_space(var); // This resource name has already been added. auto name = to_name(var_id); @@ -10158,8 +10158,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) enclose_expression(to_unpacked_row_major_matrix_expression(ops[2])), ")"); bool forward = should_forward(ops[2]) && should_forward(ops[3]); - auto &e = emit_op(ops[0], ops[1], expr, forward); - e.need_transpose = true; + emit_transposed_op(ops[0], ops[1], expr, forward); a->need_transpose = true; b->need_transpose = true; inherit_expression_dependencies(ops[1], ops[2]); @@ -11141,7 +11140,7 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id, } else if (var && ptr_type.storage != StorageClassPhysicalStorageBuffer) { - exp += get_argument_address_space(*var); + exp += get_variable_address_space(*var); } else { @@ -14034,11 +14033,17 @@ bool CompilerMSL::uses_explicit_early_fragment_test() } // In MSL, address space qualifiers are required for all pointer or reference variables -string CompilerMSL::get_argument_address_space(const SPIRVariable &argument) +string CompilerMSL::get_variable_address_space(const SPIRVariable &argument) { const auto &type = get(argument.basetype); - // BDA is always passed around by value. There is no storage class for the argument itself. - if (is_physical_pointer(type)) + return get_type_address_space(type, argument.self, true); +} + +string CompilerMSL::get_leaf_argument_address_space(const SPIRVariable &argument) +{ + const auto &type = get(argument.basetype); + // BDA and variable buffer pointer is always passed around by (pointer) value. There is no storage class for the argument itself. + if (is_physical_or_buffer_pointer(type)) return ""; return get_type_address_space(type, argument.self, true); } @@ -14115,6 +14120,7 @@ string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id, bo } else if (!argument) { + // This is used for helper UBOs we insert ourselves. addr_space = "constant"; } else if (type_is_msl_framebuffer_fetch(type)) @@ -14122,6 +14128,7 @@ string CompilerMSL::get_type_address_space(const SPIRType &type, uint32_t id, bo // Subpass inputs are passed around by value. addr_space = ""; } + break; case StorageClassFunction: @@ -14638,7 +14645,7 @@ string CompilerMSL::entry_point_args_argument_buffer(bool append_comma) claimed_bindings.set(buffer_binding); - ep_args += get_argument_address_space(var) + " "; + ep_args += get_variable_address_space(var) + " "; if (recursive_inputs.count(type.self)) ep_args += string("void* ") + to_restrict(id, true) + to_name(id) + "_vp"; @@ -14852,7 +14859,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) // Declare the primary alias as void* if (!ep_args.empty()) ep_args += ", "; - ep_args += get_argument_address_space(var) + " void* " + primary_name; + ep_args += get_variable_address_space(var) + " void* " + primary_name; ep_args += " [[buffer(" + convert_to_string(r.index) + ")"; if (interlocked_resources.count(var_id)) ep_args += ", raster_order_group(0)"; @@ -14900,7 +14907,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) { if (!ep_args.empty()) ep_args += ", "; - ep_args += get_argument_address_space(var) + " " + type_to_glsl(type) + "* " + + ep_args += get_variable_address_space(var) + " " + type_to_glsl(type) + "* " + to_restrict(var_id, true) + r.name + "_" + convert_to_string(i); ep_args += " [[buffer(" + convert_to_string(r.index + i) + ")"; if (interlocked_resources.count(var_id)) @@ -14913,7 +14920,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) { if (!ep_args.empty()) ep_args += ", "; - ep_args += get_argument_address_space(var) + " "; + ep_args += get_variable_address_space(var) + " "; if (recursive_inputs.count(type.self)) ep_args += string("void* ") + to_restrict(var_id, true) + r.name + "_vp"; @@ -15105,7 +15112,7 @@ void CompilerMSL::fix_up_shader_inputs_outputs() { recursive_inputs.insert(type.self); entry_func.fixup_hooks_in.push_back([this, &type, &var, var_id]() { - auto addr_space = get_argument_address_space(var); + auto addr_space = get_variable_address_space(var); auto var_name = to_name(var_id); statement(addr_space, " auto& ", to_restrict(var_id, true), var_name, " = *(", addr_space, " ", type_to_glsl(type), "*)", var_name, "_vp;"); @@ -15802,7 +15809,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg) // Physical pointer types are passed by pointer, not reference. auto &data_type = get_variable_data_type(var); - bool passed_by_value = is_physical_pointer(var_type); + bool passed_by_value = arg.alias_global_variable ? false : is_physical_or_buffer_pointer(var_type); auto &type = passed_by_value ? var_type : data_type; // If we need to modify the name of the variable, make sure we use the original variable. @@ -15845,7 +15852,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg) spv_function_implementations.count(SPVFuncImplDynamicImageSampler); // Allow Metal to use the array template to make arrays a value type - string address_space = get_argument_address_space(var); + string address_space = arg.alias_global_variable ? get_variable_address_space(var) : get_leaf_argument_address_space(var); bool builtin = has_decoration(var.self, DecorationBuiltIn); auto builtin_type = BuiltIn(get_decoration(arg.id, DecorationBuiltIn)); @@ -15942,10 +15949,29 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg) } else { - if (!address_space.empty()) - decl = join(address_space, " ", decl); - decl += " "; - decl += to_expression(name_id); + // Variable pointer to array is kinda awkward ... + bool pointer_to_logical_buffer_array = + !is_physical_pointer(type) && is_pointer(type) && + has_decoration(type.parent_type, DecorationArrayStride); + + if (pointer_to_logical_buffer_array) + { + decl.pop_back(); + decl += " (*"; + decl += to_expression(name_id); + decl += ")"; + bool old_is_using_builtin_array = is_using_builtin_array; + is_using_builtin_array = true; + decl += type_to_array_glsl(type, name_id); + is_using_builtin_array = old_is_using_builtin_array; + } + else + { + if (!address_space.empty()) + decl = join(address_space, " ", decl); + decl += " "; + decl += to_expression(name_id); + } } } else if (is_array(type) && !type_is_image) @@ -16252,6 +16278,7 @@ const std::unordered_set &CompilerMSL::get_reserved_keyword_set() "quad_broadcast", "thread", "threadgroup", + "signed", }; return keywords; @@ -16395,6 +16422,7 @@ const std::unordered_set &CompilerMSL::get_illegal_func_names() "uint16", "float8", "float16", + "signed", }; return illegal_func_names; @@ -16563,8 +16591,11 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member) // the C-style nesting works right. // FIXME: This is somewhat of a hack. bool old_is_using_builtin_array = is_using_builtin_array; + bool pointer_to_buffer_array = is_pointer(type) && has_decoration(type.parent_type, DecorationArrayStride); if (is_physical_pointer(type)) is_using_builtin_array = false; + else if (pointer_to_buffer_array) + is_using_builtin_array = true; type_name = join(type_address_space, " ", type_to_glsl(*p_parent_type, id)); @@ -17930,7 +17961,10 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin) { // thread_execution_width is an alias for threads_per_simdgroup, and it's only available since 1.0, // but not in fragment. - return "thread_execution_width"; + if (msl_options.supports_msl_version(3, 0)) + return "threads_per_simdgroup"; + else + return "thread_execution_width"; } case BuiltInNumSubgroups: @@ -17960,6 +17994,7 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin) return "thread_index_in_simdgroup"; } else if (execution.model == ExecutionModelKernel || execution.model == ExecutionModelGLCompute || + execution.model == ExecutionModelTaskEXT || execution.model == ExecutionModelMeshEXT || execution.model == ExecutionModelTessellationControl || (execution.model == ExecutionModelVertex && msl_options.vertex_for_tessellation)) { @@ -19898,7 +19933,7 @@ void CompilerMSL::analyze_argument_buffers() { recursive_inputs.insert(type_id); auto &entry_func = this->get(ir.default_entry_point); - auto addr_space = get_argument_address_space(buffer_var); + auto addr_space = get_variable_address_space(buffer_var); entry_func.fixup_hooks_in.push_back([this, addr_space, buffer_name, buffer_type_name]() { statement(addr_space, " auto& ", buffer_name, " = *(", addr_space, " ", buffer_type_name, "*)", buffer_name, "_vp;"); }); diff --git a/libs/bgfx/3rdparty/spirv-cross/spirv_msl.hpp b/libs/bgfx/3rdparty/spirv-cross/spirv_msl.hpp index 75d3aa7..f63f5a2 100644 --- a/libs/bgfx/3rdparty/spirv-cross/spirv_msl.hpp +++ b/libs/bgfx/3rdparty/spirv-cross/spirv_msl.hpp @@ -1126,7 +1126,9 @@ class CompilerMSL : public CompilerGLSL void mark_struct_members_packed(const SPIRType &type); void ensure_member_packing_rules_msl(SPIRType &ib_type, uint32_t index); bool validate_member_packing_rules_msl(const SPIRType &type, uint32_t index) const; - std::string get_argument_address_space(const SPIRVariable &argument); + std::string get_variable_address_space(const SPIRVariable &argument); + // Special case of get_variable_address_space which is only used for leaf functions. + std::string get_leaf_argument_address_space(const SPIRVariable &argument); std::string get_type_address_space(const SPIRType &type, uint32_t id, bool argument = false); bool decoration_flags_signal_volatile(const Bitset &flags) const; bool decoration_flags_signal_coherent(const Bitset &flags) const; diff --git a/libs/bgfx/3rdparty/spirv-headers/include/spirv/spir-v.xml b/libs/bgfx/3rdparty/spirv-headers/include/spirv/spir-v.xml index ba45e3d..605bc47 100644 --- a/libs/bgfx/3rdparty/spirv-headers/include/spirv/spir-v.xml +++ b/libs/bgfx/3rdparty/spirv-headers/include/spirv/spir-v.xml @@ -140,13 +140,14 @@ + - + @@ -174,13 +175,14 @@ + - + diff --git a/libs/bgfx/3rdparty/spirv-tools/include/generated/build-version.inc b/libs/bgfx/3rdparty/spirv-tools/include/generated/build-version.inc index 81c26ac..3977361 100644 --- a/libs/bgfx/3rdparty/spirv-tools/include/generated/build-version.inc +++ b/libs/bgfx/3rdparty/spirv-tools/include/generated/build-version.inc @@ -1 +1 @@ -"v2025.5", "SPIRV-Tools v2025.5 v2025.4-64-gd2a11ec9" +"v2025.5", "SPIRV-Tools v2025.5 v2025.5.rc1-32-g6e7423bc" diff --git a/libs/bgfx/3rdparty/spirv-tools/source/diff/diff.cpp b/libs/bgfx/3rdparty/spirv-tools/source/diff/diff.cpp index 7fd21f3..d548aea 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/diff/diff.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/diff/diff.cpp @@ -1219,6 +1219,7 @@ bool Differ::DoDebugAndAnnotationInstructionsMatch( case spv::Op::OpMemberDecorate: return DoOperandsMatch(src_inst, dst_inst, 0, 3); case spv::Op::OpExtInst: + return DoOperandsMatch(src_inst, dst_inst, 0, 2); case spv::Op::OpDecorationGroup: case spv::Op::OpGroupDecorate: case spv::Op::OpGroupMemberDecorate: @@ -2612,6 +2613,9 @@ void Differ::MatchExtInstDebugInfo() { // This section includes OpExtInst for DebugInfo extension MatchDebugAndAnnotationInstructions(src_->ext_inst_debuginfo(), dst_->ext_inst_debuginfo()); + // OpExtInst can exist in other sections too, such as with non-semantic info. + MatchDebugAndAnnotationInstructions(src_->types_values(), + dst_->types_values()); } void Differ::MatchAnnotations() { diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/aggressive_dead_code_elim_pass.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/aggressive_dead_code_elim_pass.cpp index 51dc68e..54b1002 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/aggressive_dead_code_elim_pass.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/aggressive_dead_code_elim_pass.cpp @@ -44,9 +44,10 @@ constexpr uint32_t kExtInstSetInIdx = 0; constexpr uint32_t kExtInstOpInIdx = 1; constexpr uint32_t kInterpolantInIdx = 2; constexpr uint32_t kCooperativeMatrixLoadSourceAddrInIdx = 0; -constexpr uint32_t kDebugValueLocalVariable = 2; -constexpr uint32_t kDebugValueValue = 3; -constexpr uint32_t kDebugValueExpression = 4; +constexpr uint32_t kDebugDeclareVariableInIdx = 3; +constexpr uint32_t kDebugValueLocalVariableInIdx = 2; +constexpr uint32_t kDebugValueValueInIdx = 3; +constexpr uint32_t kDebugValueExpressionInIdx = 4; // Sorting functor to present annotation instructions in an easy-to-process // order. The functor orders by opcode first and falls back on unique id @@ -290,39 +291,94 @@ Pass::Status AggressiveDCEPass::ProcessDebugInformation( std::list& structured_order) { for (auto bi = structured_order.begin(); bi != structured_order.end(); bi++) { bool succeeded = (*bi)->WhileEachInst([this](Instruction* inst) { - // DebugDeclare is not dead. It must be converted to DebugValue in a - // later pass - if (inst->IsNonSemanticInstruction() && - inst->GetShader100DebugOpcode() == - NonSemanticShaderDebugInfo100DebugDeclare) { - AddToWorklist(inst); - return true; - } + if (!inst->IsNonSemanticInstruction()) return true; - // If the Value of a DebugValue is killed, set Value operand to Undef - if (inst->IsNonSemanticInstruction() && - inst->GetShader100DebugOpcode() == - NonSemanticShaderDebugInfo100DebugValue) { - uint32_t id = inst->GetSingleWordInOperand(kDebugValueValue); - auto def = get_def_use_mgr()->GetDef(id); - if (!IsLive(def)) { + if (inst->GetShader100DebugOpcode() == + NonSemanticShaderDebugInfo100DebugDeclare) { + if (IsLive(inst)) return true; + + uint32_t var_id = + inst->GetSingleWordInOperand(kDebugDeclareVariableInIdx); + auto var_def = get_def_use_mgr()->GetDef(var_id); + + if (IsLive(var_def)) { AddToWorklist(inst); - uint32_t undef_id = Type2Undef(def->type_id()); - if (undef_id == 0) { - return false; + return true; + } + + // DebugDeclare Variable is not live. Find the value that was being + // stored to this variable. If it's live then create a new DebugValue + // with this value. Otherwise let it die in peace. + get_def_use_mgr()->ForEachUser(var_id, [this, var_id, + inst](Instruction* user) { + if (user->opcode() == spv::Op::OpStore) { + uint32_t stored_value_id = 0; + const uint32_t kStoreValueInIdx = 1; + stored_value_id = user->GetSingleWordInOperand(kStoreValueInIdx); + if (!IsLive(get_def_use_mgr()->GetDef(stored_value_id))) { + return true; + } + + // value being stored is still live + Instruction* next_inst = inst->NextNode(); + bool added = + context()->get_debug_info_mgr()->AddDebugValueForVariable( + user, var_id, stored_value_id, inst); + if (added && next_inst) { + auto new_debug_value = next_inst->PreviousNode(); + live_insts_.Set(new_debug_value->unique_id()); + } } - inst->SetInOperand(kDebugValueValue, {undef_id}); - context()->get_def_use_mgr()->UpdateDefUse(inst); - id = inst->GetSingleWordInOperand(kDebugValueLocalVariable); - auto localVar = get_def_use_mgr()->GetDef(id); - AddToWorklist(localVar); - context()->get_def_use_mgr()->UpdateDefUse(localVar); - AddOperandsToWorkList(localVar); - id = inst->GetSingleWordInOperand(kDebugValueExpression); - auto expression = get_def_use_mgr()->GetDef(id); - AddToWorklist(expression); - context()->get_def_use_mgr()->UpdateDefUse(expression); return true; + }); + } else if (inst->GetShader100DebugOpcode() == + NonSemanticShaderDebugInfo100DebugValue) { + uint32_t var_operand_idx = kDebugValueValueInIdx; + uint32_t id = inst->GetSingleWordInOperand(var_operand_idx); + auto def = get_def_use_mgr()->GetDef(id); + + if (IsLive(def)) { + AddToWorklist(inst); + return true; + } + + // Value operand of DebugValue is not live + // Set Value to Undef of appropriate type + live_insts_.Set(inst->unique_id()); + + uint32_t type_id = def->type_id(); + auto type_def = get_def_use_mgr()->GetDef(type_id); + AddToWorklist(type_def); + + uint32_t undef_id = Type2Undef(type_id); + if (undef_id == 0) return false; + + auto undef_inst = get_def_use_mgr()->GetDef(undef_id); + live_insts_.Set(undef_inst->unique_id()); + inst->SetInOperand(var_operand_idx, {undef_id}); + context()->get_def_use_mgr()->AnalyzeInstUse(inst); + + id = inst->GetSingleWordInOperand(kDebugValueLocalVariableInIdx); + auto localVar = get_def_use_mgr()->GetDef(id); + AddToWorklist(localVar); + + uint32_t expr_idx = kDebugValueExpressionInIdx; + id = inst->GetSingleWordInOperand(expr_idx); + auto expression = get_def_use_mgr()->GetDef(id); + AddToWorklist(expression); + + for (uint32_t i = expr_idx + 1; i < inst->NumInOperands(); ++i) { + id = inst->GetSingleWordInOperand(i); + auto index_def = get_def_use_mgr()->GetDef(id); + if (index_def) { + AddToWorklist(index_def); + } + } + + for (auto& line_inst : inst->dbg_line_insts()) { + if (line_inst.IsDebugLineInst()) { + AddToWorklist(&line_inst); + } } } return true; @@ -731,13 +787,16 @@ Pass::Status AggressiveDCEPass::InitializeModuleScopeLiveInstructions() { AddToWorklist(dbg_none); } - // Add top level DebugInfo to worklist + // Add DebugInfo which should never be eliminated to worklist for (auto& dbg : get_module()->ext_inst_debuginfo()) { auto op = dbg.GetShader100DebugOpcode(); if (op == NonSemanticShaderDebugInfo100DebugCompilationUnit || op == NonSemanticShaderDebugInfo100DebugEntryPoint || op == NonSemanticShaderDebugInfo100DebugSource || - op == NonSemanticShaderDebugInfo100DebugSourceContinued) { + op == NonSemanticShaderDebugInfo100DebugSourceContinued || + op == NonSemanticShaderDebugInfo100DebugLocalVariable || + op == NonSemanticShaderDebugInfo100DebugExpression || + op == NonSemanticShaderDebugInfo100DebugOperation) { AddToWorklist(&dbg); } } @@ -813,7 +872,9 @@ Pass::Status AggressiveDCEPass::ProcessImpl() { // Cleanup all CFG including all unreachable blocks. for (Function& fp : *context()->module()) { - modified |= CFGCleanup(&fp); + auto status = CFGCleanup(&fp); + if (status == Status::Failure) return Status::Failure; + if (status == Status::SuccessWithChange) modified = true; } return modified ? Status::SuccessWithChange : Status::SuccessWithoutChange; diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/cfg_cleanup_pass.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/cfg_cleanup_pass.cpp index 26fed89..6cd0479 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/cfg_cleanup_pass.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/cfg_cleanup_pass.cpp @@ -25,8 +25,17 @@ namespace opt { Pass::Status CFGCleanupPass::Process() { // Process all entry point functions. - ProcessFunction pfn = [this](Function* fp) { return CFGCleanup(fp); }; + bool failure = false; + ProcessFunction pfn = [this, &failure](Function* fp) { + auto status = CFGCleanup(fp); + if (status == Status::Failure) { + failure = true; + return false; + } + return status == Status::SuccessWithChange; + }; bool modified = context()->ProcessReachableCallTree(pfn); + if (failure) return Pass::Status::Failure; return modified ? Pass::Status::SuccessWithChange : Pass::Status::SuccessWithoutChange; } diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.cpp index 734c967..ec90d97 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.cpp @@ -27,36 +27,48 @@ Pass::Status CombineAccessChains::Process() { bool modified = false; for (auto& function : *get_module()) { - modified |= ProcessFunction(function); + auto status = ProcessFunction(function); + if (status == Status::Failure) return Status::Failure; + if (status == Status::SuccessWithChange) modified = true; } return (modified ? Status::SuccessWithChange : Status::SuccessWithoutChange); } -bool CombineAccessChains::ProcessFunction(Function& function) { +Pass::Status CombineAccessChains::ProcessFunction(Function& function) { if (function.IsDeclaration()) { - return false; + return Status::SuccessWithoutChange; } bool modified = false; + bool failure = false; cfg()->ForEachBlockInReversePostOrder( - function.entry().get(), [&modified, this](BasicBlock* block) { - block->ForEachInst([&modified, this](Instruction* inst) { + function.entry().get(), [&modified, &failure, this](BasicBlock* block) { + if (failure) return; + block->ForEachInst([&modified, &failure, this](Instruction* inst) { + if (failure) return; switch (inst->opcode()) { case spv::Op::OpAccessChain: case spv::Op::OpInBoundsAccessChain: case spv::Op::OpPtrAccessChain: - case spv::Op::OpInBoundsPtrAccessChain: - modified |= CombineAccessChain(inst); + case spv::Op::OpInBoundsPtrAccessChain: { + auto status = CombineAccessChain(inst); + if (status == Status::Failure) { + failure = true; + } else if (status == Status::SuccessWithChange) { + modified = true; + } break; + } default: break; } }); }); - return modified; + if (failure) return Status::Failure; + return modified ? Status::SuccessWithChange : Status::SuccessWithoutChange; } uint32_t CombineAccessChains::GetConstantValue( @@ -121,9 +133,9 @@ const analysis::Type* CombineAccessChains::GetIndexedType(Instruction* inst) { return type; } -bool CombineAccessChains::CombineIndices(Instruction* ptr_input, - Instruction* inst, - std::vector* new_operands) { +Pass::Status CombineAccessChains::CombineIndices( + Instruction* ptr_input, Instruction* inst, + std::vector* new_operands) { analysis::DefUseManager* def_use_mgr = context()->get_def_use_mgr(); analysis::ConstantManager* constant_mgr = context()->get_constant_mgr(); @@ -150,28 +162,30 @@ bool CombineAccessChains::CombineIndices(Instruction* ptr_input, GetConstantValue(element_constant); const analysis::Constant* new_value_constant = constant_mgr->GetConstant(last_index_constant->type(), {new_value}); + if (!new_value_constant) return Status::Failure; Instruction* new_value_inst = constant_mgr->GetDefiningInstruction(new_value_constant); + if (!new_value_inst) return Status::Failure; new_value_id = new_value_inst->result_id(); } else if (!type->AsStruct() || combining_element_operands) { // Generate an addition of the two indices. InstructionBuilder builder( context(), inst, IRContext::kAnalysisDefUse | IRContext::kAnalysisInstrToBlockMapping); - // TODO(1841): Handle id overflow. Instruction* addition = builder.AddIAdd(last_index_inst->type_id(), last_index_inst->result_id(), element_inst->result_id()); + if (!addition) return Status::Failure; new_value_id = addition->result_id(); } else { // Indexing into structs must be constant, so bail out here. - return false; + return Status::SuccessWithoutChange; } new_operands->push_back({SPV_OPERAND_TYPE_ID, {new_value_id}}); - return true; + return Status::SuccessWithChange; } -bool CombineAccessChains::CreateNewInputOperands( +Pass::Status CombineAccessChains::CreateNewInputOperands( Instruction* ptr_input, Instruction* inst, std::vector* new_operands) { // Start by copying all the input operands of the feeder access chain. @@ -183,7 +197,8 @@ bool CombineAccessChains::CreateNewInputOperands( if (IsPtrAccessChain(inst->opcode())) { // The last index of the feeder should be combined with the element operand // of |inst|. - if (!CombineIndices(ptr_input, inst, new_operands)) return false; + auto status = CombineIndices(ptr_input, inst, new_operands); + if (status != Status::SuccessWithChange) return status; } else { // The indices aren't being combined so now add the last index operand of // |ptr_input|. @@ -197,10 +212,10 @@ bool CombineAccessChains::CreateNewInputOperands( new_operands->push_back(inst->GetInOperand(i)); } - return true; + return Status::SuccessWithChange; } -bool CombineAccessChains::CombineAccessChain(Instruction* inst) { +Pass::Status CombineAccessChains::CombineAccessChain(Instruction* inst) { assert((inst->opcode() == spv::Op::OpPtrAccessChain || inst->opcode() == spv::Op::OpAccessChain || inst->opcode() == spv::Op::OpInBoundsAccessChain || @@ -213,10 +228,11 @@ bool CombineAccessChains::CombineAccessChain(Instruction* inst) { ptr_input->opcode() != spv::Op::OpInBoundsAccessChain && ptr_input->opcode() != spv::Op::OpPtrAccessChain && ptr_input->opcode() != spv::Op::OpInBoundsPtrAccessChain) { - return false; + return Status::SuccessWithoutChange; } - if (Has64BitIndices(inst) || Has64BitIndices(ptr_input)) return false; + if (Has64BitIndices(inst) || Has64BitIndices(ptr_input)) + return Status::SuccessWithoutChange; // Handles the following cases: // 1. |ptr_input| is an index-less access chain. Replace the pointer @@ -238,7 +254,7 @@ bool CombineAccessChains::CombineAccessChain(Instruction* inst) { // size/alignment of the type and converting the stride into an element // index. uint32_t array_stride = GetArrayStride(ptr_input); - if (array_stride != 0) return false; + if (array_stride != 0) return Status::SuccessWithoutChange; if (ptr_input->NumInOperands() == 1) { // The input is effectively a no-op. @@ -250,14 +266,15 @@ bool CombineAccessChains::CombineAccessChain(Instruction* inst) { inst->SetOpcode(spv::Op::OpCopyObject); } else { std::vector new_operands; - if (!CreateNewInputOperands(ptr_input, inst, &new_operands)) return false; + auto status = CreateNewInputOperands(ptr_input, inst, &new_operands); + if (status != Status::SuccessWithChange) return status; // Update the instruction. inst->SetOpcode(UpdateOpcode(inst->opcode(), ptr_input->opcode())); inst->SetInOperands(std::move(new_operands)); context()->AnalyzeUses(inst); } - return true; + return Status::SuccessWithChange; } spv::Op CombineAccessChains::UpdateOpcode(spv::Op base_opcode, diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.h b/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.h index 32ee50d..1872720 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.h +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/combine_access_chains.h @@ -40,12 +40,12 @@ class CombineAccessChains : public Pass { private: // Combine access chains in |function|. Blocks are processed in reverse // post-order. Returns true if the function is modified. - bool ProcessFunction(Function& function); + Status ProcessFunction(Function& function); // Combines an access chain (normal, in bounds or pointer) |inst| if its base // pointer is another access chain. Returns true if the access chain was // modified. - bool CombineAccessChain(Instruction* inst); + Status CombineAccessChain(Instruction* inst); // Returns the value of |constant_inst| as a uint32_t. uint32_t GetConstantValue(const analysis::Constant* constant_inst); @@ -59,13 +59,13 @@ class CombineAccessChains : public Pass { // Populates |new_operands| with the operands for the combined access chain. // Returns false if the access chains cannot be combined. - bool CreateNewInputOperands(Instruction* ptr_input, Instruction* inst, - std::vector* new_operands); + Status CreateNewInputOperands(Instruction* ptr_input, Instruction* inst, + std::vector* new_operands); // Combines the last index of |ptr_input| with the element operand of |inst|. // Adds the combined operand to |new_operands|. - bool CombineIndices(Instruction* ptr_input, Instruction* inst, - std::vector* new_operands); + Status CombineIndices(Instruction* ptr_input, Instruction* inst, + std::vector* new_operands); // Returns the opcode to use for the combined access chain. spv::Op UpdateOpcode(spv::Op base_opcode, spv::Op input_opcode); diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/const_folding_rules.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/const_folding_rules.cpp index 0f4e440..b7a69bc 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/const_folding_rules.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/const_folding_rules.cpp @@ -1126,6 +1126,26 @@ ConstantFoldingRule FoldFUnordGreaterThanEqual() { return FoldFPBinaryOp(FOLD_FPCMP_OP(>=, false)); } +ConstantFoldingRule FoldInvariantSelect() { + return [](IRContext*, Instruction* inst, + const std::vector& constants) + -> const analysis::Constant* { + assert(inst->opcode() == spv::Op::OpSelect); + (void)inst; + + if (!constants[1] || !constants[2]) { + return nullptr; + } + if (constants[1] == constants[2]) { + return constants[1]; + } + if (constants[1]->IsZero() && constants[2]->IsZero()) { + return constants[1]; + } + return nullptr; + }; +} + // Folds an OpDot where all of the inputs are constants to a // constant. A new constant is created if necessary. ConstantFoldingRule FoldOpDotWithConstants() { @@ -1435,6 +1455,18 @@ ConstantFoldingRule FoldFMix() { }; } +template +static bool NegZeroAwareLessThan(FloatType a, FloatType b) { + if (a == 0.0 && b == 0.0) { + bool sba = std::signbit(a); + bool sbb = std::signbit(b); + if (sba && !sbb) { + return true; + } + } + return a < b; +} + const analysis::Constant* FoldMin(const analysis::Type* result_type, const analysis::Constant* a, const analysis::Constant* b, @@ -1480,11 +1512,11 @@ const analysis::Constant* FoldMin(const analysis::Type* result_type, if (float_type->width() == 32) { float va = a->GetFloat(); float vb = b->GetFloat(); - return (va < vb ? a : b); + return NegZeroAwareLessThan(va, vb) ? a : b; } else if (float_type->width() == 64) { double va = a->GetDouble(); double vb = b->GetDouble(); - return (va < vb ? a : b); + return NegZeroAwareLessThan(va, vb) ? a : b; } } return nullptr; @@ -1535,11 +1567,71 @@ const analysis::Constant* FoldMax(const analysis::Type* result_type, if (float_type->width() == 32) { float va = a->GetFloat(); float vb = b->GetFloat(); - return (va > vb ? a : b); + return NegZeroAwareLessThan(vb, va) ? a : b; } else if (float_type->width() == 64) { double va = a->GetDouble(); double vb = b->GetDouble(); - return (va > vb ? a : b); + return NegZeroAwareLessThan(vb, va) ? a : b; + } + } + return nullptr; +} + +const analysis::Constant* FoldNMin(const analysis::Type* result_type, + const analysis::Constant* a, + const analysis::Constant* b, + analysis::ConstantManager*) { + if (const analysis::Float* float_type = result_type->AsFloat()) { + if (float_type->width() == 32) { + float va = a->GetFloat(); + float vb = b->GetFloat(); + if (std::isnan(va)) { + return b; + } + if (std::isnan(vb)) { + return a; + } + return NegZeroAwareLessThan(va, vb) ? a : b; + } else if (float_type->width() == 64) { + double va = a->GetDouble(); + double vb = b->GetDouble(); + if (std::isnan(va)) { + return b; + } + if (std::isnan(vb)) { + return a; + } + return NegZeroAwareLessThan(va, vb) ? a : b; + } + } + return nullptr; +} + +const analysis::Constant* FoldNMax(const analysis::Type* result_type, + const analysis::Constant* a, + const analysis::Constant* b, + analysis::ConstantManager*) { + if (const analysis::Float* float_type = result_type->AsFloat()) { + if (float_type->width() == 32) { + float va = a->GetFloat(); + float vb = b->GetFloat(); + if (std::isnan(va)) { + return b; + } + if (std::isnan(vb)) { + return a; + } + return NegZeroAwareLessThan(vb, va) ? a : b; + } else if (float_type->width() == 64) { + double va = a->GetDouble(); + double vb = b->GetDouble(); + if (std::isnan(va)) { + return b; + } + if (std::isnan(vb)) { + return a; + } + return NegZeroAwareLessThan(vb, va) ? a : b; } } return nullptr; @@ -1627,6 +1719,88 @@ const analysis::Constant* FoldClamp3( return nullptr; } +// Fold an clamp instruction when all three operands are constant. +const analysis::Constant* FoldNClamp1( + IRContext* context, Instruction* inst, + const std::vector& constants) { + assert(inst->opcode() == spv::Op::OpExtInst && + "Expecting an extended instruction."); + assert(inst->GetSingleWordInOperand(0) == + context->get_feature_mgr()->GetExtInstImportId_GLSLstd450() && + "Expecting a GLSLstd450 extended instruction."); + + // Make sure all Clamp operands are constants. + for (uint32_t i = 1; i < 4; i++) { + if (constants[i] == nullptr) { + return nullptr; + } + } + + const analysis::Constant* temp = FoldFPBinaryOp( + FoldNMax, inst->type_id(), {constants[1], constants[2]}, context); + if (temp == nullptr) { + return nullptr; + } + return FoldFPBinaryOp(FoldNMin, inst->type_id(), {temp, constants[3]}, + context); +} + +// Fold a clamp instruction when |x <= min_val|. +const analysis::Constant* FoldNClamp2( + IRContext* context, Instruction* inst, + const std::vector& constants) { + assert(inst->opcode() == spv::Op::OpExtInst && + "Expecting an extended instruction."); + assert(inst->GetSingleWordInOperand(0) == + context->get_feature_mgr()->GetExtInstImportId_GLSLstd450() && + "Expecting a GLSLstd450 extended instruction."); + + const analysis::Constant* x = constants[1]; + const analysis::Constant* min_val = constants[2]; + + if (x == nullptr || min_val == nullptr) { + return nullptr; + } + + const analysis::Constant* temp = + FoldFPBinaryOp(FoldNMax, inst->type_id(), {x, min_val}, context); + if (temp == min_val) { + // We can assume that |min_val| is less than |max_val|. Therefore, if the + // result of the max operation is |min_val|, we know the result of the min + // operation, even if |max_val| is not a constant. + return min_val; + } + return nullptr; +} + +// Fold a clamp instruction when |x >= max_val|. +const analysis::Constant* FoldNClamp3( + IRContext* context, Instruction* inst, + const std::vector& constants) { + assert(inst->opcode() == spv::Op::OpExtInst && + "Expecting an extended instruction."); + assert(inst->GetSingleWordInOperand(0) == + context->get_feature_mgr()->GetExtInstImportId_GLSLstd450() && + "Expecting a GLSLstd450 extended instruction."); + + const analysis::Constant* x = constants[1]; + const analysis::Constant* max_val = constants[3]; + + if (x == nullptr || max_val == nullptr) { + return nullptr; + } + + const analysis::Constant* temp = + FoldFPBinaryOp(FoldNMin, inst->type_id(), {x, max_val}, context); + if (temp == max_val) { + // We can assume that |min_val| is less than |max_val|. Therefore, if the + // result of the max operation is |min_val|, we know the result of the min + // operation, even if |max_val| is not a constant. + return max_val; + } + return nullptr; +} + UnaryScalarFoldingRule FoldFTranscendentalUnary(double (*fp)(double)) { return [fp](const analysis::Type* result_type, const analysis::Constant* a, @@ -1775,6 +1949,8 @@ void ConstantFoldingRules::AddFoldingRules() { rules_[spv::Op::OpFMul].push_back(FoldFMul()); rules_[spv::Op::OpFSub].push_back(FoldFSub()); + rules_[spv::Op::OpSelect].push_back(FoldInvariantSelect()); + rules_[spv::Op::OpFOrdEqual].push_back(FoldFOrdEqual()); rules_[spv::Op::OpFUnordEqual].push_back(FoldFUnordEqual()); @@ -1878,12 +2054,16 @@ void ConstantFoldingRules::AddFoldingRules() { FoldFPBinaryOp(FoldMin)); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450FMin}].push_back( FoldFPBinaryOp(FoldMin)); + ext_rules_[{ext_inst_glslstd450_id, GLSLstd450NMin}].push_back( + FoldFPBinaryOp(FoldNMin)); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450SMax}].push_back( FoldFPBinaryOp(FoldMax)); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450UMax}].push_back( FoldFPBinaryOp(FoldMax)); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450FMax}].push_back( FoldFPBinaryOp(FoldMax)); + ext_rules_[{ext_inst_glslstd450_id, GLSLstd450NMax}].push_back( + FoldFPBinaryOp(FoldNMax)); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450UClamp}].push_back( FoldClamp1); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450UClamp}].push_back( @@ -1902,6 +2082,12 @@ void ConstantFoldingRules::AddFoldingRules() { FoldClamp2); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450FClamp}].push_back( FoldClamp3); + ext_rules_[{ext_inst_glslstd450_id, GLSLstd450NClamp}].push_back( + FoldNClamp1); + ext_rules_[{ext_inst_glslstd450_id, GLSLstd450NClamp}].push_back( + FoldNClamp2); + ext_rules_[{ext_inst_glslstd450_id, GLSLstd450NClamp}].push_back( + FoldNClamp3); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450Sin}].push_back( FoldFPUnaryOp(FoldFTranscendentalUnary(std::sin))); ext_rules_[{ext_inst_glslstd450_id, GLSLstd450Cos}].push_back( diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.cpp index 3078a7c..547a5e4 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.cpp @@ -104,10 +104,17 @@ Pass::Status CopyPropagateArrays::Process() { continue; } - if (CanUpdateUses(&*var_inst, source_object->GetPointerTypeId(this))) { + uint32_t pointer_type_id = source_object->GetPointerTypeId(this); + if (pointer_type_id == 0) { + return Status::Failure; + } + + if (CanUpdateUses(&*var_inst, pointer_type_id)) { modified = true; - PropagateObject(&*var_inst, source_object.get(), store_inst); + if (!PropagateObject(&*var_inst, source_object.get(), store_inst)) { + return Status::Failure; + } } } @@ -170,15 +177,16 @@ Instruction* CopyPropagateArrays::FindStoreInstruction( return store_inst; } -void CopyPropagateArrays::PropagateObject(Instruction* var_inst, +bool CopyPropagateArrays::PropagateObject(Instruction* var_inst, MemoryObject* source, Instruction* insertion_point) { assert(var_inst->opcode() == spv::Op::OpVariable && "This function propagates variables."); Instruction* new_access_chain = BuildNewAccessChain(insertion_point, source); + if (!new_access_chain) return false; context()->KillNamesAndDecorates(var_inst); - UpdateUses(var_inst, new_access_chain); + return UpdateUses(var_inst, new_access_chain); } Instruction* CopyPropagateArrays::BuildNewAccessChain( @@ -192,7 +200,7 @@ Instruction* CopyPropagateArrays::BuildNewAccessChain( return source->GetVariable(); } - source->BuildConstants(); + if (!source->BuildConstants()) return nullptr; std::vector access_ids(source->AccessChain().size()); std::transform( source->AccessChain().cbegin(), source->AccessChain().cend(), @@ -642,7 +650,7 @@ bool CopyPropagateArrays::CanUpdateUses(Instruction* original_ptr_inst, }); } -void CopyPropagateArrays::UpdateUses(Instruction* original_ptr_inst, +bool CopyPropagateArrays::UpdateUses(Instruction* original_ptr_inst, Instruction* new_ptr_inst) { analysis::TypeManager* type_mgr = context()->get_type_mgr(); analysis::ConstantManager* const_mgr = context()->get_constant_mgr(); @@ -699,6 +707,7 @@ void CopyPropagateArrays::UpdateUses(Instruction* original_ptr_inst, def_use_mgr->GetDef(use->GetSingleWordOperand(index + 1)); auto* deref_expr_instr = context()->get_debug_info_mgr()->DerefDebugExpression(dbg_expr); + if (!deref_expr_instr) return false; use->SetOperand(index + 1, {deref_expr_instr->result_id()}); context()->AnalyzeUses(deref_expr_instr); @@ -783,6 +792,8 @@ void CopyPropagateArrays::UpdateUses(Instruction* original_ptr_inst, uint32_t new_pointer_type_id = type_mgr->FindPointerToType(new_pointee_type_id, storage_class); + if (new_pointer_type_id == 0) return false; + if (new_pointer_type_id != use->type_id()) { use->SetResultType(new_pointer_type_id); context()->AnalyzeUses(use); @@ -829,8 +840,7 @@ void CopyPropagateArrays::UpdateUses(Instruction* original_ptr_inst, uint32_t pointee_type_id = pointer_type->GetSingleWordInOperand(kTypePointerPointeeInIdx); uint32_t copy = GenerateCopy(original_ptr_inst, pointee_type_id, use); - assert(copy != 0 && - "Should not be updating uses unless we know it can be done."); + if (copy == 0) return false; context()->ForgetUses(use); use->SetInOperand(index, {copy}); @@ -852,6 +862,7 @@ void CopyPropagateArrays::UpdateUses(Instruction* original_ptr_inst, break; } } + return true; } uint32_t CopyPropagateArrays::GetMemberTypeId( @@ -955,7 +966,7 @@ bool CopyPropagateArrays::MemoryObject::Contains( return true; } -void CopyPropagateArrays::MemoryObject::BuildConstants() { +bool CopyPropagateArrays::MemoryObject::BuildConstants() { for (auto& entry : access_chain_) { if (entry.is_result_id) { continue; @@ -968,10 +979,13 @@ void CopyPropagateArrays::MemoryObject::BuildConstants() { analysis::ConstantManager* const_mgr = context->get_constant_mgr(); const analysis::Constant* index_const = const_mgr->GetConstant(uint32_type, {entry.immediate}); - entry.result_id = - const_mgr->GetDefiningInstruction(index_const)->result_id(); + if (!index_const) return false; + Instruction* constant_inst = const_mgr->GetDefiningInstruction(index_const); + if (!constant_inst) return false; + entry.result_id = constant_inst->result_id(); entry.is_result_id = true; } + return true; } } // namespace opt diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.h b/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.h index bf4bfb5..cb04a14 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.h +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/copy_prop_arrays.h @@ -118,7 +118,8 @@ class CopyPropagateArrays : public MemPass { // Converts all immediate values in the AccessChain their OpConstant // equivalent. - void BuildConstants(); + // Returns false if the constants could not be created. + bool BuildConstants(); // Returns the type id of the pointer type that can be used to point to this // memory object. @@ -175,7 +176,8 @@ class CopyPropagateArrays : public MemPass { // Replaces all loads of |var_inst| with a load from |source| instead. // |insertion_pos| is a position where it is possible to construct the // address of |source| and also dominates all of the loads of |var_inst|. - void PropagateObject(Instruction* var_inst, MemoryObject* source, + // Returns false if the propagation failed. + bool PropagateObject(Instruction* var_inst, MemoryObject* source, Instruction* insertion_pos); // Returns true if all of the references to |ptr_inst| can be rewritten and @@ -241,7 +243,7 @@ class CopyPropagateArrays : public MemPass { // types of other instructions as needed. This function should not be called // if |CanUpdateUses(original_ptr_inst, new_pointer_inst->type_id())| returns // false. - void UpdateUses(Instruction* original_ptr_inst, + bool UpdateUses(Instruction* original_ptr_inst, Instruction* new_pointer_inst); // Return true if |UpdateUses| is able to change all of the uses of diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/debug_info_manager.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/debug_info_manager.cpp index c084a6c..4570113 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/debug_info_manager.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/debug_info_manager.cpp @@ -331,6 +331,7 @@ Instruction* DebugInfoManager::GetDebugOperationWithDeref() { if (deref_operation_ != nullptr) return deref_operation_; uint32_t result_id = context()->TakeNextId(); + if (result_id == 0) return nullptr; std::unique_ptr deref_operation; if (context()->get_feature_mgr()->GetExtInstImportId_OpenCL100DebugInfo()) { @@ -374,10 +375,13 @@ Instruction* DebugInfoManager::GetDebugOperationWithDeref() { Instruction* DebugInfoManager::DerefDebugExpression(Instruction* dbg_expr) { assert(dbg_expr->GetCommonDebugOpcode() == CommonDebugInfoDebugExpression); std::unique_ptr deref_expr(dbg_expr->Clone(context())); - deref_expr->SetResultId(context()->TakeNextId()); - deref_expr->InsertOperand( - kDebugExpressOperandOperationIndex, - {SPV_OPERAND_TYPE_ID, {GetDebugOperationWithDeref()->result_id()}}); + uint32_t result_id = context()->TakeNextId(); + if (result_id == 0) return nullptr; + deref_expr->SetResultId(result_id); + Instruction* deref_op = GetDebugOperationWithDeref(); + if (!deref_op) return nullptr; + deref_expr->InsertOperand(kDebugExpressOperandOperationIndex, + {SPV_OPERAND_TYPE_ID, {deref_op->result_id()}}); auto* deref_expr_instr = context()->ext_inst_debuginfo_end()->InsertBefore(std::move(deref_expr)); AnalyzeDebugInst(deref_expr_instr); diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/folding_rules.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/folding_rules.cpp index 6690381..ecdbf85 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/folding_rules.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/folding_rules.cpp @@ -16,6 +16,7 @@ #include #include +#include #include #include "ir_builder.h" @@ -933,6 +934,42 @@ FoldingRule MergeMulNegateArithmetic() { }; } +// Returns true if |inst| is negation op and is safe to fold. +static bool IsFoldableNegation(const Instruction* inst) { + return (inst->opcode() == spv::Op::OpSNegate || + (inst->opcode() == spv::Op::OpFNegate && + inst->IsFloatingPointFoldingAllowed())); +} + +// Merges multiplies of two negations. +// Cases: +// (-x) * (-y) = x * y +FoldingRule MergeMulDoubleNegative() { + return [](IRContext* context, Instruction* inst, + const std::vector&) { + assert(inst->opcode() == spv::Op::OpFMul || + inst->opcode() == spv::Op::OpIMul); + + const analysis::Type* type = + context->get_type_mgr()->GetType(inst->type_id()); + + bool uses_float = HasFloatingPoint(type); + if (uses_float && !inst->IsFloatingPointFoldingAllowed()) return false; + + analysis::DefUseManager* def_use_mgr = context->get_def_use_mgr(); + Instruction* lhs = def_use_mgr->GetDef(inst->GetSingleWordInOperand(0)); + Instruction* rhs = def_use_mgr->GetDef(inst->GetSingleWordInOperand(1)); + + if (IsFoldableNegation(lhs) && IsFoldableNegation(rhs)) { + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {lhs->GetSingleWordInOperand(0u)}}, + {SPV_OPERAND_TYPE_ID, {rhs->GetSingleWordInOperand(0u)}}}); + return true; + } + return false; + }; +} + // Merges consecutive divides if each instruction contains one constant operand. // Does not support integer division. // Cases: @@ -1125,13 +1162,12 @@ FoldingRule MergeDivNegateArithmetic() { }; } -// Folds addition of a constant and a negation. -// Cases: -// (-x) + 2 = 2 - x -// 2 + (-x) = 2 - x +// Folds addition, where one side is a negation. +// (-x) + y = y - x +// y + (-x) = y - x FoldingRule MergeAddNegateArithmetic() { return [](IRContext* context, Instruction* inst, - const std::vector& constants) { + const std::vector&) { assert(inst->opcode() == spv::Op::OpFAdd || inst->opcode() == spv::Op::OpIAdd); const analysis::Type* type = @@ -1139,73 +1175,65 @@ FoldingRule MergeAddNegateArithmetic() { bool uses_float = HasFloatingPoint(type); if (uses_float && !inst->IsFloatingPointFoldingAllowed()) return false; - const analysis::Constant* const_input1 = ConstInput(constants); - if (!const_input1) return false; - Instruction* other_inst = NonConstInput(context, constants[0], inst); - if (uses_float && !other_inst->IsFloatingPointFoldingAllowed()) + analysis::DefUseManager* def_use_mgr = context->get_def_use_mgr(); + Instruction* lhs = def_use_mgr->GetDef(inst->GetSingleWordInOperand(0)); + Instruction* rhs = def_use_mgr->GetDef(inst->GetSingleWordInOperand(1)); + + auto TrySubstitute = [inst, uses_float](Instruction* first, + Instruction* second) { + if (IsFoldableNegation(first)) { + inst->SetOpcode(uses_float ? spv::Op::OpFSub : spv::Op::OpISub); + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {second->result_id()}}, + {SPV_OPERAND_TYPE_ID, {first->GetSingleWordInOperand(0u)}}}); + return true; + } return false; + }; - if (other_inst->opcode() == spv::Op::OpSNegate || - other_inst->opcode() == spv::Op::OpFNegate) { - inst->SetOpcode(HasFloatingPoint(type) ? spv::Op::OpFSub - : spv::Op::OpISub); - uint32_t const_id = constants[0] ? inst->GetSingleWordInOperand(0u) - : inst->GetSingleWordInOperand(1u); - inst->SetInOperands( - {{SPV_OPERAND_TYPE_ID, {const_id}}, - {SPV_OPERAND_TYPE_ID, {other_inst->GetSingleWordInOperand(0u)}}}); - return true; - } - return false; + return TrySubstitute(lhs, rhs) || TrySubstitute(rhs, lhs); }; } -// Folds subtraction of a constant and a negation. +// Folds subtraction, where one side is a negation. // Cases: // (-x) - 2 = -2 - x -// 2 - (-x) = x + 2 +// y - (-x) = x + y FoldingRule MergeSubNegateArithmetic() { return [](IRContext* context, Instruction* inst, const std::vector& constants) { assert(inst->opcode() == spv::Op::OpFSub || inst->opcode() == spv::Op::OpISub); - analysis::ConstantManager* const_mgr = context->get_constant_mgr(); const analysis::Type* type = context->get_type_mgr()->GetType(inst->type_id()); - if (IsCooperativeMatrix(type)) { - return false; - } - bool uses_float = HasFloatingPoint(type); if (uses_float && !inst->IsFloatingPointFoldingAllowed()) return false; - uint32_t width = ElementWidth(type); - if (width != 32 && width != 64) return false; + analysis::DefUseManager* def_use_mgr = context->get_def_use_mgr(); + Instruction* lhs = def_use_mgr->GetDef(inst->GetSingleWordInOperand(0)); + Instruction* rhs = def_use_mgr->GetDef(inst->GetSingleWordInOperand(1)); - const analysis::Constant* const_input1 = ConstInput(constants); - if (!const_input1) return false; - Instruction* other_inst = NonConstInput(context, constants[0], inst); - if (uses_float && !other_inst->IsFloatingPointFoldingAllowed()) + if (IsFoldableNegation(rhs)) { + inst->SetOpcode(uses_float ? spv::Op::OpFAdd : spv::Op::OpIAdd); + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {lhs->result_id()}}, + {SPV_OPERAND_TYPE_ID, {rhs->GetSingleWordInOperand(0)}}}); + return true; + } + + if (IsCooperativeMatrix(type)) { return false; + } - if (other_inst->opcode() == spv::Op::OpSNegate || - other_inst->opcode() == spv::Op::OpFNegate) { - uint32_t op1 = 0; - uint32_t op2 = 0; - spv::Op opcode = inst->opcode(); - if (constants[0] != nullptr) { - op1 = other_inst->GetSingleWordInOperand(0u); - op2 = inst->GetSingleWordInOperand(0u); - opcode = HasFloatingPoint(type) ? spv::Op::OpFAdd : spv::Op::OpIAdd; - } else { - op1 = NegateConstant(const_mgr, const_input1); - op2 = other_inst->GetSingleWordInOperand(0u); - } + uint32_t width = ElementWidth(type); + if (width != 32 && width != 64) return false; - inst->SetOpcode(opcode); + if (constants[1] && IsFoldableNegation(lhs)) { inst->SetInOperands( - {{SPV_OPERAND_TYPE_ID, {op1}}, {SPV_OPERAND_TYPE_ID, {op2}}}); + {{SPV_OPERAND_TYPE_ID, + {NegateConstant(context->get_constant_mgr(), constants[1])}}, + {SPV_OPERAND_TYPE_ID, {lhs->GetSingleWordInOperand(0)}}}); return true; } return false; @@ -1530,11 +1558,13 @@ FoldingRule MergeGenericAddSubArithmetic() { }; } -// Helper function for FactorAddMuls. If |factor0_0| is the same as |factor1_0|, -// generate |factor0_0| * (|factor0_1| + |factor1_1|). -bool FactorAddMulsOpnds(uint32_t factor0_0, uint32_t factor0_1, - uint32_t factor1_0, uint32_t factor1_1, - Instruction* inst) { +// Helper function for FactorAddSubMuls. +// If |factor0_0| is the same as |factor1_0|, generate: +// |factor0_0| * (|factor0_1| + |factor1_1|) +// |factor0_0| * (|factor0_1| - |factor1_1|) +bool FactorAddSubMulsOpnds(uint32_t factor0_0, uint32_t factor0_1, + uint32_t factor1_0, uint32_t factor1_1, + Instruction* inst) { IRContext* context = inst->context(); if (factor0_0 != factor1_0) return false; InstructionBuilder ir_builder( @@ -1545,8 +1575,10 @@ bool FactorAddMulsOpnds(uint32_t factor0_0, uint32_t factor0_1, if (!new_add_inst) { return false; } - inst->SetOpcode(inst->opcode() == spv::Op::OpFAdd ? spv::Op::OpFMul - : spv::Op::OpIMul); + + bool is_float = + inst->opcode() == spv::Op::OpFAdd || inst->opcode() == spv::Op::OpFSub; + inst->SetOpcode(is_float ? spv::Op::OpFMul : spv::Op::OpIMul); inst->SetInOperands({{SPV_OPERAND_TYPE_ID, {factor0_0}}, {SPV_OPERAND_TYPE_ID, {new_add_inst->result_id()}}}); context->UpdateDefUse(inst); @@ -1554,12 +1586,16 @@ bool FactorAddMulsOpnds(uint32_t factor0_0, uint32_t factor0_1, } // Perform the following factoring identity, handling all operand order -// combinations: (a * b) + (a * c) = a * (b + c) -FoldingRule FactorAddMuls() { +// combinations: +// (a * b) + (a * c) = a * (b + c) +// (a * b) - (a * c) = a * (b - c) +FoldingRule FactorAddSubMuls() { return [](IRContext* context, Instruction* inst, const std::vector&) { assert(inst->opcode() == spv::Op::OpFAdd || - inst->opcode() == spv::Op::OpIAdd); + inst->opcode() == spv::Op::OpFSub || + inst->opcode() == spv::Op::OpIAdd || + inst->opcode() == spv::Op::OpISub); const analysis::Type* type = context->get_type_mgr()->GetType(inst->type_id()); bool uses_float = HasFloatingPoint(type); @@ -1590,11 +1626,11 @@ FoldingRule FactorAddMuls() { for (int i = 0; i < 2; i++) { for (int j = 0; j < 2; j++) { // Check if operand i in add_op0_inst matches operand j in add_op1_inst. - if (FactorAddMulsOpnds(add_op0_inst->GetSingleWordInOperand(i), - add_op0_inst->GetSingleWordInOperand(1 - i), - add_op1_inst->GetSingleWordInOperand(j), - add_op1_inst->GetSingleWordInOperand(1 - j), - inst)) + if (FactorAddSubMulsOpnds(add_op0_inst->GetSingleWordInOperand(i), + add_op0_inst->GetSingleWordInOperand(1 - i), + add_op1_inst->GetSingleWordInOperand(j), + add_op1_inst->GetSingleWordInOperand(1 - j), + inst)) return true; } } @@ -2296,6 +2332,39 @@ FoldingRule BitCastScalarOrVector() { }; } +// Remove indirect bitcasts which have no effect. +// uint32 x; asuint32(x) => x +// uint32 x; asuint32(asint32(x)) => x +// float32 x; asuint32(asint32(x)) => asuint32(x) +FoldingRule RedundantBitcast() { + return [](IRContext* context, Instruction* inst, + const std::vector&) { + assert(inst->opcode() == spv::Op::OpBitcast); + + analysis::DefUseManager* def_mgr = context->get_def_use_mgr(); + Instruction* child = def_mgr->GetDef(inst->GetSingleWordInOperand(0)); + + if (inst->type_id() == child->type_id()) { + inst->SetOpcode(spv::Op::OpCopyObject); + inst->SetInOperands({{SPV_OPERAND_TYPE_ID, {child->result_id()}}}); + return true; + } + + if (child->opcode() != spv::Op::OpBitcast) { + return false; + } + + if (def_mgr->GetDef(child->GetSingleWordInOperand(0))->type_id() == + inst->type_id()) { + inst->SetOpcode(spv::Op::OpCopyObject); + } + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {child->GetSingleWordInOperand(0)}}}); + + return true; + }; +} + FoldingRule BitReverseScalarOrVector() { return [](IRContext* context, Instruction* inst, const std::vector& constants) { @@ -2410,6 +2479,250 @@ FoldingRule RedundantSelect() { }; } +std::optional GetBoolConstantKind(const analysis::Constant* c) { + if (!c) { + return {}; + } + if (auto composite = c->AsCompositeConstant()) { + auto& components = composite->GetComponents(); + if (components.empty()) { + return {}; + } + auto first = GetBoolConstantKind(components[0]); + if (!first) { + return {}; + } + if (std::all_of(std::begin(components) + 1, std::end(components), + [first](const analysis::Constant* c2) { + return GetBoolConstantKind(c2) == first; + })) { + return first; + } + return {}; + } else if (c->AsNullConstant()) { + return false; + } else if (c->AsBoolConstant()) { + return c->AsBoolConstant()->value(); + } + return {}; +} + +// Fold OpSelect instructions which have constant booleans as their result. +// x ? true : false = x +// x ? false : true = !x +FoldingRule FoldConstantBooleanSelect() { + return [](IRContext* context, Instruction* inst, + const std::vector& constants) { + assert(inst->opcode() == spv::Op::OpSelect); + assert(inst->NumInOperands() == 3); + assert(constants.size() == 3); + + if (!constants[1] || !constants[2]) { + return false; + } + + analysis::DefUseManager* def_mgr = context->get_def_use_mgr(); + if (inst->type_id() != + def_mgr->GetDef(inst->GetSingleWordInOperand(0))->type_id()) { + return false; + } + + std::optional uniform_true = GetBoolConstantKind(constants[1]); + std::optional uniform_false = GetBoolConstantKind(constants[2]); + + if (!uniform_true || !uniform_false) { + return false; + } + + if (uniform_true.value() && !uniform_false.value()) { + inst->SetOpcode(spv::Op::OpCopyObject); + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {inst->GetSingleWordInOperand(0)}}}); + return true; + } else if (!uniform_true.value() && uniform_false.value()) { + inst->SetOpcode(spv::Op::OpLogicalNot); + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {inst->GetSingleWordInOperand(0)}}}); + return true; + } + return false; + }; +} + +// Fold OpLogicalAnd instructions which have a constant true on one side. +// x && true = x +// true && x = x +FoldingRule RedundantLogicalAnd() { + return [](IRContext* context, Instruction* inst, + const std::vector& constants) { + assert(inst->opcode() == spv::Op::OpLogicalAnd); + + if (GetBoolConstantKind(ConstInput(constants)) == + std::optional(true)) { + Instruction* other_inst = NonConstInput(context, constants[0], inst); + inst->SetOpcode(spv::Op::OpCopyObject); + inst->SetInOperands({{SPV_OPERAND_TYPE_ID, {other_inst->result_id()}}}); + return true; + } + return false; + }; +} + +// Fold OpLogicalOr instructions which have a constant false on one side. +// x || false = x +// false || x = x +FoldingRule RedundantLogicalOr() { + return [](IRContext* context, Instruction* inst, + const std::vector& constants) { + assert(inst->opcode() == spv::Op::OpLogicalOr); + + if (GetBoolConstantKind(ConstInput(constants)) == + std::optional(false)) { + Instruction* other_inst = NonConstInput(context, constants[0], inst); + inst->SetOpcode(spv::Op::OpCopyObject); + inst->SetInOperands({{SPV_OPERAND_TYPE_ID, {other_inst->result_id()}}}); + return true; + } + return false; + }; +} + +// Fold concurrent OpLogicalNot instructions: +// !!x = x +FoldingRule RedundantLogicalNot() { + return [](IRContext* context, Instruction* inst, + const std::vector&) { + assert(inst->opcode() == spv::Op::OpLogicalNot); + Instruction* child = + context->get_def_use_mgr()->GetDef(inst->GetSingleWordInOperand(0)); + if (child->opcode() == spv::Op::OpLogicalNot) { + inst->SetOpcode(spv::Op::OpCopyObject); + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {child->GetSingleWordInOperand(0)}}}); + return true; + } + return false; + }; +} + +// Fold OpLogicalNot instructions that follow a comparison, +// if the comparison is only used by that instruction. +// +// !(a == b) = (a != b) +// !(a != b) = (a == b) +// !(a < b) = (a >= b) +// !(a >= b) = (a < b) +// !(a > b) = (a <= b) +// !(a <= b) = (a > b) +FoldingRule FoldLogicalNotComparison() { + return [](IRContext* context, Instruction* inst, + const std::vector&) { + assert(inst->opcode() == spv::Op::OpLogicalNot); + analysis::DefUseManager* def_mgr = context->get_def_use_mgr(); + Instruction* child = + context->get_def_use_mgr()->GetDef(inst->GetSingleWordInOperand(0)); + + if (def_mgr->NumUses(child) > 1) { + return false; + } + + spv::Op new_opcode = spv::Op::OpNop; + switch (child->opcode()) { + // (a == b) <=> (a != b) + case spv::Op::OpIEqual: + new_opcode = spv::Op::OpINotEqual; + break; + case spv::Op::OpINotEqual: + new_opcode = spv::Op::OpIEqual; + break; + case spv::Op::OpFOrdEqual: + new_opcode = spv::Op::OpFUnordNotEqual; + break; + case spv::Op::OpFOrdNotEqual: + new_opcode = spv::Op::OpFUnordEqual; + break; + case spv::Op::OpFUnordEqual: + new_opcode = spv::Op::OpFOrdNotEqual; + break; + case spv::Op::OpFUnordNotEqual: + new_opcode = spv::Op::OpFOrdEqual; + break; + case spv::Op::OpLogicalEqual: + new_opcode = spv::Op::OpLogicalNotEqual; + break; + case spv::Op::OpLogicalNotEqual: + new_opcode = spv::Op::OpLogicalEqual; + break; + + // (a > b) <=> (a <= b) + case spv::Op::OpUGreaterThan: + new_opcode = spv::Op::OpULessThanEqual; + break; + case spv::Op::OpULessThanEqual: + new_opcode = spv::Op::OpUGreaterThan; + break; + case spv::Op::OpSGreaterThan: + new_opcode = spv::Op::OpSLessThanEqual; + break; + case spv::Op::OpSLessThanEqual: + new_opcode = spv::Op::OpSGreaterThan; + break; + case spv::Op::OpFOrdGreaterThan: + new_opcode = spv::Op::OpFUnordLessThanEqual; + break; + case spv::Op::OpFOrdLessThanEqual: + new_opcode = spv::Op::OpFUnordGreaterThan; + break; + case spv::Op::OpFUnordGreaterThan: + new_opcode = spv::Op::OpFOrdLessThanEqual; + break; + case spv::Op::OpFUnordLessThanEqual: + new_opcode = spv::Op::OpFOrdGreaterThan; + break; + + // (a < b) <=> (a >= b) + case spv::Op::OpULessThan: + new_opcode = spv::Op::OpUGreaterThanEqual; + break; + case spv::Op::OpUGreaterThanEqual: + new_opcode = spv::Op::OpULessThan; + break; + case spv::Op::OpSLessThan: + new_opcode = spv::Op::OpSGreaterThanEqual; + break; + case spv::Op::OpSGreaterThanEqual: + new_opcode = spv::Op::OpSLessThan; + break; + case spv::Op::OpFOrdLessThan: + new_opcode = spv::Op::OpFUnordGreaterThanEqual; + break; + case spv::Op::OpFOrdGreaterThanEqual: + new_opcode = spv::Op::OpFUnordLessThan; + break; + case spv::Op::OpFUnordLessThan: + new_opcode = spv::Op::OpFOrdGreaterThanEqual; + break; + case spv::Op::OpFUnordGreaterThanEqual: + new_opcode = spv::Op::OpFOrdLessThan; + break; + + default: + break; + } + + if (new_opcode == spv::Op::OpNop) { + return false; + } + + inst->SetOpcode(new_opcode); + inst->SetInOperands( + {{SPV_OPERAND_TYPE_ID, {child->GetSingleWordInOperand(0)}}, + {SPV_OPERAND_TYPE_ID, {child->GetSingleWordInOperand(1)}}}); + + return true; + }; +} + enum class FloatConstantKind { Unknown, Zero, One }; FloatConstantKind getFloatConstantKind(const analysis::Constant* constant) { @@ -3394,6 +3707,8 @@ void FoldingRules::AddFoldingRules() { rules_[spv::Op::OpUMod].push_back(RedundantSUMod()); rules_[spv::Op::OpBitcast].push_back(BitCastScalarOrVector()); + rules_[spv::Op::OpBitcast].push_back(RedundantBitcast()); + rules_[spv::Op::OpBitReverse].push_back(BitReverseScalarOrVector()); rules_[spv::Op::OpCompositeConstruct].push_back( @@ -3417,7 +3732,7 @@ void FoldingRules::AddFoldingRules() { rules_[spv::Op::OpFAdd].push_back(MergeAddAddArithmetic()); rules_[spv::Op::OpFAdd].push_back(MergeAddSubArithmetic()); rules_[spv::Op::OpFAdd].push_back(MergeGenericAddSubArithmetic()); - rules_[spv::Op::OpFAdd].push_back(FactorAddMuls()); + rules_[spv::Op::OpFAdd].push_back(FactorAddSubMuls()); rules_[spv::Op::OpFDiv].push_back(RedundantFDiv()); rules_[spv::Op::OpFDiv].push_back(ReciprocalFDiv()); @@ -3431,6 +3746,7 @@ void FoldingRules::AddFoldingRules() { rules_[spv::Op::OpFMul].push_back(MergeMulMulArithmetic()); rules_[spv::Op::OpFMul].push_back(MergeMulDivArithmetic()); rules_[spv::Op::OpFMul].push_back(MergeMulNegateArithmetic()); + rules_[spv::Op::OpFMul].push_back(MergeMulDoubleNegative()); rules_[spv::Op::OpFNegate].push_back(MergeNegateArithmetic()); rules_[spv::Op::OpFNegate].push_back(MergeNegateAddSubArithmetic()); @@ -3440,20 +3756,23 @@ void FoldingRules::AddFoldingRules() { rules_[spv::Op::OpFSub].push_back(MergeSubNegateArithmetic()); rules_[spv::Op::OpFSub].push_back(MergeSubAddArithmetic()); rules_[spv::Op::OpFSub].push_back(MergeSubSubArithmetic()); + rules_[spv::Op::OpFSub].push_back(FactorAddSubMuls()); rules_[spv::Op::OpIAdd].push_back(MergeAddNegateArithmetic()); rules_[spv::Op::OpIAdd].push_back(MergeAddAddArithmetic()); rules_[spv::Op::OpIAdd].push_back(MergeAddSubArithmetic()); rules_[spv::Op::OpIAdd].push_back(MergeGenericAddSubArithmetic()); - rules_[spv::Op::OpIAdd].push_back(FactorAddMuls()); + rules_[spv::Op::OpIAdd].push_back(FactorAddSubMuls()); rules_[spv::Op::OpIMul].push_back(IntMultipleBy1()); rules_[spv::Op::OpIMul].push_back(MergeMulMulArithmetic()); rules_[spv::Op::OpIMul].push_back(MergeMulNegateArithmetic()); + rules_[spv::Op::OpIMul].push_back(MergeMulDoubleNegative()); rules_[spv::Op::OpISub].push_back(MergeSubNegateArithmetic()); rules_[spv::Op::OpISub].push_back(MergeSubAddArithmetic()); rules_[spv::Op::OpISub].push_back(MergeSubSubArithmetic()); + rules_[spv::Op::OpISub].push_back(FactorAddSubMuls()); rules_[spv::Op::OpBitwiseAnd].push_back(RedundantAndOrXor()); rules_[spv::Op::OpBitwiseAnd].push_back(RedundantAndAddSub()); @@ -3466,6 +3785,14 @@ void FoldingRules::AddFoldingRules() { rules_[spv::Op::OpSNegate].push_back(MergeNegateAddSubArithmetic()); rules_[spv::Op::OpSelect].push_back(RedundantSelect()); + rules_[spv::Op::OpSelect].push_back(FoldConstantBooleanSelect()); + + rules_[spv::Op::OpLogicalAnd].push_back(RedundantLogicalAnd()); + + rules_[spv::Op::OpLogicalOr].push_back(RedundantLogicalOr()); + + rules_[spv::Op::OpLogicalNot].push_back(RedundantLogicalNot()); + rules_[spv::Op::OpLogicalNot].push_back(FoldLogicalNotComparison()); rules_[spv::Op::OpStore].push_back(StoringUndef()); diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.cpp index 2ce9e85..88b1f2e 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.cpp @@ -183,6 +183,8 @@ Instruction* IRContext::KillInst(Instruction* inst) { KillOperandFromDebugInstructions(inst); + KillRelatedDebugScopes(inst); + if (AreAnalysesValid(kAnalysisDefUse)) { analysis::DefUseManager* def_use_mgr = get_def_use_mgr(); def_use_mgr->ClearInst(inst); @@ -532,6 +534,20 @@ void IRContext::KillOperandFromDebugInstructions(Instruction* inst) { } } +void IRContext::KillRelatedDebugScopes(Instruction* inst) { + // Extension has been fully unloaded, remove debug scope from every + // instruction. + if (inst->opcode() == spv::Op::OpExtInstImport) { + const std::string extension_name = inst->GetInOperand(0).AsString(); + if (extension_name == "NonSemantic.Shader.DebugInfo.100" || + extension_name == "OpenCL.DebugInfo.100") { + module()->ForEachInst([](Instruction* child) { + child->SetDebugScope(DebugScope(kNoDebugScope, kNoInlinedAt)); + }); + } + } +} + void IRContext::AddCombinatorsForCapability(uint32_t capability) { spv::Capability cap = spv::Capability(capability); if (cap == spv::Capability::Shader) { diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.h b/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.h index 89e8cb0..f4a69fc 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.h +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/ir_context.h @@ -508,6 +508,9 @@ class IRContext { // Change operands of debug instruction to DebugInfoNone. void KillOperandFromDebugInstructions(Instruction* inst); + // Remove the debug scope from any instruction related to |inst|. + void KillRelatedDebugScopes(Instruction* inst); + // Returns the next unique id for use by an instruction. inline uint32_t TakeNextUniqueId() { assert(unique_id_ != std::numeric_limits::max()); diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.cpp index e4eb751..8da0668 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.cpp @@ -340,7 +340,7 @@ bool MemPass::IsTargetVar(uint32_t varId) { // %50 = OpUndef %int // [ ... ] // %30 = OpPhi %int %int_42 %13 %50 %14 %50 %15 -void MemPass::RemovePhiOperands( +bool MemPass::RemovePhiOperands( Instruction* phi, const std::unordered_set& reachable_blocks) { std::vector keep_operands; uint32_t type_id = 0; @@ -382,6 +382,7 @@ void MemPass::RemovePhiOperands( if (!undef_id) { type_id = arg_def_instr->type_id(); undef_id = Type2Undef(type_id); + if (undef_id == 0) return false; } keep_operands.push_back( Operand(spv_operand_type_t::SPV_OPERAND_TYPE_ID, {undef_id})); @@ -400,6 +401,7 @@ void MemPass::RemovePhiOperands( context()->ForgetUses(phi); phi->ReplaceOperands(keep_operands); context()->AnalyzeUses(phi); + return true; } void MemPass::RemoveBlock(Function::iterator* bi) { @@ -422,8 +424,8 @@ void MemPass::RemoveBlock(Function::iterator* bi) { *bi = bi->Erase(); } -bool MemPass::RemoveUnreachableBlocks(Function* func) { - if (func->IsDeclaration()) return false; +Pass::Status MemPass::RemoveUnreachableBlocks(Function* func) { + if (func->IsDeclaration()) return Status::SuccessWithoutChange; bool modified = false; // Mark reachable all blocks reachable from the function's entry block. @@ -469,9 +471,11 @@ bool MemPass::RemoveUnreachableBlocks(Function* func) { // If the block is reachable and has Phi instructions, remove all // operands from its Phi instructions that reference unreachable blocks. // If the block has no Phi instructions, this is a no-op. - block.ForEachPhiInst([&reachable_blocks, this](Instruction* phi) { - RemovePhiOperands(phi, reachable_blocks); - }); + bool success = + block.WhileEachPhiInst([&reachable_blocks, this](Instruction* phi) { + return RemovePhiOperands(phi, reachable_blocks); + }); + if (!success) return Status::Failure; } // Erase unreachable blocks. @@ -484,13 +488,11 @@ bool MemPass::RemoveUnreachableBlocks(Function* func) { } } - return modified; + return modified ? Status::SuccessWithChange : Status::SuccessWithoutChange; } -bool MemPass::CFGCleanup(Function* func) { - bool modified = false; - modified |= RemoveUnreachableBlocks(func); - return modified; +Pass::Status MemPass::CFGCleanup(Function* func) { + return RemoveUnreachableBlocks(func); } void MemPass::CollectTargetVars(Function* func) { diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.h b/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.h index aef9e5f..496286b 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.h +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/mem_pass.h @@ -114,7 +114,7 @@ class MemPass : public Pass { void DCEInst(Instruction* inst, const std::function&); // Call all the cleanup helper functions on |func|. - bool CFGCleanup(Function* func); + Status CFGCleanup(Function* func); // Return true if |op| is supported decorate. inline bool IsNonTypeDecorate(spv::Op op) const { @@ -142,15 +142,15 @@ class MemPass : public Pass { bool HasOnlySupportedRefs(uint32_t varId); // Remove all the unreachable basic blocks in |func|. - bool RemoveUnreachableBlocks(Function* func); + Status RemoveUnreachableBlocks(Function* func); // Remove the block pointed by the iterator |*bi|. This also removes // all the instructions in the pointed-to block. void RemoveBlock(Function::iterator* bi); // Remove Phi operands in |phi| that are coming from blocks not in - // |reachable_blocks|. - void RemovePhiOperands( + // |reachable_blocks|. Returns false if it fails. + bool RemovePhiOperands( Instruction* phi, const std::unordered_set& reachable_blocks); diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/pass.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/pass.cpp index 08d76b5..ce37f36 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/pass.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/pass.cpp @@ -117,6 +117,9 @@ uint32_t Pass::GenerateCopy(Instruction* object_to_copy, uint32_t new_type_id, // TODO(1841): Handle id overflow. Instruction* extract = ir_builder.AddCompositeExtract( original_element_type_id, object_to_copy->result_id(), {i}); + if (extract == nullptr) { + return 0; + } uint32_t new_id = GenerateCopy(extract, new_element_type_id, insertion_position); if (new_id == 0) { @@ -125,8 +128,12 @@ uint32_t Pass::GenerateCopy(Instruction* object_to_copy, uint32_t new_type_id, element_ids.push_back(new_id); } - return ir_builder.AddCompositeConstruct(new_type_id, element_ids) - ->result_id(); + Instruction* construct = + ir_builder.AddCompositeConstruct(new_type_id, element_ids); + if (construct == nullptr) { + return 0; + } + return construct->result_id(); } case spv::Op::OpTypeStruct: { std::vector element_ids; @@ -136,6 +143,9 @@ uint32_t Pass::GenerateCopy(Instruction* object_to_copy, uint32_t new_type_id, // TODO(1841): Handle id overflow. Instruction* extract = ir_builder.AddCompositeExtract( orig_member_type_id, object_to_copy->result_id(), {i}); + if (extract == nullptr) { + return 0; + } uint32_t new_id = GenerateCopy(extract, new_member_type_id, insertion_position); if (new_id == 0) { @@ -143,8 +153,12 @@ uint32_t Pass::GenerateCopy(Instruction* object_to_copy, uint32_t new_type_id, } element_ids.push_back(new_id); } - return ir_builder.AddCompositeConstruct(new_type_id, element_ids) - ->result_id(); + Instruction* construct = + ir_builder.AddCompositeConstruct(new_type_id, element_ids); + if (construct == nullptr) { + return 0; + } + return construct->result_id(); } default: // If we do not have an aggregate type, then we have a problem. Either we diff --git a/libs/bgfx/3rdparty/spirv-tools/source/opt/value_number_table.cpp b/libs/bgfx/3rdparty/spirv-tools/source/opt/value_number_table.cpp index 8c33ab7..a93d33c 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/opt/value_number_table.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/opt/value_number_table.cpp @@ -183,29 +183,13 @@ uint32_t ValueNumberTable::AssignValueNumber(Instruction* inst) { } // Apply normal form, so a+b == b+a - switch (value_ins.opcode()) { - case spv::Op::OpIAdd: - case spv::Op::OpFAdd: - case spv::Op::OpIMul: - case spv::Op::OpFMul: - case spv::Op::OpDot: - case spv::Op::OpLogicalEqual: - case spv::Op::OpLogicalNotEqual: - case spv::Op::OpLogicalOr: - case spv::Op::OpLogicalAnd: - case spv::Op::OpIEqual: - case spv::Op::OpINotEqual: - case spv::Op::OpBitwiseOr: - case spv::Op::OpBitwiseXor: - case spv::Op::OpBitwiseAnd: - if (value_ins.GetSingleWordInOperand(0) > - value_ins.GetSingleWordInOperand(1)) { - value_ins.SetInOperands( - {{SPV_OPERAND_TYPE_ID, {value_ins.GetSingleWordInOperand(1)}}, - {SPV_OPERAND_TYPE_ID, {value_ins.GetSingleWordInOperand(0)}}}); - } - default: - break; + if (spvOpcodeIsCommutativeBinaryOperator(value_ins.opcode())) { + if (value_ins.GetSingleWordInOperand(0) > + value_ins.GetSingleWordInOperand(1)) { + value_ins.SetInOperands( + {{SPV_OPERAND_TYPE_ID, {value_ins.GetSingleWordInOperand(1)}}, + {SPV_OPERAND_TYPE_ID, {value_ins.GetSingleWordInOperand(0)}}}); + } } // Otherwise, we check if this value has been computed before. diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_atomics.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_atomics.cpp index 510960b..8cda07e 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_atomics.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_atomics.cpp @@ -235,7 +235,9 @@ spv_result_t AtomicsPass(ValidationState_t& _, const Instruction* inst) { if (!IsStorageClassAllowedByUniversalRules(storage_class)) { return _.diag(SPV_ERROR_INVALID_DATA, inst) << spvOpcodeString(opcode) - << ": storage class forbidden by universal validation rules."; + << ": Can not be used with storage class " + << spvtools::StorageClassToString(storage_class) + << " by universal validation rules"; } // Then Shader rules @@ -249,8 +251,10 @@ spv_result_t AtomicsPass(ValidationState_t& _, const Instruction* inst) { (storage_class != spv::StorageClass::PhysicalStorageBuffer) && (storage_class != spv::StorageClass::TaskPayloadWorkgroupEXT)) { return _.diag(SPV_ERROR_INVALID_DATA, inst) - << _.VkErrorID(4686) << spvOpcodeString(opcode) - << ": Vulkan spec only allows storage classes for atomic to " + << _.VkErrorID(4686) << spvOpcodeString(opcode) << ": " + << spvtools::StorageClassToString(storage_class) + << " is not allowed, the Vulkan spec only allows storage " + "classes for atomic to " "be: Uniform, Workgroup, Image, StorageBuffer, " "PhysicalStorageBuffer or TaskPayloadWorkgroupEXT."; } @@ -335,8 +339,9 @@ spv_result_t AtomicsPass(ValidationState_t& _, const Instruction* inst) { (storage_class != spv::StorageClass::CrossWorkgroup) && (storage_class != spv::StorageClass::Generic)) { return _.diag(SPV_ERROR_INVALID_DATA, inst) - << spvOpcodeString(opcode) - << ": storage class must be Function, Workgroup, " + << spvOpcodeString(opcode) << ": storage class is " + << spvtools::StorageClassToString(storage_class) + << ", but must be Function, Workgroup, " "CrossWorkGroup or Generic in the OpenCL environment."; } diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_builtins.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_builtins.cpp index c8586a7..f75a707 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_builtins.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_builtins.cpp @@ -2955,9 +2955,24 @@ spv_result_t BuiltInsValidator::ValidateMeshBuiltinInterfaceRules( const Decoration& decoration, const Instruction& inst, spv::Op scalar_type, const Instruction& referenced_from_inst) { if (function_id_) { - if (execution_models_.count(spv::ExecutionModel::MeshEXT)) { + if (!execution_models_.count(spv::ExecutionModel::MeshEXT)) { + return SPV_SUCCESS; + } + + const spv::BuiltIn builtin = decoration.builtin(); + const bool is_topology = + builtin == spv::BuiltIn::PrimitiveTriangleIndicesEXT || + builtin == spv::BuiltIn::PrimitiveLineIndicesEXT || + builtin == spv::BuiltIn::PrimitivePointIndicesEXT; + + // These builtin have the ability to be an array with MeshEXT + // When an array, we need to make sure the array size lines up + std::map entry_interface_id_map; + const bool is_interface_var = + IsMeshInterfaceVar(inst, entry_interface_id_map); + + if (!is_topology) { bool is_block = false; - const spv::BuiltIn builtin = decoration.builtin(); static const std::unordered_map mesh_vuid_map = {{ @@ -2997,12 +3012,7 @@ spv_result_t BuiltInsValidator::ValidateMeshBuiltinInterfaceRules( << " within the MeshEXT Execution Model must also be " << "decorated with the PerPrimitiveEXT decoration. "; } - - // These builtin have the ability to be an array with MeshEXT - // When an array, we need to make sure the array size lines up - std::map entry_interface_id_map; - bool found = IsMeshInterfaceVar(inst, entry_interface_id_map); - if (found) { + if (is_interface_var) { for (const auto& id : entry_interface_id_map) { uint32_t entry_point_id = id.first; uint32_t interface_var_id = id.second; @@ -3025,6 +3035,86 @@ spv_result_t BuiltInsValidator::ValidateMeshBuiltinInterfaceRules( } } } + + if (is_interface_var && is_topology) { + for (const auto& id : entry_interface_id_map) { + uint32_t entry_point_id = id.first; + + uint64_t max_output_primitives = + _.GetOutputPrimitivesEXT(entry_point_id); + uint32_t underlying_type = 0; + if (spv_result_t error = + GetUnderlyingType(_, decoration, inst, &underlying_type)) { + return error; + } + + uint64_t primitive_array_dim = 0; + if (_.GetIdOpcode(underlying_type) == spv::Op::OpTypeArray) { + underlying_type = _.FindDef(underlying_type)->word(3u); + if (!_.EvalConstantValUint64(underlying_type, &primitive_array_dim)) { + assert(0 && "Array type definition is corrupt"); + } + } + + const auto* modes = _.GetExecutionModes(entry_point_id); + if (builtin == spv::BuiltIn::PrimitiveTriangleIndicesEXT) { + if (!modes || !modes->count(spv::ExecutionMode::OutputTrianglesEXT)) { + return _.diag(SPV_ERROR_INVALID_DATA, &inst) + << _.VkErrorID(7054) + << "The PrimitiveTriangleIndicesEXT decoration must be used " + "with the OutputTrianglesEXT Execution Mode. "; + } + if (primitive_array_dim && + primitive_array_dim != max_output_primitives) { + return _.diag(SPV_ERROR_INVALID_DATA, &inst) + << _.VkErrorID(7058) + << "The size of the array decorated with " + "PrimitiveTriangleIndicesEXT (" + << primitive_array_dim + << ") must match the value specified " + "by OutputPrimitivesEXT (" + << max_output_primitives << "). "; + } + } else if (builtin == spv::BuiltIn::PrimitiveLineIndicesEXT) { + if (!modes || !modes->count(spv::ExecutionMode::OutputLinesEXT)) { + return _.diag(SPV_ERROR_INVALID_DATA, &inst) + << _.VkErrorID(7048) + << "The PrimitiveLineIndicesEXT decoration must be used " + "with the OutputLinesEXT Execution Mode. "; + } + if (primitive_array_dim && + primitive_array_dim != max_output_primitives) { + return _.diag(SPV_ERROR_INVALID_DATA, &inst) + << _.VkErrorID(7052) + << "The size of the array decorated with " + "PrimitiveLineIndicesEXT (" + << primitive_array_dim + << ") must match the value specified " + "by OutputPrimitivesEXT (" + << max_output_primitives << "). "; + } + + } else if (builtin == spv::BuiltIn::PrimitivePointIndicesEXT) { + if (!modes || !modes->count(spv::ExecutionMode::OutputPoints)) { + return _.diag(SPV_ERROR_INVALID_DATA, &inst) + << _.VkErrorID(7042) + << "The PrimitivePointIndicesEXT decoration must be used " + "with the OutputPoints Execution Mode. "; + } + if (primitive_array_dim && + primitive_array_dim != max_output_primitives) { + return _.diag(SPV_ERROR_INVALID_DATA, &inst) + << _.VkErrorID(7046) + << "The size of the array decorated with " + "PrimitivePointIndicesEXT (" + << primitive_array_dim + << ") must match the value specified " + "by OutputPrimitivesEXT (" + << max_output_primitives << "). "; + } + } + } + } } else { // Propagate this rule to all dependant ids in the global scope. id_to_at_reference_checks_[referenced_from_inst.id()].push_back( @@ -4650,12 +4740,6 @@ spv_result_t BuiltInsValidator::ValidateMeshShadingEXTBuiltinsAtDefinition( } break; case spv::BuiltIn::CullPrimitiveEXT: { - // We know this only allowed for Mesh Execution Model - if (spv_result_t error = ValidateMeshBuiltinInterfaceRules( - decoration, inst, spv::Op::OpTypeBool, inst)) { - return error; - } - for (const uint32_t entry_point : _.entry_points()) { auto* models = _.GetExecutionModels(entry_point); if (models->find(spv::ExecutionModel::MeshEXT) == models->end() && @@ -4683,88 +4767,19 @@ spv_result_t BuiltInsValidator::ValidateMeshShadingEXTBuiltinsAtDefinition( default: assert(0 && "Unexpected mesh EXT builtin"); } - for (const uint32_t entry_point : _.entry_points()) { - // execution modes and builtin are both global, so only check these - // buildit definitions if we know the entrypoint is Mesh - auto* models = _.GetExecutionModels(entry_point); - if (models->find(spv::ExecutionModel::MeshEXT) == models->end() && - models->find(spv::ExecutionModel::MeshNV) == models->end()) { - continue; - } - const auto* modes = _.GetExecutionModes(entry_point); - uint64_t max_output_primitives = _.GetOutputPrimitivesEXT(entry_point); - uint32_t underlying_type = 0; - if (spv_result_t error = - GetUnderlyingType(_, decoration, inst, &underlying_type)) { - return error; - } - - uint64_t primitive_array_dim = 0; - if (_.GetIdOpcode(underlying_type) == spv::Op::OpTypeArray) { - underlying_type = _.FindDef(underlying_type)->word(3u); - if (!_.EvalConstantValUint64(underlying_type, &primitive_array_dim)) { - assert(0 && "Array type definition is corrupt"); - } - } - switch (builtin) { - case spv::BuiltIn::PrimitivePointIndicesEXT: - if (!modes || !modes->count(spv::ExecutionMode::OutputPoints)) { - return _.diag(SPV_ERROR_INVALID_DATA, &inst) - << _.VkErrorID(7042) - << "The PrimitivePointIndicesEXT decoration must be used " - "with " - "the OutputPoints Execution Mode. "; - } - if (primitive_array_dim && - primitive_array_dim != max_output_primitives) { - return _.diag(SPV_ERROR_INVALID_DATA, &inst) - << _.VkErrorID(7046) - << "The size of the array decorated with " - "PrimitivePointIndicesEXT must match the value specified " - "by OutputPrimitivesEXT. "; - } - break; - case spv::BuiltIn::PrimitiveLineIndicesEXT: - if (!modes || !modes->count(spv::ExecutionMode::OutputLinesEXT)) { - return _.diag(SPV_ERROR_INVALID_DATA, &inst) - << _.VkErrorID(7048) - << "The PrimitiveLineIndicesEXT decoration must be used " - "with " - "the OutputLinesEXT Execution Mode. "; - } - if (primitive_array_dim && - primitive_array_dim != max_output_primitives) { - return _.diag(SPV_ERROR_INVALID_DATA, &inst) - << _.VkErrorID(7052) - << "The size of the array decorated with " - "PrimitiveLineIndicesEXT must match the value specified " - "by OutputPrimitivesEXT. "; - } - break; - case spv::BuiltIn::PrimitiveTriangleIndicesEXT: - if (!modes || !modes->count(spv::ExecutionMode::OutputTrianglesEXT)) { - return _.diag(SPV_ERROR_INVALID_DATA, &inst) - << _.VkErrorID(7054) - << "The PrimitiveTriangleIndicesEXT decoration must be used " - "with " - "the OutputTrianglesEXT Execution Mode. "; - } - if (primitive_array_dim && - primitive_array_dim != max_output_primitives) { - return _.diag(SPV_ERROR_INVALID_DATA, &inst) - << _.VkErrorID(7058) - << "The size of the array decorated with " - "PrimitiveTriangleIndicesEXT must match the value " - "specified " - "by OutputPrimitivesEXT. "; - } - break; - default: - break; // no validation rules - } + // - We know this only allowed for Mesh Execution Model. + // - The Scalar type is is boolean for CullPrimitiveEXT, the other 3 builtin + // (topology) don't need this type. + // - It is possible to have multiple mesh + // shaders (https://github.com/KhronosGroup/SPIRV-Tools/issues/6320) and we + // need to validate these at reference time. + if (spv_result_t error = ValidateMeshBuiltinInterfaceRules( + decoration, inst, spv::Op::OpTypeBool, inst)) { + return error; } } + // Seed at reference checks with this built-in. return ValidateMeshShadingEXTBuiltinsAtReference(decoration, inst, inst, inst); diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_decorations.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_decorations.cpp index 2c8ef95..5d6dc55 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_decorations.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_decorations.cpp @@ -398,24 +398,6 @@ bool IsAlignedTo(uint32_t offset, uint32_t alignment) { return 0 == (offset % alignment); } -std::string getStorageClassString(spv::StorageClass sc) { - switch (sc) { - case spv::StorageClass::Uniform: - return "Uniform"; - case spv::StorageClass::UniformConstant: - return "UniformConstant"; - case spv::StorageClass::PushConstant: - return "PushConstant"; - case spv::StorageClass::Workgroup: - return "Workgroup"; - case spv::StorageClass::PhysicalStorageBuffer: - return "PhysicalStorageBuffer"; - default: - // Only other valid storage class in these checks - return "StorageBuffer"; - } -} - // Returns SPV_SUCCESS if the given struct satisfies standard layout rules for // Block or BufferBlocks in Vulkan. Otherwise emits a diagnostic and returns // something other than SPV_SUCCESS. Matrices inherit the specified column @@ -442,7 +424,7 @@ spv_result_t checkLayout(uint32_t struct_id, spv::StorageClass storage_class, DiagnosticStream ds = std::move( vstate.diag(SPV_ERROR_INVALID_ID, vstate.FindDef(struct_id)) << "Structure id " << struct_id << " decorated as " << decoration_str - << " for variable in " << getStorageClassString(storage_class) + << " for variable in " << StorageClassToString(storage_class) << " storage class must follow " << (scalar_block_layout ? "scalar " @@ -1282,7 +1264,7 @@ spv_result_t CheckDecorationsOfBuffers(ValidationState_t& vstate) { if (!entry_points.empty() && !hasDecoration(var_id, spv::Decoration::Binding, vstate)) { return vstate.diag(SPV_ERROR_INVALID_ID, vstate.FindDef(var_id)) - << getStorageClassString(storageClass) << " id '" << var_id + << StorageClassToString(storageClass) << " id '" << var_id << "' is missing Binding decoration.\n" << "From ARB_gl_spirv extension:\n" << "Uniform and shader storage block variables must " diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_logical_pointers.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_logical_pointers.cpp index 1528701..6f510fc 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_logical_pointers.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_logical_pointers.cpp @@ -50,6 +50,9 @@ bool IsVariablePointer(const ValidationState_t& _, return iter->second; } + // Temporarily mark the instruction as NOT a variable pointer. + variable_pointers[inst->id()] = false; + bool is_var_ptr = false; switch (inst->opcode()) { case spv::Op::OpPtrAccessChain: @@ -625,7 +628,7 @@ spv_result_t TraceVariablePointers( trace_inst->uses()); std::unordered_set store_seen; while (!store_stack.empty()) { - const auto& use = store_stack.back(); + const auto use = store_stack.back(); store_stack.pop_back(); if (!store_seen.insert(use.first).second) { @@ -766,7 +769,7 @@ spv_result_t TraceUnmodifiedVariablePointers( trace_inst->uses()); std::unordered_set store_seen; while (!store_stack.empty()) { - const auto& use = store_stack.back(); + const auto use = store_stack.back(); store_stack.pop_back(); if (!store_seen.insert(use.first).second) { diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_memory.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_memory.cpp index 9372f5c..d33c3c2 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_memory.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_memory.cpp @@ -15,6 +15,7 @@ // limitations under the License. #include +#include #include #include @@ -773,16 +774,17 @@ spv_result_t ValidateVariable(ValidationState_t& _, const Instruction* inst) { if (spvIsVulkanEnv(_.context()->target_env)) { // OpTypeRuntimeArray should only ever be in a container like OpTypeStruct, // so should never appear as a bare variable. - // Unless the module has the RuntimeDescriptorArrayEXT capability. + // Unless the module has the RuntimeDescriptorArray capability. if (value_type && value_type->opcode() == spv::Op::OpTypeRuntimeArray) { - if (!_.HasCapability(spv::Capability::RuntimeDescriptorArrayEXT)) { + if (!_.HasCapability(spv::Capability::RuntimeDescriptorArray)) { return _.diag(SPV_ERROR_INVALID_ID, inst) << _.VkErrorID(4680) << "OpVariable, " << _.getIdName(inst->id()) << ", is attempting to create memory for an illegal type, " << "OpTypeRuntimeArray.\nFor Vulkan OpTypeRuntimeArray can only " << "appear as the final member of an OpTypeStruct, thus cannot " - << "be instantiated via OpVariable"; + << "be instantiated via OpVariable, unless the " + "RuntimeDescriptorArray Capability is declared"; } else { // A bare variable OpTypeRuntimeArray is allowed in this context, but // still need to check the storage class. @@ -791,7 +793,7 @@ spv_result_t ValidateVariable(ValidationState_t& _, const Instruction* inst) { storage_class != spv::StorageClass::UniformConstant) { return _.diag(SPV_ERROR_INVALID_ID, inst) << _.VkErrorID(4680) - << "For Vulkan with RuntimeDescriptorArrayEXT, a variable " + << "For Vulkan with RuntimeDescriptorArray, a variable " << "containing OpTypeRuntimeArray must have storage class of " << "StorageBuffer, Uniform, or UniformConstant."; } @@ -1118,6 +1120,29 @@ spv_result_t ValidateLoad(ValidationState_t& _, const Instruction* inst) { } } + // Skip checking if there is zero chance for this having a mesh shader + // entrypoint + if (_.HasCapability(spv::Capability::MeshShadingEXT) && + pointer_type->GetOperandAs(1) == + spv::StorageClass::Output) { + std::string errorVUID = _.VkErrorID(7107); + _.function(inst->function()->id()) + ->RegisterExecutionModelLimitation( + [errorVUID](spv::ExecutionModel model, std::string* message) { + // Seems the NV Mesh extension was less strict and allowed + // writting to outputs + if (model == spv::ExecutionModel::MeshEXT) { + if (message) { + *message = errorVUID + + "The Output Storage Class in a Mesh Execution " + "Model must not be read from"; + } + return false; + } + return true; + }); + } + _.RegisterQCOMImageProcessingTextureConsumer(pointer_id, inst, nullptr); return SPV_SUCCESS; @@ -1822,13 +1847,19 @@ spv_result_t ValidateAccessChain(ValidationState_t& _, // At this point, we have fully walked down from the base using the indeces. // The type being pointed to should be the same as the result type. if (type_pointee->id() != result_type_pointee->id()) { + bool same_type = result_type_pointee->opcode() == type_pointee->opcode(); return _.diag(SPV_ERROR_INVALID_ID, inst) - << "Op" << spvOpcodeString(opcode) << " result type (Op" + << "Op" << spvOpcodeString(opcode) << " result type " + << _.getIdName(result_type_pointee->id()) << " (Op" << spvOpcodeString(result_type_pointee->opcode()) << ") does not match the type that results from indexing into the " "base " - " (Op" - << spvOpcodeString(type_pointee->opcode()) << ")."; + " " + << _.getIdName(type_pointee->id()) << " (Op" + << spvOpcodeString(type_pointee->opcode()) << ")." + << (same_type ? " (The types must be the exact same Id, so the " + "two types referenced are slighlty different)" + : ""); } } diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mesh_shading.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mesh_shading.cpp index 3bd1dbd..d7352eb 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mesh_shading.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mesh_shading.cpp @@ -132,9 +132,9 @@ spv_result_t MeshShadingPass(ValidationState_t& _, const Instruction* inst) { } case spv::Op::OpVariable: { if (_.HasCapability(spv::Capability::MeshShadingEXT)) { - bool meshInterfaceVar = + bool is_mesh_interface_var = IsInterfaceVariable(_, inst, spv::ExecutionModel::MeshEXT); - bool fragInterfaceVar = + bool is_frag_interface_var = IsInterfaceVariable(_, inst, spv::ExecutionModel::Fragment); const spv::StorageClass storage_class = @@ -143,14 +143,14 @@ spv_result_t MeshShadingPass(ValidationState_t& _, const Instruction* inst) { bool storage_input = (storage_class == spv::StorageClass::Input); if (_.HasDecoration(inst->id(), spv::Decoration::PerPrimitiveEXT)) { - if (fragInterfaceVar && !storage_input) { + if (is_frag_interface_var && !storage_input) { return _.diag(SPV_ERROR_INVALID_DATA, inst) << "PerPrimitiveEXT decoration must be applied only to " "variables in the Input Storage Class in the Fragment " "Execution Model."; } - if (meshInterfaceVar && !storage_output) { + if (is_mesh_interface_var && !storage_output) { return _.diag(SPV_ERROR_INVALID_DATA, inst) << _.VkErrorID(4336) << "PerPrimitiveEXT decoration must be applied only to " @@ -158,6 +158,20 @@ spv_result_t MeshShadingPass(ValidationState_t& _, const Instruction* inst) { "Storage Class in the MeshEXT Execution Model."; } } + + // This only applies to user interface variables, not built-ins (they + // are validated with the rest of the builtin) + if (is_mesh_interface_var && storage_output && + !_.HasDecoration(inst->id(), spv::Decoration::BuiltIn)) { + const Instruction* pointer_inst = _.FindDef(inst->type_id()); + if (pointer_inst->opcode() == spv::Op::OpTypePointer) { + if (!_.IsArrayType(pointer_inst->word(3))) { + return _.diag(SPV_ERROR_INVALID_DATA, inst) + << "In the MeshEXT Execution Mode, all Output Variables " + "must contain an Array."; + } + } + } } break; } diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mode_setting.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mode_setting.cpp index f2b43b6..22d464f 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mode_setting.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_mode_setting.cpp @@ -27,6 +27,48 @@ namespace spvtools { namespace val { namespace { +// TODO - Make a common util if someone else needs it too outside this file +const char* ExecutionModelToString(spv::ExecutionModel value) { + switch (value) { + case spv::ExecutionModel::Vertex: + return "Vertex"; + case spv::ExecutionModel::TessellationControl: + return "TessellationControl"; + case spv::ExecutionModel::TessellationEvaluation: + return "TessellationEvaluation"; + case spv::ExecutionModel::Geometry: + return "Geometry"; + case spv::ExecutionModel::Fragment: + return "Fragment"; + case spv::ExecutionModel::GLCompute: + return "GLCompute"; + case spv::ExecutionModel::Kernel: + return "Kernel"; + case spv::ExecutionModel::TaskNV: + return "TaskNV"; + case spv::ExecutionModel::MeshNV: + return "MeshNV"; + case spv::ExecutionModel::RayGenerationKHR: + return "RayGenerationKHR"; + case spv::ExecutionModel::IntersectionKHR: + return "IntersectionKHR"; + case spv::ExecutionModel::AnyHitKHR: + return "AnyHitKHR"; + case spv::ExecutionModel::ClosestHitKHR: + return "ClosestHitKHR"; + case spv::ExecutionModel::MissKHR: + return "MissKHR"; + case spv::ExecutionModel::CallableKHR: + return "CallableKHR"; + case spv::ExecutionModel::TaskEXT: + return "TaskEXT"; + case spv::ExecutionModel::MeshEXT: + return "MeshEXT"; + default: + return "Unknown"; + } +} + spv_result_t ValidateEntryPoint(ValidationState_t& _, const Instruction* inst) { const auto entry_point_id = inst->GetOperandAs(1); auto entry_point = _.FindDef(entry_point_id); @@ -306,74 +348,79 @@ spv_result_t ValidateEntryPoint(ValidationState_t& _, const Instruction* inst) { } if (spvIsVulkanEnv(_.context()->target_env)) { - switch (execution_model) { - case spv::ExecutionModel::GLCompute: - if (!has_mode(spv::ExecutionMode::LocalSize)) { - bool ok = has_workgroup_size || has_local_size_id; - if (!ok && _.HasCapability(spv::Capability::TileShadingQCOM)) { - ok = has_mode(spv::ExecutionMode::TileShadingRateQCOM); - } - if (!ok) { - return _.diag(SPV_ERROR_INVALID_DATA, inst) - << _.VkErrorID(10685) - << "In the Vulkan environment, GLCompute execution model " - "entry points require either the " - << (_.HasCapability(spv::Capability::TileShadingQCOM) - ? "TileShadingRateQCOM, " - : "") - << "LocalSize or LocalSizeId execution mode or an object " - "decorated with WorkgroupSize must be specified."; - } + // SPV_QCOM_tile_shading checks + if (execution_model == spv::ExecutionModel::GLCompute) { + if (_.HasCapability(spv::Capability::TileShadingQCOM)) { + if (has_mode(spv::ExecutionMode::TileShadingRateQCOM) && + (has_mode(spv::ExecutionMode::LocalSize) || + has_mode(spv::ExecutionMode::LocalSizeId))) { + return _.diag(SPV_ERROR_INVALID_DATA, inst) + << "If the TileShadingRateQCOM execution mode is used, " + << "LocalSize and LocalSizeId must not be specified."; } - - if (_.HasCapability(spv::Capability::TileShadingQCOM)) { - if (has_mode(spv::ExecutionMode::TileShadingRateQCOM) && - (has_mode(spv::ExecutionMode::LocalSize) || - has_mode(spv::ExecutionMode::LocalSizeId))) { - return _.diag(SPV_ERROR_INVALID_DATA, inst) - << "If the TileShadingRateQCOM execution mode is used, " - << "LocalSize and LocalSizeId must not be specified."; - } - if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) { - return _.diag(SPV_ERROR_INVALID_DATA, inst) - << "The NonCoherentTileAttachmentQCOM execution mode must " - "not be used in any stage other than fragment."; - } - } else { - if (has_mode(spv::ExecutionMode::TileShadingRateQCOM)) { - return _.diag(SPV_ERROR_INVALID_DATA, inst) - << "If the TileShadingRateQCOM execution mode is used, the " - "TileShadingQCOM capability must be enabled."; - } + if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) { + return _.diag(SPV_ERROR_INVALID_DATA, inst) + << "The NonCoherentTileAttachmentQCOM execution mode must " + "not be used in any stage other than fragment."; } - break; - default: + } else { if (has_mode(spv::ExecutionMode::TileShadingRateQCOM)) { return _.diag(SPV_ERROR_INVALID_DATA, inst) - << "The TileShadingRateQCOM execution mode must not be used " - "in any stage other than compute."; + << "If the TileShadingRateQCOM execution mode is used, the " + "TileShadingQCOM capability must be enabled."; } - if (execution_model != spv::ExecutionModel::Fragment) { - if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) { + } + } else { + if (has_mode(spv::ExecutionMode::TileShadingRateQCOM)) { + return _.diag(SPV_ERROR_INVALID_DATA, inst) + << "The TileShadingRateQCOM execution mode must not be used " + "in any stage other than compute."; + } + if (execution_model != spv::ExecutionModel::Fragment) { + if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) { + return _.diag(SPV_ERROR_INVALID_DATA, inst) + << "The NonCoherentTileAttachmentQCOM execution mode must " + "not be used in any stage other than fragment."; + } + if (_.HasCapability(spv::Capability::TileShadingQCOM)) { + return _.diag(SPV_ERROR_INVALID_CAPABILITY, inst) + << "The TileShadingQCOM capability must not be enabled in " + "any stage other than compute or fragment."; + } + } else { + if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) { + if (!_.HasCapability(spv::Capability::TileShadingQCOM)) { return _.diag(SPV_ERROR_INVALID_DATA, inst) - << "The NonCoherentTileAttachmentQCOM execution mode must " - "not be used in any stage other than fragment."; - } - if (_.HasCapability(spv::Capability::TileShadingQCOM)) { - return _.diag(SPV_ERROR_INVALID_CAPABILITY, inst) - << "The TileShadingQCOM capability must not be enabled in " - "any stage other than compute or fragment."; - } - } else { - if (has_mode(spv::ExecutionMode::NonCoherentTileAttachmentReadQCOM)) { - if (!_.HasCapability(spv::Capability::TileShadingQCOM)) { - return _.diag(SPV_ERROR_INVALID_DATA, inst) - << "If the NonCoherentTileAttachmentReadQCOM execution " - "mode is used, the TileShadingQCOM capability must be " - "enabled."; - } + << "If the NonCoherentTileAttachmentReadQCOM execution " + "mode is used, the TileShadingQCOM capability must be " + "enabled."; } } + } + } + + switch (execution_model) { + case spv::ExecutionModel::GLCompute: + case spv::ExecutionModel::MeshEXT: + case spv::ExecutionModel::MeshNV: + case spv::ExecutionModel::TaskEXT: + case spv::ExecutionModel::TaskNV: + if (!has_mode(spv::ExecutionMode::LocalSize) && !has_workgroup_size && + !has_local_size_id && + !has_mode(spv::ExecutionMode::TileShadingRateQCOM)) { + return _.diag(SPV_ERROR_INVALID_DATA, inst) + << _.VkErrorID(10685) << "In the Vulkan environment, " + << ExecutionModelToString(execution_model) + << " execution model " + "entry points require either the " + << (_.HasCapability(spv::Capability::TileShadingQCOM) + ? "TileShadingRateQCOM, " + : "") + << "LocalSize or LocalSizeId execution mode or an object " + "decorated with WorkgroupSize must be specified."; + } + break; + default: break; } } diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_type.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_type.cpp index 786a224..12c1ef0 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validate_type.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validate_type.cpp @@ -737,9 +737,10 @@ spv_result_t ValidateTypeCooperativeMatrix(ValidationState_t& _, } } - uint64_t scope_value; - if (_.EvalConstantValUint64(scope_id, &scope_value)) { - if (scope_value == static_cast(spv::Scope::Workgroup)) { + uint64_t scope_raw_value; + if (_.EvalConstantValUint64(scope_id, &scope_raw_value)) { + spv::Scope scope_value = static_cast(scope_raw_value); + if (scope_value == spv::Scope::Workgroup) { for (auto entry_point_id : _.entry_points()) { if (!_.EntryPointHasLocalSizeOrId(entry_point_id)) { return _.diag(SPV_ERROR_INVALID_ID, inst) @@ -766,6 +767,13 @@ spv_result_t ValidateTypeCooperativeMatrix(ValidationState_t& _, } } } + if (scope_value != spv::Scope::Workgroup && + scope_value != spv::Scope::Subgroup) { + return _.diag(SPV_ERROR_INVALID_DATA, inst) + << _.VkErrorID(12243) + << "OpTypeCooperativeMatrixKHR Scope is limited to Workgroup and " + "Subgroup"; + } } return SPV_SUCCESS; diff --git a/libs/bgfx/3rdparty/spirv-tools/source/val/validation_state.cpp b/libs/bgfx/3rdparty/spirv-tools/source/val/validation_state.cpp index bc9d835..cf2ff3b 100644 --- a/libs/bgfx/3rdparty/spirv-tools/source/val/validation_state.cpp +++ b/libs/bgfx/3rdparty/spirv-tools/source/val/validation_state.cpp @@ -2703,6 +2703,8 @@ std::string ValidationState_t::VkErrorID(uint32_t id, return VUID_WRAP(VUID-ViewportIndex-ViewportIndex-07060); case 7102: return VUID_WRAP(VUID-StandaloneSpirv-MeshEXT-07102); + case 7107: + return VUID_WRAP(VUID-StandaloneSpirv-MeshEXT-07107); case 7290: return VUID_WRAP(VUID-StandaloneSpirv-Input-07290); case 7320: @@ -2813,6 +2815,8 @@ std::string ValidationState_t::VkErrorID(uint32_t id, return VUID_WRAP(VUID-StandaloneSpirv-OpUntypedVariableKHR-11167); case 11805: return VUID_WRAP(VUID-StandaloneSpirv-OpArrayLength-11805); + case 12243: + return VUID_WRAP(VUID-StandaloneSpirv-Scope-12243); default: return ""; // unknown id } diff --git a/libs/bgfx/src/bgfx.cpp b/libs/bgfx/src/bgfx.cpp index a26d7f6..d899627 100644 --- a/libs/bgfx/src/bgfx.cpp +++ b/libs/bgfx/src/bgfx.cpp @@ -1575,7 +1575,7 @@ namespace bgfx void UniformBuffer::writeUniform(UniformType::Enum _type, uint16_t _loc, const void* _value, uint16_t _num) { - const uint32_t opcode = encodeOpcode(bx::narrowCast(_type), _loc, _num, true); + const uint32_t opcode = encodeOpcode(uint8_t(_type), _loc, _num, true); write(opcode); write(_value, g_uniformTypeSize[_type]*_num); } diff --git a/libs/bgfx/src/bgfx_p.h b/libs/bgfx/src/bgfx_p.h index 55e4eca..e2b3167 100644 --- a/libs/bgfx/src/bgfx_p.h +++ b/libs/bgfx/src/bgfx_p.h @@ -287,26 +287,16 @@ namespace bgfx { constexpr uint32_t kChunkMagicTex = BX_MAKEFOURCC('T', 'E', 'X', 0x0); - inline constexpr uint32_t toAbgr8(uint8_t _r, uint8_t _g, uint8_t _b, uint8_t _a = 0xff) - { - return 0 - | (uint32_t(_r) << 24) - | (uint32_t(_g) << 16) - | (uint32_t(_b) << 8) - | (uint32_t(_a)) - ; - } - // Palette: // https://colorkit.co/color-palette-generator/a8e6cf-dcedc1-ffd3b6-76b4bd-bdeaee-8874a3-ff0000-ff8b94/ - constexpr uint32_t kColorFrame = toAbgr8(0xa8, 0xe6, 0xcf); - constexpr uint32_t kColorSubmit = toAbgr8(0xdc, 0xed, 0xc1); - constexpr uint32_t kColorView = toAbgr8(0xff, 0xd3, 0xb6); - constexpr uint32_t kColorDraw = toAbgr8(0x76, 0xb4, 0xbd); - constexpr uint32_t kColorCompute = toAbgr8(0xbd, 0xea, 0xee); - constexpr uint32_t kColorResource = toAbgr8(0x88, 0x74, 0xa3); - constexpr uint32_t kColorMarker = toAbgr8(0xff, 0x00, 0x00); - constexpr uint32_t kColorWait = toAbgr8(0xff, 0x8b, 0x94); + constexpr uint32_t kColorFrame = 0xa8'e6'cf'ff; + constexpr uint32_t kColorSubmit = 0xdc'ed'c1'ff; + constexpr uint32_t kColorView = 0xff'd3'b6'ff; + constexpr uint32_t kColorDraw = 0x76'b4'bd'ff; + constexpr uint32_t kColorCompute = 0xbd'ea'ee'ff; + constexpr uint32_t kColorResource = 0x88'74'a3'ff; + constexpr uint32_t kColorMarker = 0xff'00'00'ff; + constexpr uint32_t kColorWait = 0xff'8b'94'ff; extern InternalData g_internalData; extern PlatformData g_platformData; @@ -627,17 +617,17 @@ namespace bgfx release( (const Memory*)_mem); } - inline uint64_t packStencil(uint32_t _fstencil, uint32_t _bstencil) + inline constexpr uint64_t packStencil(uint32_t _fstencil, uint32_t _bstencil) { return (uint64_t(_bstencil)<<32)|uint64_t(_fstencil); } - inline uint32_t unpackStencil(uint8_t _0or1, uint64_t _stencil) + inline constexpr uint32_t unpackStencil(uint8_t _0or1, uint64_t _stencil) { return uint32_t( (_stencil >> (32*_0or1) ) ); } - inline bool needBorderColor(uint64_t _flags) + inline constexpr bool needBorderColor(uint64_t _flags) { return BGFX_SAMPLER_U_BORDER == (_flags & BGFX_SAMPLER_U_BORDER) || BGFX_SAMPLER_V_BORDER == (_flags & BGFX_SAMPLER_V_BORDER) @@ -645,7 +635,7 @@ namespace bgfx ; } - inline uint8_t calcNumMips(bool _hasMips, uint16_t _width, uint16_t _height, uint16_t _depth = 1) + inline constexpr uint8_t calcNumMips(bool _hasMips, uint16_t _width, uint16_t _height, uint16_t _depth = 1) { if (_hasMips) { @@ -2652,7 +2642,10 @@ namespace bgfx // clear all bytes (inclusively the padding) before we start. bx::memSet(&m_bind, 0, sizeof(m_bind) ); - discard(BGFX_DISCARD_ALL); + m_discard = false; + m_draw.clear(BGFX_DISCARD_ALL); + m_compute.clear(BGFX_DISCARD_ALL); + m_bind.clear(BGFX_DISCARD_ALL); } void begin(Frame* _frame, uint8_t _idx) @@ -2991,6 +2984,13 @@ namespace bgfx m_draw.clear(_flags); m_compute.clear(_flags); m_bind.clear(_flags); + + if (_flags & BGFX_DISCARD_STATE) + { + UniformBuffer* uniformBuffer = m_frame->m_uniformBuffer[m_uniformIdx]; + m_uniformEnd = uniformBuffer->getPos(); + m_uniformBegin = m_uniformEnd; + } } void submit(ViewId _id, ProgramHandle _program, OcclusionQueryHandle _occlusionQuery, uint32_t _depth, uint8_t _flags); diff --git a/libs/bgfx/src/renderer_gl.cpp b/libs/bgfx/src/renderer_gl.cpp index 1f14995..aa08e24 100644 --- a/libs/bgfx/src/renderer_gl.cpp +++ b/libs/bgfx/src/renderer_gl.cpp @@ -5336,7 +5336,7 @@ namespace bgfx { namespace gl } UniformType::Enum type = convertGlType(gltype); - m_constantBuffer->writeUniformHandle(bx::narrowCast(type), 0, info->m_handle, uint16_t(num) ); + m_constantBuffer->writeUniformHandle(uint8_t(type), 0, info->m_handle, uint16_t(num) ); m_constantBuffer->write(loc); BX_TRACE("store %s %d", name, info->m_handle); } diff --git a/libs/bx/include/bx/bx.h b/libs/bx/include/bx/bx.h index 83b721f..a45057d 100644 --- a/libs/bx/include/bx/bx.h +++ b/libs/bx/include/bx/bx.h @@ -224,6 +224,16 @@ namespace bx template constexpr Ty bitCast(const FromT& _from); + /// Performs `static_cast` of value `_from`, and if value doesn't fit result type `Ty` it clamps + /// the value to `Ty` min/max. + template + constexpr Ty saturateCast(FromT _from); + + /// Performs `static_cast` of value `_from`, and returns true if the value `_from` is + /// representable as `Ty`. + template + constexpr bool narrowCastTest(Ty* _out, const FromT& _from); + /// Performs `static_cast` of value `_from`, and in debug build runtime verifies/asserts /// that the value didn't change. template diff --git a/libs/bx/include/bx/inline/bx.inl b/libs/bx/include/bx/inline/bx.inl index e5f1214..1bd971c 100644 --- a/libs/bx/include/bx/inline/bx.inl +++ b/libs/bx/include/bx/inline/bx.inl @@ -171,9 +171,62 @@ namespace bx return __builtin_bit_cast(Ty, _from); } + template + requires (isInteger< Ty>() || isFloatingPoint< Ty>() ) + && (isInteger() || isFloatingPoint() ) + inline constexpr Ty saturateCast(FromT _from) + { + if constexpr (isSame, RemoveCvType >() ) + { + return _from; + } + + constexpr Ty mx = max(); + + if constexpr (isSigned() && isSigned() ) + { + if constexpr (sizeof(Ty) < sizeof(FromT) ) + { + constexpr FromT mn = min(); + + if (_from < mn) + { + return mn; + } + else if (_from > mx) + { + return mx; + } + } + } + else if constexpr (isSigned() ) + { + if (_from < FromT(0) ) + { + return Ty(0); + } + else if (asUnsigned(_from) > mx) + { + return mx; + } + } + else if (_from > asUnsigned(max() ) ) + { + return mx; + } + + return static_cast(_from); + } + template inline constexpr bool narrowCastTest(Ty* _out, const FromT& _from) { + if constexpr (isSame() ) + { + *_out = _from; + return true; + } + *_out = static_cast(_from); return static_cast(*_out) == _from; } @@ -181,10 +234,13 @@ namespace bx template inline Ty narrowCast(const FromT& _from, Location _location) { - Ty to = static_cast(_from); - BX_ASSERT_LOC(_location, static_cast(to) == _from + Ty to; + const bool result = narrowCastTest(&to, _from); + + BX_ASSERT_LOC(_location, result , "bx::narrowCast failed! Value is truncated!" ); + BX_UNUSED(result); return to; } diff --git a/libs/bx/include/bx/thread.h b/libs/bx/include/bx/thread.h index 62b1a3e..57ae304 100644 --- a/libs/bx/include/bx/thread.h +++ b/libs/bx/include/bx/thread.h @@ -6,8 +6,8 @@ #ifndef BX_THREAD_H_HEADER_GUARD #define BX_THREAD_H_HEADER_GUARD -#include "allocator.h" #include "mpscqueue.h" +#include "string.h" #if BX_CONFIG_SUPPORTS_THREADING @@ -39,7 +39,7 @@ namespace bx /// @param[in] _name Thread name used by debugger. /// @returns True if thread is created, otherwise returns false. /// - bool init(ThreadFn _fn, void* _userData = NULL, uint32_t _stackSize = 0, const char* _name = NULL); + bool init(ThreadFn _fn, void* _userData = NULL, uint32_t _stackSize = 0, const StringView& _name = ""); /// void shutdown(); @@ -51,7 +51,7 @@ namespace bx int32_t getExitCode() const; /// - void setThreadName(const char* _name); + void setThreadName(const StringView& _name); /// void push(void* _ptr); @@ -72,7 +72,7 @@ namespace bx uint32_t m_stackSize; int32_t m_exitCode; bool m_running; - char m_name[64]; + FixedString64 m_name; }; /// diff --git a/libs/bx/src/thread.cpp b/libs/bx/src/thread.cpp index f3a31a7..53ce517 100644 --- a/libs/bx/src/thread.cpp +++ b/libs/bx/src/thread.cpp @@ -109,23 +109,14 @@ namespace bx } } - bool Thread::init(ThreadFn _fn, void* _userData, uint32_t _stackSize, const char* _name) + bool Thread::init(ThreadFn _fn, void* _userData, uint32_t _stackSize, const StringView& _name) { BX_ASSERT(!m_running, "Already running!"); - m_fn = _fn; - m_userData = _userData; + m_fn = _fn; + m_userData = _userData; m_stackSize = _stackSize; - - if (NULL != _name) - { - BX_WARN(strLen(_name) < int32_t(BX_COUNTOF(m_name) )-1, "Truncating thread name."); - strCopy(m_name, BX_COUNTOF(m_name), _name); - } - else - { - m_name[0] = '\0'; - } + m_name = _name; ThreadInternal* ti = (ThreadInternal*)m_internal; #if BX_CRT_NONE @@ -223,40 +214,40 @@ namespace bx return m_exitCode; } - void Thread::setThreadName(const char* _name) + void Thread::setThreadName(const StringView& _name) { - if (NULL == _name - || 0 == strLen(_name) ) + if (_name.isEmpty() ) { return; } + m_name = _name; + ThreadInternal* ti = (ThreadInternal*)m_internal; BX_UNUSED(ti); -#if BX_CRT_NONE - BX_UNUSED(_name); -#elif BX_PLATFORM_OSX \ - || BX_PLATFORM_IOS \ +#if BX_PLATFORM_OSX \ + || BX_PLATFORM_IOS \ || BX_PLATFORM_VISIONOS - pthread_setname_np(_name); + pthread_setname_np(m_name.getCPtr() ); #elif BX_CRT_GLIBC - pthread_setname_np(ti->m_handle, _name); + pthread_setname_np(ti->m_handle, m_name.getCPtr() ); #elif BX_PLATFORM_LINUX - prctl(PR_SET_NAME,_name, 0, 0, 0); + prctl(PR_SET_NAME, m_name.getCPtr(), 0, 0, 0); #elif BX_PLATFORM_WINDOWS - // Try to use the new thread naming API from Win10 Creators update onwards if we have it typedef HRESULT (WINAPI *SetThreadDescriptionProc)(HANDLE, PCWSTR); - SetThreadDescriptionProc SetThreadDescription = dlsym((void*)GetModuleHandleA("Kernel32.dll"), "SetThreadDescription"); + SetThreadDescriptionProc SetThreadDescription = dlsym( (void*)GetModuleHandleA("Kernel32.dll"), "SetThreadDescription"); if (NULL != SetThreadDescription) { - uint32_t length = (uint32_t)strLen(_name)+1; - uint32_t size = length*sizeof(wchar_t); - wchar_t* name = (wchar_t*)BX_STACK_ALLOC(size); - mbstowcs(name, _name, size-2); - name[size-2] = 0; + const uint32_t length = m_name.getLength(); + const uint32_t max = (length+1)*sizeof(wchar_t); + wchar_t* name = (wchar_t*)BX_STACK_ALLOC(max); + mbstowcs(name, m_name.getCPtr(), length); + name[length] = 0; SetThreadDescription(ti->m_handle, name); } + else + { # if BX_COMPILER_MSVC # pragma pack(push, 8) struct ThreadName @@ -269,24 +260,23 @@ namespace bx # pragma pack(pop) ThreadName tn; tn.type = 0x1000; - tn.name = _name; + tn.name = m_name.getCPtr(); tn.id = ti->m_threadId; tn.flags = 0; __try { RaiseException(0x406d1388 - , 0 - , sizeof(tn)/4 - , reinterpret_cast(&tn) - ); + , 0 + , sizeof(tn)/4 + , reinterpret_cast(&tn) + ); } __except(EXCEPTION_EXECUTE_HANDLER) { } # endif // BX_COMPILER_MSVC -#else - BX_UNUSED(_name); + } #endif // BX_PLATFORM_ }