From e74182c01e67646336095c26e3652c0276a660d8 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 21 Jan 2025 18:52:59 -0800 Subject: [PATCH 1/4] [SYCLomatic] Fix not migrated file in the codpein cuda folder. Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/FileGenerator/GenFiles.cpp | 161 +++++++++++----------- clang/lib/Tooling/Tooling.cpp | 2 +- 2 files changed, 79 insertions(+), 84 deletions(-) diff --git a/clang/lib/DPCT/FileGenerator/GenFiles.cpp b/clang/lib/DPCT/FileGenerator/GenFiles.cpp index b53e55fef8a2..d23adcafaf45 100644 --- a/clang/lib/DPCT/FileGenerator/GenFiles.cpp +++ b/clang/lib/DPCT/FileGenerator/GenFiles.cpp @@ -1087,97 +1087,92 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, "continues.\n"); } - // The necessary header files which have no replacements will be copied to - // "-out-root" directory. - for (const auto &Entry : IncludeFileMap) { - // Generated SYCL file in outroot. E.g., /path/to/outroot/a.dp.cpp - clang::tooling::UnifiedPath FilePath = Entry.first; - // Generated CUDA file in outroot_debug. E.g., /path/to/outroot_debug/a.cu - clang::tooling::UnifiedPath DebugFilePath = Entry.first; - // Original CUDA file in inroot. E.g., /path/to/inroot/a.cu - clang::tooling::UnifiedPath OriginalFilePath = Entry.first; - if (!Entry.second) { - bool IsExcluded = DpctGlobalInfo::isExcluded(FilePath); - if (IsExcluded) { - continue; - } - // Always migrate *.cuh files to *.dp.hpp files, - // Always migrate *.cu files to *.dp.cpp files. - SourceProcessType FileType = GetSourceFileType(FilePath); - SmallString<512> TempFilePath(FilePath.getCanonicalPath()); - if (FileType & SPT_CudaHeader) { - path::replace_extension(TempFilePath, - DpctGlobalInfo::getSYCLHeaderExtension()); - } else if (FileType & SPT_CudaSource) { - path::replace_extension(TempFilePath, - DpctGlobalInfo::getSYCLSourceExtension()); - } - FilePath = TempFilePath; - - if (!rewriteCanonicalDir(FilePath, InRoot, OutRoot)) { - continue; - } - - if (dpct::DpctGlobalInfo::isCodePinEnabled() && - !rewriteCanonicalDir(DebugFilePath, InRoot, CUDAMigratedOutRoot)) { - continue; - } + // Copy the necessary files (header files, src files without replacements) to + // OutRootFolder + auto copyNecessaryFile = [&](clang::tooling::UnifiedPath OutRootFolder, + bool CUDAOutRoot = false) { + for (const auto &Entry : IncludeFileMap) { + // Generated SYCL file in outroot. E.g., /path/to/outroot/a.dp.cpp + clang::tooling::UnifiedPath FilePath = Entry.first; + // Generated CUDA file in outroot_debug. E.g., /path/to/outroot_debug/a.cu + clang::tooling::UnifiedPath DebugFilePath = Entry.first; + // Original CUDA file in inroot. E.g., /path/to/inroot/a.cu + clang::tooling::UnifiedPath OriginalFilePath = Entry.first; + if (!Entry.second) { + bool IsExcluded = DpctGlobalInfo::isExcluded(FilePath); + if (IsExcluded) { + continue; + } + // Always migrate *.cuh files to *.dp.hpp files, + // Always migrate *.cu files to *.dp.cpp files. + SourceProcessType FileType = GetSourceFileType(FilePath); + SmallString<512> TempFilePath(FilePath.getCanonicalPath()); + if ((FileType & SPT_CudaHeader) && !CUDAOutRoot) { + path::replace_extension(TempFilePath, + DpctGlobalInfo::getSYCLHeaderExtension()); + } else if ((FileType & SPT_CudaSource) && !CUDAOutRoot) { + path::replace_extension(TempFilePath, + DpctGlobalInfo::getSYCLSourceExtension()); + } + FilePath = TempFilePath; - // Check for another file with SYCL extension. For example - // In in-root we have - // * src.cpp - // * src.cu - // After migration we will end up replacing src.cpp with migrated src.cu - // when the --sycl-file-extension is cpp - // In such a case warn the user. - if (checkOverwriteAndWarn(FilePath.getCanonicalPath(), - Entry.first.getCanonicalPath())) - continue; + if (!rewriteCanonicalDir(FilePath, InRoot, OutRootFolder)) { + continue; + } - // If the file needs no replacement and it already exist, don't - // make any changes - if (fs::exists(FilePath.getCanonicalPath())) { - // A header file with this name already exists. - llvm::errs() << "File '" << FilePath - << "' already exists; skipping it.\n"; - continue; - } + // Check for another file with SYCL extension. For example + // In in-root we have + // * src.cpp + // * src.cu + // After migration we will end up replacing src.cpp with migrated src.cu + // when the --sycl-file-extension is cpp + // In such a case warn the user. + if (checkOverwriteAndWarn(FilePath.getCanonicalPath(), + Entry.first.getCanonicalPath())) + continue; - createDirectories(path::parent_path(FilePath.getCanonicalPath())); - dpct::RawFDOStream Stream(FilePath.getCanonicalPath()); + // If the file needs no replacement and it already exist, don't + // make any changes + if (fs::exists(FilePath.getCanonicalPath())) { + // A header file with this name already exists. + llvm::errs() << "File '" << FilePath + << "' already exists; skipping it.\n"; + continue; + } - llvm::Expected Result = - Tool.getFiles().getFileRef(Entry.first.getCanonicalPath()); - if (auto E = Result.takeError()) { - continue; - } - if (MapNames::PatternRewriters.empty()) { - Rewrite - .getEditBuffer(Sources.getOrCreateFileID( - *Result, clang::SrcMgr::C_User /*normal user code*/)) - .write(Stream); - } else { - std::string OutputString; - llvm::raw_string_ostream RSW(OutputString); - Rewrite - .getEditBuffer(Sources.getOrCreateFileID( - *Result, clang::SrcMgr::C_User /*normal user code*/)) - .write(RSW); - applyPatternRewriter(OutputString, Stream); - } + createDirectories(path::parent_path(FilePath.getCanonicalPath())); + dpct::RawFDOStream Stream(FilePath.getCanonicalPath()); - // This will help us to detect the same output filename - // for two different input files - OutFilePath2InFilePath[FilePath.getCanonicalPath().str()] = - Entry.first.getCanonicalPath().str(); + llvm::Expected Result = + Tool.getFiles().getFileRef(Entry.first.getCanonicalPath()); + if (auto E = Result.takeError()) { + continue; + } + if (MapNames::PatternRewriters.empty()) { + Rewrite + .getEditBuffer(Sources.getOrCreateFileID( + *Result, clang::SrcMgr::C_User /*normal user code*/)) + .write(Stream); + } else { + std::string OutputString; + llvm::raw_string_ostream RSW(OutputString); + Rewrite + .getEditBuffer(Sources.getOrCreateFileID( + *Result, clang::SrcMgr::C_User /*normal user code*/)) + .write(RSW); + applyPatternRewriter(OutputString, Stream); + } - if (dpct::DpctGlobalInfo::isCodePinEnabled()) { - // Copy non-replacement CUDA files into debug folder - fs::copy_file(OriginalFilePath.getCanonicalPath(), - DebugFilePath.getCanonicalPath()); + // This will help us to detect the same output filename + // for two different input files + OutFilePath2InFilePath[FilePath.getCanonicalPath().str()] = + Entry.first.getCanonicalPath().str(); } } - } + }; + copyNecessaryFile(OutRoot); + if (dpct::DpctGlobalInfo::isCodePinEnabled()) + copyNecessaryFile(CUDAMigratedOutRoot); std::string ScriptFineName = "Makefile.dpct"; if (!BuildScriptFile.empty()) diff --git a/clang/lib/Tooling/Tooling.cpp b/clang/lib/Tooling/Tooling.cpp index ad63d2c385e9..22ada4fb06bc 100644 --- a/clang/lib/Tooling/Tooling.cpp +++ b/clang/lib/Tooling/Tooling.cpp @@ -720,7 +720,7 @@ static void injectResourceDir(CommandLineArguments &Args, const char *Argv0, // if return value is -1, means current input file \p File is not processed, // if return value < -1, report return value to upper caller, // other values are ignored. -int ClangTool::processFiles(llvm::StringRef File,bool &ProcessingFailed, +int ClangTool::processFiles(llvm::StringRef File, bool &ProcessingFailed, bool &FileSkipped, int &StaticSymbol, ToolAction *Action) { //clear error# counter CurFileParseErrCnt=0; From bcc9837efe8e4a723ef176dd28ac38cd01f68fc1 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 21 Jan 2025 21:44:55 -0800 Subject: [PATCH 2/4] update lit Signed-off-by: Chen, Sheng S --- .../copy_file_in_db/common/test1.cpp | 5 +++++ .../copy_file_in_db/compile_commands.json | 17 +++++++++++++++++ .../dpct/debug_test/copy_file_in_db/test.h | 8 ++++++++ .../dpct/debug_test/copy_file_in_db/test2.cpp | 5 +++++ .../dpct/debug_test/copy_file_in_db/test3.cu | 19 +++++++++++++++++++ 5 files changed, 54 insertions(+) create mode 100644 clang/test/dpct/debug_test/copy_file_in_db/common/test1.cpp create mode 100644 clang/test/dpct/debug_test/copy_file_in_db/compile_commands.json create mode 100644 clang/test/dpct/debug_test/copy_file_in_db/test.h create mode 100644 clang/test/dpct/debug_test/copy_file_in_db/test2.cpp create mode 100644 clang/test/dpct/debug_test/copy_file_in_db/test3.cu diff --git a/clang/test/dpct/debug_test/copy_file_in_db/common/test1.cpp b/clang/test/dpct/debug_test/copy_file_in_db/common/test1.cpp new file mode 100644 index 000000000000..7ec25f3db0d2 --- /dev/null +++ b/clang/test/dpct/debug_test/copy_file_in_db/common/test1.cpp @@ -0,0 +1,5 @@ +// RUN: echo +// CHECK: int foo1() { +int foo1() { +return 0; +} diff --git a/clang/test/dpct/debug_test/copy_file_in_db/compile_commands.json b/clang/test/dpct/debug_test/copy_file_in_db/compile_commands.json new file mode 100644 index 000000000000..9c48df75851f --- /dev/null +++ b/clang/test/dpct/debug_test/copy_file_in_db/compile_commands.json @@ -0,0 +1,17 @@ +[ + { + "command": "g++ common/test1.cpp", + "directory": "./", + "file": "./common/test1.cpp" + }, + { + "command": "g++ test2.cpp", + "directory": "./", + "file": "./test2.cpp" + }, + { + "command": "nvcc test3.cu", + "directory": "./", + "file": "./test3.cu" + } +] diff --git a/clang/test/dpct/debug_test/copy_file_in_db/test.h b/clang/test/dpct/debug_test/copy_file_in_db/test.h new file mode 100644 index 000000000000..dbbea91048b3 --- /dev/null +++ b/clang/test/dpct/debug_test/copy_file_in_db/test.h @@ -0,0 +1,8 @@ +#include + +// CHECK: /* +// CHECK: DPCT1057:{{[0-9]+}}: Variable float_to_force was used in host code and device code. float_to_force type was updated to be used in SYCL device code and new float_to_force_host_ct1 was generated to be used in host code. You need to update the host code manually to use the new float_to_force_host_ct1. +// CHECK: */ +// CHECK: static const float float_to_force_host_ct1 = (float)(1ll << 40); +// CHECK: static dpct::constant_memory float_to_force((float)(1ll << 40)); +static __constant__ const float float_to_force = (float)(1ll << 40); diff --git a/clang/test/dpct/debug_test/copy_file_in_db/test2.cpp b/clang/test/dpct/debug_test/copy_file_in_db/test2.cpp new file mode 100644 index 000000000000..6f78ebdc92d8 --- /dev/null +++ b/clang/test/dpct/debug_test/copy_file_in_db/test2.cpp @@ -0,0 +1,5 @@ +#include "test.h" +// RUN: echo +int foo2() { + +} \ No newline at end of file diff --git a/clang/test/dpct/debug_test/copy_file_in_db/test3.cu b/clang/test/dpct/debug_test/copy_file_in_db/test3.cu new file mode 100644 index 000000000000..9cf54f799a92 --- /dev/null +++ b/clang/test/dpct/debug_test/copy_file_in_db/test3.cu @@ -0,0 +1,19 @@ +// RUN: mkdir -p %T/common +// RUN: cd %T +// RUN: cat %S/compile_commands.json > %T/compile_commands.json +// RUN: cat %S/common/test1.cpp > %T/common/test1.cpp +// RUN: cat %S/test2.cpp > %T/test2.cpp +// RUN: cat %S/test3.cu > %T/test3.cu +// RUN: cat %S/test.h > %T/test.h +// RUN: dpct --in-root=%T --out-root=%T/out -p %T --format-range=none --cuda-include-path="%cuda-path/include" --enable-codepin +// RUN: FileCheck %S/common/test1.cpp --match-full-lines --input-file %T/out_codepin_cuda/common/test1.cpp +#include "test.h" + +__global__ void kernel(){ + float a = float_to_force; +} + +int main() { + kernel<<<1, 1>>>(); + return 0; +} \ No newline at end of file From dbdefb7831489ff225b4e433572d8bf1a32b4786 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 10 Feb 2025 16:51:49 -0800 Subject: [PATCH 3/4] up Signed-off-by: Chen, Sheng S --- clang/test/dpct/debug_test/copy_file_in_db/test2.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/dpct/debug_test/copy_file_in_db/test2.cpp b/clang/test/dpct/debug_test/copy_file_in_db/test2.cpp index 6f78ebdc92d8..86c9a2f19bd8 100644 --- a/clang/test/dpct/debug_test/copy_file_in_db/test2.cpp +++ b/clang/test/dpct/debug_test/copy_file_in_db/test2.cpp @@ -2,4 +2,4 @@ // RUN: echo int foo2() { -} \ No newline at end of file +} From a92c67de9f8d37a285bd0baeb637125064192f8d Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 11 Feb 2025 23:01:43 -0800 Subject: [PATCH 4/4] update Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/FileGenerator/GenFiles.cpp | 22 +++++++++------------- 1 file changed, 9 insertions(+), 13 deletions(-) diff --git a/clang/lib/DPCT/FileGenerator/GenFiles.cpp b/clang/lib/DPCT/FileGenerator/GenFiles.cpp index d23adcafaf45..68ef9daeec98 100644 --- a/clang/lib/DPCT/FileGenerator/GenFiles.cpp +++ b/clang/lib/DPCT/FileGenerator/GenFiles.cpp @@ -1088,16 +1088,12 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, } // Copy the necessary files (header files, src files without replacements) to - // OutRootFolder - auto copyNecessaryFile = [&](clang::tooling::UnifiedPath OutRootFolder, - bool CUDAOutRoot = false) { + // OutRoot + auto copyNecessaryFiles = [&](clang::tooling::UnifiedPath OutRoot, + bool IsCUDAMigrationOutRoot = false) { for (const auto &Entry : IncludeFileMap) { // Generated SYCL file in outroot. E.g., /path/to/outroot/a.dp.cpp clang::tooling::UnifiedPath FilePath = Entry.first; - // Generated CUDA file in outroot_debug. E.g., /path/to/outroot_debug/a.cu - clang::tooling::UnifiedPath DebugFilePath = Entry.first; - // Original CUDA file in inroot. E.g., /path/to/inroot/a.cu - clang::tooling::UnifiedPath OriginalFilePath = Entry.first; if (!Entry.second) { bool IsExcluded = DpctGlobalInfo::isExcluded(FilePath); if (IsExcluded) { @@ -1107,16 +1103,16 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, // Always migrate *.cu files to *.dp.cpp files. SourceProcessType FileType = GetSourceFileType(FilePath); SmallString<512> TempFilePath(FilePath.getCanonicalPath()); - if ((FileType & SPT_CudaHeader) && !CUDAOutRoot) { + if ((FileType & SPT_CudaHeader) && !IsCUDAMigrationOutRoot) { path::replace_extension(TempFilePath, DpctGlobalInfo::getSYCLHeaderExtension()); - } else if ((FileType & SPT_CudaSource) && !CUDAOutRoot) { + } else if ((FileType & SPT_CudaSource) && !IsCUDAMigrationOutRoot) { path::replace_extension(TempFilePath, DpctGlobalInfo::getSYCLSourceExtension()); } FilePath = TempFilePath; - if (!rewriteCanonicalDir(FilePath, InRoot, OutRootFolder)) { + if (!rewriteCanonicalDir(FilePath, InRoot, OutRoot)) { continue; } @@ -1148,7 +1144,7 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, if (auto E = Result.takeError()) { continue; } - if (MapNames::PatternRewriters.empty()) { + if (MapNames::PatternRewriters.empty() || IsCUDAMigrationOutRoot) { Rewrite .getEditBuffer(Sources.getOrCreateFileID( *Result, clang::SrcMgr::C_User /*normal user code*/)) @@ -1170,9 +1166,9 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, } } }; - copyNecessaryFile(OutRoot); + copyNecessaryFiles(OutRoot); if (dpct::DpctGlobalInfo::isCodePinEnabled()) - copyNecessaryFile(CUDAMigratedOutRoot); + copyNecessaryFiles(CUDAMigratedOutRoot, true); std::string ScriptFineName = "Makefile.dpct"; if (!BuildScriptFile.empty())