ajohnson-uoregon updated this revision to Diff 415289. ajohnson-uoregon added a comment.
fixing clang-format things and adding docs; also making docs a bit more clear Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D120952/new/ https://reviews.llvm.org/D120952 Files: clang/docs/LibASTMatchersReference.html clang/include/clang/ASTMatchers/ASTMatchers.h clang/lib/ASTMatchers/Dynamic/Registry.cpp
Index: clang/lib/ASTMatchers/Dynamic/Registry.cpp =================================================================== --- clang/lib/ASTMatchers/Dynamic/Registry.cpp +++ clang/lib/ASTMatchers/Dynamic/Registry.cpp @@ -177,7 +177,11 @@ REGISTER_MATCHER(continueStmt); REGISTER_MATCHER(coreturnStmt); REGISTER_MATCHER(coyieldExpr); + REGISTER_MATCHER(cudaBlockDim); + REGISTER_MATCHER(cudaGridDim); REGISTER_MATCHER(cudaKernelCallExpr); + REGISTER_MATCHER(cudaSharedMemPerBlock); + REGISTER_MATCHER(cudaStream); REGISTER_MATCHER(cxxBaseSpecifier); REGISTER_MATCHER(cxxBindTemporaryExpr); REGISTER_MATCHER(cxxBoolLiteral); @@ -320,6 +324,7 @@ REGISTER_MATCHER(hasInit); REGISTER_MATCHER(hasInitializer); REGISTER_MATCHER(hasInitStatement); + REGISTER_MATCHER(hasKernelConfig); REGISTER_MATCHER(hasKeywordSelector); REGISTER_MATCHER(hasLHS); REGISTER_MATCHER(hasLocalQualifiers); Index: clang/include/clang/ASTMatchers/ASTMatchers.h =================================================================== --- clang/include/clang/ASTMatchers/ASTMatchers.h +++ clang/include/clang/ASTMatchers/ASTMatchers.h @@ -7848,6 +7848,80 @@ extern const internal::VariadicDynCastAllOfMatcher<Stmt, CUDAKernelCallExpr> cudaKernelCallExpr; +/// Matches the kernel launch config (in <<<>>>) on CUDA kernel calls. +/// +/// Example: cudaKernelCallExpr(hasKernelConfig()) will match <<<i,j>>> in +/// \code +/// kernel<<<i,j>>>(); +/// \endcode +AST_MATCHER_P(CUDAKernelCallExpr, hasKernelConfig, internal::Matcher<CallExpr>, + InnerMatcher) { + if (const CallExpr *Config = Node.getConfig()) { + return InnerMatcher.matches(*Config, Finder, Builder); + } + return false; +} + +/// Matches the first argument (grid dim) in <<<>>> on CUDA kernel calls. +/// +/// Example: cudaKernelCallExpr(cudaGridDim()) will match i in +/// \code +/// kernel<<<i,j>>>(); +/// \endcode +AST_MATCHER_P(CUDAKernelCallExpr, cudaGridDim, internal::Matcher<Expr>, + InnerMatcher) { + const CallExpr *Config = Node.getConfig(); + if (Config && Config->getNumArgs() > 0) { + return InnerMatcher.matches(*(Config->getArg(0)), Finder, Builder); + } + return false; +} + +/// Matches the second argument (block dim) in <<<>>> on CUDA kernel calls. +/// +/// Example: cudaKernelCallExpr(cudaBlockDim()) will match j in +/// \code +/// kernel<<<i,j>>>(); +/// \endcode +AST_MATCHER_P(CUDAKernelCallExpr, cudaBlockDim, internal::Matcher<Expr>, + InnerMatcher) { + const CallExpr *Config = Node.getConfig(); + if (Config && Config->getNumArgs() > 1) { + return InnerMatcher.matches(*(Config->getArg(1)), Finder, Builder); + } + return false; +} + +/// Matches the third argument (shared mem size) in <<<>>> on CUDA kernel calls. +/// +/// Example: cudaKernelCallExpr(cudaSharedMemPerBlock()) will match mem in +/// \code +/// kernel<<<i, j, mem, 0>>>(); +/// \endcode +AST_MATCHER_P(CUDAKernelCallExpr, cudaSharedMemPerBlock, + internal::Matcher<Expr>, InnerMatcher) { + const CallExpr *Config = Node.getConfig(); + if (Config && Config->getNumArgs() > 2) { + return InnerMatcher.matches(*(Config->getArg(2)), Finder, Builder); + } + return false; +} + +/// Matches the fourth argument (CUDA stream) in <<<>>> on CUDA kernel calls. +/// +/// Example: cudaKernelCallExpr(cudaStream()) will match 0 in +/// \code +/// kernel<<<i, j, mem, 0>>>(); +/// \endcode +AST_MATCHER_P(CUDAKernelCallExpr, cudaStream, internal::Matcher<Expr>, + InnerMatcher) { + const CallExpr *Config = Node.getConfig(); + if (Config && Config->getNumArgs() > 3) { + return InnerMatcher.matches(*(Config->getArg(3)), Finder, Builder); + } + return false; +} + /// Matches expressions that resolve to a null pointer constant, such as /// GNU's __null, C++11's nullptr, or C's NULL macro. /// Index: clang/docs/LibASTMatchersReference.html =================================================================== --- clang/docs/LibASTMatchersReference.html +++ clang/docs/LibASTMatchersReference.html @@ -6426,6 +6426,46 @@ </pre></td></tr> +<tr><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>></td><td class="name" onclick="toggle('cudaBlockDim0')"><a name="cudaBlockDim0Anchor">cudaBlockDim</a></td><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>> InnerMatcher</td></tr> +<tr><td colspan="4" class="doc" id="cudaBlockDim0"><pre>Matches the second argument (block dim) in <<<>>> on CUDA kernel calls. + +Example: cudaKernelCallExpr(cudaBlockDim()) will match j in + kernel<<<i,j>>>(); +</pre></td></tr> + + +<tr><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>></td><td class="name" onclick="toggle('cudaGridDim0')"><a name="cudaGridDim0Anchor">cudaGridDim</a></td><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>> InnerMatcher</td></tr> +<tr><td colspan="4" class="doc" id="cudaGridDim0"><pre>Matches the first argument (grid dim) in <<<>>> on CUDA kernel calls. + +Example: cudaKernelCallExpr(cudaGridDim()) will match i in + kernel<<<i,j>>>(); +</pre></td></tr> + + +<tr><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>></td><td class="name" onclick="toggle('cudaSharedMemPerBlock0')"><a name="cudaSharedMemPerBlock0Anchor">cudaSharedMemPerBlock</a></td><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>> InnerMatcher</td></tr> +<tr><td colspan="4" class="doc" id="cudaSharedMemPerBlock0"><pre>Matches the third argument (shared mem size) in <<<>>> on CUDA kernel calls. + +Example: cudaKernelCallExpr(cudaSharedMemPerBlock()) will match mem in + kernel<<<i, j, mem, 0>>>(); +</pre></td></tr> + + +<tr><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>></td><td class="name" onclick="toggle('cudaStream0')"><a name="cudaStream0Anchor">cudaStream</a></td><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>> InnerMatcher</td></tr> +<tr><td colspan="4" class="doc" id="cudaStream0"><pre>Matches the fourth argument (CUDA stream) in <<<>>> on CUDA kernel calls. + +Example: cudaKernelCallExpr(cudaStream()) will match 0 in + kernel<<<i, j, mem, 0>>>(); +</pre></td></tr> + + +<tr><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>></td><td class="name" onclick="toggle('hasKernelConfig0')"><a name="hasKernelConfig0Anchor">hasKernelConfig</a></td><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1CallExpr.html">CallExpr</a>> InnerMatcher</td></tr> +<tr><td colspan="4" class="doc" id="hasKernelConfig0"><pre>Matches the kernel launch config (in <<<>>>) on CUDA kernel calls. + +Example: cudaKernelCallExpr(hasKernelConfig()) will match <<<i,j>>> in + kernel<<<i,j>>>(); +</pre></td></tr> + + <tr><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1CXXBaseSpecifier.html">CXXBaseSpecifier</a>></td><td class="name" onclick="toggle('hasTypeLoc1')"><a name="hasTypeLoc1Anchor">hasTypeLoc</a></td><td>Matcher<<a href="https://clang.llvm.org/doxygen/classclang_1_1TypeLoc.html">TypeLoc</a>> Inner</td></tr> <tr><td colspan="4" class="doc" id="hasTypeLoc1"><pre>Matches if the type location of a node matches the inner matcher.
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits