From: Manuel Klimek Date: Tue, 5 Aug 2014 09:45:53 +0000 (+0000) Subject: Adds AST matchers for matching CUDA declarations. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=a5052804254f4c04922b95148225527a5428fb0a;p=clang Adds AST matchers for matching CUDA declarations. Patch by Jacques Pienaar. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@214852 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/docs/LibASTMatchersReference.html b/docs/LibASTMatchersReference.html index 49880539bc..82c26a527d 100644 --- a/docs/LibASTMatchersReference.html +++ b/docs/LibASTMatchersReference.html @@ -319,6 +319,16 @@ usingDecl() matches using X::x +Matcher<Decl>usingDirectiveDeclMatcher<UsingDirectiveDecl>... +
Matches using namespace declarations.
+
+Given
+  namespace X { int x; }
+  using namespace X;
+usingDirectiveDecl()
+  matches using namespace X
+ + Matcher<Decl>varDeclMatcher<VarDecl>... 
Matches variable declarations.
 
@@ -355,6 +365,14 @@ nestedNameSpecifier()
+Matcher<Stmt>CUDAKernelCallExprMatcher<CUDAKernelCallExpr>... +
Matches CUDA kernel call expression.
+
+Example matches kernel<<<i,j>>>()
+  kernel<<<i,j>>>();
+
+ + Matcher<Stmt>arraySubscriptExprMatcher<ArraySubscriptExpr>... 
Matches array subscript expressions.
 
@@ -711,7 +729,7 @@ Given
   int a[] = { 1, 2 };
   struct B { int x, y; };
   B b = { 5, 6 };
-initList()
+initListExpr()
   matches "{ 1, 2 }" and "{ 5, 6 }"
@@ -876,6 +894,18 @@ Example matches "abcd", L"abcd" +Matcher<Stmt>substNonTypeTemplateParmExprMatcher<SubstNonTypeTemplateParmExpr>... +
Matches substitutions of non-type template parameters.
+
+Given
+  template <int N>
+  struct A { static const int n = N; };
+  struct B : public A<42> {};
+substNonTypeTemplateParmExpr()
+  matches "N" in the right-hand side of "static const int n = N;"
+
+ + Matcher<Stmt>switchCaseMatcher<SwitchCase>... 
Matches case and default statements inside switch statements.
 
@@ -1269,6 +1299,7 @@ Given
     int a[] = { 2, 3 }
     int b[42];
     int c[a[0]];
+  }
 variableArrayType()
   matches "int c[a[0]]"
@@ -1381,26 +1412,6 @@ constructorDecl(hasAnyConstructorInitializer(isWritten())) -Matcher<CXXMethodDecl>hasOverloadedOperatorNameStringRef Name -
Matches overloaded operator names.
-
-Matches overloaded operator names specified in strings without the
-"operator" prefix: e.g. "<<".
-
-Given:
-  class A { int operator*(); };
-  const A &operator<<(const A &a, const A &b);
-  A a;
-  a << a;   <-- This matches
-
-operatorCallExpr(hasOverloadedOperatorName("<<"))) matches the specified
-line and recordDecl(hasMethod(hasOverloadedOperatorName("*"))) matches
-the declaration of A.
-
-Usable as: Matcher<CXXOperatorCallExpr>, Matcher<CXXMethodDecl>
-
- - Matcher<CXXMethodDecl>isConst 
Matches if the given method declaration is const.
 
@@ -1470,7 +1481,7 @@ operatorCallExpr(hasOverloadedOperatorName("<<"))) matches the specified
 line and recordDecl(hasMethod(hasOverloadedOperatorName("*"))) matches
 the declaration of A.
 
-Usable as: Matcher<CXXOperatorCallExpr>, Matcher<CXXMethodDecl>
+Usable as: Matcher<CXXOperatorCallExpr>, Matcher<FunctionDecl>
@@ -1601,10 +1612,30 @@ and reference to that variable declaration within a compound statement. -Matcher<Decl>equalsNodeDecl *Node -
Matches if a node equals another node.
+Matcher<Decl>hasCudaDeviceAttr
+
Matches declaration that has CUDA device attribute.
+
+Given
+  __attribute__((device)) void f() { ... }
+matches the function declaration of f.
+

+
+
+Matcher<Decl>hasCudaGlobalAttr
+
Matches declaration that has CUDA global attribute.
 
-Decl has pointer identity in the AST.
+Given
+  __attribute__((global)) void f() { ... }
+matches the function declaration of f.
+

+
+
+Matcher<Decl>hasCudaHostAttr
+
Matches declaration that has CUDA host attribute.
+
+Given
+  __attribute__((host)) void f() { ... }
+matches the function declaration of f.
 

 
 
