From 1b73435a5a455e5e6fbf316fbcb3a805d0c7c8bd Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 13 Jun 2024 17:00:00 +0800 Subject: [PATCH 01/20] [SYCLomatic] Fix the migration when the file in the soft link. Signed-off-by: Chen, Sheng S --- .../include/clang/Tooling/Core/UnifiedPath.h | 11 ++- clang/lib/DPCT/AnalysisInfo.cpp | 14 +++ clang/lib/DPCT/AnalysisInfo.h | 6 ++ clang/lib/DPCT/SaveNewFiles.cpp | 90 ++++++++++++++++++- clang/lib/Tooling/Core/UnifiedPath.cpp | 25 +++++- .../dpct/soft_link/soft_link_dir/link/test | 1 + .../soft_link_dir/target/test/test.hpp | 6 ++ .../soft_link_dir/vector_add_format.cu | 53 +++++++++++ .../soft_link_file/link/test/test.hpp | 1 + .../soft_link_file/target/test/test.hpp | 6 ++ .../soft_link_file/vector_add_format.cu | 53 +++++++++++ .../soft_link_file_codepin/link/test/test.hpp | 1 + .../target/test/test.hpp | 5 ++ .../vector_add_format.cu | 52 +++++++++++ .../link/hello/target_soft_link | 1 + .../soft_link/soft_link_mul_dir2/link/link | 1 + .../soft_link_mul_dir2/target/test/test.hpp | 6 ++ .../soft_link_mul_dir2/vector_add_format.cu | 53 +++++++++++ 18 files changed, 382 insertions(+), 3 deletions(-) create mode 120000 clang/test/dpct/soft_link/soft_link_dir/link/test create mode 100644 clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu create mode 120000 clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu create mode 120000 clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu create mode 120000 clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link create mode 120000 clang/test/dpct/soft_link/soft_link_mul_dir2/link/link create mode 100644 clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu diff --git a/clang/include/clang/Tooling/Core/UnifiedPath.h b/clang/include/clang/Tooling/Core/UnifiedPath.h index 67a9e64310bc..0b63f53f47fb 100644 --- a/clang/include/clang/Tooling/Core/UnifiedPath.h +++ b/clang/include/clang/Tooling/Core/UnifiedPath.h @@ -25,23 +25,28 @@ class UnifiedPath { UnifiedPath() = default; UnifiedPath(const std::string &Path, const std::string &CWD = ".") : _Path(Path) { + makeAbsolute(CWD); makeCanonical(CWD); } UnifiedPath(const llvm::StringRef Path, const std::string &CWD = ".") : _Path(Path.str()) { + makeAbsolute(CWD); makeCanonical(CWD); } UnifiedPath(const llvm::Twine &Path, const std::string &CWD = ".") : _Path(Path.str()) { + makeAbsolute(CWD); makeCanonical(CWD); } UnifiedPath(const llvm::SmallVectorImpl &Path, const std::string &CWD = ".") { _Path = std::string(Path.data(), Path.size()); + makeAbsolute(CWD); makeCanonical(CWD); } UnifiedPath(const char *Path, const std::string &CWD = ".") { _Path = std::string(Path); + makeAbsolute(CWD); makeCanonical(CWD); } bool equalsTo(const std::string &RHS) { @@ -61,16 +66,20 @@ class UnifiedPath { } llvm::StringRef getCanonicalPath() const noexcept { return _CanonicalPath; } llvm::StringRef getPath() const noexcept { return _Path; } + llvm::StringRef getAbsolutePath() const noexcept { return _AbsolutePath; } void setPath(const std::string &NewPath) { _Path = NewPath; + _AbsolutePath.clear(); _CanonicalPath.clear(); makeCanonical(); + makeAbsolute(); } - private: void makeCanonical(const std::string &CWD = "."); + void makeAbsolute(const std::string &CWD = "."); std::string _Path; std::string _CanonicalPath; + std::string _AbsolutePath; static std::unordered_map CanonicalPathCache; }; bool operator==(const clang::tooling::UnifiedPath &LHS, diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 9e4d9e72f04b..3382ad981402 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1444,6 +1444,20 @@ DpctGlobalInfo::getAbsolutePath(FileID ID) { } std::optional DpctGlobalInfo::getAbsolutePath(FileEntryRef File) { + llvm::SmallString<512> FilePathAbs(File.getName()); + SM->getFileManager().makeAbsolutePath(FilePathAbs); + return clang::tooling::UnifiedPath(FilePathAbs); +} + +std::optional +DpctGlobalInfo::getUnifiedPath(FileID ID) { + assert(SM && "SourceManager must be initialized"); + if (auto FileEntryRef = SM->getFileEntryRefForID(ID)) + return getUnifiedPath(*FileEntryRef); + return std::nullopt; +} +std::optional +DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { if (auto RealPath = File.getFileEntry().tryGetRealPathName(); !RealPath.empty()) return clang::tooling::UnifiedPath(RealPath); diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 487b1535ad07..74bbfdcbb3cc 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1044,6 +1044,12 @@ class DpctGlobalInfo { // Return the absolute path of \p File static std::optional getAbsolutePath(FileEntryRef File); + + static std::optional getUnifiedPath(FileID ID); + // Return the unified path of \p File + static std::optional + getUnifiedPath(FileEntryRef File); + static std::pair getLocInfo(SourceLocation Loc, bool *IsInvalid = nullptr /* out */); static std::string getTypeName(QualType QT, const ASTContext &Context); diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index 1be73240c0dc..38f6441616c6 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -38,6 +38,7 @@ #include #include #include +#include using namespace clang::dpct; using namespace llvm; @@ -150,6 +151,90 @@ bool rewriteDir(clang::tooling::UnifiedPath &FilePath, return true; } +bool rewriteAbsoluteDir(clang::tooling::UnifiedPath &FilePath, + const clang::tooling::UnifiedPath &InRoot, + const clang::tooling::UnifiedPath &OutRoot) { +#if defined(_WIN64) + std::string Filename = sys::path::filename(FilePath.getPath()).str(); +#endif + + if (!isChildPath(InRoot, FilePath) || DpctGlobalInfo::isExcluded(FilePath)) { + // Skip rewriting file path if FilePath is not child of InRoot + // E.g, + // FilePath : /path/to/inc/util.cuh + // InRoot : /path/to/inroot + // OutRoot : /path/to/outroot + // AnalysisScope : /path/to + return false; + } + auto PathDiff = std::mismatch(path::begin(FilePath.getAbsolutePath()), + path::end(FilePath.getAbsolutePath()), + path::begin(InRoot.getAbsolutePath())); + SmallString<512> NewFilePath = OutRoot.getAbsolutePath(); + path::append(NewFilePath, PathDiff.first, + path::end(FilePath.getAbsolutePath())); +#if defined(_WIN64) + sys::path::remove_filename(NewFilePath); + sys::path::append(NewFilePath, Filename); +#endif + + FilePath = NewFilePath; + return true; +} + +void createSymLink(const clang::tooling::UnifiedPath &FilePath, + const clang::tooling::UnifiedPath &InRoot, + const clang::tooling::UnifiedPath &OutRoot) { + llvm::outs() << "Code pin " << OutRoot.getAbsolutePath() << "\n"; + if (llvm::sys::fs::exists(FilePath.getCanonicalPath())) { + auto SourcePath = FilePath; + while ( + isChildPath(InRoot.getAbsolutePath(), SourcePath.getAbsolutePath())) { + auto AbsolutePath = SourcePath.getAbsolutePath(); + if (llvm::sys::fs::is_symlink_file(AbsolutePath)) { + std::filesystem::path RealPathLink = + std::filesystem::read_symlink( + std::filesystem::path(AbsolutePath.str())) + .string(); + + // The code is to create symbol file and link them to the target file. + // The code will iterate to get the parent path of the + // absolute path of the file. Then create the target directory if the + // canonical path is not exists in the out root path. + auto ReplPath = SourcePath; + auto ReplCanonicalPath = + tooling::UnifiedPath(SourcePath.getCanonicalPath()); + rewriteAbsoluteDir(ReplPath, InRoot, OutRoot); + rewriteDir(ReplCanonicalPath, InRoot, OutRoot); + if (llvm::sys::fs::is_directory(SourcePath.getAbsolutePath())) { + if (!llvm::sys::fs::exists(ReplCanonicalPath.getCanonicalPath())) + createDirectories(ReplCanonicalPath.getCanonicalPath()); + } + + auto AbsoluteParentPath = sys::path::parent_path(AbsolutePath); + + if (llvm::sys::fs::is_symlink_file(AbsoluteParentPath)) { + createSymLink(AbsoluteParentPath, InRoot, OutRoot); + } + auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath); + rewriteAbsoluteDir(ParentPath, InRoot, OutRoot); + llvm::outs() << "Parrrrr " << ParentPath.getCanonicalPath() <<"\n"; + if (!llvm::sys::fs::exists(ParentPath.getAbsolutePath())) { + createDirectories(ParentPath.getAbsolutePath()); + } + try { + std::filesystem::create_symlink( + std::filesystem::path(RealPathLink), + std::filesystem::path(ReplPath.getAbsolutePath().str())); + } catch (const std::filesystem::filesystem_error &e) { + std::cerr << "Error creating symbolic link: " << e.what() + << std::endl; + } + } + SourcePath = path::parent_path(AbsolutePath); + } + } +} void rewriteFileName(clang::tooling::UnifiedPath &FileName) { rewriteFileName(FileName, FileName); } @@ -976,7 +1061,7 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, 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 + // Generated CUDA file in outroot_codepin. E.g., /path/to/outroot_codepin/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; @@ -1061,6 +1146,9 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, DebugFilePath.getCanonicalPath()); } } + createSymLink(Entry.first, InRoot, SYCLMigratedOutRoot); + if (dpct::DpctGlobalInfo::isCodePinEnabled()) + createSymLink(Entry.first, InRoot, CodePinCUDAFolder); } std::string ScriptFineName = "Makefile.dpct"; diff --git a/clang/lib/Tooling/Core/UnifiedPath.cpp b/clang/lib/Tooling/Core/UnifiedPath.cpp index 12d04cc3edb7..b02080ee7859 100644 --- a/clang/lib/Tooling/Core/UnifiedPath.cpp +++ b/clang/lib/Tooling/Core/UnifiedPath.cpp @@ -46,7 +46,7 @@ void UnifiedPath::makeCanonical(const std::string &CWD) { llvm::sys::path::remove_dots(Path, /* remove_dot_dot= */ true); llvm::sys::path::native(Path); - + _AbsolutePath = Path.str(); llvm::SmallString<512> RealPath; // We need make sure the input `Path` for llvm::sys::fs::real_path is // exsiting, or else the behavior of real_path() is unexpected. @@ -82,6 +82,29 @@ void UnifiedPath::makeCanonical(const std::string &CWD) { #endif CanonicalPathCache.insert(std::pair(_Path, _CanonicalPath)); } + +void UnifiedPath::makeAbsolute(const std::string &CWD) { + if (_Path.empty()) { + return; + } + llvm::SmallString<512> Path(_Path); + llvm::sys::fs::expand_tilde(Path, Path); + if (!llvm::sys::path::is_absolute(Path)) { + llvm::SmallString<512> TempPath; + if (CWD == ".") { + llvm::sys::fs::current_path(TempPath); + } else { + UnifiedPath UnifiedCWD(CWD); + TempPath = UnifiedCWD.getCanonicalPath(); + } + llvm::sys::path::append(TempPath, llvm::sys::path::Style::native, Path); + Path = TempPath; + } + + llvm::sys::path::remove_dots(Path, /* remove_dot_dot= */ true); + llvm::sys::path::native(Path); + _AbsolutePath = Path.str(); +} std::unordered_map UnifiedPath::CanonicalPathCache; bool operator==(const clang::tooling::UnifiedPath &LHS, const clang::tooling::UnifiedPath &RHS) { diff --git a/clang/test/dpct/soft_link/soft_link_dir/link/test b/clang/test/dpct/soft_link/soft_link_dir/link/test new file mode 120000 index 000000000000..78bc33729bac --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_dir/link/test @@ -0,0 +1 @@ +../target \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu new file mode 100644 index 000000000000..9e0058e66fb1 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu @@ -0,0 +1,53 @@ +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp new file mode 120000 index 000000000000..6c8ad69f6c14 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp @@ -0,0 +1 @@ +../../target/test/test.hpp \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu new file mode 100644 index 000000000000..9e0058e66fb1 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu @@ -0,0 +1,53 @@ +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp new file mode 120000 index 000000000000..6c8ad69f6c14 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp @@ -0,0 +1 @@ +../../target/test/test.hpp \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp new file mode 100644 index 000000000000..c56c85ed7a4c --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp @@ -0,0 +1,5 @@ +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu new file mode 100644 index 000000000000..1fd7dfd06bc2 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu @@ -0,0 +1,52 @@ +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link new file mode 120000 index 000000000000..6bcd2fc5d22f --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link @@ -0,0 +1 @@ +../../target \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/link b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/link new file mode 120000 index 000000000000..b6fc4c620b67 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/link @@ -0,0 +1 @@ +hello \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu new file mode 100644 index 000000000000..258f690c84b3 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu @@ -0,0 +1,53 @@ +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/hello/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/hello/target_soft_link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} From 0a69bf45a4ea04118bc4cde7334a46aafa098521 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 14 Jun 2024 13:56:15 +0800 Subject: [PATCH 02/20] update Signed-off-by: Chen, Sheng S --- .../soft_link_dir/link/test | 0 .../soft_link_dir/target/test/test.hpp | 0 .../soft_link_dir/vector_add_format.cu | 54 ++++++++++++++++++ .../soft_link_file/link/test/test.hpp | 0 .../soft_link_file/target/test/test.hpp | 0 .../soft_link_file}/vector_add_format.cu | 1 + .../soft_link_file_codepin/link/test/test.hpp | 0 .../target/test/test.hpp | 0 .../vector_add_format.cu | 1 + .../link/hello/target_soft_link | 0 .../soft_link_mul_dir2/link/link | 0 .../soft_link_mul_dir2/target/test/test.hpp | 0 .../soft_link_mul_dir2/vector_add_format.cu | 1 + .../soft_link_win/soft_link_dir/link/temp | 0 .../soft_link_dir/target/test/test.hpp | 6 ++ .../soft_link_dir/vector_add_format.cu | 55 +++++++++++++++++++ .../soft_link_file/link/test/temp | 0 .../soft_link_file/target/test/test.hpp | 6 ++ .../soft_link_file/vector_add_format.cu | 2 + .../soft_link_file_codepin/link/test/temp | 0 .../target/test/test.hpp | 5 ++ .../vector_add_format.cu | 54 ++++++++++++++++++ .../soft_link_mul_dir2/link/hello/temp | 0 .../soft_link_mul_dir2/target/test/test.hpp | 6 ++ .../soft_link_mul_dir2/vector_add_format.cu | 55 +++++++++++++++++++ 25 files changed, 246 insertions(+) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_dir/link/test (100%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_dir/target/test/test.hpp (100%) create mode 100644 clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_file/link/test/test.hpp (100%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_file/target/test/test.hpp (100%) rename clang/test/dpct/{soft_link/soft_link_dir => soft_link_lin/soft_link_file}/vector_add_format.cu (98%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_file_codepin/link/test/test.hpp (100%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_file_codepin/target/test/test.hpp (100%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_file_codepin/vector_add_format.cu (98%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_mul_dir2/link/hello/target_soft_link (100%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_mul_dir2/link/link (100%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_mul_dir2/target/test/test.hpp (100%) rename clang/test/dpct/{soft_link => soft_link_lin}/soft_link_mul_dir2/vector_add_format.cu (98%) create mode 100644 clang/test/dpct/soft_link_win/soft_link_dir/link/temp create mode 100644 clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu create mode 100644 clang/test/dpct/soft_link_win/soft_link_file/link/test/temp create mode 100644 clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp rename clang/test/dpct/{soft_link => soft_link_win}/soft_link_file/vector_add_format.cu (94%) create mode 100644 clang/test/dpct/soft_link_win/soft_link_file_codepin/link/test/temp create mode 100644 clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu create mode 100644 clang/test/dpct/soft_link_win/soft_link_mul_dir2/link/hello/temp create mode 100644 clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu diff --git a/clang/test/dpct/soft_link/soft_link_dir/link/test b/clang/test/dpct/soft_link_lin/soft_link_dir/link/test similarity index 100% rename from clang/test/dpct/soft_link/soft_link_dir/link/test rename to clang/test/dpct/soft_link_lin/soft_link_dir/link/test diff --git a/clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_dir/target/test/test.hpp similarity index 100% rename from clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp rename to clang/test/dpct/soft_link_lin/soft_link_dir/target/test/test.hpp diff --git a/clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu new file mode 100644 index 000000000000..dd79f21ff4b9 --- /dev/null +++ b/clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu @@ -0,0 +1,54 @@ +// UNSUPPORTED: system-windows +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/test -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/test %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file/link/test/test.hpp similarity index 100% rename from clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp rename to clang/test/dpct/soft_link_lin/soft_link_file/link/test/test.hpp diff --git a/clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file/target/test/test.hpp similarity index 100% rename from clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp rename to clang/test/dpct/soft_link_lin/soft_link_file/target/test/test.hpp diff --git a/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_file/vector_add_format.cu similarity index 98% rename from clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu rename to clang/test/dpct/soft_link_lin/soft_link_file/vector_add_format.cu index 9e0058e66fb1..4c10fd9ca9b5 100644 --- a/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu +++ b/clang/test/dpct/soft_link_lin/soft_link_file/vector_add_format.cu @@ -1,3 +1,4 @@ +// UNSUPPORTED: system-windows // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file_codepin/link/test/test.hpp similarity index 100% rename from clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp rename to clang/test/dpct/soft_link_lin/soft_link_file_codepin/link/test/test.hpp diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file_codepin/target/test/test.hpp similarity index 100% rename from clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp rename to clang/test/dpct/soft_link_lin/soft_link_file_codepin/target/test/test.hpp diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_file_codepin/vector_add_format.cu similarity index 98% rename from clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu rename to clang/test/dpct/soft_link_lin/soft_link_file_codepin/vector_add_format.cu index 1fd7dfd06bc2..ccb6d4a5ce2b 100644 --- a/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu +++ b/clang/test/dpct/soft_link_lin/soft_link_file_codepin/vector_add_format.cu @@ -1,3 +1,4 @@ +// UNSUPPORTED: system-windows // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/hello/target_soft_link similarity index 100% rename from clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link rename to clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/hello/target_soft_link diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/link b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/link similarity index 100% rename from clang/test/dpct/soft_link/soft_link_mul_dir2/link/link rename to clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/link diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/target/test/test.hpp similarity index 100% rename from clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp rename to clang/test/dpct/soft_link_lin/soft_link_mul_dir2/target/test/test.hpp diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/vector_add_format.cu similarity index 98% rename from clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu rename to clang/test/dpct/soft_link_lin/soft_link_mul_dir2/vector_add_format.cu index 258f690c84b3..48c009409609 100644 --- a/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu +++ b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/vector_add_format.cu @@ -1,3 +1,4 @@ +// UNSUPPORTED: system-windows // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/hello/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp diff --git a/clang/test/dpct/soft_link_win/soft_link_dir/link/temp b/clang/test/dpct/soft_link_win/soft_link_dir/link/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu new file mode 100644 index 000000000000..3a371361606b --- /dev/null +++ b/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu @@ -0,0 +1,55 @@ +// UNSUPPORTED: system-linux +// RUN: cd %S && mklink /D link\test target +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/test -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/test %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link_win/soft_link_file/link/test/temp b/clang/test/dpct/soft_link_win/soft_link_file/link/test/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu similarity index 94% rename from clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu rename to clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu index 9e0058e66fb1..2113265ac381 100644 --- a/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu +++ b/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu @@ -1,3 +1,5 @@ +// UNSUPPORTED: system-linux +// RUN: cd %S && mklink link\test\test.hpp target\test\test.hpp // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp diff --git a/clang/test/dpct/soft_link_win/soft_link_file_codepin/link/test/temp b/clang/test/dpct/soft_link_win/soft_link_file_codepin/link/test/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp new file mode 100644 index 000000000000..c56c85ed7a4c --- /dev/null +++ b/clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp @@ -0,0 +1,5 @@ +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu new file mode 100644 index 000000000000..9efcfdde82f6 --- /dev/null +++ b/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu @@ -0,0 +1,54 @@ +// UNSUPPORTED: system-linux +// RUN: cd %S && mklink link\test\test.hpp target\test\test.hpp +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/link/hello/temp b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/link/hello/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu new file mode 100644 index 000000000000..5381136756a6 --- /dev/null +++ b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu @@ -0,0 +1,55 @@ +// UNSUPPORTED: system-linux +// RUN: cd %S && mklink /D link\hello\target_soft_link %S\target && mklink /D link\link %S\hello +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/hello/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/hello/target_soft_link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} From e10dea08ec2e5330b3a7f08be8d02e1512f7e756 Mon Sep 17 00:00:00 2001 From: Chen Date: Fri, 14 Jun 2024 14:39:04 +0800 Subject: [PATCH 03/20] update Signed-off-by: Chen --- .../soft_link_lin/soft_link_dir/link/test | 1 - .../soft_link_dir/target/test/test.hpp | 6 --- .../soft_link_dir/vector_add_format.cu | 54 ------------------- .../soft_link_file/link/test/test.hpp | 1 - .../soft_link_file/target/test/test.hpp | 6 --- .../soft_link_file/vector_add_format.cu | 54 ------------------- .../soft_link_file_codepin/link/test/test.hpp | 1 - .../target/test/test.hpp | 5 -- .../vector_add_format.cu | 53 ------------------ .../link/hello/target_soft_link | 1 - .../soft_link_mul_dir2/link/link | 1 - .../soft_link_mul_dir2/target/test/test.hpp | 6 --- .../soft_link_mul_dir2/vector_add_format.cu | 54 ------------------- .../soft_link_dir/vector_add_format.cu | 3 +- .../soft_link_file/vector_add_format.cu | 3 +- .../vector_add_format.cu | 3 +- .../soft_link_mul_dir2/vector_add_format.cu | 5 +- 17 files changed, 5 insertions(+), 252 deletions(-) delete mode 120000 clang/test/dpct/soft_link_lin/soft_link_dir/link/test delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_dir/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu delete mode 120000 clang/test/dpct/soft_link_lin/soft_link_file/link/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_file/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_file/vector_add_format.cu delete mode 120000 clang/test/dpct/soft_link_lin/soft_link_file_codepin/link/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_file_codepin/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_file_codepin/vector_add_format.cu delete mode 120000 clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/hello/target_soft_link delete mode 120000 clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/link delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_mul_dir2/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_lin/soft_link_mul_dir2/vector_add_format.cu diff --git a/clang/test/dpct/soft_link_lin/soft_link_dir/link/test b/clang/test/dpct/soft_link_lin/soft_link_dir/link/test deleted file mode 120000 index 78bc33729bac..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_dir/link/test +++ /dev/null @@ -1 +0,0 @@ -../target \ No newline at end of file diff --git a/clang/test/dpct/soft_link_lin/soft_link_dir/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_dir/target/test/test.hpp deleted file mode 100644 index cdf9d602ebaa..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_dir/target/test/test.hpp +++ /dev/null @@ -1,6 +0,0 @@ -// CHECK: #include -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu deleted file mode 100644 index dd79f21ff4b9..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_dir/vector_add_format.cu +++ /dev/null @@ -1,54 +0,0 @@ -// UNSUPPORTED: system-windows -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/test -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s -// RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/test %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} diff --git a/clang/test/dpct/soft_link_lin/soft_link_file/link/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file/link/test/test.hpp deleted file mode 120000 index 6c8ad69f6c14..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_file/link/test/test.hpp +++ /dev/null @@ -1 +0,0 @@ -../../target/test/test.hpp \ No newline at end of file diff --git a/clang/test/dpct/soft_link_lin/soft_link_file/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file/target/test/test.hpp deleted file mode 100644 index cdf9d602ebaa..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_file/target/test/test.hpp +++ /dev/null @@ -1,6 +0,0 @@ -// CHECK: #include -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_lin/soft_link_file/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_file/vector_add_format.cu deleted file mode 100644 index 4c10fd9ca9b5..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_file/vector_add_format.cu +++ /dev/null @@ -1,54 +0,0 @@ -// UNSUPPORTED: system-windows -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s -// RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} diff --git a/clang/test/dpct/soft_link_lin/soft_link_file_codepin/link/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file_codepin/link/test/test.hpp deleted file mode 120000 index 6c8ad69f6c14..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_file_codepin/link/test/test.hpp +++ /dev/null @@ -1 +0,0 @@ -../../target/test/test.hpp \ No newline at end of file diff --git a/clang/test/dpct/soft_link_lin/soft_link_file_codepin/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_file_codepin/target/test/test.hpp deleted file mode 100644 index c56c85ed7a4c..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_file_codepin/target/test/test.hpp +++ /dev/null @@ -1,5 +0,0 @@ -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_lin/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_file_codepin/vector_add_format.cu deleted file mode 100644 index ccb6d4a5ce2b..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_file_codepin/vector_add_format.cu +++ /dev/null @@ -1,53 +0,0 @@ -// UNSUPPORTED: system-windows -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} diff --git a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/hello/target_soft_link b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/hello/target_soft_link deleted file mode 120000 index 6bcd2fc5d22f..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/hello/target_soft_link +++ /dev/null @@ -1 +0,0 @@ -../../target \ No newline at end of file diff --git a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/link b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/link deleted file mode 120000 index b6fc4c620b67..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/link/link +++ /dev/null @@ -1 +0,0 @@ -hello \ No newline at end of file diff --git a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/target/test/test.hpp b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/target/test/test.hpp deleted file mode 100644 index cdf9d602ebaa..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/target/test/test.hpp +++ /dev/null @@ -1,6 +0,0 @@ -// CHECK: #include -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/vector_add_format.cu deleted file mode 100644 index 48c009409609..000000000000 --- a/clang/test/dpct/soft_link_lin/soft_link_mul_dir2/vector_add_format.cu +++ /dev/null @@ -1,54 +0,0 @@ -// UNSUPPORTED: system-windows -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s -// RUN: FileCheck --input-file %T/out/link/hello/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/hello/target_soft_link %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu index 3a371361606b..c954afb0bc44 100644 --- a/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu +++ b/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu @@ -1,5 +1,4 @@ -// UNSUPPORTED: system-linux -// RUN: cd %S && mklink /D link\test target +// RUN: cd %S/link && rm -rf test && ln -nfs ../target test // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/test -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp diff --git a/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu index 2113265ac381..c4210c03bdd6 100644 --- a/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu +++ b/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu @@ -1,5 +1,4 @@ -// UNSUPPORTED: system-linux -// RUN: cd %S && mklink link\test\test.hpp target\test\test.hpp +// RUN: cd %S && rm link/test/test.hpp && ln -nfs %S/target/test/test.hpp link/test/test.hpp // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp diff --git a/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu index 9efcfdde82f6..c7a48c5da802 100644 --- a/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu +++ b/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu @@ -1,5 +1,4 @@ -// UNSUPPORTED: system-linux -// RUN: cd %S && mklink link\test\test.hpp target\test\test.hpp +// RUN: cd %S/link/test && rm -rf test.hpp && ln -nfs ../../target/test/test.hpp test.hpp // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} diff --git a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu index 5381136756a6..2a4b620d151e 100644 --- a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu +++ b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu @@ -1,8 +1,7 @@ -// UNSUPPORTED: system-linux -// RUN: cd %S && mklink /D link\hello\target_soft_link %S\target && mklink /D link\link %S\hello +// RUN: cd %S/link/hello && rm -rf target_soft_link && ln -nfs ../../target target_soft_link && cd %S/link && rm -rf link && ln -nfs hello link // RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s -// RUN: FileCheck --input-file %T/out/link/hello/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp +// RUN: FileCheck --input-file %T/out/link/link/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/hello/target_soft_link %} #include From 3b2351b2c81866512ee629966d714fde0de4a198 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 14 Jun 2024 14:44:14 +0800 Subject: [PATCH 04/20] update Signed-off-by: Chen, Sheng S --- .../dpct/soft_link/soft_link_dir/link/temp | 0 .../dpct/soft_link/soft_link_dir/link/test | 1 + .../soft_link_dir/target/test/test.hpp | 6 +++ .../soft_link_dir/vector_add_format.cu | 54 +++++++++++++++++++ .../soft_link/soft_link_file/link/test/temp | 0 .../soft_link_file/link/test/test.hpp | 1 + .../soft_link_file/target/test/test.hpp | 6 +++ .../soft_link_file/vector_add_format.cu | 54 +++++++++++++++++++ .../soft_link_file_codepin/link/test/temp | 0 .../soft_link_file_codepin/link/test/test.hpp | 1 + .../target/test/test.hpp | 5 ++ .../vector_add_format.cu | 53 ++++++++++++++++++ .../link/hello/target_soft_link | 1 + .../soft_link_mul_dir2/link/hello/temp | 0 .../soft_link/soft_link_mul_dir2/link/link | 1 + .../soft_link_mul_dir2/target/test/test.hpp | 6 +++ .../soft_link_mul_dir2/vector_add_format.cu | 54 +++++++++++++++++++ 17 files changed, 243 insertions(+) create mode 100644 clang/test/dpct/soft_link/soft_link_dir/link/temp create mode 120000 clang/test/dpct/soft_link/soft_link_dir/link/test create mode 100644 clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu create mode 100644 clang/test/dpct/soft_link/soft_link_file/link/test/temp create mode 120000 clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu create mode 100644 clang/test/dpct/soft_link/soft_link_file_codepin/link/test/temp create mode 120000 clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu create mode 120000 clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link create mode 100644 clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/temp create mode 120000 clang/test/dpct/soft_link/soft_link_mul_dir2/link/link create mode 100644 clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp create mode 100644 clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu diff --git a/clang/test/dpct/soft_link/soft_link_dir/link/temp b/clang/test/dpct/soft_link/soft_link_dir/link/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link/soft_link_dir/link/test b/clang/test/dpct/soft_link/soft_link_dir/link/test new file mode 120000 index 000000000000..78bc33729bac --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_dir/link/test @@ -0,0 +1 @@ +../target \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_dir/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu new file mode 100644 index 000000000000..c954afb0bc44 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu @@ -0,0 +1,54 @@ +// RUN: cd %S/link && rm -rf test && ln -nfs ../target test +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/test -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/test %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file/link/test/temp b/clang/test/dpct/soft_link/soft_link_file/link/test/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp new file mode 120000 index 000000000000..6c8ad69f6c14 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file/link/test/test.hpp @@ -0,0 +1 @@ +../../target/test/test.hpp \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu new file mode 100644 index 000000000000..c553b6871615 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu @@ -0,0 +1,54 @@ +// RUN: cd %S/link/test && rm test.hpp && ln -nfs ../../target/test/test.hpp test.hpp +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/temp b/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp new file mode 120000 index 000000000000..6c8ad69f6c14 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp @@ -0,0 +1 @@ +../../target/test/test.hpp \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp new file mode 100644 index 000000000000..c56c85ed7a4c --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp @@ -0,0 +1,5 @@ +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu new file mode 100644 index 000000000000..c7a48c5da802 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu @@ -0,0 +1,53 @@ +// RUN: cd %S/link/test && rm -rf test.hpp && ln -nfs ../../target/test/test.hpp test.hpp +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link new file mode 120000 index 000000000000..6bcd2fc5d22f --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/target_soft_link @@ -0,0 +1 @@ +../../target \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/temp b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/temp new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/link b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/link new file mode 120000 index 000000000000..b6fc4c620b67 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/link @@ -0,0 +1 @@ +hello \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/target/test/test.hpp @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu new file mode 100644 index 000000000000..2a4b620d151e --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu @@ -0,0 +1,54 @@ +// RUN: cd %S/link/hello && rm -rf target_soft_link && ln -nfs ../../target target_soft_link && cd %S/link && rm -rf link && ln -nfs hello link +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/link/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/hello/target_soft_link %} + +#include +#include +#include "test/test.hpp" +#define VECTOR_SIZE 256 + +__global__ void VectorAddKernel(float* A, float* B, float* C) +{ + A[threadIdx.x] = threadIdx.x + 1.0f; + B[threadIdx.x] = threadIdx.x + 1.0f; + C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; +} + + + +int main() +{ + // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); + // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); + float *d_A, *d_B, *d_C; + + + cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); + cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); + + + VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); + + + float Result[VECTOR_SIZE] = {}; + cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); + + // CHECK: dpct::dpct_free(d_A, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); + // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); + cudaFree(d_A); + cudaFree(d_B); + cudaFree(d_C); + + for (int i = 0; i < VECTOR_SIZE; i++) { + if (i % 16 == 0) { + printf("\n"); + } + printf("%f ", Result[i]); + } + + return 0; +} From 954cc13ef33ff529870b13b7eeaaf74212364365 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 14 Jun 2024 14:45:23 +0800 Subject: [PATCH 05/20] remove win Signed-off-by: Chen, Sheng S --- .../soft_link_win/soft_link_dir/link/temp | 0 .../soft_link_dir/target/test/test.hpp | 6 --- .../soft_link_dir/vector_add_format.cu | 54 ------------------- .../soft_link_file/link/test/temp | 0 .../soft_link_file/target/test/test.hpp | 6 --- .../soft_link_file/vector_add_format.cu | 54 ------------------- .../soft_link_file_codepin/link/test/temp | 0 .../target/test/test.hpp | 5 -- .../vector_add_format.cu | 53 ------------------ .../soft_link_mul_dir2/link/hello/temp | 0 .../soft_link_mul_dir2/target/test/test.hpp | 6 --- .../soft_link_mul_dir2/vector_add_format.cu | 54 ------------------- 12 files changed, 238 deletions(-) delete mode 100644 clang/test/dpct/soft_link_win/soft_link_dir/link/temp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu delete mode 100644 clang/test/dpct/soft_link_win/soft_link_file/link/test/temp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu delete mode 100644 clang/test/dpct/soft_link_win/soft_link_file_codepin/link/test/temp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu delete mode 100644 clang/test/dpct/soft_link_win/soft_link_mul_dir2/link/hello/temp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp delete mode 100644 clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu diff --git a/clang/test/dpct/soft_link_win/soft_link_dir/link/temp b/clang/test/dpct/soft_link_win/soft_link_dir/link/temp deleted file mode 100644 index e69de29bb2d1..000000000000 diff --git a/clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp deleted file mode 100644 index cdf9d602ebaa..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_dir/target/test/test.hpp +++ /dev/null @@ -1,6 +0,0 @@ -// CHECK: #include -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu deleted file mode 100644 index c954afb0bc44..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_dir/vector_add_format.cu +++ /dev/null @@ -1,54 +0,0 @@ -// RUN: cd %S/link && rm -rf test && ln -nfs ../target test -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/test -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s -// RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/test %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_file/link/test/temp b/clang/test/dpct/soft_link_win/soft_link_file/link/test/temp deleted file mode 100644 index e69de29bb2d1..000000000000 diff --git a/clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp deleted file mode 100644 index cdf9d602ebaa..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_file/target/test/test.hpp +++ /dev/null @@ -1,6 +0,0 @@ -// CHECK: #include -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu deleted file mode 100644 index c4210c03bdd6..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_file/vector_add_format.cu +++ /dev/null @@ -1,54 +0,0 @@ -// RUN: cd %S && rm link/test/test.hpp && ln -nfs %S/target/test/test.hpp link/test/test.hpp -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s -// RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_file_codepin/link/test/temp b/clang/test/dpct/soft_link_win/soft_link_file_codepin/link/test/temp deleted file mode 100644 index e69de29bb2d1..000000000000 diff --git a/clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp deleted file mode 100644 index c56c85ed7a4c..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_file_codepin/target/test/test.hpp +++ /dev/null @@ -1,5 +0,0 @@ -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu deleted file mode 100644 index c7a48c5da802..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_file_codepin/vector_add_format.cu +++ /dev/null @@ -1,53 +0,0 @@ -// RUN: cd %S/link/test && rm -rf test.hpp && ln -nfs ../../target/test/test.hpp test.hpp -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/link/hello/temp b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/link/hello/temp deleted file mode 100644 index e69de29bb2d1..000000000000 diff --git a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp deleted file mode 100644 index cdf9d602ebaa..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/target/test/test.hpp +++ /dev/null @@ -1,6 +0,0 @@ -// CHECK: #include -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu deleted file mode 100644 index 2a4b620d151e..000000000000 --- a/clang/test/dpct/soft_link_win/soft_link_mul_dir2/vector_add_format.cu +++ /dev/null @@ -1,54 +0,0 @@ -// RUN: cd %S/link/hello && rm -rf target_soft_link && ln -nfs ../../target target_soft_link && cd %S/link && rm -rf link && ln -nfs hello link -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s -// RUN: FileCheck --input-file %T/out/link/link/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/hello/target_soft_link %} - -#include -#include -#include "test/test.hpp" -#define VECTOR_SIZE 256 - -__global__ void VectorAddKernel(float* A, float* B, float* C) -{ - A[threadIdx.x] = threadIdx.x + 1.0f; - B[threadIdx.x] = threadIdx.x + 1.0f; - C[threadIdx.x] = A[threadIdx.x] + B[threadIdx.x]; -} - - - -int main() -{ - // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); - // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); - float *d_A, *d_B, *d_C; - - - cudaMalloc(&d_A, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_B, VECTOR_SIZE * sizeof(float)); - cudaMalloc(&d_C, VECTOR_SIZE * sizeof(float)); - - - VectorAddKernel<<<1, VECTOR_SIZE>>>(d_A, d_B, d_C); - - - float Result[VECTOR_SIZE] = {}; - cudaMemcpy(Result, d_C, VECTOR_SIZE * sizeof(float), cudaMemcpyDeviceToHost); - - // CHECK: dpct::dpct_free(d_A, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_B, q_ct1); - // CHECK-NEXT: dpct::dpct_free(d_C, q_ct1); - cudaFree(d_A); - cudaFree(d_B); - cudaFree(d_C); - - for (int i = 0; i < VECTOR_SIZE; i++) { - if (i % 16 == 0) { - printf("\n"); - } - printf("%f ", Result[i]); - } - - return 0; -} From 1f701d7044294edf46abbb76a3db71a14d049489 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 14 Jun 2024 14:46:30 +0800 Subject: [PATCH 06/20] update Signed-off-by: Chen, Sheng S --- clang/test/dpct/soft_link/soft_link_file/link/test/temp | 0 1 file changed, 0 insertions(+), 0 deletions(-) delete mode 100644 clang/test/dpct/soft_link/soft_link_file/link/test/temp diff --git a/clang/test/dpct/soft_link/soft_link_file/link/test/temp b/clang/test/dpct/soft_link/soft_link_file/link/test/temp deleted file mode 100644 index e69de29bb2d1..000000000000 From 55e41fb6681c9dc73c2139ad9907e78855967637 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 20 Jun 2024 09:02:36 +0800 Subject: [PATCH 07/20] merge pulldown code Signed-off-by: Chen, Sheng S --- clang/include/clang/Tooling/Core/UnifiedPath.h | 1 + clang/lib/DPCT/SaveNewFiles.cpp | 7 +++---- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Tooling/Core/UnifiedPath.h b/clang/include/clang/Tooling/Core/UnifiedPath.h index 0b63f53f47fb..894e4986329d 100644 --- a/clang/include/clang/Tooling/Core/UnifiedPath.h +++ b/clang/include/clang/Tooling/Core/UnifiedPath.h @@ -74,6 +74,7 @@ class UnifiedPath { makeCanonical(); makeAbsolute(); } + private: void makeCanonical(const std::string &CWD = "."); void makeAbsolute(const std::string &CWD = "."); diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index 38f6441616c6..69b6c63c9b0d 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -36,9 +36,9 @@ #include #include +#include #include #include -#include using namespace clang::dpct; using namespace llvm; @@ -185,7 +185,6 @@ bool rewriteAbsoluteDir(clang::tooling::UnifiedPath &FilePath, void createSymLink(const clang::tooling::UnifiedPath &FilePath, const clang::tooling::UnifiedPath &InRoot, const clang::tooling::UnifiedPath &OutRoot) { - llvm::outs() << "Code pin " << OutRoot.getAbsolutePath() << "\n"; if (llvm::sys::fs::exists(FilePath.getCanonicalPath())) { auto SourcePath = FilePath; while ( @@ -218,7 +217,6 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, } auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath); rewriteAbsoluteDir(ParentPath, InRoot, OutRoot); - llvm::outs() << "Parrrrr " << ParentPath.getCanonicalPath() <<"\n"; if (!llvm::sys::fs::exists(ParentPath.getAbsolutePath())) { createDirectories(ParentPath.getAbsolutePath()); } @@ -1061,7 +1059,8 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, 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_codepin. E.g., /path/to/outroot_codepin/a.cu + // Generated CUDA file in outroot_codepin. E.g., + // /path/to/outroot_codepin/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; From f5ebea558dcc1887b55eb5f2df18d81ba2cdfd5f Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 21 Jun 2024 09:00:22 +0800 Subject: [PATCH 08/20] debug Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/SaveNewFiles.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index 69b6c63c9b0d..b39d3f250f7c 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -213,6 +213,7 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, auto AbsoluteParentPath = sys::path::parent_path(AbsolutePath); if (llvm::sys::fs::is_symlink_file(AbsoluteParentPath)) { + llvm::outs() << "SYmbol link is " << AbsoluteParentPath << "\n"; createSymLink(AbsoluteParentPath, InRoot, OutRoot); } auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath); From 521ddeeda5a9551520229731d47aec94fd1fb269 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 21 Jun 2024 10:15:31 +0800 Subject: [PATCH 09/20] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 1 + clang/lib/DPCT/InclusionHeaders.cpp | 1 + clang/lib/DPCT/MigrationAction.cpp | 2 +- clang/lib/DPCT/SaveNewFiles.cpp | 6 +++--- 4 files changed, 6 insertions(+), 4 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index d4c23c5ab3f3..63e313e0a0af 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1475,6 +1475,7 @@ DpctGlobalInfo::getLocInfo(SourceLocation Loc, bool *IsInvalid) { } auto LocInfo = SM->getDecomposedLoc(SM->getExpansionLoc(Loc)); auto AbsPath = getAbsolutePath(LocInfo.first); + // auto AbsPath = getUnifiedPath(LocInfo.first); if (AbsPath) return std::make_pair(AbsPath.value(), LocInfo.second); if (IsInvalid) diff --git a/clang/lib/DPCT/InclusionHeaders.cpp b/clang/lib/DPCT/InclusionHeaders.cpp index 71ce7c1cebab..0b756c5db3c7 100644 --- a/clang/lib/DPCT/InclusionHeaders.cpp +++ b/clang/lib/DPCT/InclusionHeaders.cpp @@ -132,6 +132,7 @@ void IncludesCallbacks::InclusionDirective( LastInclusionLocationUpdater Updater(FileInfo, FilenameRange.getEnd()); clang::tooling::UnifiedPath IncludedFile; + // if (auto OptionalAbs = Global.getUnifiedPath(*File)) if (auto OptionalAbs = Global.getAbsolutePath(*File)) IncludedFile = OptionalAbs.value(); diff --git a/clang/lib/DPCT/MigrationAction.cpp b/clang/lib/DPCT/MigrationAction.cpp index 8e4bb0e5cc2e..189b750dcb23 100644 --- a/clang/lib/DPCT/MigrationAction.cpp +++ b/clang/lib/DPCT/MigrationAction.cpp @@ -95,7 +95,7 @@ void DpctConsumer::Initialize(ASTContext &Context) { // Set Context for build information DpctGlobalInfo::setContext(Context); auto &SM = Context.getSourceManager(); - auto Path = DpctGlobalInfo::getAbsolutePath(SM.getMainFileID()); + auto Path = DpctGlobalInfo::getUnifiedPath(SM.getMainFileID()); assert(Path && "Can not find absolute path"); DpctGlobalInfo::getInstance().setMainFile( Info->MainFile = DpctGlobalInfo::getInstance().insertFile(Path.value())); diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index b39d3f250f7c..ece6fb642166 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -1146,9 +1146,9 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, DebugFilePath.getCanonicalPath()); } } - createSymLink(Entry.first, InRoot, SYCLMigratedOutRoot); - if (dpct::DpctGlobalInfo::isCodePinEnabled()) - createSymLink(Entry.first, InRoot, CodePinCUDAFolder); + // createSymLink(Entry.first, InRoot, SYCLMigratedOutRoot); + // if (dpct::DpctGlobalInfo::isCodePinEnabled()) + // createSymLink(Entry.first, InRoot, CodePinCUDAFolder); } std::string ScriptFineName = "Makefile.dpct"; From 59d804f1560c9938c6eaa72e8c3c5f0a503a8ff9 Mon Sep 17 00:00:00 2001 From: Chen Date: Mon, 24 Jun 2024 20:29:49 +0800 Subject: [PATCH 10/20] update Signed-off-by: Chen --- clang/lib/DPCT/AnalysisInfo.cpp | 28 ++++++++++++++-------------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index d4c23c5ab3f3..57b210b19549 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1446,20 +1446,20 @@ DpctGlobalInfo::getAbsolutePath(FileID ID) { } std::optional DpctGlobalInfo::getAbsolutePath(FileEntryRef File) { - llvm::SmallString<512> FilePathAbs(File.getName()); - SM->getFileManager().makeAbsolutePath(FilePathAbs); - return clang::tooling::UnifiedPath(FilePathAbs); -} - -std::optional -DpctGlobalInfo::getUnifiedPath(FileID ID) { - assert(SM && "SourceManager must be initialized"); - if (auto FileEntryRef = SM->getFileEntryRefForID(ID)) - return getUnifiedPath(*FileEntryRef); - return std::nullopt; -} -std::optional -DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { +// llvm::SmallString<512> FilePathAbs(File.getName()); +// SM->getFileManager().makeAbsolutePath(FilePathAbs); +// return clang::tooling::UnifiedPath(FilePathAbs); +// } + +// std::optional +// DpctGlobalInfo::getUnifiedPath(FileID ID) { +// assert(SM && "SourceManager must be initialized"); +// if (auto FileEntryRef = SM->getFileEntryRefForID(ID)) +// return getUnifiedPath(*FileEntryRef); +// return std::nullopt; +// } +// std::optional +// DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { if (auto RealPath = File.getFileEntry().tryGetRealPathName(); !RealPath.empty()) return clang::tooling::UnifiedPath(RealPath); From 2bf2b01b4a96d959d6a421d3ee02339c6028ce1a Mon Sep 17 00:00:00 2001 From: shengchenJ Date: Wed, 26 Jun 2024 21:52:20 +0800 Subject: [PATCH 11/20] test Signed-off-by: shengchenJ --- clang/lib/DPCT/AnalysisInfo.cpp | 29 ++++++++++++----------- clang/lib/Tooling/CommonOptionsParser.cpp | 2 +- 2 files changed, 16 insertions(+), 15 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 6af380fec542..0e371a23424b 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1448,20 +1448,21 @@ DpctGlobalInfo::getAbsolutePath(FileID ID) { } std::optional DpctGlobalInfo::getAbsolutePath(FileEntryRef File) { -// llvm::SmallString<512> FilePathAbs(File.getName()); -// SM->getFileManager().makeAbsolutePath(FilePathAbs); -// return clang::tooling::UnifiedPath(FilePathAbs); -// } - -// std::optional -// DpctGlobalInfo::getUnifiedPath(FileID ID) { -// assert(SM && "SourceManager must be initialized"); -// if (auto FileEntryRef = SM->getFileEntryRefForID(ID)) -// return getUnifiedPath(*FileEntryRef); -// return std::nullopt; -// } -// std::optional -// DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { + llvm::SmallString<512> FilePathAbs(File.getName()); + SM->getFileManager().makeAbsolutePath(FilePathAbs); + llvm::outs() << " Unifiedpath is " << FilePathAbs.c_str() << "\n"; + return clang::tooling::UnifiedPath(FilePathAbs); +} + +std::optional +DpctGlobalInfo::getUnifiedPath(FileID ID) { + assert(SM && "SourceManager must be initialized"); + if (auto FileEntryRef = SM->getFileEntryRefForID(ID)) + return getUnifiedPath(*FileEntryRef); + return std::nullopt; +} +std::optional +DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { if (auto RealPath = File.getFileEntry().tryGetRealPathName(); !RealPath.empty()) return clang::tooling::UnifiedPath(RealPath); diff --git a/clang/lib/Tooling/CommonOptionsParser.cpp b/clang/lib/Tooling/CommonOptionsParser.cpp index aab7926ca5e4..862561f070f6 100644 --- a/clang/lib/Tooling/CommonOptionsParser.cpp +++ b/clang/lib/Tooling/CommonOptionsParser.cpp @@ -237,7 +237,7 @@ llvm::Error CommonOptionsParser::init( #ifdef SYCLomatic_CUSTOMIZATION Compilations = CompilationDatabase::autoDetectFromDirectory( BuildPath, ErrorMessage, ErrCode, CompilationsDir); - clang::tooling::FormatSearchPath = BuildPath; + clang::tooling::FormatSearchPath = BuildPath; //Path #else Compilations = CompilationDatabase::autoDetectFromDirectory( BuildPath, ErrorMessage, ErrCode); From 2d758b1a420607d37be91bdb2ba1f67e497b4d02 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 27 Jun 2024 10:59:09 +0800 Subject: [PATCH 12/20] Update test Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 30 ++++++----------------------- clang/lib/DPCT/AnalysisInfo.h | 9 ++------- clang/lib/DPCT/InclusionHeaders.cpp | 3 +-- clang/lib/DPCT/MigrationAction.cpp | 2 +- clang/lib/DPCT/SaveNewFiles.cpp | 9 ++++----- 5 files changed, 14 insertions(+), 39 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 4dee6dabd6dc..5f5572d6339c 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1439,21 +1439,6 @@ std::string DpctGlobalInfo::getStringForRegexReplacement(StringRef MatchedStr) { return MatchedStr.str(); } } -std::optional -DpctGlobalInfo::getAbsolutePath(FileID ID) { - assert(SM && "SourceManager must be initialized"); - if (auto FileEntryRef = SM->getFileEntryRefForID(ID)) - return getAbsolutePath(*FileEntryRef); - return std::nullopt; -} -std::optional -DpctGlobalInfo::getAbsolutePath(FileEntryRef File) { - llvm::SmallString<512> FilePathAbs(File.getName()); - SM->getFileManager().makeAbsolutePath(FilePathAbs); - llvm::outs() << " Unifiedpath is " << FilePathAbs.c_str() << "\n"; - return clang::tooling::UnifiedPath(FilePathAbs); -} - std::optional DpctGlobalInfo::getUnifiedPath(FileID ID) { assert(SM && "SourceManager must be initialized"); @@ -1463,22 +1448,19 @@ DpctGlobalInfo::getUnifiedPath(FileID ID) { } std::optional DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { - if (auto RealPath = File.getFileEntry().tryGetRealPathName(); - !RealPath.empty()) - return clang::tooling::UnifiedPath(RealPath); - - llvm::SmallString<512> FilePathAbs(File.getName()); - SM->getFileManager().makeAbsolutePath(FilePathAbs); - return clang::tooling::UnifiedPath(FilePathAbs); + if (!File.getName().empty()) { + return clang::tooling::UnifiedPath(File.getName()); + } + return clang::tooling::UnifiedPath("."); } + std::pair DpctGlobalInfo::getLocInfo(SourceLocation Loc, bool *IsInvalid) { if (SM->isMacroArgExpansion(Loc)) { Loc = SM->getImmediateSpellingLoc(Loc); } auto LocInfo = SM->getDecomposedLoc(SM->getExpansionLoc(Loc)); - auto AbsPath = getAbsolutePath(LocInfo.first); - // auto AbsPath = getUnifiedPath(LocInfo.first); + auto AbsPath = getUnifiedPath(LocInfo.first); if (AbsPath) return std::make_pair(AbsPath.value(), LocInfo.second); if (IsInvalid) diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index c57525886006..07b264f509e2 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1039,14 +1039,9 @@ class DpctGlobalInfo { getLocInfo(const TypeLoc &TL, bool *IsInvalid = nullptr /*out*/) { return getLocInfo(TL.getBeginLoc(), IsInvalid); } - // Return the absolute path of \p ID - static std::optional getAbsolutePath(FileID ID); - // Return the absolute path of \p File - static std::optional - getAbsolutePath(FileEntryRef File); - + // Return the Unified path of \p ID static std::optional getUnifiedPath(FileID ID); - // Return the unified path of \p File + // Return the Unified path of \p File static std::optional getUnifiedPath(FileEntryRef File); diff --git a/clang/lib/DPCT/InclusionHeaders.cpp b/clang/lib/DPCT/InclusionHeaders.cpp index 0b756c5db3c7..a23114b03be0 100644 --- a/clang/lib/DPCT/InclusionHeaders.cpp +++ b/clang/lib/DPCT/InclusionHeaders.cpp @@ -132,8 +132,7 @@ void IncludesCallbacks::InclusionDirective( LastInclusionLocationUpdater Updater(FileInfo, FilenameRange.getEnd()); clang::tooling::UnifiedPath IncludedFile; - // if (auto OptionalAbs = Global.getUnifiedPath(*File)) - if (auto OptionalAbs = Global.getAbsolutePath(*File)) + if (auto OptionalAbs = Global.getUnifiedPath(*File)) IncludedFile = OptionalAbs.value(); if (Global.isExcluded(IncludedFile)) diff --git a/clang/lib/DPCT/MigrationAction.cpp b/clang/lib/DPCT/MigrationAction.cpp index 189b750dcb23..e7f17c43e079 100644 --- a/clang/lib/DPCT/MigrationAction.cpp +++ b/clang/lib/DPCT/MigrationAction.cpp @@ -96,7 +96,7 @@ void DpctConsumer::Initialize(ASTContext &Context) { DpctGlobalInfo::setContext(Context); auto &SM = Context.getSourceManager(); auto Path = DpctGlobalInfo::getUnifiedPath(SM.getMainFileID()); - assert(Path && "Can not find absolute path"); + assert(Path && "Can not find Unified path"); DpctGlobalInfo::getInstance().setMainFile( Info->MainFile = DpctGlobalInfo::getInstance().insertFile(Path.value())); } diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index ece6fb642166..da62dd092f79 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -196,7 +196,7 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, std::filesystem::path(AbsolutePath.str())) .string(); - // The code is to create symbol file and link them to the target file. + // The code is to create symbol file and link the file to target file. // The code will iterate to get the parent path of the // absolute path of the file. Then create the target directory if the // canonical path is not exists in the out root path. @@ -213,7 +213,6 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, auto AbsoluteParentPath = sys::path::parent_path(AbsolutePath); if (llvm::sys::fs::is_symlink_file(AbsoluteParentPath)) { - llvm::outs() << "SYmbol link is " << AbsoluteParentPath << "\n"; createSymLink(AbsoluteParentPath, InRoot, OutRoot); } auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath); @@ -1146,9 +1145,9 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, DebugFilePath.getCanonicalPath()); } } - // createSymLink(Entry.first, InRoot, SYCLMigratedOutRoot); - // if (dpct::DpctGlobalInfo::isCodePinEnabled()) - // createSymLink(Entry.first, InRoot, CodePinCUDAFolder); + createSymLink(Entry.first, InRoot, SYCLMigratedOutRoot); + if (dpct::DpctGlobalInfo::isCodePinEnabled()) + createSymLink(Entry.first, InRoot, CodePinCUDAFolder); } std::string ScriptFineName = "Makefile.dpct"; From 3409fdc07244565dcab5e0985d31cd5c99cc4557 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 27 Jun 2024 11:00:47 +0800 Subject: [PATCH 13/20] remove the comment. Signed-off-by: Chen, Sheng S --- clang/lib/Tooling/CommonOptionsParser.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Tooling/CommonOptionsParser.cpp b/clang/lib/Tooling/CommonOptionsParser.cpp index 862561f070f6..aab7926ca5e4 100644 --- a/clang/lib/Tooling/CommonOptionsParser.cpp +++ b/clang/lib/Tooling/CommonOptionsParser.cpp @@ -237,7 +237,7 @@ llvm::Error CommonOptionsParser::init( #ifdef SYCLomatic_CUSTOMIZATION Compilations = CompilationDatabase::autoDetectFromDirectory( BuildPath, ErrorMessage, ErrCode, CompilationsDir); - clang::tooling::FormatSearchPath = BuildPath; //Path + clang::tooling::FormatSearchPath = BuildPath; #else Compilations = CompilationDatabase::autoDetectFromDirectory( BuildPath, ErrorMessage, ErrCode); From 727adff8dd019ee6ba92512fc4ac0532892732c0 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 27 Jun 2024 12:48:31 +0800 Subject: [PATCH 14/20] test Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 5f5572d6339c..d883c7d394ce 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1449,6 +1449,7 @@ DpctGlobalInfo::getUnifiedPath(FileID ID) { std::optional DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { if (!File.getName().empty()) { + llvm::outs() << " File name " << File.getName() << "\n"; return clang::tooling::UnifiedPath(File.getName()); } return clang::tooling::UnifiedPath("."); From 79316d1fe99823815008e20f499af727b121811a Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 27 Jun 2024 14:53:31 +0800 Subject: [PATCH 15/20] fix. Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index d883c7d394ce..bc210a215b1e 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1448,11 +1448,13 @@ DpctGlobalInfo::getUnifiedPath(FileID ID) { } std::optional DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { - if (!File.getName().empty()) { - llvm::outs() << " File name " << File.getName() << "\n"; + if (!File.getName().empty() && llvm::sys::fs::exists(File.getName())) { return clang::tooling::UnifiedPath(File.getName()); + } else if (auto RealPath = File.getFileEntry().tryGetRealPathName(); + !RealPath.empty()) { + return clang::tooling::UnifiedPath(RealPath); } - return clang::tooling::UnifiedPath("."); + return std::nullopt; } std::pair From 7815c7bb5ec0f0a07148488604bfbfb73ff65b0b Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 2 Jul 2024 14:31:46 +0800 Subject: [PATCH 16/20] update Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index bc210a215b1e..c663381a302b 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1448,13 +1448,16 @@ DpctGlobalInfo::getUnifiedPath(FileID ID) { } std::optional DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { - if (!File.getName().empty() && llvm::sys::fs::exists(File.getName())) { - return clang::tooling::UnifiedPath(File.getName()); - } else if (auto RealPath = File.getFileEntry().tryGetRealPathName(); - !RealPath.empty()) { - return clang::tooling::UnifiedPath(RealPath); - } - return std::nullopt; + llvm::SmallString<512> FilePathAbs(File.getName()); + SM->getFileManager().makeAbsolutePath(FilePathAbs); + auto UnifiedFilePath = clang::tooling::UnifiedPath(FilePathAbs); + if (!llvm::sys::fs::exists(UnifiedFilePath.getCanonicalPath())) { + if (auto RealPath = File.getFileEntry().tryGetRealPathName(); + !RealPath.empty()) + return clang::tooling::UnifiedPath(RealPath); + return std::nullopt; + } + return UnifiedFilePath; } std::pair From b262b77c942e9d62a9b38f4473cfc0604410dc74 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 2 Jul 2024 15:39:22 +0800 Subject: [PATCH 17/20] format Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index c663381a302b..7171ac1adedb 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1455,7 +1455,7 @@ DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { if (auto RealPath = File.getFileEntry().tryGetRealPathName(); !RealPath.empty()) return clang::tooling::UnifiedPath(RealPath); - return std::nullopt; + return std::nullopt; } return UnifiedFilePath; } From 34befd98b175b92e0925a7945e57391eefc9db2b Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 8 Jul 2024 22:50:57 +0800 Subject: [PATCH 18/20] update Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 8 ++-- clang/lib/DPCT/AnalysisInfo.h | 8 ++-- clang/lib/DPCT/InclusionHeaders.cpp | 2 +- clang/lib/DPCT/MigrationAction.cpp | 4 +- clang/lib/DPCT/SaveNewFiles.cpp | 40 ++++++++++++++----- clang/lib/Tooling/Tooling.cpp | 8 ++-- .../soft_link_dir/vector_add_format.cu | 2 +- .../soft_link_file/vector_add_format.cu | 2 +- .../soft_link_file_codepin/link/test/temp | 0 .../soft_link_file_codepin/link/test/test.hpp | 1 - .../target/test/test.hpp | 5 --- .../soft_link_file_rename/link/test/test.cuh | 1 + .../target/test/test.cuh | 6 +++ .../vector_add_format.cu | 11 ++--- .../soft_link_mul_dir2/vector_add_format.cu | 2 +- 15 files changed, 63 insertions(+), 37 deletions(-) delete mode 100644 clang/test/dpct/soft_link/soft_link_file_codepin/link/test/temp delete mode 120000 clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp delete mode 100644 clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp create mode 120000 clang/test/dpct/soft_link/soft_link_file_rename/link/test/test.cuh create mode 100644 clang/test/dpct/soft_link/soft_link_file_rename/target/test/test.cuh rename clang/test/dpct/soft_link/{soft_link_file_codepin => soft_link_file_rename}/vector_add_format.cu (69%) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 7171ac1adedb..2452cf89ac5d 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1440,14 +1440,14 @@ std::string DpctGlobalInfo::getStringForRegexReplacement(StringRef MatchedStr) { } } std::optional -DpctGlobalInfo::getUnifiedPath(FileID ID) { +DpctGlobalInfo::getAbsolutePath(FileID ID) { assert(SM && "SourceManager must be initialized"); if (auto FileEntryRef = SM->getFileEntryRefForID(ID)) - return getUnifiedPath(*FileEntryRef); + return getAbsolutePath(*FileEntryRef); return std::nullopt; } std::optional -DpctGlobalInfo::getUnifiedPath(FileEntryRef File) { +DpctGlobalInfo::getAbsolutePath(FileEntryRef File) { llvm::SmallString<512> FilePathAbs(File.getName()); SM->getFileManager().makeAbsolutePath(FilePathAbs); auto UnifiedFilePath = clang::tooling::UnifiedPath(FilePathAbs); @@ -1466,7 +1466,7 @@ DpctGlobalInfo::getLocInfo(SourceLocation Loc, bool *IsInvalid) { Loc = SM->getImmediateSpellingLoc(Loc); } auto LocInfo = SM->getDecomposedLoc(SM->getExpansionLoc(Loc)); - auto AbsPath = getUnifiedPath(LocInfo.first); + auto AbsPath = getAbsolutePath(LocInfo.first); if (AbsPath) return std::make_pair(AbsPath.value(), LocInfo.second); if (IsInvalid) diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 07b264f509e2..c757d2515b3d 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1039,11 +1039,11 @@ class DpctGlobalInfo { getLocInfo(const TypeLoc &TL, bool *IsInvalid = nullptr /*out*/) { return getLocInfo(TL.getBeginLoc(), IsInvalid); } - // Return the Unified path of \p ID - static std::optional getUnifiedPath(FileID ID); - // Return the Unified path of \p File + // Return the absolute path of \p ID + static std::optional getAbsolutePath(FileID ID); + // Return the absolute path of \p File static std::optional - getUnifiedPath(FileEntryRef File); + getAbsolutePath(FileEntryRef File); static std::pair getLocInfo(SourceLocation Loc, bool *IsInvalid = nullptr /* out */); diff --git a/clang/lib/DPCT/InclusionHeaders.cpp b/clang/lib/DPCT/InclusionHeaders.cpp index a23114b03be0..71ce7c1cebab 100644 --- a/clang/lib/DPCT/InclusionHeaders.cpp +++ b/clang/lib/DPCT/InclusionHeaders.cpp @@ -132,7 +132,7 @@ void IncludesCallbacks::InclusionDirective( LastInclusionLocationUpdater Updater(FileInfo, FilenameRange.getEnd()); clang::tooling::UnifiedPath IncludedFile; - if (auto OptionalAbs = Global.getUnifiedPath(*File)) + if (auto OptionalAbs = Global.getAbsolutePath(*File)) IncludedFile = OptionalAbs.value(); if (Global.isExcluded(IncludedFile)) diff --git a/clang/lib/DPCT/MigrationAction.cpp b/clang/lib/DPCT/MigrationAction.cpp index e7f17c43e079..8e4bb0e5cc2e 100644 --- a/clang/lib/DPCT/MigrationAction.cpp +++ b/clang/lib/DPCT/MigrationAction.cpp @@ -95,8 +95,8 @@ void DpctConsumer::Initialize(ASTContext &Context) { // Set Context for build information DpctGlobalInfo::setContext(Context); auto &SM = Context.getSourceManager(); - auto Path = DpctGlobalInfo::getUnifiedPath(SM.getMainFileID()); - assert(Path && "Can not find Unified path"); + auto Path = DpctGlobalInfo::getAbsolutePath(SM.getMainFileID()); + assert(Path && "Can not find absolute path"); DpctGlobalInfo::getInstance().setMainFile( Info->MainFile = DpctGlobalInfo::getInstance().insertFile(Path.value())); } diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index da62dd092f79..1641b8270173 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -184,18 +184,22 @@ bool rewriteAbsoluteDir(clang::tooling::UnifiedPath &FilePath, void createSymLink(const clang::tooling::UnifiedPath &FilePath, const clang::tooling::UnifiedPath &InRoot, - const clang::tooling::UnifiedPath &OutRoot) { + const clang::tooling::UnifiedPath &OutRoot, + bool RewriteFileName) { if (llvm::sys::fs::exists(FilePath.getCanonicalPath())) { auto SourcePath = FilePath; while ( isChildPath(InRoot.getAbsolutePath(), SourcePath.getAbsolutePath())) { auto AbsolutePath = SourcePath.getAbsolutePath(); if (llvm::sys::fs::is_symlink_file(AbsolutePath)) { - std::filesystem::path RealPathLink = + tooling::UnifiedPath SymbolPath = std::filesystem::read_symlink( std::filesystem::path(AbsolutePath.str())) .string(); - + if (RewriteFileName) { + rewriteFileName(SymbolPath); + rewriteFileName(SourcePath); + } // The code is to create symbol file and link the file to target file. // The code will iterate to get the parent path of the // absolute path of the file. Then create the target directory if the @@ -213,7 +217,7 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, auto AbsoluteParentPath = sys::path::parent_path(AbsolutePath); if (llvm::sys::fs::is_symlink_file(AbsoluteParentPath)) { - createSymLink(AbsoluteParentPath, InRoot, OutRoot); + createSymLink(AbsoluteParentPath, InRoot, OutRoot, RewriteFileName); } auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath); rewriteAbsoluteDir(ParentPath, InRoot, OutRoot); @@ -222,7 +226,7 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, } try { std::filesystem::create_symlink( - std::filesystem::path(RealPathLink), + std::filesystem::path(SymbolPath.getPath().str()), std::filesystem::path(ReplPath.getAbsolutePath().str())); } catch (const std::filesystem::filesystem_error &e) { std::cerr << "Error creating symbolic link: " << e.what() @@ -230,6 +234,9 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, } } SourcePath = path::parent_path(AbsolutePath); + if (!isChildOrSamePath(OutRoot, SourcePath)) { + return; + } } } } @@ -303,6 +310,25 @@ static bool checkOverwriteAndWarn(StringRef OutFilePath, StringRef InFilePath) { void processallOptionAction(clang::tooling::UnifiedPath &InRoot, clang::tooling::UnifiedPath &OutRoot, bool IsForSYCL) { + extern DpctOption ProcessAll; + if (ProcessAll) { + std::error_code EC; + for (fs::recursive_directory_iterator Iter(Twine(InRoot.getPath()), EC), + End; + Iter != End; Iter.increment(EC)) { + if (Iter->type() == fs::file_type::symlink_file) { + + tooling::UnifiedPath RootSource = Iter->path(); + if (IncludeFleMap.find(Iter->path()) != IncludeFileMap.end()) { + if (IncludeFileMap[RootSource]) { + rewriteFileName(RootSource); + } + } + createSymLink(Iter->path(), InRoot, OutRoot, + IncludeFileMap[Iter->path()]); + } + } + } for (const auto &File : FilesNotInCompilationDB) { if (IncludeFileMap.find(File) != IncludeFileMap.end()) { // Skip the files parsed by dpct parser. @@ -355,7 +381,6 @@ void processAllFiles(StringRef InRoot, StringRef OutRoot, } auto FilePath = Iter->path(); - // Skip output directory if it is in the in-root directory. if (isChildOrSamePath(OutRoot.str(), FilePath)) continue; @@ -1145,9 +1170,6 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, DebugFilePath.getCanonicalPath()); } } - createSymLink(Entry.first, InRoot, SYCLMigratedOutRoot); - if (dpct::DpctGlobalInfo::isCodePinEnabled()) - createSymLink(Entry.first, InRoot, CodePinCUDAFolder); } std::string ScriptFineName = "Makefile.dpct"; diff --git a/clang/lib/Tooling/Tooling.cpp b/clang/lib/Tooling/Tooling.cpp index 54b31b51a81d..9c918dcf5245 100644 --- a/clang/lib/Tooling/Tooling.cpp +++ b/clang/lib/Tooling/Tooling.cpp @@ -1120,15 +1120,17 @@ int ClangTool::run(ToolAction *Action) { // if input file(s) is not specified in command line, and the process-all // option is given in the comomand line, dpct tries to migrate or copy all // files from -in-root to the output directory. - if(SourcePaths.size() == 0 && DoGetRunRound() == 0) { + // if(SourcePaths.size() == 0 && DoGetRunRound() == 0) { + if(DoGetRunRound() == 0) { std::vector FilesNotProcessed; - + llvm::outs() << "HHHHHHHHHHHHHHHHH \n"; // To traverse all the files in the directory specified by // -in-root, collecting *.cu files not processed by the first loop of // calling processFiles() into FilesNotProcessed, and copies the rest // files to the output directory. - DoFileProcessHandle(FilesNotProcessed); + DoFileProcessHandle(FilesNotProcessed);// Compilation. for (auto &Entry : FilesNotProcessed) { + llvm::outs() << "VVVVV " << Entry << "\n"; auto File = llvm::StringRef(Entry); int Ret = processFilesWithCrashGuard(this, File, ProcessingFailed, diff --git a/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu index c954afb0bc44..f884121a8b0c 100644 --- a/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu +++ b/clang/test/dpct/soft_link/soft_link_dir/vector_add_format.cu @@ -1,5 +1,5 @@ // RUN: cd %S/link && rm -rf test && ln -nfs ../target test -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/test -x cuda --cuda-host-only +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --process-all -- -I %S/link/test -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/test/test/test.hpp --match-full-lines %S/link/test/test/test.hpp // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/test %} diff --git a/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu index c553b6871615..657d00af0bce 100644 --- a/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu +++ b/clang/test/dpct/soft_link/soft_link_file/vector_add_format.cu @@ -1,5 +1,5 @@ // RUN: cd %S/link/test && rm test.hpp && ln -nfs ../../target/test/test.hpp test.hpp -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link -x cuda --cuda-host-only +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --process-all -- -I %S/link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/test/test.hpp --match-full-lines %S/link/test/test.hpp // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %} diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/temp b/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/temp deleted file mode 100644 index e69de29bb2d1..000000000000 diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp deleted file mode 120000 index 6c8ad69f6c14..000000000000 --- a/clang/test/dpct/soft_link/soft_link_file_codepin/link/test/test.hpp +++ /dev/null @@ -1 +0,0 @@ -../../target/test/test.hpp \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp b/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp deleted file mode 100644 index c56c85ed7a4c..000000000000 --- a/clang/test/dpct/soft_link/soft_link_file_codepin/target/test/test.hpp +++ /dev/null @@ -1,5 +0,0 @@ -// CHECK: #include -#include -int test() { - return 0; -} diff --git a/clang/test/dpct/soft_link/soft_link_file_rename/link/test/test.cuh b/clang/test/dpct/soft_link/soft_link_file_rename/link/test/test.cuh new file mode 120000 index 000000000000..d6ec8088d6da --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_rename/link/test/test.cuh @@ -0,0 +1 @@ +../../target/test/test.cuh \ No newline at end of file diff --git a/clang/test/dpct/soft_link/soft_link_file_rename/target/test/test.cuh b/clang/test/dpct/soft_link/soft_link_file_rename/target/test/test.cuh new file mode 100644 index 000000000000..cdf9d602ebaa --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_rename/target/test/test.cuh @@ -0,0 +1,6 @@ +// CHECK: #include +// CHECK: #include +#include +int test() { + return 0; +} diff --git a/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_file_rename/vector_add_format.cu similarity index 69% rename from clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu rename to clang/test/dpct/soft_link/soft_link_file_rename/vector_add_format.cu index c7a48c5da802..13e80231b645 100644 --- a/clang/test/dpct/soft_link/soft_link_file_codepin/vector_add_format.cu +++ b/clang/test/dpct/soft_link/soft_link_file_rename/vector_add_format.cu @@ -1,11 +1,12 @@ -// RUN: cd %S/link/test && rm -rf test.hpp && ln -nfs ../../target/test/test.hpp test.hpp -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --enable-codepin -- -I %S/link -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/out_codepin_cuda/link/test/test.hpp --match-full-lines %S/link/test/test.hpp -// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out_codepin_sycl/vector_add_format.dp.cpp -o %T/out_codepin_sycl/vector_add_format.dp.o -I %T/out_codepin_sycl/link %} +// RUN: cd %S/link/test && rm test.cuh && ln -nfs ../../target/test/test.cuh test.cuh +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --process-all -- -I %S/link -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s +// RUN: FileCheck --input-file %T/out/link/test/test.dp.hpp --match-full-lines %S/link/test/test.cuh +// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link %} #include #include -#include "test/test.hpp" +#include "test/test.cuh" #define VECTOR_SIZE 256 __global__ void VectorAddKernel(float* A, float* B, float* C) diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu index 2a4b620d151e..26fe471e71b9 100644 --- a/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu +++ b/clang/test/dpct/soft_link/soft_link_mul_dir2/vector_add_format.cu @@ -1,5 +1,5 @@ // RUN: cd %S/link/hello && rm -rf target_soft_link && ln -nfs ../../target target_soft_link && cd %S/link && rm -rf link && ln -nfs hello link -// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only +// RUN: dpct --in-root=%S --out-root=%T/out %s --cuda-include-path="%cuda-path/include" --process-all -- -I %S/link/link/target_soft_link -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/out/vector_add_format.dp.cpp --match-full-lines %s // RUN: FileCheck --input-file %T/out/link/link/target_soft_link/test/test.hpp --match-full-lines %S/target/test/test.hpp // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/out/vector_add_format.dp.cpp -o %T/out/vector_add_format.dp.o -I %T/out/link/hello/target_soft_link %} From 7b970de9a1ea00c45c399378669fd0b5534b792c Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 9 Jul 2024 09:12:34 +0800 Subject: [PATCH 19/20] update Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/SaveNewFiles.cpp | 9 ++++++--- clang/lib/Tooling/Core/UnifiedPath.cpp | 2 +- clang/lib/Tooling/Tooling.cpp | 7 ++----- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index 1641b8270173..10e18eb30102 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -316,10 +316,13 @@ void processallOptionAction(clang::tooling::UnifiedPath &InRoot, for (fs::recursive_directory_iterator Iter(Twine(InRoot.getPath()), EC), End; Iter != End; Iter.increment(EC)) { + // Skip output directory if it is in the in-root directory. + if (isChildOrSamePath(OutRoot.str(), FilePath)) + continue; if (Iter->type() == fs::file_type::symlink_file) { tooling::UnifiedPath RootSource = Iter->path(); - if (IncludeFleMap.find(Iter->path()) != IncludeFileMap.end()) { + if (IncludeFileMap.find(Iter->path()) != IncludeFileMap.end()) { if (IncludeFileMap[RootSource]) { rewriteFileName(RootSource); } @@ -381,6 +384,7 @@ void processAllFiles(StringRef InRoot, StringRef OutRoot, } auto FilePath = Iter->path(); + // Skip output directory if it is in the in-root directory. if (isChildOrSamePath(OutRoot.str(), FilePath)) continue; @@ -1084,8 +1088,7 @@ int saveNewFiles(clang::tooling::RefactoringTool &Tool, 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_codepin. E.g., - // /path/to/outroot_codepin/a.cu + // 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; diff --git a/clang/lib/Tooling/Core/UnifiedPath.cpp b/clang/lib/Tooling/Core/UnifiedPath.cpp index b02080ee7859..313a1f96536a 100644 --- a/clang/lib/Tooling/Core/UnifiedPath.cpp +++ b/clang/lib/Tooling/Core/UnifiedPath.cpp @@ -46,7 +46,7 @@ void UnifiedPath::makeCanonical(const std::string &CWD) { llvm::sys::path::remove_dots(Path, /* remove_dot_dot= */ true); llvm::sys::path::native(Path); - _AbsolutePath = Path.str(); + llvm::SmallString<512> RealPath; // We need make sure the input `Path` for llvm::sys::fs::real_path is // exsiting, or else the behavior of real_path() is unexpected. diff --git a/clang/lib/Tooling/Tooling.cpp b/clang/lib/Tooling/Tooling.cpp index 9c918dcf5245..8eb52349c55e 100644 --- a/clang/lib/Tooling/Tooling.cpp +++ b/clang/lib/Tooling/Tooling.cpp @@ -1120,17 +1120,14 @@ int ClangTool::run(ToolAction *Action) { // if input file(s) is not specified in command line, and the process-all // option is given in the comomand line, dpct tries to migrate or copy all // files from -in-root to the output directory. - // if(SourcePaths.size() == 0 && DoGetRunRound() == 0) { - if(DoGetRunRound() == 0) { + if (SourcePaths.size() == 0 && DoGetRunRound() == 0) { std::vector FilesNotProcessed; - llvm::outs() << "HHHHHHHHHHHHHHHHH \n"; // To traverse all the files in the directory specified by // -in-root, collecting *.cu files not processed by the first loop of // calling processFiles() into FilesNotProcessed, and copies the rest // files to the output directory. - DoFileProcessHandle(FilesNotProcessed);// Compilation. + DoFileProcessHandle(FilesNotProcessed); for (auto &Entry : FilesNotProcessed) { - llvm::outs() << "VVVVV " << Entry << "\n"; auto File = llvm::StringRef(Entry); int Ret = processFilesWithCrashGuard(this, File, ProcessingFailed, From 7edec150ecc2c995059a8d1757a04ccc5ab21e5d Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 9 Jul 2024 10:20:04 +0800 Subject: [PATCH 20/20] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 12 ++++-------- clang/lib/DPCT/AnalysisInfo.h | 1 - clang/lib/DPCT/SaveNewFiles.cpp | 12 +----------- clang/test/dpct/soft_link/soft_link_dir/link/temp | 0 .../soft_link/soft_link_mul_dir2/link/hello/temp | 0 5 files changed, 5 insertions(+), 20 deletions(-) delete mode 100644 clang/test/dpct/soft_link/soft_link_dir/link/temp delete mode 100644 clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/temp diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 2452cf89ac5d..48415fb1e288 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1448,16 +1448,12 @@ DpctGlobalInfo::getAbsolutePath(FileID ID) { } std::optional DpctGlobalInfo::getAbsolutePath(FileEntryRef File) { + if (auto RealPath = File.getFileEntry().tryGetRealPathName(); + !RealPath.empty()) + return clang::tooling::UnifiedPath(RealPath); llvm::SmallString<512> FilePathAbs(File.getName()); SM->getFileManager().makeAbsolutePath(FilePathAbs); - auto UnifiedFilePath = clang::tooling::UnifiedPath(FilePathAbs); - if (!llvm::sys::fs::exists(UnifiedFilePath.getCanonicalPath())) { - if (auto RealPath = File.getFileEntry().tryGetRealPathName(); - !RealPath.empty()) - return clang::tooling::UnifiedPath(RealPath); - return std::nullopt; - } - return UnifiedFilePath; + return clang::tooling::UnifiedPath(FilePathAbs); } std::pair diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index c757d2515b3d..ee1e787318ca 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1044,7 +1044,6 @@ class DpctGlobalInfo { // Return the absolute path of \p File static std::optional getAbsolutePath(FileEntryRef File); - static std::pair getLocInfo(SourceLocation Loc, bool *IsInvalid = nullptr /* out */); static std::string getTypeName(QualType QT, const ASTContext &Context); diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index 10e18eb30102..cd9d3e69930c 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -205,19 +205,11 @@ void createSymLink(const clang::tooling::UnifiedPath &FilePath, // absolute path of the file. Then create the target directory if the // canonical path is not exists in the out root path. auto ReplPath = SourcePath; - auto ReplCanonicalPath = - tooling::UnifiedPath(SourcePath.getCanonicalPath()); rewriteAbsoluteDir(ReplPath, InRoot, OutRoot); - rewriteDir(ReplCanonicalPath, InRoot, OutRoot); - if (llvm::sys::fs::is_directory(SourcePath.getAbsolutePath())) { - if (!llvm::sys::fs::exists(ReplCanonicalPath.getCanonicalPath())) - createDirectories(ReplCanonicalPath.getCanonicalPath()); - } auto AbsoluteParentPath = sys::path::parent_path(AbsolutePath); - if (llvm::sys::fs::is_symlink_file(AbsoluteParentPath)) { - createSymLink(AbsoluteParentPath, InRoot, OutRoot, RewriteFileName); + createSymLink(AbsoluteParentPath, InRoot, OutRoot, false); } auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath); rewriteAbsoluteDir(ParentPath, InRoot, OutRoot); @@ -317,8 +309,6 @@ void processallOptionAction(clang::tooling::UnifiedPath &InRoot, End; Iter != End; Iter.increment(EC)) { // Skip output directory if it is in the in-root directory. - if (isChildOrSamePath(OutRoot.str(), FilePath)) - continue; if (Iter->type() == fs::file_type::symlink_file) { tooling::UnifiedPath RootSource = Iter->path(); diff --git a/clang/test/dpct/soft_link/soft_link_dir/link/temp b/clang/test/dpct/soft_link/soft_link_dir/link/temp deleted file mode 100644 index e69de29bb2d1..000000000000 diff --git a/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/temp b/clang/test/dpct/soft_link/soft_link_mul_dir2/link/hello/temp deleted file mode 100644 index e69de29bb2d1..000000000000