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
