From e2a868ee118a37c79914fd8663f5259e04ce41e4 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 15 Apr 2025 14:05:09 +0800 Subject: [PATCH 01/17] shared Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 1 + clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 81 ++++++++++++++++++++++- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 32 ++++++++- 3 files changed, 112 insertions(+), 2 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index f596578e1296..84f4f37b51b1 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2811,6 +2811,7 @@ void CtTypeInfo::setArrayInfo(const DependentSizedArrayTypeLoc &TL, bool NeedSizeFold) { ContainSizeofType = containSizeOfType(TL.getSizeExpr()); ExprAnalysis EA; + // EA.analyze(TL.getSizeExpr()); auto TDSI = EA.getTemplateDependentStringInfo(); if (TDSI->containsTemplateDependentMacro()) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 41aa106e5060..f13dc8dcc996 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -421,6 +421,74 @@ void StringReplacements::replaceString() { ReplMap.clear(); } +void StringReplacements::adjustSourceStrAndTDRs() { + std::vector Occurrences; + size_t Pos = 0; + while ((Pos = SourceStr.find(ConstExprExpansionInfoVDStr, Pos)) != + std::string::npos) { + Occurrences.push_back(Pos); + Pos += ConstExprExpansionInfoVDStr.length(); + } + + int LengthDiff = ConstExprExpansionInfoVDInitStr.length() - + ConstExprExpansionInfoVDStr.length(); + + for (auto Occ : Occurrences) { + const auto &Iter = TDRs.find(Occ); + if (Iter != TDRs.end()) { + TDRs.erase(Iter); + } + } + std::map> NewTDRs; + for (const auto &[Pos, Value] : TDRs) { + int NewPos = Pos; + for (auto Occ : Occurrences) { + if (Occ < static_cast(Pos)) { + NewPos += LengthDiff; + } else { + break; + } + } + auto NewValue = std::make_shared( + SourceStr, NewPos, Value->getLength(), Value->getTemplateIndex()); + NewTDRs[NewPos] = NewValue; + } + TDRs = std::move(NewTDRs); + + std::map> + NewTDRsInConstExprExpansion; + for (const auto &[Pos, Value] : TDRsInConstExprExpansion) { + for (auto Occ : Occurrences) { + int NewPos = Occ + Pos; + for (auto PrevOcc : Occurrences) { + if (PrevOcc < static_cast(Occ)) { + NewPos += LengthDiff; + } else { + break; + } + } + auto NewValue = std::make_shared( + SourceStr, NewPos, Value->getLength(), Value->getTemplateIndex()); + NewTDRsInConstExprExpansion[NewPos] = NewValue; + break; + } + } + TDRsInConstExprExpansion = std::move(NewTDRsInConstExprExpansion); + + Pos = 0; + while ((Pos = SourceStr.find(ConstExprExpansionInfoVDStr, Pos)) != + std::string::npos) { + SourceStr.replace(Pos, ConstExprExpansionInfoVDStr.length(), + ConstExprExpansionInfoVDInitStr); + Pos += ConstExprExpansionInfoVDInitStr.length(); + } + + TDRs.insert(TDRsInConstExprExpansion.begin(), TDRsInConstExprExpansion.end()); + for (const auto &[Pos, Value] : TDRs) { + Value->alterSource(SourceStr); + } +} + ExprAnalysis::ExprAnalysis(const Expr *Expression) : Context(DpctGlobalInfo::getContext()), SM(DpctGlobalInfo::getSourceManager()) { @@ -553,7 +621,18 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { } if (auto TemplateDecl = dyn_cast(DRE->getDecl())) addReplacement(DRE, TemplateDecl->getIndex()); - else if (auto ECD = dyn_cast(DRE->getDecl())) { + else if (const auto *VD = dyn_cast(DRE->getDecl()); + VD && VD->isConstexpr()) { + if (VD->getInit() && VD->getInit()->getBeginLoc().isValid() && + !ConstExprExpansion) { + std::string VDInitStr = ExprAnalysis::ref(VD->getInit()); + std::string VDStr = VD->getNameAsString(); + ReplSet.addConstExprExpansionInfo(VDStr, VDInitStr); + ConstExprExpansion = true; + dispatch(VD->getInit()); + ConstExprExpansion = false; + } + } else if (auto ECD = dyn_cast(DRE->getDecl())) { std::unordered_set targetStr = { "thread_scope_system", "thread_scope_device", "thread_scope_block", "memory_order_relaxed", "memory_order_acq_rel", "memory_order_release", diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 2576b1089a7f..4cb0eb08623b 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -77,6 +77,7 @@ class TemplateDependentReplacement { } inline size_t getOffset() const { return Offset; } inline size_t getLength() const { return Length; } + inline unsigned getTemplateIndex() const { return TemplateIndex; } const TemplateArgumentInfo & getTargetArgument(const std::vector &TemplateList); void replace(const std::vector &TemplateList); @@ -136,6 +137,20 @@ class StringReplacements { std::make_pair(Offset, std::make_shared( SourceStr, Offset, Length, TemplateIndex))); } + + inline void addTemplateDependentReplacementInConstExprExpansion( + size_t Offset, size_t Length, unsigned TemplateIndex) { + TDRsInConstExprExpansion.insert( + std::make_pair(Offset, std::make_shared( + SourceStr, Offset, Length, TemplateIndex))); + } + + inline void addConstExprExpansionInfo(std::string VDStr, + std::string VDInitStr) { + ConstExprExpansionInfoVDStr = VDStr; + ConstExprExpansionInfoVDInitStr = VDInitStr; + } + // Add a string replacement void addStringReplacement(size_t Offset, size_t Length, std::string Text) { auto Result = ReplMap.insert(std::make_pair( @@ -155,6 +170,11 @@ class StringReplacements { // Generate replacement text info which dependent on template args. std::shared_ptr getTemplateDependentStringInfo() { + if (!ConstExprExpansionInfoVDStr.empty() && + !ConstExprExpansionInfoVDInitStr.empty() && + !TDRsInConstExprExpansion.empty()) { + adjustSourceStrAndTDRs(); + } replaceString(); return std::make_shared(SourceStr, TDRs); } @@ -170,11 +190,16 @@ class StringReplacements { StringReplacements &operator=(StringReplacements) = delete; void replaceString(); + void adjustSourceStrAndTDRs(); unsigned ShiftLength; std::string SourceStr; std::map> ReplMap; std::map> TDRs; + std::map> + TDRsInConstExprExpansion; + std::string ConstExprExpansionInfoVDStr; + std::string ConstExprExpansionInfoVDInitStr; }; /// Analyze expression and generate its migrated string @@ -583,7 +608,11 @@ class ExprAnalysis { inline void addReplacement(size_t Offset, size_t Length, unsigned TemplateIndex) { - ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); + if (ConstExprExpansion) + ReplSet.addTemplateDependentReplacementInConstExprExpansion( + Offset, Length, TemplateIndex); + else + ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); } // Analyze the expression, jump to corresponding analysis function according @@ -691,6 +720,7 @@ class ExprAnalysis { std::string RewritePrefix; std::string RewritePostfix; std::set HelperFeatureSet; + bool ConstExprExpansion = false; }; // Analyze pointer allocated by cudaMallocManaged. From 5a48a524eaac1981f5383b94d5bfba1641847ced Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 15 Apr 2025 14:42:06 +0800 Subject: [PATCH 02/17] Add test Signed-off-by: Jiang, Zhiwei --- clang/test/dpct/sharedmem_var_static.cu | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/clang/test/dpct/sharedmem_var_static.cu b/clang/test/dpct/sharedmem_var_static.cu index c6d9832187eb..98f876f2fe9b 100644 --- a/clang/test/dpct/sharedmem_var_static.cu +++ b/clang/test/dpct/sharedmem_var_static.cu @@ -4,11 +4,12 @@ // RUN: dpct --format-range=none --usm-level=none -out-root %T/sharedmem_var_static %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only // RUN: FileCheck %s --match-full-lines --input-file %T/sharedmem_var_static/sharedmem_var_static.dp.cpp // RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST %T/sharedmem_var_static/sharedmem_var_static.dp.cpp -o %T/sharedmem_var_static/sharedmem_var_static.dp.o %} -#ifndef NO_BUILD_TEST + #include #include #define SIZE 64 +#ifndef NO_BUILD_TEST class TestObject{ public: // CHECK: static void run(int *in, int *out, const sycl::nd_item<3> &item_ct1, int &a0) { @@ -226,3 +227,15 @@ void fooh() { fook<<<1, 1>>>(); } #endif + +constexpr int kWarpSize = 32; + +template __global__ void kerfunc() { + constexpr int kNumWarps = ThreadsPerBlock / kWarpSize; + __shared__ int smem[kNumWarps * NumWarpQ]; +} + +void foo2() { + // CHECK: sycl::local_accessor smem_acc_ct1(sycl::range<1>(ThreadsPerBlock / kWarpSize * 8), cgh); + kerfunc<128, 8><<<32, 32>>>(); +} From 0bbb4b087d389e0d6a0fabc44f8008400fcf8578 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 16 Apr 2025 13:36:46 +0800 Subject: [PATCH 03/17] Fix Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 1 - clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 8 +++++--- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 21 ++++++++++++--------- clang/test/dpct/sharedmem_var_static.cu | 2 +- 4 files changed, 18 insertions(+), 14 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 602dc6137da3..855d6414a7d9 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2811,7 +2811,6 @@ void CtTypeInfo::setArrayInfo(const DependentSizedArrayTypeLoc &TL, bool NeedSizeFold) { ContainSizeofType = containSizeOfType(TL.getSizeExpr()); ExprAnalysis EA; - // EA.analyze(TL.getSizeExpr()); auto TDSI = EA.getTemplateDependentStringInfo(); if (TDSI->containsTemplateDependentMacro()) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index f13dc8dcc996..fe4a65af8822 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -625,11 +625,13 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { VD && VD->isConstexpr()) { if (VD->getInit() && VD->getInit()->getBeginLoc().isValid() && !ConstExprExpansion) { - std::string VDInitStr = ExprAnalysis::ref(VD->getInit()); + ConstExprExpansion = true; + ExprAnalysis EA(VD->getInit()); + std::string VDInitStr = EA.getReplacedString(); std::string VDStr = VD->getNameAsString(); ReplSet.addConstExprExpansionInfo(VDStr, VDInitStr); - ConstExprExpansion = true; - dispatch(VD->getInit()); + ReplSet.addTemplateDependentReplacementInConstExprExpansion( + EA.getReplSetTDRs()); ConstExprExpansion = false; } } else if (auto ECD = dyn_cast(DRE->getDecl())) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 4cb0eb08623b..7bca464e0437 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -139,10 +139,9 @@ class StringReplacements { } inline void addTemplateDependentReplacementInConstExprExpansion( - size_t Offset, size_t Length, unsigned TemplateIndex) { - TDRsInConstExprExpansion.insert( - std::make_pair(Offset, std::make_shared( - SourceStr, Offset, Length, TemplateIndex))); + const std::map> + &InTDRs) { + TDRsInConstExprExpansion.insert(InTDRs.begin(), InTDRs.end()); } inline void addConstExprExpansionInfo(std::string VDStr, @@ -183,6 +182,10 @@ class StringReplacements { replaceString(); return SourceStr; } + inline std::map> + getTDRs() { + return TDRs; + } private: StringReplacements(const StringReplacements &) = delete; @@ -255,6 +258,10 @@ class ExprAnalysis { inline const std::string &getReplacedString() { return ReplSet.getReplacedString(); } + inline std::map> + getReplSetTDRs() { + return ReplSet.getTDRs(); + } inline std::shared_ptr getTemplateDependentStringInfo() { auto Res = ReplSet.getTemplateDependentStringInfo(); @@ -608,11 +615,7 @@ class ExprAnalysis { inline void addReplacement(size_t Offset, size_t Length, unsigned TemplateIndex) { - if (ConstExprExpansion) - ReplSet.addTemplateDependentReplacementInConstExprExpansion( - Offset, Length, TemplateIndex); - else - ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); + ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); } // Analyze the expression, jump to corresponding analysis function according diff --git a/clang/test/dpct/sharedmem_var_static.cu b/clang/test/dpct/sharedmem_var_static.cu index 98f876f2fe9b..736cf99a9256 100644 --- a/clang/test/dpct/sharedmem_var_static.cu +++ b/clang/test/dpct/sharedmem_var_static.cu @@ -236,6 +236,6 @@ template __global__ void kerfunc() { } void foo2() { - // CHECK: sycl::local_accessor smem_acc_ct1(sycl::range<1>(ThreadsPerBlock / kWarpSize * 8), cgh); + // CHECK: sycl::local_accessor smem_acc_ct1(sycl::range<1>(128 / kWarpSize * 8), cgh); kerfunc<128, 8><<<32, 32>>>(); } From 173781d9808b36c89ff651d178ec9dfca60bac8c Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 28 Apr 2025 16:05:48 +0800 Subject: [PATCH 04/17] WIP Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 1 + clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 80 ++--------------------- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 36 +++++----- 3 files changed, 20 insertions(+), 97 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 67fe389b928e..b9fbc01d2e7a 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2811,6 +2811,7 @@ void CtTypeInfo::setArrayInfo(const DependentSizedArrayTypeLoc &TL, bool NeedSizeFold) { ContainSizeofType = containSizeOfType(TL.getSizeExpr()); ExprAnalysis EA; + EA.IsAnalyzingCtTypeInfo = true; EA.analyze(TL.getSizeExpr()); auto TDSI = EA.getTemplateDependentStringInfo(); if (TDSI->containsTemplateDependentMacro()) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index fe4a65af8822..4ff1642cb85d 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -421,74 +421,6 @@ void StringReplacements::replaceString() { ReplMap.clear(); } -void StringReplacements::adjustSourceStrAndTDRs() { - std::vector Occurrences; - size_t Pos = 0; - while ((Pos = SourceStr.find(ConstExprExpansionInfoVDStr, Pos)) != - std::string::npos) { - Occurrences.push_back(Pos); - Pos += ConstExprExpansionInfoVDStr.length(); - } - - int LengthDiff = ConstExprExpansionInfoVDInitStr.length() - - ConstExprExpansionInfoVDStr.length(); - - for (auto Occ : Occurrences) { - const auto &Iter = TDRs.find(Occ); - if (Iter != TDRs.end()) { - TDRs.erase(Iter); - } - } - std::map> NewTDRs; - for (const auto &[Pos, Value] : TDRs) { - int NewPos = Pos; - for (auto Occ : Occurrences) { - if (Occ < static_cast(Pos)) { - NewPos += LengthDiff; - } else { - break; - } - } - auto NewValue = std::make_shared( - SourceStr, NewPos, Value->getLength(), Value->getTemplateIndex()); - NewTDRs[NewPos] = NewValue; - } - TDRs = std::move(NewTDRs); - - std::map> - NewTDRsInConstExprExpansion; - for (const auto &[Pos, Value] : TDRsInConstExprExpansion) { - for (auto Occ : Occurrences) { - int NewPos = Occ + Pos; - for (auto PrevOcc : Occurrences) { - if (PrevOcc < static_cast(Occ)) { - NewPos += LengthDiff; - } else { - break; - } - } - auto NewValue = std::make_shared( - SourceStr, NewPos, Value->getLength(), Value->getTemplateIndex()); - NewTDRsInConstExprExpansion[NewPos] = NewValue; - break; - } - } - TDRsInConstExprExpansion = std::move(NewTDRsInConstExprExpansion); - - Pos = 0; - while ((Pos = SourceStr.find(ConstExprExpansionInfoVDStr, Pos)) != - std::string::npos) { - SourceStr.replace(Pos, ConstExprExpansionInfoVDStr.length(), - ConstExprExpansionInfoVDInitStr); - Pos += ConstExprExpansionInfoVDInitStr.length(); - } - - TDRs.insert(TDRsInConstExprExpansion.begin(), TDRsInConstExprExpansion.end()); - for (const auto &[Pos, Value] : TDRs) { - Value->alterSource(SourceStr); - } -} - ExprAnalysis::ExprAnalysis(const Expr *Expression) : Context(DpctGlobalInfo::getContext()), SM(DpctGlobalInfo::getSourceManager()) { @@ -622,17 +554,13 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { if (auto TemplateDecl = dyn_cast(DRE->getDecl())) addReplacement(DRE, TemplateDecl->getIndex()); else if (const auto *VD = dyn_cast(DRE->getDecl()); - VD && VD->isConstexpr()) { - if (VD->getInit() && VD->getInit()->getBeginLoc().isValid() && - !ConstExprExpansion) { - ConstExprExpansion = true; + VD && VD->isConstexpr() && + IsAnalyzingCtTypeInfo /*&& IsDependent*/) { + if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) { ExprAnalysis EA(VD->getInit()); std::string VDInitStr = EA.getReplacedString(); std::string VDStr = VD->getNameAsString(); - ReplSet.addConstExprExpansionInfo(VDStr, VDInitStr); - ReplSet.addTemplateDependentReplacementInConstExprExpansion( - EA.getReplSetTDRs()); - ConstExprExpansion = false; + addReplacement(0, VDStr.size(), VDInitStr); } } else if (auto ECD = dyn_cast(DRE->getDecl())) { std::unordered_set targetStr = { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 7bca464e0437..017af48dc36a 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -138,16 +138,12 @@ class StringReplacements { SourceStr, Offset, Length, TemplateIndex))); } - inline void addTemplateDependentReplacementInConstExprExpansion( - const std::map> - &InTDRs) { - TDRsInConstExprExpansion.insert(InTDRs.begin(), InTDRs.end()); - } - - inline void addConstExprExpansionInfo(std::string VDStr, - std::string VDInitStr) { - ConstExprExpansionInfoVDStr = VDStr; - ConstExprExpansionInfoVDInitStr = VDInitStr; + inline void addTemplateDependentReplacement( + size_t Offset, size_t Length, + std::shared_ptr TDSI) { + // TDRs.insert( + // std::make_pair(Offset, std::make_shared( + // SourceStr, Offset, Length, TemplateIndex))); } // Add a string replacement @@ -169,11 +165,6 @@ class StringReplacements { // Generate replacement text info which dependent on template args. std::shared_ptr getTemplateDependentStringInfo() { - if (!ConstExprExpansionInfoVDStr.empty() && - !ConstExprExpansionInfoVDInitStr.empty() && - !TDRsInConstExprExpansion.empty()) { - adjustSourceStrAndTDRs(); - } replaceString(); return std::make_shared(SourceStr, TDRs); } @@ -193,16 +184,11 @@ class StringReplacements { StringReplacements &operator=(StringReplacements) = delete; void replaceString(); - void adjustSourceStrAndTDRs(); unsigned ShiftLength; std::string SourceStr; std::map> ReplMap; std::map> TDRs; - std::map> - TDRsInConstExprExpansion; - std::string ConstExprExpansionInfoVDStr; - std::string ConstExprExpansionInfoVDInitStr; }; /// Analyze expression and generate its migrated string @@ -618,6 +604,12 @@ class ExprAnalysis { ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); } + inline void + addReplacement(size_t Offset, size_t Length, + std::shared_ptr TDSI) { + ReplSet.addTemplateDependentReplacement(Offset, Length, TDSI); + } + // Analyze the expression, jump to corresponding analysis function according // to its class // Precondition: Expression != nullptr @@ -723,7 +715,9 @@ class ExprAnalysis { std::string RewritePrefix; std::string RewritePostfix; std::set HelperFeatureSet; - bool ConstExprExpansion = false; + +public: + bool IsAnalyzingCtTypeInfo = false; }; // Analyze pointer allocated by cudaMallocManaged. From 43a81245cf4b04ec82f562501e4a5dcc31a4d614 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 28 Apr 2025 16:08:30 +0800 Subject: [PATCH 05/17] wip Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 9 --------- 1 file changed, 9 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 017af48dc36a..e0b9e4d9b009 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -77,7 +77,6 @@ class TemplateDependentReplacement { } inline size_t getOffset() const { return Offset; } inline size_t getLength() const { return Length; } - inline unsigned getTemplateIndex() const { return TemplateIndex; } const TemplateArgumentInfo & getTargetArgument(const std::vector &TemplateList); void replace(const std::vector &TemplateList); @@ -173,10 +172,6 @@ class StringReplacements { replaceString(); return SourceStr; } - inline std::map> - getTDRs() { - return TDRs; - } private: StringReplacements(const StringReplacements &) = delete; @@ -244,10 +239,6 @@ class ExprAnalysis { inline const std::string &getReplacedString() { return ReplSet.getReplacedString(); } - inline std::map> - getReplSetTDRs() { - return ReplSet.getTDRs(); - } inline std::shared_ptr getTemplateDependentStringInfo() { auto Res = ReplSet.getTemplateDependentStringInfo(); From d01ce43f995e2b3ab6cedb65739a3129265b2f7a Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 29 Apr 2025 11:40:53 +0800 Subject: [PATCH 06/17] WIP Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 3 ++ clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 39 +++++++++++++++-------- 2 files changed, 29 insertions(+), 13 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 4ff1642cb85d..1d399af5be78 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -560,7 +560,10 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { ExprAnalysis EA(VD->getInit()); std::string VDInitStr = EA.getReplacedString(); std::string VDStr = VD->getNameAsString(); + // This offset is relative to the original str addReplacement(0, VDStr.size(), VDInitStr); + // + addReplacement("ThreadsPerBlock", 0); } } else if (auto ECD = dyn_cast(DRE->getDecl())) { std::unordered_set targetStr = { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index e0b9e4d9b009..b85ef5c49cae 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -40,6 +40,7 @@ class StringReplacement { } inline const std::string &getReplacedText() { return Text; } + inline size_t getLength() { return Length; } private: // SourceStr is the string which need replaced. @@ -132,17 +133,30 @@ class StringReplacements { // Add a template dependent replacement inline void addTemplateDependentReplacement(size_t Offset, size_t Length, unsigned TemplateIndex) { - TDRs.insert( - std::make_pair(Offset, std::make_shared( - SourceStr, Offset, Length, TemplateIndex))); + // Find items in ReplMap whose offset <= Offset. + // Then check length, is there is overlap, ignore this insertion. + // Finally calculate the shift length. + int Shift = 0; + auto UpperBound = ReplMap.upper_bound(Offset); + for (auto It = ReplMap.begin(); It != UpperBound; ++It) { + if ((It->first + It->second->getLength()) > Offset) { + // overlap + return; + } + Shift += + (It->second->getReplacedText().length() - It->second->getLength()); + } + auto TDR = std::make_shared( + SourceStr, Offset, Length, TemplateIndex); + TDR->shift(Shift); + TDRs.insert(std::make_pair(Offset + Shift, TDR)); } - inline void addTemplateDependentReplacement( - size_t Offset, size_t Length, - std::shared_ptr TDSI) { - // TDRs.insert( - // std::make_pair(Offset, std::make_shared( - // SourceStr, Offset, Length, TemplateIndex))); + inline void addTemplateDependentReplacement(std::string String, + unsigned TemplateIndex) { + auto TDR = std::make_shared( + String, String.size(), TemplateIndex); + TDRs.insert(std::make_pair(0, TDR)); } // Add a string replacement @@ -595,10 +609,9 @@ class ExprAnalysis { ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); } - inline void - addReplacement(size_t Offset, size_t Length, - std::shared_ptr TDSI) { - ReplSet.addTemplateDependentReplacement(Offset, Length, TDSI); + inline void addReplacement(std::string String, + unsigned TemplateIndex) { + ReplSet.addTemplateDependentReplacement(String, TemplateIndex); } // Analyze the expression, jump to corresponding analysis function according From f2dde676e6e303d66a3e5cc8db717a5b998ee59a Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 29 Apr 2025 14:54:11 +0800 Subject: [PATCH 07/17] Fix Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 18 +++++++++++++----- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 4 +++- 2 files changed, 16 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 1d399af5be78..f2e8ae0518bb 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -552,7 +552,10 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { RefString = DRE->getNameInfo().getAsString(); } if (auto TemplateDecl = dyn_cast(DRE->getDecl())) - addReplacement(DRE, TemplateDecl->getIndex()); + if (FFFFF) + addReplacement("ThreadsPerBlock", TemplateDecl->getIndex()); + else + addReplacement(DRE, TemplateDecl->getIndex()); else if (const auto *VD = dyn_cast(DRE->getDecl()); VD && VD->isConstexpr() && IsAnalyzingCtTypeInfo /*&& IsDependent*/) { @@ -560,10 +563,15 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { ExprAnalysis EA(VD->getInit()); std::string VDInitStr = EA.getReplacedString(); std::string VDStr = VD->getNameAsString(); - // This offset is relative to the original str - addReplacement(0, VDStr.size(), VDInitStr); - // - addReplacement("ThreadsPerBlock", 0); + + auto Loc = ReplSet.getSourceStr().find(VDStr); + // TODO: more than 1 substrings matched + if (Loc != std::string::npos) { + addReplacement(Loc, VDStr.size(), VDInitStr); + FFFFF = true; + dispatch(VD->getInit()); + FFFFF = false; + } } } else if (auto ECD = dyn_cast(DRE->getDecl())) { std::unordered_set targetStr = { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index b85ef5c49cae..554b950fc76c 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -155,7 +155,7 @@ class StringReplacements { inline void addTemplateDependentReplacement(std::string String, unsigned TemplateIndex) { auto TDR = std::make_shared( - String, String.size(), TemplateIndex); + String, 0, String.size(), TemplateIndex); TDRs.insert(std::make_pair(0, TDR)); } @@ -186,6 +186,7 @@ class StringReplacements { replaceString(); return SourceStr; } + inline const std::string &getSourceStr() { return SourceStr; } private: StringReplacements(const StringReplacements &) = delete; @@ -202,6 +203,7 @@ class StringReplacements { /// Analyze expression and generate its migrated string class ExprAnalysis { + bool FFFFF = false; public: inline std::string getRewritePrefix() { return RewritePrefix; } From 63cd2179b209eb067f714e85353e9e24722e72f0 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Tue, 29 Apr 2025 16:10:26 +0800 Subject: [PATCH 08/17] FIX Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 13 +++++++------ clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 13 ++++++------- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index f2e8ae0518bb..8593f22b28fd 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -552,13 +552,14 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { RefString = DRE->getNameInfo().getAsString(); } if (auto TemplateDecl = dyn_cast(DRE->getDecl())) - if (FFFFF) - addReplacement("ThreadsPerBlock", TemplateDecl->getIndex()); - else + if (ConstExprExpanding) { + addReplacement(0 /*FIXME*/, TemplateDecl->getNameAsString(), + TemplateDecl->getIndex()); + } else addReplacement(DRE, TemplateDecl->getIndex()); else if (const auto *VD = dyn_cast(DRE->getDecl()); VD && VD->isConstexpr() && - IsAnalyzingCtTypeInfo /*&& IsDependent*/) { + IsAnalyzingCtTypeInfo /*FIXME: && IsDependent*/) { if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) { ExprAnalysis EA(VD->getInit()); std::string VDInitStr = EA.getReplacedString(); @@ -568,9 +569,9 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { // TODO: more than 1 substrings matched if (Loc != std::string::npos) { addReplacement(Loc, VDStr.size(), VDInitStr); - FFFFF = true; + ConstExprExpanding = true; dispatch(VD->getInit()); - FFFFF = false; + ConstExprExpanding = false; } } } else if (auto ECD = dyn_cast(DRE->getDecl())) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 554b950fc76c..c077e981e1c2 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -152,13 +152,12 @@ class StringReplacements { TDRs.insert(std::make_pair(Offset + Shift, TDR)); } - inline void addTemplateDependentReplacement(std::string String, + inline void addTemplateDependentReplacement(size_t Offset, std::string String, unsigned TemplateIndex) { auto TDR = std::make_shared( - String, 0, String.size(), TemplateIndex); - TDRs.insert(std::make_pair(0, TDR)); + String, Offset, String.size(), TemplateIndex); + TDRs.insert(std::make_pair(Offset, TDR)); } - // Add a string replacement void addStringReplacement(size_t Offset, size_t Length, std::string Text) { auto Result = ReplMap.insert(std::make_pair( @@ -203,7 +202,6 @@ class StringReplacements { /// Analyze expression and generate its migrated string class ExprAnalysis { - bool FFFFF = false; public: inline std::string getRewritePrefix() { return RewritePrefix; } @@ -611,9 +609,9 @@ class ExprAnalysis { ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); } - inline void addReplacement(std::string String, + inline void addReplacement(size_t Offset, std::string String, unsigned TemplateIndex) { - ReplSet.addTemplateDependentReplacement(String, TemplateIndex); + ReplSet.addTemplateDependentReplacement(Offset, String, TemplateIndex); } // Analyze the expression, jump to corresponding analysis function according @@ -721,6 +719,7 @@ class ExprAnalysis { std::string RewritePrefix; std::string RewritePostfix; std::set HelperFeatureSet; + bool ConstExprExpanding = false; public: bool IsAnalyzingCtTypeInfo = false; From 5fd4cb606704bfcd468e117d2d80f09bbfeae880 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 30 Apr 2025 09:23:46 +0800 Subject: [PATCH 09/17] fix Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 18 +++++++++++++----- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 4 +++- clang/test/dpct/sharedmem_var_static.cu | 4 ++-- 3 files changed, 18 insertions(+), 8 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 8593f22b28fd..9828f0e3a7e3 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -25,6 +25,7 @@ #include "clang/AST/StmtCXX.h" #include "clang/AST/TypeLoc.h" #include "llvm/Support/raw_ostream.h" +#include extern clang::tooling::UnifiedPath DpctInstallPath; namespace clang { @@ -552,9 +553,16 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { RefString = DRE->getNameInfo().getAsString(); } if (auto TemplateDecl = dyn_cast(DRE->getDecl())) - if (ConstExprExpanding) { - addReplacement(0 /*FIXME*/, TemplateDecl->getNameAsString(), - TemplateDecl->getIndex()); + if (ConstExprExpansionInfo) { + auto Loc = ConstExprExpansionInfo.value().first.find( + TemplateDecl->getNameAsString()); + // TODO: more than 1 substrings matched + if (Loc != std::string::npos) { + // Offset is relative to the final migrated string + addReplacement(Loc + ConstExprExpansionInfo.value().second, + TemplateDecl->getNameAsString(), + TemplateDecl->getIndex()); + } } else addReplacement(DRE, TemplateDecl->getIndex()); else if (const auto *VD = dyn_cast(DRE->getDecl()); @@ -569,9 +577,9 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { // TODO: more than 1 substrings matched if (Loc != std::string::npos) { addReplacement(Loc, VDStr.size(), VDInitStr); - ConstExprExpanding = true; + ConstExprExpansionInfo = std::make_pair(VDInitStr, Loc); dispatch(VD->getInit()); - ConstExprExpanding = false; + ConstExprExpansionInfo = std::nullopt; } } } else if (auto ECD = dyn_cast(DRE->getDecl())) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index c077e981e1c2..60399f71d5e3 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -719,7 +719,9 @@ class ExprAnalysis { std::string RewritePrefix; std::string RewritePostfix; std::set HelperFeatureSet; - bool ConstExprExpanding = false; + std::optional> + ConstExprExpansionInfo = std::nullopt; public: bool IsAnalyzingCtTypeInfo = false; diff --git a/clang/test/dpct/sharedmem_var_static.cu b/clang/test/dpct/sharedmem_var_static.cu index 5ac1aa5d47eb..fc87433fee7e 100644 --- a/clang/test/dpct/sharedmem_var_static.cu +++ b/clang/test/dpct/sharedmem_var_static.cu @@ -229,11 +229,11 @@ void fooh() { constexpr int kWarpSize = 32; template __global__ void kerfunc() { - constexpr int kNumWarps = ThreadsPerBlock / kWarpSize; + constexpr int kNumWarps = (2 * ThreadsPerBlock / kWarpSize); __shared__ int smem[kNumWarps * NumWarpQ]; } void foo2() { - // CHECK: sycl::local_accessor smem_acc_ct1(sycl::range<1>(128 / kWarpSize * 8), cgh); + // CHECK: sycl::local_accessor smem_acc_ct1(sycl::range<1>((2 * 128 / kWarpSize) * 8), cgh); kerfunc<128, 8><<<32, 32>>>(); } From 929e86e1aa99cc2bf79077e35ba91647a6ed1c5a Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 7 May 2025 11:32:15 +0800 Subject: [PATCH 10/17] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 21 ++++-------------- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 27 ++++++++++++++++------- 2 files changed, 23 insertions(+), 25 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 9828f0e3a7e3..869f4c7af142 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -553,33 +553,20 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { RefString = DRE->getNameInfo().getAsString(); } if (auto TemplateDecl = dyn_cast(DRE->getDecl())) - if (ConstExprExpansionInfo) { - auto Loc = ConstExprExpansionInfo.value().first.find( - TemplateDecl->getNameAsString()); - // TODO: more than 1 substrings matched - if (Loc != std::string::npos) { - // Offset is relative to the final migrated string - addReplacement(Loc + ConstExprExpansionInfo.value().second, - TemplateDecl->getNameAsString(), - TemplateDecl->getIndex()); - } - } else - addReplacement(DRE, TemplateDecl->getIndex()); + addReplacement(DRE, TemplateDecl->getIndex()); else if (const auto *VD = dyn_cast(DRE->getDecl()); VD && VD->isConstexpr() && IsAnalyzingCtTypeInfo /*FIXME: && IsDependent*/) { if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) { ExprAnalysis EA(VD->getInit()); - std::string VDInitStr = EA.getReplacedString(); + auto TDSI = EA.getTemplateDependentStringInfo(); std::string VDStr = VD->getNameAsString(); - + std::string VDInitStr = EA.getReplacedString(); auto Loc = ReplSet.getSourceStr().find(VDStr); // TODO: more than 1 substrings matched if (Loc != std::string::npos) { addReplacement(Loc, VDStr.size(), VDInitStr); - ConstExprExpansionInfo = std::make_pair(VDInitStr, Loc); - dispatch(VD->getInit()); - ConstExprExpansionInfo = std::nullopt; + addReplacement(Loc, "ThreadsPerBlock" /*FIXME*/, TDSI); } } } else if (auto ECD = dyn_cast(DRE->getDecl())) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 60399f71d5e3..0998d9e16525 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -78,6 +78,7 @@ class TemplateDependentReplacement { } inline size_t getOffset() const { return Offset; } inline size_t getLength() const { return Length; } + inline size_t getTemplateIndex() const { return TemplateIndex; } const TemplateArgumentInfo & getTargetArgument(const std::vector &TemplateList); void replace(const std::vector &TemplateList); @@ -118,6 +119,10 @@ class TemplateDependentStringInfo { HelperFeatureSet = Set; } bool containsTemplateDependentMacro() const { return ContainsTemplateDependentMacro; } + const std::vector> & + getTDRs() const { + return TDRs; + } }; /// Store an expr source string which may need replaced and its replacements @@ -152,12 +157,17 @@ class StringReplacements { TDRs.insert(std::make_pair(Offset + Shift, TDR)); } - inline void addTemplateDependentReplacement(size_t Offset, std::string String, - unsigned TemplateIndex) { - auto TDR = std::make_shared( - String, Offset, String.size(), TemplateIndex); - TDRs.insert(std::make_pair(Offset, TDR)); + inline void addTemplateDependentReplacement( + size_t Offset, std::string String, + std::shared_ptr TDSI) { + for (const auto &Item : TDSI->getTDRs()) { + Offset += Item->getOffset(); + auto TDR = std::make_shared( + String, Offset, String.size(), Item->getTemplateIndex()); + TDRs.insert(std::make_pair(Offset, TDR)); + } } + // Add a string replacement void addStringReplacement(size_t Offset, size_t Length, std::string Text) { auto Result = ReplMap.insert(std::make_pair( @@ -609,9 +619,10 @@ class ExprAnalysis { ReplSet.addTemplateDependentReplacement(Offset, Length, TemplateIndex); } - inline void addReplacement(size_t Offset, std::string String, - unsigned TemplateIndex) { - ReplSet.addTemplateDependentReplacement(Offset, String, TemplateIndex); + inline void + addReplacement(size_t Offset, std::string String, + std::shared_ptr TDSI) { + ReplSet.addTemplateDependentReplacement(Offset, String, TDSI); } // Analyze the expression, jump to corresponding analysis function according From aa7aad607cd94f51071055696db0d6825f5db479 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 7 May 2025 12:57:09 +0800 Subject: [PATCH 11/17] Fix Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 1 - clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 3 --- 2 files changed, 4 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 869f4c7af142..49363be99c95 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -25,7 +25,6 @@ #include "clang/AST/StmtCXX.h" #include "clang/AST/TypeLoc.h" #include "llvm/Support/raw_ostream.h" -#include extern clang::tooling::UnifiedPath DpctInstallPath; namespace clang { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 0998d9e16525..8cc735527455 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -730,9 +730,6 @@ class ExprAnalysis { std::string RewritePrefix; std::string RewritePostfix; std::set HelperFeatureSet; - std::optional> - ConstExprExpansionInfo = std::nullopt; public: bool IsAnalyzingCtTypeInfo = false; From fc43abde886941c07be1adf41e590af6b1ef25c3 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Wed, 7 May 2025 13:56:09 +0800 Subject: [PATCH 12/17] Fix Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 19 ++++++++++--------- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 16 +++++++++------- 2 files changed, 19 insertions(+), 16 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 49363be99c95..8e600b2fb7d5 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -554,18 +554,19 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { if (auto TemplateDecl = dyn_cast(DRE->getDecl())) addReplacement(DRE, TemplateDecl->getIndex()); else if (const auto *VD = dyn_cast(DRE->getDecl()); - VD && VD->isConstexpr() && - IsAnalyzingCtTypeInfo /*FIXME: && IsDependent*/) { + VD && VD->isConstexpr() && IsAnalyzingCtTypeInfo) { if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) { ExprAnalysis EA(VD->getInit()); auto TDSI = EA.getTemplateDependentStringInfo(); - std::string VDStr = VD->getNameAsString(); - std::string VDInitStr = EA.getReplacedString(); - auto Loc = ReplSet.getSourceStr().find(VDStr); - // TODO: more than 1 substrings matched - if (Loc != std::string::npos) { - addReplacement(Loc, VDStr.size(), VDInitStr); - addReplacement(Loc, "ThreadsPerBlock" /*FIXME*/, TDSI); + if (!TDSI->getTDRs().empty()) { + std::string VDStr = VD->getNameAsString(); + std::string VDInitStr = EA.getReplacedString(); + auto Loc = ReplSet.getSourceStr().find(VDStr); + // TODO: more than 1 substrings matched + if (Loc != std::string::npos) { + addReplacement(Loc, VDStr.size(), VDInitStr); + addReplacement(Loc, TDSI); + } } } } else if (auto ECD = dyn_cast(DRE->getDecl())) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 8cc735527455..dc55ec99724a 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -76,6 +76,7 @@ class TemplateDependentReplacement { return std::make_shared( SrcStr, Offset, Length, TemplateIndex); } + inline const std::string & getSourceStr() const { return SourceStr; } inline size_t getOffset() const { return Offset; } inline size_t getLength() const { return Length; } inline size_t getTemplateIndex() const { return TemplateIndex; } @@ -158,13 +159,14 @@ class StringReplacements { } inline void addTemplateDependentReplacement( - size_t Offset, std::string String, - std::shared_ptr TDSI) { + size_t Offset, std::shared_ptr TDSI) { for (const auto &Item : TDSI->getTDRs()) { - Offset += Item->getOffset(); + std::string String = + Item->getSourceStr().substr(Item->getOffset(), Item->getLength()); + size_t NewOffset = Offset + Item->getOffset(); auto TDR = std::make_shared( - String, Offset, String.size(), Item->getTemplateIndex()); - TDRs.insert(std::make_pair(Offset, TDR)); + String, NewOffset, String.size(), Item->getTemplateIndex()); + TDRs.insert(std::make_pair(NewOffset, TDR)); } } @@ -620,9 +622,9 @@ class ExprAnalysis { } inline void - addReplacement(size_t Offset, std::string String, + addReplacement(size_t Offset, std::shared_ptr TDSI) { - ReplSet.addTemplateDependentReplacement(Offset, String, TDSI); + ReplSet.addTemplateDependentReplacement(Offset, TDSI); } // Analyze the expression, jump to corresponding analysis function according From ce1138abcc571d27e435c1991d9aa1273f2f3bc8 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Thu, 8 May 2025 15:43:53 +0800 Subject: [PATCH 13/17] Format Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index dc55ec99724a..c29d6154279d 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -76,7 +76,7 @@ class TemplateDependentReplacement { return std::make_shared( SrcStr, Offset, Length, TemplateIndex); } - inline const std::string & getSourceStr() const { return SourceStr; } + inline const std::string &getSourceStr() const { return SourceStr; } inline size_t getOffset() const { return Offset; } inline size_t getLength() const { return Length; } inline size_t getTemplateIndex() const { return TemplateIndex; } From 0a5baa7b408589802f63bc6d199647600c245c45 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 12 May 2025 14:33:07 +0800 Subject: [PATCH 14/17] 123 Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 48 +++++++++++++--- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 70 +++++++---------------- 2 files changed, 60 insertions(+), 58 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 8e600b2fb7d5..6394cbdf2f6a 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -421,6 +421,42 @@ void StringReplacements::replaceString() { ReplMap.clear(); } +void StringReplacements::addStringReplacement(size_t Offset, size_t Length, + std::string Text) { + auto Result = ReplMap.insert(std::make_pair( + Offset, + std::make_shared(SourceStr, Offset, Length, Text))); + if (Result.second) { + auto Shift = Result.first->second->getReplacedText().length() - Length; + ShiftLength += Shift; + auto TDRItr = TDRs.upper_bound(Result.first->first); + while (TDRItr != TDRs.end()) { + TDRItr->second->shift(Shift); + ++TDRItr; + } + } +} + +void StringReplacements::addTemplateDependentReplacement( + size_t Offset, size_t Length, unsigned TemplateIndex) { + // Find items in ReplMap whose offset <= Offset. + // Then check length, is there is overlap, ignore this insertion. + // Finally calculate the shift length. + int Shift = 0; + auto UpperBound = ReplMap.upper_bound(Offset); + for (auto It = ReplMap.begin(); It != UpperBound; ++It) { + if ((It->first + It->second->getLength()) > Offset) { + // overlap + return; + } + Shift += (It->second->getReplacedText().length() - It->second->getLength()); + } + auto TDR = std::make_shared( + SourceStr, Offset, Length, TemplateIndex); + TDR->shift(Shift); + TDRs.insert(std::make_pair(Offset + Shift, TDR)); +} + ExprAnalysis::ExprAnalysis(const Expr *Expression) : Context(DpctGlobalInfo::getContext()), SM(DpctGlobalInfo::getSourceManager()) { @@ -558,15 +594,9 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) { ExprAnalysis EA(VD->getInit()); auto TDSI = EA.getTemplateDependentStringInfo(); - if (!TDSI->getTDRs().empty()) { - std::string VDStr = VD->getNameAsString(); - std::string VDInitStr = EA.getReplacedString(); - auto Loc = ReplSet.getSourceStr().find(VDStr); - // TODO: more than 1 substrings matched - if (Loc != std::string::npos) { - addReplacement(Loc, VDStr.size(), VDInitStr); - addReplacement(Loc, TDSI); - } + if (TDSI->getTDRs().size() == 1) { + auto LocInfo = getOffsetAndLength(DRE); + addReplacement(LocInfo.first, LocInfo.second, TDSI); } } } else if (auto ECD = dyn_cast(DRE->getDecl())) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index c29d6154279d..5b0f6e47f830 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -57,14 +57,14 @@ class TemplateArgumentInfo; /// Store replacement dependent on template args class TemplateDependentReplacement { - std::string &SourceStr; + std::string SourceStr; size_t Offset; size_t Length; unsigned TemplateIndex; public: - TemplateDependentReplacement(std::string &SrcStr, size_t Offset, - size_t Length, unsigned TemplateIndex) + TemplateDependentReplacement(std::string SrcStr, size_t Offset, size_t Length, + unsigned TemplateIndex) : SourceStr(SrcStr), Offset(Offset), Length(Length), TemplateIndex(TemplateIndex) {} TemplateDependentReplacement(const TemplateDependentReplacement &rhs) @@ -72,7 +72,7 @@ class TemplateDependentReplacement { rhs.TemplateIndex) {} inline std::shared_ptr - alterSource(std::string &SrcStr) { + alterSource(std::string SrcStr) { return std::make_shared( SrcStr, Offset, Length, TemplateIndex); } @@ -137,54 +137,26 @@ class StringReplacements { inline void reset() { ReplMap.clear(); } // Add a template dependent replacement - inline void addTemplateDependentReplacement(size_t Offset, size_t Length, - unsigned TemplateIndex) { - // Find items in ReplMap whose offset <= Offset. - // Then check length, is there is overlap, ignore this insertion. - // Finally calculate the shift length. - int Shift = 0; - auto UpperBound = ReplMap.upper_bound(Offset); - for (auto It = ReplMap.begin(); It != UpperBound; ++It) { - if ((It->first + It->second->getLength()) > Offset) { - // overlap - return; - } - Shift += - (It->second->getReplacedText().length() - It->second->getLength()); - } - auto TDR = std::make_shared( - SourceStr, Offset, Length, TemplateIndex); - TDR->shift(Shift); - TDRs.insert(std::make_pair(Offset + Shift, TDR)); - } + void addTemplateDependentReplacement(size_t Offset, size_t Length, + unsigned TemplateIndex); inline void addTemplateDependentReplacement( - size_t Offset, std::shared_ptr TDSI) { - for (const auto &Item : TDSI->getTDRs()) { - std::string String = - Item->getSourceStr().substr(Item->getOffset(), Item->getLength()); - size_t NewOffset = Offset + Item->getOffset(); - auto TDR = std::make_shared( - String, NewOffset, String.size(), Item->getTemplateIndex()); - TDRs.insert(std::make_pair(NewOffset, TDR)); - } + size_t Offset, size_t Length, + std::shared_ptr TDSI) { + if (TDSI->getTDRs().size() != 1) + return; + const auto &Item = TDSI->getTDRs()[0]; + addStringReplacement(Offset, Length, Item->getSourceStr()); + std::string String = + Item->getSourceStr().substr(Item->getOffset(), Item->getLength()); + size_t NewOffset = Offset + Item->getOffset(); + auto TDR = std::make_shared( + String, NewOffset, String.size(), Item->getTemplateIndex()); + TDRs.insert(std::make_pair(NewOffset, TDR)); } // Add a string replacement - void addStringReplacement(size_t Offset, size_t Length, std::string Text) { - auto Result = ReplMap.insert(std::make_pair( - Offset, - std::make_shared(SourceStr, Offset, Length, Text))); - if (Result.second) { - auto Shift = Result.first->second->getReplacedText().length() - Length; - ShiftLength += Shift; - auto TDRItr = TDRs.upper_bound(Result.first->first); - while (TDRItr != TDRs.end()) { - TDRItr->second->shift(Shift); - ++TDRItr; - } - } - } + void addStringReplacement(size_t Offset, size_t Length, std::string Text); // Generate replacement text info which dependent on template args. std::shared_ptr @@ -622,9 +594,9 @@ class ExprAnalysis { } inline void - addReplacement(size_t Offset, + addReplacement(size_t Offset, size_t Length, std::shared_ptr TDSI) { - ReplSet.addTemplateDependentReplacement(Offset, TDSI); + ReplSet.addTemplateDependentReplacement(Offset, Length, TDSI); } // Analyze the expression, jump to corresponding analysis function according From 106a0a668d0567bf302ec90c832f8720348a3a94 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 12 May 2025 16:21:26 +0800 Subject: [PATCH 15/17] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 5b0f6e47f830..e1232f060e29 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -169,7 +169,6 @@ class StringReplacements { replaceString(); return SourceStr; } - inline const std::string &getSourceStr() { return SourceStr; } private: StringReplacements(const StringReplacements &) = delete; From 7730bf1e6d863ee937bef461226e4ea2bc566e20 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 19 May 2025 11:37:34 +0800 Subject: [PATCH 16/17] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/AnalysisInfo.cpp | 2 +- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 64 +++++++++++------------ clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 25 ++++----- 3 files changed, 44 insertions(+), 47 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 6f7a9bace642..310834529023 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -2826,7 +2826,7 @@ void CtTypeInfo::setArrayInfo(const DependentSizedArrayTypeLoc &TL, auto TDSI = EA.getTemplateDependentStringInfo(); if (TDSI->containsTemplateDependentMacro()) TemplateDependentMacro = true; - Range.emplace_back(EA.getTemplateDependentStringInfo()); + Range.emplace_back(TDSI); setTypeInfo(TL.getElementLoc(), NeedSizeFold); } void CtTypeInfo::setArrayInfo(const IncompleteArrayTypeLoc &TL, diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index f00a78d31773..5e71fad1407a 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -412,7 +412,27 @@ void ExprAnalysis::initSourceRange(const SourceRange &Range) { } void StringReplacements::replaceString() { - SourceStr.reserve(SourceStr.length() + ShiftLength); + for (auto &TDR : TDRs2) { + // Find items in ReplMap whose offset <= TDR.first. + // Then check length, is there is overlap, ignore this insertion. + // Finally calculate the shift length. + int Shift = 0; + auto UpperBound = ReplMap.upper_bound(TDR.first); + for (auto It = ReplMap.begin(); It != UpperBound; ++It) { + if ((It->first + It->second->getLength()) > TDR.first) { + // overlap + continue; + } + Shift += + (It->second->getReplacedText().length() - It->second->getLength()); + } + auto NewTDR = std::make_shared( + TDR.second->SourceStr, TDR.first, TDR.second->getLength(), + TDR.second->getTemplateIndex()); + NewTDR->shift(Shift); + TDRs.insert(std::make_pair(TDR.first + Shift, NewTDR)); + } + auto Itr = ReplMap.rbegin(); while (Itr != ReplMap.rend()) { Itr->second->replaceString(); @@ -423,38 +443,15 @@ void StringReplacements::replaceString() { void StringReplacements::addStringReplacement(size_t Offset, size_t Length, std::string Text) { - auto Result = ReplMap.insert(std::make_pair( - Offset, - std::make_shared(SourceStr, Offset, Length, Text))); - if (Result.second) { - auto Shift = Result.first->second->getReplacedText().length() - Length; - ShiftLength += Shift; - auto TDRItr = TDRs.upper_bound(Result.first->first); - while (TDRItr != TDRs.end()) { - TDRItr->second->shift(Shift); - ++TDRItr; - } - } + ReplMap.insert(std::make_pair(Offset, std::make_shared( + SourceStr, Offset, Length, Text))); } void StringReplacements::addTemplateDependentReplacement( size_t Offset, size_t Length, unsigned TemplateIndex) { - // Find items in ReplMap whose offset <= Offset. - // Then check length, is there is overlap, ignore this insertion. - // Finally calculate the shift length. - int Shift = 0; - auto UpperBound = ReplMap.upper_bound(Offset); - for (auto It = ReplMap.begin(); It != UpperBound; ++It) { - if ((It->first + It->second->getLength()) > Offset) { - // overlap - return; - } - Shift += (It->second->getReplacedText().length() - It->second->getLength()); - } - auto TDR = std::make_shared( - SourceStr, Offset, Length, TemplateIndex); - TDR->shift(Shift); - TDRs.insert(std::make_pair(Offset + Shift, TDR)); + TDRs2.insert( + std::make_pair(Offset, std::make_shared( + SourceStr, Offset, Length, TemplateIndex))); } ExprAnalysis::ExprAnalysis(const Expr *Expression) @@ -609,13 +606,12 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) { addReplacement(DRE, TemplateDecl->getIndex()); else if (const auto *VD = dyn_cast(DRE->getDecl()); VD && VD->isConstexpr() && IsAnalyzingCtTypeInfo) { - if (VD->getInit() && VD->getInit()->getBeginLoc().isValid()) { + if (VD->getInit() && VD->getInit()->getBeginLoc().isValid() && + (VD->getInit()->getDependence() != ExprDependence::None)) { ExprAnalysis EA(VD->getInit()); auto TDSI = EA.getTemplateDependentStringInfo(); - if (TDSI->getTDRs().size() == 1) { - auto LocInfo = getOffsetAndLength(DRE); - addReplacement(LocInfo.first, LocInfo.second, TDSI); - } + auto LocInfo = getOffsetAndLength(DRE); + addReplacement(LocInfo.first, LocInfo.second, TDSI); } } else if (auto ECD = dyn_cast(DRE->getDecl())) { std::unordered_set targetStr = { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 540fc81f0151..7488b7ec25e3 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -43,6 +43,7 @@ class StringReplacement { inline size_t getLength() { return Length; } private: +public: // SourceStr is the string which need replaced. // Offset is the position where replacement happen. // Length is the replaced substring length @@ -57,6 +58,7 @@ class TemplateArgumentInfo; /// Store replacement dependent on template args class TemplateDependentReplacement { +public: std::string SourceStr; size_t Offset; size_t Length; @@ -129,7 +131,7 @@ class TemplateDependentStringInfo { /// Store an expr source string which may need replaced and its replacements class StringReplacements { public: - StringReplacements() : ShiftLength(0) {} + StringReplacements() {} inline void init(std::string &&SrcStr) { SourceStr = std::move(SrcStr); ReplMap.clear(); @@ -143,16 +145,15 @@ class StringReplacements { inline void addTemplateDependentReplacement( size_t Offset, size_t Length, std::shared_ptr TDSI) { - if (TDSI->getTDRs().size() != 1) - return; - const auto &Item = TDSI->getTDRs()[0]; - addStringReplacement(Offset, Length, Item->getSourceStr()); - std::string String = - Item->getSourceStr().substr(Item->getOffset(), Item->getLength()); - size_t NewOffset = Offset + Item->getOffset(); - auto TDR = std::make_shared( - String, NewOffset, String.size(), Item->getTemplateIndex()); - TDRs.insert(std::make_pair(NewOffset, TDR)); + addStringReplacement(Offset, Length, TDSI->getSourceString()); + for (const auto &Item : TDSI->getTDRs()) { + std::string String = + TDSI->getSourceString().substr(Item->getOffset(), Item->getLength()); + size_t NewOffset = Offset + Item->getOffset(); + auto TDR = std::make_shared( + String, NewOffset, String.size(), Item->getTemplateIndex()); + TDRs.insert(std::make_pair(NewOffset, TDR)); + } } // Add a string replacement @@ -177,10 +178,10 @@ class StringReplacements { void replaceString(); - unsigned ShiftLength; std::string SourceStr; std::map> ReplMap; std::map> TDRs; + std::map> TDRs2; }; /// Analyze expression and generate its migrated string From 701eb23a1bb495fde9c8dc6b45c2ebe3b62eab58 Mon Sep 17 00:00:00 2001 From: "Jiang, Zhiwei" Date: Mon, 19 May 2025 13:14:23 +0800 Subject: [PATCH 17/17] Update Signed-off-by: Jiang, Zhiwei --- clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp | 15 +-------------- clang/lib/DPCT/RuleInfra/ExprAnalysis.h | 14 ++++++++++---- clang/test/dpct/sharedmem_var_static.cu | 8 ++++---- 3 files changed, 15 insertions(+), 22 deletions(-) diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp index 5e71fad1407a..da94e4dcbc23 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.cpp @@ -427,7 +427,7 @@ void StringReplacements::replaceString() { (It->second->getReplacedText().length() - It->second->getLength()); } auto NewTDR = std::make_shared( - TDR.second->SourceStr, TDR.first, TDR.second->getLength(), + TDR.second->getSourceStr(), TDR.first, TDR.second->getLength(), TDR.second->getTemplateIndex()); NewTDR->shift(Shift); TDRs.insert(std::make_pair(TDR.first + Shift, NewTDR)); @@ -441,19 +441,6 @@ void StringReplacements::replaceString() { ReplMap.clear(); } -void StringReplacements::addStringReplacement(size_t Offset, size_t Length, - std::string Text) { - ReplMap.insert(std::make_pair(Offset, std::make_shared( - SourceStr, Offset, Length, Text))); -} - -void StringReplacements::addTemplateDependentReplacement( - size_t Offset, size_t Length, unsigned TemplateIndex) { - TDRs2.insert( - std::make_pair(Offset, std::make_shared( - SourceStr, Offset, Length, TemplateIndex))); -} - ExprAnalysis::ExprAnalysis(const Expr *Expression) : Context(DpctGlobalInfo::getContext()), SM(DpctGlobalInfo::getSourceManager()) { diff --git a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h index 7488b7ec25e3..e7fa74174f21 100644 --- a/clang/lib/DPCT/RuleInfra/ExprAnalysis.h +++ b/clang/lib/DPCT/RuleInfra/ExprAnalysis.h @@ -43,7 +43,6 @@ class StringReplacement { inline size_t getLength() { return Length; } private: -public: // SourceStr is the string which need replaced. // Offset is the position where replacement happen. // Length is the replaced substring length @@ -58,7 +57,6 @@ class TemplateArgumentInfo; /// Store replacement dependent on template args class TemplateDependentReplacement { -public: std::string SourceStr; size_t Offset; size_t Length; @@ -140,7 +138,11 @@ class StringReplacements { // Add a template dependent replacement void addTemplateDependentReplacement(size_t Offset, size_t Length, - unsigned TemplateIndex); + unsigned TemplateIndex) { + TDRs2.insert( + std::make_pair(Offset, std::make_shared( + SourceStr, Offset, Length, TemplateIndex))); + } inline void addTemplateDependentReplacement( size_t Offset, size_t Length, @@ -157,7 +159,11 @@ class StringReplacements { } // Add a string replacement - void addStringReplacement(size_t Offset, size_t Length, std::string Text); + void addStringReplacement(size_t Offset, size_t Length, std::string Text) { + ReplMap.insert(std::make_pair( + Offset, + std::make_shared(SourceStr, Offset, Length, Text))); + } // Generate replacement text info which dependent on template args. std::shared_ptr diff --git a/clang/test/dpct/sharedmem_var_static.cu b/clang/test/dpct/sharedmem_var_static.cu index fc87433fee7e..913c066e4d38 100644 --- a/clang/test/dpct/sharedmem_var_static.cu +++ b/clang/test/dpct/sharedmem_var_static.cu @@ -228,12 +228,12 @@ void fooh() { constexpr int kWarpSize = 32; -template __global__ void kerfunc() { - constexpr int kNumWarps = (2 * ThreadsPerBlock / kWarpSize); +template __global__ void kerfunc() { + constexpr int kNumWarps = (2 * ThreadsPerBlock / kWarpSize * ccc); __shared__ int smem[kNumWarps * NumWarpQ]; } void foo2() { - // CHECK: sycl::local_accessor smem_acc_ct1(sycl::range<1>((2 * 128 / kWarpSize) * 8), cgh); - kerfunc<128, 8><<<32, 32>>>(); + // CHECK: sycl::local_accessor smem_acc_ct1(sycl::range<1>((2 * 128 / kWarpSize * 16) * 8), cgh); + kerfunc<128, 8, 16><<<32, 32>>>(); }