Good point :) Added a negative test for matchesWithCuda and notMatchesWithCuda. And will use phab in future.
On Mon, Aug 4, 2014 at 4:35 AM, Manuel Klimek <[email protected]> wrote: > On Mon, Aug 4, 2014 at 1:34 PM, Manuel Klimek <[email protected]> wrote: > >> On Fri, Aug 1, 2014 at 7:43 PM, Jacques Pienaar <[email protected]> >> wrote: >> >>> Hi, >>> >>> The patch adds ASTMatchers for matching CUDA declarations. >>> >>> Kindly asking for review. >>> >> >> (just for the future: it'll make it easier for me to review if you run it >> through phab - this change is small enough that it doesn't make a big >> difference though) >> >> Can you please run: >> $ cd tools/clang/docs/tools >> $ python ./dump_ast_matchers.py >> (this will automatically update the documentation) >> > > Also, you might want to add a reverse test (otherwise we have a) code > that's not executed and b) we could implement matchesWithCuda as "return > true" :) > >> >> >>> >>> Jacques >>> >>> -- >>> Jacques Pienaar | Platforms | [email protected] | 765-430-6883 >>> >>> _______________________________________________ >>> cfe-commits mailing list >>> [email protected] >>> http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits >>> >>> >> > -- Jacques Pienaar | Platforms | [email protected] | 765-430-6883
Index: docs/LibASTMatchersReference.html =================================================================== --- docs/LibASTMatchersReference.html (revision 214724) +++ docs/LibASTMatchersReference.html (working copy) @@ -319,6 +319,16 @@ 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 @@ </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 @@ 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 @@ </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 @@ int a[] = { 2, 3 } int b[42]; int c[a[0]]; + } variableArrayType() matches "int c[a[0]]" </pre></td></tr> @@ -1381,26 +1412,6 @@ </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 @@ 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,13 +1612,33 @@ </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> + + <tr><td>Matcher<<a href="http://clang.llvm.org/doxygen/classclang_1_1Decl.html";>Decl</a>></td><td class="name" onclick="toggle('isImplicit0')"><a name="isImplicit0Anchor">isImplicit</a></td><td></td></tr> <tr><td colspan="4" class="doc" id="isImplicit0"><pre>Matches a declaration that has been implicitly added by the compiler (eg. implicit defaultcopy constructors). @@ -1667,6 +1698,26 @@ </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 @@ </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 @@ 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 @@ void a(X b) { X &x = b; const X &y = b; + } }; </pre></td></tr> Index: include/clang/ASTMatchers/ASTMatchers.h =================================================================== --- include/clang/ASTMatchers/ASTMatchers.h (revision 214724) +++ include/clang/ASTMatchers/ASTMatchers.h (working copy) @@ -1844,7 +1844,7 @@ /// 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 @@ /// void a(X b) { /// X &x = b; /// const X &y = b; +/// } /// }; /// \endcode AST_MATCHER_P(QualType, references, internal::Matcher<QualType>, @@ -3094,6 +3095,7 @@ /// int a[] = { 2, 3 } /// int b[42]; /// int c[a[0]]; +/// } /// \endcode /// variableArrayType() /// matches "int c[a[0]]" @@ -3643,6 +3645,48 @@ 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 Index: unittests/ASTMatchers/ASTMatchersTest.cpp =================================================================== --- unittests/ASTMatchers/ASTMatchersTest.cpp (revision 214724) +++ unittests/ASTMatchers/ASTMatchersTest.cpp (working copy) @@ -649,6 +649,24 @@ 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. Index: unittests/ASTMatchers/ASTMatchersTest.h =================================================================== --- unittests/ASTMatchers/ASTMatchersTest.h (revision 214724) +++ unittests/ASTMatchers/ASTMatchersTest.h (working copy) @@ -103,7 +103,73 @@ 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, BoundNodesCallback *FindResultVerifier,
_______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