@@ -1667,6 +1698,26 @@ Usable as: Matcher<FunctionDecl>hasOverloadedOperatorNameStringRef Name
+
Matches overloaded operator names.
+
+Matches overloaded operator names specified in strings without the
+"operator" prefix: e.g. "<<".
+
+Given:
+  class A { int operator*(); };
+  const A &operator<<(const A &a, const A &b);
+  A a;
+  a << a;   <-- This matches
+
+operatorCallExpr(hasOverloadedOperatorName("<<"))) matches the specified
+line and recordDecl(hasMethod(hasOverloadedOperatorName("*"))) matches
+the declaration of A.
+
+Usable as: Matcher<CXXOperatorCallExpr>, Matcher<FunctionDecl>
+

+
+
 Matcher<FunctionDecl>isDefinition
 
Matches if a declaration has a body attached.
 
@@ -1900,14 +1951,6 @@ and reference to that variable declaration within a compound statement.
 

 
 
-Matcher<Stmt>equalsNodeStmt *Node
-
Matches if a node equals another node.
-
-Stmt has pointer identity in the AST.
-
-

-
-
 Matcher<TagDecl>isDefinition
 
Matches if a declaration has a body attached.
 
@@ -2550,7 +2593,7 @@ given matcher.
 
 Example matches y.x() (matcher = callExpr(callee(methodDecl(hasName("x")))))
   class Y { public: void x(); };
-  void z() { Y y; y.x();
+  void z() { Y y; y.x(); }
 

 
 
@@ -3448,6 +3491,7 @@ Example matches X &x and const X &y
     void a(X b) {
       X &x = b;
       const X &y = b;
+    }
   };
