diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4223,7 +4223,9 @@ return; } const auto *FD = cast(D); - if (!FD->getReturnType()->isVoidType()) { + if (!FD->getReturnType()->isVoidType() && + !FD->getReturnType()->getAs() && + !FD->getReturnType()->isInstantiationDependentType()) { SourceRange RTRange = FD->getReturnTypeSourceRange(); S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return) << FD->getType() diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -5891,7 +5891,9 @@ << FDecl << Fn->getSourceRange()); // CUDA: Kernel function must have 'void' return type - if (!FuncT->getReturnType()->isVoidType()) + if (!FuncT->getReturnType()->isVoidType() && + !FuncT->getReturnType()->getAs() && + !FuncT->getReturnType()->isInstantiationDependentType()) return ExprError(Diag(LParenLoc, diag::err_kern_type_not_void_return) << Fn->getType() << Fn->getSourceRange()); } else { diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp --- a/clang/lib/Sema/SemaStmt.cpp +++ b/clang/lib/Sema/SemaStmt.cpp @@ -3500,6 +3500,14 @@ return true; } + // CUDA: Kernel function must have 'void' return type. + if (getLangOpts().CUDA) + if (FD->hasAttr() && !Deduced->isVoidType()) { + Diag(FD->getLocation(), diag::err_kern_type_not_void_return) + << FD->getType() << FD->getSourceRange(); + return true; + } + // If a function with a declared return type that contains a placeholder type // has multiple return statements, the return type is deduced for each return // statement. [...] if the type deduced is not the same in each deduction, diff --git a/clang/lib/Sema/SemaTemplateDeduction.cpp b/clang/lib/Sema/SemaTemplateDeduction.cpp --- a/clang/lib/Sema/SemaTemplateDeduction.cpp +++ b/clang/lib/Sema/SemaTemplateDeduction.cpp @@ -3093,6 +3093,13 @@ Function->getTypeSpecStartLoc(), Function->getDeclName()); if (ResultType.isNull() || Trap.hasErrorOccurred()) return TDK_SubstitutionFailure; + // CUDA: Kernel function must have 'void' return type. + if (getLangOpts().CUDA) + if (Function->hasAttr() && !ResultType->isVoidType()) { + Diag(Function->getLocation(), diag::err_kern_type_not_void_return) + << Function->getType() << Function->getSourceRange(); + return TDK_SubstitutionFailure; + } } // Instantiate the types of each of the function parameters given the diff --git a/clang/test/SemaCUDA/autoret-global.cu b/clang/test/SemaCUDA/autoret-global.cu new file mode 100644 --- /dev/null +++ b/clang/test/SemaCUDA/autoret-global.cu @@ -0,0 +1,44 @@ +// RUN: %clang_cc1 -std=c++14 -fsyntax-only -verify %s + +#include "Inputs/cuda.h" + +template +__global__ T foo() { + // expected-note@-1 {{kernel function type 'T ()' must have void return type}} +} + +void f0() { + foo<<<0, 0>>>(); + foo<<<0, 0>>>(); + // expected-error@-1 {{no matching function for call to 'foo'}} +} + +__global__ auto f1() { +} + +__global__ auto f2(int x) { + return x + 1; + // expected-error@-2 {{kernel function type 'auto (int)' must have void return type}} +} + +template struct enable_if { typedef T type; }; +template struct enable_if {}; + +template +__global__ +auto bar() -> typename enable_if::type { + // expected-note@-1 {{requirement '3 == 1' was not satisfied [with N = 3]}} +} + +template +__global__ +auto bar() -> typename enable_if::type { + // expected-note@-1 {{requirement '3 == 2' was not satisfied [with N = 3]}} +} + +void f3() { + bar<1><<<0, 0>>>(); + bar<2><<<0, 0>>>(); + bar<3><<<0, 0>>>(); + // expected-error@-1 {{no matching function for call to 'bar'}} +}