barcisz updated this revision to Diff 460522.
barcisz added a comment.
Herald added a subscriber: mgorny.
rebase
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D133942/new/
https://reviews.llvm.org/D133942
Files:
clang-tools-extra/clang-tidy/cuda/CMakeLists.txt
clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp
clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.cpp
clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.h
clang-tools-extra/docs/ReleaseNotes.rst
clang-tools-extra/docs/clang-tidy/checks/cuda/unsafe-kernel-call.rst
clang-tools-extra/docs/clang-tidy/checks/list.rst
clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h
clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-function-handler.cu
clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-macro-handler.cu
Index: clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-macro-handler.cu
===================================================================
--- /dev/null
+++ clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-macro-handler.cu
@@ -0,0 +1,116 @@
+// RUN: %check_clang_tidy %s cuda-unsafe-kernel-call %t -- \
+// RUN: -config="{CheckOptions: \
+// RUN: [{key: cuda-unsafe-kernel-call.HandlerName, \
+// RUN: value: 'CUDA_CHECK_KERNEL'}, \
+// RUN: {key: cuda-unsafe-kernel-call.AcceptedHandlers, \
+// RUN: value: 'ALTERNATIVE_CUDA_CHECK_KERNEL, cudaCheckKernel, \
+// RUN: alternative::alternativeCudaCheckKernel, \
+// RUN: otherAlternativeCudaCheckKernel'}] \
+// RUN: }" \
+// RUN: -- -isystem %clang_tidy_headers
+
+#include <cuda/cuda_runtime.h>
+
+#define CUDA_CHECK_KERNEL() do {} while(0)
+
+#define ALTERNATIVE_CUDA_CHECK_KERNEL() CUDA_CHECK_KERNEL()
+
+void cudaCheckKernel();
+
+namespace alternative {
+
+void alternativeCudaCheckKernel();
+void otherAlternativeCudaCheckKernel();
+
+}
+
+__global__
+void b();
+
+#define KERNEL_CALL() do {b<<<1, 2>>>();} while(0)
+
+void errorCheck() {
+ auto err = cudaGetLastError();
+}
+
+void bad() {
+ b<<<1, 2>>>(); // sample comment
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+
+ KERNEL_CALL(); // sample comment
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // There isn't supposed to be a fix here since it's a macro call
+
+ if(true)
+ b<<<1, 2>>>() ; // Brackets omitted purposefully, since they create an additional AST node
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ else {
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ }
+ auto err = cudaGetLastError();
+
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ if (true)
+ cudaGetLastError();
+
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ for(;;)
+ auto err2 = cudaGetLastError(); // Brackets omitted purposefully, since they create an additional AST node
+
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ auto err3 = true ? 1 : cudaGetLastError();
+
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ auto err4 = cudaDeviceReset() + cudaGetLastError();
+
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // Calling an error-checking function after a kernel is not considered safe.
+ errorCheck();
+}
+
+void good() {
+ b<<<1, 2>>>();; /* The semicolons are here because the
+ detection of the macro is done with a lexer */ ;
+ CUDA_CHECK_KERNEL();
+
+ b<<<1, 2>>>();
+ ALTERNATIVE_CUDA_CHECK_KERNEL();
+
+ b<<<1, 2>>>();
+ alternative::alternativeCudaCheckKernel();
+
+ b<<<1, 2>>>();
+ alternative::otherAlternativeCudaCheckKernel();
+
+ b<<<1, 2>>>();
+ switch(1 + cudaGetLastError()) {
+ default:;
+ }
+
+ b<<<1, 2>>>();
+ if(3 < cudaGetLastError()) {
+ 1;
+ } else {
+ 2;
+ }
+
+ b<<<1, 2>>>();
+ for(int i = cudaGetLastError();;);
+
+ b<<<1, 2>>>();
+ do {
+ do {
+ do {
+ auto err2 = cudaGetLastError();
+ } while(0);
+ } while(0);
+ } while(0);
+}
Index: clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-function-handler.cu
===================================================================
--- /dev/null
+++ clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-function-handler.cu
@@ -0,0 +1,150 @@
+// RUN: %check_clang_tidy %s cuda-unsafe-kernel-call %t -- \
+// RUN: -config="{CheckOptions: \
+// RUN: [{key: cuda-unsafe-kernel-call.HandlerName, \
+// RUN: value: 'errorCheck'}] \
+// RUN: }" \
+// RUN: -- -isystem %clang_tidy_headers
+
+#include <cuda/cuda_runtime.h>
+
+__global__
+void b();
+
+void general();
+
+void errorCheck() {
+ auto err = cudaGetLastError();
+}
+
+void bad_next_line_stmt() {
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ general();
+
+ b<<<1, 2>>>(); /* some */ /* comments */ // present
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ general();
+
+ if (true) // Dummy comment
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} if (true) { // Dummy comment{{$}}
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>();{{$}}
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ else // Dummy comment
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} } else { // Dummy comment{{$}}
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>();{{$}}
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ // CHECK-FIXES: {{^}} }{{$}}
+ general();
+
+ while (true) b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} while (true) { b<<<1, 2>>>();{{$}}
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ // CHECK-FIXES: {{^}} }{{$}}
+ general();
+
+ for (;;) // Dummy comment
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} for (;;) { // Dummy comment{{$}}
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>();{{$}}
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ // CHECK-FIXES: {{^}} }{{$}}
+ general();
+
+ if (true) {
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ general();
+ } else {
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ general();
+ }
+
+ while(true) {
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ general();
+ }
+
+ for (;;) {
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ general();
+ }
+
+ do {
+ b<<<1, 2>>>();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} errorCheck();{{$}}
+ general();
+ } while(true);
+}
+
+void bad_same_line_stmt() {
+ b<<<1, 2>>>(); general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>(); errorCheck(); general();{{$}}
+
+ b<<<1, 2>>>(); /* hello */ /* there */ general(); // kenobi
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>(); errorCheck(); /* hello */ /* there */ general(); // kenobi{{$}}
+
+ if (true) // Dummy comment
+ b<<<1, 2>>>(); /* comment */ general(); // comment
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} {b<<<1, 2>>>(); errorCheck();} /* comment */ general(); // comment{{$}}
+
+ while (true) b<<<1, 2>>>(); general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} while (true) {b<<<1, 2>>>(); errorCheck();} general();{{$}}
+
+ for (;;) // Dummy comment
+ b<<<1, 2>>>(); /* comment */ general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} {b<<<1, 2>>>(); errorCheck();} /* comment */ general();{{$}}
+
+ if (true) {
+ b<<<1, 2>>>(); general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>(); errorCheck(); general();{{$}}
+ } else {
+ b<<<1, 2>>>(); /* comment */ general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>(); errorCheck(); /* comment */ general();{{$}}
+ }
+
+ while(true) {
+ b<<<1, 2>>>(); /* comment */ general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>(); errorCheck(); /* comment */ general();{{$}}
+ }
+
+ for (;;) {
+ b<<<1, 2>>>(); general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>(); errorCheck(); general();{{$}}
+ }
+
+ do {
+ b<<<1, 2>>>(); /* comment */ general();
+ // CHECK-MESSAGES: :[[@LINE-1]]:{{[0-9]+}}: warning: Possible unchecked error after a kernel launch.
+ // CHECK-FIXES: {{^}} b<<<1, 2>>>(); errorCheck(); /* comment */ general();{{$}}
+ } while(true);
+}
+
+void good() {
+ b<<<1, 2>>>();
+ errorCheck(); // Here the function call works because the handler is set to its name
+}
Index: clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h
===================================================================
--- clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h
+++ clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h
@@ -1,3 +1,4 @@
#include "cuda.h"
cudaError_t cudaDeviceReset();
+cudaError_t cudaGetLastError();
Index: clang-tools-extra/docs/clang-tidy/checks/list.rst
===================================================================
--- clang-tools-extra/docs/clang-tidy/checks/list.rst
+++ clang-tools-extra/docs/clang-tidy/checks/list.rst
@@ -199,8 +199,9 @@
`cppcoreguidelines-pro-type-vararg <cppcoreguidelines/pro-type-vararg.html>`_,
`cppcoreguidelines-slicing <cppcoreguidelines/slicing.html>`_,
`cppcoreguidelines-special-member-functions <cppcoreguidelines/special-member-functions.html>`_,
- `cuda-unsafe-api-call <cuda/unsafe-api-call.html>`_, "Yes"
`cppcoreguidelines-virtual-class-destructor <cppcoreguidelines/virtual-class-destructor.html>`_, "Yes"
+ `cuda-unsafe-api-call <cuda/unsafe-api-call.html>`_, "Yes"
+ `cuda-unsafe-kernel-call <cuda/unsafe-kernel-call.html>`_, "Yes"
`darwin-avoid-spinlock <darwin/avoid-spinlock.html>`_,
`darwin-dispatch-once-nonstatic <darwin/dispatch-once-nonstatic.html>`_, "Yes"
`fuchsia-default-arguments-calls <fuchsia/default-arguments-calls.html>`_,
Index: clang-tools-extra/docs/clang-tidy/checks/cuda/unsafe-kernel-call.rst
===================================================================
--- /dev/null
+++ clang-tools-extra/docs/clang-tidy/checks/cuda/unsafe-kernel-call.rst
@@ -0,0 +1,6 @@
+.. title:: clang-tidy - cuda-unsafe-kernel-call
+
+cuda-unsafe-kernel-call
+=======================
+
+FIXME: Describe what patterns does the check detect and why. Give examples.
Index: clang-tools-extra/docs/ReleaseNotes.rst
===================================================================
--- clang-tools-extra/docs/ReleaseNotes.rst
+++ clang-tools-extra/docs/ReleaseNotes.rst
@@ -111,6 +111,12 @@
Warns whenever the error from CUDA API call is ignored/not handled with a set handler
and provides fixes for it.
+- New :doc:`cuda-unsafe-kernel-call
+ <clang-tidy/checks/cuda/unsafe-kernel-call>` check.
+
+ Warns whenever the possible error after launchign a CUDA kernel is not checked
+ (with a `cudaGetLastError()` function).
+
New check aliases
^^^^^^^^^^^^^^^^^
Index: clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.h
===================================================================
--- /dev/null
+++ clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.h
@@ -0,0 +1,81 @@
+//===--- UnsafeKernelCallCheck.h - clang-tidy -------------------*- 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 LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_CUDA_UNSAFEKERNELCALLCHECK_H
+#define LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_CUDA_UNSAFEKERNELCALLCHECK_H
+
+#include "../ClangTidyCheck.h"
+#include "llvm/ADT/StringSet.h"
+#include <unordered_set>
+
+namespace clang {
+namespace tidy {
+namespace cuda {
+
+/// Checks for whether the possible errors with kernel launches are handled.
+///
+/// CUDA kernels do not always launch correctly. This may happen due to a driver
+/// malfunction, lack of permissions, lack of a GPU, or a multitude of other
+/// reasons. Such errors should be detected by calling the cudaGetLastError()
+/// function following the kernel invocation. The invocation of the error should
+/// be the the first side-effectful AST node after the invocation of the kernel
+/// call (traversing the AST post-order) and a part of the first non-expression
+/// statement after the kernel call. More precisely, it should be the first CFG
+/// statement produced in line after the kernel call using the default options
+/// for CFG building. This is because having the error checks closer to the
+/// kernel invocation makes it easier to debug the code.
+///
+/// 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. The handler will also be accepted
+/// even if it does not actually call cudaGetLastError().
+/// - "AcceptedHandlers" (optional):
+/// a comma-separated list specifying the only accepted handling
+/// functions/macros that can alternatively handle the kernel error besides
+/// the handler specified in HandlerName. 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).
+class UnsafeKernelCallCheck : public ClangTidyCheck {
+ class PPCallback;
+
+public:
+ UnsafeKernelCallCheck(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;
+ void reportIssue(const Stmt &Stmt, ASTContext &Context);
+ bool checkHandlerMacro(const Stmt &Stmt, ASTContext &Context);
+
+ const std::string AcceptedHandlersList;
+ const llvm::StringSet<llvm::MallocAllocator> AcceptedHandlersSet;
+ bool isAcceptedHandler(const StringRef &Name);
+ static llvm::StringSet<llvm::MallocAllocator>
+ splitAcceptedHandlers(const llvm::StringRef &AcceptedHandlers,
+ const llvm::StringRef &HandlerName);
+
+ std::unordered_set<SourceLocation,
+ std::function<unsigned(const SourceLocation &)>>
+ HandlerMacroLocations;
+};
+
+} // namespace cuda
+} // namespace tidy
+} // namespace clang
+
+#endif // LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_CUDA_UNSAFEKERNELCALLCHECK_H
Index: clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.cpp
===================================================================
--- /dev/null
+++ clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.cpp
@@ -0,0 +1,357 @@
+//===--- UnsafeKernelCallCheck.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 "UnsafeKernelCallCheck.h"
+#include "../utils/FixItHintUtils.h"
+#include "../utils/LexerUtils.h"
+#include "clang/Analysis/CFG.h"
+#include "clang/Basic/SourceManagerInternals.h"
+#include "clang/Lex/PPCallbacks.h"
+#include "clang/Lex/Preprocessor.h"
+#include "clang/Tooling/FixIt.h"
+#include <cctype>
+
+using namespace clang::ast_matchers;
+
+namespace clang {
+namespace tidy {
+namespace cuda {
+
+namespace {
+
+constexpr auto HandlerNameOptionName = "HandlerName";
+constexpr auto AcceptedHandlersOptionName = "AcceptedHandlers";
+
+} // namespace
+
+UnsafeKernelCallCheck::UnsafeKernelCallCheck(
+ llvm::StringRef Name, clang::tidy::ClangTidyContext *Context)
+ : ClangTidyCheck(Name, Context),
+ HandlerName(Options.get(HandlerNameOptionName, "")),
+ AcceptedHandlersList(Options.get(AcceptedHandlersOptionName, "")),
+ AcceptedHandlersSet(
+ splitAcceptedHandlers(AcceptedHandlersList, HandlerName)),
+ HandlerMacroLocations(
+ 8, [](const SourceLocation &sLoc) { return sLoc.getHashValue(); }) {
+ if (AcceptedHandlersSet.find("") != AcceptedHandlersSet.end()) {
+ configurationDiag(
+ "Empty handler name found in the list of accepted handlers",
+ DiagnosticIDs::Error);
+ }
+}
+
+llvm::StringSet<llvm::MallocAllocator>
+UnsafeKernelCallCheck::splitAcceptedHandlers(
+ const llvm::StringRef &AcceptedHandlers,
+ const llvm::StringRef &HandlerName) {
+ if (AcceptedHandlers.trim().empty()) {
+ return HandlerName.empty()
+ ? llvm::StringSet<llvm::MallocAllocator>()
+ : llvm::StringSet<llvm::MallocAllocator>{HandlerName};
+ }
+ llvm::SmallVector<llvm::StringRef> AcceptedHandlersVector;
+ AcceptedHandlers.split(AcceptedHandlersVector, ',');
+
+ llvm::StringSet<llvm::MallocAllocator> AcceptedHandlersSet;
+ for (auto AcceptedHandler : AcceptedHandlersVector) {
+ AcceptedHandlersSet.insert(AcceptedHandler.trim());
+ }
+ if (!AcceptedHandlersSet.empty() && !HandlerName.empty()) {
+ AcceptedHandlersSet.insert(HandlerName);
+ }
+
+ return AcceptedHandlersSet;
+}
+
+void UnsafeKernelCallCheck::storeOptions(ClangTidyOptions::OptionMap &Opts) {
+ Options.store(Opts, HandlerNameOptionName, HandlerName);
+ Options.store(Opts, AcceptedHandlersOptionName, AcceptedHandlersList);
+}
+
+bool UnsafeKernelCallCheck::isAcceptedHandler(const StringRef &Name) {
+ return AcceptedHandlersSet.contains(Name);
+}
+
+// Gathers the instances of the handler as a macro being used
+class UnsafeKernelCallCheck::PPCallback : public PPCallbacks {
+public:
+ PPCallback(UnsafeKernelCallCheck &Check) : Check(Check) {}
+
+ void MacroExpands(const Token &MacroNameTok, const MacroDefinition &MD,
+ SourceRange Range, const MacroArgs *Args) override {
+ if (Check.isAcceptedHandler(MacroNameTok.getIdentifierInfo()->getName())) {
+ Check.HandlerMacroLocations.insert(MacroNameTok.getLocation());
+ }
+ }
+
+private:
+ UnsafeKernelCallCheck &Check;
+};
+
+void UnsafeKernelCallCheck::registerPPCallbacks(
+ const SourceManager &SM, Preprocessor *PP, Preprocessor *ModuleExpanderPP) {
+ ModuleExpanderPP->addPPCallbacks(
+ std::make_unique<UnsafeKernelCallCheck::PPCallback>(*this));
+}
+
+void UnsafeKernelCallCheck::registerMatchers(MatchFinder *Finder) {
+ Finder->addMatcher(functionDecl(hasBody(hasDescendant(cudaKernelCallExpr())))
+ .bind("function"),
+ this);
+}
+
+namespace {
+
+// Fetches the first parent available. Should be used
+// for things that are common for the parents, like the location,
+// since the only way a node can have multiple parents is with templates
+template <typename Node, typename Parent = Node>
+inline const Parent *getParent(const Node &Stmt, ASTContext &Context) {
+ auto parents = Context.getParents(Stmt);
+
+ return parents.empty() ? nullptr : parents.begin()->template get<Parent>();
+}
+
+bool isKernelCall(const Stmt *Stmt) {
+ return Stmt->getStmtClass() == Stmt::CUDAKernelCallExprClass;
+}
+
+bool isInCudaRuntimeHeader(SourceLocation Loc, const SourceManager &SM) {
+ constexpr auto CudaHeaderNameSuffix = "cuda_runtime.h";
+ while (Loc.isValid()) {
+ if (SM.getFilename(Loc).endswith(CudaHeaderNameSuffix)) {
+ return true;
+ }
+ Loc = SM.getIncludeLoc(SM.getFileID(Loc));
+ }
+ return false;
+}
+
+bool isCudaGetLastErrorCall(const Stmt *const Stmt, const SourceManager &SM) {
+ constexpr auto GetLastErrorFunctionName = "cudaGetLastError";
+ constexpr auto GetLastErrorFunctionScopedType = "::cudaError_t";
+ constexpr auto GetLastErrorFunctionType = GetLastErrorFunctionScopedType + 2;
+ if (Stmt->getStmtClass() != Stmt::CallExprClass) {
+ return false;
+ }
+ auto CallExprNode = static_cast<const CallExpr *>(Stmt);
+
+ if (!CallExprNode->getCalleeDecl() ||
+ CallExprNode->getCalleeDecl()->getKind() != Decl::Function) {
+ return false;
+ }
+ const auto FunctionDeclNode =
+ static_cast<const FunctionDecl *>(CallExprNode->getCalleeDecl());
+
+ const auto ReturnTypeName = FunctionDeclNode->getReturnType().getAsString();
+ return FunctionDeclNode->getName() == GetLastErrorFunctionName &&
+ (ReturnTypeName == GetLastErrorFunctionType ||
+ StringRef(ReturnTypeName).endswith(GetLastErrorFunctionScopedType)) &&
+ isInCudaRuntimeHeader(FunctionDeclNode->getLocation(), SM);
+}
+
+bool isHandlerCall(
+ const Stmt *const Stmt,
+ std::function<bool(const llvm::StringRef &)> HandlerNamePredicate) {
+ if (Stmt->getStmtClass() != Stmt::CallExprClass) {
+ return false;
+ }
+ auto CallExprNode = static_cast<const CallExpr *>(Stmt);
+
+ if (!CallExprNode->getCalleeDecl() ||
+ CallExprNode->getCalleeDecl()->getKind() != Decl::Function) {
+ return false;
+ }
+ const auto FunctionDeclNode =
+ static_cast<const FunctionDecl *>(CallExprNode->getCalleeDecl());
+
+ return HandlerNamePredicate(FunctionDeclNode->getName()) ||
+ HandlerNamePredicate(FunctionDeclNode->getQualifiedNameAsString());
+}
+
+/// Searches for the closest CFGElement that is an instance of CFGStmt. Does not
+/// increment the index if it already indexes a CFGStmt.
+const Stmt *findStmt(const CFGBlock *const Block, size_t &Idx) {
+ while (Idx < Block->size() && !(*Block)[Idx].getAs<CFGStmt>().has_value()) {
+ Idx++;
+ }
+ if (Idx < Block->size()) {
+ return (*Block)[Idx].castAs<CFGStmt>().getStmt();
+ }
+ return nullptr;
+}
+
+inline bool isBlockReachable(const CFGBlock::AdjacentBlock &Block) {
+ return Block && Block.isReachable();
+}
+
+template <typename Iter>
+inline size_t countReachableBlocks(llvm::iterator_range<Iter> Range) {
+ return std::count_if(Range.begin(), Range.end(), isBlockReachable);
+}
+
+template <typename Iter>
+inline Iter findReachableBlock(llvm::iterator_range<Iter> Range) {
+ return std::find_if(Range.begin(), Range.end(), isBlockReachable);
+}
+
+/// Searches for a next statement from this successor block as if all the empty
+/// blocks were removed and all blocks that could be merged were merged. For
+/// instance, in the following code the call to b() should be found assuming the
+/// `block` argument is set to the first CFG block after the first block:
+/// int foo() {
+/// a();
+/// do {
+/// do {
+/// b()
+/// } while(0);
+/// } while(0);
+/// }
+const Stmt *findNextStmtNonEmptyBlock(const CFGBlock *const Block) {
+ // Enforce that the next block could be mergeable with the next block, i.e.
+ // has no non-trivial predecesors. Trivial predecessors here are chains of
+ // empty predecessors that have up to one predecessor that is itself a trivial
+ // predecessor.
+ int PrunedPredCount = 0;
+ for (auto Pred : Block->preds()) {
+ while (Pred && Pred.isReachable() && Pred->empty() &&
+ countReachableBlocks(Pred->preds()) == 1) {
+ Pred = *findReachableBlock(Pred->preds());
+ }
+ if (Pred && (!Pred->empty() || countReachableBlocks(Pred->preds()) > 1)) {
+ ++PrunedPredCount;
+ }
+ }
+ if (PrunedPredCount > 1) {
+ return nullptr;
+ }
+
+ // Check if there is any statement in this block that we could return
+ size_t Idx = 0;
+ if (const auto Stmt = findStmt(Block, Idx)) {
+ return Stmt;
+ }
+
+ // If the block is empty then try our luck with the next block, provided there
+ // is only one
+ if (countReachableBlocks(Block->succs()) != 1) {
+ return nullptr;
+ }
+ const auto NextBlock = *findReachableBlock(Block->succs());
+ return findNextStmtNonEmptyBlock(NextBlock);
+}
+
+} // namespace
+
+void UnsafeKernelCallCheck::check(const MatchFinder::MatchResult &Result) {
+ const auto FunctionDeclNode =
+ Result.Nodes.getNodeAs<FunctionDecl>("function");
+ const auto Cfg = CFG::buildCFG(FunctionDeclNode, FunctionDeclNode->getBody(),
+ Result.Context, CFG::BuildOptions());
+
+ for (const auto &block : *Cfg) {
+ size_t Idx = 0;
+ while (const auto Stmt = findStmt(block, Idx)) {
+ ++Idx;
+ if (!isKernelCall(Stmt)) {
+ continue;
+ }
+ if (checkHandlerMacro(*Stmt, *Result.Context)) {
+ continue;
+ }
+
+ auto NextStmt = findStmt(block, Idx);
+ // Workaround for the do {...} while(0) not being erased out during
+ // pruning
+ if (!NextStmt) {
+ if (countReachableBlocks(block->succs()) != 1) {
+ reportIssue(*Stmt, *Result.Context);
+ continue;
+ }
+ const auto NextBlock = findReachableBlock(block->succs());
+ NextStmt = findNextStmtNonEmptyBlock(*NextBlock);
+ }
+
+ if (NextStmt && isCudaGetLastErrorCall(NextStmt, *Result.SourceManager)) {
+ continue;
+ }
+ if (NextStmt &&
+ isHandlerCall(NextStmt, [this](const llvm::StringRef &Name) {
+ return isAcceptedHandler(Name);
+ })) {
+ continue;
+ }
+ reportIssue(*Stmt, *Result.Context);
+ }
+ }
+}
+
+// Searches for a handler macro being used right after the kernel call
+bool UnsafeKernelCallCheck::checkHandlerMacro(const Stmt &Stmt,
+ ASTContext &Context) {
+ llvm::Optional<Token> Token = Lexer::findNextToken(
+ Stmt.getEndLoc(), Context.getSourceManager(), Context.getLangOpts());
+ if (!Token.has_value()) {
+ return false;
+ }
+ while (Token->isOneOf(tok::semi, tok::comment)) {
+ Token =
+ Lexer::findNextToken(Token->getLocation(), Context.getSourceManager(),
+ Context.getLangOpts());
+ if (!Token.has_value()) {
+ return false;
+ }
+ }
+ return HandlerMacroLocations.find(Token->getLocation()) !=
+ HandlerMacroLocations.end();
+}
+
+void UnsafeKernelCallCheck::reportIssue(const Stmt &Stmt, ASTContext &Context) {
+ // Get the wrapping expression
+ const clang::Stmt *ExprWithCleanups =
+ getParent<clang::Stmt, clang::ExprWithCleanups>(Stmt, Context);
+
+ // Under certain compilation options kernel calls may not be wrapped
+ // in cleanups
+ if (!ExprWithCleanups) {
+ ExprWithCleanups = &Stmt;
+ }
+
+ const bool IsInMacro = ExprWithCleanups->getBeginLoc().isInvalid() ||
+ ExprWithCleanups->getBeginLoc().isMacroID() ||
+ ExprWithCleanups->getEndLoc().isInvalid() ||
+ ExprWithCleanups->getEndLoc().isMacroID();
+
+ if (!HandlerName.empty()) {
+ const auto DiagnosticBuilder = diag(
+ Stmt.getEndLoc(), (llvm::Twine("Possible unchecked error after a "
+ "kernel launch. Try adding the `") +
+ HandlerName + "()` macro after the kernel call:")
+ .str());
+ if (IsInMacro) {
+ return;
+ }
+ const auto ExprTerminator = utils::lexer::findNextTerminator(
+ ExprWithCleanups->getEndLoc(), Context.getSourceManager(),
+ Context.getLangOpts());
+ const auto ParentStmt = getParent<clang::Stmt>(*ExprWithCleanups, Context);
+ assert(ParentStmt);
+ DiagnosticBuilder << utils::fixit::addSubsequentStatement(
+ SourceRange(ExprWithCleanups->getBeginLoc(), ExprTerminator),
+ *ParentStmt, HandlerName + "()", Context);
+ } else {
+ diag(Stmt.getEndLoc(),
+ "Possible unchecked error after a kernel launch. Try using "
+ "`cudaGetLastError()` right after the kernel call to get the error or "
+ "specify a project-wide kernel call error handler.");
+ }
+}
+
+} // namespace cuda
+} // namespace tidy
+} // namespace clang
Index: clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp
===================================================================
--- clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp
+++ clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp
@@ -10,6 +10,7 @@
#include "../ClangTidyModule.h"
#include "../ClangTidyModuleRegistry.h"
#include "UnsafeApiCallCheck.h"
+#include "UnsafeKernelCallCheck.h"
using namespace clang::ast_matchers;
@@ -21,6 +22,8 @@
public:
void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {
CheckFactories.registerCheck<UnsafeApiCallCheck>("cuda-unsafe-api-call");
+ CheckFactories.registerCheck<UnsafeKernelCallCheck>(
+ "cuda-unsafe-kernel-call");
}
};
Index: clang-tools-extra/clang-tidy/cuda/CMakeLists.txt
===================================================================
--- clang-tools-extra/clang-tidy/cuda/CMakeLists.txt
+++ clang-tools-extra/clang-tidy/cuda/CMakeLists.txt
@@ -1,6 +1,7 @@
add_clang_library(clangTidyCudaModule
CudaTidyModule.cpp
UnsafeApiCallCheck.cpp
+ UnsafeKernelCallCheck.cpp
LINK_LIBS
clangTidy
clangTidyUtils
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
