r214852 - Adds AST matchers for matching CUDA declarations.
Manuel Klimek
klimek at google.com
Tue Aug 5 02:45:54 PDT 2014
Author: klimek
Date: Tue Aug 5 04:45:53 2014
New Revision: 214852
URL: http://llvm.org/viewvc/llvm-project?rev=214852&view=rev
Log:
Adds AST matchers for matching CUDA declarations.
Patch by Jacques Pienaar.
Modified:
cfe/trunk/docs/LibASTMatchersReference.html
cfe/trunk/include/clang/ASTMatchers/ASTMatchers.h
cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.cpp
cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.h
Modified: cfe/trunk/docs/LibASTMatchersReference.html
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/docs/LibASTMatchersReference.html?rev=214852&r1=214851&r2=214852&view=diff
==============================================================================
--- cfe/trunk/docs/LibASTMatchersReference.html (original)
+++ cfe/trunk/docs/LibASTMatchersReference.html Tue Aug 5 04:45:53 2014
@@ -319,6 +319,16 @@ usingDecl()
matches using X::x </pre></td></tr>
+<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>></td><td class="name" onclick="toggle('usingDirectiveDecl0')"><a name="usingDirectiveDecl0Anchor">usingDirectiveDecl</a></td><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1UsingDirectiveDecl.html">UsingDirectiveDecl</a>>...</td></tr>
+<tr><td colspan="4" class="doc" id="usingDirectiveDecl0"><pre>Matches using namespace declarations.
+
+Given
+ namespace X { int x; }
+ using namespace X;
+usingDirectiveDecl()
+ matches using namespace X </pre></td></tr>
+
+
<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>></td><td class="name" onclick="toggle('varDecl0')"><a name="varDecl0Anchor">varDecl</a></td><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1VarDecl.html">VarDecl</a>>...</td></tr>
<tr><td colspan="4" class="doc" id="varDecl0"><pre>Matches variable declarations.
@@ -355,6 +365,14 @@ nestedNameSpecifier()
</pre></td></tr>
+<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Stmt.html">Stmt</a>></td><td class="name" onclick="toggle('CUDAKernelCallExpr0')"><a name="CUDAKernelCallExpr0Anchor">CUDAKernelCallExpr</a></td><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>>...</td></tr>
+<tr><td colspan="4" class="doc" id="CUDAKernelCallExpr0"><pre>Matches CUDA kernel call expression.
+
+Example matches kernel<<<i,j>>>()
+ kernel<<<i,j>>>();
+</pre></td></tr>
+
+
<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Stmt.html">Stmt</a>></td><td class="name" onclick="toggle('arraySubscriptExpr0')"><a name="arraySubscriptExpr0Anchor">arraySubscriptExpr</a></td><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1ArraySubscriptExpr.html">ArraySubscriptExpr</a>>...</td></tr>
<tr><td colspan="4" class="doc" id="arraySubscriptExpr0"><pre>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 }"
</pre></td></tr>
@@ -876,6 +894,18 @@ Example matches "abcd", L"abcd"
</pre></td></tr>
+<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Stmt.html">Stmt</a>></td><td class="name" onclick="toggle('substNonTypeTemplateParmExpr0')"><a name="substNonTypeTemplateParmExpr0Anchor">substNonTypeTemplateParmExpr</a></td><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1SubstNonTypeTemplateParmExpr.html">SubstNonTypeTemplateParmExpr</a>>...</td></tr>
+<tr><td colspan="4" class="doc" id="substNonTypeTemplateParmExpr0"><pre>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;"
+</pre></td></tr>
+
+
<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Stmt.html">Stmt</a>></td><td class="name" onclick="toggle('switchCase0')"><a name="switchCase0Anchor">switchCase</a></td><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1SwitchCase.html">SwitchCase</a>>...</td></tr>
<tr><td colspan="4" class="doc" id="switchCase0"><pre>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]]"
</pre></td></tr>
@@ -1381,26 +1412,6 @@ constructorDecl(hasAnyConstructorInitial
</pre></td></tr>
-<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXMethodDecl.html">CXXMethodDecl</a>></td><td class="name" onclick="toggle('hasOverloadedOperatorName0')"><a name="hasOverloadedOperatorName0Anchor">hasOverloadedOperatorName</a></td><td>StringRef Name</td></tr>
-<tr><td colspan="4" class="doc" id="hasOverloadedOperatorName0"><pre>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<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXOperatorCallExpr.html">CXXOperatorCallExpr</a>>, Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXMethodDecl.html">CXXMethodDecl</a>>
-</pre></td></tr>
-
-
<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXMethodDecl.html">CXXMethodDecl</a>></td><td class="name" onclick="toggle('isConst0')"><a name="isConst0Anchor">isConst</a></td><td></td></tr>
<tr><td colspan="4" class="doc" id="isConst0"><pre>Matches if the given method declaration is const.
@@ -1470,7 +1481,7 @@ operatorCallExpr(hasOverloadedOperatorNa
line and recordDecl(hasMethod(hasOverloadedOperatorName("*"))) matches
the declaration of A.
-Usable as: Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXOperatorCallExpr.html">CXXOperatorCallExpr</a>>, Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXMethodDecl.html">CXXMethodDecl</a>>
+Usable as: Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXOperatorCallExpr.html">CXXOperatorCallExpr</a>>, Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1FunctionDecl.html">FunctionDecl</a>>
</pre></td></tr>
@@ -1601,10 +1612,30 @@ and reference to that variable declarati
</pre></td></tr>
-<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>></td><td class="name" onclick="toggle('equalsNode0')"><a name="equalsNode0Anchor">equalsNode</a></td><td>Decl *Node</td></tr>
-<tr><td colspan="4" class="doc" id="equalsNode0"><pre>Matches if a node equals another node.
+<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>></td><td class="name" onclick="toggle('hasCudaDeviceAttr0')"><a name="hasCudaDeviceAttr0Anchor">hasCudaDeviceAttr</a></td><td></td></tr>
+<tr><td colspan="4" class="doc" id="hasCudaDeviceAttr0"><pre>Matches declaration that has CUDA device attribute.
-Decl has pointer identity in the AST.
+Given
+ __attribute__((device)) void f() { ... }
+matches the function declaration of f.
+</pre></td></tr>
+
+
+<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>></td><td class="name" onclick="toggle('hasCudaGlobalAttr0')"><a name="hasCudaGlobalAttr0Anchor">hasCudaGlobalAttr</a></td><td></td></tr>
+<tr><td colspan="4" class="doc" id="hasCudaGlobalAttr0"><pre>Matches declaration that has CUDA global attribute.
+
+Given
+ __attribute__((global)) void f() { ... }
+matches the function declaration of f.
+</pre></td></tr>
+
+
+<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html">Decl</a>></td><td class="name" onclick="toggle('hasCudaHostAttr0')"><a name="hasCudaHostAttr0Anchor">hasCudaHostAttr</a></td><td></td></tr>
+<tr><td colspan="4" class="doc" id="hasCudaHostAttr0"><pre>Matches declaration that has CUDA host attribute.
+
+Given
+ __attribute__((host)) void f() { ... }
+matches the function declaration of f.
</pre></td></tr>
@@ -1667,6 +1698,26 @@ Usable as: Matcher<<a href="http://cla
</pre></td></tr>
+<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1FunctionDecl.html">FunctionDecl</a>></td><td class="name" onclick="toggle('hasOverloadedOperatorName0')"><a name="hasOverloadedOperatorName0Anchor">hasOverloadedOperatorName</a></td><td>StringRef Name</td></tr>
+<tr><td colspan="4" class="doc" id="hasOverloadedOperatorName0"><pre>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<<a href="http://clang.llvm.org/doxygen/classclang_1_1CXXOperatorCallExpr.html">CXXOperatorCallExpr</a>>, Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1FunctionDecl.html">FunctionDecl</a>>
+</pre></td></tr>
+
+
<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1FunctionDecl.html">FunctionDecl</a>></td><td class="name" onclick="toggle('isDefinition2')"><a name="isDefinition2Anchor">isDefinition</a></td><td></td></tr>
<tr><td colspan="4" class="doc" id="isDefinition2"><pre>Matches if a declaration has a body attached.
@@ -1900,14 +1951,6 @@ and reference to that variable declarati
</pre></td></tr>
-<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Stmt.html">Stmt</a>></td><td class="name" onclick="toggle('equalsNode1')"><a name="equalsNode1Anchor">equalsNode</a></td><td>Stmt *Node</td></tr>
-<tr><td colspan="4" class="doc" id="equalsNode1"><pre>Matches if a node equals another node.
-
-Stmt has pointer identity in the AST.
-
-</pre></td></tr>
-
-
<tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1TagDecl.html">TagDecl</a>></td><td class="name" onclick="toggle('isDefinition0')"><a name="isDefinition0Anchor">isDefinition</a></td><td></td></tr>
<tr><td colspan="4" class="doc" id="isDefinition0"><pre>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(); }
</pre></td></tr>
@@ -3448,6 +3491,7 @@ Example matches X &x and const X &am
void a(X b) {
X &x = b;
const X &y = b;
+ }
};
</pre></td></tr>
Modified: cfe/trunk/include/clang/ASTMatchers/ASTMatchers.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/ASTMatchers/ASTMatchers.h?rev=214852&r1=214851&r2=214852&view=diff
==============================================================================
--- cfe/trunk/include/clang/ASTMatchers/ASTMatchers.h (original)
+++ cfe/trunk/include/clang/ASTMatchers/ASTMatchers.h Tue Aug 5 04:45:53 2014
@@ -1844,7 +1844,7 @@ AST_MATCHER_P(CallExpr, callee, internal
/// 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<Decl>, InnerMatcher,
1) {
@@ -1952,6 +1952,7 @@ AST_MATCHER_P_OVERLOAD(QualType, pointsT
/// void a(X b) {
/// X &x = b;
/// const X &y = b;
+/// }
/// };
/// \endcode
AST_MATCHER_P(QualType, references, internal::Matcher<QualType>,
@@ -3094,6 +3095,7 @@ AST_TYPE_MATCHER(IncompleteArrayType, in
/// 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,
return InnerMatcher.matches(*Node.getLHS(), Finder, Builder);
}
+/// \brief Matches CUDA kernel call expression.
+///
+/// Example matches kernel<<<i,j>>>()
+/// \code
+/// kernel<<<i,j>>>();
+/// \endcode
+const internal::VariadicDynCastAllOfMatcher<Stmt, CUDAKernelCallExpr>
+ 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<clang::CUDADeviceAttr>();
+}
+
+/// \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<clang::CUDAHostAttr>();
+}
+
+/// \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<clang::CUDAGlobalAttr>();
+}
+
} // end namespace ast_matchers
} // end namespace clang
Modified: cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.cpp?rev=214852&r1=214851&r2=214852&view=diff
==============================================================================
--- cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.cpp (original)
+++ cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.cpp Tue Aug 5 04:45:53 2014
@@ -649,6 +649,24 @@ TEST(DeclarationMatcher, HasDescendantMe
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.
Modified: cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.h?rev=214852&r1=214851&r2=214852&view=diff
==============================================================================
--- cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.h (original)
+++ cfe/trunk/unittests/ASTMatchers/ASTMatchersTest.h Tue Aug 5 04:45:53 2014
@@ -103,6 +103,72 @@ testing::AssertionResult notMatches(cons
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 <typename T>
+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<FrontendActionFactory> Factory(
+ newFrontendActionFactory(&Finder));
+ // Some tests use typeof, which is a gnu extension.
+ std::vector<std::string> 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 <typename T>
+testing::AssertionResult matchesWithCuda(const std::string &Code,
+ const T &AMatcher) {
+ return matchesConditionallyWithCuda(Code, AMatcher, true, "-std=c++11");
+}
+
+template <typename T>
+testing::AssertionResult notMatchesWithCuda(const std::string &Code,
+ const T &AMatcher) {
+ return matchesConditionallyWithCuda(Code, AMatcher, false, "-std=c++11");
+}
+
template <typename T>
testing::AssertionResult
matchAndVerifyResultConditionally(const std::string &Code, const T &AMatcher,
More information about the cfe-commits
mailing list