barcisz created this revision.
Herald added subscribers: mattd, carlosgalvezp, yaxunl, mgorny.
Herald added a project: All.
barcisz requested review of this revision.
Herald added a project: clang-tools-extra.
Herald added a subscriber: cfe-commits.
Add cuda-unchecked-kernel-call check
====================================
Motivation
----------
Calls to CUDA kernels can yield errors after their invocation. These errors can
be obtained by calling `cudaGetLastError()`, which also resets CUDA’s error
state. There is a non error-resetting version of this function called
`cudaPeekAtLastError()`, but the lint check does not accept this (see
below). A limited set of errors can block a kernel from launching including
driver malfunctions, **trying to allocate too much shared memory**, or using
too many threads or blocks. Since those errors can cause unexpected behavior
that blocks subsequent computation, they should be caught as close
to the launch point as possible. The lint check enforces this by requiring that
every kernel be immediately followed by an error check.
Behavior
--------
The **cuda-unchecked-kernel-call** checks whether there is a call to
`cudaGetLastError()` directly after each kernel call. To be precise, there can
be no side-effecting or branching code between the kernel call and the call to
`cudaGetLastError()`, such as branching due to the `?:` operator or due
to a call to a function. This is because a more complicated behavior is likely
to be harder for humans to read and would would be significantly slower to
automatically check. We want to encourage well-designed, multi-line macros that
check for errors, so we explicitly allow macros whose content is
`do { /* error check */ } while(false)`, since this is the recommended way
<https://wiki.sei.cmu.edu/confluence/display/c/PRE10-C.+Wrap+multistatement+macros+in+a+do-while+loop>
of making multi-line macros.
The check does also accept the handler it was provided as a valid way to handle
the error, even if the handler does not comply with the rule above (or is a
function which cannot be easily and quickly checked). However, it is still
encouraged to call `cudaGetLastError()` early in the handler’s code
for the code to be readable.
Automatic fixes
---------------
The lint check can be configured to automatically fix the issue by adding an
error handling macro right after the kernel launch. You can specify the error
handler for your project by setting the **HandlerName** option for the
**cuda-unchecked-kernel-call**. Here is an example of how this fix can
transform unhandled code from:
void foo(bool b) {
if (b)
kernel<<<x, y>>>();
}
to
void foo(bool b) {
if(b)
{kernel<<<x, y>>>(); `C10_CUDA_KERNEL_LAUNCH_CHECK`();}
}
The specific handler used for this example is taken from PyTorch and its
definition can be found here
<https://github.com/pytorch/pytorch/blob/master/c10/cuda/CUDAException.h>.
Known Limitations
-----------------
Using cudaPeekAtLastError()
---------------------------
`cudaPeekAtLastError()` can also be used to check for CUDA kernel launch
errors. However, there are several reasons why this is not and will most likely
not be considered as a valid way to check for errors after kernel invocations.
This all has to do with the purpose of the function, which is to
not reset the internal error variable:
- Subsequent kernel calls, even if they don’t produce any errors, will seem as
if they produced an error due to the error not being reset. This behavior is
easy to overlook and may cause he significant difficulty in debugging.
- Our linter cannot easily check whether the error was reset before subsequent
kernel calls. It might even be impossible to do so due to the error leaking
inter-procedurally from functions whose code we can’t access.
Checking for errors that occurred while a kernel was running
------------------------------------------------------------
Our linter does not check whether errors occurred while a kernel was running.
The linter only enforces checks that a kernel launched correctly.
`cudaDeviceSynchronize()` and similar API calls can be used to see that a
kernel’s computation was successful, but these are blocking calls, so we are
not
able to suggest where they should go automatically.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D133956
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/clang-tidy/utils/FixItHintUtils.cpp
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.h
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/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda.h
===================================================================
--- clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda.h
+++ clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda.h
@@ -1,5 +1,7 @@
/* Minimal declarations for CUDA support. Testing purposes only. */
+using size_t = long long unsigned;
+
#define __constant__ __attribute__((constant))
#define __device__ __attribute__((device))
#define __global__ __attribute__((global))
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/utils/FixItHintUtils.cpp
===================================================================
--- clang-tools-extra/clang-tidy/utils/FixItHintUtils.cpp
+++ clang-tools-extra/clang-tidy/utils/FixItHintUtils.cpp
@@ -11,6 +11,8 @@
#include "clang/AST/ASTContext.h"
#include "clang/AST/Type.h"
+#include <iostream>
+
namespace clang {
namespace tidy {
namespace utils {
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
