diff --git a/clang/include/clang/Tooling/Core/UnifiedPath.h b/clang/include/clang/Tooling/Core/UnifiedPath.h index 67a9e64310bc..894e4986329d 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,21 @@ 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 c56015bd3e69..48415fb1e288 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -1451,11 +1451,11 @@ 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); return clang::tooling::UnifiedPath(FilePathAbs); } + std::pair DpctGlobalInfo::getLocInfo(SourceLocation Loc, bool *IsInvalid) { if (SM->isMacroArgExpansion(Loc)) { diff --git a/clang/lib/DPCT/SaveNewFiles.cpp b/clang/lib/DPCT/SaveNewFiles.cpp index 1be73240c0dc..cd9d3e69930c 100644 --- a/clang/lib/DPCT/SaveNewFiles.cpp +++ b/clang/lib/DPCT/SaveNewFiles.cpp @@ -36,6 +36,7 @@ #include #include +#include #include #include @@ -150,6 +151,87 @@ 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, + 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)) { + 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 + // canonical path is not exists in the out root path. + auto ReplPath = SourcePath; + rewriteAbsoluteDir(ReplPath, InRoot, OutRoot); + + auto AbsoluteParentPath = sys::path::parent_path(AbsolutePath); + if (llvm::sys::fs::is_symlink_file(AbsoluteParentPath)) { + createSymLink(AbsoluteParentPath, InRoot, OutRoot, false); + } + auto ParentPath = tooling::UnifiedPath(AbsoluteParentPath); + rewriteAbsoluteDir(ParentPath, InRoot, OutRoot); + if (!llvm::sys::fs::exists(ParentPath.getAbsolutePath())) { + createDirectories(ParentPath.getAbsolutePath()); + } + try { + std::filesystem::create_symlink( + 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() + << std::endl; + } + } + SourcePath = path::parent_path(AbsolutePath); + if (!isChildOrSamePath(OutRoot, SourcePath)) { + return; + } + } + } +} void rewriteFileName(clang::tooling::UnifiedPath &FileName) { rewriteFileName(FileName, FileName); } @@ -220,6 +302,26 @@ 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)) { + // Skip output directory if it is in the in-root directory. + if (Iter->type() == fs::file_type::symlink_file) { + + tooling::UnifiedPath RootSource = Iter->path(); + if (IncludeFileMap.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. diff --git a/clang/lib/Tooling/Core/UnifiedPath.cpp b/clang/lib/Tooling/Core/UnifiedPath.cpp index 12d04cc3edb7..313a1f96536a 100644 --- a/clang/lib/Tooling/Core/UnifiedPath.cpp +++ b/clang/lib/Tooling/Core/UnifiedPath.cpp @@ -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/lib/Tooling/Tooling.cpp b/clang/lib/Tooling/Tooling.cpp index 54b31b51a81d..8eb52349c55e 100644 --- a/clang/lib/Tooling/Tooling.cpp +++ b/clang/lib/Tooling/Tooling.cpp @@ -1120,9 +1120,8 @@ 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) { std::vector FilesNotProcessed; - // 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 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..f884121a8b0c --- /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" --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 %} + +#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..657d00af0bce --- /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" --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 %} + +#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_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_rename/vector_add_format.cu b/clang/test/dpct/soft_link/soft_link_file_rename/vector_add_format.cu new file mode 100644 index 000000000000..13e80231b645 --- /dev/null +++ b/clang/test/dpct/soft_link/soft_link_file_rename/vector_add_format.cu @@ -0,0 +1,54 @@ +// 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.cuh" +#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..26fe471e71b9 --- /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" --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 %} + +#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; +}