diff --git a/modules/ocarina/ast/function.cpp b/modules/ocarina/ast/function.cpp index bf36c468..5d3ec3d3 100644 --- a/modules/ocarina/ast/function.cpp +++ b/modules/ocarina/ast/function.cpp @@ -3,6 +3,7 @@ // #include "function.h" +#include "core/util/switch_profile.h" #include "core/type_system/type_desc.h" #include "core/util/util.h" #include "generator/source_emitter.h" @@ -45,6 +46,7 @@ ocarina::list Function::resolve_ast_templates(ocarina::list< } void Function::correct() noexcept { + switch_profile::Scope profile{"Function::correct", "ast_correct"}; TIMER(FunctionCorrect) FunctionCorrector().apply(this); } @@ -622,4 +624,4 @@ uint64_t Function::compute_hash() const noexcept { return ret; } -}// namespace ocarina \ No newline at end of file +}// namespace ocarina diff --git a/modules/ocarina/ast/function.h b/modules/ocarina/ast/function.h index 92f2d5c8..fef8a8f6 100644 --- a/modules/ocarina/ast/function.h +++ b/modules/ocarina/ast/function.h @@ -9,6 +9,7 @@ #include "core/header.h" #include "core/type.h" #include "core/type_system/precision_policy.h" +#include "core/util/switch_profile.h" #include "expression.h" #include "statement.h" #include "ast_node.h" @@ -257,6 +258,7 @@ class OC_AST_API Function : public concepts::Noncopyable, public Hashable { } template static shared_ptr define_kernel(Func &&func) noexcept { + switch_profile::Scope profile{"Function::define_kernel", "dsl"}; shared_ptr function = _define(Tag::KERNEL, std::forward(func)); function->correct(); return function; @@ -339,4 +341,4 @@ class OC_AST_API Function : public concepts::Noncopyable, public Hashable { [[nodiscard]] constexpr const Type *return_type() const noexcept { return ret_; } constexpr void set_raytracing(bool val) const noexcept { raytracing_ = val; } }; -}// namespace ocarina \ No newline at end of file +}// namespace ocarina diff --git a/modules/ocarina/backends/cuda/cuda_compiler.cpp b/modules/ocarina/backends/cuda/cuda_compiler.cpp index 270e9a71..b05e1af4 100644 --- a/modules/ocarina/backends/cuda/cuda_compiler.cpp +++ b/modules/ocarina/backends/cuda/cuda_compiler.cpp @@ -3,6 +3,7 @@ // #include "cuda_compiler.h" +#include "core/util/switch_profile.h" #include "cuda_device.h" #include "ast/function.h" #include "ast_to_cuda_source.h" @@ -30,6 +31,7 @@ constexpr unsigned ast_to_cuda_source_cache_version = 1u; } [[nodiscard]] string emit_cuda_source(const Function &function) noexcept { + switch_profile::Scope profile{"emit_cuda_source", "codegen"}; switch (Env::shader_codegen_path()) { case ShaderCodegenPath::EAstToSource: { AstToCudaSource emitter{Env::code_obfuscation()}; @@ -66,89 +68,107 @@ CUDACompiler::CUDACompiler(CUDADevice *device) : device_(device) {} ocarina::string CUDACompiler::compile(const Function &function, int sm) const noexcept { - - fs::path cuda_path = get_cuda_path(); - cuda_path = cuda_path / "include"; - - string header_path = "-I" + cuda_path.string(); - int ver_major = 0; - int ver_minor = 0; - OC_NVRTC_CHECK(nvrtcVersion(&ver_major, &ver_minor)); - int nvrtc_version = ver_major * 10000 + ver_minor * 100; - const auto storage_policy = function.storage_policy(); - const int use_float32 = storage_policy.policy == PrecisionPolicy::force_f16 ? 0 : 1; - auto nvrtc_option = fmt::format("-DLC_NVRTC_VERSION={}", nvrtc_version); - auto real_option = fmt::format("-DOCARINA_CUDA_USE_FLOAT32={}", use_float32); - std::vector header_names{"cuda_device_std.h", "cuda_device_scalar.h", "cuda_device_vector.h", - "cuda_device_matrix.h", "cuda_device_builtin.h", "cuda_tensor_builtin.h", "cuda_vector_func.h", - "cuda_device_math.h", "cuda_matrix_func.h", "cuda_device_resource.h"}; + switch_profile::Scope profile{"CUDACompiler::compile", "compiler_setup"}; + // These strings own the storage referenced by NVRTC's pointer arrays. + // Keep them alive after the options phase ends, through the compile call. + string header_path, nvrtc_option, real_option, compute_sm, rt_option, const_option; + std::vector header_names; std::vector header_sources; std::vector header_sources_ptr; - - for (auto fn : header_names) { - string source = RHIContext::read_file(string("cuda/") + fn); - header_sources_ptr.push_back(source.c_str()); - header_sources.push_back(ocarina::move(source)); - } - - // NVRTC will embed the chosen architecture into the PTX (e.g. .target sm_120). - // OptiX's PTX frontend can lag behind the newest GPU targets; if we hand it PTX - // with an unsupported .target, optixModuleCreate() fails with "Invalid PTX input". - // - // PTX is ABI-compatible by the driver for the real GPU anyway, so we can OptiX modules - // It's safer to cap the PTX target to a widely supported SM. - int ptx_sm = sm; - if (function.is_raytracing()) { - // keep this conservative(for compatibility) given OptiX releases. - ptx_sm = std::min(ptx_sm, 89); - } - auto compute_sm = ocarina::format("-arch=compute_{}", ptx_sm); - auto rt_option = fmt::format("-DLC_OPTIX_VERSION={}", OPTIX_VERSION); - auto const_option = fmt::format("-Dlc_constant={}", nvrtc_version <= 110200 ? "const" : "constexpr"); - ocarina::vector compile_option = { - "--std=c++17", - compute_sm.c_str(), - const_option.c_str(), - rt_option.c_str(), - nvrtc_option.c_str(), - real_option.c_str(), - "-default-device", - "--use_fast_math", - "-restrict", - header_path.c_str(), - //#ifndef NDEBUG - "-lineinfo", - //#endif - "-extra-device-vectorization", - "-dw", - "-w"}; + ocarina::vector compile_option; ocarina::list includes; - if (function.is_raytracing()) { - static string inc_path = (fs::current_path() / "cuda" / "optix").string(); - includes.push_back(ocarina::format("-I {}", inc_path)); - compile_option.push_back(includes.back().c_str()); - header_names.push_back("optix_device_header.h"); - string source = RHIContext::read_file(string("cuda/optix_device_header.h")); - header_sources_ptr.push_back(source.c_str()); - header_sources.push_back(ocarina::move(source)); + int ptx_sm = sm; + { + switch_profile::Scope options_profile{"CUDACompiler::compile.headers_options", "compiler_setup"}; + fs::path cuda_path = get_cuda_path(); + cuda_path = cuda_path / "include"; + + header_path = "-I" + cuda_path.string(); + int ver_major = 0; + int ver_minor = 0; + OC_NVRTC_CHECK(nvrtcVersion(&ver_major, &ver_minor)); + int nvrtc_version = ver_major * 10000 + ver_minor * 100; + const auto storage_policy = function.storage_policy(); + const int use_float32 = storage_policy.policy == PrecisionPolicy::force_f16 ? 0 : 1; + nvrtc_option = fmt::format("-DLC_NVRTC_VERSION={}", nvrtc_version); + real_option = fmt::format("-DOCARINA_CUDA_USE_FLOAT32={}", use_float32); + header_names = {"cuda_device_std.h", "cuda_device_scalar.h", "cuda_device_vector.h", + "cuda_device_matrix.h", "cuda_device_builtin.h", "cuda_tensor_builtin.h", "cuda_vector_func.h", + "cuda_device_math.h", "cuda_matrix_func.h", "cuda_device_resource.h"}; + + for (auto fn : header_names) { + string source = RHIContext::read_file(string("cuda/") + fn); + header_sources_ptr.push_back(source.c_str()); + header_sources.push_back(ocarina::move(source)); + } + + // NVRTC will embed the chosen architecture into the PTX (e.g. .target sm_120). + // OptiX's PTX frontend can lag behind the newest GPU targets; if we hand it PTX + // with an unsupported .target, optixModuleCreate() fails with "Invalid PTX input". + // + // PTX is ABI-compatible by the driver for the real GPU anyway, so we can OptiX modules + // It's safer to cap the PTX target to a widely supported SM. + if (function.is_raytracing()) { + // keep this conservative(for compatibility) given OptiX releases. + ptx_sm = std::min(ptx_sm, 89); + } + compute_sm = ocarina::format("-arch=compute_{}", ptx_sm); + rt_option = fmt::format("-DLC_OPTIX_VERSION={}", OPTIX_VERSION); + const_option = fmt::format("-Dlc_constant={}", nvrtc_version <= 110200 ? "const" : "constexpr"); + compile_option = { + "--std=c++17", + compute_sm.c_str(), + const_option.c_str(), + rt_option.c_str(), + nvrtc_option.c_str(), + real_option.c_str(), + "-default-device", + "--use_fast_math", + "-restrict", + header_path.c_str(), + //#ifndef NDEBUG + "-lineinfo", + //#endif + "-extra-device-vectorization", + "-dw", + "-w"}; + if (function.is_raytracing()) { + static string inc_path = (fs::current_path() / "cuda" / "optix").string(); + includes.push_back(ocarina::format("-I {}", inc_path)); + compile_option.push_back(includes.back().c_str()); + header_names.push_back("optix_device_header.h"); + string source = RHIContext::read_file(string("cuda/optix_device_header.h")); + header_sources_ptr.push_back(source.c_str()); + header_sources.push_back(ocarina::move(source)); + } + for (const auto &header_name : header_names) { + includes.push_back(ocarina::format("-include={}", header_name)); + compile_option.push_back(includes.back().c_str()); + } } - for (const auto &header_name : header_names) { - includes.push_back(ocarina::format("-include={}", header_name)); - compile_option.push_back(includes.back().c_str()); + string fn; + { + switch_profile::Scope hash_profile{"CUDACompiler::compile.cache_key", "cache_key"}; + uint64_t ext_hash = hash64(hash64_list(compile_option), + hash64_list(header_sources_ptr), + ast_to_cuda_source_cache_version, + shader_codegen_path_version(Env::shader_codegen_path())); + fn = function.func_name(ext_hash, function.description()); } - uint64_t ext_hash = hash64(hash64_list(compile_option), - hash64_list(header_sources_ptr), - ast_to_cuda_source_cache_version, - shader_codegen_path_version(Env::shader_codegen_path())); - auto compile = [&](const string &cu, const string &fn, int sm) -> string { - TIMER_TAG(compile, "compile " + fn); + OC_INFO_FORMAT("compiling shader with NVRTC: {}", fn); + TIMER_TAG(compile, "NVRTC compile " + fn); + switch_profile::Scope nvrtc_profile{"CUDACompiler::compile.nvrtc", "nvrtc_setup"}; nvrtcProgram program{}; OC_NVRTC_CHECK(nvrtcCreateProgram(&program, cu.c_str(), fn.c_str(), header_names.size(), header_sources_ptr.data(), header_names.data())); - const nvrtcResult compile_res = nvrtcCompileProgram(program, compile_option.size(), compile_option.data()); + nvrtcResult compile_res; + { + switch_profile::Scope compile_profile{"nvrtcCompileProgram", "nvrtc"}; + compile_res = nvrtcCompileProgram(program, compile_option.size(), compile_option.data()); + } size_t log_size = 0; OC_NVRTC_CHECK(nvrtcGetProgramLogSize(program, &log_size)); string log; @@ -169,14 +189,12 @@ ocarina::string CUDACompiler::compile(const Function &function, int sm) const no return ptx; }; - string fn = function.func_name(ext_hash, function.description()); - ocarina::string ptx_fn = fn + ".ptx"; string cu_fn = fn + ".cu"; ocarina::string ptx; RHIContext *context = device_->context(); if (!context->is_exist_cache(ptx_fn)) { - OC_INFO_FORMAT("miss ptx file {}", ptx_fn); + OC_INFO_FORMAT("shader PTX cache miss: {}", ptx_fn); if (!context->is_exist_cache(cu_fn)) { const ocarina::string cu = emit_cuda_source(function); context->write_global_cache(cu_fn, cu); @@ -188,7 +206,7 @@ ocarina::string CUDACompiler::compile(const Function &function, int sm) const no context->write_global_cache(ptx_fn, ptx); } } else { - OC_INFO_FORMAT("find ptx file {}", ptx_fn); + OC_INFO_FORMAT("shader PTX cache hit: {}", ptx_fn); ptx = context->read_global_cache(ptx_fn); } return ptx; diff --git a/modules/ocarina/backends/cuda/cuda_device.cpp b/modules/ocarina/backends/cuda/cuda_device.cpp index 15c705bf..9fb76c9a 100644 --- a/modules/ocarina/backends/cuda/cuda_device.cpp +++ b/modules/ocarina/backends/cuda/cuda_device.cpp @@ -3,6 +3,7 @@ // #include "cuda_device.h" +#include "core/util/switch_profile.h" #include "cuda_stream.h" #include "cuda_texture.h" #include "cuda_shader.h" @@ -239,6 +240,7 @@ handle_ty CUDADevice::create_texture_from_external(ocarina::uint tex_handle) noe } handle_ty CUDADevice::create_shader(const Function &function) noexcept { + switch_profile::Scope profile{"CUDADevice::create_shader", "backend"}; CUDACompiler compiler(this); ocarina::string ptx = compiler.compile(function, compute_capability_); @@ -502,4 +504,4 @@ OC_EXPORT_API ocarina::CUDADevice *create_device(ocarina::RHIContext *context) { OC_EXPORT_API void destroy(ocarina::CUDADevice *device) { ocarina::delete_with_allocator(device); OC_INFO("cuda device is destroy!"); -} \ No newline at end of file +} diff --git a/modules/ocarina/backends/cuda/cuda_shader.cpp b/modules/ocarina/backends/cuda/cuda_shader.cpp index 3c24ff9f..cfd8f8b6 100644 --- a/modules/ocarina/backends/cuda/cuda_shader.cpp +++ b/modules/ocarina/backends/cuda/cuda_shader.cpp @@ -3,6 +3,7 @@ // #include "cuda_shader.h" +#include "core/util/switch_profile.h" #include "util.h" #include "cuda_device.h" #include @@ -153,8 +154,18 @@ class CUDASimpleShader : public CUDAShader { CUDASimpleShader(Device::Impl *device, const ocarina::string &ptx, const Function &f) : CUDAShader(device, f) { - OC_CU_CHECK(cuModuleLoadData(&module_, ptx.c_str())); - OC_CU_CHECK(cuModuleGetFunction(&func_handle_, module_, function_.func_name().c_str())); + CUresult load_result; + { + switch_profile::Scope profile{"cuModuleLoadData", "cuda_module"}; + load_result = cuModuleLoadData(&module_, ptx.c_str()); + } + OC_CU_CHECK(load_result); + CUresult function_result; + { + switch_profile::Scope profile{"cuModuleGetFunction", "cuda_module"}; + function_result = cuModuleGetFunction(&func_handle_, module_, function_.func_name().c_str()); + } + OC_CU_CHECK(function_result); } ~CUDASimpleShader() override { OC_CU_CHECK(cuModuleUnload(module_)); @@ -227,6 +238,7 @@ class OptixShader : public CUDAShader { public: void init_module(const string_view &ptx_code) { + switch_profile::Scope profile{"OptixShader::init_module", "optix_module"}; OptixModuleCompileOptions module_compile_options = {}; static constexpr std::array ray_trace_payload_semantics{ OPTIX_PAYLOAD_SEMANTICS_TRACE_CALLER_WRITE | OPTIX_PAYLOAD_SEMANTICS_CH_READ, @@ -271,17 +283,21 @@ class OptixShader : public CUDAShader { pipeline_compile_options_.pipelineLaunchParamsVariableName = "params"; char log[2048]; size_t log_size = sizeof(log); - OC_OPTIX_CHECK_WITH_LOG(optixModuleCreate( - device_->optix_device_context(), - &module_compile_options, - &pipeline_compile_options_, - ptx_code.data(), ptx_code.size(), - log, &log_size, &optix_module_), - log); + OptixResult module_result; + { + switch_profile::Scope api_profile{"optixModuleCreate", "optix_module"}; + module_result = optixModuleCreate( + device_->optix_device_context(), + &module_compile_options, + &pipeline_compile_options_, + ptx_code.data(), ptx_code.size(), + log, &log_size, &optix_module_); + } + OC_OPTIX_CHECK_WITH_LOG(module_result, log); } void build_pipeline(OptixDeviceContext optix_device_context) noexcept { - + switch_profile::Scope profile{"OptixShader::build_pipeline", "optix_pipeline"}; constexpr int max_trace_depth = 1; OptixPipelineLinkOptions pipeline_link_options = {}; @@ -290,15 +306,19 @@ class OptixShader : public CUDAShader { char log[2048]; size_t sizeof_log = sizeof(log); - OC_OPTIX_CHECK_WITH_LOG(optixPipelineCreate( - optix_device_context, - &pipeline_compile_options_, - &pipeline_link_options, - (OptixProgramGroup *)&program_group_table_, - program_group_table_.size(), - log, &sizeof_log, - &optix_pipeline_), - log); + OptixResult pipeline_result; + { + switch_profile::Scope api_profile{"optixPipelineCreate", "optix_pipeline"}; + pipeline_result = optixPipelineCreate( + optix_device_context, + &pipeline_compile_options_, + &pipeline_link_options, + (OptixProgramGroup *)&program_group_table_, + program_group_table_.size(), + log, &sizeof_log, + &optix_pipeline_); + } + OC_OPTIX_CHECK_WITH_LOG(pipeline_result, log); // Set shaders stack sizes. OptixStackSizes stack_sizes = {}; @@ -326,6 +346,7 @@ class OptixShader : public CUDAShader { } void build_sbt(ProgramGroupTable program_group_table) { + switch_profile::Scope profile{"OptixShader::build_sbt", "optix_sbt"}; SBTRecord sbt[ProgramGroupTable::size()] = {}; sbt_records_ = Buffer(device_, ProgramGroupTable::size(), "OptixShader::sbt_records_"); OC_OPTIX_CHECK(optixSbtRecordPackHeader(program_group_table.raygen_group, &sbt[0])); @@ -345,6 +366,7 @@ class OptixShader : public CUDAShader { ProgramGroupTable create_program_groups(OptixDeviceContext optix_device_context, const ProgramName &program_name) { + switch_profile::Scope profile{"OptixShader::create_program_groups", "optix_program_group"}; OptixProgramGroupOptions program_group_options = {}; // program_group_options.payloadType = payload_types_.data(); char log[2048]; @@ -355,15 +377,16 @@ class OptixShader : public CUDAShader { raygen_prog_group_desc.kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; raygen_prog_group_desc.raygen.module = optix_module_; raygen_prog_group_desc.raygen.entryFunctionName = program_name.raygen; - OC_OPTIX_CHECK_WITH_LOG(optixProgramGroupCreate( - optix_device_context, - &raygen_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.raygen_group)), - log); + OptixResult raygen_result; + { + switch_profile::Scope api_profile{"optixProgramGroupCreate.raygen", "optix_program_group"}; + raygen_result = optixProgramGroupCreate( + optix_device_context, &raygen_prog_group_desc, + 1,// num program groups + &program_group_options, log, &sizeof_log, + &(program_group_table.raygen_group)); + } + OC_OPTIX_CHECK_WITH_LOG(raygen_result, log); } { OptixProgramGroupDesc hit_prog_group_desc = {}; @@ -372,15 +395,16 @@ class OptixShader : public CUDAShader { hit_prog_group_desc.hitgroup.entryFunctionNameCH = program_name.closesthit_closest; sizeof_log = sizeof(log); - OC_OPTIX_CHECK_WITH_LOG(optixProgramGroupCreate( - optix_device_context, - &hit_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.hit_closest_group)), - log); + OptixResult closest_result; + { + switch_profile::Scope api_profile{"optixProgramGroupCreate.closesthit", "optix_program_group"}; + closest_result = optixProgramGroupCreate( + optix_device_context, &hit_prog_group_desc, + 1,// num program groups + &program_group_options, log, &sizeof_log, + &(program_group_table.hit_closest_group)); + } + OC_OPTIX_CHECK_WITH_LOG(closest_result, log); memset(&hit_prog_group_desc, 0, sizeof(OptixProgramGroupDesc)); hit_prog_group_desc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; @@ -388,29 +412,31 @@ class OptixShader : public CUDAShader { hit_prog_group_desc.hitgroup.entryFunctionNameCH = program_name.closesthit_occlusion; sizeof_log = sizeof(log); - OC_OPTIX_CHECK_WITH_LOG(optixProgramGroupCreate( - optix_device_context, - &hit_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.hit_any_group)), - log); + OptixResult anyhit_result; + { + switch_profile::Scope api_profile{"optixProgramGroupCreate.anyhit", "optix_program_group"}; + anyhit_result = optixProgramGroupCreate( + optix_device_context, &hit_prog_group_desc, + 1,// num program groups + &program_group_options, log, &sizeof_log, + &(program_group_table.hit_any_group)); + } + OC_OPTIX_CHECK_WITH_LOG(anyhit_result, log); } { OptixProgramGroupDesc miss_prog_group_desc = {}; miss_prog_group_desc.kind = OPTIX_PROGRAM_GROUP_KIND_MISS; sizeof_log = sizeof(log); - OC_OPTIX_CHECK_WITH_LOG(optixProgramGroupCreate( - optix_device_context, - &miss_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.miss_closest_group)), - log); + OptixResult miss_result; + { + switch_profile::Scope api_profile{"optixProgramGroupCreate.miss", "optix_program_group"}; + miss_result = optixProgramGroupCreate( + optix_device_context, &miss_prog_group_desc, + 1,// num program groups + &program_group_options, log, &sizeof_log, + &(program_group_table.miss_closest_group)); + } + OC_OPTIX_CHECK_WITH_LOG(miss_result, log); } return program_group_table; @@ -476,6 +502,7 @@ class OptixShader : public CUDAShader { }; CUDAShader *CUDAShader::create(Device::Impl *device, const string &ptx, const Function &f) { + switch_profile::Scope profile{"CUDAShader::create", "backend"}; if (f.is_raytracing()) { return ocarina::new_with_allocator(device, ptx, f); } else { @@ -483,4 +510,4 @@ CUDAShader *CUDAShader::create(Device::Impl *device, const string &ptx, const Fu } } -}// namespace ocarina \ No newline at end of file +}// namespace ocarina diff --git a/modules/ocarina/core/util/switch_profile.h b/modules/ocarina/core/util/switch_profile.h new file mode 100644 index 00000000..0ba90078 --- /dev/null +++ b/modules/ocarina/core/util/switch_profile.h @@ -0,0 +1,59 @@ +#pragma once + +#include +#include +#include +#include +#include +#include + +// Opt-in, flushed wall-clock events shared by the renderer and its backend. +namespace ocarina::switch_profile { +inline bool enabled() noexcept { + static const bool value = [] { + const char* env = std::getenv("CORONA_SWITCH_PROFILE"); + return env && std::string(env) == "1"; + }(); + return value; +} +inline std::string quoted(const char* value) { + std::string result = "\""; + for (const unsigned char c : std::string(value)) { + if (c == '"' || c == '\\') result += '\\'; + if (c >= 32) result += static_cast(c); + else result += '?'; + } + return result + '"'; +} +// Keep a named instance in the measured function or block. Destruction closes +// the event on normal return, early return and exception unwinding. +class [[nodiscard]] Scope { + const char* name_; + const char* category_; + const bool active_; + void emit(char phase) const noexcept { + try { + const auto time = std::chrono::duration_cast( + std::chrono::steady_clock::now().time_since_epoch()).count(); + const auto thread = std::hash{}(std::this_thread::get_id()); + const auto line = std::string("CORONA_PROFILE {\"ph\":\"") + phase + + "\",\"name\":" + quoted(name_) + ",\"cat\":" + quoted(category_) + + ",\"ts\":" + std::to_string(time) + ",\"tid\":" + std::to_string(thread) + "}\n"; + std::fwrite(line.data(), 1, line.size(), stdout); + std::fflush(stdout); + } catch (...) { + // Profiling must not change the renderer's error behavior. + } + } +public: + Scope(const char* name, const char* category) noexcept + : name_(name), category_(category), active_(enabled()) { + if (active_) emit('B'); + } + ~Scope() noexcept { + if (active_) emit('E'); + } + Scope(const Scope&) = delete; + Scope& operator=(const Scope&) = delete; +}; +} // namespace ocarina::switch_profile diff --git a/modules/ocarina/rhi/context.cpp b/modules/ocarina/rhi/context.cpp index 25723e94..f31e9362 100644 --- a/modules/ocarina/rhi/context.cpp +++ b/modules/ocarina/rhi/context.cpp @@ -3,6 +3,7 @@ // #include "context.h" +#include "core/util/switch_profile.h" #include "core/runtime/dynamic_module.h" #include "rhi/device.h" #include "core/runtime/platform.h" @@ -85,6 +86,7 @@ bool RHIContext::create_directory_if_necessary(const fs::path &path) { } string RHIContext::read_file(const fs::path &fn) { + switch_profile::Scope profile{"RHIContext::read_file", "file_io"}; std::ifstream fst; fst.open(fn.c_str()); std::stringstream buffer; @@ -100,10 +102,12 @@ void RHIContext::write_file(const fs::path &fn, const std::string &text) { } void RHIContext::write_global_cache(const string &fn, const string &text) const noexcept { + switch_profile::Scope profile{"RHIContext::write_global_cache", "cache_io"}; write_file(cache_directory() / fn, text); } string RHIContext::read_global_cache(const string &fn) const noexcept { + switch_profile::Scope profile{"RHIContext::read_global_cache", "cache_io"}; return read_file(cache_directory() / fn); } @@ -114,6 +118,7 @@ void RHIContext::clear_cache() const noexcept { } bool RHIContext::is_exist_cache(const string &fn) const noexcept { + switch_profile::Scope profile{"RHIContext::is_exist_cache", "cache_io"}; if (!impl_->use_cache) { return false; } @@ -168,4 +173,4 @@ Device RHIContext::create_device(const string &backend_name) noexcept { return Device{Device::Handle{create_device(this), destroy_func}}; } -}// namespace ocarina \ No newline at end of file +}// namespace ocarina diff --git a/modules/ocarina/rhi/device.h b/modules/ocarina/rhi/device.h index 506e9ed2..f1d2f0b3 100644 --- a/modules/ocarina/rhi/device.h +++ b/modules/ocarina/rhi/device.h @@ -231,7 +231,7 @@ class OC_RHI_API Device : public concepts::Noncopyable { [[nodiscard]] Texture3D create_texture(Image *image_resource, const TextureViewCreation &texture_view, const string &desc = "") const noexcept; template [[nodiscard]] auto compile(const Kernel &kernel, const string &shader_desc = "", ShaderTag tag = CS) const noexcept { - OC_INFO_FORMAT("compile shader : {}", shader_desc.c_str()); + OC_DEBUG_FORMAT("prepare shader: {}", shader_desc.c_str()); kernel.function()->set_description(shader_desc); return create>(kernel.function(), tag); } @@ -310,4 +310,4 @@ class OC_RHI_API Device : public concepts::Noncopyable { DynamicBufferLayout layout = DynamicBufferLayout::AOS) const noexcept; }; -}// namespace ocarina \ No newline at end of file +}// namespace ocarina