Ground work for cuda-related checks in clang-tidy

[barcisz]

Summary

This diff adds support for adding CUDA checks to clang-tidy, as well as tests for those checks. This is in preparation for the checks introduced in D133804 and D133956, both of which were
developed by Meta (formerly Facebook). The cuda-unused

[tschuett]

clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp

26

Is Google a copy and paste error?

33

Is Google a copy and paste error?

[ivanmurashko]

clang-tools-extra/test/lit.cfg.py

19 ↗

(On Diff #458497)

.cu is specified several times there. It looks like the change is not necessary at the line

[barcisz]

Fixed a copy-paste error with Google->Cuda in CudaTidyModule.cpp

[barcisz]

removed duplication of '.cu' in the lit config for clang-tidy tests

[barcisz]

clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp

26

Done

33

Done

clang-tools-extra/test/lit.cfg.py

19 ↗

(On Diff #458497)

Done

[barcisz]

@tschuett does it look alright now?

[barcisz]

Better explanation of cuda-related flags in tests

[barcisz]

Common documentation for cuda checks

[njames93]

What's with the .keep files. Will lit or sphinx fail if the cuda directory doesn't exist, if not I'd be happier removing them(they'd get removed once the first cuda check lands anyway).

clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp

2

Google...

It may be easier/safer to use the add_new_check script to generate this, then just remove all the new check specific files.

clang-tools-extra/docs/ReleaseNotes.rst

98 ↗

(On Diff #460070)

Put it in here.

152 ↗

(On Diff #460070)

This line doesn't belong in the new checks portion of the release notes, just stick it in the general clang-tidy section.
Also maybe just say "Introduce the cuda module for checks specific to cuda code" and remove all the documentation below, That will be filled in when the first check lands.

clang-tools-extra/docs/clang-tidy/index.rst

143 ↗

(On Diff #460070)

Unnecessary edit can be removed.

241 ↗

(On Diff #460070)

Ditto.

clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda-initializers.h

2

What's the purpose of this file here, do you envision many checks requiring to include this?

clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda.h

4

This looks worrying as test cases don't run with a standard library

[barcisz]

Changes suggested by njames93

[barcisz]

Dummy definition for size_t

[barcisz]

moved size_t definition to an stddef.h stub

[njames93]

clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/stddef.h

12 ↗

(On Diff #460520)

If this is all the file is used for, and it's only used for this one test file, can probably remove this header and inline the definition.
I also feel this definition may be fragile on certain platforms.

[barcisz]

clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/stddef.h

12 ↗

(On Diff #460520)

I wanted to add it in a separate header for future proofness; as for the definition, I cannot really use uint64_t here since I cannot include the real stddef.h

[aaron.ballman]

clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/stddef.h

12 ↗

(On Diff #460520)

using size_t = decltype(sizeof(0)); is a better way to spell this, and I'd probably inline it rather than add a new header file.

[barcisz]

Inlined definition of size_t

clang-tools-extra/clang-tidy/CMakeLists.txt

 add_subdirectory(android)
 add_subdirectory(abseil)
 add_subdirectory(altera)
 add_subdirectory(boost)
 add_subdirectory(bugprone)
 add_subdirectory(cert)
 add_subdirectory(concurrency)
 add_subdirectory(cppcoreguidelines)
+add_subdirectory(cuda)
 add_subdirectory(darwin)
 add_subdirectory(fuchsia)
 add_subdirectory(google)
 add_subdirectory(hicpp)
 add_subdirectory(linuxkernel)
 add_subdirectory(llvm)
 add_subdirectory(llvmlibc)
 add_subdirectory(misc)
set(ALL_CLANG_TIDY_CHECKS
   clangTidyAndroidModule
   clangTidyAbseilModule
   clangTidyAlteraModule
   clangTidyBoostModule
   clangTidyBugproneModule
   clangTidyCERTModule
   clangTidyConcurrencyModule
   clangTidyCppCoreGuidelinesModule
+  clangTidyCudaModule
   clangTidyDarwinModule
   clangTidyFuchsiaModule
   clangTidyGoogleModule
   clangTidyHICPPModule
   clangTidyLinuxKernelModule
   clangTidyLLVMModule
   clangTidyLLVMLibcModule
   clangTidyMiscModule

clang-tools-extra/clang-tidy/ClangTidyForceLinker.h

 static int LLVM_ATTRIBUTE_UNUSED ConcurrencyModuleAnchorDestination =
     ConcurrencyModuleAnchorSource;
 
 // This anchor is used to force the linker to link the CppCoreGuidelinesModule.
 extern volatile int CppCoreGuidelinesModuleAnchorSource;
 static int LLVM_ATTRIBUTE_UNUSED CppCoreGuidelinesModuleAnchorDestination =
     CppCoreGuidelinesModuleAnchorSource;
 
+// This anchor is used to force the linker to link the CudaModule.
+extern volatile int CudaModuleAnchorSource;
+static int LLVM_ATTRIBUTE_UNUSED CudaModuleAnchorDestination =
+    CudaModuleAnchorSource;
+
 // This anchor is used to force the linker to link the DarwinModule.
 extern volatile int DarwinModuleAnchorSource;
 static int LLVM_ATTRIBUTE_UNUSED DarwinModuleAnchorDestination =
     DarwinModuleAnchorSource;
 
 // This anchor is used to force the linker to link the FuchsiaModule.
 extern volatile int FuchsiaModuleAnchorSource;
 static int LLVM_ATTRIBUTE_UNUSED FuchsiaModuleAnchorDestination =

clang-tools-extra/clang-tidy/cuda/CMakeLists.txt

+add_clang_library(clangTidyCudaModule
+  CudaTidyModule.cpp
+  LINK_LIBS
+  clangTidy
+  clangTidyUtils
+  )
+
+clang_target_link_libraries(clangTidyAlteraModule
+  PRIVATE
+  clangAnalysis
+  clangAST
+  clangASTMatchers
+  clangBasic
+  clangLex
+  )

clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp

+//===--- GoogleTidyModule.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 "../ClangTidy.h"
+#include "../ClangTidyCheck.h"
+#include "../ClangTidyModule.h"
+#include "../ClangTidyModuleRegistry.h"
+
+using namespace clang::ast_matchers;
+
+namespace clang {
+namespace tidy {
+namespace cuda {
+
+class CudaModule : public ClangTidyModule {
+public:
+  void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {}
+};
+
+// Register the CudaTidyModule using this statically initialized variable.
+static ClangTidyModuleRegistry::Add<CudaModule>
+    X("cuda-module", "Adds Cuda-related lint checks.");
+
+} // namespace cuda
+
+// This anchor is used to force the linker to link in the generated object file
+// and thus register the CudaModule.
+volatile int CudaModuleAnchorSource = 0;
+
+} // namespace tidy
+} // namespace clang

clang-tools-extra/test/clang-tidy/check_clang_tidy.py

def __init__(self, args, extra_args):
     self.has_check_messages = False
     self.has_check_notes = False
     self.fixes = MessagePrefix('CHECK-FIXES')
     self.messages = MessagePrefix('CHECK-MESSAGES')
     self.notes = MessagePrefix('CHECK-NOTES')
 
     file_name_with_extension = self.assume_file_name or self.input_file_name
     _, extension = os.path.splitext(file_name_with_extension)
-    if extension not in ['.c', '.hpp', '.m', '.mm']:
+    if extension not in ['.c', '.cu', '.hpp', '.m', '.mm']:
       extension = '.cpp'
     self.temp_file_name = self.temp_file_name + extension
 
     self.clang_extra_args = []
     self.clang_tidy_extra_args = extra_args
     if '--' in extra_args:
       i = self.clang_tidy_extra_args.index('--')
       self.clang_extra_args = self.clang_tidy_extra_args[i + 1:]
       self.clang_tidy_extra_args = self.clang_tidy_extra_args[:i]
 
     # If the test does not specify a config style, force an empty one; otherwise
     # auto-detection logic can discover a ".clang-tidy" file that is not related to
     # the test.
     if not any(
         [arg.startswith('-config=') for arg in self.clang_tidy_extra_args]):
       self.clang_tidy_extra_args.append('-config={}')
 
     if extension in ['.m', '.mm']:
       self.clang_extra_args = ['-fobjc-abi-version=2', '-fobjc-arc', '-fblocks'] + \
           self.clang_extra_args
 
-    if extension in ['.cpp', '.hpp', '.mm']:
+    if extension in ['.cpp', '.cu', '.hpp', '.mm']:
       self.clang_extra_args.append('-std=' + self.std)
 
+    # Tests should not rely on a certain cuda device being available on the machine,
+    # or a certain version of it
+    if extension == '.cu':
+      self.clang_extra_args.extend(["--no-cuda-version-check", "-nocudalib", "-nocudainc"])
+
     # Tests should not rely on STL being available, and instead provide mock
     # implementations of relevant APIs.
     self.clang_extra_args.append('-nostdinc++')
 
     if self.resource_dir is not None:
       self.clang_extra_args.append('-resource-dir=%s' % self.resource_dir)
 
   def read_input(self):

clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda-initializers.h

+// CUDA struct types with interesting initialization properties.
+// Keep in sync with clang/test/SemaCUDA/Inputs/cuda-initializers.h.
+
+// Base classes with different initializer variants.
+
+// trivial constructor -- allowed
+struct T {
+  int t;
+};
+
+// empty constructor
+struct EC {
+  int ec;
+  __device__ EC() {}     // -- allowed
+  __device__ EC(int) {}  // -- not allowed
+};
+
+// empty destructor
+struct ED {
+  __device__ ~ED() {}     // -- allowed
+};
+
+struct ECD {
+  __device__ ECD() {}     // -- allowed
+  __device__ ~ECD() {}    // -- allowed
+};
+
+// empty templated constructor -- allowed with no arguments
+struct ETC {
+  template <typename... T> __device__ ETC(T...) {}
+};
+
+// undefined constructor -- not allowed
+struct UC {
+  int uc;
+  __device__ UC();
+};
+
+// undefined destructor -- not allowed
+struct UD {
+  int ud;
+  __device__ ~UD();
+};
+
+// empty constructor w/ initializer list -- not allowed
+struct ECI {
+  int eci;
+  __device__ ECI() : eci(1) {}
+};
+
+// non-empty constructor -- not allowed
+struct NEC {
+  int nec;
+  __device__ NEC() { nec = 1; }
+};
+
+// non-empty destructor -- not allowed
+struct NED {
+  int ned;
+  __device__ ~NED() { ned = 1; }
+};
+
+// no-constructor,  virtual method -- not allowed
+struct NCV {
+  int ncv;
+  __device__ virtual void vm() {}
+};
+
+// virtual destructor -- not allowed.
+struct VD {
+  __device__ virtual ~VD() {}
+};
+
+// dynamic in-class field initializer -- not allowed
+__device__ int f();
+struct NCF {
+  int ncf = f();
+};
+
+// static in-class field initializer.  NVCC does not allow it, but
+// clang generates static initializer for this, so we'll accept it.
+// We still can't use it on __shared__ vars as they don't allow *any*
+// initializers.
+struct NCFS {
+  int ncfs = 3;
+};
+
+// undefined templated constructor -- not allowed
+struct UTC {
+  template <typename... T> __device__ UTC(T...);
+};
+
+// non-empty templated constructor -- not allowed
+struct NETC {
+  int netc;
+  template <typename... T> __device__ NETC(T...) { netc = 1; }
+};
+
+// Regular base class -- allowed
+struct T_B_T : T {};
+
+// Incapsulated object of allowed class -- allowed
+struct T_F_T {
+  T t;
+};
+
+// array of allowed objects -- allowed
+struct T_FA_T {
+  T t[2];
+};
+
+
+// Calling empty base class initializer is OK
+struct EC_I_EC : EC {
+  __device__ EC_I_EC() : EC() {}
+};
+
+// .. though passing arguments is not allowed.
+struct EC_I_EC1 : EC {
+  __device__ EC_I_EC1() : EC(1) {}
+};
+
+// Virtual base class -- not allowed
+struct T_V_T : virtual T {};
+
+// Inherited from or incapsulated class with non-empty constructor --
+// not allowed
+struct T_B_NEC : NEC {};
+struct T_F_NEC {
+  NEC nec;
+};
+struct T_FA_NEC {
+  NEC nec[2];
+};
+
+
+// Inherited from or incapsulated class with non-empty desstructor --
+// not allowed
+struct T_B_NED : NED {};
+struct T_F_NED {
+  NED ned;
+};
+struct T_FA_NED {
+  NED ned[2];
+};

clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda.h

+/* Minimal declarations for CUDA support.  Testing purposes only. */
+
+#include <stddef.h>
+
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+typedef struct cudaStream *cudaStream_t;
+typedef enum cudaError {} cudaError_t;
+extern "C" int cudaConfigureCall(dim3 gridSize, dim3 blockSize,
+                                 size_t sharedSize = 0,
+                                 cudaStream_t stream = 0);
+extern "C" int __cudaPushCallConfiguration(dim3 gridSize, dim3 blockSize,
+                                           size_t sharedSize = 0,
+                                           cudaStream_t stream = 0);
+extern "C" cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim,
+                                        dim3 blockDim, void **args,
+                                        size_t sharedMem, cudaStream_t stream);
+
+extern "C" __device__ int printf(const char*, ...);