Cuda Check for ignored return errors from api calls to cuda

[barcisz]

Summary

Add cuda-unchecked-api-call check

Motivation

cuda_runtime.h header is included by default in all CUDA files. Functions defined there provide an interface between regular c/c++ and the CUDA driver. Most of them return errors in the form of cudaError_t enum. The specific errors returned for each call are described in the CUDA documentation. However, it is really easy to ignore those errors and that the calls can fail. This can happen for call-related reasons, such as the address provided to the call being in the wrong address space (host vs. device), or for external reasons, such as problems with the CUDA device or with the driver. In any case, those can and will impact your execution, in the best case leading to a crash or, in the worst case, resulting in an incorrect program state.

Behavior

The cuda-unchecked-api-call check checks whether the value returned by a call to the CUDA API is unused, similar to the bugprone-unused-return-value check (however, it is more specific and it’s more likely to be used by the people that need it). It defines a CUDA API function as a function that:

• returns a type cudaError_t
• Is included through a header whose suffix is cuda_runtime.h (this allows for cuda_runtime.h to be replaced by for example _cuda_runtime.h in case some buck setup is configured like that)

Automatic fixes

The lint check can be configured to produce a FixItHint that puts the value from the CUDA API call inside a macro or a function handler. You can specify the error handler for your project by setting the HandlerName option for the cuda-unchecked-api-call. Here is an example of how this fix can transform unhandled code from:

void foo() {
  cudaDeviceReset();
}

to

void foo() {
  C10_CUDA_CHECK(cudaDeviceReset());
}

The specific handler used for this example is taken from PyTorch and its definition can be found here.

Limiting the allowed handlers

Since the projects may only have a limited set of handlers for the errors thrown by CUDA, there is also an option to limit the allowed ways to handle the value of the error returned by the check by setting the AcceptedHandlers option to a comma-separated list of names (which can be scoped) of those allowed error handlers. If HandlerName is set then it will also be implicitly added to that list. This option also works with dummy macros that pass the error through and do not do anything (which may be present in the code for performance reasons).

Parent diffs

This diff relies on D133801 and D133436 to properly run, so feel free to take a look at those as well

[barcisz]

Added some explanation comments

[tra]

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

121 ↗

