[clang-tidy] Suppress readability-static-access-through-instance for CUDA built-in variables

[carlosgalvezp]

Summary

clang-tidy can be used to statically analyze CUDA code,
thanks to clang being able to compile CUDA code natively.
This makes clang-tidy the one and only open-source
static analyzer for CUDA.

However it currently warns for native CUDA built-in
variables, like threadIdx, due to the way they
are implemented in clang.

Users don't need to know the details of the clang
implementation, and they should continue to write
idiomatic code. Therefore, suppress the warning
if a CUDA built-in variable is encountered.

Fixes https://bugs.llvm.org/show_bug.cgi?id=48758

[carlosgalvezp]

Btw, regarding this CHECK-MESSAGES-NOT, how does it work? I can't find it in check_clang_tidy.py. Wouldn't the test fail anyway if unexpected warnings are printed?

[carlosgalvezp]

Replace search with StringRef::contains.

[aaron.ballman]

In D112334#3081213, @carlosgalvezp wrote:

Btw, regarding this CHECK-MESSAGES-NOT, how does it work? I can't find it in check_clang_tidy.py. Wouldn't the test fail anyway if unexpected warnings are printed?

FileCheck matches input CHECK lines with output lines, so if there's no CHECK line for a particular output, FileCheck won't report any problems with it. (-verify is used by the frontend to verify diagnostics, but that's not something that can be used within clang-tidy currently.) That said, I'm not entirely certain of the mechanisms behind CHECK-MESSAGES-NOT in the python scripts.

LGTM aside from a nit.

clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp

75

I presume we want this to strictly be a prefix?

[carlosgalvezp]

Addressed comments.

[carlosgalvezp]

Thanks for the review!

I'm not super happy with the commit in the sense that it feels more like a workaround than a proper fix, with the hardcoded name __cuda_built_in and so on. But honestly I don't know how implement it more properly, for example check if the BaseTypeName is a CUDA built-in type. Another issue is that currently one can't ignore warnings from system headers (like __clang_cuda_builtin_vars.h)if they come from a macro defined there. But I suppose that's a harder problem to tackle.

But if you are happy with the patch I can go ahead and merge, and come back to it if I learn how to improve it later on as I get to learn more about the codebase.

clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp

75

Sure, but then I can't put the test in a named namespace. I changed it to anonymous namespace + link to Bugzilla, hope it works.

[aaron.ballman]

In D112334#3084668, @carlosgalvezp wrote:

Thanks for the review!

I'm not super happy with the commit in the sense that it feels more like a workaround than a proper fix, with the hardcoded name __cuda_built_in and so on. But honestly I don't know how implement it more properly, for example check if the BaseTypeName is a CUDA built-in type. Another issue is that currently one can't ignore warnings from system headers (like __clang_cuda_builtin_vars.h)if they come from a macro defined there. But I suppose that's a harder problem to tackle.

But if you are happy with the patch I can go ahead and merge, and come back to it if I learn how to improve it later on as I get to learn more about the codebase.

I'm also a bit uneasy about it because I suspect this may apply to more than just CUDA (e.g., I would imagine this applies to any builtin type). However, there's no easy way to tell that these are builtin types because they're exposed via a Headers include, and not via one of the intrinsic types in Types.h. So I don't know of a better way to fix this and this seems like an incremental improvement.

However, this won't be fool-proof either. A user could make a typedef to one of these types; access through the typedef will still warn because of the name-based matching. You could look through type sugar for that issue, but I'm not certain how critical that is to solve.

clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp

75

These types are not defined within a namespace in __clang_cuda_builtin_vars.h, so I wouldn't expect them to be in the test file TBH.

[carlosgalvezp]

Move CUDA built-in declarations to a separate header, and create a macro to define the variables just like it's done in the real version. Back to named namespace.

[tra]

Is there a way to annotate the builtin vars in the header to let clang-tidy know explicitly that the code must be handled in a special way? That would avoid the guesswork.

[aaron.ballman]

In D112334#3085119, @tra wrote:

Is there a way to annotate the builtin vars in the header to let clang-tidy know explicitly that the code must be handled in a special way? That would avoid the guesswork.

Hmmm, there's https://github.com/llvm/llvm-project/commit/e09107ab80dced55414fa458cf78e6cdfe90da6e but that's for identifying builtin functions, not builtin types. I don't think we have an attribute for identifying builtin type declarations. It might be worthwhile for someone to explore whether it would make sense for BuiltinAttr to be used here or not, but that might be a heavy lift to make this clang-tidy check a bit more robust.

Patch LGTM again. :-)

