From 69ea7968fce6353eabe1dde53f42603a6e1902cf Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 13 May 2025 14:06:03 +0800 Subject: [PATCH 1/9] [SYCLomatic] Cast the CUDA enum type to int if the type is scoped. Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/ASTTraversal.cpp | 1 + clang/lib/DPCT/RulesLang/RulesLang.cpp | 34 ++++++++++++++++++++++++++ clang/lib/DPCT/RulesLang/RulesLang.h | 7 ++++++ clang/test/dpct/enum_type.cu | 21 ++++++++++++++++ 4 files changed, 63 insertions(+) create mode 100644 clang/test/dpct/enum_type.cu diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 427f0c4fcc8c..e4e81e4de9e2 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -93,6 +93,7 @@ void MigrationRule::emplaceTransformation(TextModification *TM) { // RuleLang REGISTER_RULE(IterationSpaceBuiltinRule, PassKind::PK_Analysis) REGISTER_RULE(ErrorHandlingIfStmtRule, PassKind::PK_Migration) +REGISTER_RULE(BinaryOperatorCallRule, PassKind::PK_Migration) REGISTER_RULE(ErrorHandlingHostAPIRule, PassKind::PK_Migration) REGISTER_RULE(AtomicFunctionRule, PassKind::PK_Migration) REGISTER_RULE(ZeroLengthArrayRule, PassKind::PK_Migration) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index d40666018a31..eebfa90686b3 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4534,6 +4534,40 @@ void StreamAPICallRule::runRule(const MatchFinder::MatchResult &Result) { } } +void BinaryOperatorCallRule::registerMatcher(ast_matchers::MatchFinder &MF) { + MF.addMatcher(binaryOperator(isComparisonOperator()).bind("binOp"), this); +} +void BinaryOperatorCallRule::runRule( + const ast_matchers::MatchFinder::MatchResult &Result) { + auto BO = getNodeAsType(Result, "binOp"); + if (!BO) + return; + auto InsertEnumCast = [&](const Expr *E) { + QualType EType = E->getType(); + if (EType->isEnumeralType()) { + if (const auto *EnumType = EType->getAs()) { + const clang::EnumDecl *EnumDecl = EnumType->getDecl(); + clang::SourceLocation EnumLoc = EnumDecl->getLocation(); + if (dpct::DpctGlobalInfo::isInCudaPath(EnumLoc) && + !EnumDecl->isScoped()) { + SourceLocation EndLoc = Lexer::getLocForEndOfToken( + E->getEndLoc(), 0, DpctGlobalInfo::getSourceManager(), + LangOptions()); + DpctGlobalInfo::getInstance().addReplacement( + std::make_shared( + DpctGlobalInfo::getSourceManager(), E->getBeginLoc(), 0, + "static_cast(", nullptr)); + DpctGlobalInfo::getInstance().addReplacement( + std::make_shared( + DpctGlobalInfo::getSourceManager(), EndLoc, 0, ")", nullptr)); + } + } + } + }; + InsertEnumCast(BO->getLHS()->IgnoreImpCasts()); + InsertEnumCast(BO->getRHS()->IgnoreImpCasts()); +} + void KernelCallRefRule::registerMatcher(ast_matchers::MatchFinder &MF) { MF.addMatcher( functionDecl( diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index a9e83884103d..66231a59452b 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -434,6 +434,13 @@ class StreamAPICallRule : public NamedMigrationRule { void runRule(const ast_matchers::MatchFinder::MatchResult &Result); }; +/// Migration rule for binary operator calls +class BinaryOperatorCallRule : public NamedMigrationRule { +public: + void registerMatcher(ast_matchers::MatchFinder &MF) override; + void runRule(const ast_matchers::MatchFinder::MatchResult &Result); +}; + /// Migration rule for kernel API calls class KernelCallRule : public NamedMigrationRule { std::unordered_set Insertions; diff --git a/clang/test/dpct/enum_type.cu b/clang/test/dpct/enum_type.cu new file mode 100644 index 000000000000..55216ac820d6 --- /dev/null +++ b/clang/test/dpct/enum_type.cu @@ -0,0 +1,21 @@ +// RUN: dpct --format-range=none --use-experimental-features=matrix -out-root %T/enum_type %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/enum_type/enum_type.dp.cpp --match-full-lines %s +// RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST %T/enum_type/enum_type.dp.cpp -o %T/enum_type/enum_type.dp.o %} +#include +int main() { +// CHECK: sycl::usm::alloc mem_type; + CUmemorytype mem_type; +// CHECK: if (static_cast(mem_type) == 0) + if (mem_type == 0) + ; +// CHECK: if (0 == static_cast(mem_type)) + if (0 == mem_type) + ; +// CHECK: if (0 <= static_cast(mem_type)) + if (0 <= mem_type) + ; +// CHECK: if (static_cast(mem_type) > 0) + if (mem_type > 0) + ; + return 0; +} From e2e89e848119b24543b7db3087520d414e59b6d7 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Wed, 14 May 2025 08:07:08 -0700 Subject: [PATCH 2/9] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 18 +++++++++++++++++- 1 file changed, 17 insertions(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index eebfa90686b3..7cf87a9f5fe7 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4542,12 +4542,23 @@ void BinaryOperatorCallRule::runRule( auto BO = getNodeAsType(Result, "binOp"); if (!BO) return; + std::vector SkipType = {"cudaMemoryType", "cufftResult_t", + "cudaError", "CUresult", "cudaError_enum"}; + auto InsertEnumCast = [&](const Expr *E) { QualType EType = E->getType(); if (EType->isEnumeralType()) { - if (const auto *EnumType = EType->getAs()) { + + if (const auto *EnumType = + EType.getCanonicalType()->getAs()) { + const clang::EnumDecl *EnumDecl = EnumType->getDecl(); + std::string EnumName = EnumDecl->getNameAsString(); clang::SourceLocation EnumLoc = EnumDecl->getLocation(); + auto it = std::find(SkipType.begin(), SkipType.end(), EnumName); + if (it != SkipType.end() || + EnumName.empty()) // Empty means the enum is Anonymous + return; if (dpct::DpctGlobalInfo::isInCudaPath(EnumLoc) && !EnumDecl->isScoped()) { SourceLocation EndLoc = Lexer::getLocForEndOfToken( @@ -4564,6 +4575,11 @@ void BinaryOperatorCallRule::runRule( } } }; + auto LHSType = BO->getLHS()->IgnoreCasts()->getType(); + auto RHSType = BO->getRHS()->IgnoreCasts()->getType(); + if (LHSType->isEnumeralType() && RHSType->isEnumeralType()) { + return; + } InsertEnumCast(BO->getLHS()->IgnoreImpCasts()); InsertEnumCast(BO->getRHS()->IgnoreImpCasts()); } From 5cce188acbd7e8bf00187f9a8e83c949b55fa562 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Wed, 14 May 2025 18:45:53 -0700 Subject: [PATCH 3/9] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 7cf87a9f5fe7..5ebffb85bf05 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4542,8 +4542,11 @@ void BinaryOperatorCallRule::runRule( auto BO = getNodeAsType(Result, "binOp"); if (!BO) return; - std::vector SkipType = {"cudaMemoryType", "cufftResult_t", - "cudaError", "CUresult", "cudaError_enum"}; + + // The migration rule covered these test type, no need to cast. + std::vector MigratedEnumType = {"cudaMemoryType", "cufftResult_t", + "cudaError", "CUresult", + "cudaError_enum", "cudaComputeMode"}; auto InsertEnumCast = [&](const Expr *E) { QualType EType = E->getType(); @@ -4555,8 +4558,9 @@ void BinaryOperatorCallRule::runRule( const clang::EnumDecl *EnumDecl = EnumType->getDecl(); std::string EnumName = EnumDecl->getNameAsString(); clang::SourceLocation EnumLoc = EnumDecl->getLocation(); - auto it = std::find(SkipType.begin(), SkipType.end(), EnumName); - if (it != SkipType.end() || + auto it = std::find(MigratedEnumType.begin(), MigratedEnumType.end(), + EnumName); + if (it != MigratedEnumType.end() || EnumName.empty()) // Empty means the enum is Anonymous return; if (dpct::DpctGlobalInfo::isInCudaPath(EnumLoc) && @@ -4575,8 +4579,8 @@ void BinaryOperatorCallRule::runRule( } } }; - auto LHSType = BO->getLHS()->IgnoreCasts()->getType(); - auto RHSType = BO->getRHS()->IgnoreCasts()->getType(); + auto LHSType = BO->getLHS()->IgnoreImpCasts()->getType(); + auto RHSType = BO->getRHS()->IgnoreImpCasts()->getType(); if (LHSType->isEnumeralType() && RHSType->isEnumeralType()) { return; } From f8f4575e043d6e8c535a5f00c05ec539ad610b52 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 15 May 2025 22:53:03 -0700 Subject: [PATCH 4/9] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/ASTTraversal.cpp | 2 +- clang/lib/DPCT/RulesLang/RulesLang.cpp | 28 +++++++++++--------------- clang/lib/DPCT/RulesLang/RulesLang.h | 2 +- 3 files changed, 14 insertions(+), 18 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index e4e81e4de9e2..97127494bbb4 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -93,7 +93,7 @@ void MigrationRule::emplaceTransformation(TextModification *TM) { // RuleLang REGISTER_RULE(IterationSpaceBuiltinRule, PassKind::PK_Analysis) REGISTER_RULE(ErrorHandlingIfStmtRule, PassKind::PK_Migration) -REGISTER_RULE(BinaryOperatorCallRule, PassKind::PK_Migration) +REGISTER_RULE(CastScopedEnumTypeRule, PassKind::PK_Migration) REGISTER_RULE(ErrorHandlingHostAPIRule, PassKind::PK_Migration) REGISTER_RULE(AtomicFunctionRule, PassKind::PK_Migration) REGISTER_RULE(ZeroLengthArrayRule, PassKind::PK_Migration) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 5ebffb85bf05..9a5136916a48 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4534,33 +4534,28 @@ void StreamAPICallRule::runRule(const MatchFinder::MatchResult &Result) { } } -void BinaryOperatorCallRule::registerMatcher(ast_matchers::MatchFinder &MF) { +void CastScopedEnumTypeRule::registerMatcher(ast_matchers::MatchFinder &MF) { MF.addMatcher(binaryOperator(isComparisonOperator()).bind("binOp"), this); } -void BinaryOperatorCallRule::runRule( +void CastScopedEnumTypeRule::runRule( const ast_matchers::MatchFinder::MatchResult &Result) { auto BO = getNodeAsType(Result, "binOp"); if (!BO) return; - // The migration rule covered these test type, no need to cast. - std::vector MigratedEnumType = {"cudaMemoryType", "cufftResult_t", - "cudaError", "CUresult", - "cudaError_enum", "cudaComputeMode"}; - auto InsertEnumCast = [&](const Expr *E) { QualType EType = E->getType(); if (EType->isEnumeralType()) { - if (const auto *EnumType = EType.getCanonicalType()->getAs()) { const clang::EnumDecl *EnumDecl = EnumType->getDecl(); std::string EnumName = EnumDecl->getNameAsString(); clang::SourceLocation EnumLoc = EnumDecl->getLocation(); - auto it = std::find(MigratedEnumType.begin(), MigratedEnumType.end(), - EnumName); - if (it != MigratedEnumType.end() || + std::string ReplacedName = + MapNames::findReplacedName(MapNames::TypeNamesMap, EnumName); + + if (!ReplacedName.empty() || ReplacedName == EnumName || EnumName.empty()) // Empty means the enum is Anonymous return; if (dpct::DpctGlobalInfo::isInCudaPath(EnumLoc) && @@ -4579,13 +4574,14 @@ void BinaryOperatorCallRule::runRule( } } }; - auto LHSType = BO->getLHS()->IgnoreImpCasts()->getType(); - auto RHSType = BO->getRHS()->IgnoreImpCasts()->getType(); - if (LHSType->isEnumeralType() && RHSType->isEnumeralType()) { + auto LHSExpr = BO->getLHS()->IgnoreImpCasts(); + auto RHSExpr = BO->getRHS()->IgnoreImpCasts(); + if (LHSExpr->getType()->isEnumeralType() && + RHSExpr->getType()->isEnumeralType()) { return; } - InsertEnumCast(BO->getLHS()->IgnoreImpCasts()); - InsertEnumCast(BO->getRHS()->IgnoreImpCasts()); + InsertEnumCast(LHSExpr); + InsertEnumCast(RHSExpr); } void KernelCallRefRule::registerMatcher(ast_matchers::MatchFinder &MF) { diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index 66231a59452b..d79be82a47d0 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -435,7 +435,7 @@ class StreamAPICallRule : public NamedMigrationRule { }; /// Migration rule for binary operator calls -class BinaryOperatorCallRule : public NamedMigrationRule { +class CastScopedEnumTypeRule : public NamedMigrationRule { public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); From dff923b580696d658b38e4f4e82c42d9d0a318cd Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 16 May 2025 01:51:11 -0700 Subject: [PATCH 5/9] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 64 +++++++++++++------------- clang/test/dpct/enum_type.cu | 13 ++++-- 2 files changed, 39 insertions(+), 38 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 9a5136916a48..8f0b49d75bba 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4543,45 +4543,43 @@ void CastScopedEnumTypeRule::runRule( if (!BO) return; + // List the types don't need to explicit cast type after migration. + const std::unordered_set TypeNoCast = { + "int", MapNames::getDpctNamespace() + "err0", + MapNames::getDpctNamespace() + "err1", + MapNames::getDpctNamespace() + "pointer_attributes"}; + + auto IsReplacedTypeNoCast = [&](const std::string &Type) { + return TypeNoCast.count(Type) > 0; + }; + auto InsertEnumCast = [&](const Expr *E) { - QualType EType = E->getType(); - if (EType->isEnumeralType()) { - if (const auto *EnumType = - EType.getCanonicalType()->getAs()) { - - const clang::EnumDecl *EnumDecl = EnumType->getDecl(); - std::string EnumName = EnumDecl->getNameAsString(); - clang::SourceLocation EnumLoc = EnumDecl->getLocation(); - std::string ReplacedName = - MapNames::findReplacedName(MapNames::TypeNamesMap, EnumName); - - if (!ReplacedName.empty() || ReplacedName == EnumName || - EnumName.empty()) // Empty means the enum is Anonymous - return; - if (dpct::DpctGlobalInfo::isInCudaPath(EnumLoc) && - !EnumDecl->isScoped()) { - SourceLocation EndLoc = Lexer::getLocForEndOfToken( - E->getEndLoc(), 0, DpctGlobalInfo::getSourceManager(), - LangOptions()); - DpctGlobalInfo::getInstance().addReplacement( - std::make_shared( - DpctGlobalInfo::getSourceManager(), E->getBeginLoc(), 0, - "static_cast(", nullptr)); - DpctGlobalInfo::getInstance().addReplacement( - std::make_shared( - DpctGlobalInfo::getSourceManager(), EndLoc, 0, ")", nullptr)); - } - } + const clang::EnumDecl *EnumDecl = + E->getType().getCanonicalType()->getAs()->getDecl(); + + std::string EnumName = EnumDecl->getNameAsString(); + std::string ReplacedName = + MapNames::findReplacedName(MapNames::TypeNamesMap, EnumName); + + if (IsReplacedTypeNoCast(ReplacedName) || ReplacedName == EnumName || + EnumName.empty()) // Empty means the enum is Anonymous + return; + if (dpct::DpctGlobalInfo::isInCudaPath(EnumDecl->getLocation()) && + !EnumDecl->isScoped()) { + + insertAroundStmt(E, "static_cast(", ")"); } }; auto LHSExpr = BO->getLHS()->IgnoreImpCasts(); auto RHSExpr = BO->getRHS()->IgnoreImpCasts(); - if (LHSExpr->getType()->isEnumeralType() && - RHSExpr->getType()->isEnumeralType()) { - return; + if (LHSExpr->getType()->isEnumeralType() && !dyn_cast(LHSExpr) && + !RHSExpr->getType()->isEnumeralType()) { + InsertEnumCast(LHSExpr); + } + if (!LHSExpr->getType()->isEnumeralType() && + RHSExpr->getType()->isEnumeralType() && !dyn_cast(RHSExpr)) { + InsertEnumCast(RHSExpr); } - InsertEnumCast(LHSExpr); - InsertEnumCast(RHSExpr); } void KernelCallRefRule::registerMatcher(ast_matchers::MatchFinder &MF) { diff --git a/clang/test/dpct/enum_type.cu b/clang/test/dpct/enum_type.cu index 55216ac820d6..3b0654a715cf 100644 --- a/clang/test/dpct/enum_type.cu +++ b/clang/test/dpct/enum_type.cu @@ -3,18 +3,21 @@ // RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST %T/enum_type/enum_type.dp.cpp -o %T/enum_type/enum_type.dp.o %} #include int main() { -// CHECK: sycl::usm::alloc mem_type; + // CHECK: sycl::usm::alloc mem_type; CUmemorytype mem_type; -// CHECK: if (static_cast(mem_type) == 0) + // CHECK: if (static_cast(mem_type) == 0) if (mem_type == 0) ; -// CHECK: if (0 == static_cast(mem_type)) + // CHECK: if (0 == static_cast(mem_type)) if (0 == mem_type) ; -// CHECK: if (0 <= static_cast(mem_type)) + // CHECK: if (sycl::usm::alloc::host == mem_type) + if (CU_MEMORYTYPE_HOST == mem_type) + ; + // CHECK: if (0 <= static_cast(mem_type)) if (0 <= mem_type) ; -// CHECK: if (static_cast(mem_type) > 0) + // CHECK: if (static_cast(mem_type) > 0) if (mem_type > 0) ; return 0; From cc75d54eebb24383e8cbb2ce5e93f10a09e9002e Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Sun, 18 May 2025 19:27:33 -0700 Subject: [PATCH 6/9] update Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 8f0b49d75bba..1ce21811eadd 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4562,7 +4562,8 @@ void CastScopedEnumTypeRule::runRule( MapNames::findReplacedName(MapNames::TypeNamesMap, EnumName); if (IsReplacedTypeNoCast(ReplacedName) || ReplacedName == EnumName || - EnumName.empty()) // Empty means the enum is Anonymous + EnumName.empty() || + ReplacedName.empty()) // EnumName Empty means the enum is Anonymous return; if (dpct::DpctGlobalInfo::isInCudaPath(EnumDecl->getLocation()) && !EnumDecl->isScoped()) { From 198dcbcaf8b77296036673abb52e58c869961b7f Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 19 May 2025 01:19:00 -0700 Subject: [PATCH 7/9] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/RulesLang.cpp | 17 ++++++----------- 1 file changed, 6 insertions(+), 11 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 1ce21811eadd..6a56daedebab 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -4549,10 +4549,6 @@ void CastScopedEnumTypeRule::runRule( MapNames::getDpctNamespace() + "err1", MapNames::getDpctNamespace() + "pointer_attributes"}; - auto IsReplacedTypeNoCast = [&](const std::string &Type) { - return TypeNoCast.count(Type) > 0; - }; - auto InsertEnumCast = [&](const Expr *E) { const clang::EnumDecl *EnumDecl = E->getType().getCanonicalType()->getAs()->getDecl(); @@ -4561,24 +4557,23 @@ void CastScopedEnumTypeRule::runRule( std::string ReplacedName = MapNames::findReplacedName(MapNames::TypeNamesMap, EnumName); - if (IsReplacedTypeNoCast(ReplacedName) || ReplacedName == EnumName || + if (TypeNoCast.count(ReplacedName) || ReplacedName == EnumName || EnumName.empty() || ReplacedName.empty()) // EnumName Empty means the enum is Anonymous return; - if (dpct::DpctGlobalInfo::isInCudaPath(EnumDecl->getLocation()) && - !EnumDecl->isScoped()) { - + if (dpct::DpctGlobalInfo::isInCudaPath(EnumDecl->getLocation())) { insertAroundStmt(E, "static_cast(", ")"); } }; auto LHSExpr = BO->getLHS()->IgnoreImpCasts(); auto RHSExpr = BO->getRHS()->IgnoreImpCasts(); + if (LHSExpr->getType()->isEnumeralType() && !dyn_cast(LHSExpr) && !RHSExpr->getType()->isEnumeralType()) { InsertEnumCast(LHSExpr); - } - if (!LHSExpr->getType()->isEnumeralType() && - RHSExpr->getType()->isEnumeralType() && !dyn_cast(RHSExpr)) { + } else if (!LHSExpr->getType()->isEnumeralType() && + RHSExpr->getType()->isEnumeralType() && + !dyn_cast(RHSExpr)) { InsertEnumCast(RHSExpr); } } From 859267abca874ba2b47ff55d6890791afa1daf5f Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 19 May 2025 23:58:06 -0700 Subject: [PATCH 8/9] format Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/RulesLang.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h index d79be82a47d0..cba7ca4859a4 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.h +++ b/clang/lib/DPCT/RulesLang/RulesLang.h @@ -435,7 +435,8 @@ class StreamAPICallRule : public NamedMigrationRule { }; /// Migration rule for binary operator calls -class CastScopedEnumTypeRule : public NamedMigrationRule { +class CastScopedEnumTypeRule + : public NamedMigrationRule { public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); From e6601dcad8a7e79f1bea1ca8c9dda6ac2317645a Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Wed, 28 May 2025 01:05:21 -0700 Subject: [PATCH 9/9] Remove the experimental option. Signed-off-by: Chen, Sheng S --- clang/test/dpct/enum_type.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/dpct/enum_type.cu b/clang/test/dpct/enum_type.cu index 3b0654a715cf..87b5bdf39c89 100644 --- a/clang/test/dpct/enum_type.cu +++ b/clang/test/dpct/enum_type.cu @@ -1,4 +1,4 @@ -// RUN: dpct --format-range=none --use-experimental-features=matrix -out-root %T/enum_type %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: dpct --format-range=none -out-root %T/enum_type %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only // RUN: FileCheck --input-file %T/enum_type/enum_type.dp.cpp --match-full-lines %s // RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST %T/enum_type/enum_type.dp.cpp -o %T/enum_type/enum_type.dp.o %} #include