Author: jlebar Date: Sat Jan 23 15:28:17 2016 New Revision: 258643 URL: http://llvm.org/viewvc/llvm-project?rev=258643&view=rev Log: [CUDA] Disallow variadic functions other than printf in device code.
Reviewers: tra Subscribers: cfe-commits, echristo, jhen Differential Revision: http://reviews.llvm.org/D16484 Added: cfe/trunk/test/SemaCUDA/vararg.cu - copied, changed from r258642, cfe/trunk/test/SemaCUDA/va-arg.cu Removed: cfe/trunk/test/SemaCUDA/va-arg.cu Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/lib/Sema/SemaDecl.cpp Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=258643&r1=258642&r2=258643&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Sat Jan 23 15:28:17 2016 @@ -6423,6 +6423,8 @@ def warn_kern_is_method : Extension< def warn_kern_is_inline : Warning< "ignored 'inline' attribute on kernel function %0">, InGroup<CudaCompat>; +def err_variadic_device_fn : Error< + "CUDA device code does not support variadic functions">; def err_va_arg_in_device : Error< "CUDA device code does not support va_arg">; def err_alias_not_supported_on_nvptx : Error<"CUDA does not support aliases">; Modified: cfe/trunk/lib/Sema/SemaDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=258643&r1=258642&r2=258643&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) +++ cfe/trunk/lib/Sema/SemaDecl.cpp Sat Jan 23 15:28:17 2016 @@ -8276,18 +8276,26 @@ Sema::ActOnFunctionDeclarator(Scope *S, MarkUnusedFileScopedDecl(NewFD); - if (getLangOpts().CUDA) - if (IdentifierInfo *II = NewFD->getIdentifier()) - if (!NewFD->isInvalidDecl() && - NewFD->getDeclContext()->getRedeclContext()->isTranslationUnit()) { - if (II->isStr("cudaConfigureCall")) { - if (!R->getAs<FunctionType>()->getReturnType()->isScalarType()) - Diag(NewFD->getLocation(), diag::err_config_scalar_return); + if (getLangOpts().CUDA) { + IdentifierInfo *II = NewFD->getIdentifier(); + if (II && II->isStr("cudaConfigureCall") && !NewFD->isInvalidDecl() && + NewFD->getDeclContext()->getRedeclContext()->isTranslationUnit()) { + if (!R->getAs<FunctionType>()->getReturnType()->isScalarType()) + Diag(NewFD->getLocation(), diag::err_config_scalar_return); + + Context.setcudaConfigureCallDecl(NewFD); + } + + // Variadic functions, other than a *declaration* of printf, are not allowed + // in device-side CUDA code. + if (NewFD->isVariadic() && (NewFD->hasAttr<CUDADeviceAttr>() || + NewFD->hasAttr<CUDAGlobalAttr>()) && + !(II && II->isStr("printf") && NewFD->isExternC() && + !D.isFunctionDefinition())) { + Diag(NewFD->getLocation(), diag::err_variadic_device_fn); + } + } - Context.setcudaConfigureCallDecl(NewFD); - } - } - // Here we have an function template explicit specialization at class scope. // The actually specialization will be postponed to template instatiation // time via the ClassScopeFunctionSpecializationDecl node. Removed: cfe/trunk/test/SemaCUDA/va-arg.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/va-arg.cu?rev=258642&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/va-arg.cu (original) +++ cfe/trunk/test/SemaCUDA/va-arg.cu (removed) @@ -1,28 +0,0 @@ -// REQUIRES: x86-registered-target -// REQUIRES: nvptx-registered-target -// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \ -// RUN: -verify -DEXPECT_ERR %s -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only %s - -#include <stdarg.h> -#include "Inputs/cuda.h" - -__device__ void foo() { - va_list list; - va_arg(list, int); -#ifdef EXPECT_ERR - // expected-error@-2 {{CUDA device code does not support va_arg}} -#endif -} - -void bar() { - va_list list; - va_arg(list, int); // OK: host-only -} - -__device__ void baz() { -#if !defined(__CUDA_ARCH__) - va_list list; - va_arg(list, int); // OK: only seen when compiling for host -#endif -} Copied: cfe/trunk/test/SemaCUDA/vararg.cu (from r258642, cfe/trunk/test/SemaCUDA/va-arg.cu) URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/vararg.cu?p2=cfe/trunk/test/SemaCUDA/vararg.cu&p1=cfe/trunk/test/SemaCUDA/va-arg.cu&r1=258642&r2=258643&rev=258643&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/va-arg.cu (original) +++ cfe/trunk/test/SemaCUDA/vararg.cu Sat Jan 23 15:28:17 2016 @@ -2,7 +2,7 @@ // REQUIRES: nvptx-registered-target // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \ // RUN: -verify -DEXPECT_ERR %s -// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s #include <stdarg.h> #include "Inputs/cuda.h" @@ -26,3 +26,17 @@ __device__ void baz() { va_arg(list, int); // OK: only seen when compiling for host #endif } + +__device__ void vararg(const char* x, ...) {} +// expected-error@-1 {{CUDA device code does not support variadic functions}} + +extern "C" __device__ int printf(const char* fmt, ...); // OK, special case. + +// Definition of printf not allowed. +extern "C" __device__ int printf(const char* fmt, ...) { return 0; } +// expected-error@-1 {{CUDA device code does not support variadic functions}} + +namespace ns { +__device__ int printf(const char* fmt, ...); +// expected-error@-1 {{CUDA device code does not support variadic functions}} +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
