This revision was automatically updated to reflect the committed changes.
Closed by commit rL264106: [CUDA] Don't allow templated variadic functions.
(authored by jlebar).
Changed prior to commit:
http://reviews.llvm.org/D18373?vs=51332&id=51354#toc
Repository:
rL LLVM
http://reviews.llvm.org/D18373
Files:
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/test/SemaCUDA/vararg.cu
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -8341,6 +8341,28 @@
isExplicitSpecialization || isFunctionTemplateSpecialization);
}
+ 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, unless someone passed
+ // -fcuda-allow-variadic-functions.
+ if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() &&
+ (NewFD->hasAttr<CUDADeviceAttr>() ||
+ NewFD->hasAttr<CUDAGlobalAttr>()) &&
+ !(II && II->isStr("printf") && NewFD->isExternC() &&
+ !D.isFunctionDefinition())) {
+ Diag(NewFD->getLocation(), diag::err_variadic_device_fn);
+ }
+ }
+
if (getLangOpts().CPlusPlus) {
if (FunctionTemplate) {
if (NewFD->isInvalidDecl())
@@ -8390,28 +8412,6 @@
MarkUnusedFileScopedDecl(NewFD);
- 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, unless someone passed
- // -fcuda-allow-variadic-functions.
- if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() &&
- (NewFD->hasAttr<CUDADeviceAttr>() ||
- NewFD->hasAttr<CUDAGlobalAttr>()) &&
- !(II && II->isStr("printf") && NewFD->isExternC() &&
- !D.isFunctionDefinition())) {
- Diag(NewFD->getLocation(), diag::err_variadic_device_fn);
- }
- }
-
// 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.
Index: cfe/trunk/test/SemaCUDA/vararg.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/vararg.cu
+++ cfe/trunk/test/SemaCUDA/vararg.cu
@@ -35,6 +35,12 @@
// expected-error@-2 {{CUDA device code does not support variadic functions}}
#endif
+template <typename T>
+__device__ void vararg(T t, ...) {}
+#ifdef EXPECT_VARARG_ERR
+// expected-error@-2 {{CUDA device code does not support variadic functions}}
+#endif
+
extern "C" __device__ int printf(const char* fmt, ...); // OK, special case.
// Definition of printf not allowed.
Index: cfe/trunk/lib/Sema/SemaDecl.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp
+++ cfe/trunk/lib/Sema/SemaDecl.cpp
@@ -8341,6 +8341,28 @@
isExplicitSpecialization || isFunctionTemplateSpecialization);
}
+ 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, unless someone passed
+ // -fcuda-allow-variadic-functions.
+ if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() &&
+ (NewFD->hasAttr<CUDADeviceAttr>() ||
+ NewFD->hasAttr<CUDAGlobalAttr>()) &&
+ !(II && II->isStr("printf") && NewFD->isExternC() &&
+ !D.isFunctionDefinition())) {
+ Diag(NewFD->getLocation(), diag::err_variadic_device_fn);
+ }
+ }
+
if (getLangOpts().CPlusPlus) {
if (FunctionTemplate) {
if (NewFD->isInvalidDecl())
@@ -8390,28 +8412,6 @@
MarkUnusedFileScopedDecl(NewFD);
- 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, unless someone passed
- // -fcuda-allow-variadic-functions.
- if (!getLangOpts().CUDAAllowVariadicFunctions && NewFD->isVariadic() &&
- (NewFD->hasAttr<CUDADeviceAttr>() ||
- NewFD->hasAttr<CUDAGlobalAttr>()) &&
- !(II && II->isStr("printf") && NewFD->isExternC() &&
- !D.isFunctionDefinition())) {
- Diag(NewFD->getLocation(), diag::err_variadic_device_fn);
- }
- }
-
// 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.
Index: cfe/trunk/test/SemaCUDA/vararg.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/vararg.cu
+++ cfe/trunk/test/SemaCUDA/vararg.cu
@@ -35,6 +35,12 @@
// expected-error@-2 {{CUDA device code does not support variadic functions}}
#endif
+template <typename T>
+__device__ void vararg(T t, ...) {}
+#ifdef EXPECT_VARARG_ERR
+// expected-error@-2 {{CUDA device code does not support variadic functions}}
+#endif
+
extern "C" __device__ int printf(const char* fmt, ...); // OK, special case.
// Definition of printf not allowed.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
