[SYCL] Fix debug info generation when integration footer is present

clang/include/clang/Basic/CodeGenOptions.def

@@ -485,6 +485,10 @@ CODEGENOPT(SkipRaxSetup, 1, 0)
 ENUM_CODEGENOPT(ZeroCallUsedRegs, llvm::ZeroCallUsedRegs::ZeroCallUsedRegsKind,
                 5, llvm::ZeroCallUsedRegs::ZeroCallUsedRegsKind::Skip)
 
+/// Whether to expect -main-file-name to be an absolute path to use it for
+/// checksum calculations or not.
+CODEGENOPT(SYCLUseMainFileName, 1, 0)
+
 /// Whether to use opaque pointers.
 CODEGENOPT(OpaquePointers, 1, 0)
 

clang/include/clang/Basic/CodeGenOptions.h

@@ -220,6 +220,9 @@ class CodeGenOptions : public CodeGenOptionsBase {
   /// file, for example with -save-temps.
   std::string MainFileName;
 
+  /// The user provided name for the "main file", with its full path.
+  std::string FullMainFileName;
+
   /// The name for the split debug info file used for the DW_AT_[GNU_]dwo_name
   /// attribute in the skeleton CU.
   std::string SplitDwarfFile;

clang/include/clang/Driver/Options.td

@@ -6151,6 +6151,10 @@ def main_file_name : Separate<["-"], "main-file-name">,
   HelpText<"Main file name to use for debug info and source if missing">,
   Flags<[CC1Option, CC1AsOption, NoDriverOption]>,
   MarshallingInfoString<CodeGenOpts<"MainFileName">>;
+def full_main_file_name : Separate<["-"], "full-main-file-name">,
+  HelpText<"File name with full path to use for debug info during host and device compile">,
+  Flags<[CC1Option, CC1AsOption, NoDriverOption]>,
+  MarshallingInfoString<CodeGenOpts<"FullMainFileName">>;
 def split_dwarf_output : Separate<["-"], "split-dwarf-output">,
   HelpText<"File name to use for split dwarf debug info output">,
   Flags<[CC1Option, CC1AsOption, NoDriverOption]>,
@@ -6510,6 +6514,10 @@ def fsycl_disable_range_rounding : Flag<["-"], "fsycl-disable-range-rounding">,
 def fsycl_enable_int_header_diags: Flag<["-"], "fsycl-enable-int-header-diags">,
   HelpText<"Enable diagnostics that require the SYCL integration header.">,
   MarshallingInfoFlag<LangOpts<"SYCLEnableIntHeaderDiags">>;
+def fsycl_use_main_file_name : Flag<["-"], "fsycl-use-main-file-name">,
+  HelpText<"Tells compiler that -main-file-name contains an absolute path and "
+           "file specified there should be used for checksum calculation.">,
+  MarshallingInfoFlag<CodeGenOpts<"SYCLUseMainFileName">>;
 
 } // let Flags = [CC1Option, NoDriverOption]
 

clang/lib/CodeGen/CGDebugInfo.cpp

@@ -375,6 +375,19 @@ Optional<StringRef> CGDebugInfo::getSource(const SourceManager &SM,
   return Source;
 }
 