[carlosgalvezp]

Thanks for the review! I forgot to mention the Differential Revision in the commit message after pushing so this review stays open, is there any way I can add it now in some other way? I suppose we don't want force-push?

[whisperity]

In D112334#3084533, @aaron.ballman wrote:

In D112334#3081213, @carlosgalvezp wrote:

Btw, regarding this CHECK-MESSAGES-NOT, how does it work? I can't find it in check_clang_tidy.py. Wouldn't the test fail anyway if unexpected warnings are printed?

FileCheck matches input CHECK lines with output lines, so if there's no CHECK line for a particular output, FileCheck won't report any problems with it. (-verify is used by the frontend to verify diagnostics, but that's not something that can be used within clang-tidy currently.) That said, I'm not entirely certain of the mechanisms behind CHECK-MESSAGES-NOT in the python scripts.

Just a clarification for the record, as I came across this too. There is no business logic per se associated with CHECK-MESSAGES-NOT. I think it's some sort of a bad example that we observed and taken from each other and sometimes overusing. CHECK-MESSAGES-NOT is useful in the context where there is a warning that currently isn't emitted (because the check isn't fully finished, or depends on a FIXME or something) but you want the message to appear in the code for later development as a placeholder. It indicates what should be there, but currently isn't there and we are not expecting it. The moment the check is further developed, it's easy to cut back the -NOT part (2t-ld2w if you're using Vim) and turn it into a proper check that expects the warning to be there. (And of course, hopefully also emits it!)

Semantically, it is equivalent to saying // NO-WARN: Clang is awesome! or // MY-LITTLE-PONY or whatever in the text. It will be ignored by the script. 😉 The script only triggers for CHECK-MESSAGES:, CHECK-FIXES: and CHECK-NOTES: and the prefixes you may or may not be able to specify at invocation explicitly.

IMHO if we never want that warning to pop it is misleading and noise to add the would-be text of a warning there. If you want to indicate that there is no need for a warning (any warning emitted by the check!), I believe it is better to use // NO-WARN. That indicates both the intent that the code in that line or nearby SHOULD PASS but also clearly documents that it wasn't a developer oversight that the test appears as a "negative case" (i.e. you did not forget to add the CHECK- lines in).

[whisperity]

@carlosgalvezp Force pushes are prevented by the serverside integration so we couldn't force push even if we wanted to. What one could do is to revert the patch and land it again (or land an empty patch that has the association line in the message...) but both of those would be just excess noise in the (already waaaay too noisy!) history.

Hopefully this will help.

[carlosgalvezp]

@whisperity Understood, thanks for the help! I'll take more care to remember the reference next time :) I added it as a comment in Github for reference so there's at least some traceability.

clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp

 //===--- StaticAccessedThroughInstanceCheck.cpp - clang-tidy---------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "StaticAccessedThroughInstanceCheck.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/ASTMatchers/ASTMatchFinder.h"
+#include "llvm/ADT/StringRef.h"
 
 using namespace clang::ast_matchers;
 
 namespace clang {
 namespace tidy {
 namespace readability {
 
 static unsigned getNameSpecifierNestingLevel(const QualType &QType) {
void StaticAccessedThroughInstanceCheck::check(
   const auto *MemberExpression =
       Result.Nodes.getNodeAs<MemberExpr>("memberExpression");
 
   if (MemberExpression->getBeginLoc().isMacroID())
     return;
 
   const Expr *BaseExpr = MemberExpression->getBase();
 
-  // Do not warn for overlaoded -> operators.
+  // Do not warn for overloaded -> operators.
   if (isa<CXXOperatorCallExpr>(BaseExpr))
     return;
 
   QualType BaseType =
       BaseExpr->getType()->isPointerType()
           ? BaseExpr->getType()->getPointeeType().getUnqualifiedType()
           : BaseExpr->getType().getUnqualifiedType();
 
   const ASTContext *AstContext = Result.Context;
   PrintingPolicy PrintingPolicyWithSupressedTag(AstContext->getLangOpts());
   PrintingPolicyWithSupressedTag.SuppressTagKeyword = true;
   PrintingPolicyWithSupressedTag.SuppressUnwrittenScope = true;
   std::string BaseTypeName =
       BaseType.getAsString(PrintingPolicyWithSupressedTag);
 
+  // Do not warn for CUDA built-in variables.
+  if (StringRef(BaseTypeName).startswith("__cuda_builtin_"))
+    return;
+
   SourceLocation MemberExprStartLoc = MemberExpression->getBeginLoc();
   auto Diag =
       diag(MemberExprStartLoc, "static member accessed through instance");
 
   if (BaseExpr->HasSideEffects(*AstContext) ||
       getNameSpecifierNestingLevel(BaseType) > NameSpecifierNestingThreshold)
     return;
 

clang-tools-extra/test/clang-tidy/checkers/Inputs/readability-static-accessed-through-instance/__clang_cuda_builtin_vars.h

+//===--- __clang_cuda_builtin_vars.h - Stub header for tests ----*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef ___CLANG_CUDA_BUILTIN_VARS_H_
+#define ___CLANG_CUDA_BUILTIN_VARS_H_
+
+#define __CUDA_DEVICE_BUILTIN(FIELD) \
+  static unsigned int FIELD;
+
+struct __cuda_builtin_threadIdx_t {
+  __CUDA_DEVICE_BUILTIN(x);
+};
+
+struct __cuda_builtin_blockIdx_t {
+  __CUDA_DEVICE_BUILTIN(x);
+};
+
+struct __cuda_builtin_blockDim_t {
+  __CUDA_DEVICE_BUILTIN(x);
+};
+
+struct __cuda_builtin_gridDim_t {
+  __CUDA_DEVICE_BUILTIN(x);
+};
+
+__cuda_builtin_threadIdx_t threadIdx;
+__cuda_builtin_blockIdx_t blockIdx;
+__cuda_builtin_blockDim_t blockDim;
+__cuda_builtin_gridDim_t gridDim;
+
+#endif // ___CLANG_CUDA_BUILTIN_VARS_H_

clang-tools-extra/test/clang-tidy/checkers/readability-static-accessed-through-instance.cpp

-// RUN: %check_clang_tidy %s readability-static-accessed-through-instance %t
+// RUN: %check_clang_tidy %s readability-static-accessed-through-instance %t -- -- -isystem %S/Inputs/readability-static-accessed-through-instance
+#include <__clang_cuda_builtin_vars.h>
 
 struct C {
   static void foo();
   static int x;
   int nsx;
   void mf() {
     (void)&x;    // OK, x is accessed inside the struct.
     (void)&C::x; // OK, x is accessed using a qualified-id.

 }
 
 void use_inline() {
   Outer::S V;
   V.I;
   // CHECK-MESSAGES: :[[@LINE-1]]:3: warning: static member
   // CHECK-FIXES: {{^}}  Outer::S::I;{{$}}
 }
+
+// https://bugs.llvm.org/show_bug.cgi?id=48758
+namespace Bugzilla_48758 {
+
+unsigned int x1 = threadIdx.x;
+// CHECK-MESSAGES-NOT: :[[@LINE-1]]:10: warning: static member
+unsigned int x2 = blockIdx.x;
+// CHECK-MESSAGES-NOT: :[[@LINE-1]]:10: warning: static member
+unsigned int x3 = blockDim.x;
+// CHECK-MESSAGES-NOT: :[[@LINE-1]]:10: warning: static member
+unsigned int x4 = gridDim.x;
+// CHECK-MESSAGES-NOT: :[[@LINE-1]]:10: warning: static member
+
+} // namespace Bugzilla_48758