diff --git a/.clang-tidy b/.clang-tidy --- a/.clang-tidy +++ b/.clang-tidy @@ -1,4 +1,4 @@ -Checks: '-*,clang-diagnostic-*,llvm-*,misc-*,-misc-const-correctness,-misc-unused-parameters,-misc-non-private-member-variables-in-classes,-misc-no-recursion,readability-identifier-naming' +Checks: '-*,cuda-*,clang-diagnostic-*,llvm-*,misc-*,-misc-const-correctness,-misc-unused-parameters,-misc-non-private-member-variables-in-classes,-misc-no-recursion,readability-identifier-naming' CheckOptions: - key: readability-identifier-naming.ClassCase value: CamelCase diff --git a/clang-tools-extra/clang-tidy/cuda/CMakeLists.txt b/clang-tools-extra/clang-tidy/cuda/CMakeLists.txt --- a/clang-tools-extra/clang-tidy/cuda/CMakeLists.txt +++ b/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 diff --git a/clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp b/clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp --- a/clang-tools-extra/clang-tidy/cuda/CudaTidyModule.cpp +++ b/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("cuda-unsafe-api-call"); + CheckFactories.registerCheck( + "cuda-unsafe-kernel-call"); } }; diff --git a/clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.h b/clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.h new file mode 100644 --- /dev/null +++ b/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 + +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 AcceptedHandlersSet; + bool isAcceptedHandler(const StringRef &Name); + static llvm::StringSet + splitAcceptedHandlers(const llvm::StringRef &AcceptedHandlers, + const llvm::StringRef &HandlerName); + + std::unordered_set> + HandlerMacroLocations; +}; + +} // namespace cuda +} // namespace tidy +} // namespace clang + +#endif // LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_CUDA_UNSAFEKERNELCALLCHECK_H diff --git a/clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.cpp b/clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.cpp new file mode 100644 --- /dev/null +++ b/clang-tools-extra/clang-tidy/cuda/UnsafeKernelCallCheck.cpp @@ -0,0 +1,358 @@ +//===--- 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 + +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 +UnsafeKernelCallCheck::splitAcceptedHandlers( + const llvm::StringRef &AcceptedHandlers, + const llvm::StringRef &HandlerName) { + if (AcceptedHandlers.trim().empty()) { + return HandlerName.empty() + ? llvm::StringSet() + : llvm::StringSet{HandlerName}; + } + llvm::SmallVector AcceptedHandlersVector; + AcceptedHandlers.split(AcceptedHandlersVector, ','); + + llvm::StringSet 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(*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 +inline const Parent *getParent(const Node &Stmt, ASTContext &Context) { + auto parents = Context.getParents(Stmt); + + return parents.empty() ? nullptr : parents.begin()->template get(); +} + +bool isKernelCall(const Stmt *Stmt) { + return Stmt->getStmtClass() == Stmt::CUDAKernelCallExprClass; +} + +bool isInCudaRuntimeHeader(SourceLocation Loc, const SourceManager &SM) { + constexpr auto CudaHeaderNameSuffix = "cuda_runtime.h"; + constexpr auto CudaWrapperHeaderNameSuffix = "cuda_runtime_wrapper.h"; + while (Loc.isValid()) { + if (SM.getFilename(Loc).endswith(CudaHeaderNameSuffix) || SM.getFilename(Loc).endswith(CudaWrapperHeaderNameSuffix)) { + 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(Stmt); + + if (!CallExprNode->getCalleeDecl() || + CallExprNode->getCalleeDecl()->getKind() != Decl::Function) { + return false; + } + const auto FunctionDeclNode = + static_cast(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 HandlerNamePredicate) { + if (Stmt->getStmtClass() != Stmt::CallExprClass) { + return false; + } + auto CallExprNode = static_cast(Stmt); + + if (!CallExprNode->getCalleeDecl() || + CallExprNode->getCalleeDecl()->getKind() != Decl::Function) { + return false; + } + const auto FunctionDeclNode = + static_cast(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().has_value()) { + Idx++; + } + if (Idx < Block->size()) { + return (*Block)[Idx].castAs().getStmt(); + } + return nullptr; +} + +inline bool isBlockReachable(const CFGBlock::AdjacentBlock &Block) { + return Block && Block.isReachable(); +} + +template +inline size_t countReachableBlocks(llvm::iterator_range Range) { + return std::count_if(Range.begin(), Range.end(), isBlockReachable); +} + +template +inline Iter findReachableBlock(llvm::iterator_range 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("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 = 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(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.getBeginLoc(), (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(*ExprWithCleanups, Context); + assert(ParentStmt); + DiagnosticBuilder << utils::fixit::addSubsequentStatement( + SourceRange(ExprWithCleanups->getBeginLoc(), ExprTerminator), + *ParentStmt, HandlerName + "()", Context); + } else { + diag(Stmt.getBeginLoc(), + "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 diff --git a/clang-tools-extra/docs/ReleaseNotes.rst b/clang-tools-extra/docs/ReleaseNotes.rst --- a/clang-tools-extra/docs/ReleaseNotes.rst +++ b/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 + ` check. + + Warns whenever the possible error after launchign a CUDA kernel is not checked + (with a `cudaGetLastError()` function). + New check aliases ^^^^^^^^^^^^^^^^^ diff --git a/clang-tools-extra/docs/clang-tidy/checks/cuda/unsafe-kernel-call.rst b/clang-tools-extra/docs/clang-tidy/checks/cuda/unsafe-kernel-call.rst new file mode 100644 --- /dev/null +++ b/clang-tools-extra/docs/clang-tidy/checks/cuda/unsafe-kernel-call.rst @@ -0,0 +1,69 @@ +.. title:: clang-tidy - cuda-unsafe-kernel-call + +cuda-unsafe-kernel-call +======================= + +Finds CUDA kernel calls which do not have any post-invocation error handling +implemented for them. It expects to capture the error after the kernel +invocation using a call to ``cudaGetLastError()``. + +Specification +------------- + +The check finds the declaration of ``cudaGetLastError()`` by checking that: + + - tt has the expected name + + - its return type is ``cudaError_t`` + + - it is included in a file that ends with either ``cuda_runtime.h`` or + ``cuda_runtime_wrapper.h`` (those headers are automatically included from the + during CUDA code compilation) + +The check then generates a Control Flow Graph for the program. To check that the +kernel call is error-handled in a valid way it expects that the first expression +tree or function call in the control flow graph, right after the kernel call, is +the call to ``cudaGetLastError()``. This call must also happen in a direct line +from the kernel call, i.e. no node on the path from the kernel call to call to +``cudaGetLastError()`` can have more than 1 incoming or outgoing control flow +branches. + +Example: + +.. code-block:: c++ + + __global__ + void kernel(); + + void foo() { + kernel<<<64, 128>>>(); + } + +results in the following warnings:: + + 1 warning generated when compiling for host. + test.cu:5:3: warning: 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. [cuda-unsafe-kernel-call] + kernel<<<64, 128>>>(); + ^ + +Options +------- + +.. option:: HandlerName + + The name of the function or macro that should be used in fix it hints to + check the error after a kernel invocation. It will be placed as the next + statement after the kernel call. Even if it is a function call or a macro + that does not call ``cudaGetLastError()``, it will be accepted as a valid + way to handle a kernel call. If the specified handler is a function name + then it can be scoped; however, for performance reasons, if the function + name is scoped then it has to be its fully scoped name. + +.. option:: AcceptedHandlers + + The list of handler functions or macros that are allowed for the specific + project. Just like the handler specified in HandlerName, be it a macro or + a function, they will also be allowed as a valid way to handle the kernel + call even if they would not be accepted otherwise. If the specified handler + is a function name then it can be scoped; however, for performance reasons, + if the function name is scoped then it has to be its fully scoped name. diff --git a/clang-tools-extra/docs/clang-tidy/checks/list.rst b/clang-tools-extra/docs/clang-tidy/checks/list.rst --- a/clang-tools-extra/docs/clang-tidy/checks/list.rst +++ b/clang-tools-extra/docs/clang-tidy/checks/list.rst @@ -199,8 +199,9 @@ `cppcoreguidelines-pro-type-vararg `_, `cppcoreguidelines-slicing `_, `cppcoreguidelines-special-member-functions `_, - `cuda-unsafe-api-call `_, "Yes" `cppcoreguidelines-virtual-class-destructor `_, "Yes" + `cuda-unsafe-api-call `_, "Yes" + `cuda-unsafe-kernel-call `_, "Yes" `darwin-avoid-spinlock `_, `darwin-dispatch-once-nonstatic `_, "Yes" `fuchsia-default-arguments-calls `_, diff --git a/clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h b/clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h --- a/clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h +++ b/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(); diff --git a/clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-function-handler.cu b/clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-function-handler.cu new file mode 100644 --- /dev/null +++ b/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 + +__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 +} diff --git a/clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-macro-handler.cu b/clang-tools-extra/test/clang-tidy/checkers/cuda/unsafe-kernel-call-macro-handler.cu new file mode 100644 --- /dev/null +++ b/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 + +#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); +} diff --git a/test.cu b/test.cu new file mode 100644 --- /dev/null +++ b/test.cu @@ -0,0 +1,8 @@ +#include "clang-tools-extra/test/clang-tidy/checkers/Inputs/Headers/cuda/cuda_runtime.h" + +__global__ +void kernel(); + +void foo() { + kernel<<<64, 128>>>(); +}