+// Compute valid FID for FileName.
+FileID ComputeValidFileID(SourceManager &SM, StringRef FileName) {
+  FileID MainFileID = SM.getMainFileID();
+  // Find the filename FileName and load it.
+  llvm::Expected<FileEntryRef>  ExpectedFileRef =
+    SM.getFileManager().getFileRef(FileName);
+  if (ExpectedFileRef) {
+    MainFileID = SM.getOrCreateFileID(ExpectedFileRef.get(),
+                                      SrcMgr::CharacteristicKind::C_User);
+  }
+  return MainFileID;
+}
+
 llvm::DIFile *CGDebugInfo::getOrCreateFile(SourceLocation Loc) {
   SourceManager &SM = CGM.getContext().getSourceManager();
   StringRef FileName;
@@ -407,6 +420,13 @@ llvm::DIFile *CGDebugInfo::getOrCreateFile(SourceLocation Loc) {
 
   SmallString<32> Checksum;
 
+  if (SM.getFileEntryForID(SM.getMainFileID()) &&
+      CGM.getCodeGenOpts().SYCLUseMainFileName && FID.isInvalid())
+    // When an integration footer is involved, the main file is a temporary
+    // file generated by the compiler. FileName is pointing to original user
+    // source file. We use it here to properly calculate its checksum.
+    FID = ComputeValidFileID(SM, FileName);
+
   Optional<llvm::DIFile::ChecksumKind> CSKind = computeChecksum(FID, Checksum);
   Optional<llvm::DIFile::ChecksumInfo<StringRef>> CSInfo;
   if (CSKind)
@@ -513,22 +533,37 @@ void CGDebugInfo::CreateCompileUnit() {
   // Get absolute path name.
   SourceManager &SM = CGM.getContext().getSourceManager();
   std::string MainFileName = CGM.getCodeGenOpts().MainFileName;
+  std::string FullMainFileName = CGM.getCodeGenOpts().FullMainFileName;
   if (MainFileName.empty())
     MainFileName = "<stdin>";
 
   // The main file name provided via the "-main-file-name" option contains just
   // the file name itself with no path information. This file name may have had
   // a relative path, so we look into the actual file entry for the main
   // file to determine the real absolute path for the file.
+  // An exception here is workflow when integration footer is involved: in that
+  // case driver passes an absolute path to the original user-provided source
+  // file, whilst main file corresponds to a temporary file generated by the
+  // compiler.
   std::string MainFileDir;
   if (Optional<FileEntryRef> MainFile =
           SM.getFileEntryRefForID(SM.getMainFileID())) {
     MainFileDir = std::string(MainFile->getDir().getName());
-    if (!llvm::sys::path::is_absolute(MainFileName)) {
+    FileID MainFileID = SM.getMainFileID();
+    if (!llvm::sys::path::is_absolute(MainFileName) &&
+        !CGM.getCodeGenOpts().SYCLUseMainFileName) {
       llvm::SmallString<1024> MainFileDirSS(MainFileDir);
       llvm::sys::path::append(MainFileDirSS, MainFileName);
       MainFileName =
           std::string(llvm::sys::path::remove_leading_dotslash(MainFileDirSS));
+    } else if (CGM.getCodeGenOpts().SYCLUseMainFileName) {
+      // When an integration footer is involved, the main file is a temporary
+      // file generated by the compiler. FullMainFileName is pointing to
+      // original user source file. We use it here to properly calculate its
+      // checksum.
+      MainFileID = ComputeValidFileID(SM, FullMainFileName);
+      // Make sure the filename points to the original user source filename.
+      MainFileName = FullMainFileName;
     }
     // If the main file name provided is identical to the input file name, and
     // if the input file is a preprocessed source, use the module name for
@@ -540,7 +575,7 @@ void CGDebugInfo::CreateCompileUnit() {
             .isPreprocessed())
       MainFileName = CGM.getModule().getName().str();
 
-    CSKind = computeChecksum(SM.getMainFileID(), Checksum);
+    CSKind = computeChecksum(MainFileID, Checksum);
   }
 
   llvm::dwarf::SourceLanguage LangTag;

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5438,8 +5438,20 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   // Set the main file name, so that debug info works even with
   // -save-temps.
   CmdArgs.push_back("-main-file-name");
-  CmdArgs.push_back(getBaseInputName(Args, Input));
+  if (!IsSYCL || Args.hasArg(options::OPT_fno_sycl_use_footer)) {
+    CmdArgs.push_back(getBaseInputName(Args, Input));
+  } else {
+    SmallString<256> AbsPath = llvm::StringRef(Input.getBaseInput());
+    D.getVFS().makeAbsolute(AbsPath);
+    CmdArgs.push_back(
+        Args.MakeArgString(llvm::sys::path::filename(Input.getBaseInput())));
+    CmdArgs.push_back("-fsycl-use-main-file-name");
+  }
 
+  if (IsSYCL || Args.hasArg(options::OPT_fsycl_footer_path_EQ)) {
+    CmdArgs.push_back("-full-main-file-name");
+    CmdArgs.push_back(Input.getBaseInput());
+  }
   // Some flags which affect the language (via preprocessor
   // defines).
   if (Args.hasArg(options::OPT_static))

clang/test/CodeGenSYCL/Inputs/checksum.cpp

@@ -0,0 +1,4 @@
+int main() {
+  int x = 0;
+  return x + 1;
+}

clang/test/CodeGenSYCL/Inputs/debug-info-checksum.cpp

@@ -0,0 +1,9 @@
+#include "Inputs/sycl.hpp"
+
+int main() {
+  sycl::sampler Sampler;
+  sycl::kernel_single_task<class use_kernel_for_test>([=]() {
+    Sampler.use();
+  });
+  return 0;
+}

clang/test/CodeGenSYCL/debug-info-checksum-temp-name.cpp

@@ -0,0 +1,29 @@
+// This file attempts to emulate content of a file, which is passed to host
+// compiler when integration footer is used: this file is considered to be a
+// main file and Inputs/debug-info-checksum.cpp serves as integration footer.
+//
+// The file is specifically named debug-info-checksum-temp-name.cpp to emulate
+// the temporary file name, which is generated by the compiler driver for the
+// host .cpp file after appending footer to it.
+//
+// The command line executed is based on what is actually invoked by the
+// compiler driver and we explicitly pass 'Inputs/debug-info-checksum.cpp' as
+// a main file name to ensure that we can instruct the compiler to emit the
+// correct debug info (paths and checksums), even though the input file is not
+// exactly what user specified on the command line.
+//
+// RUN: %clang_cc1 -fsycl-is-host -I %S %S/Inputs/debug-info-checksum.cpp \
+// RUN:   -triple x86_64-unknown-linux-gpu \
+// RUN:   -main-file-name "%S/Inputs/debug-info-checksum.cpp" \
+// RUN:   -full-main-file-name "%S/Inputs/debug-info-checksum.cpp" \
+// RUN:   -fsycl-use-main-file-name -dwarf-version=5 -S -emit-llvm \
+// RUN:   -O0 -debug-info-kind=constructor -o - | FileCheck %s
+//
+// Verify that DICompileUnit points to a correct file and that checksum is also
+// correct.
+//
+// CHECK: !DICompileUnit({{.*}} file: ![[#FILE:]]
+// CHECK: ![[#FILE]] = !DIFile(filename: "{{.*}}clang{{.+}}test{{.+}}CodeGenSYCL{{.+}}Inputs{{.+}}debug-info-checksum.cpp"
+// CHECK-SAME: checksumkind: CSK_MD5, checksum: "f1fb5d68350b47d90a53968ac8c40529"
+
+#include "Inputs/debug-info-checksum.cpp"

clang/test/CodeGenSYCL/debug-info-file-checksum.cpp

@@ -0,0 +1,33 @@
+// This test checks that a checksum is created correctly for the compiled file,
+// and that the same checksum is generated for host and target compilation.
+// It also checks that DICompileUnit in host and target compilation is referring
+// to the original source file name (not the temporary file created by the
+// compilation process) .
+
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -fsycl-is-device \
+// RUN: -fsycl-int-header=%t.header.h -fsycl-int-footer=%t.footer.h \
+// RUN: -main-file-name %S/Inputs/checksum.cpp -fsycl-use-main-file-name \
+// RUN: -full-main-file-name "%S/Inputs/checksum.cpp" \
+// RUN: -gcodeview -debug-info-kind=limited -emit-llvm -O0 -o - "%S/Inputs/checksum.cpp" \
+// RUN: | FileCheck %s -check-prefix=COMP1
+
+// RUN: append-file "%S/Inputs/checksum.cpp" \
+// RUN: --append=%t.footer.h \
+// RUN: --orig-filename="%S/Inputs/checksum.cpp" \
+// RUN: --output=%t.checksum.cpp --use-include
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsycl-is-host \
+// RUN: -include %t.header.h -dependency-filter %t.header.h \
+// RUN: -main-file-name %S/Inputs/checksum.cpp -fsycl-use-main-file-name \
+// RUN: -full-main-file-name %S/Inputs/checksum.cpp \
+// RUN: -gcodeview -debug-info-kind=limited -emit-llvm -O0 -o - \
+// RUN: %t.checksum.cpp \
+// RUN: | FileCheck %s -check-prefix=COMP2
+
+// COMP1: !DICompileUnit({{.*}} file: ![[#FILE1:]]
+// COMP1: ![[#FILE1]] = !DIFile(filename: "{{.*}}clang{{.+}}test{{.+}}CodeGenSYCL{{.+}}checksum.cpp"
+// COMP1-SAME: checksumkind: CSK_MD5, checksum: "259269f735d83ec32c46a11352458493")
+
+// COMP2: !DICompileUnit({{.*}} file: ![[#FILE2:]]
+// COMP2: ![[#FILE2]] = !DIFile(filename: "{{.*}}clang{{.+}}test{{.+}}CodeGenSYCL{{.+}}checksum.cpp"
+// COMP2-SAME: checksumkind: CSK_MD5, checksum: "259269f735d83ec32c46a11352458493")

clang/test/Driver/sycl-int-footer.cpp

@@ -3,7 +3,7 @@
 // RUN:   | FileCheck -check-prefix FOOTER %s -DSRCDIR=%/S -DCMDDIR=cmdline/dir
 // FOOTER: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]" "-fsycl-int-footer=[[INTFOOTER:.+\h]]" "-sycl-std={{.*}}"{{.*}} "-include" "dummy.h"
 // FOOTER: append-file{{.*}} "[[INPUTFILE:.+\.cpp]]" "--append=[[INTFOOTER]]" "--orig-filename=[[INPUTFILE]]" "--output=[[APPENDEDSRC:.+\.cpp]]"
-// FOOTER: clang{{.*}} "-include" "[[INTHEADER]]"{{.*}} "-fsycl-is-host"{{.*}} "-include" "dummy.h"{{.*}} "-iquote" "[[SRCDIR]]" "-I" "cmdline/dir"
+// FOOTER: clang{{.*}} "-include" "[[INTHEADER]]"{{.*}} "-fsycl-is-host"{{.*}} "-main-file-name" "[[SRCFILE:.+\cpp]]" "-fsycl-use-main-file-name{{.*}} "-include" "dummy.h"{{.*}} "-iquote" "[[SRCDIR]]" "-I" "cmdline/dir"
 // FOOTER-NOT: "-include" "[[INTHEADER]]"
 
 /// Preprocessed file creation with integration footer
@@ -24,9 +24,15 @@
 
 /// Check that integration footer can be disabled
 // RUN:  %clangxx -fsycl -fno-sycl-use-footer %s -### 2>&1 \
-// RUN:   | FileCheck -check-prefix NO-FOOTER --implicit-check-not "-fsycl-int-footer" %s
+// RUN:   | FileCheck -check-prefix NO-FOOTER --implicit-check-not "-fsycl-int-footer" --implicit-check-not "-fsycl-use-main-file-name" %s
 // NO-FOOTER: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]" "-sycl-std={{.*}}"
-// NO-FOOTER: clang{{.*}} "-include" "[[INTHEADER]]"{{.*}} "-fsycl-is-host"{{.*}} "-o"
+// NO-FOOTER: clang{{.*}} "-include" "[[INTHEADER]]"{{.*}} "-fsycl-is-host"{{.*}} "-main-file-name" "sycl-int-footer.cpp"{{.*}} "-o"
+
+// Test that -fsycl-use-main-file-name is not passed if -fsycl is not passed.
+// This test is located here, because -fsycl-use-main-file-name is tightly
+// connected to the integration footer.
+// RUN: %clangxx %s -### 2>&1 | FileCheck %s --check-prefix NO-FSYCL --implicit-check-not "-fsycl-use-main-file-name"
+// NO-FSYCL: clang{{.*}} "-main-file-name" "sycl-int-footer.cpp"
 
 /// Check phases without integration footer
 // RUN: %clangxx -fsycl -fno-sycl-instrument-device-code -fno-sycl-device-lib=all -fno-sycl-use-footer -target x86_64-unknown-linux-gnu %s -ccc-print-phases 2>&1 \