From 4cf1758f805a1c945aa7603c679dbf66145d22db Mon Sep 17 00:00:00 2001 From: GraphZLL <249002031@qq.com> Date: Mon, 5 Oct 2026 11:00:21 +0800 Subject: [PATCH 1/2] =?UTF-8?q?perf(ocarina):=20=E5=A2=9E=E5=8A=A0?= =?UTF-8?q?=E7=9D=80=E8=89=B2=E5=99=A8=E5=88=9B=E5=BB=BA=E4=B8=8E=E7=BC=93?= =?UTF-8?q?=E5=AD=98=E8=80=97=E6=97=B6=E5=88=86=E6=9E=90?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- modules/ocarina/ast/function.cpp | 4 +- .../ocarina/backends/cuda/cuda_compiler.cpp | 20 ++++-- modules/ocarina/backends/cuda/cuda_device.cpp | 4 +- modules/ocarina/backends/cuda/cuda_shader.cpp | 47 +++++++++----- modules/ocarina/core/util/switch_profile.h | 64 +++++++++++++++++++ modules/ocarina/rhi/context.cpp | 7 +- modules/ocarina/rhi/device.h | 4 +- 7 files changed, 124 insertions(+), 26 deletions(-) create mode 100644 modules/ocarina/core/util/switch_profile.h 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/backends/cuda/cuda_compiler.cpp b/modules/ocarina/backends/cuda/cuda_compiler.cpp index 270e9a71..6b037376 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,7 +68,8 @@ CUDACompiler::CUDACompiler(CUDADevice *device) : device_(device) {} ocarina::string CUDACompiler::compile(const Function &function, int sm) const noexcept { - + switch_profile::Scope profile{"CUDACompiler::compile", "compiler_setup"}; + switch_profile::Scope options_profile{"CUDACompiler::compile.headers_options", "compiler_setup"}; fs::path cuda_path = get_cuda_path(); cuda_path = cuda_path / "include"; @@ -137,18 +140,24 @@ ocarina::string CUDACompiler::compile(const Function &function, int sm) const no compile_option.push_back(includes.back().c_str()); } + options_profile.finish(); + 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())); 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()); + const nvrtcResult compile_res = switch_profile::measure("nvrtcCompileProgram", "nvrtc", [&] { + return nvrtcCompileProgram(program, compile_option.size(), compile_option.data()); + }); size_t log_size = 0; OC_NVRTC_CHECK(nvrtcGetProgramLogSize(program, &log_size)); string log; @@ -170,13 +179,14 @@ ocarina::string CUDACompiler::compile(const Function &function, int sm) const no }; string fn = function.func_name(ext_hash, function.description()); + hash_profile.finish(); 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 +198,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..7f79cc94 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,14 @@ 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())); + const auto load_result = switch_profile::measure("cuModuleLoadData", "cuda_module", [&] { + return cuModuleLoadData(&module_, ptx.c_str()); + }); + OC_CU_CHECK(load_result); + const auto function_result = switch_profile::measure("cuModuleGetFunction", "cuda_module", [&] { + return cuModuleGetFunction(&func_handle_, module_, function_.func_name().c_str()); + }); + OC_CU_CHECK(function_result); } ~CUDASimpleShader() override { OC_CU_CHECK(cuModuleUnload(module_)); @@ -227,6 +234,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,6 +279,7 @@ class OptixShader : public CUDAShader { pipeline_compile_options_.pipelineLaunchParamsVariableName = "params"; char log[2048]; size_t log_size = sizeof(log); + switch_profile::Scope api_profile{"optixModuleCreate", "optix_module"}; OC_OPTIX_CHECK_WITH_LOG(optixModuleCreate( device_->optix_device_context(), &module_compile_options, @@ -281,7 +290,7 @@ class OptixShader : public CUDAShader { } 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,6 +299,8 @@ class OptixShader : public CUDAShader { char log[2048]; size_t sizeof_log = sizeof(log); + { + switch_profile::Scope api_profile{"optixPipelineCreate", "optix_pipeline"}; OC_OPTIX_CHECK_WITH_LOG(optixPipelineCreate( optix_device_context, &pipeline_compile_options_, @@ -299,6 +310,7 @@ class OptixShader : public CUDAShader { log, &sizeof_log, &optix_pipeline_), log); + } // Set shaders stack sizes. OptixStackSizes stack_sizes = {}; @@ -326,6 +338,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 +358,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 +369,15 @@ 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( + const auto raygen_result = switch_profile::measure("optixProgramGroupCreate.raygen", "optix_program_group", [&] { return optixProgramGroupCreate( optix_device_context, &raygen_prog_group_desc, 1,// num program groups &program_group_options, log, &sizeof_log, - &(program_group_table.raygen_group)), - log); + &(program_group_table.raygen_group)); }); + OC_OPTIX_CHECK_WITH_LOG(raygen_result, log); } { OptixProgramGroupDesc hit_prog_group_desc = {}; @@ -372,15 +386,15 @@ class OptixShader : public CUDAShader { hit_prog_group_desc.hitgroup.entryFunctionNameCH = program_name.closesthit_closest; sizeof_log = sizeof(log); - OC_OPTIX_CHECK_WITH_LOG(optixProgramGroupCreate( + const auto closest_result = switch_profile::measure("optixProgramGroupCreate.closesthit", "optix_program_group", [&] { return 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); + &(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 +402,29 @@ class OptixShader : public CUDAShader { hit_prog_group_desc.hitgroup.entryFunctionNameCH = program_name.closesthit_occlusion; sizeof_log = sizeof(log); - OC_OPTIX_CHECK_WITH_LOG(optixProgramGroupCreate( + const auto anyhit_result = switch_profile::measure("optixProgramGroupCreate.anyhit", "optix_program_group", [&] { return 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); + &(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( + const auto miss_result = switch_profile::measure("optixProgramGroupCreate.miss", "optix_program_group", [&] { return 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); + &(program_group_table.miss_closest_group)); }); + OC_OPTIX_CHECK_WITH_LOG(miss_result, log); } return program_group_table; @@ -476,6 +490,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 +498,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..dea318ac --- /dev/null +++ b/modules/ocarina/core/util/switch_profile.h @@ -0,0 +1,64 @@ +#pragma once + +#include +#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 + '"'; +} +class Scope { + const char* name_; + const char* category_; + 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'); + } + void finish() noexcept { + if (active_) { emit('E'); active_ = false; } + } + ~Scope() { finish(); } + Scope(const Scope&) = delete; + Scope& operator=(const Scope&) = delete; +}; +template +decltype(auto) measure(const char* name, const char* category, F&& action) { + Scope scope{name, category}; + return std::forward(action)(); +} +} // 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 From e3edffc6d2803c33821e6d43f5badc062dd718cd Mon Sep 17 00:00:00 2001 From: GraphZLL <249002031@qq.com> Date: Mon, 5 Oct 2026 12:17:31 +0800 Subject: [PATCH 2/2] =?UTF-8?q?refactor(ocarina):=20=E5=B0=86=E5=88=87?= =?UTF-8?q?=E6=8D=A2=E8=80=97=E6=97=B6=E5=88=86=E6=9E=90=E6=94=B9=E4=B8=BA?= =?UTF-8?q?=20RAII=20=E4=BD=9C=E7=94=A8=E5=9F=9F?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- modules/ocarina/ast/function.h | 4 +- .../ocarina/backends/cuda/cuda_compiler.cpp | 164 +++++++++--------- modules/ocarina/backends/cuda/cuda_shader.cpp | 124 +++++++------ modules/ocarina/core/util/switch_profile.h | 17 +- 4 files changed, 163 insertions(+), 146 deletions(-) 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 6b037376..b05e1af4 100644 --- a/modules/ocarina/backends/cuda/cuda_compiler.cpp +++ b/modules/ocarina/backends/cuda/cuda_compiler.cpp @@ -69,84 +69,93 @@ CUDACompiler::CUDACompiler(CUDADevice *device) ocarina::string CUDACompiler::compile(const Function &function, int sm) const noexcept { switch_profile::Scope profile{"CUDACompiler::compile", "compiler_setup"}; - switch_profile::Scope options_profile{"CUDACompiler::compile.headers_options", "compiler_setup"}; - 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"}; + // 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()); } - options_profile.finish(); - 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())); - auto compile = [&](const string &cu, const string &fn, int sm) -> string { OC_INFO_FORMAT("compiling shader with NVRTC: {}", fn); TIMER_TAG(compile, "NVRTC compile " + fn); @@ -155,9 +164,11 @@ ocarina::string CUDACompiler::compile(const Function &function, int sm) const no 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 = switch_profile::measure("nvrtcCompileProgram", "nvrtc", [&] { - return 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; @@ -178,9 +189,6 @@ ocarina::string CUDACompiler::compile(const Function &function, int sm) const no return ptx; }; - string fn = function.func_name(ext_hash, function.description()); - hash_profile.finish(); - ocarina::string ptx_fn = fn + ".ptx"; string cu_fn = fn + ".cu"; ocarina::string ptx; diff --git a/modules/ocarina/backends/cuda/cuda_shader.cpp b/modules/ocarina/backends/cuda/cuda_shader.cpp index 7f79cc94..cfd8f8b6 100644 --- a/modules/ocarina/backends/cuda/cuda_shader.cpp +++ b/modules/ocarina/backends/cuda/cuda_shader.cpp @@ -154,13 +154,17 @@ class CUDASimpleShader : public CUDAShader { CUDASimpleShader(Device::Impl *device, const ocarina::string &ptx, const Function &f) : CUDAShader(device, f) { - const auto load_result = switch_profile::measure("cuModuleLoadData", "cuda_module", [&] { - return cuModuleLoadData(&module_, ptx.c_str()); - }); + CUresult load_result; + { + switch_profile::Scope profile{"cuModuleLoadData", "cuda_module"}; + load_result = cuModuleLoadData(&module_, ptx.c_str()); + } OC_CU_CHECK(load_result); - const auto function_result = switch_profile::measure("cuModuleGetFunction", "cuda_module", [&] { - return cuModuleGetFunction(&func_handle_, module_, function_.func_name().c_str()); - }); + 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 { @@ -279,14 +283,17 @@ class OptixShader : public CUDAShader { pipeline_compile_options_.pipelineLaunchParamsVariableName = "params"; char log[2048]; size_t log_size = sizeof(log); - switch_profile::Scope api_profile{"optixModuleCreate", "optix_module"}; - 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 { @@ -299,18 +306,19 @@ class OptixShader : public CUDAShader { char log[2048]; size_t sizeof_log = sizeof(log); + OptixResult pipeline_result; { - switch_profile::Scope api_profile{"optixPipelineCreate", "optix_pipeline"}; - 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); + 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 = {}; @@ -369,14 +377,15 @@ 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; - const auto raygen_result = switch_profile::measure("optixProgramGroupCreate.raygen", "optix_program_group", [&] { return optixProgramGroupCreate( - optix_device_context, - &raygen_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.raygen_group)); }); + 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); } { @@ -386,14 +395,15 @@ class OptixShader : public CUDAShader { hit_prog_group_desc.hitgroup.entryFunctionNameCH = program_name.closesthit_closest; sizeof_log = sizeof(log); - const auto closest_result = switch_profile::measure("optixProgramGroupCreate.closesthit", "optix_program_group", [&] { return optixProgramGroupCreate( - optix_device_context, - &hit_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.hit_closest_group)); }); + 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)); @@ -402,28 +412,30 @@ class OptixShader : public CUDAShader { hit_prog_group_desc.hitgroup.entryFunctionNameCH = program_name.closesthit_occlusion; sizeof_log = sizeof(log); - const auto anyhit_result = switch_profile::measure("optixProgramGroupCreate.anyhit", "optix_program_group", [&] { return optixProgramGroupCreate( - optix_device_context, - &hit_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.hit_any_group)); }); + 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); - const auto miss_result = switch_profile::measure("optixProgramGroupCreate.miss", "optix_program_group", [&] { return optixProgramGroupCreate( - optix_device_context, - &miss_prog_group_desc, - 1,// num program groups - &program_group_options, - log, - &sizeof_log, - &(program_group_table.miss_closest_group)); }); + 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); } diff --git a/modules/ocarina/core/util/switch_profile.h b/modules/ocarina/core/util/switch_profile.h index dea318ac..0ba90078 100644 --- a/modules/ocarina/core/util/switch_profile.h +++ b/modules/ocarina/core/util/switch_profile.h @@ -6,7 +6,6 @@ #include #include #include -#include // Opt-in, flushed wall-clock events shared by the renderer and its backend. namespace ocarina::switch_profile { @@ -26,10 +25,12 @@ inline std::string quoted(const char* value) { } return result + '"'; } -class Scope { +// 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_; - bool active_; + const bool active_; void emit(char phase) const noexcept { try { const auto time = std::chrono::duration_cast( @@ -49,16 +50,10 @@ class Scope { : name_(name), category_(category), active_(enabled()) { if (active_) emit('B'); } - void finish() noexcept { - if (active_) { emit('E'); active_ = false; } + ~Scope() noexcept { + if (active_) emit('E'); } - ~Scope() { finish(); } Scope(const Scope&) = delete; Scope& operator=(const Scope&) = delete; }; -template -decltype(auto) measure(const char* name, const char* category, F&& action) { - Scope scope{name, category}; - return std::forward(action)(); -} } // namespace ocarina::switch_profile