From 768b42436358f98037798a0c5be0e31b98955c32 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Fri, 31 May 2024 15:44:26 +0800 Subject: [PATCH 1/7] [SYCLomatic] Fix migration bug of dim3 ctor in class. Signed-off-by: Tang, Jiajun jiajun.tang@intel.com --- clang/lib/DPCT/ASTTraversal.cpp | 21 +++++++++------------ clang/test/dpct/replace-dim3.cu | 13 ++++++++++++- 2 files changed, 21 insertions(+), 13 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 4dc3658e6228..d9facf0d0824 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -3022,18 +3022,15 @@ void ReplaceDim3CtorRule::registerMatcher(MatchFinder &MF) { .bind("dim3CtorDecl"), this); - MF.addMatcher( - cxxConstructExpr(hasType(namedDecl(hasName("dim3"))), argumentCountIs(3), - // skip fields in a struct. The source loc is - // messed up (points to the start of the struct) - unless(hasParent(initListExpr())), - unless(hasAncestor(cxxRecordDecl())), - unless(hasParent(varDecl())), - unless(hasParent(exprWithCleanups())), - unless(hasAncestor(cxxConstructExpr( - hasType(namedDecl(hasName("dim3"))))))) - .bind("dim3CtorNoDecl"), - this); + MF.addMatcher(cxxConstructExpr(hasType(namedDecl(hasName("dim3"))), + argumentCountIs(3), + unless(hasParent(initListExpr())), + unless(hasParent(varDecl())), + unless(hasParent(exprWithCleanups())), + unless(hasAncestor(cxxConstructExpr( + hasType(namedDecl(hasName("dim3"))))))) + .bind("dim3CtorNoDecl"), + this); MF.addMatcher( typeLoc(loc(qualType(hasDeclaration(anyOf( diff --git a/clang/test/dpct/replace-dim3.cu b/clang/test/dpct/replace-dim3.cu index 268eb166c2f3..e7a728c9d488 100644 --- a/clang/test/dpct/replace-dim3.cu +++ b/clang/test/dpct/replace-dim3.cu @@ -4,10 +4,10 @@ // RUN: FileCheck --input-file %T/replace-dim3/replace-dim3.dp.cpp --match-full-lines %s // RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %T/replace-dim3/replace-dim3.dp.cpp -o %T/replace-dim3/replace-dim3.dp.o %} -#ifndef BUILD_TEST #include #include +#ifndef BUILD_TEST #define NUM 23 #define CALL_FUNC(func) func() @@ -284,3 +284,14 @@ void dim3_foo() { }); } #endif + +// CHECK: class Dim3Struct { +// CHECK-NEXT: Dim3Struct() : x(sycl::range<3>(1, 2, 1)) {} +// CHECK-NEXT: sycl::range<3> x = sycl::range<3>(1, 4, 3); +// CHECK-NEXT: void f() { sycl::range<3>(1, 6, 5); } +// CHECK-NEXT: }; +class Dim3Struct { + Dim3Struct() : x(dim3(1, 2)) {} + dim3 x = dim3(3, 4); + void f() { dim3(5, 6); } +}; From 3bf571f2f05fce07f077385929f0d4ae6077e1c4 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Tue, 4 Jun 2024 18:34:03 +0800 Subject: [PATCH 2/7] Move dim3 ctor rule to EA. --- clang/lib/DPCT/ASTTraversal.cpp | 58 ++---------- clang/lib/DPCT/ASTTraversal.h | 3 - clang/lib/DPCT/ExprAnalysis.cpp | 104 +++++++++++++++++---- clang/lib/DPCT/TextModification.cpp | 137 ---------------------------- clang/lib/DPCT/TextModification.h | 33 ------- 5 files changed, 94 insertions(+), 241 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index d9facf0d0824..2fbaccf60dba 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -3010,26 +3010,14 @@ void ReplaceDim3CtorRule::registerMatcher(MatchFinder &MF) { argumentCountIs(1), unless(hasAncestor(cxxConstructExpr( hasType(namedDecl(hasName("dim3"))))))) - .bind("dim3Top"), - this); - - MF.addMatcher(cxxConstructExpr( - hasType(namedDecl(hasName("dim3"))), argumentCountIs(3), - anyOf(hasParent(varDecl()), hasParent(exprWithCleanups())), - unless(hasParent(initListExpr())), - unless(hasAncestor( - cxxConstructExpr(hasType(namedDecl(hasName("dim3"))))))) - .bind("dim3CtorDecl"), + .bind("dim3Ctor"), this); MF.addMatcher(cxxConstructExpr(hasType(namedDecl(hasName("dim3"))), - argumentCountIs(3), unless(hasParent(initListExpr())), - unless(hasParent(varDecl())), - unless(hasParent(exprWithCleanups())), unless(hasAncestor(cxxConstructExpr( hasType(namedDecl(hasName("dim3"))))))) - .bind("dim3CtorNoDecl"), + .bind("dim3Ctor"), this); MF.addMatcher( @@ -3040,41 +3028,15 @@ void ReplaceDim3CtorRule::registerMatcher(MatchFinder &MF) { this); } -ReplaceDim3Ctor *ReplaceDim3CtorRule::getReplaceDim3Modification( - const MatchFinder::MatchResult &Result) { - if (auto Ctor = getNodeAsType(Result, "dim3CtorDecl")) { - if(getParentKernelCall(Ctor)) - return nullptr; - // dim3 a; or dim3 a(1); - return new ReplaceDim3Ctor(Ctor, true /*isDecl*/); - } else if (auto Ctor = - getNodeAsType(Result, "dim3CtorNoDecl")) { - if(getParentKernelCall(Ctor)) - return nullptr; - // deflt = dim3(3); - return new ReplaceDim3Ctor(Ctor, false /*isDecl*/); - } else if (auto Ctor = getNodeAsType(Result, "dim3Top")) { - if(getParentKernelCall(Ctor)) - return nullptr; - // dim3 d3_6_3 = dim3(ceil(test.x + NUM), NUM + test.y, NUM + test.z + NUM); - if (auto A = ReplaceDim3Ctor::getConstructExpr(Ctor->getArg(0))) { - // strip the top CXXConstructExpr, if there's a CXXConstructExpr further - // down - return new ReplaceDim3Ctor(Ctor, A); - } else { - // Copy constructor case: dim3 a(copyfrom) - // No replacements are needed - return nullptr; - } - } - - return nullptr; -} - void ReplaceDim3CtorRule::runRule(const MatchFinder::MatchResult &Result) { - ReplaceDim3Ctor *R = getReplaceDim3Modification(Result); - if (R) { - emplaceTransformation(R); + if (auto Ctor = getNodeAsType(Result, "dim3Ctor")) { + if (getParentKernelCall(Ctor)) + return; + ExprAnalysis EA; + EA.analyze(Ctor); + emplaceTransformation(EA.getReplacement()); + EA.applyAllSubExprRepl(); + return; } if (auto TL = getNodeAsType(Result, "dim3Type")) { diff --git a/clang/lib/DPCT/ASTTraversal.h b/clang/lib/DPCT/ASTTraversal.h index 694fc93988a2..c66f2defabb5 100644 --- a/clang/lib/DPCT/ASTTraversal.h +++ b/clang/lib/DPCT/ASTTraversal.h @@ -581,9 +581,6 @@ class VectorTypeOperatorRule }; class ReplaceDim3CtorRule : public NamedMigrationRule { - ReplaceDim3Ctor *getReplaceDim3Modification( - const ast_matchers::MatchFinder::MatchResult &Result); - public: void registerMatcher(ast_matchers::MatchFinder &MF) override; void runRule(const ast_matchers::MatchFinder::MatchResult &Result); diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index 7310975ce944..0efdc55247ea 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -628,18 +628,44 @@ void ExprAnalysis::analyzeExpr(const CXXUnresolvedConstructExpr *Ctor) { } void ExprAnalysis::analyzeExpr(const CXXTemporaryObjectExpr *Temp) { - if (Temp->getConstructor()->getDeclName().getAsString() != "dim3") { - analyzeType(Temp->getTypeSourceInfo()->getTypeLoc()); - } + analyzeType(Temp->getTypeSourceInfo()->getTypeLoc()); analyzeExpr(static_cast(Temp)); } +const CXXConstructExpr *getConstructExpr(const Expr *E) { + if (const auto *C = dyn_cast_or_null(E)) { + return C; + } else if (isa(E)) { + return getConstructExpr( + dyn_cast(E)->getSubExpr()); + } else if (isa(E)) { + return getConstructExpr(dyn_cast(E)->getSubExpr()); + } else { + return nullptr; + } +} + void ExprAnalysis::analyzeExpr(const CXXConstructExpr *Ctor) { if (Ctor->getConstructor()->getDeclName().getAsString() == "dim3") { + auto NeedChangeName = dyn_cast(Ctor) == nullptr; + // strip the top CXXConstructExpr, if there's a CXXConstructExpr further + // down + if (Ctor->getNumArgs() == 1) { + Ctor = getConstructExpr(Ctor->getArg(0)); + } + if (Ctor == nullptr) { + return; + } + auto Parents = dpct::DpctGlobalInfo::getContext().getParents(*Ctor); + NeedChangeName = NeedChangeName && Parents.size() == 1 && + Parents[0].get() == nullptr && + Parents[0].get() == nullptr; std::string ArgsString; llvm::raw_string_ostream OS(ArgsString); - DpctGlobalInfo::printCtadClass(OS, MapNames::getClNamespace() + "range", 3) - << "("; + if (NeedChangeName) + DpctGlobalInfo::printCtadClass(OS, MapNames::getClNamespace() + "range", + 3) + << "("; ArgumentAnalysis A; std::string ArgStr = ""; for (auto Arg : Ctor->arguments()) { @@ -647,23 +673,61 @@ void ExprAnalysis::analyzeExpr(const CXXConstructExpr *Ctor) { ArgStr = ", " + A.getReplacedString() + ArgStr; } ArgStr.replace(0, 2, ""); - OS << ArgStr << ")"; + OS << ArgStr; + if (NeedChangeName) + OS << ")"; OS.flush(); - // Special handling for implicit ctor. - // #define GET_BLOCKS(a) a - // dim3 A = GET_BLOCKS(1); - // Result if using SM.getExpansionRange: - // sycl::range<3> A = sycl::range<3>(1, 1, GET_BLOCKS(1)); - // Result if using addReplacement(E): - // #define GET_BLOCKS(a) sycl::range<3>(1, 1, a) - // sycl::range<3> A = GET_BLOCKS(1); - if (Ctor->getParenOrBraceRange().isInvalid() && isOuterMostMacro(Ctor)) { - return addReplacement( - SM.getExpansionRange(Ctor->getBeginLoc()).getBegin(), - SM.getExpansionRange(Ctor->getEndLoc()).getEnd(), ArgsString); - } - addReplacement(Ctor, ArgsString); + CharSourceRange CSR; + if (!NeedChangeName) { + SourceRange SR = Ctor->getParenOrBraceRange(); + if (SR.isInvalid()) { + // convert to spelling location if the dim3 constructor is in a macro + // otherwise, Lexer::getLocForEndOfToken returns invalid source location + auto CtorLoc = Ctor->getLocation().isMacroID() + ? SM.getSpellingLoc(Ctor->getLocation()) + : Ctor->getLocation(); + // dim3 a; + // MACRO(... dim3 a; ...) + auto CtorEndLoc = Lexer::getLocForEndOfToken( + CtorLoc, 0, SM, DpctGlobalInfo::getContext().getLangOpts()); + CSR = CharSourceRange(SourceRange(CtorEndLoc, CtorEndLoc), false); + ArgsString = "(" + ArgsString + ")"; + } else { + SourceRange SR1 = + SourceRange(SR.getBegin().getLocWithOffset(1), SR.getEnd()); + CSR = CharSourceRange(SR1, false); + } + } else { + // adjust the statement to replace if top-level constructor includes the + // variable being defined + const Stmt *S = Ctor; + if (!S) { + return; + } + if (S->getBeginLoc().isMacroID() && !isOuterMostMacro(S)) { + auto Range = getDefinitionRange(S->getBeginLoc(), S->getEndLoc()); + auto Begin = Range.getBegin(); + auto End = Range.getEnd(); + End = End.getLocWithOffset(Lexer::MeasureTokenLength( + End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts())); + CSR = CharSourceRange::getTokenRange(Begin, End); + } else { + // Use getStmtExpansionSourceRange(S) to support cases like + // dim3 a = MACRO; + auto Range = getStmtExpansionSourceRange(S); + auto Begin = Range.getBegin(); + auto End = Range.getEnd(); + CSR = CharSourceRange::getTokenRange( + Begin, + End.getLocWithOffset(Lexer::MeasureTokenLength( + End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts()))); + } + } + auto Range = getDefinitionRange(CSR.getBegin(), CSR.getEnd()); + auto Length = SM.getDecomposedLoc(Range.getEnd()).second - + SM.getDecomposedLoc(Range.getBegin()).second; + addReplacement(Range.getBegin(), Length, ArgsString); return; } for (auto It = Ctor->arg_begin(); It != Ctor->arg_end(); It++) { diff --git a/clang/lib/DPCT/TextModification.cpp b/clang/lib/DPCT/TextModification.cpp index 0010962e66da..a11c402a744c 100644 --- a/clang/lib/DPCT/TextModification.cpp +++ b/clang/lib/DPCT/TextModification.cpp @@ -523,135 +523,6 @@ ReplaceInclude::getReplacement(const ASTContext &Context) const { this); } -void ReplaceDim3Ctor::setRange() { - auto &SM = DpctGlobalInfo::getSourceManager(); - if (isDecl) { - SourceRange SR = Ctor->getParenOrBraceRange(); - if (SR.isInvalid()) { - // convert to spelling location if the dim3 constructor is in a macro - // otherwise, Lexer::getLocForEndOfToken returns invalid source location - auto CtorLoc = Ctor->getLocation().isMacroID() - ? SM.getSpellingLoc(Ctor->getLocation()) - : Ctor->getLocation(); - // dim3 a; - // MACRO(... dim3 a; ...) - auto CtorEndLoc = Lexer::getLocForEndOfToken( - CtorLoc, 0, SM, DpctGlobalInfo::getContext().getLangOpts()); - CSR = CharSourceRange(SourceRange(CtorEndLoc, CtorEndLoc), false); - } else { - SourceRange SR1 = - SourceRange(SR.getBegin().getLocWithOffset(1), SR.getEnd()); - CSR = CharSourceRange(SR1, false); - } - } else { - // adjust the statement to replace if top-level constructor includes the - // variable being defined - const Stmt *S = getReplaceStmt(Ctor); - if (!S) { - return; - } - if (S->getBeginLoc().isMacroID() && !isOuterMostMacro(S)) { - auto Range = getDefinitionRange(S->getBeginLoc(), S->getEndLoc()); - auto Begin = Range.getBegin(); - auto End = Range.getEnd(); - End = End.getLocWithOffset(Lexer::MeasureTokenLength( - End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts())); - CSR = CharSourceRange::getTokenRange(Begin, End); - } else { - // Use getStmtExpansionSourceRange(S) to support cases like - // dim3 a = MACRO; - auto Range = getStmtExpansionSourceRange(S); - auto Begin = Range.getBegin(); - auto End = Range.getEnd(); - CSR = CharSourceRange::getTokenRange( - Begin, - End.getLocWithOffset(Lexer::MeasureTokenLength( - End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts()))); - } - } -} - -ReplaceInclude *ReplaceDim3Ctor::getEmpty() { - return new ReplaceInclude(CSR, ""); -} - -// Strips possible Materialize and Cast operators from CXXConstructor -const CXXConstructExpr *ReplaceDim3Ctor::getConstructExpr(const Expr *E) { - if (auto C = dyn_cast_or_null(E)) { - return C; - } else if (isa(E)) { - return getConstructExpr( - dyn_cast(E)->getSubExpr()); - } else if (isa(E)) { - return getConstructExpr(dyn_cast(E)->getSubExpr()); - } else { - return nullptr; - } -} - -// Returns the full replacement string for the CXXConstructorExpr -std::string -ReplaceDim3Ctor::getSyclRangeCtor(const CXXConstructExpr *Ctor) const { - ExprAnalysis Analysis(Ctor); - return Analysis.getReplacedString(); -} - -const Stmt *ReplaceDim3Ctor::getReplaceStmt(const Stmt *S) const { - if (auto Ctor = dyn_cast_or_null(S)) { - if (Ctor->getNumArgs() == 1) { - return getConstructExpr(Ctor->getArg(0)); - } - } - return S; -} - -std::string ReplaceDim3Ctor::getReplaceString() const { - if (isDecl) { - // Get the new parameter list for the replaced constructor, without the - // parens - std::string ReplacedString; - llvm::raw_string_ostream OS(ReplacedString); - ArgumentAnalysis AA; - std::string ArgStr = ""; - for (auto Arg : Ctor->arguments()) { - AA.analyze(Arg); - ArgStr = ", " + AA.getReplacedString() + ArgStr; - } - ArgStr.replace(0, 2, ""); - OS << ArgStr; - OS.flush(); - if (Ctor->getParenOrBraceRange().isInvalid()) { - // dim3 = a; - ReplacedString = "(" + ReplacedString + ")"; - } - return ReplacedString; - } else { - std::string S; - if (FinalCtor) { - S = getSyclRangeCtor(FinalCtor); - } else { - S = getSyclRangeCtor(Ctor); - } - return S; - } -} - -std::shared_ptr -ReplaceDim3Ctor::getReplacement(const ASTContext &Context) const { - if (this->isIgnoreTM()) - return nullptr; - // Use getDefinitionRange in general cases, - // For cases like dim3 a = MACRO; - // CSR is already set to the expansion range. - auto &SM = dpct::DpctGlobalInfo::getSourceManager(); - ReplacementString = getReplaceString(); - auto Range = getDefinitionRange(CSR.getBegin(), CSR.getEnd()); - auto Length = SM.getDecomposedLoc(Range.getEnd()).second - - SM.getDecomposedLoc(Range.getBegin()).second; - return std::make_shared(SM, Range.getBegin(), Length, - getReplaceString(), this); -} - std::shared_ptr InsertComment::getReplacement(const ASTContext &Context) const { if (this->isIgnoreTM()) @@ -930,14 +801,6 @@ void ReplaceInclude::print(llvm::raw_ostream &OS, ASTContext &Context, printReplacement(OS, T); } -void ReplaceDim3Ctor::print(llvm::raw_ostream &OS, ASTContext &Context, - const bool PrintDetail) const { - printHeader(OS, getID(), PrintDetail ? getParentRuleName() : StringRef()); - printLocation(OS, CSR.getBegin(), Context, PrintDetail); - Ctor->printPretty(OS, nullptr, PrintingPolicy(Context.getLangOpts())); - printReplacement(OS, ReplacementString); -} - void InsertComment::print(llvm::raw_ostream &OS, ASTContext &Context, const bool PrintDetail) const { printHeader(OS, getID(), PrintDetail ? getParentRuleName() : StringRef()); diff --git a/clang/lib/DPCT/TextModification.h b/clang/lib/DPCT/TextModification.h index 6549bc8be1a4..2435b1fe805f 100644 --- a/clang/lib/DPCT/TextModification.h +++ b/clang/lib/DPCT/TextModification.h @@ -529,39 +529,6 @@ class ReplaceInclude : public TextModification { const bool PrintDetail = true) const override; }; -/// Replace Dim3 constructors -class ReplaceDim3Ctor : public TextModification { - bool isDecl; - const CXXConstructExpr *Ctor; - const CXXConstructExpr *FinalCtor; - CharSourceRange CSR; - mutable std::string ReplacementString; - - void setRange(); - const Stmt *getReplaceStmt(const Stmt *S) const; - std::string getSyclRangeCtor(const CXXConstructExpr *Ctor) const; - std::string getReplaceString() const; - -public: - ReplaceDim3Ctor(const CXXConstructExpr *_Ctor, bool _isDecl = false) - : TextModification(TMID::ReplaceDim3Ctor, G2), isDecl(_isDecl), - Ctor(_Ctor), FinalCtor(nullptr) { - setRange(); - } - ReplaceDim3Ctor(const CXXConstructExpr *_Ctor, - const CXXConstructExpr *_FinalCtor) - : TextModification(TMID::ReplaceDim3Ctor, G2), isDecl(false), Ctor(_Ctor), - FinalCtor(_FinalCtor) { - setRange(); - } - static const CXXConstructExpr *getConstructExpr(const Expr *E); - ReplaceInclude *getEmpty(); - std::shared_ptr - getReplacement(const ASTContext &Context) const override; - void print(llvm::raw_ostream &OS, ASTContext &Context, - const bool PrintDetail = true) const override; -}; - class InsertBeforeStmt : public TextModification { const Stmt *S; std::string T; From 33c5b601c58bb48af3e7d1cb5dd56d6fb443f8fd Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Thu, 13 Jun 2024 13:43:31 +0800 Subject: [PATCH 3/7] Siplify dim3 ctor EA logic. --- clang/lib/DPCT/ASTTraversal.cpp | 98 +++--------- clang/lib/DPCT/ExprAnalysis.cpp | 145 ++++++++---------- clang/lib/DPCT/ExprAnalysis.h | 1 + clang/lib/Sema/SemaInit.cpp | 12 +- clang/test/dpct/checkFormatAll.cu | 4 +- clang/test/dpct/checkFormatMigrated.cu | 4 +- clang/test/dpct/ctad.cu | 18 +-- clang/test/dpct/datatypes_test_part2.cu | 4 +- clang/test/dpct/dim3.cu | 8 +- clang/test/dpct/enable-all-extensions.cu | 2 +- clang/test/dpct/formatIndent.cu | 6 +- clang/test/dpct/formatMigratedExplicitly.cu | 4 +- clang/test/dpct/formatMigratedGoogle.cu | 4 +- clang/test/dpct/formatMigratedLLVM.cu | 4 +- .../dpct/kernel-call-origcode-embedded.cu | 4 +- clang/test/dpct/math_functions_std.cu | 2 +- clang/test/dpct/replace-dim3.cu | 36 ++--- .../dpct/thrust/thrust_testing/source/foo.cu | 2 +- clang/test/dpct/types001.cu | 2 +- 19 files changed, 148 insertions(+), 212 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 2fbaccf60dba..4a85bb8cc241 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -1688,7 +1688,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { MF.addMatcher( typeLoc( loc(qualType(hasDeclaration(namedDecl(hasAnyName( - "cudaError", "curandStatus", "cublasStatus", "CUstream", + "dim3", "cudaError", "curandStatus", "cublasStatus", "CUstream", "CUstream_st", "thrust::complex", "thrust::device_vector", "thrust::device_ptr", "thrust::device_reference", "thrust::host_vector", "cublasHandle_t", "CUevent_st", "__half", @@ -3006,6 +3006,12 @@ void ReplaceDim3CtorRule::registerMatcher(MatchFinder &MF) { // Find dim3 constructors which are part of different casts (representing // different syntaxes). This includes copy constructors. All constructors // will be visited once. + MF.addMatcher( + cxxNewExpr(hasType(pointsTo(namedDecl(hasName("dim3"))))).bind("dim3New"), + this); + MF.addMatcher( + explicitCastExpr(hasType(namedDecl(hasName("dim3")))).bind("dim3Cast"), + this); MF.addMatcher(cxxConstructExpr(hasType(namedDecl(hasName("dim3"))), argumentCountIs(1), unless(hasAncestor(cxxConstructExpr( @@ -3019,90 +3025,32 @@ void ReplaceDim3CtorRule::registerMatcher(MatchFinder &MF) { hasType(namedDecl(hasName("dim3"))))))) .bind("dim3Ctor"), this); - - MF.addMatcher( - typeLoc(loc(qualType(hasDeclaration(anyOf( - namedDecl(hasAnyName("dim3")), - typedefDecl(hasAnyName("dim3"))))))) - .bind("dim3Type"), - this); } void ReplaceDim3CtorRule::runRule(const MatchFinder::MatchResult &Result) { - if (auto Ctor = getNodeAsType(Result, "dim3Ctor")) { + const Expr *E = nullptr; + if (const auto *New = getNodeAsType(Result, "dim3New")) { + if (getParentKernelCall(New->getConstructExpr())) + return; + E = New; + } + if (const auto *Cast = getNodeAsType(Result, "dim3Cast")) { + if (getParentKernelCall(Cast->getSubExprAsWritten())) + return; + E = Cast; + } + if (const auto *Ctor = getNodeAsType(Result, "dim3Ctor")) { if (getParentKernelCall(Ctor)) return; + E = Ctor; + } + if (E) { ExprAnalysis EA; - EA.analyze(Ctor); + EA.analyze(E); emplaceTransformation(EA.getReplacement()); EA.applyAllSubExprRepl(); return; } - - if (auto TL = getNodeAsType(Result, "dim3Type")) { - if (TL->getBeginLoc().isInvalid()) - return; - - auto BeginLoc = - getDefinitionRange(TL->getBeginLoc(), TL->getEndLoc()).getBegin(); - SourceManager *SM = Result.SourceManager; - - // WA for concatenated macro token - if (SM->isWrittenInScratchSpace(SM->getSpellingLoc(TL->getBeginLoc()))) { - BeginLoc = SM->getExpansionLoc(TL->getBeginLoc()); - } - - Token Tok; - auto LOpts = Result.Context->getLangOpts(); - Lexer::getRawToken(BeginLoc, Tok, *SM, LOpts, true); - if (Tok.isAnyIdentifier()) { - if (TL->getType()->isElaboratedTypeSpecifier()) { - // To handle case like "struct cudaExtent extent;" - auto ETC = TL->getUnqualifiedLoc().getAs(); - auto NTL = ETC.getNamedTypeLoc(); - - if (NTL.getTypeLocClass() == clang::TypeLoc::Record) { - auto TSL = NTL.getUnqualifiedLoc().getAs(); - - const std::string TyName = - dpct::DpctGlobalInfo::getTypeName(TSL.getType()); - std::string Str = - MapNames::findReplacedName(MapNames::TypeNamesMap, TyName); - insertHeaderForTypeRule(TyName, BeginLoc); - requestHelperFeatureForTypeNames(TyName); - - if (!Str.empty()) { - emplaceTransformation( - new ReplaceToken(BeginLoc, TSL.getEndLoc(), std::move(Str))); - return; - } - } - } - - std::string TypeName = Tok.getRawIdentifier().str(); - std::string Str = - MapNames::findReplacedName(MapNames::TypeNamesMap, TypeName); - insertHeaderForTypeRule(TypeName, BeginLoc); - requestHelperFeatureForTypeNames(TypeName); - if (auto VD = DpctGlobalInfo::findAncestor(TL)) { - auto TypeStr = VD->getType().getAsString(); - if (VD->getKind() == Decl::Var && TypeStr == "dim3") { - std::string Replacement; - std::string ReplacedType = "range"; - llvm::raw_string_ostream OS(Replacement); - DpctGlobalInfo::printCtadClass( - OS, buildString(MapNames::getClNamespace(), ReplacedType), 3); - Str = OS.str(); - } - } - - if (!Str.empty()) { - SrcAPIStaticsMap[TypeName]++; - emplaceTransformation(new ReplaceToken(BeginLoc, std::move(Str))); - return; - } - } - } } REGISTER_RULE(ReplaceDim3CtorRule, PassKind::PK_Migration) diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index 0efdc55247ea..259afa090bf3 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -632,102 +632,79 @@ void ExprAnalysis::analyzeExpr(const CXXTemporaryObjectExpr *Temp) { analyzeExpr(static_cast(Temp)); } -const CXXConstructExpr *getConstructExpr(const Expr *E) { - if (const auto *C = dyn_cast_or_null(E)) { - return C; - } else if (isa(E)) { - return getConstructExpr( - dyn_cast(E)->getSubExpr()); - } else if (isa(E)) { - return getConstructExpr(dyn_cast(E)->getSubExpr()); - } else { - return nullptr; - } +void ExprAnalysis::analyzeExpr(const CXXNewExpr *New) { + analyzeType(New->getAllocatedTypeSourceInfo()->getTypeLoc()); + analyzeExpr(New->getConstructExpr()); } void ExprAnalysis::analyzeExpr(const CXXConstructExpr *Ctor) { if (Ctor->getConstructor()->getDeclName().getAsString() == "dim3") { - auto NeedChangeName = dyn_cast(Ctor) == nullptr; - // strip the top CXXConstructExpr, if there's a CXXConstructExpr further - // down + // Only handle the param of dim3 here. if (Ctor->getNumArgs() == 1) { - Ctor = getConstructExpr(Ctor->getArg(0)); - } - if (Ctor == nullptr) { - return; - } - auto Parents = dpct::DpctGlobalInfo::getContext().getParents(*Ctor); - NeedChangeName = NeedChangeName && Parents.size() == 1 && - Parents[0].get() == nullptr && - Parents[0].get() == nullptr; - std::string ArgsString; - llvm::raw_string_ostream OS(ArgsString); - if (NeedChangeName) - DpctGlobalInfo::printCtadClass(OS, MapNames::getClNamespace() + "range", - 3) - << "("; - ArgumentAnalysis A; - std::string ArgStr = ""; - for (auto Arg : Ctor->arguments()) { - A.analyze(Arg); - ArgStr = ", " + A.getReplacedString() + ArgStr; - } - ArgStr.replace(0, 2, ""); - OS << ArgStr; - if (NeedChangeName) - OS << ")"; - OS.flush(); - - CharSourceRange CSR; - if (!NeedChangeName) { + dispatch(Ctor->getArg(0)); + } else { + std::string ArgsString; + llvm::raw_string_ostream OS(ArgsString); + ArgumentAnalysis A; + std::string ArgStr = ""; + for (auto Arg : Ctor->arguments()) { + A.analyze(Arg); + ArgStr = ", " + A.getReplacedString() + ArgStr; + } + ArgStr.replace(0, 2, ""); + OS << ArgStr; + OS.flush(); + + CharSourceRange CSR; SourceRange SR = Ctor->getParenOrBraceRange(); + if (SR.isInvalid()) { - // convert to spelling location if the dim3 constructor is in a macro - // otherwise, Lexer::getLocForEndOfToken returns invalid source location - auto CtorLoc = Ctor->getLocation().isMacroID() - ? SM.getSpellingLoc(Ctor->getLocation()) - : Ctor->getLocation(); - // dim3 a; - // MACRO(... dim3 a; ...) - auto CtorEndLoc = Lexer::getLocForEndOfToken( - CtorLoc, 0, SM, DpctGlobalInfo::getContext().getLangOpts()); - CSR = CharSourceRange(SourceRange(CtorEndLoc, CtorEndLoc), false); - ArgsString = "(" + ArgsString + ")"; + auto CtorLoc = Ctor->getLocation(); + if (Ctor->getLocation().isMacroID()) { + if (isOuterMostMacro(Ctor)) { + // #define NUM 1 + // dim3 a = NUM; + auto Parens = SourceRange( + SM.getExpansionRange(Ctor->getBeginLoc()).getBegin(), + SM.getExpansionRange(Ctor->getEndLoc()).getEnd()); + CtorLoc = getRangeInRange(Parens, CallSpellingBegin, + CallSpellingEnd, false) + .first; + } else { + // convert to spelling location if the dim3 constructor is in a + // macro otherwise, Lexer::getLocForEndOfToken returns invalid + // source location + CtorLoc = SM.getSpellingLoc(Ctor->getLocation()); + } + } + if (Ctor->getArg(0)->isDefaultArgument()) { + // dim3 a; + // MACRO(... dim3 a; ...) + auto CtorEndLoc = Lexer::getLocForEndOfToken( + CtorLoc, 0, SM, DpctGlobalInfo::getContext().getLangOpts()); + CSR = CharSourceRange(SourceRange(CtorEndLoc, CtorEndLoc), false); + } else { + // Mesure the whole expression of arguments: + // dim3 a = 1 + 1; + auto Range = getStmtExpansionSourceRange(Ctor); + auto Begin = Range.getBegin(); + auto End = Range.getEnd(); + CSR = CharSourceRange::getTokenRange( + Begin, + End.getLocWithOffset(Lexer::MeasureTokenLength( + End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts()))); + } + ArgsString = "{" + ArgsString + "}"; } else { SourceRange SR1 = SourceRange(SR.getBegin().getLocWithOffset(1), SR.getEnd()); CSR = CharSourceRange(SR1, false); } - } else { - // adjust the statement to replace if top-level constructor includes the - // variable being defined - const Stmt *S = Ctor; - if (!S) { - return; - } - if (S->getBeginLoc().isMacroID() && !isOuterMostMacro(S)) { - auto Range = getDefinitionRange(S->getBeginLoc(), S->getEndLoc()); - auto Begin = Range.getBegin(); - auto End = Range.getEnd(); - End = End.getLocWithOffset(Lexer::MeasureTokenLength( - End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts())); - CSR = CharSourceRange::getTokenRange(Begin, End); - } else { - // Use getStmtExpansionSourceRange(S) to support cases like - // dim3 a = MACRO; - auto Range = getStmtExpansionSourceRange(S); - auto Begin = Range.getBegin(); - auto End = Range.getEnd(); - CSR = CharSourceRange::getTokenRange( - Begin, - End.getLocWithOffset(Lexer::MeasureTokenLength( - End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts()))); - } + auto Range = getDefinitionRange(CSR.getBegin(), CSR.getEnd()); + auto Length = SM.getDecomposedLoc(Range.getEnd()).second - + SM.getDecomposedLoc(Range.getBegin()).second; + addReplacement(Range.getBegin(), Length, ArgsString); } - auto Range = getDefinitionRange(CSR.getBegin(), CSR.getEnd()); - auto Length = SM.getDecomposedLoc(Range.getEnd()).second - - SM.getDecomposedLoc(Range.getBegin()).second; - addReplacement(Range.getBegin(), Length, ArgsString); return; } for (auto It = Ctor->arg_begin(); It != Ctor->arg_end(); It++) { @@ -998,12 +975,12 @@ inline void ExprAnalysis::analyzeExpr(const UnresolvedLookupExpr *ULE) { } void ExprAnalysis::analyzeExpr(const ExplicitCastExpr *Cast) { + analyzeType(Cast->getTypeInfoAsWritten(), Cast); if (Cast->getCastKind() == CastKind::CK_ConstructorConversion) { if (DpctGlobalInfo::getUnqualifiedTypeName(Cast->getTypeAsWritten()) == "dim3") return dispatch(Cast->getSubExpr()); } - analyzeType(Cast->getTypeInfoAsWritten(), Cast); dispatch(Cast->getSubExprAsWritten()); } diff --git a/clang/lib/DPCT/ExprAnalysis.h b/clang/lib/DPCT/ExprAnalysis.h index 54a35863ba25..e95b49f6b555 100644 --- a/clang/lib/DPCT/ExprAnalysis.h +++ b/clang/lib/DPCT/ExprAnalysis.h @@ -632,6 +632,7 @@ class ExprAnalysis { } void analyzeExpr(const CXXConstructExpr *Ctor); + void analyzeExpr(const CXXNewExpr *New); void analyzeExpr(const CXXTemporaryObjectExpr *Temp); void analyzeExpr(const CXXUnresolvedConstructExpr *Ctor); void analyzeExpr(const MemberExpr *ME); diff --git a/clang/lib/Sema/SemaInit.cpp b/clang/lib/Sema/SemaInit.cpp index 396f9fe7a01f..27665066a17e 100644 --- a/clang/lib/Sema/SemaInit.cpp +++ b/clang/lib/Sema/SemaInit.cpp @@ -7059,6 +7059,12 @@ PerformConstructorInitialization(Sema &S, ? Kind.getEqualLoc() : Kind.getLocation(); +#ifdef SYCLomatic_CUSTOMIZATION + if (Kind.isCStyleCast()) { + Loc = Kind.getRange().getEnd(); + } +#endif + if (Kind.getKind() == InitializationKind::IK_Default) { // Force even a trivial, implicit default constructor to be // semantically checked. We do this explicitly because we don't build @@ -7141,7 +7147,11 @@ PerformConstructorInitialization(Sema &S, SourceRange ParenOrBraceRange; if (IsListInitialization) ParenOrBraceRange = SourceRange(LBraceLoc, RBraceLoc); - else if (Kind.getKind() == InitializationKind::IK_Direct) + else if (Kind.getKind() == InitializationKind::IK_Direct +#ifdef SYCLomatic_CUSTOMIZATION + && !Kind.isCStyleCast() +#endif + ) ParenOrBraceRange = Kind.getParenOrBraceRange(); // If the entity allows NRVO, mark the construction as elidable diff --git a/clang/test/dpct/checkFormatAll.cu b/clang/test/dpct/checkFormatAll.cu index f5f870a4b061..ed593641fe03 100644 --- a/clang/test/dpct/checkFormatAll.cu +++ b/clang/test/dpct/checkFormatAll.cu @@ -33,8 +33,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK-NEXT: sycl::device dev_ct1; //CHECK-NEXT: sycl::queue q_ct1(dev_ct1, //CHECK-NEXT: sycl::property_list{sycl::property::queue::in_order()}); -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); -//CHECK-NEXT: sycl::range<3> threaddim = sycl::range<3>(1, 1, 32); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/checkFormatMigrated.cu b/clang/test/dpct/checkFormatMigrated.cu index 730f565dfcbe..55e56aeab3da 100644 --- a/clang/test/dpct/checkFormatMigrated.cu +++ b/clang/test/dpct/checkFormatMigrated.cu @@ -34,8 +34,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); -//CHECK-NEXT: sycl::range<3> threaddim = sycl::range<3>(1, 1, 32); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/ctad.cu b/clang/test/dpct/ctad.cu index 7da089990679..a2102b0fc758 100644 --- a/clang/test/dpct/ctad.cu +++ b/clang/test/dpct/ctad.cu @@ -20,7 +20,7 @@ int main() { // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); // range default constructor does the right thing. - // CHECK: sycl::range deflt(1, 1, 1); + // CHECK: sycl::range<3> deflt{1, 1, 1}; dim3 deflt; // CHECK: sycl::range deflt_1{0, 0, 0}; @@ -28,12 +28,12 @@ int main() { cudaExtent deflt_1; cudaPos deflt_2; - // CHECK: sycl::range round1_1(1, 1, NUM); + // CHECK: sycl::range<3> round1_1(1, 1, NUM); dim3 round1_1(NUM); cudaExtent exten = make_cudaExtent(1,1,1);; - // CHECK: sycl::range castini = sycl::range(1, 1, 4); + // CHECK: sycl::range<3> castini = (sycl::range<3>){1, 1, 4}; dim3 castini = (dim3)4; // CHECK: sycl::range castini_1 = exten; @@ -41,14 +41,14 @@ int main() { cudaExtent castini_1 = exten; cudaPos castini_2 = deflt_2; - // CHECK: sycl::range copyctor1 = sycl::range(sycl::range(1, 1, 33)); + // CHECK: sycl::range<3> copyctor1 = sycl::range<3>((sycl::range<3>){1, 1, 33}); dim3 copyctor1 = dim3((dim3)33); - // CHECK: sycl::range copyctor2 = sycl::range(copyctor1); + // CHECK: sycl::range<3> copyctor2 = sycl::range<3>(copyctor1); dim3 copyctor2 = dim3(copyctor1); - // CHECK: sycl::range copyctor3(copyctor1); + // CHECK: sycl::range<3> copyctor3(copyctor1); dim3 copyctor3(copyctor1); // CHECK: sycl::range copyctor31(exten); @@ -56,9 +56,9 @@ int main() { cudaExtent copyctor31(exten); cudaPos copyctor32(deflt_2); - // CHECK: func(sycl::range(1, 1, 1), sycl::range(1, 1, 1), sycl::range(1, 1, 2), sycl::range(1, 2, 3)); + // CHECK: func((sycl::range<3>){1, 1, 1}, sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 2), sycl::range<3>(1, 2, 3)); func((dim3)1, dim3(1), dim3(2, 1), dim3(3, 2, 1)); - // CHECK: func(deflt, sycl::range(deflt), sycl::range(deflt), sycl::range(1, 1, 2 + 3 * 3)); + // CHECK: func(deflt, sycl::range<3>(deflt), (sycl::range<3>)deflt, {1, 1, 2 + 3 * 3}); func(deflt, dim3(deflt), (dim3)deflt, 2 + 3 * 3); // CHECK: sycl::range<3> *p_extent = nullptr; @@ -85,7 +85,7 @@ int main() { dim3 **ppw; }; - // CHECK: sycl::range gpu_blocks(1, 1, 1 / (castini[2] * 200)); + // CHECK: sycl::range<3> gpu_blocks(1, 1, 1 / (castini[2] * 200)); dim3 gpu_blocks(1 / (castini.x * 200)); // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { diff --git a/clang/test/dpct/datatypes_test_part2.cu b/clang/test/dpct/datatypes_test_part2.cu index 9fdaa3f64900..872059048675 100644 --- a/clang/test/dpct/datatypes_test_part2.cu +++ b/clang/test/dpct/datatypes_test_part2.cu @@ -12,7 +12,7 @@ void case_1(void) { { -// CHECK: sycl::range<3> var1(1, 1, 1); +// CHECK: sycl::range<3> var1{1, 1, 1}; // CHECK-NEXT: sycl::range<3> *var2; // CHECK-NEXT: sycl::range<3> &var3 = var1; // CHECK-NEXT: sycl::range<3> &&var4 = std::move(var1); @@ -1293,7 +1293,7 @@ __device__ void foo_t(){ // CHECK-NEXT: #define T8_1 sycl::range<3> * // CHECK-NEXT: #define T8_2 sycl::range<3> & // CHECK-NEXT: #define T8_3 sycl::range<3> && -// CHECK-NEXT: T8_0 a1(1, 1, 1); +// CHECK-NEXT: T8_0 a1{1, 1, 1}; // CHECK-NEXT: T8_1 a2; // CHECK-NEXT: T8_2 a3=a1; // CHECK-NEXT: T8_3 a4=std::move(a1); diff --git a/clang/test/dpct/dim3.cu b/clang/test/dpct/dim3.cu index 544d39526b6a..e8e3579716bd 100644 --- a/clang/test/dpct/dim3.cu +++ b/clang/test/dpct/dim3.cu @@ -19,12 +19,12 @@ int main() { // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam2 = {0, sycl::range<3>(1, 1, 0)}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam2 = {0, {1, 1, 0}}; cudaKernelNodeParams kernelNodeParam2 = {0, 0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam3 = {0, sycl::range<3>(1, 1, 0), sycl::range<3>(1, 1, 0)}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam3 = {0, {1, 1, 0}, {1, 1, 0}}; cudaKernelNodeParams kernelNodeParam3 = {0, 0, 0}; // CHECK: /* @@ -40,11 +40,11 @@ int main() { // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam6{0, sycl::range<3>(1, 1, 0)}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam6{0, {1, 1, 0}}; cudaKernelNodeParams kernelNodeParam6{0, 0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam7{0, sycl::range<3>(1, 1, 0), sycl::range<3>(1, 1, 0)}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam7{0, {1, 1, 0}, {1, 1, 0}}; cudaKernelNodeParams kernelNodeParam7{0, 0, 0}; } diff --git a/clang/test/dpct/enable-all-extensions.cu b/clang/test/dpct/enable-all-extensions.cu index 7b24ea6f968f..53d99f6aa958 100644 --- a/clang/test/dpct/enable-all-extensions.cu +++ b/clang/test/dpct/enable-all-extensions.cu @@ -67,7 +67,7 @@ void h() { void foo1() { int n; - // CHECK: sycl::range<3> abc(1, 1, 1); + // CHECK: sycl::range<3> abc{1, 1, 1}; // CHECK-NEXT: abc[1] = std::min(std::max(512 / (unsigned int)abc[2], 1u), (unsigned int)n); // CHECK-NEXT: abc[0] = std::min(std::max(512 / ((unsigned int)abc[2] * (unsigned int)abc[1]), 1u), (unsigned int)n); dim3 abc; diff --git a/clang/test/dpct/formatIndent.cu b/clang/test/dpct/formatIndent.cu index 21d988802a53..abcc83672bb4 100644 --- a/clang/test/dpct/formatIndent.cu +++ b/clang/test/dpct/formatIndent.cu @@ -17,7 +17,7 @@ void foo(){ //CHECK:void foo1(){ //CHECK-NEXT: //some comments //CHECK-NEXT: //some comments -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; //CHECK-NEXT:} void foo1(){ //some comments @@ -27,7 +27,7 @@ void foo1(){ //CHECK:void foo2(){ //CHECK-NEXT: //some comments -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; //CHECK-NEXT:} void foo2(){ //some comments @@ -36,7 +36,7 @@ void foo2(){ //CHECK:void foo3(){ //CHECK-NEXT: int test; -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; //CHECK-NEXT:} void foo3(){ int test; diff --git a/clang/test/dpct/formatMigratedExplicitly.cu b/clang/test/dpct/formatMigratedExplicitly.cu index 77bf7b84a15d..c06203df0074 100644 --- a/clang/test/dpct/formatMigratedExplicitly.cu +++ b/clang/test/dpct/formatMigratedExplicitly.cu @@ -32,8 +32,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); -//CHECK-NEXT: sycl::range<3> threaddim = sycl::range<3>(1, 1, 32); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/formatMigratedGoogle.cu b/clang/test/dpct/formatMigratedGoogle.cu index 2563eb7d529e..7d0daa859478 100644 --- a/clang/test/dpct/formatMigratedGoogle.cu +++ b/clang/test/dpct/formatMigratedGoogle.cu @@ -32,8 +32,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); -//CHECK-NEXT: sycl::range<3> threaddim = sycl::range<3>(1, 1, 32); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/formatMigratedLLVM.cu b/clang/test/dpct/formatMigratedLLVM.cu index 3d48029e8708..3555eaa69d3d 100644 --- a/clang/test/dpct/formatMigratedLLVM.cu +++ b/clang/test/dpct/formatMigratedLLVM.cu @@ -32,8 +32,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); -//CHECK-NEXT: sycl::range<3> threaddim = sycl::range<3>(1, 1, 32); +//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/kernel-call-origcode-embedded.cu b/clang/test/dpct/kernel-call-origcode-embedded.cu index 0deffd8d922a..4e9fd6f3ea86 100644 --- a/clang/test/dpct/kernel-call-origcode-embedded.cu +++ b/clang/test/dpct/kernel-call-origcode-embedded.cu @@ -54,11 +54,11 @@ int main() { // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); // CHECK: /* DPCT_ORIG dim3 griddim = 2;*/ - // CHECK-NEXT: sycl::range<3> griddim = sycl::range<3>(1, 1, 2); + // CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; dim3 griddim = 2; // CHECK: /* DPCT_ORIG dim3 threaddim = 32;*/ - // CHECK-NEXT: sycl::range<3> threaddim = sycl::range<3>(1, 1, 32); + // CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; dim3 threaddim = 32; void *karg1 = 0; diff --git a/clang/test/dpct/math_functions_std.cu b/clang/test/dpct/math_functions_std.cu index 386f7ff98a6f..84c35cb95616 100644 --- a/clang/test/dpct/math_functions_std.cu +++ b/clang/test/dpct/math_functions_std.cu @@ -64,7 +64,7 @@ void h() { void foo1() { int n; - //CHECK: sycl::range<3> abc(1, 1, 1); + //CHECK: sycl::range<3> abc{1, 1, 1}; //CHECK-NEXT: abc[1] = std::min(std::max(512 / (unsigned int)abc[2], 1u), (unsigned int) n); //CHECK-NEXT: abc[0] = std::min(std::max(512 / ((unsigned int)abc[2] * (unsigned int)abc[1]), 1u), (unsigned int)n); dim3 abc; diff --git a/clang/test/dpct/replace-dim3.cu b/clang/test/dpct/replace-dim3.cu index e7a728c9d488..a5a0c79169ec 100644 --- a/clang/test/dpct/replace-dim3.cu +++ b/clang/test/dpct/replace-dim3.cu @@ -40,7 +40,7 @@ int main() { // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); // range default constructor does the right thing. - // CHECK: sycl::range<3> deflt(1, 1, 1); + // CHECK: sycl::range<3> deflt{1, 1, 1}; dim3 deflt; // CHECK: sycl::range<3> round1(1, 1, 1); @@ -53,14 +53,14 @@ int main() { // CHECK: sycl::range<3> round2_1(1, NUM, NUM); dim3 round2_1(NUM, NUM); - // CHECK: sycl::range<3> assign = sycl::range<3>(1, 1, 32); + // CHECK: sycl::range<3> assign = {1, 1, 32}; dim3 assign = 32; - // CHECK: sycl::range<3> assign_1 = sycl::range<3>(1, 1, NUM); + // CHECK: sycl::range<3> assign_1 = {1, 1, NUM}; dim3 assign_1 = NUM; - // CHECK: sycl::range<3> castini = sycl::range<3>(1, 1, 4); + // CHECK: sycl::range<3> castini = (sycl::range<3>){1, 1, 4}; dim3 castini = (dim3)4; - // CHECK: sycl::range<3> castini_1 = sycl::range<3>(1, 1, NUM); + // CHECK: sycl::range<3> castini_1 = (sycl::range<3>){1, 1, NUM}; dim3 castini_1 = (dim3)NUM; // CHECK: sycl::range<3> castini2 = sycl::range<3>(1, 2, 2); @@ -77,14 +77,14 @@ int main() { deflt = dim3(3); // CHECK: deflt = sycl::range<3>(1, 1, NUM); deflt = dim3(NUM); - // CHECK: deflt = sycl::range<3>(1, 1, 5); + // CHECK: deflt = {1, 1, 5}; deflt = 5; - // CHECK: deflt = sycl::range<3>(1, 1, ((NUM%32 == 0) ? NUM/32 : (NUM/32 + 1))); + // CHECK: deflt = {1, 1, ((NUM%32 == 0) ? NUM/32 : (NUM/32 + 1))}; deflt = ((NUM%32 == 0) ? NUM/32 : (NUM/32 + 1)); - // CHECK: sycl::range<3> copyctor1 = sycl::range<3>(sycl::range<3>(1, 1, 33)); + // CHECK: sycl::range<3> copyctor1 = sycl::range<3>((sycl::range<3>){1, 1, 33}); dim3 copyctor1 = dim3((dim3)33); - // CHECK: sycl::range<3> copyctor1_1 = sycl::range<3>(sycl::range<3>(1, 1, NUM)); + // CHECK: sycl::range<3> copyctor1_1 = sycl::range<3>((sycl::range<3>){1, 1, NUM}); dim3 copyctor1_1 = dim3((dim3)NUM); // CHECK: sycl::range<3> copyctor2 = sycl::range<3>(copyctor1); @@ -93,17 +93,17 @@ int main() { // CHECK: sycl::range<3> copyctor3(copyctor1); dim3 copyctor3(copyctor1); - // CHECK: func(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 2), sycl::range<3>(1, 2, 3)); + // CHECK: func((sycl::range<3>){1, 1, 1}, sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 2), sycl::range<3>(1, 2, 3)); func((dim3)1, dim3(1), dim3(2, 1), dim3(3, 2, 1)); - // CHECK: func(sycl::range<3>(1, 1, NUM), sycl::range<3>(1, 1, NUM), sycl::range<3>(1, NUM, NUM), sycl::range<3>(NUM, NUM, NUM)); + // CHECK: func((sycl::range<3>){1, 1, NUM}, sycl::range<3>(1, 1, NUM), sycl::range<3>(1, NUM, NUM), sycl::range<3>(NUM, NUM, NUM)); func((dim3)NUM, dim3(NUM), dim3(NUM, NUM), dim3(NUM, NUM, NUM)); - // CHECK: func(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 2), sycl::range<3>(1, 1, 3), sycl::range<3>(1, 1, 4)); + // CHECK: func({1, 1, 1}, {1, 1, 2}, {1, 1, 3}, {1, 1, 4}); func(1, 2, 3, 4); - // CHECK: func(sycl::range<3>(1, 1, NUM), sycl::range<3>(1, 1, NUM), sycl::range<3>(1, 1, NUM), sycl::range<3>(1, 1, NUM)); + // CHECK: func({1, 1, NUM}, {1, 1, NUM}, {1, 1, NUM}, {1, 1, NUM}); func(NUM, NUM, NUM, NUM); - // CHECK: func(deflt, sycl::range<3>(deflt), sycl::range<3>(deflt), sycl::range<3>(1, 1, 2 + 3 * 3)); + // CHECK: func(deflt, sycl::range<3>(deflt), (sycl::range<3>)deflt, {1, 1, 2 + 3 * 3}); func(deflt, dim3(deflt), (dim3)deflt, 2 + 3 * 3); - // CHECK: func(deflt, sycl::range<3>(deflt), sycl::range<3>(deflt), sycl::range<3>(1, 1, NUM + NUM * NUM)); + // CHECK: func(deflt, sycl::range<3>(deflt), (sycl::range<3>)deflt, {1, 1, NUM + NUM * NUM}); func(deflt, dim3(deflt), (dim3)deflt, NUM + NUM * NUM); // CHECK: sycl::range<3> test(3, 2, 1); @@ -262,14 +262,14 @@ __global__ void kernel_foo(float *a, wrap *mt, unsigned int N) { } // CHECK: void dim3_foo() { -// CHECK-NEXT: DIM3_DEFAULT_VAR(block0(1, 1, 1)); +// CHECK-NEXT: DIM3_DEFAULT_VAR(block0{1, 1, 1}); // CHECK-NEXT: CALL_FUNC( []() { -// CHECK-NEXT: sycl::range<3> block1(1, 1, 1); +// CHECK-NEXT: sycl::range<3> block1{1, 1, 1}; // CHECK-NEXT: sycl::range<3> block2{1, 1, 1}; // CHECK-NEXT: sycl::range<3> block3(1, 1, 2); // CHECK-NEXT: sycl::range<3> block4(1, 3, 2); // CHECK-NEXT: sycl::range<3> block5(4, 3, 2); -// CHECK-NEXT: DIM3_DEFAULT_VAR(block6(1, 1, 1)); +// CHECK-NEXT: DIM3_DEFAULT_VAR(block6{1, 1, 1}); // CHECK-NEXT: }); // CHECK-NEXT: } void dim3_foo() { diff --git a/clang/test/dpct/thrust/thrust_testing/source/foo.cu b/clang/test/dpct/thrust/thrust_testing/source/foo.cu index c8efd013b9be..8d4992ca5469 100644 --- a/clang/test/dpct/thrust/thrust_testing/source/foo.cu +++ b/clang/test/dpct/thrust/thrust_testing/source/foo.cu @@ -33,7 +33,7 @@ void baz(ForwardIterator1 first1, ForwardIterator1 last1, int main() { - // CHECK: sycl::range<3> t(1, 1, 1); + // CHECK: sycl::range<3> t{1, 1, 1}; dim3 t; return 0; } diff --git a/clang/test/dpct/types001.cu b/clang/test/dpct/types001.cu index 84d6eee49838..ff0000d22173 100644 --- a/clang/test/dpct/types001.cu +++ b/clang/test/dpct/types001.cu @@ -111,7 +111,7 @@ void my_error_checker(T ReturnValue, char const *const FuncName) { #define MY_ERROR_CHECKER(CALL) my_error_checker((CALL), #CALL) int main(int argc, char **argv) { - //CHECK:sycl::range<3> d3(1, 1, 1); + //CHECK:sycl::range<3> d3{1, 1, 1}; //CHECK-NEXT:int a = sizeof(sycl::range<3>); //CHECK-NEXT:a = sizeof(d3); //CHECK-NEXT:a = sizeof d3; From bc0fbbbc79fd52ca0f762de986d5a0dfc0fe3046 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Fri, 14 Jun 2024 17:47:59 +0800 Subject: [PATCH 4/7] Add InitListAnalysis. --- clang/lib/DPCT/ASTTraversal.cpp | 105 +++++++++++++++++++++++++------- clang/lib/DPCT/ExprAnalysis.cpp | 81 ++++++++++++++++++++++-- clang/lib/DPCT/ExprAnalysis.h | 16 +++++ clang/test/dpct/ctad.cu | 14 ++--- clang/test/dpct/dim3.cu | 12 ++-- clang/test/dpct/replace-dim3.cu | 21 +++++++ 6 files changed, 207 insertions(+), 42 deletions(-) diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 4a85bb8cc241..7a996592c60a 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -1688,7 +1688,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { MF.addMatcher( typeLoc( loc(qualType(hasDeclaration(namedDecl(hasAnyName( - "dim3", "cudaError", "curandStatus", "cublasStatus", "CUstream", + "cudaError", "curandStatus", "cublasStatus", "CUstream", "CUstream_st", "thrust::complex", "thrust::device_vector", "thrust::device_ptr", "thrust::device_reference", "thrust::host_vector", "cublasHandle_t", "CUevent_st", "__half", @@ -3012,39 +3012,98 @@ void ReplaceDim3CtorRule::registerMatcher(MatchFinder &MF) { MF.addMatcher( explicitCastExpr(hasType(namedDecl(hasName("dim3")))).bind("dim3Cast"), this); - MF.addMatcher(cxxConstructExpr(hasType(namedDecl(hasName("dim3"))), - argumentCountIs(1), - unless(hasAncestor(cxxConstructExpr( - hasType(namedDecl(hasName("dim3"))))))) - .bind("dim3Ctor"), - this); + MF.addMatcher( + cxxConstructExpr(hasType(namedDecl(hasName("dim3")))).bind("dim3Ctor"), + this); - MF.addMatcher(cxxConstructExpr(hasType(namedDecl(hasName("dim3"))), - unless(hasParent(initListExpr())), - unless(hasAncestor(cxxConstructExpr( - hasType(namedDecl(hasName("dim3"))))))) - .bind("dim3Ctor"), - this); + MF.addMatcher( + typeLoc(loc(qualType(hasDeclaration(anyOf( + namedDecl(hasAnyName("dim3")), + typedefDecl(hasAnyName("dim3"))))))) + .bind("dim3Type"), + this); } void ReplaceDim3CtorRule::runRule(const MatchFinder::MatchResult &Result) { + if (auto TL = getNodeAsType(Result, "dim3Type")) { + if (TL->getBeginLoc().isInvalid()) + return; + + auto BeginLoc = + getDefinitionRange(TL->getBeginLoc(), TL->getEndLoc()).getBegin(); + SourceManager *SM = Result.SourceManager; + + // WA for concatenated macro token + if (SM->isWrittenInScratchSpace(SM->getSpellingLoc(TL->getBeginLoc()))) { + BeginLoc = SM->getExpansionLoc(TL->getBeginLoc()); + } + + Token Tok; + auto LOpts = Result.Context->getLangOpts(); + Lexer::getRawToken(BeginLoc, Tok, *SM, LOpts, true); + if (Tok.isAnyIdentifier()) { + if (TL->getType()->isElaboratedTypeSpecifier()) { + // To handle case like "struct cudaExtent extent;" + auto ETC = TL->getUnqualifiedLoc().getAs(); + auto NTL = ETC.getNamedTypeLoc(); + + if (NTL.getTypeLocClass() == clang::TypeLoc::Record) { + auto TSL = NTL.getUnqualifiedLoc().getAs(); + + const std::string TyName = + dpct::DpctGlobalInfo::getTypeName(TSL.getType()); + std::string Str = + MapNames::findReplacedName(MapNames::TypeNamesMap, TyName); + insertHeaderForTypeRule(TyName, BeginLoc); + requestHelperFeatureForTypeNames(TyName); + + if (!Str.empty()) { + emplaceTransformation( + new ReplaceToken(BeginLoc, TSL.getEndLoc(), std::move(Str))); + return; + } + } + } + + std::string TypeName = Tok.getRawIdentifier().str(); + std::string Str = + MapNames::findReplacedName(MapNames::TypeNamesMap, TypeName); + insertHeaderForTypeRule(TypeName, BeginLoc); + requestHelperFeatureForTypeNames(TypeName); + if (auto VD = DpctGlobalInfo::findAncestor(TL)) { + auto TypeStr = VD->getType().getAsString(); + if (VD->getKind() == Decl::Var && TypeStr == "dim3") { + std::string Replacement; + std::string ReplacedType = "range"; + llvm::raw_string_ostream OS(Replacement); + DpctGlobalInfo::printCtadClass( + OS, buildString(MapNames::getClNamespace(), ReplacedType), 3); + Str = OS.str(); + } + } + + if (!Str.empty()) { + SrcAPIStaticsMap[TypeName]++; + emplaceTransformation(new ReplaceToken(BeginLoc, std::move(Str))); + return; + } + } + return; + } + const Expr *E = nullptr; if (const auto *New = getNodeAsType(Result, "dim3New")) { - if (getParentKernelCall(New->getConstructExpr())) - return; E = New; - } - if (const auto *Cast = getNodeAsType(Result, "dim3Cast")) { - if (getParentKernelCall(Cast->getSubExprAsWritten())) - return; + } else if (const auto *Cast = + getNodeAsType(Result, "dim3Cast")) { E = Cast; - } - if (const auto *Ctor = getNodeAsType(Result, "dim3Ctor")) { - if (getParentKernelCall(Ctor)) - return; + } else if (const auto *Ctor = + getNodeAsType(Result, "dim3Ctor")) { E = Ctor; } if (E) { + if (getParentKernelCall(E)) + return; ExprAnalysis EA; EA.analyze(E); emplaceTransformation(EA.getReplacement()); diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index 259afa090bf3..6e3304690132 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -639,6 +639,17 @@ void ExprAnalysis::analyzeExpr(const CXXNewExpr *New) { void ExprAnalysis::analyzeExpr(const CXXConstructExpr *Ctor) { if (Ctor->getConstructor()->getDeclName().getAsString() == "dim3") { + const auto *InitList = DpctGlobalInfo::findAncestor(Ctor); + if (InitList && + (Ctor->getParenOrBraceRange().isInvalid() || + Ctor->getBeginLoc() == Ctor->getEndLoc()) && + Ctor->getArg(0)->isDefaultArgument()) { + // Handle implicit ctor in linit list: cudaKernelNodeParams p = {0}; + InitListAnalysis ILA(InitList); + addReplacement(InitList->getBeginLoc(), InitList->getEndLoc(), + ILA.getReplacedInitListStr()); + return; + } // Only handle the param of dim3 here. if (Ctor->getNumArgs() == 1) { dispatch(Ctor->getArg(0)); @@ -976,12 +987,7 @@ inline void ExprAnalysis::analyzeExpr(const UnresolvedLookupExpr *ULE) { void ExprAnalysis::analyzeExpr(const ExplicitCastExpr *Cast) { analyzeType(Cast->getTypeInfoAsWritten(), Cast); - if (Cast->getCastKind() == CastKind::CK_ConstructorConversion) { - if (DpctGlobalInfo::getUnqualifiedTypeName(Cast->getTypeAsWritten()) == - "dim3") - return dispatch(Cast->getSubExpr()); - } - dispatch(Cast->getSubExprAsWritten()); + dispatch(Cast->getSubExpr()); } // Precondition: CE != nullptr @@ -2412,5 +2418,68 @@ void IndexAnalysis::analyzeExpr(const ParenExpr *PE) { } void IndexAnalysis::analyzeExpr(const IntegerLiteral *IL) { return; } +InitListAnalysis::InitListAnalysis(const InitListExpr *ILE) : ExprAnalysis() { + int LastDim3ImplicitArg = ILE->getNumInits() - 1; + while (LastDim3ImplicitArg >= 0) { + const auto *Init = ILE->getInit(LastDim3ImplicitArg); + const auto *Ctor = dyn_cast(Init); + if (Init->getBeginLoc() == Init->getEndLoc() && Ctor && + Ctor->getConstructor()->getDeclName().getAsString() == "dim3") + break; + --LastDim3ImplicitArg; + } + if (LastDim3ImplicitArg < 0) { + return; + } + ReplaceLoc = ILE->getBeginLoc().getLocWithOffset(1); + for (int Index = 0; Index <= LastDim3ImplicitArg; ++Index) { + const auto *Init = ILE->getInit(Index); + dispatch(Init); + } + ReplacedInitListStr.replace(1, 2, ""); + ReplacedInitListStr += "}"; +} + +void InitListAnalysis::dispatch(const Stmt *Expression) { + switch (Expression->getStmtClass()) { + ANALYZE_EXPR(ImplicitValueInitExpr) + ANALYZE_EXPR(InitListExpr) + ANALYZE_EXPR(CXXConstructExpr) + default: + if (const auto *E = dyn_cast(Expression)) { + ArgumentAnalysis A; + A.analyze(E); + ReplacedInitListStr += ", " + A.getReplacedString(); + } + } +} + +void InitListAnalysis::analyzeExpr(const ImplicitValueInitExpr *IVIE) { + ReplacedInitListStr += ", {}"; +} + +void InitListAnalysis::analyzeExpr(const InitListExpr *ILE) { + if (ILE->getBeginLoc() == ILE->getEndLoc()) { + InitListAnalysis ILA(ILE); + ReplacedInitListStr += ", " + ILA.getReplacedInitListStr(); + } else { + ArgumentAnalysis A; + A.analyze(ILE); + ReplacedInitListStr += ", " + A.getReplacedString(); + } +} + +void InitListAnalysis::analyzeExpr(const CXXConstructExpr *Ctor) { + if ((Ctor->getParenOrBraceRange().isInvalid() || + Ctor->getBeginLoc() == Ctor->getEndLoc()) && + Ctor->getArg(0)->isDefaultArgument()) { + ReplacedInitListStr += ", {1, 1, 1}"; + } else { + ArgumentAnalysis A; + A.analyze(Ctor); + ReplacedInitListStr += ", " + A.getReplacedString(); + } +} + } // namespace dpct } // namespace clang diff --git a/clang/lib/DPCT/ExprAnalysis.h b/clang/lib/DPCT/ExprAnalysis.h index e95b49f6b555..1c033bcfcd5d 100644 --- a/clang/lib/DPCT/ExprAnalysis.h +++ b/clang/lib/DPCT/ExprAnalysis.h @@ -969,6 +969,22 @@ class IndexAnalysis : public ExprAnalysis { std::stack ContainNonAdditiveOp; }; +class InitListAnalysis : private ExprAnalysis { +public: + explicit InitListAnalysis(const InitListExpr *ILE); + SourceLocation getReplaceLoc() const { return ReplaceLoc; } + std::string getReplacedInitListStr() const { return ReplacedInitListStr; } + +private: + void dispatch(const Stmt *Expression) override; + void analyzeExpr(const ImplicitValueInitExpr *IVIE); + void analyzeExpr(const InitListExpr *ILE); + void analyzeExpr(const CXXConstructExpr *Ctor); + + SourceLocation ReplaceLoc; + std::string ReplacedInitListStr = "{"; +}; + } // namespace dpct } // namespace clang diff --git a/clang/test/dpct/ctad.cu b/clang/test/dpct/ctad.cu index a2102b0fc758..641414b4f222 100644 --- a/clang/test/dpct/ctad.cu +++ b/clang/test/dpct/ctad.cu @@ -20,7 +20,7 @@ int main() { // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); // range default constructor does the right thing. - // CHECK: sycl::range<3> deflt{1, 1, 1}; + // CHECK: sycl::range deflt{1, 1, 1}; dim3 deflt; // CHECK: sycl::range deflt_1{0, 0, 0}; @@ -28,12 +28,12 @@ int main() { cudaExtent deflt_1; cudaPos deflt_2; - // CHECK: sycl::range<3> round1_1(1, 1, NUM); + // CHECK: sycl::range round1_1(1, 1, NUM); dim3 round1_1(NUM); cudaExtent exten = make_cudaExtent(1,1,1);; - // CHECK: sycl::range<3> castini = (sycl::range<3>){1, 1, 4}; + // CHECK: sycl::range castini = (sycl::range<3>){1, 1, 4}; dim3 castini = (dim3)4; // CHECK: sycl::range castini_1 = exten; @@ -41,14 +41,14 @@ int main() { cudaExtent castini_1 = exten; cudaPos castini_2 = deflt_2; - // CHECK: sycl::range<3> copyctor1 = sycl::range<3>((sycl::range<3>){1, 1, 33}); + // CHECK: sycl::range copyctor1 = sycl::range<3>((sycl::range<3>){1, 1, 33}); dim3 copyctor1 = dim3((dim3)33); - // CHECK: sycl::range<3> copyctor2 = sycl::range<3>(copyctor1); + // CHECK: sycl::range copyctor2 = sycl::range<3>(copyctor1); dim3 copyctor2 = dim3(copyctor1); - // CHECK: sycl::range<3> copyctor3(copyctor1); + // CHECK: sycl::range copyctor3(copyctor1); dim3 copyctor3(copyctor1); // CHECK: sycl::range copyctor31(exten); @@ -85,7 +85,7 @@ int main() { dim3 **ppw; }; - // CHECK: sycl::range<3> gpu_blocks(1, 1, 1 / (castini[2] * 200)); + // CHECK: sycl::range gpu_blocks(1, 1, 1 / (castini[2] * 200)); dim3 gpu_blocks(1 / (castini.x * 200)); // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { diff --git a/clang/test/dpct/dim3.cu b/clang/test/dpct/dim3.cu index e8e3579716bd..ac5ad303519c 100644 --- a/clang/test/dpct/dim3.cu +++ b/clang/test/dpct/dim3.cu @@ -9,17 +9,17 @@ int main() { // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam0 = {}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam0 = {{{\{}}}, {1, 1, 1}, {1, 1, 1}}; cudaKernelNodeParams kernelNodeParam0 = {}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam1 = {0}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam1 = {0, {1, 1, 1}, {1, 1, 1}}; cudaKernelNodeParams kernelNodeParam1 = {0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam2 = {0, {1, 1, 0}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam2 = {0, {1, 1, 0}, {1, 1, 1}}; cudaKernelNodeParams kernelNodeParam2 = {0, 0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. @@ -30,17 +30,17 @@ int main() { // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam4{}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam4{{{\{}}}, {1, 1, 1}, {1, 1, 1}}; cudaKernelNodeParams kernelNodeParam4{}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam5{0}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam5{0, {1, 1, 1}, {1, 1, 1}}; cudaKernelNodeParams kernelNodeParam5{0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam6{0, {1, 1, 0}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam6{0, {1, 1, 0}, {1, 1, 1}}; cudaKernelNodeParams kernelNodeParam6{0, 0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. diff --git a/clang/test/dpct/replace-dim3.cu b/clang/test/dpct/replace-dim3.cu index a5a0c79169ec..625886035a49 100644 --- a/clang/test/dpct/replace-dim3.cu +++ b/clang/test/dpct/replace-dim3.cu @@ -295,3 +295,24 @@ class Dim3Struct { dim3 x = dim3(3, 4); void f() { dim3(5, 6); } }; + +struct A { + int x; + dim3 y; + int z; +}; +struct B { + int x; + A y; + dim3 z; +}; + +int dim3_implicit_ctor() { + dim3 d; + d.x = 5; + B b1 = {}; + B b2 = {0}; + B b3 = {0, {}}; + B b4 = {0, {1}}; + B b5 = {0, {1, {1}}}; +} From a1efa3c63d8b0d28b7134a3d264b718bfa900f6f Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Fri, 21 Jun 2024 14:59:36 +0800 Subject: [PATCH 5/7] Add dpct::dim3. --- clang/lib/DPCT/APINamesCooperativeGroups.inc | 10 +- clang/lib/DPCT/ASTTraversal.cpp | 164 +------------ clang/lib/DPCT/ASTTraversal.h | 13 - clang/lib/DPCT/ExprAnalysis.cpp | 224 +----------------- clang/lib/DPCT/ExprAnalysis.h | 17 -- clang/lib/DPCT/MapNames.cpp | 2 +- clang/lib/Sema/SemaInit.cpp | 12 +- clang/runtime/dpct-rt/include/dpct/math.hpp | 15 ++ clang/test/dpct/checkFormatAll.cu | 4 +- clang/test/dpct/checkFormatMigrated.cu | 4 +- clang/test/dpct/compat_with_clang.cu | 2 +- clang/test/dpct/cooperative_groups2.cu | 4 +- .../dpct/cooperative_groups_thread_group.cu | 4 +- ...ative_groups_thread_group_no_free_query.cu | 4 +- clang/test/dpct/ctad.cu | 30 +-- clang/test/dpct/curand-device-usm.cu | 4 +- clang/test/dpct/datatypes_test_part2.cu | 68 +++--- .../dpct/decltype_of_vector_type_field.cu | 6 +- clang/test/dpct/device001.cu | 2 +- clang/test/dpct/dim3.cu | 16 +- .../dpct/enable-all-experimental-features.cu | 6 +- clang/test/dpct/enable-all-extensions.cu | 6 +- clang/test/dpct/formatIndent.cu | 6 +- clang/test/dpct/formatMigratedExplicitly.cu | 4 +- clang/test/dpct/formatMigratedGoogle.cu | 4 +- clang/test/dpct/formatMigratedLLVM.cu | 4 +- .../dpct/kernel-call-origcode-embedded.cu | 6 +- clang/test/dpct/kernel-call.cu | 10 +- clang/test/dpct/kernel-usm.cu | 2 +- clang/test/dpct/kernel_1d_range.cu | 4 +- clang/test/dpct/kernel_without_name.cu | 4 +- clang/test/dpct/macro_test.cu | 12 +- clang/test/dpct/math_functions_std.cu | 6 +- clang/test/dpct/replace-dim3.cu | 159 +++++++------ clang/test/dpct/template-instantiation.cu | 6 +- .../dpct/thrust/thrust_testing/source/foo.cu | 2 +- clang/test/dpct/types001.cu | 46 ++-- 37 files changed, 246 insertions(+), 646 deletions(-) diff --git a/clang/lib/DPCT/APINamesCooperativeGroups.inc b/clang/lib/DPCT/APINamesCooperativeGroups.inc index 4f2cc052ff2b..142bfe30989d 100644 --- a/clang/lib/DPCT/APINamesCooperativeGroups.inc +++ b/clang/lib/DPCT/APINamesCooperativeGroups.inc @@ -970,10 +970,12 @@ MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_group.num_threads", MemberExprBase(), false, "get_local_linear_range") MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_group.get_type", MemberExprBase(), false, "get_type") -MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.group_index", MemberExprBase(), - false, "get_group_id") -MEMBER_CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.thread_index", MemberExprBase(), - false, "get_local_id") +CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.group_index", + CALL(MapNames::getDpctNamespace() + "dim3", + MEMBER_CALL(MemberExprBase(), false, "get_group_id"))) +CALL_FACTORY_ENTRY("cooperative_groups::__v1::thread_block.thread_index", + CALL(MapNames::getDpctNamespace() + "dim3", + MEMBER_CALL(MemberExprBase(), false, "get_local_id"))) CONDITIONAL_FACTORY_ENTRY( UseNonUniformGroups, diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 7a996592c60a..c7fb20a202e0 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -1688,7 +1688,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { MF.addMatcher( typeLoc( loc(qualType(hasDeclaration(namedDecl(hasAnyName( - "cudaError", "curandStatus", "cublasStatus", "CUstream", + "dim3", "cudaError", "curandStatus", "cublasStatus", "CUstream", "CUstream_st", "thrust::complex", "thrust::device_vector", "thrust::device_ptr", "thrust::device_reference", "thrust::host_vector", "cublasHandle_t", "CUevent_st", "__half", @@ -3002,164 +3002,6 @@ void VectorTypeOperatorRule::runRule(const MatchFinder::MatchResult &Result) { REGISTER_RULE(VectorTypeOperatorRule, PassKind::PK_Migration) -void ReplaceDim3CtorRule::registerMatcher(MatchFinder &MF) { - // Find dim3 constructors which are part of different casts (representing - // different syntaxes). This includes copy constructors. All constructors - // will be visited once. - MF.addMatcher( - cxxNewExpr(hasType(pointsTo(namedDecl(hasName("dim3"))))).bind("dim3New"), - this); - MF.addMatcher( - explicitCastExpr(hasType(namedDecl(hasName("dim3")))).bind("dim3Cast"), - this); - MF.addMatcher( - cxxConstructExpr(hasType(namedDecl(hasName("dim3")))).bind("dim3Ctor"), - this); - - MF.addMatcher( - typeLoc(loc(qualType(hasDeclaration(anyOf( - namedDecl(hasAnyName("dim3")), - typedefDecl(hasAnyName("dim3"))))))) - .bind("dim3Type"), - this); -} - -void ReplaceDim3CtorRule::runRule(const MatchFinder::MatchResult &Result) { - if (auto TL = getNodeAsType(Result, "dim3Type")) { - if (TL->getBeginLoc().isInvalid()) - return; - - auto BeginLoc = - getDefinitionRange(TL->getBeginLoc(), TL->getEndLoc()).getBegin(); - SourceManager *SM = Result.SourceManager; - - // WA for concatenated macro token - if (SM->isWrittenInScratchSpace(SM->getSpellingLoc(TL->getBeginLoc()))) { - BeginLoc = SM->getExpansionLoc(TL->getBeginLoc()); - } - - Token Tok; - auto LOpts = Result.Context->getLangOpts(); - Lexer::getRawToken(BeginLoc, Tok, *SM, LOpts, true); - if (Tok.isAnyIdentifier()) { - if (TL->getType()->isElaboratedTypeSpecifier()) { - // To handle case like "struct cudaExtent extent;" - auto ETC = TL->getUnqualifiedLoc().getAs(); - auto NTL = ETC.getNamedTypeLoc(); - - if (NTL.getTypeLocClass() == clang::TypeLoc::Record) { - auto TSL = NTL.getUnqualifiedLoc().getAs(); - - const std::string TyName = - dpct::DpctGlobalInfo::getTypeName(TSL.getType()); - std::string Str = - MapNames::findReplacedName(MapNames::TypeNamesMap, TyName); - insertHeaderForTypeRule(TyName, BeginLoc); - requestHelperFeatureForTypeNames(TyName); - - if (!Str.empty()) { - emplaceTransformation( - new ReplaceToken(BeginLoc, TSL.getEndLoc(), std::move(Str))); - return; - } - } - } - - std::string TypeName = Tok.getRawIdentifier().str(); - std::string Str = - MapNames::findReplacedName(MapNames::TypeNamesMap, TypeName); - insertHeaderForTypeRule(TypeName, BeginLoc); - requestHelperFeatureForTypeNames(TypeName); - if (auto VD = DpctGlobalInfo::findAncestor(TL)) { - auto TypeStr = VD->getType().getAsString(); - if (VD->getKind() == Decl::Var && TypeStr == "dim3") { - std::string Replacement; - std::string ReplacedType = "range"; - llvm::raw_string_ostream OS(Replacement); - DpctGlobalInfo::printCtadClass( - OS, buildString(MapNames::getClNamespace(), ReplacedType), 3); - Str = OS.str(); - } - } - - if (!Str.empty()) { - SrcAPIStaticsMap[TypeName]++; - emplaceTransformation(new ReplaceToken(BeginLoc, std::move(Str))); - return; - } - } - return; - } - - const Expr *E = nullptr; - if (const auto *New = getNodeAsType(Result, "dim3New")) { - E = New; - } else if (const auto *Cast = - getNodeAsType(Result, "dim3Cast")) { - E = Cast; - } else if (const auto *Ctor = - getNodeAsType(Result, "dim3Ctor")) { - E = Ctor; - } - if (E) { - if (getParentKernelCall(E)) - return; - ExprAnalysis EA; - EA.analyze(E); - emplaceTransformation(EA.getReplacement()); - EA.applyAllSubExprRepl(); - return; - } -} - -REGISTER_RULE(ReplaceDim3CtorRule, PassKind::PK_Migration) - -// rule for dim3 types member fields replacements. -void Dim3MemberFieldsRule::registerMatcher(MatchFinder &MF) { - // dim3->x/y/z => (*dim3)[0]/[1]/[2] - // dim3.x/y/z => dim3[0]/[1]/[2] - // int64_t{dim3->x/y/z} => int64_t((*dim3)[0]/[1]/[2]) - // int64_t{dim3.x/y/z} => int64_t(dim3[0]/[1]/[2]) - auto Dim3MemberExpr = [&]() { - return memberExpr(anyOf( - has(implicitCastExpr(hasType(pointsTo(typedefDecl(hasName("dim3")))))), - hasObjectExpression(hasType(qualType(hasCanonicalType( - recordType(hasDeclaration(cxxRecordDecl(hasName("dim3")))))))))); - }; - MF.addMatcher(Dim3MemberExpr().bind("Dim3MemberExpr"), this); - MF.addMatcher( - cxxFunctionalCastExpr( - allOf(hasTypeLoc(loc(isSignedInteger())), - hasDescendant( - initListExpr(hasInit(0, ignoringImplicit(Dim3MemberExpr()))) - .bind("InitListExpr")))), - this); -} - -void Dim3MemberFieldsRule::runRule(const MatchFinder::MatchResult &Result) { - // E.g. - // dim3 *pd3, d3; - // pd3->z; d3.z; - // int64_t{d3.x}, int64_t{pd3->x}; - // will migrate to: - // (*pd3)[0]; d3[0]; - // sycl::range<3> *pd3, d3; - // int64_t(d3[0]), int64_t((*pd3)[0]); - ExprAnalysis EA; - if (const auto *ILE = getNodeAsType(Result, "InitListExpr")) { - EA.analyze(ILE); - } else if (const auto *ME = - getNodeAsType(Result, "Dim3MemberExpr")) { - EA.analyze(ME); - } else { - return; - } - emplaceTransformation(EA.getReplacement()); - EA.applyAllSubExprRepl(); -} - -REGISTER_RULE(Dim3MemberFieldsRule, PassKind::PK_Migration) - void DeviceInfoVarRule::registerMatcher(MatchFinder &MF) { MF.addMatcher( memberExpr( @@ -11854,9 +11696,7 @@ void MathFunctionsRule::registerMatcher(MatchFinder &MF) { internal::Matcher( new internal::HasNameMatcher(MathFunctionsCallExpr)), anyOf(unless(hasDeclContext(namespaceDecl(anything()))), - hasDeclContext(namespaceDecl(hasName("std")))))), - unless(hasAncestor( - cxxConstructExpr(hasType(typedefDecl(hasName("dim3"))))))) + hasDeclContext(namespaceDecl(hasName("std"))))))) .bind("math"), this); diff --git a/clang/lib/DPCT/ASTTraversal.h b/clang/lib/DPCT/ASTTraversal.h index c66f2defabb5..590db5596806 100644 --- a/clang/lib/DPCT/ASTTraversal.h +++ b/clang/lib/DPCT/ASTTraversal.h @@ -580,19 +580,6 @@ class VectorTypeOperatorRule static const char NamespaceName[]; }; -class ReplaceDim3CtorRule : public NamedMigrationRule { -public: - void registerMatcher(ast_matchers::MatchFinder &MF) override; - void runRule(const ast_matchers::MatchFinder::MatchResult &Result); -}; - -/// Migration rule for dim3 types member fields replacements. -class Dim3MemberFieldsRule : public NamedMigrationRule { -public: - void registerMatcher(ast_matchers::MatchFinder &MF) override; - void runRule(const ast_matchers::MatchFinder::MatchResult &Result); -}; - class CudaExtentRule : public NamedMigrationRule { CharSourceRange getConstructorRange(const CXXConstructExpr *Ctor); void replaceConstructor(const CXXConstructExpr *Ctor); diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index 6e3304690132..7b695fa49747 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -600,19 +600,6 @@ void ExprAnalysis::analyzeExpr(const InitListExpr *ILE) { if (QT->isPointerType()) { QT = QT->getPointeeType(); } - if (DpctGlobalInfo::getUnqualifiedTypeName( - QT->getCanonicalTypeUnqualified()) == "dim3") { - // Replace initializer list with explicit type conversion (e.g., - // 'int64_t{d3[2]}' to 'int64_t(d3[2])') to slience narrowing - // error (e.g., 'size_t -> int64_t') for - // non-constant-expression in int64_t initializer list. - // E.g., - // dim3 d3; int64_t{d3.x}; - // will be migratd to - // sycl::range<3> d3; int64_t(d3[2]); - addReplacement(ILE->getLBraceLoc(), "("); - addReplacement(ILE->getRBraceLoc(), ")"); - } } } } @@ -632,92 +619,7 @@ void ExprAnalysis::analyzeExpr(const CXXTemporaryObjectExpr *Temp) { analyzeExpr(static_cast(Temp)); } -void ExprAnalysis::analyzeExpr(const CXXNewExpr *New) { - analyzeType(New->getAllocatedTypeSourceInfo()->getTypeLoc()); - analyzeExpr(New->getConstructExpr()); -} - void ExprAnalysis::analyzeExpr(const CXXConstructExpr *Ctor) { - if (Ctor->getConstructor()->getDeclName().getAsString() == "dim3") { - const auto *InitList = DpctGlobalInfo::findAncestor(Ctor); - if (InitList && - (Ctor->getParenOrBraceRange().isInvalid() || - Ctor->getBeginLoc() == Ctor->getEndLoc()) && - Ctor->getArg(0)->isDefaultArgument()) { - // Handle implicit ctor in linit list: cudaKernelNodeParams p = {0}; - InitListAnalysis ILA(InitList); - addReplacement(InitList->getBeginLoc(), InitList->getEndLoc(), - ILA.getReplacedInitListStr()); - return; - } - // Only handle the param of dim3 here. - if (Ctor->getNumArgs() == 1) { - dispatch(Ctor->getArg(0)); - } else { - std::string ArgsString; - llvm::raw_string_ostream OS(ArgsString); - ArgumentAnalysis A; - std::string ArgStr = ""; - for (auto Arg : Ctor->arguments()) { - A.analyze(Arg); - ArgStr = ", " + A.getReplacedString() + ArgStr; - } - ArgStr.replace(0, 2, ""); - OS << ArgStr; - OS.flush(); - - CharSourceRange CSR; - SourceRange SR = Ctor->getParenOrBraceRange(); - - if (SR.isInvalid()) { - auto CtorLoc = Ctor->getLocation(); - if (Ctor->getLocation().isMacroID()) { - if (isOuterMostMacro(Ctor)) { - // #define NUM 1 - // dim3 a = NUM; - auto Parens = SourceRange( - SM.getExpansionRange(Ctor->getBeginLoc()).getBegin(), - SM.getExpansionRange(Ctor->getEndLoc()).getEnd()); - CtorLoc = getRangeInRange(Parens, CallSpellingBegin, - CallSpellingEnd, false) - .first; - } else { - // convert to spelling location if the dim3 constructor is in a - // macro otherwise, Lexer::getLocForEndOfToken returns invalid - // source location - CtorLoc = SM.getSpellingLoc(Ctor->getLocation()); - } - } - if (Ctor->getArg(0)->isDefaultArgument()) { - // dim3 a; - // MACRO(... dim3 a; ...) - auto CtorEndLoc = Lexer::getLocForEndOfToken( - CtorLoc, 0, SM, DpctGlobalInfo::getContext().getLangOpts()); - CSR = CharSourceRange(SourceRange(CtorEndLoc, CtorEndLoc), false); - } else { - // Mesure the whole expression of arguments: - // dim3 a = 1 + 1; - auto Range = getStmtExpansionSourceRange(Ctor); - auto Begin = Range.getBegin(); - auto End = Range.getEnd(); - CSR = CharSourceRange::getTokenRange( - Begin, - End.getLocWithOffset(Lexer::MeasureTokenLength( - End, SM, dpct::DpctGlobalInfo::getContext().getLangOpts()))); - } - ArgsString = "{" + ArgsString + "}"; - } else { - SourceRange SR1 = - SourceRange(SR.getBegin().getLocWithOffset(1), SR.getEnd()); - CSR = CharSourceRange(SR1, false); - } - auto Range = getDefinitionRange(CSR.getBegin(), CSR.getEnd()); - auto Length = SM.getDecomposedLoc(Range.getEnd()).second - - SM.getDecomposedLoc(Range.getBegin()).second; - addReplacement(Range.getBegin(), Length, ArgsString); - } - return; - } for (auto It = Ctor->arg_begin(); It != Ctor->arg_end(); It++) { dispatch(*It); } @@ -806,63 +708,6 @@ void ExprAnalysis::analyzeExpr(const MemberExpr *ME) { } } } - } else if (BaseType == "dim3") { - if (ME->isArrow()) { - addReplacement(ME->getBase(), "(" + getDrefName(ME->getBase()) + ")"); - } - addReplacement( - ME->getOperatorLoc(), ME->getMemberLoc(), - MapNames::findReplacedName(MapNames::Dim3MemberNamesMap, - ME->getMemberNameInfo().getAsString())); - - auto needAddTypecast = [](const Expr *E) -> bool { - auto &Context = DpctGlobalInfo::getContext(); - clang::DynTypedNodeList Parents = Context.getParents(*E); - bool hasCast = false; - while (!Parents.empty()) { - auto &Cur = Parents[0]; - if (const auto ICE = Cur.get()) { - CastKind CK = ICE->getCastKind(); - if (CK == CastKind::CK_FloatingCast || - CK == CastKind::CK_IntegralCast) { - hasCast = true; - Parents = Context.getParents(Cur); - continue; - } - } else if (Cur.get() || Cur.get() || - Cur.get()) { - hasCast = true; - Parents = Context.getParents(Cur); - continue; - } else if (const auto CE = Cur.get()) { - if (hasCast) - return false; - auto *Callee = - dyn_cast(CE->getCallee()->IgnoreParenImpCasts()); - if (!Callee) - return false; - if (CE->getDirectCallee()->isTemplateInstantiation()) - return true; - if (!Callee->getQualifier()) - return false; - if (Callee->getQualifier()->getKind() != - NestedNameSpecifier::SpecifierKind::Namespace) - return false; - if (Callee->getQualifier()->getAsNamespace()->getNameAsString() != - "std") - return false; - if (Callee->getNameInfo().getAsString() == "max" || - Callee->getNameInfo().getAsString() == "min") - return true; - return false; - } - Parents = Context.getParents(Cur); - } - return false; - }; - if (needAddTypecast(ME)) { - addReplacement(ME->getBeginLoc(), 0, "(unsigned int)"); - } } else if (BaseType == "cudaDeviceProp") { auto MemberName = ME->getMemberNameInfo().getAsString(); @@ -987,7 +832,7 @@ inline void ExprAnalysis::analyzeExpr(const UnresolvedLookupExpr *ULE) { void ExprAnalysis::analyzeExpr(const ExplicitCastExpr *Cast) { analyzeType(Cast->getTypeInfoAsWritten(), Cast); - dispatch(Cast->getSubExpr()); + dispatch(Cast->getSubExprAsWritten()); } // Precondition: CE != nullptr @@ -1352,10 +1197,6 @@ void ExprAnalysis::analyzeDecltypeType(DecltypeTypeLoc TL) { return; auto Name = getNestedNameSpecifierString(Qualifier); auto Range = getDefinitionRange(SR.getBegin(), SR.getEnd()); - // Types like 'dim3::x' should be migrated to 'size_t'. - if (Name == "dim3::") { - addReplacement(Range.getBegin(), Range.getEnd(), "size_t"); - } Name.resize(Name.length() - 2); // Remove the "::". if (MapNames::SupportedVectorTypes.count(Name)) { auto ReplacedStr = @@ -2418,68 +2259,5 @@ void IndexAnalysis::analyzeExpr(const ParenExpr *PE) { } void IndexAnalysis::analyzeExpr(const IntegerLiteral *IL) { return; } -InitListAnalysis::InitListAnalysis(const InitListExpr *ILE) : ExprAnalysis() { - int LastDim3ImplicitArg = ILE->getNumInits() - 1; - while (LastDim3ImplicitArg >= 0) { - const auto *Init = ILE->getInit(LastDim3ImplicitArg); - const auto *Ctor = dyn_cast(Init); - if (Init->getBeginLoc() == Init->getEndLoc() && Ctor && - Ctor->getConstructor()->getDeclName().getAsString() == "dim3") - break; - --LastDim3ImplicitArg; - } - if (LastDim3ImplicitArg < 0) { - return; - } - ReplaceLoc = ILE->getBeginLoc().getLocWithOffset(1); - for (int Index = 0; Index <= LastDim3ImplicitArg; ++Index) { - const auto *Init = ILE->getInit(Index); - dispatch(Init); - } - ReplacedInitListStr.replace(1, 2, ""); - ReplacedInitListStr += "}"; -} - -void InitListAnalysis::dispatch(const Stmt *Expression) { - switch (Expression->getStmtClass()) { - ANALYZE_EXPR(ImplicitValueInitExpr) - ANALYZE_EXPR(InitListExpr) - ANALYZE_EXPR(CXXConstructExpr) - default: - if (const auto *E = dyn_cast(Expression)) { - ArgumentAnalysis A; - A.analyze(E); - ReplacedInitListStr += ", " + A.getReplacedString(); - } - } -} - -void InitListAnalysis::analyzeExpr(const ImplicitValueInitExpr *IVIE) { - ReplacedInitListStr += ", {}"; -} - -void InitListAnalysis::analyzeExpr(const InitListExpr *ILE) { - if (ILE->getBeginLoc() == ILE->getEndLoc()) { - InitListAnalysis ILA(ILE); - ReplacedInitListStr += ", " + ILA.getReplacedInitListStr(); - } else { - ArgumentAnalysis A; - A.analyze(ILE); - ReplacedInitListStr += ", " + A.getReplacedString(); - } -} - -void InitListAnalysis::analyzeExpr(const CXXConstructExpr *Ctor) { - if ((Ctor->getParenOrBraceRange().isInvalid() || - Ctor->getBeginLoc() == Ctor->getEndLoc()) && - Ctor->getArg(0)->isDefaultArgument()) { - ReplacedInitListStr += ", {1, 1, 1}"; - } else { - ArgumentAnalysis A; - A.analyze(Ctor); - ReplacedInitListStr += ", " + A.getReplacedString(); - } -} - } // namespace dpct } // namespace clang diff --git a/clang/lib/DPCT/ExprAnalysis.h b/clang/lib/DPCT/ExprAnalysis.h index 1c033bcfcd5d..54a35863ba25 100644 --- a/clang/lib/DPCT/ExprAnalysis.h +++ b/clang/lib/DPCT/ExprAnalysis.h @@ -632,7 +632,6 @@ class ExprAnalysis { } void analyzeExpr(const CXXConstructExpr *Ctor); - void analyzeExpr(const CXXNewExpr *New); void analyzeExpr(const CXXTemporaryObjectExpr *Temp); void analyzeExpr(const CXXUnresolvedConstructExpr *Ctor); void analyzeExpr(const MemberExpr *ME); @@ -969,22 +968,6 @@ class IndexAnalysis : public ExprAnalysis { std::stack ContainNonAdditiveOp; }; -class InitListAnalysis : private ExprAnalysis { -public: - explicit InitListAnalysis(const InitListExpr *ILE); - SourceLocation getReplaceLoc() const { return ReplaceLoc; } - std::string getReplacedInitListStr() const { return ReplacedInitListStr; } - -private: - void dispatch(const Stmt *Expression) override; - void analyzeExpr(const ImplicitValueInitExpr *IVIE); - void analyzeExpr(const InitListExpr *ILE); - void analyzeExpr(const CXXConstructExpr *Ctor); - - SourceLocation ReplaceLoc; - std::string ReplacedInitListStr = "{"; -}; - } // namespace dpct } // namespace clang diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index d4a1abc8b14c..2df22dd9a9e2 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -111,7 +111,7 @@ void MapNames::setExplicitNamespaceMap() { {"cudaPointerAttributes", std::make_shared(getDpctNamespace() + "pointer_attributes", HelperFeatureEnum::device_ext)}, - {"dim3", std::make_shared(getClNamespace() + "range<3>")}, + {"dim3", std::make_shared(getDpctNamespace() + "dim3")}, {"int2", std::make_shared(getClNamespace() + "int2")}, {"double2", std::make_shared(getClNamespace() + "double2")}, {"__half", std::make_shared(getClNamespace() + "half")}, diff --git a/clang/lib/Sema/SemaInit.cpp b/clang/lib/Sema/SemaInit.cpp index 27665066a17e..396f9fe7a01f 100644 --- a/clang/lib/Sema/SemaInit.cpp +++ b/clang/lib/Sema/SemaInit.cpp @@ -7059,12 +7059,6 @@ PerformConstructorInitialization(Sema &S, ? Kind.getEqualLoc() : Kind.getLocation(); -#ifdef SYCLomatic_CUSTOMIZATION - if (Kind.isCStyleCast()) { - Loc = Kind.getRange().getEnd(); - } -#endif - if (Kind.getKind() == InitializationKind::IK_Default) { // Force even a trivial, implicit default constructor to be // semantically checked. We do this explicitly because we don't build @@ -7147,11 +7141,7 @@ PerformConstructorInitialization(Sema &S, SourceRange ParenOrBraceRange; if (IsListInitialization) ParenOrBraceRange = SourceRange(LBraceLoc, RBraceLoc); - else if (Kind.getKind() == InitializationKind::IK_Direct -#ifdef SYCLomatic_CUSTOMIZATION - && !Kind.isCStyleCast() -#endif - ) + else if (Kind.getKind() == InitializationKind::IK_Direct) ParenOrBraceRange = Kind.getParenOrBraceRange(); // If the entity allows NRVO, mark the construction as elidable diff --git a/clang/runtime/dpct-rt/include/dpct/math.hpp b/clang/runtime/dpct-rt/include/dpct/math.hpp index 6aeba145521f..c4b5f592419d 100644 --- a/clang/runtime/dpct-rt/include/dpct/math.hpp +++ b/clang/runtime/dpct-rt/include/dpct/math.hpp @@ -15,6 +15,21 @@ #include namespace dpct { +class dim3 { +public: + unsigned x, y, z; + + constexpr dim3(size_t x = 1, size_t y = 1, size_t z = 1) : x(x), y(y), z(z) {} + + dim3(const sycl::id<3> &r) : dim3(r[2], r[1], r[0]) {} + + operator sycl::range<3>() const { return sycl::range<3>(z, y, x); } +}; + +inline dim3 operator*(const dim3 &a, const dim3 &b) { + return dim3{a.x * b.x, a.y * b.y, a.z * b.z}; +} + namespace detail { template class vectorized_binary { diff --git a/clang/test/dpct/checkFormatAll.cu b/clang/test/dpct/checkFormatAll.cu index ed593641fe03..f84b301ae225 100644 --- a/clang/test/dpct/checkFormatAll.cu +++ b/clang/test/dpct/checkFormatAll.cu @@ -33,8 +33,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK-NEXT: sycl::device dev_ct1; //CHECK-NEXT: sycl::queue q_ct1(dev_ct1, //CHECK-NEXT: sycl::property_list{sycl::property::queue::in_order()}); -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; -//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; +//CHECK-NEXT: dpct::dim3 griddim = 2; +//CHECK-NEXT: dpct::dim3 threaddim = 32; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/checkFormatMigrated.cu b/clang/test/dpct/checkFormatMigrated.cu index 55e56aeab3da..c880b2af0d79 100644 --- a/clang/test/dpct/checkFormatMigrated.cu +++ b/clang/test/dpct/checkFormatMigrated.cu @@ -34,8 +34,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; -//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; +//CHECK-NEXT: dpct::dim3 griddim = 2; +//CHECK-NEXT: dpct::dim3 threaddim = 32; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/compat_with_clang.cu b/clang/test/dpct/compat_with_clang.cu index bdf14197a321..0cfbdfb2eb44 100644 --- a/clang/test/dpct/compat_with_clang.cu +++ b/clang/test/dpct/compat_with_clang.cu @@ -14,7 +14,7 @@ __device__ inline void foo1(__half2 *array, __half a) { } // CHECK: void foo2(int a, int b) { -// CHECK-NEXT: sycl::range<3> block{1, 1, dpct::min(512, uint32_t(a * b))}; +// CHECK-NEXT: dpct::dim3 block{dpct::min(512, uint32_t(a * b))}; // CHECK-NEXT: } void foo2(int a, int b) { dim3 block{min(512, uint32_t(a * b))}; diff --git a/clang/test/dpct/cooperative_groups2.cu b/clang/test/dpct/cooperative_groups2.cu index b60fc4efabf2..851c56b2d688 100644 --- a/clang/test/dpct/cooperative_groups2.cu +++ b/clang/test/dpct/cooperative_groups2.cu @@ -17,8 +17,8 @@ __device__ void foo() { // CHECK: auto block = item_ct1.get_group(); auto block = cg::this_thread_block(); - // CHECK: auto group_x = block.get_group_id()[2]; - // CHECK-NEXT: auto thread_x = block.get_local_id()[2]; + // CHECK: auto group_x = dpct::dim3(block.get_group_id()).x; + // CHECK-NEXT: auto thread_x = dpct::dim3(block.get_local_id()).x; auto group_x = block.group_index().x; auto thread_x = block.thread_index().x; diff --git a/clang/test/dpct/cooperative_groups_thread_group.cu b/clang/test/dpct/cooperative_groups_thread_group.cu index 8ab84d4ea5dd..d4148da5794f 100644 --- a/clang/test/dpct/cooperative_groups_thread_group.cu +++ b/clang/test/dpct/cooperative_groups_thread_group.cu @@ -23,13 +23,13 @@ __device__ void testThreadGroup(cg::thread_group g) { g.size(); auto block = cg::this_thread_block(); - // CHECK: block.get_local_id(); + // CHECK: dpct::dim3(block.get_local_id()); block.thread_index(); } __global__ void kernelFunc() { auto block = cg::this_thread_block(); - // CHECK: block.get_local_id(); + // CHECK: dpct::dim3(block.get_local_id()); block.thread_index(); // CHECK: auto threadBlockGroup = sycl::ext::oneapi::experimental::this_group<3>(); auto threadBlockGroup = cg::this_thread_block(); diff --git a/clang/test/dpct/cooperative_groups_thread_group_no_free_query.cu b/clang/test/dpct/cooperative_groups_thread_group_no_free_query.cu index 1bdb01c6e256..c28883535bc8 100644 --- a/clang/test/dpct/cooperative_groups_thread_group_no_free_query.cu +++ b/clang/test/dpct/cooperative_groups_thread_group_no_free_query.cu @@ -25,13 +25,13 @@ __device__ void testThreadGroup(cg::thread_group g) { g.size(); auto block = cg::this_thread_block(); - // CHECK: block.get_local_id(); + // CHECK: dpct::dim3(block.get_local_id()); block.thread_index(); } __global__ void kernelFunc() { auto block = cg::this_thread_block(); - // CHECK: block.get_local_id(); + // CHECK: dpct::dim3(block.get_local_id()); block.thread_index(); // CHECK: auto threadBlockGroup = item_ct1.get_group(); auto threadBlockGroup = cg::this_thread_block(); diff --git a/clang/test/dpct/ctad.cu b/clang/test/dpct/ctad.cu index 641414b4f222..c5b1abd2ff61 100644 --- a/clang/test/dpct/ctad.cu +++ b/clang/test/dpct/ctad.cu @@ -8,7 +8,7 @@ #define NUM 23 -// CHECK: void func(sycl::range<3> a, sycl::range<3> b, sycl::range<3> c, sycl::range<3> d) { +// CHECK: void func(dpct::dim3 a, dpct::dim3 b, dpct::dim3 c, dpct::dim3 d) { void func(dim3 a, dim3 b, dim3 c, dim3 d) { } @@ -20,7 +20,7 @@ int main() { // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); // range default constructor does the right thing. - // CHECK: sycl::range deflt{1, 1, 1}; + // CHECK: dpct::dim3 deflt; dim3 deflt; // CHECK: sycl::range deflt_1{0, 0, 0}; @@ -28,12 +28,12 @@ int main() { cudaExtent deflt_1; cudaPos deflt_2; - // CHECK: sycl::range round1_1(1, 1, NUM); + // CHECK: dpct::dim3 round1_1(NUM); dim3 round1_1(NUM); cudaExtent exten = make_cudaExtent(1,1,1);; - // CHECK: sycl::range castini = (sycl::range<3>){1, 1, 4}; + // CHECK: dpct::dim3 castini = (dpct::dim3)4; dim3 castini = (dim3)4; // CHECK: sycl::range castini_1 = exten; @@ -41,14 +41,14 @@ int main() { cudaExtent castini_1 = exten; cudaPos castini_2 = deflt_2; - // CHECK: sycl::range copyctor1 = sycl::range<3>((sycl::range<3>){1, 1, 33}); + // CHECK: dpct::dim3 copyctor1 = dpct::dim3((dpct::dim3)33); dim3 copyctor1 = dim3((dim3)33); - // CHECK: sycl::range copyctor2 = sycl::range<3>(copyctor1); + // CHECK: dpct::dim3 copyctor2 = dpct::dim3(copyctor1); dim3 copyctor2 = dim3(copyctor1); - // CHECK: sycl::range copyctor3(copyctor1); + // CHECK: dpct::dim3 copyctor3(copyctor1); dim3 copyctor3(copyctor1); // CHECK: sycl::range copyctor31(exten); @@ -56,17 +56,17 @@ int main() { cudaExtent copyctor31(exten); cudaPos copyctor32(deflt_2); - // CHECK: func((sycl::range<3>){1, 1, 1}, sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 2), sycl::range<3>(1, 2, 3)); + // CHECK: func((dpct::dim3)1, dpct::dim3(1), dpct::dim3(2, 1), dpct::dim3(3, 2, 1)); func((dim3)1, dim3(1), dim3(2, 1), dim3(3, 2, 1)); - // CHECK: func(deflt, sycl::range<3>(deflt), (sycl::range<3>)deflt, {1, 1, 2 + 3 * 3}); + // CHECK: func(deflt, dpct::dim3(deflt), (dpct::dim3)deflt, 2 + 3 * 3); func(deflt, dim3(deflt), (dim3)deflt, 2 + 3 * 3); // CHECK: sycl::range<3> *p_extent = nullptr; cudaExtent *p_extent = nullptr; - // CHECK: sycl::range<3> *p = &deflt; + // CHECK: dpct::dim3 *p = &deflt; dim3 *p = &deflt; - // CHECK: sycl::range<3> **pp = &p; + // CHECK: dpct::dim3 **pp = &p; dim3 **pp = &p; // CHECK: sycl::range<3> *p_1 = &deflt_1; @@ -77,15 +77,15 @@ int main() { struct container { unsigned int x, y, z; - // CHECK: sycl::range<3> w; + // CHECK: dpct::dim3 w; dim3 w; - // CHECK: sycl::range<3> *pw; + // CHECK: dpct::dim3 *pw; dim3 *pw; - // CHECK: sycl::range<3> **ppw; + // CHECK: dpct::dim3 **ppw; dim3 **ppw; }; - // CHECK: sycl::range gpu_blocks(1, 1, 1 / (castini[2] * 200)); + // CHECK: dpct::dim3 gpu_blocks(1 / (castini.x * 200)); dim3 gpu_blocks(1 / (castini.x * 200)); // CHECK: q_ct1.submit( // CHECK-NEXT: [&](sycl::handler &cgh) { diff --git a/clang/test/dpct/curand-device-usm.cu b/clang/test/dpct/curand-device-usm.cu index 454fef9fd442..a35ad49db26c 100644 --- a/clang/test/dpct/curand-device-usm.cu +++ b/clang/test/dpct/curand-device-usm.cu @@ -110,9 +110,9 @@ int main(int argc, char **argv) { CHECK(cudaMalloc((void **)&dOut, sizeof(int) * 10)); //CHECK: CHECK(DPCT_CHECK_ERROR(RandomStates = (dpct::rng::device::rng_generator> *)sycl::malloc_device(sizeof(dpct::rng::device::rng_generator>) * 10 * 10, q_ct1))); CHECK(cudaMalloc((void **)&RandomStates, sizeof(curandState) * 10 * 10)); - //CHECK: sycl::range<3> grid(1, 1, 10); + //CHECK: dpct::dim3 grid(10, 1); dim3 grid(10, 1); - //CHECK: CHECK(DPCT_CHECK_ERROR(dOut = sycl::malloc_device(grid[2], q_ct1))); + //CHECK: CHECK(DPCT_CHECK_ERROR(dOut = sycl::malloc_device(grid.x, q_ct1))); CHECK(cudaMalloc((void **)&dOut, sizeof(int) * grid.x)); return 0; diff --git a/clang/test/dpct/datatypes_test_part2.cu b/clang/test/dpct/datatypes_test_part2.cu index 872059048675..19235fb7c168 100644 --- a/clang/test/dpct/datatypes_test_part2.cu +++ b/clang/test/dpct/datatypes_test_part2.cu @@ -12,10 +12,10 @@ void case_1(void) { { -// CHECK: sycl::range<3> var1{1, 1, 1}; -// CHECK-NEXT: sycl::range<3> *var2; -// CHECK-NEXT: sycl::range<3> &var3 = var1; -// CHECK-NEXT: sycl::range<3> &&var4 = std::move(var1); +// CHECK: dpct::dim3 var1; +// CHECK-NEXT: dpct::dim3 *var2; +// CHECK-NEXT: dpct::dim3 &var3 = var1; +// CHECK-NEXT: dpct::dim3 &&var4 = std::move(var1); dim3 var1; dim3 *var2; dim3 &var3 = var1; @@ -240,8 +240,8 @@ CUstream_st *var2; // case 2 void case_2(void) { { -// CHECK: new sycl::range<3>(1, 1, 1); -// CHECK-NEXT: new sycl::range<3> *(); +// CHECK: new dpct::dim3(); +// CHECK-NEXT: new dpct::dim3 *(); new dim3(); new dim3 *(); } @@ -386,9 +386,9 @@ void case_2(void) { } // case 3 -// CHECK: sycl::range<3> foo0(); -// CHECK-NEXT: sycl::range<3> *foo1(); -// CHECK-NEXT: sycl::range<3> &foo2(); +// CHECK: dpct::dim3 foo0(); +// CHECK-NEXT: dpct::dim3 *foo1(); +// CHECK-NEXT: dpct::dim3 &foo2(); dim3 foo0(); dim3 *foo1(); dim3 &foo2(); @@ -535,10 +535,10 @@ CUstream_st *foo_2(); // case 4 template struct S {}; -// CHECK: template <> struct S> {}; -// CHECK-NEXT: template <> struct S *> {}; -// CHECK-NEXT: template <> struct S &> {}; -// CHECK-NEXT: template <> struct S &&> {}; +// CHECK: template <> struct S {}; +// CHECK-NEXT: template <> struct S {}; +// CHECK-NEXT: template <> struct S {}; +// CHECK-NEXT: template <> struct S {}; template <> struct S {}; template <> struct S {}; template <> struct S {}; @@ -710,10 +710,10 @@ template <> struct S {}; template void template_foo() {} void case_5(){ -// CHECK: template_foo>(); -// CHECK-NEXT: template_foo *>(); -// CHECK-NEXT: template_foo &>(); -// CHECK-NEXT: template_foo &&>(); +// CHECK: template_foo(); +// CHECK-NEXT: template_foo(); +// CHECK-NEXT: template_foo(); +// CHECK-NEXT: template_foo(); template_foo(); template_foo(); template_foo(); @@ -903,10 +903,10 @@ template_foo(); // case 6 -// CHECK: using UT0 = sycl::range<3>; -// CHECK-NEXT: using UT1 = sycl::range<3> *; -// CHECK-NEXT: using UT2 = sycl::range<3> &; -// CHECK-NEXT: using UT3 = sycl::range<3> &&; +// CHECK: using UT0 = dpct::dim3; +// CHECK-NEXT: using UT1 = dpct::dim3 *; +// CHECK-NEXT: using UT2 = dpct::dim3 &; +// CHECK-NEXT: using UT3 = dpct::dim3 &&; using UT0 = dim3; using UT1 = dim3 *; using UT2 = dim3 &; @@ -1095,10 +1095,10 @@ using UT_4 = CUstream_st &&; // case 7 -// CHECK: typedef sycl::range<3> T0; -// CHECK-NEXT: typedef sycl::range<3>* T1; -// CHECK-NEXT: typedef sycl::range<3>& T2; -// CHECK-NEXT: typedef sycl::range<3>&& T3; +// CHECK: typedef dpct::dim3 T0; +// CHECK-NEXT: typedef dpct::dim3* T1; +// CHECK-NEXT: typedef dpct::dim3& T2; +// CHECK-NEXT: typedef dpct::dim3&& T3; typedef dim3 T0; typedef dim3* T1; typedef dim3& T2; @@ -1289,11 +1289,11 @@ typedef CUstream_st&& T_4; __device__ void foo_t(){ { -// CHECK: #define T8_0 sycl::range<3> -// CHECK-NEXT: #define T8_1 sycl::range<3> * -// CHECK-NEXT: #define T8_2 sycl::range<3> & -// CHECK-NEXT: #define T8_3 sycl::range<3> && -// CHECK-NEXT: T8_0 a1{1, 1, 1}; +// CHECK: #define T8_0 dpct::dim3 +// CHECK-NEXT: #define T8_1 dpct::dim3 * +// CHECK-NEXT: #define T8_2 dpct::dim3 & +// CHECK-NEXT: #define T8_3 dpct::dim3 && +// CHECK-NEXT: T8_0 a1; // CHECK-NEXT: T8_1 a2; // CHECK-NEXT: T8_2 a3=a1; // CHECK-NEXT: T8_3 a4=std::move(a1); @@ -1716,10 +1716,10 @@ template void template_foo(T var) {} #define foo3(DataType) template_foo(DataType & varname) #define foo4(DataType) template_foo(DataType && varname) -// CHECK: template <> void foo1(sycl::range<3>){} -// CHECK-NEXT: template <> void foo2(sycl::range<3>){} -// CHECK-NEXT: template <> void foo3(sycl::range<3>){} -// CHECK-NEXT: template <> void foo4(sycl::range<3>){} +// CHECK: template <> void foo1(dpct::dim3){} +// CHECK-NEXT: template <> void foo2(dpct::dim3){} +// CHECK-NEXT: template <> void foo3(dpct::dim3){} +// CHECK-NEXT: template <> void foo4(dpct::dim3){} template <> void foo1(dim3){} template <> void foo2(dim3){} template <> void foo3(dim3){} diff --git a/clang/test/dpct/decltype_of_vector_type_field.cu b/clang/test/dpct/decltype_of_vector_type_field.cu index f670a837c996..6560c2a0f277 100644 --- a/clang/test/dpct/decltype_of_vector_type_field.cu +++ b/clang/test/dpct/decltype_of_vector_type_field.cu @@ -3,11 +3,11 @@ // RUN: %if build_lit %{icpx -c -fsycl %T/decltype_of_vector_type_field/decltype_of_vector_type_field.dp.cpp -o %T/decltype_of_vector_type_field/decltype_of_vector_type_field.dp.o %} void f() { - // CHECK: using dim3_x_type = size_t; + // CHECK: using dim3_x_type = decltype(dpct::dim3::x); using dim3_x_type = decltype(dim3::x); - // CHECK: using dim3_y_type = size_t; + // CHECK: using dim3_y_type = decltype(dpct::dim3::y); using dim3_y_type = decltype(dim3::y); - // CHECK: using dim3_z_type = size_t; + // CHECK: using dim3_z_type = decltype(dpct::dim3::z); using dim3_z_type = decltype(dim3::z); // CHECK: using int1_x_type = int32_t; using int1_x_type = decltype(int1::x); diff --git a/clang/test/dpct/device001.cu b/clang/test/dpct/device001.cu index ae4cf24fdb5b..2267b885cfcb 100644 --- a/clang/test/dpct/device001.cu +++ b/clang/test/dpct/device001.cu @@ -142,7 +142,7 @@ int main(int argc, char **argv) { // CHECK-NEXT:size_t share_multi_proc_mem_size = deviceProp.get_local_mem_size(); size_t share_multi_proc_mem_size = deviceProp.sharedMemPerMultiprocessor; - // CHECK: sycl::range<3> grid(1, 1, deviceProp.get_max_compute_units() * (deviceProp.get_max_work_items_per_compute_unit() / deviceProp.get_max_sub_group_size())); + // CHECK: dpct::dim3 grid(deviceProp.get_max_compute_units() * (deviceProp.get_max_work_items_per_compute_unit() / deviceProp.get_max_sub_group_size())); dim3 grid(deviceProp.multiProcessorCount * (deviceProp.maxThreadsPerMultiProcessor / deviceProp.warpSize)); // CHECK:/* diff --git a/clang/test/dpct/dim3.cu b/clang/test/dpct/dim3.cu index ac5ad303519c..e93fe3e3bb44 100644 --- a/clang/test/dpct/dim3.cu +++ b/clang/test/dpct/dim3.cu @@ -9,42 +9,42 @@ int main() { // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam0 = {{{\{}}}, {1, 1, 1}, {1, 1, 1}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam0 = {}; cudaKernelNodeParams kernelNodeParam0 = {}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam1 = {0, {1, 1, 1}, {1, 1, 1}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam1 = {0}; cudaKernelNodeParams kernelNodeParam1 = {0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam2 = {0, {1, 1, 0}, {1, 1, 1}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam2 = {0, 0}; cudaKernelNodeParams kernelNodeParam2 = {0, 0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam3 = {0, {1, 1, 0}, {1, 1, 0}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam3 = {0, 0, 0}; cudaKernelNodeParams kernelNodeParam3 = {0, 0, 0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam4{{{\{}}}, {1, 1, 1}, {1, 1, 1}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam4{}; cudaKernelNodeParams kernelNodeParam4{}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam5{0, {1, 1, 1}, {1, 1, 1}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam5{0}; cudaKernelNodeParams kernelNodeParam5{0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam6{0, {1, 1, 0}, {1, 1, 1}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam6{0, 0}; cudaKernelNodeParams kernelNodeParam6{0, 0}; // CHECK: /* // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of cudaKernelNodeParams type is not supported. // CHECK-NEXT: */ - // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam7{0, {1, 1, 0}, {1, 1, 0}}; + // CHECK-NEXT: cudaKernelNodeParams kernelNodeParam7{0, 0, 0}; cudaKernelNodeParams kernelNodeParam7{0, 0, 0}; } diff --git a/clang/test/dpct/enable-all-experimental-features.cu b/clang/test/dpct/enable-all-experimental-features.cu index 272253b6950e..0a24e9bb89a8 100644 --- a/clang/test/dpct/enable-all-experimental-features.cu +++ b/clang/test/dpct/enable-all-experimental-features.cu @@ -29,13 +29,13 @@ __device__ void testThreadGroup(cg::thread_group g) { g.size(); auto block = cg::this_thread_block(); - // CHECK: block.get_local_id(); + // CHECK: dpct::dim3(block.get_local_id()); block.thread_index(); } __global__ void kernelFunc() { auto block = cg::this_thread_block(); - // CHECK: block.get_local_id(); + // CHECK: dpct::dim3(block.get_local_id()); block.thread_index(); // CHECK: auto threadBlockGroup = sycl::ext::oneapi::experimental::this_group<3>(); auto threadBlockGroup = cg::this_thread_block(); @@ -94,7 +94,7 @@ namespace cg = cooperative_groups; __global__ void kernelFunc1() { auto block = cg::this_thread_block(); - // CHECK: block.get_local_id(); + // CHECK: dpct::dim3(block.get_local_id()); block.thread_index(); // CHECK: auto threadBlockGroup = sycl::ext::oneapi::experimental::this_group<3>(); auto threadBlockGroup = cg::this_thread_block(); diff --git a/clang/test/dpct/enable-all-extensions.cu b/clang/test/dpct/enable-all-extensions.cu index 53d99f6aa958..720504a58102 100644 --- a/clang/test/dpct/enable-all-extensions.cu +++ b/clang/test/dpct/enable-all-extensions.cu @@ -67,9 +67,9 @@ void h() { void foo1() { int n; - // CHECK: sycl::range<3> abc{1, 1, 1}; - // CHECK-NEXT: abc[1] = std::min(std::max(512 / (unsigned int)abc[2], 1u), (unsigned int)n); - // CHECK-NEXT: abc[0] = std::min(std::max(512 / ((unsigned int)abc[2] * (unsigned int)abc[1]), 1u), (unsigned int)n); + // CHECK: dpct::dim3 abc; + // CHECK-NEXT: abc.y = std::min(std::max(512 / abc.x, 1u), (unsigned int)n); + // CHECK-NEXT: abc.z = std::min(std::max(512 / (abc.x * abc.y), 1u), (unsigned int)n); dim3 abc; abc.y = std::min(std::max(512 / abc.x, 1u), (unsigned int)n); abc.z = std::min(std::max(512 / (abc.x * abc.y), 1u), (unsigned int)n); diff --git a/clang/test/dpct/formatIndent.cu b/clang/test/dpct/formatIndent.cu index abcc83672bb4..233f127d6fc7 100644 --- a/clang/test/dpct/formatIndent.cu +++ b/clang/test/dpct/formatIndent.cu @@ -17,7 +17,7 @@ void foo(){ //CHECK:void foo1(){ //CHECK-NEXT: //some comments //CHECK-NEXT: //some comments -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: dpct::dim3 griddim = 2; //CHECK-NEXT:} void foo1(){ //some comments @@ -27,7 +27,7 @@ void foo1(){ //CHECK:void foo2(){ //CHECK-NEXT: //some comments -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: dpct::dim3 griddim = 2; //CHECK-NEXT:} void foo2(){ //some comments @@ -36,7 +36,7 @@ void foo2(){ //CHECK:void foo3(){ //CHECK-NEXT: int test; -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; +//CHECK-NEXT: dpct::dim3 griddim = 2; //CHECK-NEXT:} void foo3(){ int test; diff --git a/clang/test/dpct/formatMigratedExplicitly.cu b/clang/test/dpct/formatMigratedExplicitly.cu index c06203df0074..2fe138312d2a 100644 --- a/clang/test/dpct/formatMigratedExplicitly.cu +++ b/clang/test/dpct/formatMigratedExplicitly.cu @@ -32,8 +32,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; -//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; +//CHECK-NEXT: dpct::dim3 griddim = 2; +//CHECK-NEXT: dpct::dim3 threaddim = 32; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/formatMigratedGoogle.cu b/clang/test/dpct/formatMigratedGoogle.cu index 7d0daa859478..f7a79791a44b 100644 --- a/clang/test/dpct/formatMigratedGoogle.cu +++ b/clang/test/dpct/formatMigratedGoogle.cu @@ -32,8 +32,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; -//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; +//CHECK-NEXT: dpct::dim3 griddim = 2; +//CHECK-NEXT: dpct::dim3 threaddim = 32; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/formatMigratedLLVM.cu b/clang/test/dpct/formatMigratedLLVM.cu index 3555eaa69d3d..4d5c25111e97 100644 --- a/clang/test/dpct/formatMigratedLLVM.cu +++ b/clang/test/dpct/formatMigratedLLVM.cu @@ -32,8 +32,8 @@ __global__ void testKernelPtr(const int *L, const int *M, int N) { //CHECK:int main() { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.in_order_queue(); -//CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; -//CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; +//CHECK-NEXT: dpct::dim3 griddim = 2; +//CHECK-NEXT: dpct::dim3 threaddim = 32; //CHECK-NEXT: int *karg1, *karg2; //CHECK-NEXT: karg1 = sycl::malloc_device(32, q_ct1); //CHECK-NEXT: karg2 = sycl::malloc_device(32, q_ct1); diff --git a/clang/test/dpct/kernel-call-origcode-embedded.cu b/clang/test/dpct/kernel-call-origcode-embedded.cu index 4e9fd6f3ea86..bdb3874e21e8 100644 --- a/clang/test/dpct/kernel-call-origcode-embedded.cu +++ b/clang/test/dpct/kernel-call-origcode-embedded.cu @@ -54,11 +54,11 @@ int main() { // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); // CHECK: /* DPCT_ORIG dim3 griddim = 2;*/ - // CHECK-NEXT: sycl::range<3> griddim = {1, 1, 2}; + // CHECK-NEXT: dpct::dim3 griddim = 2; dim3 griddim = 2; // CHECK: /* DPCT_ORIG dim3 threaddim = 32;*/ - // CHECK-NEXT: sycl::range<3> threaddim = {1, 1, 32}; + // CHECK-NEXT: dpct::dim3 threaddim = 32; dim3 threaddim = 32; void *karg1 = 0; @@ -130,7 +130,7 @@ int main() { // CHECK-NEXT: DPCT1049:{{[0-9]+}}: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed. // CHECK-NEXT: */ // CHECK-NEXT: q_ct1.parallel_for>( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, griddim[2]) * sycl::range<3>(1, 1, griddim[1] + 2), sycl::range<3>(1, 1, griddim[1] + 2)), + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, griddim.x) * sycl::range<3>(1, 1, griddim.y + 2), sycl::range<3>(1, 1, griddim.y + 2)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { // CHECK-NEXT: testKernel(karg1int, karg2int, karg3int, item_ct1); // CHECK-NEXT: }); diff --git a/clang/test/dpct/kernel-call.cu b/clang/test/dpct/kernel-call.cu index cf6d04f6426e..9764f08d7d39 100644 --- a/clang/test/dpct/kernel-call.cu +++ b/clang/test/dpct/kernel-call.cu @@ -198,7 +198,7 @@ int main() { // CHECK-NEXT: auto arr_karg3int_ct2 = arr[karg3int]; // CHECK-EMPTY: // CHECK-NEXT: cgh.parallel_for>( - // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, griddim[2]) * sycl::range<3>(1, 1, griddim[1] + 2), sycl::range<3>(1, 1, griddim[1] + 2)), + // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, griddim.x) * sycl::range<3>(1, 1, griddim.y + 2), sycl::range<3>(1, 1, griddim.y + 2)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { // CHECK-NEXT: testKernel(karg1int, karg2int, item_ct1, arr_karg3int_ct2); // CHECK-NEXT: }); @@ -293,7 +293,7 @@ int *g_a; __global__ void foo_kernel3(int *d) { d[0]; } -//CHECK:void run_foo(sycl::range<3> c, sycl::range<3> d) { +//CHECK:void run_foo(dpct::dim3 c, dpct::dim3 d) { //CHECK-NEXT: if (1) //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { @@ -310,7 +310,7 @@ void run_foo(dim3 c, dim3 d) { if (1) foo_kernel3<<>>(&g_a[0]); } -//CHECK:void run_foo2(sycl::range<3> c, sycl::range<3> d) { +//CHECK:void run_foo2(dpct::dim3 c, dpct::dim3 d) { //CHECK-NEXT: dpct::device_ext &dev_ct1 = dpct::get_current_device(); //CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); //CHECK-NEXT: if (1) @@ -345,7 +345,7 @@ void run_foo2(dim3 c, dim3 d) { else foo_kernel3<<>>(g_a); } -//CHECK:void run_foo3(sycl::range<3> c, sycl::range<3> d) { +//CHECK:void run_foo3(dpct::dim3 c, dpct::dim3 d) { //CHECK-NEXT: for (;;) //CHECK-NEXT: /* //CHECK-NEXT: DPCT1049:{{[0-9]+}}: The work-group size passed to the SYCL kernel may exceed the limit. To get the device limit, query info::device::max_work_group_size. Adjust the work-group size if needed. @@ -365,7 +365,7 @@ void run_foo3(dim3 c, dim3 d) { for (;;) foo_kernel3<<>>(g_a); } -//CHECK:void run_foo4(sycl::range<3> c, sycl::range<3> d) { +//CHECK:void run_foo4(dpct::dim3 c, dpct::dim3 d) { //CHECK-NEXT: while (1) //CHECK-NEXT: dpct::get_out_of_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { diff --git a/clang/test/dpct/kernel-usm.cu b/clang/test/dpct/kernel-usm.cu index 4959a17296c1..670d26d76583 100644 --- a/clang/test/dpct/kernel-usm.cu +++ b/clang/test/dpct/kernel-usm.cu @@ -247,7 +247,7 @@ int *g_a; __global__ void foo_kernel3(int *d) { } -//CHECK:void run_foo(sycl::range<3> c, sycl::range<3> d) { +//CHECK:void run_foo(dpct::dim3 c, dpct::dim3 d) { //CHECK-NEXT: if (1) //CHECK-NEXT: dpct::get_in_order_queue().submit( //CHECK-NEXT: [&](sycl::handler &cgh) { diff --git a/clang/test/dpct/kernel_1d_range.cu b/clang/test/dpct/kernel_1d_range.cu index 76ce2dde45df..df8d01e20e31 100644 --- a/clang/test/dpct/kernel_1d_range.cu +++ b/clang/test/dpct/kernel_1d_range.cu @@ -647,8 +647,8 @@ int query_block(const int x) { void foo7() { int n = 128; - //CHECK:sycl::range<3> block(1, 1, n); - //CHECK-NEXT:sycl::range<3> grid(1, 1, query_block(n)); + //CHECK:dpct::dim3 block(n); + //CHECK-NEXT:dpct::dim3 grid(query_block(n)); dim3 block(n); dim3 grid(query_block(n)); //CHECK:dpct::get_in_order_queue().parallel_for( diff --git a/clang/test/dpct/kernel_without_name.cu b/clang/test/dpct/kernel_without_name.cu index 7c3cc8196a67..f3a484c95ecf 100644 --- a/clang/test/dpct/kernel_without_name.cu +++ b/clang/test/dpct/kernel_without_name.cu @@ -224,7 +224,7 @@ void run_foo6() { dim3 grid; //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: auto grid_x_grid_y_ct0 = grid[2] * grid[1]; + //CHECK-NEXT: auto grid_x_grid_y_ct0 = grid.x * grid.y; //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), @@ -235,7 +235,7 @@ void run_foo6() { foo_kernel5<<<1, 1>>>(grid.x * grid.y); //CHECK:q_ct1.submit( //CHECK-NEXT: [&](sycl::handler &cgh) { - //CHECK-NEXT: auto grid_x_ct0 = ++grid[2]; + //CHECK-NEXT: auto grid_x_ct0 = ++grid.x; //CHECK-EMPTY: //CHECK-NEXT: cgh.parallel_for( //CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), diff --git a/clang/test/dpct/macro_test.cu b/clang/test/dpct/macro_test.cu index ebf7ac0ce590..6862447c1a70 100644 --- a/clang/test/dpct/macro_test.cu +++ b/clang/test/dpct/macro_test.cu @@ -58,7 +58,7 @@ public: #define CALL(x) x; #define EMPTY_MACRO(x) x -//CHECK:#define GET_MEMBER_MACRO(x) x[1] = 5 +//CHECK:#define GET_MEMBER_MACRO(x) x.y = 5 #define GET_MEMBER_MACRO(x) x.y = 5 __global__ void foo_kernel() {} @@ -99,9 +99,9 @@ void foo() { #endif - // CHECK: (*d3.A)[2] = 3; - // CHECK-NEXT: d3.B[2] = 2; - // CHECK-NEXT: EMPTY_MACRO(d3.B[2]); + // CHECK: d3.A->x = 3; + // CHECK-NEXT: d3.B.x = 2; + // CHECK-NEXT: EMPTY_MACRO(d3.B.x); // CHECK-NEXT: GET_MEMBER_MACRO(d3.B); d3.A->x = 3; d3.B.x = 2; @@ -268,7 +268,7 @@ MACRO_KC2(griddim,threaddim,1,0) // CHECK: MACRO_KC2(3,2,1,0) MACRO_KC2(3,2,1,0) -// CHECK: MACRO_KC2(sycl::range<3>(5, 4, 3), 2, 1, 0) +// CHECK: MACRO_KC2(dpct::dim3(5, 4, 3), 2, 1, 0) MACRO_KC2(dim3(5,4,3),2,1,0) int *a; @@ -1355,7 +1355,7 @@ void foo38() { template void foo38(T *t); -//CHECK: #define GRID grid[2] = 3; +//CHECK: #define GRID grid.x = 3; #define GRID grid.x = 3; template diff --git a/clang/test/dpct/math_functions_std.cu b/clang/test/dpct/math_functions_std.cu index 84c35cb95616..2d4433cf2dfb 100644 --- a/clang/test/dpct/math_functions_std.cu +++ b/clang/test/dpct/math_functions_std.cu @@ -64,9 +64,9 @@ void h() { void foo1() { int n; - //CHECK: sycl::range<3> abc{1, 1, 1}; - //CHECK-NEXT: abc[1] = std::min(std::max(512 / (unsigned int)abc[2], 1u), (unsigned int) n); - //CHECK-NEXT: abc[0] = std::min(std::max(512 / ((unsigned int)abc[2] * (unsigned int)abc[1]), 1u), (unsigned int)n); + //CHECK: dpct::dim3 abc; + //CHECK-NEXT: abc.y = std::min(std::max(512 / abc.x, 1u), (unsigned int) n); + //CHECK-NEXT: abc.z = std::min(std::max(512 / (abc.x * abc.y), 1u), (unsigned int)n); dim3 abc; abc.y = std::min(std::max(512 / abc.x, 1u), (unsigned int) n); abc.z = std::min(std::max(512 / (abc.x * abc.y), 1u), (unsigned int)n); diff --git a/clang/test/dpct/replace-dim3.cu b/clang/test/dpct/replace-dim3.cu index 625886035a49..8ac34b85e002 100644 --- a/clang/test/dpct/replace-dim3.cu +++ b/clang/test/dpct/replace-dim3.cu @@ -11,26 +11,26 @@ #define NUM 23 #define CALL_FUNC(func) func() -// CHECK: #define DIM3_DEFAULT_VAR(name) sycl::range<3> name +// CHECK: #define DIM3_DEFAULT_VAR(name) dpct::dim3 name #define DIM3_DEFAULT_VAR(name) dim3 name -// CHECK: void func(sycl::range<3> a, sycl::range<3> b, sycl::range<3> c, sycl::range<3> d) { +// CHECK: void func(dpct::dim3 a, dpct::dim3 b, dpct::dim3 c, dpct::dim3 d) { void func(dim3 a, dim3 b, dim3 c, dim3 d) { } -// CHECK: void test(const sycl::range<3>& a, const sycl::range<3>& b) { +// CHECK: void test(const dpct::dim3& a, const dpct::dim3& b) { void test(const dim3& a, const dim3& b) { } -// CHECK: void test(sycl::range<3>&& a, sycl::range<3>&& b) { +// CHECK: void test(dpct::dim3&& a, dpct::dim3&& b) { void test(dim3&& a, dim3&& b) { } -// CHECK: void test(const sycl::range<3>* a, const sycl::range<3>* b) { +// CHECK: void test(const dpct::dim3* a, const dpct::dim3* b) { void test(const dim3* a, const dim3* b) { } -// CHECK: void test(const sycl::range<3>** a, const sycl::range<3>** b) { +// CHECK: void test(const dpct::dim3** a, const dpct::dim3** b) { void test(const dim3** a, const dim3** b) { } @@ -40,148 +40,148 @@ int main() { // CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device(); // CHECK-NEXT: sycl::queue &q_ct1 = dev_ct1.out_of_order_queue(); // range default constructor does the right thing. - // CHECK: sycl::range<3> deflt{1, 1, 1}; + // CHECK: dpct::dim3 deflt; dim3 deflt; - // CHECK: sycl::range<3> round1(1, 1, 1); + // CHECK: dpct::dim3 round1(1); dim3 round1(1); - // CHECK: sycl::range<3> round1_1(1, 1, NUM); + // CHECK: dpct::dim3 round1_1(NUM); dim3 round1_1(NUM); - // CHECK: sycl::range<3> round2(1, 1, 2); + // CHECK: dpct::dim3 round2(2, 1); dim3 round2(2, 1); - // CHECK: sycl::range<3> round2_1(1, NUM, NUM); + // CHECK: dpct::dim3 round2_1(NUM, NUM); dim3 round2_1(NUM, NUM); - // CHECK: sycl::range<3> assign = {1, 1, 32}; + // CHECK: dpct::dim3 assign = 32; dim3 assign = 32; - // CHECK: sycl::range<3> assign_1 = {1, 1, NUM}; + // CHECK: dpct::dim3 assign_1 = NUM; dim3 assign_1 = NUM; - // CHECK: sycl::range<3> castini = (sycl::range<3>){1, 1, 4}; + // CHECK: dpct::dim3 castini = (dpct::dim3)4; dim3 castini = (dim3)4; - // CHECK: sycl::range<3> castini_1 = (sycl::range<3>){1, 1, NUM}; + // CHECK: dpct::dim3 castini_1 = (dpct::dim3)NUM; dim3 castini_1 = (dim3)NUM; - // CHECK: sycl::range<3> castini2 = sycl::range<3>(1, 2, 2); + // CHECK: dpct::dim3 castini2 = dpct::dim3(2, 2); dim3 castini2 = dim3(2, 2); - // CHECK: sycl::range<3> castini2_1 = sycl::range<3>(1, NUM, NUM); + // CHECK: dpct::dim3 castini2_1 = dpct::dim3(NUM, NUM); dim3 castini2_1 = dim3(NUM, NUM); - // CHECK: sycl::range<3> castini3 = sycl::range<3>(10, 1, 3); + // CHECK: dpct::dim3 castini3 = dpct::dim3(3, 1, 10); dim3 castini3 = dim3(3, 1, 10); - // CHECK: sycl::range<3> castini3_1 = sycl::range<3>(NUM, NUM, NUM); + // CHECK: dpct::dim3 castini3_1 = dpct::dim3(NUM, NUM, NUM); dim3 castini3_1 = dim3(NUM, NUM, NUM); - // CHECK: deflt = sycl::range<3>(1, 1, 3); + // CHECK: deflt = dpct::dim3(3); deflt = dim3(3); - // CHECK: deflt = sycl::range<3>(1, 1, NUM); + // CHECK: deflt = dpct::dim3(NUM); deflt = dim3(NUM); - // CHECK: deflt = {1, 1, 5}; + // CHECK: deflt = 5; deflt = 5; - // CHECK: deflt = {1, 1, ((NUM%32 == 0) ? NUM/32 : (NUM/32 + 1))}; + // CHECK: deflt = ((NUM%32 == 0) ? NUM/32 : (NUM/32 + 1)); deflt = ((NUM%32 == 0) ? NUM/32 : (NUM/32 + 1)); - // CHECK: sycl::range<3> copyctor1 = sycl::range<3>((sycl::range<3>){1, 1, 33}); + // CHECK: dpct::dim3 copyctor1 = dpct::dim3((dpct::dim3)33); dim3 copyctor1 = dim3((dim3)33); - // CHECK: sycl::range<3> copyctor1_1 = sycl::range<3>((sycl::range<3>){1, 1, NUM}); + // CHECK: dpct::dim3 copyctor1_1 = dpct::dim3((dpct::dim3)NUM); dim3 copyctor1_1 = dim3((dim3)NUM); - // CHECK: sycl::range<3> copyctor2 = sycl::range<3>(copyctor1); + // CHECK: dpct::dim3 copyctor2 = dpct::dim3(copyctor1); dim3 copyctor2 = dim3(copyctor1); - // CHECK: sycl::range<3> copyctor3(copyctor1); + // CHECK: dpct::dim3 copyctor3(copyctor1); dim3 copyctor3(copyctor1); - // CHECK: func((sycl::range<3>){1, 1, 1}, sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 2), sycl::range<3>(1, 2, 3)); + // CHECK: func((dpct::dim3)1, dpct::dim3(1), dpct::dim3(2, 1), dpct::dim3(3, 2, 1)); func((dim3)1, dim3(1), dim3(2, 1), dim3(3, 2, 1)); - // CHECK: func((sycl::range<3>){1, 1, NUM}, sycl::range<3>(1, 1, NUM), sycl::range<3>(1, NUM, NUM), sycl::range<3>(NUM, NUM, NUM)); + // CHECK: func((dpct::dim3)NUM, dpct::dim3(NUM), dpct::dim3(NUM, NUM), dpct::dim3(NUM, NUM, NUM)); func((dim3)NUM, dim3(NUM), dim3(NUM, NUM), dim3(NUM, NUM, NUM)); - // CHECK: func({1, 1, 1}, {1, 1, 2}, {1, 1, 3}, {1, 1, 4}); + // CHECK: func(1, 2, 3, 4); func(1, 2, 3, 4); - // CHECK: func({1, 1, NUM}, {1, 1, NUM}, {1, 1, NUM}, {1, 1, NUM}); + // CHECK: func(NUM, NUM, NUM, NUM); func(NUM, NUM, NUM, NUM); - // CHECK: func(deflt, sycl::range<3>(deflt), (sycl::range<3>)deflt, {1, 1, 2 + 3 * 3}); + // CHECK: func(deflt, dpct::dim3(deflt), (dpct::dim3)deflt, 2 + 3 * 3); func(deflt, dim3(deflt), (dim3)deflt, 2 + 3 * 3); - // CHECK: func(deflt, sycl::range<3>(deflt), (sycl::range<3>)deflt, {1, 1, NUM + NUM * NUM}); + // CHECK: func(deflt, dpct::dim3(deflt), (dpct::dim3)deflt, NUM + NUM * NUM); func(deflt, dim3(deflt), (dim3)deflt, NUM + NUM * NUM); - // CHECK: sycl::range<3> test(3, 2, 1); + // CHECK: dpct::dim3 test(1, 2, 3); dim3 test(1, 2, 3); - // CHECK: sycl::range<3> test_1(NUM, NUM, NUM); + // CHECK: dpct::dim3 test_1(NUM, NUM, NUM); dim3 test_1(NUM, NUM, NUM); - // CHECK: int b = test[2] + test[1] + test [0]; + // CHECK: int b = test.x + test. y + test .z; int b = test.x + test. y + test .z; - // CHECK: sycl::range<3> *p = &test; + // CHECK: dpct::dim3 *p = &test; dim3 *p = &test; - // CHECK: sycl::range<3> **pp = &p; + // CHECK: dpct::dim3 **pp = &p; dim3 **pp = &p; - // CHECK: int a = (*p)[2] + (*p)[1] + (*p)[0]; + // CHECK: int a = p->x + p->y + p->z; int a = p->x + p->y + p->z; - // CHECK: int aa = (*(*pp))[2] + (*(*pp))[1] + (*(*pp))[0]; + // CHECK: int aa = (*pp)->x + (*pp)->y + (*pp)->z; int aa = (*pp)->x + (*pp)->y + (*pp)->z; struct container { unsigned int x, y, z; - // CHECK: sycl::range<3> w; + // CHECK: dpct::dim3 w; dim3 w; - // CHECK: sycl::range<3> *pw; + // CHECK: dpct::dim3 *pw; dim3 *pw; - // CHECK: sycl::range<3> **ppw; + // CHECK: dpct::dim3 **ppw; dim3 **ppw; }; typedef struct container container; container t; - // CHECK: int c = t.w[2] + t.w[1] + t.w[0]; + // CHECK: int c = t.w.x + t.w.y + t.w.z; int c = t.w.x + t.w.y + t.w.z; - // CHECK: int c2 = (*t.pw)[2] + (*t.pw)[1] + (*t.pw)[0]; + // CHECK: int c2 = t.pw->x + t.pw->y + t.pw->z; int c2 = t.pw->x + t.pw->y + t.pw->z; - // CHECK: int c3 = (*(*t.ppw))[2] + (*(*t.ppw))[1] + (*(*t.ppw))[0]; + // CHECK: int c3 = (*t.ppw)->x + (*t.ppw)->y + (*t.ppw)->z; int c3 = (*t.ppw)->x + (*t.ppw)->y + (*t.ppw)->z; - // CHECK: sycl::range<3> d3_1(1, 1, test[2]); + // CHECK: dpct::dim3 d3_1(test.x); dim3 d3_1(test.x); - // CHECK: sycl::range<3> d3_2(1, 1, test[2] + 1); + // CHECK: dpct::dim3 d3_2(test.x + 1); dim3 d3_2(test.x + 1); - // CHECK: sycl::range<3> d3_2_1(1, 1, static_cast(test[2] + 32)); + // CHECK: dpct::dim3 d3_2_1(static_cast(test.x + 32)); dim3 d3_2_1(static_cast(test.x + 32)); - // CHECK: sycl::range<3> d3_2_2(1, 1, test[2] + NUM); + // CHECK: dpct::dim3 d3_2_2(test.x + NUM); dim3 d3_2_2(test.x + NUM); - // CHECK: sycl::range<3> d3_3(1, 1, 2 + test[2] + 1); + // CHECK: dpct::dim3 d3_3(2 + test.x + 1); dim3 d3_3(2 + test.x + 1); - // CHECK: sycl::range<3> d3_3_1(1, 1, 32 + test[2] + 64); + // CHECK: dpct::dim3 d3_3_1(32 + test.x + 64); dim3 d3_3_1(32 + test.x + 64); - // CHECK: sycl::range<3> d3_3_2(1, 1, NUM + test[2] + NUM); + // CHECK: dpct::dim3 d3_3_2(NUM + test.x + NUM); dim3 d3_3_2(NUM + test.x + NUM); - // CHECK: sycl::range<3> d3_4(1, test[1], test[2]); + // CHECK: dpct::dim3 d3_4(test.x, test.y); dim3 d3_4(test.x, test.y); - // CHECK: sycl::range<3> d3_5(test[0], test[1], test[2]); + // CHECK: dpct::dim3 d3_5(test.x, test.y, test.z); dim3 d3_5(test.x, test.y, test.z); - // CHECK: sycl::range<3> d3_6 = sycl::range<3>(3 + test[0] + 4, 2 + test[1], test[2] + 1); + // CHECK: dpct::dim3 d3_6 = dpct::dim3(test.x + 1, 2 + test.y, 3 + test.z + 4); dim3 d3_6 = dim3(test.x + 1, 2 + test.y, 3 + test.z + 4); - // CHECK: sycl::range<3> d3_6_1 = sycl::range<3>(113 + test[0] + 114, 112 + test[1], test[2] + 111); + // CHECK: dpct::dim3 d3_6_1 = dpct::dim3(test.x + 111, 112 + test.y, 113 + test.z + 114); dim3 d3_6_1 = dim3(test.x + 111, 112 + test.y, 113 + test.z + 114); - // CHECK: sycl::range<3> d3_6_2 = sycl::range<3>(NUM + test[0] + NUM, NUM + test[1], test[2] + NUM); + // CHECK: dpct::dim3 d3_6_2 = dpct::dim3(test.x + NUM, NUM + test.y, NUM + test.z + NUM); dim3 d3_6_2 = dim3(test.x + NUM, NUM + test.y, NUM + test.z + NUM); - // todoCHECK: sycl::range<3> d3_6_3 = sycl::range<3>(NUM + test[0] + NUM, NUM + test[1], sycl::ceil(test[2] + NUM)); + // todoCHECK: dpct::dim3 d3_6_3 = dpct::dim3(ceil(test.x + NUM), NUM + test.y, NUM + test.z + NUM); dim3 d3_6_3 = dim3(ceil(test.x + NUM), NUM + test.y, NUM + test.z + NUM); - // CHECK: sycl::range<3> gpu_blocks(1, 1, 1 / (d3_6_3[2] * 200)); + // CHECK: dpct::dim3 gpu_blocks(1 / (d3_6_3.x * 200)); dim3 gpu_blocks(1 / (d3_6_3.x * 200)); // CHECK: q_ct1.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel(d3_6[2]); + // CHECK-NEXT: kernel(d3_6.x); // CHECK-NEXT: }); kernel<<<1, 1>>>(d3_6.x); // CHECK: q_ct1.parallel_for>( // CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, NUM) * sycl::range<3>(1, 1, NUM), sycl::range<3>(1, 1, NUM)), // CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - // CHECK-NEXT: kernel(d3_6[2]); + // CHECK-NEXT: kernel(d3_6.x); // CHECK-NEXT: }); kernel<<>>(d3_6.x); } @@ -195,11 +195,11 @@ __host__ __device__ T getgriddim(T totallen, T blockdim) template static void memsetCuda(T * d_mem, T v, int n) { - // CHECK: sycl::range<3> dimBlock(1, 1, 256); - // CHECK: sycl::range<3> dimGrid_2(1, 1, std::max(2048, 3)); - // CHECK: sycl::range<3> dimGrid_1(1, 1, std::max(2048, 3)); - // CHECK: std::min(2048, getgriddim(n, dimBlock[2])); - // CHECK: sycl::range<3> dimGrid(1, 1, std::min(2048, getgriddim(n, dimBlock[2]))); + // CHECK: dpct::dim3 dimBlock(256); + // CHECK: dpct::dim3 dimGrid_2(std::max(2048, 3)); + // CHECK: dpct::dim3 dimGrid_1(std::max(2048, 3)); + // CHECK: std::min(2048, getgriddim(n, dimBlock.x)); + // CHECK: dpct::dim3 dimGrid(std::min(2048, getgriddim(n, dimBlock.x))); dim3 dimBlock(256); dim3 dimGrid_2(max(2048, 3)); dim3 dimGrid_1(std::max(2048, 3)); @@ -262,14 +262,14 @@ __global__ void kernel_foo(float *a, wrap *mt, unsigned int N) { } // CHECK: void dim3_foo() { -// CHECK-NEXT: DIM3_DEFAULT_VAR(block0{1, 1, 1}); +// CHECK-NEXT: DIM3_DEFAULT_VAR(block0); // CHECK-NEXT: CALL_FUNC( []() { -// CHECK-NEXT: sycl::range<3> block1{1, 1, 1}; -// CHECK-NEXT: sycl::range<3> block2{1, 1, 1}; -// CHECK-NEXT: sycl::range<3> block3(1, 1, 2); -// CHECK-NEXT: sycl::range<3> block4(1, 3, 2); -// CHECK-NEXT: sycl::range<3> block5(4, 3, 2); -// CHECK-NEXT: DIM3_DEFAULT_VAR(block6{1, 1, 1}); +// CHECK-NEXT: dpct::dim3 block1; +// CHECK-NEXT: dpct::dim3 block2{}; +// CHECK-NEXT: dpct::dim3 block3(2); +// CHECK-NEXT: dpct::dim3 block4(2,3); +// CHECK-NEXT: dpct::dim3 block5(2,3,4); +// CHECK-NEXT: DIM3_DEFAULT_VAR(block6); // CHECK-NEXT: }); // CHECK-NEXT: } void dim3_foo() { @@ -286,9 +286,9 @@ void dim3_foo() { #endif // CHECK: class Dim3Struct { -// CHECK-NEXT: Dim3Struct() : x(sycl::range<3>(1, 2, 1)) {} -// CHECK-NEXT: sycl::range<3> x = sycl::range<3>(1, 4, 3); -// CHECK-NEXT: void f() { sycl::range<3>(1, 6, 5); } +// CHECK-NEXT: Dim3Struct() : x(dpct::dim3(1, 2)) {} +// CHECK-NEXT: dpct::dim3 x = dpct::dim3(3, 4); +// CHECK-NEXT: void f() { dpct::dim3(5, 6); } // CHECK-NEXT: }; class Dim3Struct { Dim3Struct() : x(dim3(1, 2)) {} @@ -310,9 +310,14 @@ struct B { int dim3_implicit_ctor() { dim3 d; d.x = 5; + // CHECK: B b1 = {}; B b1 = {}; + // CHECK: B b2 = {0}; B b2 = {0}; + // CHECK: B b3 = {0, {}}; B b3 = {0, {}}; + // CHECK: B b4 = {0, {1}}; B b4 = {0, {1}}; + // CHECK: B b5 = {0, {1, {1}}}; B b5 = {0, {1, {1}}}; } diff --git a/clang/test/dpct/template-instantiation.cu b/clang/test/dpct/template-instantiation.cu index 858702345dbb..0baa365c1145 100644 --- a/clang/test/dpct/template-instantiation.cu +++ b/clang/test/dpct/template-instantiation.cu @@ -136,11 +136,11 @@ int main() { unsigned u; dim3 dim; - // CHECK: func_2_same_pram(u, (unsigned int)dim[1]); + // CHECK: func_2_same_pram(u, dim.y); func_2_same_pram(u, dim.y); - // CHECK: func_2_same_pram(u, (unsigned int)dim[1] + 1); + // CHECK: func_2_same_pram(u, dim.y + 1); func_2_same_pram(u, dim.y + 1); - // CHECK: func_2_same_pram(u, func_same_return((unsigned int)dim[1])); + // CHECK: func_2_same_pram(u, func_same_return(dim.y)); func_2_same_pram(u, func_same_return(dim.y)); } diff --git a/clang/test/dpct/thrust/thrust_testing/source/foo.cu b/clang/test/dpct/thrust/thrust_testing/source/foo.cu index 8d4992ca5469..916c26d46242 100644 --- a/clang/test/dpct/thrust/thrust_testing/source/foo.cu +++ b/clang/test/dpct/thrust/thrust_testing/source/foo.cu @@ -33,7 +33,7 @@ void baz(ForwardIterator1 first1, ForwardIterator1 last1, int main() { - // CHECK: sycl::range<3> t{1, 1, 1}; + // CHECK: dpct::dim3 t; dim3 t; return 0; } diff --git a/clang/test/dpct/types001.cu b/clang/test/dpct/types001.cu index ff0000d22173..fda32b0c145a 100644 --- a/clang/test/dpct/types001.cu +++ b/clang/test/dpct/types001.cu @@ -41,11 +41,11 @@ const cudaError *perrors1[23]; // CHECK: const dpct::err0 **pperrors1[23]; const cudaError **pperrors1[23]; -// CHECK: sycl::range<3> dims[23]; +// CHECK: dpct::dim3 dims[23]; dim3 dims[23]; -// CHECK: const sycl::range<3> *pdims[23]; +// CHECK: const dpct::dim3 *pdims[23]; const dim3 *pdims[23]; -// CHECK: const sycl::range<3> **ppdims[23]; +// CHECK: const dpct::dim3 **ppdims[23]; const dim3 **ppdims[23]; struct s { @@ -70,11 +70,11 @@ struct s { // CHECK: const dpct::err0 **pperrors1[23]; const cudaError **pperrors1[23]; - // CHECK: sycl::range<3> dims[23]; + // CHECK: dpct::dim3 dims[23]; dim3 dims[23]; - // CHECK: const sycl::range<3> *pdims[23]; + // CHECK: const dpct::dim3 *pdims[23]; const dim3 *pdims[23]; - // CHECK: const sycl::range<3> **ppdims[23]; + // CHECK: const dpct::dim3 **ppdims[23]; const dim3 **ppdims[23]; }; @@ -111,8 +111,8 @@ void my_error_checker(T ReturnValue, char const *const FuncName) { #define MY_ERROR_CHECKER(CALL) my_error_checker((CALL), #CALL) int main(int argc, char **argv) { - //CHECK:sycl::range<3> d3{1, 1, 1}; - //CHECK-NEXT:int a = sizeof(sycl::range<3>); + //CHECK:dpct::dim3 d3; + //CHECK-NEXT:int a = sizeof(dpct::dim3); //CHECK-NEXT:a = sizeof(d3); //CHECK-NEXT:a = sizeof d3; dim3 d3; @@ -609,15 +609,15 @@ void foo_2(cudaDataType_t a1, cudaDataType a2, cublasDataType_t a3) { } __device__ void foo_3() { - // CHECK: sycl::range<3> d3 = {3, 2, 1}, *pd3 = &d3; + // CHECK: dpct::dim3 d3 = {1, 2, 3}, *pd3 = &d3; dim3 d3 = {1, 2, 3}, *pd3 = &d3; int64_t m = 0; - // CHECK: m = std::min(m, int64_t((*pd3)[2])); - // CHECK-NEXT: m = std::min(m, int64_t((*pd3)[1])); - // CHECK-NEXT: m = std::min(m, int64_t((*pd3)[0])); - // CHECK-NEXT: m = std::min(m, int64_t(d3[2])); - // CHECK-NEXT: m = std::min(m, int64_t(d3[1])); - // CHECK-NEXT: m = std::min(m, int64_t(d3[0])); + // CHECK: m = std::min(m, int64_t{pd3->x}); + // CHECK-NEXT: m = std::min(m, int64_t{pd3->y}); + // CHECK-NEXT: m = std::min(m, int64_t{pd3->z}); + // CHECK-NEXT: m = std::min(m, int64_t{d3.x}); + // CHECK-NEXT: m = std::min(m, int64_t{d3.y}); + // CHECK-NEXT: m = std::min(m, int64_t{d3.z}); m = std::min(m, int64_t{pd3->x}); m = std::min(m, int64_t{pd3->y}); m = std::min(m, int64_t{pd3->z}); @@ -634,28 +634,28 @@ constexpr inline integer ceil_div(integer n, integer m) { void foo_4() { const int64_t num_irows = 32; const int64_t num_orows = 32; - // CHECK: sycl::range<3> threads(1, 1, 32); + // CHECK: dpct::dim3 threads(32); dim3 threads(32); int64_t maxGridDim = 1024; - // CHECK: sycl::range<3> grid_1(1, std::min(maxGridDim, ceil_div(num_irows, int64_t(threads[2]))), std::min(maxGridDim, num_orows)); + // CHECK: dpct::dim3 grid_1(std::min(maxGridDim, num_orows), std::min(maxGridDim, ceil_div(num_irows, int64_t{threads.x}))); dim3 grid_1(std::min(maxGridDim, num_orows), std::min(maxGridDim, ceil_div(num_irows, int64_t{threads.x}))); int row_size = 16; - // CHECK: sycl::range<3> grid_2(1, 1, std::min(maxGridDim, ceil_div(row_size, int(threads[1])))); + // CHECK: dpct::dim3 grid_2(std::min(maxGridDim, ceil_div(row_size, int(threads.y)))); dim3 grid_2(std::min(maxGridDim, ceil_div(row_size, int(threads.y)))); - // CHECK: int64_t m = int64_t(threads[1]); + // CHECK: int64_t m = int64_t{threads.y}; int64_t m = int64_t{threads.y}; - // CHECK: m = int64_t(threads[1]); + // CHECK: m = int64_t{threads.y}; m = int64_t{threads.y}; typedef int64_t MY_INT64; - // CHECK: m = std::min(int64_t(threads[2]), MY_INT64(threads[0])); + // CHECK: m = std::min(int64_t{threads.x}, MY_INT64{threads.z}); m = std::min(int64_t{threads.x}, MY_INT64{threads.z}); int num = 1024; // CHECK: m = int64_t{num}; m = int64_t{num}; - // CHECK: m = std::min(int64_t(threads[2]), MY_INT64{num}); + // CHECK: m = std::min(int64_t{threads.x}, MY_INT64{num}); m = std::min(int64_t{threads.x}, MY_INT64{num}); struct CFoo { @@ -665,7 +665,7 @@ void foo_4() { }; // CHECK: CFoo cfoo{num}; CFoo cfoo{num}; - // CHECK: m = std::min(int64_t(threads[2]), int64_t{cfoo}); + // CHECK: m = std::min(int64_t{threads.x}, int64_t{cfoo}); m = std::min(int64_t{threads.x}, int64_t{cfoo}); } From 50bc68cde8c69f5ba9dd8de2d625098a86396108 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Mon, 24 Jun 2024 16:21:27 +0800 Subject: [PATCH 6/7] Move dim3 to util.hpp --- clang/runtime/dpct-rt/include/dpct/math.hpp | 15 --------------- clang/runtime/dpct-rt/include/dpct/util.hpp | 14 ++++++++++++++ 2 files changed, 14 insertions(+), 15 deletions(-) diff --git a/clang/runtime/dpct-rt/include/dpct/math.hpp b/clang/runtime/dpct-rt/include/dpct/math.hpp index c4b5f592419d..6aeba145521f 100644 --- a/clang/runtime/dpct-rt/include/dpct/math.hpp +++ b/clang/runtime/dpct-rt/include/dpct/math.hpp @@ -15,21 +15,6 @@ #include namespace dpct { -class dim3 { -public: - unsigned x, y, z; - - constexpr dim3(size_t x = 1, size_t y = 1, size_t z = 1) : x(x), y(y), z(z) {} - - dim3(const sycl::id<3> &r) : dim3(r[2], r[1], r[0]) {} - - operator sycl::range<3>() const { return sycl::range<3>(z, y, x); } -}; - -inline dim3 operator*(const dim3 &a, const dim3 &b) { - return dim3{a.x * b.x, a.y * b.y, a.z * b.z}; -} - namespace detail { template class vectorized_binary { diff --git a/clang/runtime/dpct-rt/include/dpct/util.hpp b/clang/runtime/dpct-rt/include/dpct/util.hpp index 74c979499d19..f87af05cb3da 100644 --- a/clang/runtime/dpct-rt/include/dpct/util.hpp +++ b/clang/runtime/dpct-rt/include/dpct/util.hpp @@ -31,6 +31,20 @@ T __spirv_GroupNonUniformShuffleUp(__spv::Scope::Flag, T, unsigned) noexcept; #endif namespace dpct { +class dim3 { +public: + unsigned x, y, z; + + constexpr dim3(size_t x = 1, size_t y = 1, size_t z = 1) : x(x), y(y), z(z) {} + + dim3(const sycl::id<3> &r) : dim3(r[2], r[1], r[0]) {} + + operator sycl::range<3>() const { return sycl::range<3>(z, y, x); } +}; + +inline dim3 operator*(const dim3 &a, const dim3 &b) { + return dim3{a.x * b.x, a.y * b.y, a.z * b.z}; +} namespace detail { From 35a4d6552d63fe5e1af2edfb9a3c1b29d7dcb5ab Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Thu, 4 Jul 2024 11:33:34 +0800 Subject: [PATCH 7/7] fix comment. --- clang/runtime/dpct-rt/include/dpct/util.hpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/runtime/dpct-rt/include/dpct/util.hpp b/clang/runtime/dpct-rt/include/dpct/util.hpp index f87af05cb3da..e2e21ef42a7e 100644 --- a/clang/runtime/dpct-rt/include/dpct/util.hpp +++ b/clang/runtime/dpct-rt/include/dpct/util.hpp @@ -31,11 +31,13 @@ T __spirv_GroupNonUniformShuffleUp(__spv::Scope::Flag, T, unsigned) noexcept; #endif namespace dpct { +/// dim3 is used to store 3 component dimensions. class dim3 { public: unsigned x, y, z; - constexpr dim3(size_t x = 1, size_t y = 1, size_t z = 1) : x(x), y(y), z(z) {} + constexpr dim3(unsigned x = 1, unsigned y = 1, unsigned z = 1) + : x(x), y(y), z(z) {} dim3(const sycl::id<3> &r) : dim3(r[2], r[1], r[0]) {}