Skip to content

Commit

Permalink
Fix names of printed ptx and cubin files
Browse files Browse the repository at this point in the history
This uses the function name as the filename. Since #3778 we have been
generating filenames like `__tmp_.ptx` since the `kernel_name` is not
yet filled when cubin and ptx files are output. Now we will get
`__tmp_nvfuser_none_f0_c0_r0_g0.{cu,cubin,ptx}` or if
`NVFUSER_ENABLE=static_fusion_count` is provided
`__tmp_nvfuser_1.{cu,cubin,ptx}`.
  • Loading branch information
jacobhinkle committed Jan 29, 2025
1 parent a470821 commit ec56517
Showing 1 changed file with 4 additions and 9 deletions.
13 changes: 4 additions & 9 deletions csrc/runtime/compiled_kernel.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -689,7 +689,6 @@ std::vector<char> compileNvrtcProgramToPtx(const nvrtcProgram& program) {
std::unique_ptr<executor_utils::CudaExecutable> compileSource(
const std::string& full_src_code,
const std::string& func_name,
const std::string& kernel_name,
const bool compile_to_sass,
NvrtcCompileDriver& nvrtc_compile) {
std::stringstream log;
Expand All @@ -700,7 +699,7 @@ std::unique_ptr<executor_utils::CudaExecutable> compileSource(
NVFUSER_NVRTC_SAFE_CALL(nvrtcDestroyProgram(&program));
});

createNvrtcProgram(program, kernel_name, full_src_code);
createNvrtcProgram(program, func_name, full_src_code);

NVFUSER_NVRTC_SAFE_CALL(nvrtcAddNameExpression(program, func_name.c_str()));
log << nvrtc_compile.invoke(program, full_src_code) << std::endl;
Expand All @@ -716,15 +715,15 @@ std::unique_ptr<executor_utils::CudaExecutable> compileSource(
compiled_kernel->cubin = compileNvrtcProgramToCubin(program);
if (isDebugDumpEnabled(DebugDumpOption::Cubin)) {
compiled_kernel->cubin_filename =
dumpCompiledCodeToFile(compiled_kernel->cubin, kernel_name, ".cubin");
dumpCompiledCodeToFile(compiled_kernel->cubin, func_name, ".cubin");
}
}

if (!compile_to_sass || isDebugDumpEnabled(DebugDumpOption::Ptx)) {
compiled_kernel->ptx = compileNvrtcProgramToPtx(program);
if (isDebugDumpEnabled(DebugDumpOption::Ptx)) {
compiled_kernel->ptx_filename =
dumpCompiledCodeToFile(compiled_kernel->ptx, kernel_name, ".ptx");
dumpCompiledCodeToFile(compiled_kernel->ptx, func_name, ".ptx");
}
}

Expand Down Expand Up @@ -810,11 +809,7 @@ std::unique_ptr<executor_utils::CudaExecutable> getCudaExecutable(
(compile_to_sass ? compiled_kernel->cubin
: compiled_kernel->ptx)))) {
compiled_kernel = compileSource(
full_src_code,
func_name,
compiled_kernel->kernel_name,
compile_to_sass,
nvrtc_compile_driver);
full_src_code, func_name, compile_to_sass, nvrtc_compile_driver);
log << compiled_kernel->compile_log << std::endl;
if (use_kernel_db) {
auto result = kernel_db.write(
Expand Down

0 comments on commit ec56517

Please sign in to comment.