From 707fbc067bf329340e6c73ef408b9d9d1364f483 Mon Sep 17 00:00:00 2001 From: "Wang, Yihan" Date: Tue, 14 May 2024 16:32:43 +0800 Subject: [PATCH] [SYCLomatic] Format Signed-off-by: Wang, Yihan --- clang/lib/DPCT/AnalysisInfo.h | 6 +- clang/lib/DPCT/CUBAPIMigration.cpp | 2 +- clang/lib/DPCT/ExprAnalysis.cpp | 12 +++- .../dpct/cub/blocklevel/blockreduce_p3.cu | 58 +++++++++---------- 4 files changed, 43 insertions(+), 35 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index c47b57bbe7c5..161abc7a9f2d 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -2144,9 +2144,11 @@ class TempStorageVarInfo { unsigned Offset; std::string Name; std::shared_ptr Type; + public: - TempStorageVarInfo(unsigned Off, StringRef Name, std::shared_ptr T) - : Offset(Off), Name(Name.str()), Type(T) {} + TempStorageVarInfo(unsigned Off, StringRef Name, + std::shared_ptr T) + : Offset(Off), Name(Name.str()), Type(T) {} const std::string &getName() const { return Name; } unsigned getOffset() const { return Offset; } void addAccessorDecl(StmtList &AccessorList, StringRef LocalSize) const; diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index cd6e3d92fce7..47f09d3befaa 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -1256,7 +1256,7 @@ void CubRule::processBlockLevelMemberCall(const CXXMemberCallExpr *BlockMC) { if (auto FuncInfo = DeviceFunctionDecl::LinkRedecls(FD)) { auto LocInfo = DpctGlobalInfo::getLocInfo(TempStorage); auto TypeInfo = std::make_shared(DataTypeLoc); - + FuncInfo->getVarMap().addCUBTempStorage( std::make_shared( LocInfo.second, TempStorage->getName(), TypeInfo)); diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index d4d8a9186c43..298721145d3a 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -1191,9 +1191,15 @@ void ExprAnalysis::analyzeType(TypeLoc TL, const Expr *CSCE, break; } case TypeLoc::SubstTemplateTypeParm: - TYPELOC_CAST(SubstTemplateTypeParmTypeLoc).getType()->getAs()->getAssociatedDecl(); - if (auto T = TYPELOC_CAST(SubstTemplateTypeParmTypeLoc).getType()->getAs()) - return addReplacement(TL.getBeginLoc(), TL.getEndLoc(), CSCE, T->getIndex()); + TYPELOC_CAST(SubstTemplateTypeParmTypeLoc) + .getType() + ->getAs() + ->getAssociatedDecl(); + if (auto T = TYPELOC_CAST(SubstTemplateTypeParmTypeLoc) + .getType() + ->getAs()) + return addReplacement(TL.getBeginLoc(), TL.getEndLoc(), CSCE, + T->getIndex()); break; case TypeLoc::TemplateTypeParm: if (auto D = TYPELOC_CAST(TemplateTypeParmTypeLoc).getDecl()) { diff --git a/clang/test/dpct/cub/blocklevel/blockreduce_p3.cu b/clang/test/dpct/cub/blocklevel/blockreduce_p3.cu index db54ee5dafc9..9a0574311380 100644 --- a/clang/test/dpct/cub/blocklevel/blockreduce_p3.cu +++ b/clang/test/dpct/cub/blocklevel/blockreduce_p3.cu @@ -72,30 +72,30 @@ __global__ void reduce_kernel_dependent1(T *da) { template void test() { - T *ha = (float *)malloc(N * sizeof(T)); - T *da; - cudaMalloc(&da, N * sizeof(T)); + T *ha = (float *)malloc(N * sizeof(T)); + T *da; + cudaMalloc(&da, N * sizeof(T)); - for (int i = 0; i < N; i++) { - ha[i] = i * 1.0f; - } + for (int i = 0; i < N; i++) { + ha[i] = i * 1.0f; + } - cudaMemcpy(da, ha, N * sizeof(T), cudaMemcpyHostToDevice); - // CHECK: void test() - // CHECK: q_ct1.submit( - // CHECK: [&](sycl::handler &cgh) { - // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size() * sizeof(T)), cgh); - // CHECK: cgh.parallel_for( - // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), - // CHECK: [=](sycl::nd_item<3> item_ct1) { - // CHECK: reduce_kernel_dependent1(da, item_ct1, temp_storage_acc); - // CHECK: }); - // CHECK: }); - reduce_kernel_dependent1<<<1, 32>>>(da); - cudaMemcpy(ha, da, 1 * sizeof(T), cudaMemcpyDeviceToHost); - std::cout << ha[0] << std::endl; - cudaFree(da); - free(ha); + cudaMemcpy(da, ha, N * sizeof(T), cudaMemcpyHostToDevice); + // CHECK: void test() + // CHECK: q_ct1.submit( + // CHECK: [&](sycl::handler &cgh) { + // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size() * sizeof(T)), cgh); + // CHECK: cgh.parallel_for( + // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), + // CHECK: [=](sycl::nd_item<3> item_ct1) { + // CHECK: reduce_kernel_dependent1(da, item_ct1, temp_storage_acc); + // CHECK: }); + // CHECK: }); + reduce_kernel_dependent1<<<1, 32>>>(da); + cudaMemcpy(ha, da, 1 * sizeof(T), cudaMemcpyDeviceToHost); + std::cout << ha[0] << std::endl; + cudaFree(da); + free(ha); } int main() { @@ -115,7 +115,7 @@ int main() { // CHECK: [&](sycl::handler &cgh) { // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size(float) * sizeof(float)), cgh); // CHECK: cgh.parallel_for( - // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), + // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), // CHECK: [=](sycl::nd_item<3> item_ct1) { // CHECK: reduce_kernel(da, item_ct1, temp_storage_acc); // CHECK: }); @@ -126,7 +126,7 @@ int main() { // CHECK: [&](sycl::handler &cgh) { // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size(float) * sizeof(float)), cgh); // CHECK: cgh.parallel_for( - // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), + // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), // CHECK: [=](sycl::nd_item<3> item_ct1) { // CHECK: reduce_kernel1(da, item_ct1, temp_storage_acc); // CHECK: }); @@ -137,7 +137,7 @@ int main() { // CHECK: [&](sycl::handler &cgh) { // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size(float) * sizeof(T)), cgh); // CHECK: cgh.parallel_for( - // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), + // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), // CHECK: [=](sycl::nd_item<3> item_ct1) { // CHECK: reduce_kernel_dependent(da, item_ct1, temp_storage_acc); // CHECK: }); @@ -148,7 +148,7 @@ int main() { // CHECK: [&](sycl::handler &cgh) { // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size(float) * sizeof(T)), cgh); // CHECK: cgh.parallel_for( - // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), + // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), // CHECK: [=](sycl::nd_item<3> item_ct1) { // CHECK: reduce_kernel_dependent1(da, item_ct1, temp_storage_acc); // CHECK: }); @@ -160,7 +160,7 @@ int main() { free(ha); } - { + { int *ha = (int *)malloc(N * sizeof(int)); int *da; cudaMalloc(&da, N * sizeof(int)); @@ -174,7 +174,7 @@ int main() { // CHECK: [&](sycl::handler &cgh) { // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size(int) * sizeof(T)), cgh); // CHECK: cgh.parallel_for( - // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), + // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), // CHECK: [=](sycl::nd_item<3> item_ct1) { // CHECK: reduce_kernel_dependent(da, item_ct1, temp_storage_acc); // CHECK: }); @@ -185,7 +185,7 @@ int main() { // CHECK: [&](sycl::handler &cgh) { // CHECK: sycl::local_accessor temp_storage_acc(sycl::range(sycl::range<3>(1, 1, 32).size(int) * sizeof(T)), cgh); // CHECK: cgh.parallel_for( - // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), + // CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 32), sycl::range<3>(1, 1, 32)), // CHECK: [=](sycl::nd_item<3> item_ct1) { // CHECK: reduce_kernel_dependent1(da, item_ct1, temp_storage_acc); // CHECK: });