[clang][AST matchers] adding submatchers under cudaKernelCallExpr to match kernel launch config

[ajohnson-uoregon]

Summary

adding more AST matchers for all possible launch params to a CUDA kernel, e.g. cudaGridDim and cudaSharedMemPerBlock

[aaron.ballman]

Typically, we only add AST matchers when there's an obvious need for them (we do this because compiling this file takes a *long* time and generates quite a few symbols in the executable). I'm not certain that these CUDA matchers rise to that level of need -- are you planning to write a number of clang-tidy checks that use these new interfaces (or some other in-tree use)?

(Assuming you think we should continue to progress this patch, you should fix the clang-format issues, add test coverage, and regenerate documentation.)

[ajohnson-uoregon]

fixing clang-format things and adding docs; also making docs a bit more clear

[ajohnson-uoregon]

I still need to write tests but: I do have a use case for these over here: https://github.com/ajohnson-uoregon/llvm-project/blob/feature-ajohnson/clang-tools-extra/clang-rewrite/ConstructMatchers.cpp#L472
tl;dr, we'd like to match the kernel launch arguments (i.e., the arguments to __cudaPushCallConfiguration()) and these matchers made writing the code to generate those AST matchers much easier.

Without at least the hasKernelConfig() matcher, it's actually currently impossible to match the kernel launch args. (I wasn't able to find a way after quite a while poking at the AST, at least.) As for the others, it's not clear how to match the kernel launch args without exposing the fact that there's a second CallExpr inside the CUDAKernelCallExpr to the user and writing a pretty messy matcher, along the lines of cudaKernelCallExpr(hasKernelConfig(callExpr(hasArgument(0, expr())))) for the grid dim. cudaKernelCallExpr(cudaGridDim()) is a lot cleaner and easier to understand.

[aaron.ballman]

In D120952#3381224, @ajohnson-uoregon wrote:

I still need to write tests but: I do have a use case for these over here: https://github.com/ajohnson-uoregon/llvm-project/blob/feature-ajohnson/clang-tools-extra/clang-rewrite/ConstructMatchers.cpp#L472
tl;dr, we'd like to match the kernel launch arguments (i.e., the arguments to __cudaPushCallConfiguration()) and these matchers made writing the code to generate those AST matchers much easier.

Thanks for the example use case, that helps! Do you expect to need it in more than one project though? We support defining local AST matchers so people can do one-off matching, but we typically only add AST matchers for things that we expect to be generally useful (multiple projects would benefit from it). For example, this project adds a number of local matchers: https://github.com/llvm/llvm-project/blob/main/clang-tools-extra/clang-tidy/misc/RedundantExpressionCheck.cpp#L435

Without at least the hasKernelConfig() matcher, it's actually currently impossible to match the kernel launch args. (I wasn't able to find a way after quite a while poking at the AST, at least.) As for the others, it's not clear how to match the kernel launch args without exposing the fact that there's a second CallExpr inside the CUDAKernelCallExpr to the user and writing a pretty messy matcher, along the lines of cudaKernelCallExpr(hasKernelConfig(callExpr(hasArgument(0, expr())))) for the grid dim. cudaKernelCallExpr(cudaGridDim()) is a lot cleaner and easier to understand.

Yup, I think you need these kind of matchers for what you want to do. What I'm less certain of is whether others will need them (we don't have any CUDA-specific clang-tidy modules, I don't think any of the existing coding standards we support checks for have anything to say about CUDA, etc).

[ajohnson-uoregon]

For the specific argument matchers (cudaBlockDim(), cudaStream(), etc), I agree, I'm not sure if there's any other use cases in tree. I wasn't aware you could write local definitions like that, I could definitely do that within my project for those.

I would still like to upstream hasKernelConfig() though, so it's at least possible to match the kernel config without defining a local matcher.

[aaron.ballman]

In D120952#3384239, @ajohnson-uoregon wrote:

For the specific argument matchers (cudaBlockDim(), cudaStream(), etc), I agree, I'm not sure if there's any other use cases in tree. I wasn't aware you could write local definitions like that, I could definitely do that within my project for those.

SGTM!

I would still like to upstream hasKernelConfig() though, so it's at least possible to match the kernel config without defining a local matcher.

