From 075f97f89fea7ef011a39f58da624a0b5923dda6 Mon Sep 17 00:00:00 2001
From: Jacob Hinkle <1454944+jacobhinkle@users.noreply.github.com>
Date: Wed, 29 Jan 2025 12:30:42 -0500
Subject: [PATCH] Fix names of printed ptx and cubin files (#3788)

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}`.
---
 csrc/runtime/compiled_kernel.cpp | 13 ++++---------
 1 file changed, 4 insertions(+), 9 deletions(-)

diff --git a/csrc/runtime/compiled_kernel.cpp b/csrc/runtime/compiled_kernel.cpp
index fe3464bd551..cea3f9e7ef0 100644
--- a/csrc/runtime/compiled_kernel.cpp
+++ b/csrc/runtime/compiled_kernel.cpp
@@ -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;
@@ -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;
@@ -716,7 +715,7 @@ 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");
     }
   }
 
@@ -724,7 +723,7 @@ std::unique_ptr<executor_utils::CudaExecutable> compileSource(
     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");
     }
   }
 
@@ -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(