(On Diff #459870)

The comment is incorrect. These flags have nothing to do with GPU availability, but rather with CUDA SDK which is normally expected to provide the 'standard' set of CUDA headers and libdevice bitcode.

clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-api-call-function-handler.cu

2

This does not look right.

[barcisz]

More explanation comments for the check's code

[barcisz]

Rebase and better comments for cuda-related decisions in the check

[barcisz]

Removed unnecessary cuda-related compilation flags from tests

[barcisz]

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

121 ↗

(On Diff #459870)

Yes by CUDA I meant cuda library here but I'll change it accordingly

clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-api-call-function-handler.cu

2

Fixed

[barcisz]

Update base to D133801

[barcisz]

Documentation and small check-message-related bugfixes

[barcisz]

Brought back different message prefix for when AcceptedHandlers is set

[barcisz]

Removed unneeded headers from the check

[barcisz]

Removed unneeded headers from the check

[barcisz]

Rebase

[barcisz]

Use header guards instead of pragma

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

 add_clang_library(clangTidyCudaModule
   CudaTidyModule.cpp
+  UnsafeApiCallCheck.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 --------------------------------===//
+//===--- CudaTidyModule.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"
+#include "UnsafeApiCallCheck.h"
 
 using namespace clang::ast_matchers;
 
 namespace clang {
 namespace tidy {
 namespace cuda {
 
 class CudaModule : public ClangTidyModule {
 public:
-  void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {}
+  void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {
+    CheckFactories.registerCheck<UnsafeApiCallCheck>("cuda-unsafe-api-call");
+  }
 };
 
 // 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/clang-tidy/cuda/UnsafeApiCallCheck.h

+//===--- UnsafeApiCallCheck.h - 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "../ClangTidyCheck.h"
+#include "llvm/ADT/StringSet.h"
+#include <memory>
+#include <unordered_set>
+
+namespace clang {
+namespace tidy {
+namespace cuda {
+
+/// Checks for whether the possible errors with the CUDA API invocations have
+/// been handled.
+///
+/// Calls to CUDA API can sometimes fail to perform the action. This may happen
+/// due to a driver malfunction, lack of permissions, lack of a GPU, or a
+/// multitude of other reasons. Such errors are returned by those API calls and
+/// should be handled in some way.
+/// The check provides the following options:
+///  - "HandlerName" (optional):
+///      specifies the name of the function or the macro to which the return
+///      value of the API call should be passed. This effectively automates the
+///      process of adding the error checks in question for projects that have
+///      such a mechanism implemented in them.
+///  - "AcceptedHandlers" (optional):
+///      a comma-separated list specifying the only accepted handling
+///      functions/macros into which the error from the api call can be passed.
+///      If not specified all ways to handle the error that do not just ignore
+///      the output value are accepted. The handlers may have scope specifiers
+///      included in them, but if so then the full qualified name (with all
+///      namespaces explicitly stated) has to be provided (for the performance
+///      sake). If the handler set in the "HandlerName" is not in the list of
+///      accepted handlers then it gets added to it automatially.
+///
+/// Since the behavior of the check is significantly different when the
+/// "AcceptedHandlers" option is set, the implementation is essentially split
+/// into 2 paths, as highlighted by the comments near declarations.
+class UnsafeApiCallCheck : public ClangTidyCheck {
+  class PPCallbacks;
+
+  // For gathering api calls with an unused value - only those nodes
+  // can have a FixItHint when we limit the accepted handlers.
+  //
+  // Only used when "AcceptedHandlers" is set
+  class UnusedValueCallback
+      : public clang::ast_matchers::MatchFinder::MatchCallback {
+  public:
+    UnusedValueCallback(UnsafeApiCallCheck *check) : Check(check) {}
+    void
+    run(const clang::ast_matchers::MatchFinder::MatchResult &Result) override;
+    void onStartOfTranslationUnit() override;
+
+  private:
+    UnsafeApiCallCheck *Check;
+  };
+
+public:
+  UnsafeApiCallCheck(llvm::StringRef Name,
+                     clang::tidy::ClangTidyContext *Context);
+
+  void registerPPCallbacks(const SourceManager &SM, Preprocessor *PP,
+                           Preprocessor *ModuleExpanderPP) override;
+  void registerMatchers(clang::ast_matchers::MatchFinder *Finder) override;
+  void
+  check(const clang::ast_matchers::MatchFinder::MatchResult &Result) override;
+  void storeOptions(ClangTidyOptions::OptionMap &Opts) override;
+
+private:
+  const std::string HandlerName;
+
+  // Only used when "AcceptedHandlers" is set
+  void
+  checkUnusedValue(const clang::ast_matchers::MatchFinder::MatchResult &Result);
+  // Only used when "AcceptedHandlers" is not set
+  void
+  checkBadHandler(const clang::ast_matchers::MatchFinder::MatchResult &Result);
+
+  // Only used when "AcceptedHandlers" is not set
+  void registerUnusedValueMatchers(clang::ast_matchers::MatchFinder *Finder);
+  // Only used when "AcceptedHandlers" is set
+  void registerBadlyHandledMatchers(clang::ast_matchers::MatchFinder *Finder);
+
+  const std::string AcceptedHandlersList; // Data store for AcceptedHandlersSet
+  const llvm::StringSet<llvm::MallocAllocator> AcceptedHandlersSet;
+  // Generates AcceptedHandlersSet from AcceptedHandlersList
+  static llvm::StringSet<llvm::MallocAllocator>
+  splitAcceptedHandlers(const llvm::StringRef &AcceptedHandlers,
+                        const llvm::StringRef &HandlerName);
+  bool limitAcceptedHandlers();
+
+  // Only used when "AcceptedHandlers" is set
+  std::unordered_set<SourceLocation,
+                     std::function<unsigned(const SourceLocation &)>>
+      AcceptedHandlerMacroLocations;
+  std::unordered_set<const Stmt *> UnusedValueNodes;
+  std::unique_ptr<UnusedValueCallback> UnusedValueCallbackInstance;
+};
+
+} // namespace cuda
+} // namespace tidy
+} // namespace clang

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

+//===--- UnsafeApiCallCheck.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 "UnsafeApiCallCheck.h"
+#include "../utils/Matchers.h"
+#include "clang/AST/ASTContext.h"
+#include "clang/ASTMatchers/ASTMatchFinder.h"
+#include "clang/Lex/Preprocessor.h"
+#include "clang/Tooling/FixIt.h"
+
+#include <functional>
+#include <sstream>
+
+using namespace clang::ast_matchers;
+
+namespace clang {
+namespace tidy {
+namespace cuda {
+
+namespace {
+
+constexpr auto HandlerNameOptionName = "HandlerName";
+constexpr auto AcceptedHandlersOptionName = "AcceptedHandlers";
+
+} // namespace
+
+UnsafeApiCallCheck::UnsafeApiCallCheck(llvm::StringRef Name,
+                                       clang::tidy::ClangTidyContext *Context)
+    : ClangTidyCheck(Name, Context),
+      HandlerName(Options.get(HandlerNameOptionName, "")),
+      AcceptedHandlersList(Options.get(AcceptedHandlersOptionName, "")),
+      AcceptedHandlersSet(
+          splitAcceptedHandlers(AcceptedHandlersList, HandlerName)),
+      AcceptedHandlerMacroLocations(
+          8, [](const SourceLocation &sLoc) { return sLoc.getHashValue(); }) {
+  // If an empty string was inserted then that means that there was an empty
+  // accepted handler in the list
+  if (AcceptedHandlersSet.find("") != AcceptedHandlersSet.end()) {
+    configurationDiag(
+        "Empty handler name found in the list of accepted handlers",
+        DiagnosticIDs::Error);
+  }
+}
+
+llvm::StringSet<llvm::MallocAllocator>
+UnsafeApiCallCheck::splitAcceptedHandlers(
+    const llvm::StringRef &AcceptedHandlers,
+    const llvm::StringRef &HandlerName) {
+  // Check the case for when the accepted handlers are empty since otherwise
+  // split(...) will still fill the vector with an empty element
+  if (AcceptedHandlers.trim().empty()) {
+    return llvm::StringSet<llvm::MallocAllocator>();
+  }
+  llvm::SmallVector<llvm::StringRef> AcceptedHandlersVector;
+  AcceptedHandlers.split(AcceptedHandlersVector, ',');
+
+  llvm::StringSet<llvm::MallocAllocator> AcceptedHandlersSet;
+  for (auto AcceptedHandler : AcceptedHandlersVector) {
+    AcceptedHandlersSet.insert(AcceptedHandler.trim());
+  }
+
+  // If the handler for FixItHints is set then add it to
+  if (!AcceptedHandlersSet.empty() && !HandlerName.empty()) {
+    AcceptedHandlersSet.insert(HandlerName);
+  }
+
+  return AcceptedHandlersSet;
+}
+
+void UnsafeApiCallCheck::storeOptions(ClangTidyOptions::OptionMap &Opts) {
+  Options.store(Opts, HandlerNameOptionName, HandlerName);
+  Options.store(Opts, AcceptedHandlersOptionName, AcceptedHandlersList);
+}
+
+inline bool UnsafeApiCallCheck::limitAcceptedHandlers() {
+  return !AcceptedHandlersSet.empty();
+}
+
+// Used for finding the occurences of accepted handler macros in the source
+// code.
+class UnsafeApiCallCheck::PPCallbacks : public clang::PPCallbacks {
+public:
+  PPCallbacks(UnsafeApiCallCheck *Check, const SourceManager &SM)
+      : Check(Check), SM(SM) {}
+
+  void MacroExpands(const Token &MacroNameTok, const MacroDefinition &MD,
+                    SourceRange Range, const MacroArgs *Args) override {
+    if (Check->AcceptedHandlersSet.find(
+            MacroNameTok.getIdentifierInfo()->getName()) !=
+        Check->AcceptedHandlersSet.end()) {
+      Check->AcceptedHandlerMacroLocations.insert(MacroNameTok.getLocation());
+    }
+  }
+
+private:
+  UnsafeApiCallCheck *Check;
+};
+
+void UnsafeApiCallCheck::registerPPCallbacks(const SourceManager &SM,
+                                             Preprocessor *PP,
+                                             Preprocessor *ModuleExpanderPP) {
+  if (limitAcceptedHandlers()) {
+    ModuleExpanderPP->addPPCallbacks(std::make_unique<PPCallbacks>(this, SM));
+  }
+}
+
+namespace {
+
+// Check if the declaration is in a specific header based on a condition
+AST_MATCHER_P(Decl, isInSourceFile, std::function<bool(const StringRef &)>,
+              SourceFileNameCond) {
+  auto Loc = Node.getLocation();
+  const auto &SM = Finder->getASTContext().getSourceManager();
+  while (Loc.isValid()) {
+    if (SourceFileNameCond(SM.getFilename(Loc))) {
+      return true;
+    }
+    Loc = SM.getIncludeLoc(SM.getFileID(Loc));
+  }
+  return false;
+}
+
+// Check if the name of the declaration matches a specific condition
+AST_MATCHER_P(NamedDecl, hasName, std::function<bool(const StringRef &)>,
+              DeclNameCond) {
+  return DeclNameCond(Node.getName());
+}
+
+// Check if the fully qualified name of the declaration matches a specific
+// condition
+AST_MATCHER_P(NamedDecl, hasQualName, std::function<bool(const StringRef &)>,
+              DeclNameCond) {
+  return DeclNameCond(Node.getQualifiedNameAsString());
+}
+
+constexpr auto UnusedValueBinding = "UnusedValueCall";
+constexpr auto badlyHandledBinding = "badlyHandledCall";
+
+// Common matchers for both unlimited and limited accepted handlers.
+const auto HostFunction = functionDecl(unless(anyOf(
+    hasAttr(attr::CUDADevice),
+    hasAttr(attr::CUDAGlobal)))); // CUDA API cannot be called from device code
+const auto ApiCallExpression = callExpr(
+    callee(functionDecl(isInSourceFile([](StringRef FileName) {
+                          return FileName.endswith("cuda_runtime.h") ||
+                                 FileName.endswith("cuda_runtime_wrapper.h");
+                        }), // All CUDA API is included from the cuda_runtime.h
+                            // header or __cuda_runtime_wrapper.h
+                        returns(asString("cudaError_t")))));
+
+} // namespace
+
+void UnsafeApiCallCheck::UnusedValueCallback::run(
+    const MatchFinder::MatchResult &Result) {
+  auto Node = Result.Nodes.getNodeAs<Stmt>(UnusedValueBinding);
+  assert(Node);
+  Check->UnusedValueNodes.insert(Node);
+}
+
+void UnsafeApiCallCheck::UnusedValueCallback::onStartOfTranslationUnit() {
+  Check->UnusedValueNodes.clear();
+}
+
+void UnsafeApiCallCheck::registerMatchers(MatchFinder *Finder) {
+  if (limitAcceptedHandlers()) {
+    registerBadlyHandledMatchers(Finder);
+  } else {
+    registerUnusedValueMatchers(Finder);
+  }
+}
+
+void UnsafeApiCallCheck::registerUnusedValueMatchers(MatchFinder *Finder) {
+  const auto UnusedValue =
+      matchers::isValueUnused(stmt(ApiCallExpression.bind(UnusedValueBinding)));
+  Finder->addMatcher(functionDecl(HostFunction, hasBody(UnusedValue)), this);
+}
+
+void UnsafeApiCallCheck::registerBadlyHandledMatchers(MatchFinder *Finder) {
+  const auto UnusedValue =
+      matchers::isValueUnused(stmt(ApiCallExpression.bind(UnusedValueBinding)));
+  UnusedValueCallbackInstance = std::make_unique<UnusedValueCallback>(this);
+  Finder->addMatcher(functionDecl(HostFunction, hasBody(UnusedValue)),
+                     UnusedValueCallbackInstance.get());
+
+  const auto AcceptedHandlerPred = [this](const StringRef &Name) {
+    return AcceptedHandlersSet.contains(Name);
+  };
+
+  const auto AcceptedHandlerDecl = functionDecl(
+      anyOf(hasName(AcceptedHandlerPred), hasQualName(AcceptedHandlerPred)));
+  const auto AcceptedHandlerParent = callExpr(callee(AcceptedHandlerDecl));
+
+  Finder->addMatcher(
+      functionDecl(
+          HostFunction,
+          forEachDescendant(stmt(ApiCallExpression.bind(badlyHandledBinding),
+                                 unless(hasParent(AcceptedHandlerParent))))),
+      this);
+}
+
+namespace {
+
+constexpr auto HandlerMsg =
+    "Consider wrapping it with a call to an error handler:";
+constexpr auto NoHandlerMsg =
+    "Consider adding logic to check if an error has been returned "
+    "or specify the error handler for this project.";
+constexpr auto MacroMsg =
+    "Consider adding logic to check if an error has been returned.";
+
+inline bool isStmtInMacro(const Stmt *const Stmt) {
+  return Stmt->getBeginLoc().isInvalid() || Stmt->getBeginLoc().isMacroID() ||
+         Stmt->getEndLoc().isInvalid() || Stmt->getEndLoc().isMacroID();
+}
+
+} // namespace
+
+void UnsafeApiCallCheck::check(const MatchFinder::MatchResult &Result) {
+  if (limitAcceptedHandlers()) {
+    checkBadHandler(Result);
+  } else {
+    checkUnusedValue(Result);
+  }
+}
+
+void UnsafeApiCallCheck::checkUnusedValue(
+    const MatchFinder::MatchResult &Result) {
+  const auto ApiCallNode = Result.Nodes.getNodeAs<Stmt>(UnusedValueBinding);
+  assert(ApiCallNode);
+
+  // This disables the check for arguments inside macros, since we assume that
+  // such a macro is intended as a handler (even if it just passes the argument
+  // right through)
+  if (Result.SourceManager->isMacroArgExpansion(ApiCallNode->getBeginLoc())) {
+    return;
+  }
+
+  const auto DiagnosticBuilder =
+      diag(ApiCallNode->getBeginLoc(), "Unchecked CUDA API call. ");
+  if (HandlerName.empty()) {
+    DiagnosticBuilder << NoHandlerMsg;
+  } else if (isStmtInMacro(ApiCallNode)) {
+    DiagnosticBuilder << MacroMsg;
+  } else {
+    DiagnosticBuilder << HandlerMsg
+                      << FixItHint::CreateReplacement(
+                             ApiCallNode->getSourceRange(),
+                             (HandlerName + "(" +
+                              tooling::fixit::getText(
+                                  ApiCallNode->getSourceRange(),
+                                  *Result.Context) +
+                              ")")
+                                 .str());
+  }
+}
+
+void UnsafeApiCallCheck::checkBadHandler(
+    const MatchFinder::MatchResult &Result) {
+  const auto ApiCallNode = Result.Nodes.getNodeAs<Stmt>(badlyHandledBinding);
+  assert(ApiCallNode);
+
+  const auto ApiCallNodeMacroLocation = Result.SourceManager->getExpansionLoc(
+      Result.SourceManager->getMacroArgExpandedLocation(
+          ApiCallNode->getBeginLoc()));
+
+  // This disables the check for arguments inside macros, since we assume that
+  // such a macro is intended as a handler (even if it just passes the argument
+  // right through)
+  if (Result.SourceManager->isMacroArgExpansion(ApiCallNode->getBeginLoc()) &&
+      AcceptedHandlerMacroLocations.find(ApiCallNodeMacroLocation) !=
+          AcceptedHandlerMacroLocations.end()) {
+    return;
+  }
+
+  const auto DiagnosticBuilder =
+      diag(ApiCallNode->getBeginLoc(), "CUDA API call not checked properly. ");
+
+  if (HandlerName.empty()) {
+    DiagnosticBuilder << NoHandlerMsg;
+  } else if (isStmtInMacro(ApiCallNode) ||
+             UnusedValueNodes.find(ApiCallNode) == UnusedValueNodes.end()) {
+    DiagnosticBuilder << "Consider wrapping it with a call to `" << HandlerName
+                      << '`';
+  } else {
+    DiagnosticBuilder << HandlerMsg
+                      << FixItHint::CreateReplacement(
+                             ApiCallNode->getSourceRange(),
+                             (HandlerName + "(" +
+                              tooling::fixit::getText(
+                                  ApiCallNode->getSourceRange(),
+                                  *Result.Context) +
+                              ")")
+                                 .str());
+  }
+}
+
+} // namespace cuda
+} // namespace tidy
+} // namespace clang

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;
+typedef enum cudaError {
+  cudaErrorInvalidValue,
+  cudaErrorMemoryAllocation
+} 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*, ...);

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

+#include "cuda.h"
+
+cudaError_t cudaDeviceReset();

clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-api-call-function-handler.cu

+// RUN: %check_clang_tidy %s cuda-unsafe-api-call %t -- \
+// RUN:   -config="{CheckOptions: \
+// RUN:             [{key: cuda-unsafe-api-call.HandlerName, \
+// RUN:               value: 'cudaHandler'}, \
+// RUN:              {key: cuda-unsafe-api-call.AcceptedHandlers, \
+// RUN:               value: 'CUDA_HANDLER, DUMMY_CUDA_HANDLER, \
+// RUN:                       alternative::cudaAlternativeHandler, \
+// RUN:                       cudaOtherAlternativeHandler, bad::cudaBadHandler'}] \
+// RUN:             }" \
+// RUN:   -- -isystem %clang_tidy_headers -std=c++14
+#include <cuda/cuda_runtime.h>
+
+#define DUMMY_CUDA_HANDLER(stmt) stmt
+#define CUDA_HANDLER(stmt) do {auto err = stmt;} while(0)
+#define API_CALL() do {cudaDeviceReset();} while(0)
+#define HANDLED_API_CALL() do {int err2 = cudaDeviceReset();} while(0)
+
+void cudaHandler();
+void cudaHandler(cudaError_t error);
+void badCudaHandler(cudaError_t error);
+
+namespace alternative {
+
+void cudaAlternativeHandler(cudaError_t error);
+
+void cudaOtherAlternativeHandler(cudaError_t error);
+
+} // namespace alternative
+
+void bad() {
+  API_CALL();
+  // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: CUDA API call not checked properly.
+  // There isn't supposed to be a fix here since it's a macro call
+
+  HANDLED_API_CALL();
+  // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: CUDA API call not checked properly.
+  // There isn't supposed to be a fix here since it's a macro call
+
+  cudaDeviceReset();
+  // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: CUDA API call not checked properly.
+  // CHECK-FIXES:  {{^}}  cudaHandler(cudaDeviceReset());{{$}}
+  cudaHandler();
+
+  if (true)
+    cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: CUDA API call not checked properly.
+    // CHECK-FIXES:  {{^}}    cudaHandler(cudaDeviceReset());{{$}}
+
+  badCudaHandler(cudaDeviceReset());
+  // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: CUDA API call not checked properly.
+  // There isn't supposed to be a fix here since the result value is not unused
+
+  int err = cudaDeviceReset();
+  // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: CUDA API call not checked properly.
+  // There isn't supposed to be a fix here since the result value is not unused
+
+  if (cudaDeviceReset()) {
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: CUDA API call not checked properly.
+    // There isn't supposed to be a fix here since the result value is not unused
+    return;
+  }
+
+}
+
+void good() {
+  cudaHandler(cudaDeviceReset());
+  alternative::cudaAlternativeHandler(cudaDeviceReset());
+  alternative::cudaOtherAlternativeHandler(cudaDeviceReset());
+  CUDA_HANDLER(cudaDeviceReset() + 1);
+  DUMMY_CUDA_HANDLER(cudaDeviceReset());
+}

clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-api-call-macro-handler.cu

+// RUN: %check_clang_tidy %s cuda-unsafe-api-call %t -- \
+// RUN:   -config="{CheckOptions: \
+// RUN:             [{key: cuda-unsafe-api-call.HandlerName, \
+// RUN:               value: 'CUDA_HANDLER'}] \
+// RUN:             }" \
+// RUN:   -- -isystem %clang_tidy_headers -std=c++14
+#include <cuda/cuda_runtime.h>
+
+class DummyContainer {
+ public:
+  int* begin();
+  int* end();
+};
+
+#define DUMMY_CUDA_HANDLER(stmt) stmt
+#define CUDA_HANDLER(stmt) do {auto err = stmt;} while(0)
+#define API_CALL() do {cudaDeviceReset();} while(0)
+
+void errorCheck();
+void errorCheck(cudaError_t error);
+
+void bad() {
+  API_CALL();
+  // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+  // There isn't supposed to be a fix here since it's a macro call
+
+  cudaDeviceReset();
+  // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+  // CHECK-FIXES:  {{^}}  CUDA_HANDLER(cudaDeviceReset());{{$}}
+  errorCheck();
+
+  if (true)
+    cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}    CUDA_HANDLER(cudaDeviceReset());{{$}}
+
+  while (true)
+    cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}    CUDA_HANDLER(cudaDeviceReset());{{$}}
+
+  do
+    cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}    CUDA_HANDLER(cudaDeviceReset());{{$}}
+  while(false);
+
+  switch (0) {
+    case 0:
+      cudaDeviceReset();
+      // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+      // CHECK-FIXES:  {{^}}      CUDA_HANDLER(cudaDeviceReset());{{$}}
+  }
+
+  for(
+    cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}    CUDA_HANDLER(cudaDeviceReset());{{$}}
+    ;
+    cudaDeviceReset()
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}    CUDA_HANDLER(cudaDeviceReset()){{$}}
+  ) cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}  ) CUDA_HANDLER(cudaDeviceReset());{{$}}
+
+  for(int i : DummyContainer())
+    cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}    CUDA_HANDLER(cudaDeviceReset());{{$}}
+
+  auto x = ({
+    cudaDeviceReset();
+    // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Unchecked CUDA API call.
+    // CHECK-FIXES:  {{^}}    CUDA_HANDLER(cudaDeviceReset());{{$}}
+    true;
+  });
+}
+
+int good() {
+  DUMMY_CUDA_HANDLER(cudaDeviceReset());
+
+  if (cudaDeviceReset()) {
+    return 0;
+  }
+
+  switch (cudaDeviceReset()) {
+    case cudaErrorInvalidValue: return 1;
+    case cudaErrorMemoryAllocation: return 2;
+    default: return 3;
+  }
+
+  auto err = ({cudaDeviceReset();});
+  // NOTE: We don't check that `errorCheck()` actually handles the error; we just assume it does.
+  errorCheck(cudaDeviceReset());
+}