diff options
| author | Peter Collingbourne <peter@pcc.me.uk> | 2010-12-12 23:02:57 +0000 |
|---|---|---|
| committer | Peter Collingbourne <peter@pcc.me.uk> | 2010-12-12 23:02:57 +0000 |
| commit | 2c2c8dd0acb2a51067299bfcec9ff2145f2031c8 (patch) | |
| tree | 603e35811801272283558d68bad4c2538407198a | |
| parent | aee543a1a3d70de38cd2607fd2f3179551febc93 (diff) | |
Sema: diagnose kernel functions with non-void return type
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@121653 91177308-0d34-0410-b5e6-96231b3b80d8
| -rw-r--r-- | include/clang/Basic/DiagnosticSemaKinds.td | 5 | ||||
| -rw-r--r-- | lib/Sema/SemaDeclAttr.cpp | 15 | ||||
| -rw-r--r-- | test/SemaCUDA/qualifiers.cu | 3 |
3 files changed, 22 insertions, 1 deletions
diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td index 55270b0116..1b23aaaf63 100644 --- a/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp index e5f27fcfbb..a2a7fdcdd5 100644 --- a/lib/Sema/SemaDeclAttr.cpp +++ b/lib/Sema/SemaDeclAttr.cpp @@ -2168,6 +2168,21 @@ static void HandleGlobalAttr(Decl *d, const AttributeList &Attr, Sema &S) { return; } + FunctionDecl *FD = cast<FunctionDecl>(d); + if (!FD->getResultType()->isVoidType()) { + TypeLoc TL = FD->getTypeSourceInfo()->getTypeLoc(); + if (FunctionTypeLoc* FTL = dyn_cast<FunctionTypeLoc>(&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/test/SemaCUDA/qualifiers.cu b/test/SemaCUDA/qualifiers.cu index 8d5b759a6d..1346d654b8 100644 --- a/test/SemaCUDA/qualifiers.cu +++ b/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; +} |