Author: abataev
Date: Tue Feb 26 12:52:16 2019
New Revision: 354925
URL: http://llvm.org/viewvc/llvm-project?rev=354925&view=rev
Log:
[OPENMP]Delay emission for unsupported va_arg expression.
If the OpenMP device is NVPTX and va_arg is used, delay emission of the
error for va_arg unless it is used in the device code.
Added:
cfe/trunk/test/OpenMP/nvptx_va_arg_delayed_diags.c
Modified:
cfe/trunk/lib/Sema/SemaExpr.cpp
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=354925&r1=354924&r2=354925&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Tue Feb 26 12:52:16 2019
@@ -13909,6 +13909,11 @@ ExprResult Sema::BuildVAArgExpr(SourceLo
}
}
+ // NVPTX does not support va_arg expression.
+ if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
+ Context.getTargetInfo().getTriple().isNVPTX())
+ targetDiag(E->getBeginLoc(), diag::err_va_arg_in_device);
+
// It might be a __builtin_ms_va_list. (But don't ever mark a va_arg()
// as Microsoft ABI on an actual Microsoft platform, where
// __builtin_ms_va_list and __builtin_va_list are the same.)
Added: cfe/trunk/test/OpenMP/nvptx_va_arg_delayed_diags.c
URL:
http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_va_arg_delayed_diags.c?rev=354925&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_va_arg_delayed_diags.c (added)
+++ cfe/trunk/test/OpenMP/nvptx_va_arg_delayed_diags.c Tue Feb 26 12:52:16 2019
@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown
-fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown
-fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device
-fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only
+// RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -x c -triple
nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device
-fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only
+// RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -x c -triple
nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device
-fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+#ifndef DIAGS
+// expected-no-diagnostics
+#endif // DIAGS
+
+#ifdef IMMEDIATE
+#pragma omp declare target
+#endif //IMMEDIATE
+void t1(int r, ...) {
+#ifdef DIAGS
+// expected-error@+4 {{CUDA device code does not support va_arg}}
+#endif // DIAGS
+ __builtin_va_list list;
+ __builtin_va_start(list, r);
+ (void)__builtin_va_arg(list, int);
+ __builtin_va_end(list);
+}
+
+#ifdef IMMEDIATE
+#pragma omp end declare target
+#endif //IMMEDIATE
+
+int main() {
+#ifdef DELAYED
+#pragma omp target
+#endif // DELAYED
+ {
+#ifdef DELAYED
+// expected-note@+2 {{called by 'main'}}
+#endif // DELAYED
+ t1(0);
+ }
+ return 0;
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
