Mirrors/bgfx
1
0
Fork 0
mirror of https://github.com/bkaradzic/bgfx.git synced 2026-02-17 20:52:36 +01:00
Code Issues Packages Projects Releases Wiki Activity

Updated spirv-tools.

Browse Source
This commit is contained in:
Бранимир Караџић
2025-12-13 11:48:21 -08:00
committed by Branimir Karadžić
parent 3df341ac5e
commit 1a4e866800
26 changed files with 1145 additions and 393 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
3rdparty/spirv-tools/include/generated/build-version.inc vendored
View File

@@ -1 +1 @@
"v2025.5", "SPIRV-Tools v2025.5 v2025.4-64-gd2a11ec9"
"v2025.5", "SPIRV-Tools v2025.5 v2025.5.rc1-32-g6e7423bc"

4
3rdparty/spirv-tools/source/diff/diff.cpp vendored
View File

@@ -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() {

133
3rdparty/spirv-tools/source/opt/aggressive_dead_code_elim_pass.cpp vendored
View File

@@ -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,40 +291,95 @@ Pass::Status AggressiveDCEPass::ProcessDebugInformation(
std::list<BasicBlock*>& 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;
}
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;
}
// 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());
}
}
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;

11
3rdparty/spirv-tools/source/opt/cfg_cleanup_pass.cpp vendored
View File

@@ -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;
}

63
3rdparty/spirv-tools/source/opt/combine_access_chains.cpp vendored
View File

