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

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

diff --git a/clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp b/clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp
index df4a39e..f2c1b0f 100644 (file)
--- a/clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp
+++ b/clang-tools-extra/clang-tidy/readability/StaticAccessedThroughInstanceCheck.cpp
@@ -9,6 +9,7 @@
 #include "StaticAccessedThroughInstanceCheck.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/ASTMatchers/ASTMatchFinder.h"
+#include "llvm/ADT/StringRef.h"
 
 using namespace clang::ast_matchers;
 
@@ -54,7 +55,7 @@ void StaticAccessedThroughInstanceCheck::check(
 
   const Expr *BaseExpr = MemberExpression->getBase();
 
-  // Do not warn for overlaoded -> operators.
+  // Do not warn for overloaded -> operators.
   if (isa<CXXOperatorCallExpr>(BaseExpr))
     return;
 
@@ -70,6 +71,10 @@ void StaticAccessedThroughInstanceCheck::check(
   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");

diff --git a/clang-tools-extra/test/clang-tidy/checkers/Inputs/readability-static-accessed-through-instance/__clang_cuda_builtin_vars.h b/clang-tools-extra/test/clang-tidy/checkers/Inputs/readability-static-accessed-through-instance/__clang_cuda_builtin_vars.h
new file mode 100644 (file)
index 0000000..63eb173
--- /dev/null
+++ b/clang-tools-extra/test/clang-tidy/checkers/Inputs/readability-static-accessed-through-instance/__clang_cuda_builtin_vars.h
@@ -0,0 +1,36 @@
+//===--- __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_

diff --git a/clang-tools-extra/test/clang-tidy/checkers/readability-static-accessed-through-instance.cpp b/clang-tools-extra/test/clang-tidy/checkers/readability-static-accessed-through-instance.cpp
index 9de96ae..cd8d198 100644 (file)
--- a/clang-tools-extra/test/clang-tidy/checkers/readability-static-accessed-through-instance.cpp
+++ b/clang-tools-extra/test/clang-tidy/checkers/readability-static-accessed-through-instance.cpp
@@ -1,4 +1,5 @@
-// 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();
@@ -248,3 +249,17 @@ void use_inline() {
   // 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