I think that's reasonable. One thing to keep in mind for that is there are multiple language modes that all have the notion of a kernel (OpenCL and SYCL both do as well). We should make sure the matcher can be extended for those cases someday in the future, or be given a CUDA-specific name (e.g, hasCUDAKernelConfig()) if that's a better approach. (I suspect polymorphic matchers will work fine for this, but it's worth someone considering explicitly.)

[ajohnson-uoregon]

Removing specific argument matchers and leaving just hasKernelConfig()

[aaron.ballman]

Thanks! This is missing test coverage for the new matcher, and you should also add a release note for it.

clang/include/clang/ASTMatchers/ASTMatchers.h

7859–7861

clang/docs/LibASTMatchersReference.html

 pointerType(pointee(isConstQualified(), isInteger()))
   matches "int const *b"
 
 Usable as: Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1BlockPointerType.html">BlockPointerType</a>&gt;, Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1MemberPointerType.html">MemberPointerType</a>&gt;,
   Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1PointerType.html">PointerType</a>&gt;, Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1ReferenceType.html">ReferenceType</a>&gt;
 </pre></td></tr>
 
 
+<tr><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>&gt;</td><td class="name" onclick="toggle('cudaBlockDim0')"><a name="cudaBlockDim0Anchor">cudaBlockDim</a></td><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>&gt; InnerMatcher</td></tr>
+<tr><td colspan="4" class="doc" id="cudaBlockDim0"><pre>Matches the second argument (block dim) in &lt;&lt;&lt;&gt;&gt;&gt; on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaBlockDim()) will match j in
+  kernel&lt;&lt;&lt;i,j&gt;&gt;&gt;();
+</pre></td></tr>
+
+
+<tr><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>&gt;</td><td class="name" onclick="toggle('cudaGridDim0')"><a name="cudaGridDim0Anchor">cudaGridDim</a></td><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>&gt; InnerMatcher</td></tr>
+<tr><td colspan="4" class="doc" id="cudaGridDim0"><pre>Matches the first argument (grid dim) in &lt;&lt;&lt;&gt;&gt;&gt; on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaGridDim()) will match i in
+  kernel&lt;&lt;&lt;i,j&gt;&gt;&gt;();
+</pre></td></tr>
+
+
+<tr><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>&gt;</td><td class="name" onclick="toggle('cudaSharedMemPerBlock0')"><a name="cudaSharedMemPerBlock0Anchor">cudaSharedMemPerBlock</a></td><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>&gt; InnerMatcher</td></tr>
+<tr><td colspan="4" class="doc" id="cudaSharedMemPerBlock0"><pre>Matches the third argument (shared mem size) in &lt;&lt;&lt;&gt;&gt;&gt; on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaSharedMemPerBlock()) will match mem in
+  kernel&lt;&lt;&lt;i, j, mem, 0&gt;&gt;&gt;();
+</pre></td></tr>
+
+
+<tr><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>&gt;</td><td class="name" onclick="toggle('cudaStream0')"><a name="cudaStream0Anchor">cudaStream</a></td><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1Expr.html">Expr</a>&gt; InnerMatcher</td></tr>
+<tr><td colspan="4" class="doc" id="cudaStream0"><pre>Matches the fourth argument (CUDA stream) in &lt;&lt;&lt;&gt;&gt;&gt; on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(cudaStream()) will match 0 in
+  kernel&lt;&lt;&lt;i, j, mem, 0&gt;&gt;&gt;();
+</pre></td></tr>
+
+
+<tr><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1CUDAKernelCallExpr.html">CUDAKernelCallExpr</a>&gt;</td><td class="name" onclick="toggle('hasKernelConfig0')"><a name="hasKernelConfig0Anchor">hasKernelConfig</a></td><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1CallExpr.html">CallExpr</a>&gt; InnerMatcher</td></tr>
+<tr><td colspan="4" class="doc" id="hasKernelConfig0"><pre>Matches the kernel launch config (in &lt;&lt;&lt;&gt;&gt;&gt;) on CUDA kernel calls.
+
+Example: cudaKernelCallExpr(hasKernelConfig()) will match &lt;&lt;&lt;i,j&gt;&gt;&gt; in
+  kernel&lt;&lt;&lt;i,j&gt;&gt;&gt;();
+</pre></td></tr>
+
+
 <tr><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1CXXBaseSpecifier.html">CXXBaseSpecifier</a>&gt;</td><td class="name" onclick="toggle('hasTypeLoc1')"><a name="hasTypeLoc1Anchor">hasTypeLoc</a></td><td>Matcher&lt;<a href="https://clang.llvm.org/doxygen/classclang_1_1TypeLoc.html">TypeLoc</a>&gt; Inner</td></tr>
 <tr><td colspan="4" class="doc" id="hasTypeLoc1"><pre>Matches if the type location of a node matches the inner matcher.
 
 Examples:
   int x;
 declaratorDecl(hasTypeLoc(loc(asString("int"))))
   matches int x
 

clang/include/clang/ASTMatchers/ASTMatchers.h

 ///
 /// Example matches,
 /// \code
 ///   kernel<<<i,j>>>();
 /// \endcode
 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.
 ///
 /// Given:
 /// \code
 ///   void *v1 = NULL;
 ///   void *v2 = nullptr;
 ///   void *v3 = __null; // GNU extension

clang/lib/ASTMatchers/Dynamic/Registry.cpp

RegistryMaps::RegistryMaps() {
   REGISTER_MATCHER(coawaitExpr);
   REGISTER_MATCHER(conditionalOperator);
   REGISTER_MATCHER(constantArrayType);
   REGISTER_MATCHER(constantExpr);
   REGISTER_MATCHER(containsDeclaration);
   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);
   REGISTER_MATCHER(cxxCatchStmt);
   REGISTER_MATCHER(cxxConstCastExpr);
   REGISTER_MATCHER(cxxConstructExpr);
   REGISTER_MATCHER(cxxConstructorDecl);
   REGISTER_MATCHER(cxxConversionDecl);
RegistryMaps::RegistryMaps() {
   REGISTER_MATCHER(hasGlobalStorage);
   REGISTER_MATCHER(hasImplicitDestinationType);
   REGISTER_MATCHER(hasInClassInitializer);
   REGISTER_MATCHER(hasIncrement);
   REGISTER_MATCHER(hasIndex);
   REGISTER_MATCHER(hasInit);
   REGISTER_MATCHER(hasInitializer);
   REGISTER_MATCHER(hasInitStatement);
+  REGISTER_MATCHER(hasKernelConfig);
   REGISTER_MATCHER(hasKeywordSelector);
   REGISTER_MATCHER(hasLHS);
   REGISTER_MATCHER(hasLocalQualifiers);
   REGISTER_MATCHER(hasLocalStorage);
   REGISTER_MATCHER(hasLoopInit);
   REGISTER_MATCHER(hasLoopVariable);
   REGISTER_MATCHER(hasMemberName);
   REGISTER_MATCHER(hasMethod);