From e8cfaf4258cad84d8a9c936a93b4dd7d73ab7f5e Mon Sep 17 00:00:00 2001 From: Peter Collingbourne Date: Sun, 12 Dec 2010 23:02:57 +0000 Subject: [PATCH] Sema: diagnose kernel functions with non-void return type llvm-svn: 121653 --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 5 ++++- clang/lib/Sema/SemaDeclAttr.cpp | 15 +++++++++++++++ clang/test/SemaCUDA/qualifiers.cu | 3 +++ 3 files changed, 22 insertions(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 55270b0116f5..1b23aaaf6389 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -2859,9 +2859,12 @@ def err_atomic_builtin_pointer_size : Error< "first argument to atomic builtin must be a pointer to 1,2,4,8 or 16 byte " "type (%0 invalid)">; - def err_deleted_function_use : Error<"attempt to use a deleted function">; +def err_kern_type_not_void_return : Error< + "kernel function type %0 must have void return type">; + + def err_cannot_pass_objc_interface_to_vararg : Error< "cannot pass object with interface type %0 by-value through variadic " "%select{function|block|method}1">; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index e5f27fcfbb88..a2a7fdcdd539 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -2168,6 +2168,21 @@ static void HandleGlobalAttr(Decl *d, const AttributeList &Attr, Sema &S) { return; } + FunctionDecl *FD = cast(d); + if (!FD->getResultType()->isVoidType()) { + TypeLoc TL = FD->getTypeSourceInfo()->getTypeLoc(); + if (FunctionTypeLoc* FTL = dyn_cast(&TL)) { + S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return) + << FD->getType() + << FixItHint::CreateReplacement(FTL->getResultLoc().getSourceRange(), + "void"); + } else { + S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return) + << FD->getType(); + } + return; + } + d->addAttr(::new (S.Context) CUDAGlobalAttr(Attr.getLoc(), S.Context)); } else { S.Diag(Attr.getLoc(), diag::warn_attribute_ignored) << "global"; diff --git a/clang/test/SemaCUDA/qualifiers.cu b/clang/test/SemaCUDA/qualifiers.cu index 8d5b759a6d7b..1346d654b8c4 100644 --- a/clang/test/SemaCUDA/qualifiers.cu +++ b/clang/test/SemaCUDA/qualifiers.cu @@ -3,3 +3,6 @@ #include "cuda.h" __global__ void g1(int x) {} +__global__ int g2(int x) { // expected-error {{must have void return type}} + return 1; +}