From 48520bd5e6344d792840aa37fe1ba5d564232788 Mon Sep 17 00:00:00 2001 From: dpankratz Date: Wed, 27 Jan 2021 09:18:00 -0700 Subject: [PATCH] spirv-opt: Add auto-inst passes --- Android.mk | 8 + include/spirv-tools/instrument.hpp | 39 + include/spirv-tools/optimizer.hpp | 98 ++ source/opt/CMakeLists.txt | 18 +- source/opt/auto_inst_debug_pass.cpp | 68 ++ source/opt/auto_inst_debug_pass.h | 106 ++ ..._inst_divergence_characterization_pass.cpp | 134 +++ ...to_inst_divergence_characterization_pass.h | 106 ++ .../opt/auto_inst_dyn_shader_trace_pass.cpp | 52 + source/opt/auto_inst_dyn_shader_trace_pass.h | 103 ++ .../auto_inst_dyn_trace_ray_trace_pass.cpp | 184 ++++ .../opt/auto_inst_dyn_trace_ray_trace_pass.h | 135 +++ source/opt/auto_inst_execution_trace_pass.cpp | 68 ++ source/opt/auto_inst_execution_trace_pass.h | 122 +++ source/opt/auto_inst_pass.cpp | 927 ++++++++++++++++++ source/opt/auto_inst_pass.h | 322 ++++++ source/opt/auto_inst_simt_efficiency_pass.cpp | 39 + source/opt/auto_inst_simt_efficiency_pass.h | 101 ++ .../auto_inst_warp_entry_and_exit_pass.cpp | 43 + .../opt/auto_inst_warp_entry_and_exit_pass.h | 99 ++ source/opt/instrument_pass.cpp | 5 +- source/opt/instrument_pass.h | 1 + source/opt/ir_builder.h | 26 + source/opt/ir_context.cpp | 13 + source/opt/optimizer.cpp | 60 ++ source/opt/passes.h | 7 + 26 files changed, 2882 insertions(+), 2 deletions(-) create mode 100644 source/opt/auto_inst_debug_pass.cpp create mode 100644 source/opt/auto_inst_debug_pass.h create mode 100644 source/opt/auto_inst_divergence_characterization_pass.cpp create mode 100644 source/opt/auto_inst_divergence_characterization_pass.h create mode 100644 source/opt/auto_inst_dyn_shader_trace_pass.cpp create mode 100644 source/opt/auto_inst_dyn_shader_trace_pass.h create mode 100644 source/opt/auto_inst_dyn_trace_ray_trace_pass.cpp create mode 100644 source/opt/auto_inst_dyn_trace_ray_trace_pass.h create mode 100644 source/opt/auto_inst_execution_trace_pass.cpp create mode 100644 source/opt/auto_inst_execution_trace_pass.h create mode 100644 source/opt/auto_inst_pass.cpp create mode 100644 source/opt/auto_inst_pass.h create mode 100644 source/opt/auto_inst_simt_efficiency_pass.cpp create mode 100644 source/opt/auto_inst_simt_efficiency_pass.h create mode 100644 source/opt/auto_inst_warp_entry_and_exit_pass.cpp create mode 100644 source/opt/auto_inst_warp_entry_and_exit_pass.h diff --git a/Android.mk b/Android.mk index 0b64ea6d..46728c67 100644 --- a/Android.mk +++ b/Android.mk @@ -73,6 +73,14 @@ SPVTOOLS_SRC_FILES := \ source/val/validate_type.cpp SPVTOOLS_OPT_SRC_FILES := \ + source/opt/auto_inst_pass.cpp \ + source/opt/auto_inst_debug_pass.cpp \ + source/opt/auto_inst_divergence_characterization_pass.cpp \ + source/opt/auto_inst_dyn_trace_ray_trace_pass.cpp \ + source/opt/auto_inst_dyn_shader_trace_pass.cpp \ + source/opt/auto_inst_execution_trace_pass.cpp \ + source/opt/auto_inst_simt_efficiency_pass.cpp \ + source/opt/auto_inst_warp_entry_and_exit_pass.cpp \ source/opt/aggressive_dead_code_elim_pass.cpp \ source/opt/amd_ext_to_khr.cpp \ source/opt/basic_block.cpp \ diff --git a/include/spirv-tools/instrument.hpp b/include/spirv-tools/instrument.hpp index 2b47a564..5a2b00c9 100644 --- a/include/spirv-tools/instrument.hpp +++ b/include/spirv-tools/instrument.hpp @@ -250,6 +250,45 @@ static const int kDebugInputBuffAddrPtrOffset = 1; // not a valid buffer, the length associated with the 0x0 address is zero. static const int kDebugInputBuffAddrLengthOffset = 0; +// Auto-Inst Bindings +// +// These bindings are used primarily to differentiate different instrumentation +// primitives which have meaning to the instrumentation generating them +// and the analysis parsing them. +// +// NOTE: do not set these values to 0 since, by default, the instrumentation +// buffer is 0 initialized and would cause pernicious bugs. + +static const int kAutoInstUniqueSubgroupId = 1; + +// auto_inst_divergence_characterization +// Primitive to track the beginning of a traceRay +static const uint32_t kAutoInstDivCharPreTraceRay = 2; +// Primitive to track the end of a traceRay +static const uint32_t kAutoInstDivCharPostTraceRay = 3; +// Primitive to get the active threads in a basic block +static const uint32_t kAutoInstDivCharActiveThreads = 4; +// Primtive to get the active threads at the +// beginning of a ray tracing shader +static const uint32_t kAutoInstDivCharShaderEntryPoint = 5; +// Primitive to track when threads have finished executing +// the pipeline. +static const uint32_t kAutoInstDivCharQuitPipeline = 6; + +// auto_inst_dyn_shader_trace +// Track when shader is executed +static const uint32_t kAutoInstDynShaderTraceEntryPoint = 2; + +// auto_inst_dyn_trace_ray_trace +// Track when traceRay call started +static const uint32_t kAutoInstTraceRayTracePreTraceRay = 2; +// Track when merge point after traceRay is executed +static const uint32_t kAutoInstTraceRayTraceMergePoint = 3; + +// auto_inst_warp_entry_and_exit +static const uint32_t kAutoInstWarpEntryAndExitBeginPipeline = 1; +static const uint32_t kAutoInstWarpEntryAndExitEndPipeline = 2; + } // namespace spvtools #endif // INCLUDE_SPIRV_TOOLS_INSTRUMENT_HPP_ diff --git a/include/spirv-tools/optimizer.hpp b/include/spirv-tools/optimizer.hpp index 27352b25..a8cc6d34 100644 --- a/include/spirv-tools/optimizer.hpp +++ b/include/spirv-tools/optimizer.hpp @@ -15,8 +15,10 @@ #ifndef INCLUDE_SPIRV_TOOLS_OPTIMIZER_HPP_ #define INCLUDE_SPIRV_TOOLS_OPTIMIZER_HPP_ +#include #include #include +#include #include #include #include @@ -792,6 +794,102 @@ Optimizer::PassToken CreateInstBuffAddrCheckPass(uint32_t desc_set, Optimizer::PassToken CreateInstDebugPrintfPass(uint32_t desc_set, uint32_t shader_id); +// Create a pass which will help debug problems with other autoinstrumentation +// passes. +// +// If |test_atomic_ops| is true, then an atomicAdd will be added to the shader +// and its result will be written to the buffer given by |desc_set|. If +// |test_subgroup_ops| is true, then a subgroupElect will be added to the shader +// and its result will be written to the buffer. +// If both are false then a constant will be written to the buffer. +Optimizer::PassToken CreateAutoInstDebugPass(uint32_t desc_set, + uint32_t shader_id, + bool test_atomic_ops, + bool test_subgroup_ops); + +// Create a pass which will automatically insert instrumentation to +// capture the extent of different contributors to divergence. +// +// The instrumentation will write buffers in debug descriptor set |desc_set|. +// It will write |shader_id| in each output record to identify the shader +// module which generated the record if necessary. +Optimizer::PassToken CreateAutoInstDivergenceCharacterizationPass( + uint32_t desc_set, uint32_t shader_id, + std::function< + void(std::unordered_map&& inst_id2prim_id, + std::unordered_map&& inst_id2inst_count)> + static_data_callback); + +// Create a pass which will automatically insert instrumentation to +// determine the runtime execution counts of each shader. +// +// The instrumentation will write buffers in debug descriptor set |desc_set|. +// It will write |shader_id| in each output record to identify the shader +// module which generated the record. +Optimizer::PassToken CreateAutoInstDynShaderTracePass(uint32_t desc_set, + uint32_t shader_id); + +// Create a pass which will automatically insert instrumentation to +// disambiguate runtime traceRay calls found within control-flow. +// +// The instrumentation will write buffers in debug descriptor set |desc_set|. +// It will write |shader_id| in each output record to identify the shader +// module which generated the record if necessary. +// +// The |static_data_callback| is called after the instrumentation pass has +// finished. It is populated with a mapping from instrumentation callsite id +// to instrumentation type. It is also populated with a mapping from merge point +// to all the traceRay calls sites that could have executed within the +// control-flow. +Optimizer::PassToken CreateAutoInstDynTraceRayTracePass( + uint32_t desc_set, uint32_t shader_id, + std::function&&, + std::unordered_map>&&)> + static_data_callback); + +// Create a pass which will automatically insert instrumentation to +// compute the simt efficiency of the shader module. +// +// The instrumentation will write buffers in debug descriptor set |desc_set|. +// |reserved_words_count| is the number of lower words in the buffer that +// have a fixed function and are reserved. +// +// It will write |shader_id| in each output record to identify the shader +// module which generated the record if necessary. +Optimizer::PassToken CreateAutoInstSimtEfficiencyPass( + uint32_t desc_set, uint32_t shader_id, uint32_t reserved_words_count); + +// Create a pass which will automatically insert instrumentation to +// capture the number of times the ray tracing pipeline entrypoint is executed +// vs how many times the exit is executed. On architectures with a SIMD +// execution model #entries == #exits. On MIMD (or psuedo-MIMD) execution models +// #entries != #exits. +// +// The instrumentation will write buffers in debug descriptor set |desc_set|. +// It will write |shader_id| in each output record to identify the shader +// module which generated the record if necessary. +Optimizer::PassToken CreateAutoInstWarpEntryAndExitPass(uint32_t desc_set, + uint32_t shader_id); + + +// Create a pass which will automatically insert instrumentation to +// capture the execution trace of the pipeline. +// +// The instrumentation will write buffers in debug descriptor set |desc_set|. +// It will write |shader_id| in each output record to identify the shader +// module which generated the record. +// +// The |static_data_callback| is called after the instrumentation pass has +// finished. It is populated with a mapping from instrumentation callsite id +// to the other opcodes in the basic block. This data allows the analysis +// to develop complete dynamic instruction counts of the shader module without +// needing to transfer the data at runtime. +Optimizer::PassToken CreateAutoInstExecutionTracePass( + uint32_t desc_set, uint32_t shader_id, + std::function>&&, + std::unordered_map&&)> + static_data_callback); + // Create a pass to upgrade to the VulkanKHR memory model. // This pass upgrades the Logical GLSL450 memory model to Logical VulkanKHR. // Additionally, it modifies memory, image, atomic and barrier operations to diff --git a/source/opt/CMakeLists.txt b/source/opt/CMakeLists.txt index f3ac5906..a59b18c8 100644 --- a/source/opt/CMakeLists.txt +++ b/source/opt/CMakeLists.txt @@ -14,6 +14,14 @@ set(SPIRV_TOOLS_OPT_SOURCES aggressive_dead_code_elim_pass.h amd_ext_to_khr.h + auto_inst_pass.h + auto_inst_debug_pass.h + auto_inst_divergence_characterization_pass.h + auto_inst_dyn_shader_trace_pass.h + auto_inst_dyn_trace_ray_trace_pass.h + auto_inst_execution_trace_pass.h + auto_inst_simt_efficiency_pass.h + auto_inst_warp_entry_and_exit_pass.h basic_block.h block_merge_pass.h block_merge_util.h @@ -122,6 +130,14 @@ set(SPIRV_TOOLS_OPT_SOURCES aggressive_dead_code_elim_pass.cpp amd_ext_to_khr.cpp + auto_inst_pass.cpp + auto_inst_debug_pass.cpp + auto_inst_divergence_characterization_pass.cpp + auto_inst_dyn_shader_trace_pass.cpp + auto_inst_dyn_trace_ray_trace_pass.cpp + auto_inst_execution_trace_pass.cpp + auto_inst_simt_efficiency_pass.cpp + auto_inst_warp_entry_and_exit_pass.cpp basic_block.cpp block_merge_pass.cpp block_merge_util.cpp @@ -167,7 +183,7 @@ set(SPIRV_TOOLS_OPT_SOURCES inline_pass.cpp inst_bindless_check_pass.cpp inst_buff_addr_check_pass.cpp - inst_debug_printf_pass.cpp + inst_debug_printf_pass.cpp instruction.cpp instruction_list.cpp instrument_pass.cpp diff --git a/source/opt/auto_inst_debug_pass.cpp b/source/opt/auto_inst_debug_pass.cpp new file mode 100644 index 00000000..d4321e92 --- /dev/null +++ b/source/opt/auto_inst_debug_pass.cpp @@ -0,0 +1,68 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "auto_inst_debug_pass.h" + +namespace spvtools { +namespace opt { + +bool AutoInstDebugPass::PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) { + (void)stage_idx; + + uint32_t buf_id = GetOutputBufferId(); + uint32_t buf_uint_ptr_id = GetOutputBufferPtrId(); + + if (test_atomic_ops_) { + Instruction* offset_ptr = builder->AddTernaryOp( + buf_uint_ptr_id, SpvOpAccessChain, buf_id, + builder->GetUintConstantId(kDebugOutputDataOffset), + builder->GetUintConstantId(0)); + uint32_t mask_none_id = builder->GetUintConstantId(SpvMemoryAccessMaskNone); + uint32_t scope_invok_id = builder->GetUintConstantId(SpvScopeInvocation); + (void)builder->AddQuadOp(GetUintId(), SpvOpAtomicIAdd, + offset_ptr->result_id(), scope_invok_id, + mask_none_id, builder->GetUintConstantId(1)); + } else if (test_subgroup_ops_) { + Instruction* subgroup_leader_cond = + builder->AddUnaryOp(GetBoolId(), SpvOpGroupNonUniformElect, + builder->GetUintConstantId(SpvScopeSubgroup)); + + auto active_thread_mask = + GenSubgroupBallotId(builder, subgroup_leader_cond->result_id()); + + Instruction* offset_ptr = builder->AddTernaryOp( + buf_uint_ptr_id, SpvOpAccessChain, buf_id, + builder->GetUintConstantId(kDebugOutputDataOffset), + builder->GetUintConstantId(0)); + builder->AddStore(offset_ptr->result_id(), active_thread_mask); + + } else if (!test_atomic_ops_ && !test_subgroup_ops_) { + Instruction* buffer_capacity = + builder->AddIdLiteralOp(GetUintId(), SpvOpArrayLength, + GetOutputBufferId(), kDebugOutputDataOffset); + + Instruction* offset_ptr = builder->AddTernaryOp( + buf_uint_ptr_id, SpvOpAccessChain, buf_id, + builder->GetUintConstantId(kDebugOutputDataOffset), + builder->GetUintConstantId(0)); + + builder->AddStore(offset_ptr->result_id(), buffer_capacity->result_id()); + } + + return true; +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_debug_pass.h b/source/opt/auto_inst_debug_pass.h new file mode 100644 index 00000000..7fb59430 --- /dev/null +++ b/source/opt/auto_inst_debug_pass.h @@ -0,0 +1,106 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_DEBUG_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_DEBUG_PASS_H_ + +#include "auto_inst_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstDebugPass : public AutoInstPass { + public: + AutoInstDebugPass(uint32_t desc_set, uint32_t shader_id, bool test_atomic_ops, + bool test_subgroup_ops) + : AutoInstPass(desc_set, shader_id, kDefaultReservedWordsCnt), + test_atomic_ops_(test_atomic_ops), + test_subgroup_ops_(test_subgroup_ops) {} + + const char* name() const override { return "auto-inst-debug-pass"; } + + protected: + const bool test_atomic_ops_; + const bool test_subgroup_ops_; + + private: + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + void InitializeInstrumentation() override{}; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + void FinalizeInstrumentation() override{}; + + // Any instructions added via |builder| will appear before |inst| + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)bb; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)builder; + (void)stage_idx; + return false; + } +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_AUTO_INST_DEBUG_PASS_H_ diff --git a/source/opt/auto_inst_divergence_characterization_pass.cpp b/source/opt/auto_inst_divergence_characterization_pass.cpp new file mode 100644 index 00000000..a462c05e --- /dev/null +++ b/source/opt/auto_inst_divergence_characterization_pass.cpp @@ -0,0 +1,134 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "auto_inst_divergence_characterization_pass.h" + +namespace { +std::unordered_set pseudo_ops = { + SpvOpPhi, SpvOpFunction, SpvOpFunctionParameter, + SpvOpFunctionEnd, SpvOpVariable, SpvOpLoopMerge}; + +static const std::set kAllRayTracingStages = { + SpvExecutionModelRayGenerationNV, SpvExecutionModelAnyHitNV, + SpvExecutionModelClosestHitNV, SpvExecutionModelMissNV, + SpvExecutionModelIntersectionNV, SpvExecutionModelCallableNV}; + +} // namespace + +namespace spvtools { +namespace opt { + +void AutoInstDivergenceCharacterizationPass::GenInst( + InstructionBuilder* builder, Instruction* instruction, uint32_t prim_id, + uint32_t stage_idx) { + uint32_t inst_id = GenInstCallsiteId(instruction); + + auto flat_thread_id = GenFlatRtThreadId(builder, stage_idx); + auto active_thread_mask_id = + GenSubgroupBallotId(builder, builder->GetBoolConstant(true)->result_id()); + + GenSubgroupUpdateCall(builder, {builder->GetUintConstantId(inst_id), + flat_thread_id, active_thread_mask_id}); + + inst_id2prim_id_[inst_id] = prim_id; +} + +void AutoInstDivergenceCharacterizationPass::InitializeInstrumentation() { + inst_id2inst_count_.clear(); + inst_id2prim_id_.clear(); +} + +bool AutoInstDivergenceCharacterizationPass::PreInstructionInstrument( + Instruction* inst, InstructionBuilder* builder, uint32_t stage_idx) { + if (inst->opcode() == SpvOpTraceNV || inst->opcode() == SpvOpTraceRayKHR) { + // Record when threads begin a traceRay + GenInst(builder, &*builder->GetInsertPoint(), kAutoInstDivCharPreTraceRay, + stage_idx); + return true; + } + + return false; +}; + +bool AutoInstDivergenceCharacterizationPass::PostInstructionInstrument( + Instruction* inst, InstructionBuilder* builder, uint32_t stage_idx) { + if (inst->opcode() == SpvOpTraceNV || inst->opcode() == SpvOpTraceRayKHR) { + // Record when threads end a traceRay + GenInst(builder, &*builder->GetInsertPoint(), kAutoInstDivCharPostTraceRay, + stage_idx); + return true; + } + + return false; +} + +bool AutoInstDivergenceCharacterizationPass::PreEntryPointInstrument( + InstructionBuilder* builder, uint32_t stage_idx) { + if (stage_idx == SpvExecutionModelRayGenerationKHR) { + inst_id2prim_id_[kAutoInstUniqueSubgroupId] = kAutoInstUniqueSubgroupId; + // Create an instrumentation id which will be used by the analysis + // to determine how the subsequent words should be understood. + auto unique_warp_id_inst_id = + builder->GetUintConstantId(kAutoInstUniqueSubgroupId); + + GenUniqueSubgroupIdFuncCall(builder, unique_warp_id_inst_id, stage_idx); + + return true; + } else if (kAllRayTracingStages.count(stage_idx) != 0) { + // Record when threads run a shader during a traceRay + GenInst(builder, + builder->GetIntConstant(kAutoInstDivCharShaderEntryPoint, false), + kAutoInstDivCharShaderEntryPoint, stage_idx); + return true; + } + return false; +} + +bool AutoInstDivergenceCharacterizationPass::PostEntryPointInstrument( + InstructionBuilder* builder, uint32_t stage_idx) { + if (stage_idx != SpvExecutionModelRayGenerationKHR) return false; + + // Record the threads that quit the pipeline + GenInst(builder, builder->GetIntConstant(kAutoInstDivCharQuitPipeline, false), + kAutoInstDivCharQuitPipeline, stage_idx); + + return true; +} + +bool AutoInstDivergenceCharacterizationPass::BasicBlockInstrument( + BasicBlock* bb, InstructionBuilder* builder, uint32_t stage_idx) { + if (kAllRayTracingStages.count(stage_idx) == 0) return false; + + auto inst = builder->GetInsertPoint(); + + // Record active threads in each basic block execution + GenInst(builder, &*inst, kAutoInstDivCharActiveThreads, stage_idx); + + uint32_t count = 0; + for (auto& ii : *bb) { + if (pseudo_ops.count(ii.opcode()) != 0) count++; + } + + inst_id2inst_count_[GenInstCallsiteId(&*inst)] = count; + + return true; +} + +void AutoInstDivergenceCharacterizationPass::FinalizeInstrumentation() { + static_data_callback_(std::move(inst_id2prim_id_), + std::move(inst_id2inst_count_)); +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_divergence_characterization_pass.h b/source/opt/auto_inst_divergence_characterization_pass.h new file mode 100644 index 00000000..54c967a9 --- /dev/null +++ b/source/opt/auto_inst_divergence_characterization_pass.h @@ -0,0 +1,106 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_DIVERGENCE_CHARACTERIZATION_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_DIVERGENCE_CHARACTERIZATION_PASS_H_ + +#include + +#include "auto_inst_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstDivergenceCharacterizationPass : public AutoInstPass { + public: + AutoInstDivergenceCharacterizationPass( + uint32_t desc_set, uint32_t shader_id, + std::function< + void(std::unordered_map&& inst_id2prim_id, + std::unordered_map&& inst_id2inst_count)> + static_data_callback) + : AutoInstPass(desc_set, shader_id, kDefaultReservedWordsCnt), + static_data_callback_(static_data_callback) {} + + const char* name() const override { + return "auto-inst-divergence-characterization-pass"; + } + + private: + // Mapping from inst callsite id to primitive id + std::unordered_map inst_id2prim_id_; + + // Mapping from inst callsite id to the number of insts in the bb + std::unordered_map inst_id2inst_count_; + + // Callback for sending static data + std::function&& inst_id2prim_id, + std::unordered_map&& inst_id2inst_count)> + static_data_callback_; + + // Generate instrumentation for this pass + void GenInst(InstructionBuilder* builder, Instruction* inst, + AutoInstId prim_id, AutoInstId stage_idx); + + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + void InitializeInstrumentation() override; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + void FinalizeInstrumentation() override; + + // Any instructions added via |builder| will appear before |inst| + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override; +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_AUTO_INST_DIVERGENCE_CHARACTERIZATION_PASS_H_ diff --git a/source/opt/auto_inst_dyn_shader_trace_pass.cpp b/source/opt/auto_inst_dyn_shader_trace_pass.cpp new file mode 100644 index 00000000..74416f43 --- /dev/null +++ b/source/opt/auto_inst_dyn_shader_trace_pass.cpp @@ -0,0 +1,52 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + + +#include "auto_inst_dyn_shader_trace_pass.h" + +namespace { + +static const std::set kAllRayTracingStages = { + SpvExecutionModelRayGenerationNV, SpvExecutionModelAnyHitNV, + SpvExecutionModelClosestHitNV, SpvExecutionModelMissNV, + SpvExecutionModelIntersectionNV, SpvExecutionModelCallableNV}; +} + +namespace spvtools { +namespace opt { + +bool AutoInstDynShaderTracePass::PreEntryPointInstrument( + InstructionBuilder* builder, uint32_t stage_idx) { + // Ray Generation begins and ends the ray tracing pipeline + if (stage_idx == SpvExecutionModelRayGenerationNV) { + // Create an instrumentation id which will be used by the analysis + // to determine how the subsequent words should be understood. + auto unique_warp_id_inst_id = + builder->GetUintConstantId(kAutoInstUniqueSubgroupId); + GenUniqueSubgroupIdFuncCall(builder, unique_warp_id_inst_id, stage_idx); + } + + auto prim_id = builder->GetUintConstantId(kAutoInstDynShaderTraceEntryPoint); + auto flat_thread_id = GenFlatRtThreadId(builder, stage_idx); + auto shader_id = builder->GetUintConstantId(shader_id_); + auto active_thread_mask_id = + GenSubgroupBallotId(builder, builder->GetBoolConstant(true)->result_id()); + GenSubgroupUpdateCall( + builder, {prim_id, flat_thread_id, shader_id, active_thread_mask_id}); + + return true; +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_dyn_shader_trace_pass.h b/source/opt/auto_inst_dyn_shader_trace_pass.h new file mode 100644 index 00000000..840f8942 --- /dev/null +++ b/source/opt/auto_inst_dyn_shader_trace_pass.h @@ -0,0 +1,103 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_DYN_SHADER_TRACE_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_DYN_SHADER_TRACE_PASS_H_ + +#include + +#include "auto_inst_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstDynShaderTracePass : public AutoInstPass { + public: + AutoInstDynShaderTracePass(uint32_t desc_set, uint32_t shader_id) + : AutoInstPass(desc_set, shader_id, kDefaultReservedWordsCnt) {} + + const char* name() const override { + return "auto-inst-dyn-shader-trace-pass"; + } + + private: + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + void InitializeInstrumentation() override{}; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + void FinalizeInstrumentation() override{}; + + // Any instructions added via |builder| will appear before |inst| + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)bb; + (void)builder; + (void)stage_idx; + return false; + } + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)builder; + (void)stage_idx; + return false; + } +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_AUTO_INST_DYN_SHADER_TRACE_PASS_H_ diff --git a/source/opt/auto_inst_dyn_trace_ray_trace_pass.cpp b/source/opt/auto_inst_dyn_trace_ray_trace_pass.cpp new file mode 100644 index 00000000..edd63b34 --- /dev/null +++ b/source/opt/auto_inst_dyn_trace_ray_trace_pass.cpp @@ -0,0 +1,184 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + + +#include "auto_inst_dyn_trace_ray_trace_pass.h" + +namespace { + +static const std::set kAllRayTracingStages = { + SpvExecutionModelRayGenerationNV, SpvExecutionModelAnyHitNV, + SpvExecutionModelClosestHitNV, SpvExecutionModelMissNV, + SpvExecutionModelIntersectionNV, SpvExecutionModelCallableNV}; + +static const int kEntryPointFunctionIdInIdx = 1; +} // namespace + +namespace spvtools { +namespace opt { + +void AutoInstDynTraceRayTracePass::GenInst(InstructionBuilder* builder, + uint32_t inst_id, uint32_t prim_type, + uint32_t stage_idx) { + if (inst_id2prim_type_.count(inst_id) != 0) { + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, + "Multiple instrumentation sites with the same id detected!\n"); + return; + } + + auto flat_thread_id = GenFlatRtThreadId(builder, stage_idx); + auto active_thread_mask_id = + GenSubgroupBallotId(builder, builder->GetBoolConstant(true)->result_id()); + + GenSubgroupUpdateCall(builder, {builder->GetUintConstantId(inst_id), + flat_thread_id, active_thread_mask_id}); + + inst_id2prim_type_[inst_id] = prim_type; +} + +void AutoInstDynTraceRayTracePass::InitializeInstrumentation() { + inst_id2prim_type_.clear(); + fn_id_2_divergent_ids_.clear(); + merge_id_2_divergent_ids_.clear(); + loop_begin_bb_ids_.clear(); + + Function* entry_point_func = nullptr; + + for (auto e : get_module()->entry_points()) { + auto func_id = e.GetSingleWordInOperand(kEntryPointFunctionIdInIdx); + entry_point_func = id2function_[func_id]; + } + + auto active_merge_ids = std::stack(); + PopulateDivergentLabelsInfo(entry_point_func, active_merge_ids); +} + +bool AutoInstDynTraceRayTracePass::PreEntryPointInstrument( + InstructionBuilder* builder, uint32_t stage_idx) { + if (stage_idx != SpvExecutionModelRayGenerationKHR) return false; + + // Create an instrumentation id which will be used by the analysis + // to determine how the subsequent words should be understood. + auto unique_warp_id_inst_id = + builder->GetUintConstantId(kAutoInstUniqueSubgroupId); + + // This instrumentation is unique across the pipeline so it is sufficient + // to choose an inst_id that can never alias with the other types of + // instrumentation. + inst_id2prim_type_[kAutoInstUniqueSubgroupId] = kAutoInstUniqueSubgroupId; + + GenUniqueSubgroupIdFuncCall(builder, unique_warp_id_inst_id, stage_idx); + + return true; +} + +bool AutoInstDynTraceRayTracePass::BasicBlockInstrument( + BasicBlock* bb, InstructionBuilder* builder, uint32_t stage_idx) { + if (kAllRayTracingStages.count(stage_idx) == 0) return false; + + if (merge_id_2_divergent_ids_.count(bb->id()) != 0) { + GenInst(builder, bb->id(), kAutoInstTraceRayTraceMergePoint, stage_idx); + return true; + } + + for (auto& ii : *bb) { + if (ii.opcode() == SpvOpTraceRayKHR || ii.opcode() == SpvOpTraceNV) { + builder->SetInsertPoint(&ii); + GenInst(builder, bb->id(), kAutoInstTraceRayTracePreTraceRay, stage_idx); + return true; + } + } + + return false; +} + +void AutoInstDynTraceRayTracePass::FinalizeInstrumentation() { + static_data_callback_(std::move(inst_id2prim_type_), + std::move(merge_id_2_divergent_ids_)); +} + +void AutoInstDynTraceRayTracePass::PopulateDivergentLabelsInfo( + Function* func, std::stack& active_merge_ids) { + if (fn_id_2_divergent_ids_.count(func->result_id()) > 0) return; + + bool is_divergent_control_flow = active_merge_ids.size() > 0; + std::vector divergent_labels; + std::stack active_loop_ids; + + for (auto& blk : *func) { + if (active_merge_ids.size() > 0 && blk.id() == active_merge_ids.top()) { + // Need to move the merge label to the beginning of the loop + // iteration in order to determine how many times the label was visited at + // runtime + if (active_loop_ids.size() > 0) { + if (merge_id_2_divergent_ids_.count(active_merge_ids.top()) > 0) { + loop_begin_bb_ids_.insert(active_loop_ids.top()); + } + merge_id_2_divergent_ids_[active_loop_ids.top()] = + merge_id_2_divergent_ids_[active_merge_ids.top()]; + merge_id_2_divergent_ids_.erase(active_merge_ids.top()); + active_loop_ids.pop(); + } + active_merge_ids.pop(); + } + + for (auto& inst : blk) { + // Determine divergent labels to track + if (inst.opcode() == SpvOpTraceRayKHR || inst.opcode() == SpvOpTraceNV) { + divergent_labels.push_back(blk.id()); + } else if (inst.opcode() == SpvOpFunctionCall) { + // Add divergent labels according to func being called + auto func_to_call_id = inst.GetSingleWordOperand(2); + if (fn_id_2_divergent_ids_.count(func_to_call_id) == 0) { + // recurse if fn not discovered yet + PopulateDivergentLabelsInfo(id2function_[func_to_call_id], + active_merge_ids); + } + divergent_labels = fn_id_2_divergent_ids_[func_to_call_id]; + + } else if (inst.opcode() == SpvOpSelectionMerge || + inst.opcode() == SpvOpLoopMerge) { + auto merge_id = inst.GetSingleWordOperand(0); + + if (active_merge_ids.size() == 0) { + active_merge_ids.push(merge_id); + } else if (inst.opcode() == SpvOpLoopMerge) { + active_merge_ids.push(merge_id); + active_loop_ids.push(inst.GetSingleWordOperand(1)); + } + } + + // Update datastructures with divergent labels + if (divergent_labels.size() > 0) { + fn_id_2_divergent_ids_[func->result_id()].insert( + fn_id_2_divergent_ids_[func->result_id()].end(), + divergent_labels.begin(), divergent_labels.end()); + + // Update all active to-be-merged labels + if (active_merge_ids.size() > 0 && !is_divergent_control_flow) { + auto id = active_merge_ids.top(); + + merge_id_2_divergent_ids_[id].insert( + merge_id_2_divergent_ids_[id].end(), divergent_labels.begin(), + divergent_labels.end()); + } + + divergent_labels.clear(); + } + } + } +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_dyn_trace_ray_trace_pass.h b/source/opt/auto_inst_dyn_trace_ray_trace_pass.h new file mode 100644 index 00000000..d095398d --- /dev/null +++ b/source/opt/auto_inst_dyn_trace_ray_trace_pass.h @@ -0,0 +1,135 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_DYN_TRACE_RAY_TRACE_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_DYN_TRACE_RAY_TRACE_PASS_H_ + +#include + +#include "auto_inst_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstDynTraceRayTracePass : public AutoInstPass { + public: + AutoInstDynTraceRayTracePass( + uint32_t desc_set, uint32_t shader_id, + std::function&&, + std::unordered_map>&&)> + static_data_callback) + : AutoInstPass(desc_set, shader_id, kDefaultReservedWordsCnt), + static_data_callback_(static_data_callback) {} + + const char* name() const override { + return "auto-inst-dyn-trace-ray-trace-pass"; + } + + private: + // Callback to transfer static data to creator of the pass. + std::function&&, + std::unordered_map>&&)> + static_data_callback_; + + // Static data which forms a mapping from instrumentation callsite id + // to the instructions in the basic block. + std::unordered_map inst_id2prim_type_; + + // Static data which contains the mapping of MergePoint instrumentation + // to all the traceRay callsites that must have executed. + std::unordered_map> merge_id_2_divergent_ids_; + + // Memoization table for each function + std::unordered_map> fn_id_2_divergent_ids_; + + // Which bbs to add instrumentation before to track loop iterations. + std::set loop_begin_bb_ids_; + + // Generate instrumentation for this pass + void GenInst(InstructionBuilder* builder, uint32_t inst_id, + uint32_t prim_type, uint32_t stage_idx); + + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + void InitializeInstrumentation() override; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + void FinalizeInstrumentation() override; + + // Any instructions added via |builder| will appear before |inst| + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + } + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)builder; + (void)stage_idx; + return false; + }; + + // This function generates an analysis of |func| to determine + // 1) the SpvOpTraceRays that exist within control-flow. + // 2) where to add instrumentation to detect loop iterations and + // opportunities to execute traceRays. + // + // This allows runtime traceRay calls to be disambiguated. + void PopulateDivergentLabelsInfo(Function* func, + std::stack& active_merge_ids); +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_AUTO_INST_DYN_TRACE_RAY_TRACE_PASS_H_ diff --git a/source/opt/auto_inst_execution_trace_pass.cpp b/source/opt/auto_inst_execution_trace_pass.cpp new file mode 100644 index 00000000..3eeff48e --- /dev/null +++ b/source/opt/auto_inst_execution_trace_pass.cpp @@ -0,0 +1,68 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "auto_inst_execution_trace_pass.h" + +namespace spvtools { +namespace opt { + +void AutoInstExecutionTracePass::InitializeInstrumentation() { + inst_id2bb_inst_ids_.clear(); + inst_id2opcode_.clear(); +} + +bool AutoInstExecutionTracePass::BasicBlockInstrument( + BasicBlock* bb, InstructionBuilder* builder, uint32_t stage_idx) { + (void)stage_idx; + auto module_offset = uid2offset_[bb->begin()->unique_id()]; + if (shader_id_ >= (1 << 12) || module_offset >= (1 << 20)) { + std::string message = + "Shader id count or shader module size are too large!\n"; + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, message.c_str()); + return false; + } + + // Create an instrumentation callsite id that is unique across + // the whole pipeline by including the shader id in the upper bits. + auto inst = &*builder->GetInsertPoint(); + + uint32_t inst_id = GenInstCallsiteId(inst); + + for (auto inst_itr : *bb) { + auto other_inst_id = GenInstCallsiteId(&inst_itr); + + // Create group of instructions which must execute + // if the instrumentation executed. + inst_id2bb_inst_ids_[inst_id].insert(other_inst_id); + // Add opcode to the static metadata map so it can be added + // to a dynamic opcode total. + inst_id2opcode_[other_inst_id] = inst_itr.opcode(); + } + // Write the same inst_id as in the static data so that when a + // buffer entry is parsed, the inst_id can be used to look up + // the other instructions that must have also been executed. + auto active_thread_mask_id = + GenSubgroupBallotId(builder, builder->GetBoolConstant(true)->result_id()); + GenSubgroupUpdateCall( + builder, {builder->GetUintConstantId(inst_id), active_thread_mask_id}); + return true; +} + +void AutoInstExecutionTracePass::FinalizeInstrumentation() { + static_data_callback_(std::move(inst_id2bb_inst_ids_), + std::move(inst_id2opcode_)); +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_execution_trace_pass.h b/source/opt/auto_inst_execution_trace_pass.h new file mode 100644 index 00000000..124f3ecd --- /dev/null +++ b/source/opt/auto_inst_execution_trace_pass.h @@ -0,0 +1,122 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_EXECUTION_TRACE_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_EXECUTION_TRACE_PASS_H_ + +#include + +#include "auto_inst_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstExecutionTracePass : public AutoInstPass { + public: + AutoInstExecutionTracePass( + uint32_t desc_set, uint32_t shader_id, + std::function< + void(std::unordered_map>&& + inst_id2bb_inst_ids, + std::unordered_map&& inst_id2opcode)> + static_data_callback) + : AutoInstPass(desc_set, shader_id, kDefaultReservedWordsCnt), + static_data_callback_(static_data_callback) {} + + const char* name() const override { return "auto-inst-execution-trace-pass"; } + + private: + // Callback to transfer static data to creator of the pass. + const std::function>&& inst_id2bb_inst_ids, + std::unordered_map&& inst_id2opcode)> + static_data_callback_; + + // Static data which forms a mapping from instrumentation callsite id + // to the ids of other instructions in the basic block. + std::unordered_map> inst_id2bb_inst_ids_; + + // Static data which forms a mappign from inst_id to opcode. + // This together with |inst_id2bb_inst_ids_| gives the runtime + // instruction mix. + std::unordered_map inst_id2opcode_; + + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + void InitializeInstrumentation() override; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + void FinalizeInstrumentation() override; + + // Any instructions added via |builder| will appear before |inst| + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)builder; + (void)stage_idx; + return false; + } + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)builder; + (void)stage_idx; + return false; + }; +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_AUTO_INST_EXECUTION_TRACE_PASS_H_ diff --git a/source/opt/auto_inst_pass.cpp b/source/opt/auto_inst_pass.cpp new file mode 100644 index 00000000..6a5f7f60 --- /dev/null +++ b/source/opt/auto_inst_pass.cpp @@ -0,0 +1,927 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + + +#include "auto_inst_pass.h" + +namespace { + +// Operand indices of SpvOpEntryPoint +static const int kEntryPointFunctionIdInIdx = 1; +static const int kEntryPointExecutionModelInIdx = 0; + +// Type of function from OpFunction +static const int kFunctionTypeIdx = 3; + +// UniqueSubgroupId function constants +static const int kUniqueSubgroupIdParamCnt = 1; +static const int kUniqueSubgroupIdParamInstIdIdx = 0; +} // namespace + +namespace spvtools { +namespace opt { + +bool AutoInstPass::HandleInstrumentHooks( + BasicBlock::iterator ref_inst_itr, + UptrVectorIterator ref_block_itr, uint32_t stage_idx, + InstructionBuilder* builder) { + bool is_instrumented = false; + + // Initialize instrumentation validity based on inst opcode. + bool skip_pre_instrumentation = + uninstrumentable_pre_ops.count(ref_inst_itr->opcode()) > 0; + bool skip_post_instrumentation = + uninstrumentable_post_ops.count(ref_inst_itr->opcode()) > 0; + + if (ref_block_itr->GetParent()->begin() == ref_block_itr) { + // Despite the SPIR-V specification stating: + // "All OpVariable instructions in a function must be in the first block in + // the function. " The real restriction is as follows: "All OpVariable + // instructions in a function must be the first instructions in the first + // block." Therefore this check prevents instrumenting OpVariables in the + // first block of a function. + if (ref_inst_itr->opcode() == SpvOpVariable) { + skip_pre_instrumentation = true; + } + if (std::next(ref_inst_itr) != ref_block_itr->end() && + std::next(ref_inst_itr)->opcode() == SpvOpVariable) { + skip_post_instrumentation = true; + } + } + + if (skip_pre_instrumentation && skip_post_instrumentation) return false; + + if (!skip_pre_instrumentation) { + // set insert point to immediately before current inst + builder->SetInsertPoint(&*ref_inst_itr); + + if (instrumented_bb_ids.count(ref_block_itr->id()) == 0) { + is_instrumented |= + BasicBlockInstrument(&*ref_block_itr, builder, stage_idx); + } + + if (instrumented_inst_ids.count(ref_inst_itr->unique_id()) == 0) { + is_instrumented |= + PreInstructionInstrument(&*ref_inst_itr, builder, stage_idx); + } + } + if (!skip_post_instrumentation) { + if (instrumented_inst_ids.count(ref_inst_itr->unique_id()) == 0) { + // Before next inst is after this inst + builder->SetInsertPoint(&*std::next(ref_inst_itr)); + is_instrumented |= + PostInstructionInstrument(&*ref_inst_itr, builder, stage_idx); + } + } + + if (is_instrumented) { + // Record the bb and inst that were just visited + instrumented_bb_ids.insert(ref_block_itr->id()); + instrumented_inst_ids.insert(ref_inst_itr->unique_id()); + } + + return is_instrumented; +} + +void AutoInstPass::GenInstProgrammableCode( + BasicBlock::iterator ref_inst_itr, + UptrVectorIterator ref_block_itr, uint32_t stage_idx) { + // Initialize DefUse manager before dismantling module + (void)get_def_use_mgr(); + + InstructionBuilder builder(context(), &*ref_block_itr); + + bool is_instrumented = + HandleInstrumentHooks(ref_inst_itr, ref_block_itr, stage_idx, &builder); + if (!is_instrumented) return; + has_added_instrumentation_ = true; +} + +uint32_t AutoInstPass::GenSubgroupBallotId(InstructionBuilder* builder, + uint32_t pred_id) { + if (!get_feature_mgr()->HasExtension(kSPV_KHR_subgroup_vote)) { + context()->AddExtension("SPV_KHR_subgroup_vote"); + } + + if (!get_feature_mgr()->HasCapability(SpvCapabilityGroupNonUniformBallot)) { + context()->AddCapability(SpvCapabilityGroupNonUniformBallot); + } + + uint32_t scope_ballot_idx = builder->GetUintConstantId(SpvScopeSubgroup); + Instruction* ballot_inst = builder->AddBinaryOp( + GetVec4UintId(), SpvOpGroupNonUniformBallot, scope_ballot_idx, pred_id); + + return builder + ->AddIdLiteralOp(GetUintId(), SpvOpCompositeExtract, + ballot_inst->result_id(), 0) + ->result_id(); +} + +std::pair AutoInstPass::GenReadClockIds( + InstructionBuilder* builder) { + if (!get_feature_mgr()->HasExtension(kSPV_KHR_shader_clock)) { + context()->AddExtension("SPV_KHR_shader_clock"); + } + + if (!get_feature_mgr()->HasCapability(SpvCapabilityShaderClockKHR)) { + context()->AddCapability(SpvCapabilityShaderClockKHR); + } + + auto time_inst = + builder->AddUnaryOp(GetVecUintId(2u), SpvOpReadClockKHR, + builder->GetUintConstantId(SpvScopeDevice)); + Instruction* time_lower = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, time_inst->result_id(), 0); + Instruction* time_upper = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, time_inst->result_id(), 1); + return std::make_pair(time_lower->result_id(), time_upper->result_id()); +} + +uint32_t AutoInstPass::GenFlatRtThreadId(InstructionBuilder* builder, + uint32_t stage_idx) { + switch (stage_idx) { + case SpvExecutionModelRayGenerationNV: + case SpvExecutionModelIntersectionNV: + case SpvExecutionModelAnyHitNV: + case SpvExecutionModelClosestHitNV: + case SpvExecutionModelMissNV: + case SpvExecutionModelCallableNV: { + auto launch_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInLaunchIdKHR), builder); + Instruction* launch_x = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_id, 0); + Instruction* launch_y = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_id, 1); + Instruction* launch_z = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_id, 2); + + auto launch_size_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInLaunchSizeKHR), builder); + Instruction* launch_size_x = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_size_id, 0); + Instruction* launch_size_y = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_size_id, 1); + + auto xy_size = builder->AddBinaryOp(GetUintId(), SpvOpIMul, + launch_size_x->result_id(), + launch_size_y->result_id()); + auto z_term = builder->AddBinaryOp( + GetUintId(), SpvOpIMul, launch_z->result_id(), xy_size->result_id()); + + auto y_term = + builder->AddBinaryOp(GetUintId(), SpvOpIMul, launch_y->result_id(), + launch_size_x->result_id()); + + auto flat_thread_id = builder->AddBinaryOp( + GetUintId(), SpvOpIAdd, z_term->result_id(), y_term->result_id()); + flat_thread_id = builder->AddBinaryOp(GetUintId(), SpvOpIAdd, + flat_thread_id->result_id(), + launch_x->result_id()); + return flat_thread_id->result_id(); + } + + default: + consumer()( + SPV_MSG_ERROR, 0, {0, 0, 0}, + "Cannot create a flattened rt thread id for requested shader stage! " + "Defaulting to 0.\n"); + return builder->GetUintConstantId(0); + } +} + +uint32_t AutoInstPass::GenFlatComputeThreadId(InstructionBuilder* builder, + uint32_t stage_idx) { + if (stage_idx != SpvExecutionModelGLCompute) { + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, + "Cannot create a flattened compute thread id for requested " + "shader stage! " + "Defaulting to 0.\n"); + return builder->GetUintConstantId(0); + } + auto invocation_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInGlobalInvocationId), builder); + Instruction* invocation_x = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, invocation_id, 0); + Instruction* invocation_y = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, invocation_id, 1); + Instruction* invocation_z = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, invocation_id, 2); + + auto num_workgroups_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInNumWorkgroups), builder); + Instruction* num_workgroups_x = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, num_workgroups_id, 0); + Instruction* num_workgroups_y = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, num_workgroups_id, 1); + + auto workgroup_size_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInWorkgroupSize), builder); + + Instruction* workgroup_size_x = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, workgroup_size_id, 0); + Instruction* workgroup_size_y = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, workgroup_size_id, 1); + + Instruction* global_size_x = builder->AddBinaryOp( + GetUintId(), SpvOpIMul, num_workgroups_x->result_id(), + workgroup_size_x->result_id()); + + Instruction* global_size_y = builder->AddBinaryOp( + GetUintId(), SpvOpIMul, num_workgroups_y->result_id(), + workgroup_size_y->result_id()); + + auto xy_size = + builder->AddBinaryOp(GetUintId(), SpvOpIMul, global_size_x->result_id(), + global_size_y->result_id()); + auto z_term = builder->AddBinaryOp( + GetUintId(), SpvOpIMul, invocation_z->result_id(), xy_size->result_id()); + + auto y_term = + builder->AddBinaryOp(GetUintId(), SpvOpIMul, invocation_y->result_id(), + global_size_x->result_id()); + + auto flat_thread_id = builder->AddBinaryOp( + GetUintId(), SpvOpIAdd, z_term->result_id(), y_term->result_id()); + flat_thread_id = + builder->AddBinaryOp(GetUintId(), SpvOpIAdd, flat_thread_id->result_id(), + invocation_x->result_id()); + return flat_thread_id->result_id(); +} + +std::vector AutoInstPass::GenThreadId(InstructionBuilder* builder, + uint32_t stage_idx) { + switch (stage_idx) { + case SpvExecutionModelVertex: { + // Load and store VertexId and InstanceId + auto vertex_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInVertexIndex), builder); + auto instance_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInInstanceIndex), builder); + return {vertex_id, instance_id}; + } + case SpvExecutionModelGLCompute: { + // Load and store GlobalInvocationId. + + uint32_t load_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInGlobalInvocationId), + builder); + Instruction* x_inst = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, load_id, 0); + Instruction* y_inst = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, load_id, 1); + Instruction* z_inst = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, load_id, 2); + return {x_inst->result_id(), y_inst->result_id(), z_inst->result_id()}; + } + case SpvExecutionModelTessellationControl: + case SpvExecutionModelGeometry: { + // Load and store PrimitiveId and InvocationId. + auto primitive_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInPrimitiveId), builder); + auto instance_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInInvocationId), builder); + return {primitive_id, instance_id}; + } + + case SpvExecutionModelTessellationEvaluation: { + // Load and store PrimitiveId and TessCoord.uv + auto primitive_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInPrimitiveId), builder); + uint32_t load_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInTessCoord), builder); + Instruction* uvec3_cast_inst = + builder->AddUnaryOp(GetVec3UintId(), SpvOpBitcast, load_id); + uint32_t uvec3_cast_id = uvec3_cast_inst->result_id(); + Instruction* u_inst = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, uvec3_cast_id, 0); + Instruction* v_inst = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, uvec3_cast_id, 1); + return {primitive_id, u_inst->result_id(), v_inst->result_id()}; + } + case SpvExecutionModelFragment: { + // Load FragCoord and convert to Uint + Instruction* frag_coord_inst = builder->AddUnaryOp( + GetVec4FloatId(), SpvOpLoad, + context()->GetBuiltinInputVarId(SpvBuiltInFragCoord)); + Instruction* uint_frag_coord_inst = builder->AddUnaryOp( + GetVec4UintId(), SpvOpBitcast, frag_coord_inst->result_id()); + Instruction* x_inst = + builder->AddIdLiteralOp(GetUintId(), SpvOpCompositeExtract, + uint_frag_coord_inst->result_id(), 0); + Instruction* y_inst = + builder->AddIdLiteralOp(GetUintId(), SpvOpCompositeExtract, + uint_frag_coord_inst->result_id(), 1); + Instruction* z_inst = + builder->AddIdLiteralOp(GetUintId(), SpvOpCompositeExtract, + uint_frag_coord_inst->result_id(), 2); + return {x_inst->result_id(), y_inst->result_id(), z_inst->result_id()}; + } + case SpvExecutionModelRayGenerationNV: + case SpvExecutionModelIntersectionNV: + case SpvExecutionModelAnyHitNV: + case SpvExecutionModelClosestHitNV: + case SpvExecutionModelMissNV: + case SpvExecutionModelCallableNV: { + // Load and store LaunchIdNV. + auto launch_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInLaunchIdKHR), builder); + Instruction* launch_x = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_id, 0); + Instruction* launch_y = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_id, 1); + Instruction* launch_z = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_id, 2); + + auto launch_size_id = GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInLaunchSizeKHR), builder); + Instruction* launch_size_x = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_size_id, 0); + Instruction* launch_size_y = builder->AddIdLiteralOp( + GetUintId(), SpvOpCompositeExtract, launch_size_id, 1); + + auto xy_size = builder->AddBinaryOp(GetUintId(), SpvOpIMul, + launch_size_x->result_id(), + launch_size_y->result_id()); + auto z_term = builder->AddBinaryOp( + GetUintId(), SpvOpIMul, launch_z->result_id(), xy_size->result_id()); + + auto y_term = + builder->AddBinaryOp(GetUintId(), SpvOpIMul, launch_y->result_id(), + launch_size_x->result_id()); + + auto flat_thread_id = builder->AddBinaryOp( + GetUintId(), SpvOpIAdd, z_term->result_id(), y_term->result_id()); + flat_thread_id = builder->AddBinaryOp(GetUintId(), SpvOpIAdd, + flat_thread_id->result_id(), + launch_x->result_id()); + return {flat_thread_id->result_id()}; + } + default: { + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, + "Cannot create a thread id for requested shader stage!\n"); + return {}; + } + } +} + +uint32_t AutoInstPass::GenInstCallsiteId(Instruction* inst) { + auto module_offset = uid2offset_[inst->unique_id()]; + if (shader_id_ >= (1 << 12) || module_offset >= (1 << 20)) { + std::string message = + "Shader id count or shader module size are too large!\n"; + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, message.c_str()); + return 0; + } + + // Create an instrumentation callsite id that is unique across + // the whole pipeline by including the shader id in the upper bits. + return (shader_id_ << 20) | module_offset; +} + +uint32_t AutoInstPass::GenSubgroupLocalInvocationId( + InstructionBuilder* builder) { + if (!get_feature_mgr()->HasCapability(SpvCapabilityGroupNonUniform)) { + context()->AddCapability(SpvCapabilityGroupNonUniform); + } + return GenVarLoad( + context()->GetBuiltinInputVarId(SpvBuiltInSubgroupLocalInvocationId), + builder); +} + +std::unique_ptr AutoInstPass::GenIfStatement( + uint32_t condition_id, std::unique_ptr curr_block, + std::function(InstructionBuilder* builder, + std::unique_ptr)> + inside_if_callback) { + auto output_func = curr_block->GetParent(); + InstructionBuilder builder( + context(), &*curr_block, + IRContext::kAnalysisDefUse | IRContext::kAnalysisInstrToBlockMapping); + uint32_t merge_blk_id = TakeNextId(); + uint32_t inside_if_blk_id = TakeNextId(); + std::unique_ptr merge_label(NewLabel(merge_blk_id)); + std::unique_ptr inside_if_label(NewLabel(inside_if_blk_id)); + (void)builder.AddConditionalBranch(condition_id, inside_if_blk_id, + merge_blk_id, merge_blk_id, + SpvSelectionControlMaskNone); + + output_func->AddBasicBlock(std::move(curr_block)); + curr_block = MakeUnique(std::move(inside_if_label)); + curr_block->SetParent(&*output_func); + + builder.SetInsertPoint(&*curr_block); + + curr_block = inside_if_callback(&builder, std::move(curr_block)); + + builder.SetInsertPoint(&*curr_block); + if (!curr_block->IsReturn()) (void)builder.AddBranch(merge_blk_id); + + output_func->AddBasicBlock(std::move(curr_block)); + curr_block = MakeUnique(std::move(merge_label)); + curr_block->SetParent(&*output_func); + return curr_block; +} + +std::unique_ptr AutoInstPass::GenThreadUpdate( + InstructionBuilder* builder, std::unique_ptr curr_block, + std::vector element_ids) { + uint32_t buf_id = GetOutputBufferId(); + uint32_t buf_uint_ptr_id = GetOutputBufferPtrId(); + if (element_ids.size() > 65535) { + std::string message = + "ThreadUpdate does not support more than 65535 elements in a single " + "entry!"; + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, message.c_str()); + } + // Compute size of entry that will be written to the buffer. + uint32_t entry_size = + builder->GetUintConstantId((uint32_t)element_ids.size()); + + // Update number of words written + uint32_t mask_none_id = builder->GetUintConstantId(SpvMemoryAccessMaskNone); + uint32_t scope_invok_id = builder->GetUintConstantId(SpvScopeInvocation); + Instruction* buffer_consumed_ptr = + builder->AddBinaryOp(buf_uint_ptr_id, SpvOpAccessChain, buf_id, + builder->GetUintConstantId(kDebugOutputSizeOffset)); + Instruction* write_offset = builder->AddQuadOp( + GetUintId(), SpvOpAtomicIAdd, buffer_consumed_ptr->result_id(), + scope_invok_id, mask_none_id, entry_size); + + Instruction* updated_consumed_value = builder->AddBinaryOp( + GetUintId(), SpvOpIAdd, write_offset->result_id(), entry_size); + updated_consumed_value = builder->AddBinaryOp( + GetUintId(), SpvOpIAdd, updated_consumed_value->result_id(), + builder->GetUintConstantId(reserved_words_count_)); + Instruction* buffer_capacity = + builder->AddIdLiteralOp(GetUintId(), SpvOpArrayLength, + GetOutputBufferId(), kDebugOutputDataOffset); + + Instruction* out_of_bounds_cond = builder->AddBinaryOp( + GetBoolId(), SpvOpUGreaterThanEqual, updated_consumed_value->result_id(), + buffer_capacity->result_id()); + + curr_block = + GenIfStatement(out_of_bounds_cond->result_id(), std::move(curr_block), + [this](InstructionBuilder* ir_builder, + std::unique_ptr block) { + ir_builder->AddNullaryOp(0, SpvOpReturn); + return block; + }); + builder->SetInsertPoint(&*curr_block); + + uint32_t i = reserved_words_count_; + for (auto it : element_ids) { + Instruction* element_offset = + builder->AddBinaryOp(GetUintId(), SpvOpIAdd, write_offset->result_id(), + builder->GetUintConstantId(i)); + Instruction* offset_ptr = builder->AddTernaryOp( + buf_uint_ptr_id, SpvOpAccessChain, buf_id, + builder->GetUintConstantId(kDebugOutputDataOffset), + element_offset->result_id()); + (void)builder->AddStore(offset_ptr->result_id(), it); + i++; + } + + return curr_block; +} + +uint32_t AutoInstPass::GetThreadUpdateFuncId(uint32_t num_parameters) { + if (param_cnt2thread_update_func_id_[num_parameters] == 0) { + // Create function + param_cnt2thread_update_func_id_[num_parameters] = TakeNextId(); + analysis::TypeManager* type_mgr = context()->get_type_mgr(); + std::vector param_types; + for (uint32_t c = 0; c < num_parameters; ++c) + param_types.push_back(type_mgr->GetType(GetUintId())); + analysis::Function func_ty(type_mgr->GetType(GetVoidId()), param_types); + analysis::Type* reg_func_ty = type_mgr->GetRegisteredType(&func_ty); + std::unique_ptr func_inst( + new Instruction(get_module()->context(), SpvOpFunction, GetVoidId(), + param_cnt2thread_update_func_id_[num_parameters], + {{spv_operand_type_t::SPV_OPERAND_TYPE_LITERAL_INTEGER, + {SpvFunctionControlMaskNone}}, + {spv_operand_type_t::SPV_OPERAND_TYPE_ID, + {type_mgr->GetTypeInstruction(reg_func_ty)}}})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_inst); + std::unique_ptr output_func = + MakeUnique(std::move(func_inst)); + // Add parameters + std::vector param_vec; + for (uint32_t c = 0; c < num_parameters; ++c) { + uint32_t pid = TakeNextId(); + param_vec.push_back(pid); + std::unique_ptr param_inst( + new Instruction(get_module()->context(), SpvOpFunctionParameter, + GetUintId(), pid, {})); + get_def_use_mgr()->AnalyzeInstDefUse(&*param_inst); + output_func->AddParameter(std::move(param_inst)); + } + + // Create first block + uint32_t test_blk_id = TakeNextId(); + std::unique_ptr test_label(NewLabel(test_blk_id)); + std::unique_ptr new_blk_ptr = + MakeUnique(std::move(test_label)); + new_blk_ptr->SetParent(&*output_func); + InstructionBuilder builder( + context(), &*new_blk_ptr, + IRContext::kAnalysisDefUse | IRContext::kAnalysisInstrToBlockMapping); + + builder.SetInsertPoint(&*new_blk_ptr); + + new_blk_ptr = GenThreadUpdate(&builder, std::move(new_blk_ptr), param_vec); + + builder.SetInsertPoint(&*new_blk_ptr); + + // Close merge block and function and add function to module + (void)builder.AddNullaryOp(0, SpvOpReturn); + new_blk_ptr->SetParent(&*output_func); + output_func->AddBasicBlock(std::move(new_blk_ptr)); + std::unique_ptr func_end_inst( + new Instruction(get_module()->context(), SpvOpFunctionEnd, 0, 0, {})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_end_inst); + output_func->SetFunctionEnd(std::move(func_end_inst)); + context()->AddFunction(std::move(output_func)); + } + return param_cnt2thread_update_func_id_[num_parameters]; +} + +void AutoInstPass::GenThreadUpdateCall(InstructionBuilder* builder, + std::vector param_ids) { + uint32_t func_id = GetThreadUpdateFuncId((uint32_t)param_ids.size()); + + std::vector operands = {func_id}; + operands.insert(operands.end(), param_ids.begin(), param_ids.end()); + + (void)builder->AddNaryOp(GetVoidId(), SpvOpFunctionCall, operands); +} + +std::unique_ptr AutoInstPass::GenSubgroupUpdate( + InstructionBuilder* builder, std::unique_ptr curr_block, + std::vector element_ids) { + Instruction* subgroup_leader_cond = + builder->AddUnaryOp(GetBoolId(), SpvOpGroupNonUniformElect, + builder->GetUintConstantId(SpvScopeSubgroup)); + curr_block = GenIfStatement( + subgroup_leader_cond->result_id(), std::move(curr_block), + [this, element_ids](InstructionBuilder* ir_builder, + std::unique_ptr block) { + block = GenThreadUpdate(ir_builder, std::move(block), element_ids); + ir_builder->SetInsertPoint(&*block); + return block; + }); + builder->SetInsertPoint(&*curr_block); + + return curr_block; +} + +uint32_t AutoInstPass::GetSubgroupUpdateFuncId(uint32_t num_parameters) { + if (param_cnt2subgroup_update_func_id_[num_parameters] == 0) { + // Create function + param_cnt2subgroup_update_func_id_[num_parameters] = TakeNextId(); + analysis::TypeManager* type_mgr = context()->get_type_mgr(); + std::vector param_types; + for (uint32_t c = 0; c < num_parameters; ++c) + param_types.push_back(type_mgr->GetType(GetUintId())); + analysis::Function func_ty(type_mgr->GetType(GetVoidId()), param_types); + analysis::Type* reg_func_ty = type_mgr->GetRegisteredType(&func_ty); + std::unique_ptr func_inst( + new Instruction(get_module()->context(), SpvOpFunction, GetVoidId(), + param_cnt2subgroup_update_func_id_[num_parameters], + {{spv_operand_type_t::SPV_OPERAND_TYPE_LITERAL_INTEGER, + {SpvFunctionControlMaskNone}}, + {spv_operand_type_t::SPV_OPERAND_TYPE_ID, + {type_mgr->GetTypeInstruction(reg_func_ty)}}})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_inst); + std::unique_ptr output_func = + MakeUnique(std::move(func_inst)); + // Add parameters + std::vector param_vec; + for (uint32_t c = 0; c < num_parameters; ++c) { + uint32_t pid = TakeNextId(); + param_vec.push_back(pid); + std::unique_ptr param_inst( + new Instruction(get_module()->context(), SpvOpFunctionParameter, + GetUintId(), pid, {})); + get_def_use_mgr()->AnalyzeInstDefUse(&*param_inst); + output_func->AddParameter(std::move(param_inst)); + } + + // Create first block + uint32_t test_blk_id = TakeNextId(); + std::unique_ptr test_label(NewLabel(test_blk_id)); + std::unique_ptr new_blk_ptr = + MakeUnique(std::move(test_label)); + new_blk_ptr->SetParent(&*output_func); + InstructionBuilder builder( + context(), &*new_blk_ptr, + IRContext::kAnalysisDefUse | IRContext::kAnalysisInstrToBlockMapping); + + builder.SetInsertPoint(&*new_blk_ptr); + + new_blk_ptr = GenSubgroupUpdate(&builder, std::move(new_blk_ptr), param_vec); + + builder.SetInsertPoint(&*new_blk_ptr); + + // Close merge block and function and add function to module + (void)builder.AddNullaryOp(0, SpvOpReturn); + new_blk_ptr->SetParent(&*output_func); + output_func->AddBasicBlock(std::move(new_blk_ptr)); + std::unique_ptr func_end_inst( + new Instruction(get_module()->context(), SpvOpFunctionEnd, 0, 0, {})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_end_inst); + output_func->SetFunctionEnd(std::move(func_end_inst)); + context()->AddFunction(std::move(output_func)); + } + return param_cnt2subgroup_update_func_id_[num_parameters]; +} + +void AutoInstPass::GenSubgroupUpdateCall(InstructionBuilder* builder, + std::vector param_ids) { + uint32_t func_id = GetSubgroupUpdateFuncId((uint32_t)param_ids.size()); + + std::vector operands = {func_id}; + operands.insert(operands.end(), param_ids.begin(), param_ids.end()); + + (void)builder->AddNaryOp(GetVoidId(), SpvOpFunctionCall, operands); +} + +void AutoInstPass::GenUniqueSubgroupIdFuncCall(InstructionBuilder* builder, + uint32_t inst_id, + uint32_t stage_idx) { + if (stage_idx != SpvExecutionModelRayGenerationNV && + stage_idx != SpvExecutionModelGLCompute) { + std::string message = + "Unique function id call cannot be generated unless the shader stage " + "is compute or RayGeneration\n"; + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, message.c_str()); + return; + } + + uint32_t output_func_id = GetUniqueSubgroupIdFunctionId(stage_idx); + (void)builder->AddNaryOp(GetVoidId(), SpvOpFunctionCall, + {output_func_id, inst_id}); +} + +uint32_t AutoInstPass::GetUniqueSubgroupIdFunctionId(uint32_t stage_idx) { + if (unique_subgroup_id_function_id_ == 0) { + // Create function + unique_subgroup_id_function_id_ = TakeNextId(); + analysis::TypeManager* type_mgr = context()->get_type_mgr(); + std::vector param_types; + + for (uint32_t c = 0; c < kUniqueSubgroupIdParamCnt; ++c) + param_types.push_back(type_mgr->GetType(GetUintId())); + + analysis::Function func_ty(type_mgr->GetType(GetVoidId()), param_types); + analysis::Type* reg_func_ty = type_mgr->GetRegisteredType(&func_ty); + std::unique_ptr func_inst( + new Instruction(get_module()->context(), SpvOpFunction, GetVoidId(), + unique_subgroup_id_function_id_, + {{spv_operand_type_t::SPV_OPERAND_TYPE_LITERAL_INTEGER, + {SpvFunctionControlMaskNone}}, + {spv_operand_type_t::SPV_OPERAND_TYPE_ID, + {type_mgr->GetTypeInstruction(reg_func_ty)}}})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_inst); + std::unique_ptr output_func = + MakeUnique(std::move(func_inst)); + + // Add parameters + std::vector param_vec; + for (uint32_t c = 0; c < kUniqueSubgroupIdParamCnt; ++c) { + uint32_t pid = TakeNextId(); + param_vec.push_back(pid); + std::unique_ptr param_inst( + new Instruction(get_module()->context(), SpvOpFunctionParameter, + GetUintId(), pid, {})); + get_def_use_mgr()->AnalyzeInstDefUse(&*param_inst); + output_func->AddParameter(std::move(param_inst)); + } + + // Create first block + uint32_t test_blk_id = TakeNextId(); + std::unique_ptr test_label(NewLabel(test_blk_id)); + std::unique_ptr new_blk_ptr = + MakeUnique(std::move(test_label)); + new_blk_ptr->SetParent(&*output_func); + InstructionBuilder builder( + context(), &*new_blk_ptr, + IRContext::kAnalysisDefUse | IRContext::kAnalysisInstrToBlockMapping); + + // Create variable to hold subgroup id computed by leader + + uint32_t varTyPtrId = context()->get_type_mgr()->FindPointerToType( + GetUintId(), SpvStorageClassFunction); + assert(varTyPtrId && "Cannot create uint ptr type."); + auto zero = builder.GetUintConstantId(0); + + auto new_var_op = + builder.AddUnaryOp(varTyPtrId, SpvOpVariable, SpvStorageClassFunction); + auto unique_subgroup_ptr_id = new_var_op->result_id(); + builder.AddStore(new_var_op->result_id(), zero); + + Instruction* subgroup_leader_cond = + builder.AddUnaryOp(GetBoolId(), SpvOpGroupNonUniformElect, + builder.GetUintConstantId(SpvScopeSubgroup)); + + new_blk_ptr = GenIfStatement( + subgroup_leader_cond->result_id(), std::move(new_blk_ptr), + [this, unique_subgroup_ptr_id](InstructionBuilder* ir_builder, + std::unique_ptr block) { + uint32_t mask_none_id = + ir_builder->GetUintConstantId(SpvMemoryAccessMaskNone); + uint32_t scope_invok_id = + ir_builder->GetUintConstantId(SpvScopeInvocation); + Instruction* unique_id_ptr = ir_builder->AddTernaryOp( + GetOutputBufferPtrId(), SpvOpAccessChain, GetOutputBufferId(), + ir_builder->GetUintConstantId(kDebugOutputDataOffset), + ir_builder->GetUintConstantId(0)); + Instruction* unique_id = ir_builder->AddQuadOp( + GetUintId(), SpvOpAtomicIAdd, unique_id_ptr->result_id(), + scope_invok_id, mask_none_id, ir_builder->GetUintConstantId(1)); + + ir_builder->AddStore(unique_subgroup_ptr_id, unique_id->result_id()); + return block; + }); + builder.SetInsertPoint(&*new_blk_ptr); + Instruction* broadcasted_id = + builder.AddBinaryOp(GetUintId(), SpvOpGroupNonUniformBroadcastFirst, + builder.GetUintConstantId(SpvScopeSubgroup), + GenVarLoad(unique_subgroup_ptr_id, &builder)); + + uint32_t intra_subgroup_id = GenSubgroupLocalInvocationId(&builder); + // Shift the thread id in the subgroup in to the top log2(SUBGROUP_SIZE)=5 bits + Instruction* shifted_subgroup_id = builder.AddBinaryOp( + GetUintId(), SpvOpShiftLeftLogical, intra_subgroup_id, + builder.GetUintConstantId(27 /*= 32 - log2(32) */)); + // Combine the unique subgroup id and intra subgroup id + Instruction* joined_subgroup_ids = builder.AddBinaryOp( + GetUintId(), SpvOpBitwiseOr, shifted_subgroup_id->result_id(), + broadcasted_id->result_id()); + + // Generate thread id which will be used to created thread_id -> subgroup_id + // mapping + auto flat_thread_id = (stage_idx == SpvExecutionModelRayGenerationNV) + ? GenFlatRtThreadId(&builder, stage_idx) + : GenFlatComputeThreadId(&builder, stage_idx); + + auto inst_id = param_vec[kUniqueSubgroupIdParamInstIdIdx]; + + new_blk_ptr = GenThreadUpdate( + &builder, std::move(new_blk_ptr), + {inst_id, flat_thread_id, joined_subgroup_ids->result_id()}); + + // Close merge block and function and add function to module + (void)builder.AddNullaryOp(0, SpvOpReturn); + new_blk_ptr->SetParent(&*output_func); + output_func->AddBasicBlock(std::move(new_blk_ptr)); + std::unique_ptr func_end_inst( + new Instruction(get_module()->context(), SpvOpFunctionEnd, 0, 0, {})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_end_inst); + output_func->SetFunctionEnd(std::move(func_end_inst)); + context()->AddFunction(std::move(output_func)); + } + return unique_subgroup_id_function_id_; +} + +void AutoInstPass::GenInstrumentedEntryPoints() { + for (auto entry_point_inst : get_module()->entry_points()) { + auto stage_idx = + entry_point_inst.GetSingleWordInOperand(kEntryPointExecutionModelInIdx); + auto entry_point_func_id = + entry_point_inst.GetSingleWordInOperand(kEntryPointFunctionIdInIdx); + Instruction* entry_point_func = + get_def_use_mgr()->GetDef(entry_point_func_id); + + auto dummy_func_id = TakeNextId(); + analysis::TypeManager* type_mgr = context()->get_type_mgr(); + analysis::Function func_ty(type_mgr->GetType(GetVoidId()), {}); + analysis::Type* reg_func_ty = type_mgr->GetRegisteredType(&func_ty); + auto expected_ty_id = type_mgr->GetId(reg_func_ty); + + auto entry_point_func_ty_id = + entry_point_func->GetSingleWordOperand(kFunctionTypeIdx); + + if (expected_ty_id != entry_point_func_ty_id) { + std::string message = + "Could not generate dummy entrypoint due to an unexpected EntryPoint " + "function signature."; + consumer()(SPV_MSG_ERROR, 0, {0, 0, 0}, message.c_str()); + return; + } + + // Create dummy function to original entry point + std::unique_ptr func_inst(new Instruction( + get_module()->context(), SpvOpFunction, GetVoidId(), dummy_func_id, + {{spv_operand_type_t::SPV_OPERAND_TYPE_LITERAL_INTEGER, + {SpvFunctionControlMaskNone}}, + {spv_operand_type_t::SPV_OPERAND_TYPE_ID, + {type_mgr->GetTypeInstruction(reg_func_ty)}}})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_inst); + + std::unique_ptr output_func = + MakeUnique(std::move(func_inst)); + + // Create first block + uint32_t test_blk_id = TakeNextId(); + std::unique_ptr test_label(NewLabel(test_blk_id)); + std::unique_ptr new_blk_ptr = + MakeUnique(std::move(test_label)); + new_blk_ptr->SetParent(&*output_func); + InstructionBuilder builder( + context(), &*new_blk_ptr, + IRContext::kAnalysisDefUse | IRContext::kAnalysisInstrToBlockMapping); + + builder.SetInsertPoint(&*new_blk_ptr); + + // Invoke instrumentation hook + auto is_instrumented = PreEntryPointInstrument(&builder, stage_idx); + + // Call original entrypoint + (void)builder.AddNaryOp(GetVoidId(), SpvOpFunctionCall, + {entry_point_func_id}); + + // Invoke instrumentation hook + is_instrumented |= PostEntryPointInstrument(&builder, stage_idx); + + // Close merge block and function and add function to module + (void)builder.AddNullaryOp(0, SpvOpReturn); + new_blk_ptr->SetParent(&*output_func); + output_func->AddBasicBlock(std::move(new_blk_ptr)); + std::unique_ptr func_end_inst( + new Instruction(get_module()->context(), SpvOpFunctionEnd, 0, 0, {})); + get_def_use_mgr()->AnalyzeInstDefUse(&*func_end_inst); + output_func->SetFunctionEnd(std::move(func_end_inst)); + + if (is_instrumented) { + // If the instrumentation hooks insert code then + // add dummy entrypoint and replace the original + // EntryPoint with dummy entrypoint. + context()->AddFunction(std::move(output_func)); + context()->ReplaceAllUsesWithPredicate( + entry_point_func_id, dummy_func_id, [](Instruction* inst) { + return inst->opcode() != SpvOpFunctionCall; + }); + } + } +} + +Pass::Status AutoInstPass::ProcessImpl() { + for (auto fii = get_module()->begin(); fii != get_module()->end(); ++fii) { + auto bb = fii->begin(); + bb->IsLoopHeader(); + } + + InstProcessFunction pfn = + [this](BasicBlock::iterator ref_inst_itr, + UptrVectorIterator ref_block_itr, uint32_t stage_idx, + std::vector>* new_blocks) { + (void)new_blocks; + GenInstProgrammableCode(ref_inst_itr, ref_block_itr, stage_idx); + }; + InstProcessEntryPointCallTree(pfn); + + // Add new entrypoint after other instrumentation to avoid it also being + // instrumented. + GenInstrumentedEntryPoints(); + + context()->BuildInvalidAnalyses(IRContext::kAnalysisDefUse | + IRContext::kAnalysisInstrToBlockMapping); + return has_added_instrumentation_ ? Status::SuccessWithChange + : Status::SuccessWithoutChange; +} + +Pass::Status AutoInstPass::Process() { + // Initialize base class + InitializeInstrument(); + + // init auto instrumentation metadata + instrumented_bb_ids.clear(); + instrumented_inst_ids.clear(); + has_added_instrumentation_ = false; + + // initialize inheriting class + InitializeInstrumentation(); + + auto res = ProcessImpl(); + + // finalize inheriting class + FinalizeInstrumentation(); + + // insert instrumentation + return res; +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_pass.h b/source/opt/auto_inst_pass.h new file mode 100644 index 00000000..ed91a44d --- /dev/null +++ b/source/opt/auto_inst_pass.h @@ -0,0 +1,322 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_PASS_H_ + +#include + +#include "instrument_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstPass : public InstrumentPass { + public: + AutoInstPass(uint32_t desc_set, uint32_t shader_id, + uint32_t reserved_words_count) + : InstrumentPass(desc_set, shader_id, kInstValidationIdAuto), + reserved_words_count_(reserved_words_count) {} + + ~AutoInstPass() override = default; + + // See optimizer.hpp for pass user documentation. + Status Process() override; + + const char* name() const override { return "auto-inst-pass"; } + + private: + // Track whether any instrumentation occurred + bool has_added_instrumentation_ = false; + + // Track the function ids created to support thread/subgroup update + // accepting different numbers of parameters + std::unordered_map param_cnt2thread_update_func_id_; + std::unordered_map param_cnt2subgroup_update_func_id_; + + // Track the function id for creating unique subgroup ids. + uint32_t unique_subgroup_id_function_id_ = 0; + + // Instructions whose semantics are destroyed by having instrumentation + // inserted directly before them. + const std::set uninstrumentable_pre_ops = {SpvOpPhi, + SpvOpUnreachable}; + + // Instructions whose semantics are destroyed by having instrumentation + // inserted directly after them. + const std::set uninstrumentable_post_ops = { + SpvOpSelectionMerge, // Both merge instructions must precede a branch + SpvOpLoopMerge, + SpvOpBranch, // Branch instructions end a basic block which would put the + // instrumentation in limbo + SpvOpBranchConditional, + SpvOpSwitch, + SpvOpReturn, // Instrumenting after returns/unreachable would never + // execute + SpvOpReturnValue, + SpvOpUnreachable, + }; + + // Add a function to the module that appends + // an entry to the buffer containing |num_params| words + // for every active thread invoking the function. + uint32_t GetThreadUpdateFuncId(uint32_t num_params); + + // Generate a sequence of instructions in |builder| in function |function| + // that write the value corresponding to the ids in |element_ids|. The + // output buffer will be written to once by each thread. + // The sequence of instructions will be inserted starting with |curr_block|. + // |curr_block| will be consumed by this function call due to control flow. + // + // |element_ids| is a vector of uint32_t which contains the + // id's of values that will be written to the output buffer. + std::unique_ptr GenThreadUpdate( + InstructionBuilder* builder, std::unique_ptr curr_block, + std::vector param_ids); + + // Add a function to the module that appends + // an entry to the buffer containing |num_params| words + // for every subgroup invoking the function. + uint32_t GetSubgroupUpdateFuncId(uint32_t num_params); + + // Create function containing functionality for generating a + // unique subgroup (or subgroup) id. This function should only + // be called at the beginning of a shader in uniform control flow. + // + // The buffer entry created will be of the following form: + // word 0: + // word 1: + // word 2: + // + // This information can be used by the analysis to create a mapping + // from flattened thread id (available anywhere in the rt pipeline) + // to subgroup id which allows for inter-shader subgroup tracking. + // Furthermore, the intra-subgroup-id allows for attribution of subgroup-level + // instrumentation to individual threads (i.e. for heatmap visualizations). + uint32_t GetUniqueSubgroupIdFunctionId(uint32_t stage_idx); + + // Generate a sequence of instructions in |builder| in function |function| + // that write the value corresponding to the ids in |element_ids|. The + // output buffer will only be written to by the subgroup leader. + // The sequence of instructions will be inserted starting with |curr_block|. + // |curr_block| will be consumed by this function call due to control flow. + // + // |element_ids| is a vector of uint32_t which contains the + // id's of values that will be written to the output buffer. + std::unique_ptr GenSubgroupUpdate( + InstructionBuilder* builder, std::unique_ptr curr_block, + std::vector param_ids); + + protected: + // In this class it can be very confusing differentiating + // between instruction SSA ids and ids for instrumentation. + // This type is designed to make it explicit which type of id + // it is when mixing and matching. + using AutoInstId = uint32_t; + + // For some analyses keeping track of which threads belong + // to which subgroups and also how many subgroups executed the shader stage + // is interesting. Since saving 1 word is not important, + // the default for this value is set to 1 so CreateUniquesubgroupIdCall + // works out of the box. + static const int kDefaultReservedWordsCnt = 1; + + // The number of lowers words in the instrumentation buffer that are reserved + // for fixed functions (i.e. not dynamically appended runtime entries) + // NOTE: this does not include the buffer size which is always tracked + const uint32_t reserved_words_count_; + + // Track which basic blocks and instructions the pass has + // given an opportunity to instrument to prevent reinstrumenting. + std::set instrumented_bb_ids; + std::set instrumented_inst_ids; + + // Apply GenDebugPrintfCode to every instruction in module. + Pass::Status ProcessImpl(); + + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + virtual void InitializeInstrumentation() = 0; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + virtual void FinalizeInstrumentation() = 0; + + // Any instructions added via |builder| will appear before |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + virtual bool PreInstructionInstrument(Instruction* inst, + InstructionBuilder* builder, + uint32_t stage_idx) = 0; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + virtual bool PostInstructionInstrument(Instruction* inst, + InstructionBuilder* builder, + uint32_t stage_idx) = 0; + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + virtual bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) = 0; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + virtual bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) = 0; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + virtual bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) = 0; + + // If |ref_inst_itr| is selected as an instrumentation location, + // return in |new_blocks| the result of adding instrumentation. + // The instructions write a record to the output buffer stream + // The block at |ref_block_itr| can just be replaced with the + // block in |new_blocks|. Besides the buffer writes, this + // block will comprise all instructions preceding and following + // |ref_inst_itr|. + // + // This function is designed to be passed to + // InstrumentPass::InstProcessEntryPointCallTree(), which applies the + // function to each instruction in a module and replaces the instruction + // if warranted. + // + void GenInstProgrammableCode(BasicBlock::iterator ref_inst_itr, + UptrVectorIterator ref_block_itr, + uint32_t stage_idx); + + // Generate a sequence of instructions in |builder| which + // materialize the lower and upper 32 bits of from + // OpReadClock. + // + // Returns std::pair + std::pair GenReadClockIds(InstructionBuilder* builder); + + // Generate a sequence of instructions in |builder| which + // materialize a 32-bit thread mask where each bit + // is true if the thread is active and false otherwise. + // The uint32_t value returned corresponds to the thread_mask. + uint32_t GenSubgroupBallotId(InstructionBuilder* builder, uint32_t pred_id); + + // Generate a sequence of instructions in |builder| which + // materialize the value of SpvBuiltinSubgroupLocalInvocationId. + uint32_t GenSubgroupLocalInvocationId(InstructionBuilder* builder); + + // Returns an id corresponding to a uint created in |builder| + // which contains a flattened thread id calculated from the + // unique work id of the shader stage. + // + // This is primarily useful for tracking threads' execution between + // shaders in the ray tracing pipeline. + uint32_t GenFlatRtThreadId(InstructionBuilder* builder, uint32_t stage_idx); + + // Returns an id corresponding to a uint created in |builder| + // which contains a flattened thread id calculated from the + // GlobalSize and GlobalLaunchID + // + // This is primarily useful for tracking threads' execution between + // different compute pipelines. + uint32_t GenFlatComputeThreadId(InstructionBuilder* builder, + uint32_t stage_idx); + + // Returns a vector of ids corresponding to a uint created in |builder| + // which contains a unique work id of the shader stage. + // + // This is primarily useful for tracking threads' execution behaviour over + // time. + std::vector GenThreadId(InstructionBuilder* builder, + uint32_t stage_idx); + + // Returns an identifier + // for an instrumentation callsite which is unique across the + // whole ray-tracing pipeline. + uint32_t GenInstCallsiteId(Instruction* inst); + + // Generate a sequence of instructions in function |function| that + // create an if statement where the body is executed iff the value + // corresponding to |condition_id| evaluates to true at runtime. + // |old_block| will be closed by an OpBranchConditional + // + // The callback |inside_if_callback| will be invoked in the body + // of the if statement. The |inside_if_callback| accepts an + // InstructionBuilder |builder| at the beginning of the if body. + // As well as a unique_ptr |curr_block| to the BasicBlock of the + // if body. The |inside_if_callback| may add more basic blocks + // but must return a unique_ptr to the basic block that ends the if + // body. + std::unique_ptr GenIfStatement( + uint32_t condition_id, std::unique_ptr old_block, + std::function( + InstructionBuilder* builder, std::unique_ptr curr_block)> + inside_if_callback); + + // Generates a seuqence of instructions in |builder| which invoke the + // ThreadUpdate function which writes the values that are identifier in + // |param_ids| to the StorageBuffer for each thread that invokes the call. + void GenThreadUpdateCall(InstructionBuilder* builder, + std::vector param_ids); + + // Generates a seuqence of instructions in |builder| which invoke the + // subgroupUpdate function which writes the values that are identifier in + // |param_ids| to the StorageBuffer for each subgroup that invokes the call. + void GenSubgroupUpdateCall(InstructionBuilder* builder, + std::vector param_ids); + + // Generate a function call in a block which will be appended to |new_blocks| + // This function should only bGe called at the beginning of a shader in + // uniform control flow. This ensures that every thread in the subgroup + // receives the value computed by the leader. + // + // |inst_offset_id| is used to report instrumentation metadata to validation + // layer. |stage_idx| is the current SpvExecutionMode. + void GenUniqueSubgroupIdFuncCall(InstructionBuilder* builder, + uint32_t inst_offset_id, uint32_t stage_idx); + + // Pass the current context in terms of: + // 1) instruction in |ref_inst_itr| + // 2) BB in |ref_block_itr| + // 3) shader stage in |stage_idx| + // + // This allows the instrumentation hooks to decide what + // instrumentation to add to |builder|. + // If instrumentation is added then this function returns true + // otherwise false. + // + bool HandleInstrumentHooks(BasicBlock::iterator ref_inst_itr, + UptrVectorIterator ref_block_itr, + uint32_t stage_idx, InstructionBuilder* builder); + + // Generate dummy EntryPoints which invoke the PreEntryPointInstrument + // and PostEntryPointInstrument hooks around a call to the original entrypoint + // function. + void GenInstrumentedEntryPoints(); +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_INST_PROGRAMMABLE_PASS_H_ diff --git a/source/opt/auto_inst_simt_efficiency_pass.cpp b/source/opt/auto_inst_simt_efficiency_pass.cpp new file mode 100644 index 00000000..cfb7e38a --- /dev/null +++ b/source/opt/auto_inst_simt_efficiency_pass.cpp @@ -0,0 +1,39 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + + +#include "auto_inst_simt_efficiency_pass.h" + +namespace spvtools { +namespace opt { + +bool AutoInstSimtEfficiencyPass::BasicBlockInstrument( + BasicBlock* bb, InstructionBuilder* builder, uint32_t stage_idx) { + // Suppress unused parameter warnings + (void)bb; + (void)stage_idx; + + Instruction* true_constant_op = builder->GetBoolConstant(true); + uint32_t true_constant_id = true_constant_op->result_id(); + // Create active thread mask by having all threads vote true + uint32_t active_thread_mask_id = + GenSubgroupBallotId(builder, true_constant_id); + // Write active thread mask + GenSubgroupUpdateCall(builder, {active_thread_mask_id}); + + return true; +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_simt_efficiency_pass.h b/source/opt/auto_inst_simt_efficiency_pass.h new file mode 100644 index 00000000..3d2b58cf --- /dev/null +++ b/source/opt/auto_inst_simt_efficiency_pass.h @@ -0,0 +1,101 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_SIMT_EFFICIENCY_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_SIMT_EFFICIENCY_PASS_H_ + +#include + +#include "auto_inst_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstSimtEfficiencyPass : public AutoInstPass { + public: + AutoInstSimtEfficiencyPass(uint32_t desc_set, uint32_t shader_id, + uint32_t reserved_words_count) + : AutoInstPass(desc_set, shader_id, reserved_words_count) {} + + const char* name() const override { return "auto-inst-simt-efficiency-pass"; } + + private: + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + void InitializeInstrumentation() override{}; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + void FinalizeInstrumentation() override{}; + + // Any instructions added via |builder| will appear before |inst| + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)builder; + (void)stage_idx; + return false; + } + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)builder; + (void)stage_idx; + return false; + }; +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_AUTO_INST_SIMT_EFFICIENCY_PASS_H_ diff --git a/source/opt/auto_inst_warp_entry_and_exit_pass.cpp b/source/opt/auto_inst_warp_entry_and_exit_pass.cpp new file mode 100644 index 00000000..d985a70d --- /dev/null +++ b/source/opt/auto_inst_warp_entry_and_exit_pass.cpp @@ -0,0 +1,43 @@ +// Copyright (c) 2021 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "auto_inst_warp_entry_and_exit_pass.h" + +namespace spvtools { +namespace opt { + +bool AutoInstWarpEntryAndExitPass::PreEntryPointInstrument( + InstructionBuilder* builder, uint32_t stage_idx) { + if (stage_idx != SpvExecutionModelRayGenerationKHR) return false; + + // Record every warp that began the pipeline + auto prim_id = + builder->GetUintConstantId(kAutoInstWarpEntryAndExitBeginPipeline); + GenSubgroupUpdateCall(builder, {prim_id}); + return true; +} + +bool AutoInstWarpEntryAndExitPass::PostEntryPointInstrument( + InstructionBuilder* builder, uint32_t stage_idx) { + if (stage_idx != SpvExecutionModelRayGenerationKHR) return false; + + // Record every warp that completed the the pipeline + auto prim_id = + builder->GetUintConstantId(kAutoInstWarpEntryAndExitEndPipeline); + GenSubgroupUpdateCall(builder, {prim_id}); + return true; +} + +} // namespace opt +} // namespace spvtools diff --git a/source/opt/auto_inst_warp_entry_and_exit_pass.h b/source/opt/auto_inst_warp_entry_and_exit_pass.h new file mode 100644 index 00000000..44a5d175 --- /dev/null +++ b/source/opt/auto_inst_warp_entry_and_exit_pass.h @@ -0,0 +1,99 @@ +// Copyright (c) 2021 The Khronos Group Inc. + +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef LIBSPIRV_OPT_AUTO_INST_WARP_ENTRY_AND_EXIT_PASS_H_ +#define LIBSPIRV_OPT_AUTO_INST_WARP_ENTRY_AND_EXIT_PASS_H_ + +#include + +#include "auto_inst_pass.h" + +namespace spvtools { +namespace opt { + +class AutoInstWarpEntryAndExitPass : public AutoInstPass { + public: + AutoInstWarpEntryAndExitPass(uint32_t desc_set, uint32_t shader_id) + : AutoInstPass(desc_set, shader_id, kDefaultReservedWordsCnt) {} + + const char* name() const override { + return "auto-inst-warp-entry-and-exit-pass"; + } + + private: + // Allows inheriting classes to initialize their knowledge + // of module before beginning instrumentation + void InitializeInstrumentation() override{}; + + // Allows inheriting classes to finalize before + // the pass finishes executing. + void FinalizeInstrumentation() override{}; + + // Any instructions added via |builder| will appear before |inst| + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear after |inst|. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostInstructionInstrument(Instruction* inst, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)inst; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will appear before the content of + // |bb|. |stage_idx| contains the SpvExecutionModel that builder is operating + // in. This function is expected to return true if it added instructions to + // builder, otherwise false. + bool BasicBlockInstrument(BasicBlock* bb, InstructionBuilder* builder, + uint32_t stage_idx) override { + (void)bb; + (void)builder; + (void)stage_idx; + return false; + }; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PreEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override; + + // Any instructions added via |builder| will execute before the + // entrypoint function of the shader. + // |stage_idx| contains the SpvExecutionModel that builder is operating in. + // This function is expected to return true if it added instructions to + // builder, otherwise false. + bool PostEntryPointInstrument(InstructionBuilder* builder, + uint32_t stage_idx) override; +}; + +} // namespace opt +} // namespace spvtools + +#endif // LIBSPIRV_OPT_AUTO_INST_WARP_ENTRY_AND_EXIT_PASS_H_ diff --git a/source/opt/instrument_pass.cpp b/source/opt/instrument_pass.cpp index ed34fb02..1a84a002 100644 --- a/source/opt/instrument_pass.cpp +++ b/source/opt/instrument_pass.cpp @@ -288,7 +288,9 @@ void InstrumentPass::GenStageStreamWriteCode(uint32_t stage_idx, GenDebugOutputFieldCode(base_offset_id, kInstRayTracingOutLaunchIdZ, z_launch_inst->result_id(), builder); } break; - default: { assert(false && "unsupported stage"); } break; + default: { + assert(false && "unsupported stage"); + } break; } } @@ -435,6 +437,7 @@ uint32_t InstrumentPass::GetOutputBufferBinding() { case kInstValidationIdBuffAddr: return kDebugOutputBindingStream; case kInstValidationIdDebugPrintf: + case kInstValidationIdAuto: return kDebugOutputPrintfStream; default: assert(false && "unexpected validation id"); diff --git a/source/opt/instrument_pass.h b/source/opt/instrument_pass.h index 12b939d4..5fb49f8b 100644 --- a/source/opt/instrument_pass.h +++ b/source/opt/instrument_pass.h @@ -62,6 +62,7 @@ namespace opt { static const uint32_t kInstValidationIdBindless = 0; static const uint32_t kInstValidationIdBuffAddr = 1; static const uint32_t kInstValidationIdDebugPrintf = 2; +static const uint32_t kInstValidationIdAuto = 3; class InstrumentPass : public Pass { using cbb_ptr = const BasicBlock*; diff --git a/source/opt/ir_builder.h b/source/opt/ir_builder.h index fe5feff5..b7c8e0c5 100644 --- a/source/opt/ir_builder.h +++ b/source/opt/ir_builder.h @@ -392,6 +392,10 @@ class InstructionBuilder { return uint_inst->result_id(); } + uint32_t GetIntConstantId(int value) { + return GetIntConstant(value, true)->result_id(); + } + // Adds either a signed or unsigned 32 bit integer constant to the binary // depedning on the |sign|. If |sign| is true then the value is added as a // signed constant otherwise as an unsigned constant. If |sign| is false the @@ -428,6 +432,28 @@ class InstructionBuilder { return GetContext()->get_constant_mgr()->GetDefiningInstruction(constant); } + Instruction* GetBoolConstant(bool value) { + analysis::Bool bool_type{}; + + // Get or create the integer type. This rebuilds the type and manages the + // memory for the rebuilt type. + uint32_t type_id = + GetContext()->get_type_mgr()->GetTypeInstruction(&bool_type); + + // Get the memory managed type so that it is safe to be stored by + // GetConstant. + analysis::Type* rebuilt_type = + GetContext()->get_type_mgr()->GetType(type_id); + + // Create the constant value. + const analysis::Constant* constant = + GetContext()->get_constant_mgr()->GetConstant(rebuilt_type, + {(uint32_t)value}); + + // Create the OpConstant instruction using the type and the value. + return GetContext()->get_constant_mgr()->GetDefiningInstruction(constant); + } + Instruction* AddCompositeExtract(uint32_t type, uint32_t id_of_composite, const std::vector& index_list) { std::vector operands; diff --git a/source/opt/ir_context.cpp b/source/opt/ir_context.cpp index 82107b5c..094513c3 100644 --- a/source/opt/ir_context.cpp +++ b/source/opt/ir_context.cpp @@ -805,6 +805,15 @@ uint32_t IRContext::GetBuiltinInputVarId(uint32_t builtin) { reg_type = type_mgr->GetRegisteredType(&v4float_ty); break; } + case SpvBuiltInNumSubgroups: + case SpvBuiltInWorkgroupId: + case SpvBuiltInWarpIDNV: + case SpvBuiltInWarpsPerSMNV: + case SpvBuiltInSMIDNV: + case SpvBuiltInSMCountNV: + case SpvBuiltInSubgroupId: + case SpvBuiltInLocalInvocationIndex: + case SpvBuiltInSubgroupSize: case SpvBuiltInVertexIndex: case SpvBuiltInInstanceIndex: case SpvBuiltInPrimitiveId: @@ -814,7 +823,11 @@ uint32_t IRContext::GetBuiltinInputVarId(uint32_t builtin) { reg_type = type_mgr->GetRegisteredType(&uint_ty); break; } + case SpvBuiltInLocalInvocationId: + case SpvBuiltInLaunchSizeNV: case SpvBuiltInGlobalInvocationId: + case SpvBuiltInNumWorkgroups: + case SpvBuiltInWorkgroupSize: case SpvBuiltInLaunchIdNV: { analysis::Integer uint_ty(32, false); analysis::Type* reg_uint_ty = type_mgr->GetRegisteredType(&uint_ty); diff --git a/source/opt/optimizer.cpp b/source/opt/optimizer.cpp index 8726ff93..5497fab7 100644 --- a/source/opt/optimizer.cpp +++ b/source/opt/optimizer.cpp @@ -908,6 +908,66 @@ Optimizer::PassToken CreateInstDebugPrintfPass(uint32_t desc_set, MakeUnique(desc_set, shader_id)); } +Optimizer::PassToken CreateAutoInstDebugPass(uint32_t desc_set, + uint32_t shader_id, + bool test_atomic_ops, + bool test_subgroup_ops) { + return MakeUnique( + MakeUnique(desc_set, shader_id, test_atomic_ops, + test_subgroup_ops)); +} + +Optimizer::PassToken CreateAutoInstDivergenceCharacterizationPass( + uint32_t desc_set, uint32_t shader_id, + std::function< + void(std::unordered_map&& inst_id2prim_id, + std::unordered_map&& inst_id2inst_count)> + static_data_callback) { + return MakeUnique( + MakeUnique( + desc_set, shader_id, static_data_callback)); +} + +Optimizer::PassToken CreateAutoInstDynShaderTracePass(uint32_t desc_set, + uint32_t shader_id) { + return MakeUnique( + MakeUnique(desc_set, shader_id)); +} + +Optimizer::PassToken CreateAutoInstDynTraceRayTracePass( + uint32_t desc_set, uint32_t shader_id, + std::function&&, + std::unordered_map>&&)> + static_data_callback) { + return MakeUnique( + MakeUnique(desc_set, shader_id, + static_data_callback)); +} + +Optimizer::PassToken CreateAutoInstExecutionTracePass( + uint32_t desc_set, uint32_t shader_id, + std::function< + void(std::unordered_map>&&, + std::unordered_map&& inst_id2bb_opcodes)> + static_data_callback) { + return MakeUnique( + MakeUnique(desc_set, shader_id, + static_data_callback)); +} + +Optimizer::PassToken CreateAutoInstSimtEfficiencyPass( + uint32_t desc_set, uint32_t shader_id, uint32_t reserved_words_count) { + return MakeUnique( + MakeUnique(desc_set, shader_id, + reserved_words_count)); +} + +Optimizer::PassToken CreateAutoInstWarpEntryAndExitPass(uint32_t desc_set, + uint32_t shader_id) { + return MakeUnique( + MakeUnique(desc_set, shader_id)); +} + Optimizer::PassToken CreateInstBuffAddrCheckPass(uint32_t desc_set, uint32_t shader_id) { return MakeUnique( diff --git a/source/opt/passes.h b/source/opt/passes.h index d47cc1ce..9a7c9c22 100644 --- a/source/opt/passes.h +++ b/source/opt/passes.h @@ -19,6 +19,13 @@ #include "source/opt/aggressive_dead_code_elim_pass.h" #include "source/opt/amd_ext_to_khr.h" +#include "source/opt/auto_inst_debug_pass.h" +#include "source/opt/auto_inst_divergence_characterization_pass.h" +#include "source/opt/auto_inst_dyn_trace_ray_trace_pass.h" +#include "source/opt/auto_inst_dyn_shader_trace_pass.h" +#include "source/opt/auto_inst_execution_trace_pass.h" +#include "source/opt/auto_inst_simt_efficiency_pass.h" +#include "source/opt/auto_inst_warp_entry_and_exit_pass.h" #include "source/opt/block_merge_pass.h" #include "source/opt/ccp_pass.h" #include "source/opt/cfg_cleanup_pass.h" -- 2.29.2.windows.2