@@ -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<Operand>* new_operands) {
Pass::Status CombineAccessChains::CombineIndices(
Instruction* ptr_input, Instruction* inst,
std::vector<Operand>* 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<Operand>* 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<Operand> 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,

12
3rdparty/spirv-tools/source/opt/combine_access_chains.h vendored
View File

@@ -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<Operand>* new_operands);
Status CreateNewInputOperands(Instruction* ptr_input, Instruction* inst,
std::vector<Operand>* 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<Operand>* new_operands);
Status CombineIndices(Instruction* ptr_input, Instruction* inst,
std::vector<Operand>* new_operands);
// Returns the opcode to use for the combined access chain.
spv::Op UpdateOpcode(spv::Op base_opcode, spv::Op input_opcode);

194
3rdparty/spirv-tools/source/opt/const_folding_rules.cpp vendored
View File

@@ -1126,6 +1126,26 @@ ConstantFoldingRule FoldFUnordGreaterThanEqual() {
return FoldFPBinaryOp(FOLD_FPCMP_OP(>=, false));
}
ConstantFoldingRule FoldInvariantSelect() {
return [](IRContext*, Instruction* inst,
const std::vector<const analysis::Constant*>& 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 <typename FloatType>
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<const analysis::Constant*>& 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<const analysis::Constant*>& 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<const analysis::Constant*>& 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(

36
3rdparty/spirv-tools/source/opt/copy_prop_arrays.cpp vendored
View File

@@ -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<uint32_t> 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

8
3rdparty/spirv-tools/source/opt/copy_prop_arrays.h vendored
View File

@@ -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

12
3rdparty/spirv-tools/source/opt/debug_info_manager.cpp vendored
View File

@@ -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<Instruction> 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<Instruction> 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);

461
3rdparty/spirv-tools/source/opt/folding_rules.cpp vendored
View File

@@ -16,6 +16,7 @@
#include <limits>
#include <memory>
#include <optional>
#include <utility>
#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<const analysis::Constant*>&) {
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<const analysis::Constant*>& constants) {
const std::vector<const analysis::Constant*>&) {
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())
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 (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;
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;
};
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<const analysis::Constant*>& 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());
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(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;
}
bool uses_float = HasFloatingPoint(type);
if (uses_float && !inst->IsFloatingPointFoldingAllowed()) return false;
uint32_t width = ElementWidth(type);
if (width != 32 && width != 64) 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())
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);
}
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<const analysis::Constant*>&) {
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<const analysis::Constant*>&) {
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<const analysis::Constant*>& constants) {
@@ -2410,6 +2479,250 @@ FoldingRule RedundantSelect() {
};
}
std::optional<bool> 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<const analysis::Constant*>& 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<bool> uniform_true = GetBoolConstantKind(constants[1]);
std::optional<bool> 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<const analysis::Constant*>& constants) {
assert(inst->opcode() == spv::Op::OpLogicalAnd);
if (GetBoolConstantKind(ConstInput(constants)) ==
std::optional<bool>(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<const analysis::Constant*>& constants) {
assert(inst->opcode() == spv::Op::OpLogicalOr);
if (GetBoolConstantKind(ConstInput(constants)) ==
std::optional<bool>(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<const analysis::Constant*>&) {
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<const analysis::Constant*>&) {
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());

16
3rdparty/spirv-tools/source/opt/ir_context.cpp vendored
View File

@@ -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) {

3
3rdparty/spirv-tools/source/opt/ir_context.h vendored
View File

@@ -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<uint32_t>::max());

24
3rdparty/spirv-tools/source/opt/mem_pass.cpp vendored
View File

@@ -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<BasicBlock*>& reachable_blocks) {
std::vector<Operand> 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) {

8
3rdparty/spirv-tools/source/opt/mem_pass.h vendored
View File

@@ -114,7 +114,7 @@ class MemPass : public Pass {
void DCEInst(Instruction* inst, const std::function<void(Instruction*)>&);
// 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<BasicBlock*>& reachable_blocks);

22
3rdparty/spirv-tools/source/opt/pass.cpp vendored
View File

@@ -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<uint32_t> 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

30
3rdparty/spirv-tools/source/opt/value_number_table.cpp vendored
View File

@@ -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.

15
3rdparty/spirv-tools/source/val/validate_atomics.cpp vendored
View File

@@ -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.";
}

201
3rdparty/spirv-tools/source/val/validate_builtins.cpp vendored
View File

@@ -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<uint32_t, uint32_t> 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<spv::BuiltIn, MeshBuiltinVUIDs>
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<uint32_t, uint32_t> 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);

22
3rdparty/spirv-tools/source/val/validate_decorations.cpp vendored
View File

@@ -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 "

7
3rdparty/spirv-tools/source/val/validate_logical_pointers.cpp vendored
View File

@@ -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<const Instruction*> 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<const Instruction*> 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) {

45
3rdparty/spirv-tools/source/val/validate_memory.cpp vendored
View File

@@ -15,6 +15,7 @@
// limitations under the License.
#include <algorithm>
#include <cstdint>
#include <string>
#include <vector>
@@ -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, <id> "
<< _.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<spv::StorageClass>(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 <id> "
<< _.getIdName(result_type_pointee->id()) << " (Op"
<< spvOpcodeString(result_type_pointee->opcode())
<< ") does not match the type that results from indexing into the "
"base "
"<id> (Op"
<< spvOpcodeString(type_pointee->opcode()) << ").";
"<id> "
<< _.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)"
: "");
}
}

22
3rdparty/spirv-tools/source/val/validate_mesh_shading.cpp vendored
View File

@@ -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;
}

169
3rdparty/spirv-tools/source/val/validate_mode_setting.cpp vendored
View File

@@ -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<uint32_t>(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) {
// 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 (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.";
}
}
} 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)
<< _.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.";
<< "If the NonCoherentTileAttachmentReadQCOM execution "
"mode is used, the TileShadingQCOM capability must be "
"enabled.";
}
}
}
}
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.";
}
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:
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)
<< "If the NonCoherentTileAttachmentReadQCOM execution "
"mode is used, the TileShadingQCOM capability must be "
"enabled.";
}
}
}
break;
}
}

14
3rdparty/spirv-tools/source/val/validate_type.cpp vendored
View File

@@ -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<uint32_t>(spv::Scope::Workgroup)) {
uint64_t scope_raw_value;
if (_.EvalConstantValUint64(scope_id, &scope_raw_value)) {
spv::Scope scope_value = static_cast<spv::Scope>(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;

4
3rdparty/spirv-tools/source/val/validation_state.cpp vendored
View File

@@ -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
}