"kernel function type %0 must have void return type">;
def err_config_scalar_return : Error<
"CUDA special function 'cudaConfigureCall' must have scalar return type">;
+def err_kern_call_not_global_function : Error<
+ "kernel call to non-global function %0">;
def err_cannot_pass_objc_interface_to_vararg : Error<
return ExprError(Diag(LParenLoc, diag::err_typecheck_call_not_function)
<< Fn->getType() << Fn->getSourceRange());
+ if (getLangOptions().CUDA) {
+ if (Config) {
+ // CUDA: Kernel calls must be to global functions
+ if (FDecl && !FDecl->hasAttr<CUDAGlobalAttr>())
+ return ExprError(Diag(LParenLoc,diag::err_kern_call_not_global_function)
+ << FDecl->getName() << Fn->getSourceRange());
+
+ // CUDA: Kernel function must have 'void' return type
+ if (!FuncT->getResultType()->isVoidType())
+ return ExprError(Diag(LParenLoc, diag::err_kern_type_not_void_return)
+ << Fn->getType() << Fn->getSourceRange());
+ }
+ }
+
// Check for a valid return type
if (CheckCallReturnType(FuncT->getResultType(),
Fn->getSourceRange().getBegin(), TheCall,
g1<<<arg, arg>>>(1);
}
+void h1(int x) {}
+int h2(int x) { return 1; }
+
int main(void) {
g1<<<1, 1>>>(42);
t1(1);
+
+ h1<<<1, 1>>>(42); // expected-error {{kernel call to non-global function h1}}
+
+ int (*fp)(int) = h2;
+ fp<<<1, 1>>>(42); // expected-error {{must have void return type}}
}