diff --git a/include/clang/ASTMatchers/ASTMatchers.h b/include/clang/ASTMatchers/ASTMatchers.h index 1ff4ab367f..30550e0b06 100644 --- a/include/clang/ASTMatchers/ASTMatchers.h +++ b/include/clang/ASTMatchers/ASTMatchers.h @@ -1844,7 +1844,7 @@ AST_MATCHER_P(CallExpr, callee, internal::Matcher, /// Example matches y.x() (matcher = callExpr(callee(methodDecl(hasName("x"))))) /// \code /// class Y { public: void x(); }; -/// void z() { Y y; y.x(); +/// void z() { Y y; y.x(); } /// \endcode AST_MATCHER_P_OVERLOAD(CallExpr, callee, internal::Matcher, InnerMatcher, 1) { @@ -1952,6 +1952,7 @@ AST_MATCHER_P_OVERLOAD(QualType, pointsTo, internal::Matcher, /// void a(X b) { /// X &x = b; /// const X &y = b; +/// } /// }; /// \endcode AST_MATCHER_P(QualType, references, internal::Matcher, @@ -3094,6 +3095,7 @@ AST_TYPE_MATCHER(IncompleteArrayType, incompleteArrayType); /// int a[] = { 2, 3 } /// int b[42]; /// int c[a[0]]; +/// } /// \endcode /// variableArrayType() /// matches "int c[a[0]]" @@ -3643,6 +3645,48 @@ AST_MATCHER_P(CaseStmt, hasCaseConstant, internal::Matcher, return InnerMatcher.matches(*Node.getLHS(), Finder, Builder); } +/// \brief Matches CUDA kernel call expression. +/// +/// Example matches kernel<<>>() +/// \code +/// kernel<<>>(); +/// \endcode +const internal::VariadicDynCastAllOfMatcher + CUDAKernelCallExpr; + +/// \brief Matches declaration that has CUDA device attribute. +/// +/// Given +/// \code +/// __attribute__((device)) void f() { ... } +/// \endcode +/// matches the function declaration of f. +AST_MATCHER(Decl, hasCudaDeviceAttr) { + return Node.hasAttr(); +} + +/// \brief Matches declaration that has CUDA host attribute. +/// +/// Given +/// \code +/// __attribute__((host)) void f() { ... } +/// \endcode +/// matches the function declaration of f. +AST_MATCHER(Decl, hasCudaHostAttr) { + return Node.hasAttr(); +} + +/// \brief Matches declaration that has CUDA global attribute. +/// +/// Given +/// \code +/// __attribute__((global)) void f() { ... } +/// \endcode +/// matches the function declaration of f. +AST_MATCHER(Decl, hasCudaGlobalAttr) { + return Node.hasAttr(); +} + } // end namespace ast_matchers } // end namespace clang diff --git a/unittests/ASTMatchers/ASTMatchersTest.cpp b/unittests/ASTMatchers/ASTMatchersTest.cpp index e769e887bd..bf4dee82bd 100644 --- a/unittests/ASTMatchers/ASTMatchersTest.cpp +++ b/unittests/ASTMatchers/ASTMatchersTest.cpp @@ -649,6 +649,24 @@ TEST(DeclarationMatcher, HasDescendantMemoization) { EXPECT_TRUE(matches("void f() { int i; }", CannotMemoize)); } +TEST(DeclarationMatcher, MatchCudaDecl) { + EXPECT_TRUE(matchesWithCuda("__global__ void f() { }" + "void g() { f<<<1, 2>>>(); }", + CUDAKernelCallExpr())); + EXPECT_TRUE(matchesWithCuda("__attribute__((device)) void f() {}", + hasCudaDeviceAttr())); + EXPECT_TRUE(matchesWithCuda("__attribute__((host)) void f() {}", + hasCudaHostAttr())); + EXPECT_TRUE(matchesWithCuda("__attribute__((global)) void f() {}", + hasCudaGlobalAttr())); + EXPECT_FALSE(matchesWithCuda("void f() {}", + hasCudaGlobalAttr())); + EXPECT_TRUE(notMatchesWithCuda("void f() {}", + hasCudaGlobalAttr())); + EXPECT_FALSE(notMatchesWithCuda("__attribute__((global)) void f() {}", + hasCudaGlobalAttr())); +} + // Implements a run method that returns whether BoundNodes contains a // Decl bound to Id that can be dynamically cast to T. // Optionally checks that the check succeeded a specific number of times. diff --git a/unittests/ASTMatchers/ASTMatchersTest.h b/unittests/ASTMatchers/ASTMatchersTest.h index 2e4ee2c5cb..8919ceae5d 100644 --- a/unittests/ASTMatchers/ASTMatchersTest.h +++ b/unittests/ASTMatchers/ASTMatchersTest.h @@ -103,6 +103,72 @@ testing::AssertionResult notMatches(const std::string &Code, return matchesConditionally(Code, AMatcher, false, "-std=c++11"); } +// Function based on matchesConditionally with "-x cuda" argument added and +// small CUDA header prepended to the code string. +template +testing::AssertionResult matchesConditionallyWithCuda( + const std::string &Code, const T &AMatcher, bool ExpectMatch, + llvm::StringRef CompileArg) { + const std::string CudaHeader = + "typedef unsigned int size_t;\n" + "#define __constant__ __attribute__((constant))\n" + "#define __device__ __attribute__((device))\n" + "#define __global__ __attribute__((global))\n" + "#define __host__ __attribute__((host))\n" + "#define __shared__ __attribute__((shared))\n" + "struct dim3 {" + " unsigned x, y, z;" + " __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1)" + " : x(x), y(y), z(z) {}" + "};" + "typedef struct cudaStream *cudaStream_t;" + "int cudaConfigureCall(dim3 gridSize, dim3 blockSize," + " size_t sharedSize = 0," + " cudaStream_t stream = 0);"; + + bool Found = false, DynamicFound = false; + MatchFinder Finder; + VerifyMatch VerifyFound(nullptr, &Found); + Finder.addMatcher(AMatcher, &VerifyFound); + VerifyMatch VerifyDynamicFound(nullptr, &DynamicFound); + if (!Finder.addDynamicMatcher(AMatcher, &VerifyDynamicFound)) + return testing::AssertionFailure() << "Could not add dynamic matcher"; + std::unique_ptr Factory( + newFrontendActionFactory(&Finder)); + // Some tests use typeof, which is a gnu extension. + std::vector Args({"-xcuda", CompileArg}); + if (!runToolOnCodeWithArgs(Factory->create(), + CudaHeader + Code, Args)) { + return testing::AssertionFailure() << "Parsing error in \"" << Code << "\""; + } + if (Found != DynamicFound) { + return testing::AssertionFailure() << "Dynamic match result (" + << DynamicFound + << ") does not match static result (" + << Found << ")"; + } + if (!Found && ExpectMatch) { + return testing::AssertionFailure() + << "Could not find match in \"" << Code << "\""; + } else if (Found && !ExpectMatch) { + return testing::AssertionFailure() + << "Found unexpected match in \"" << Code << "\""; + } + return testing::AssertionSuccess(); +} + +template +testing::AssertionResult matchesWithCuda(const std::string &Code, + const T &AMatcher) { + return matchesConditionallyWithCuda(Code, AMatcher, true, "-std=c++11"); +} + +template +testing::AssertionResult notMatchesWithCuda(const std::string &Code, + const T &AMatcher) { + return matchesConditionallyWithCuda(Code, AMatcher, false, "-std=c++11"); +} + template testing::AssertionResult matchAndVerifyResultConditionally(const std::string &Code, const T &AMatcher,