maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[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.

llvm-svn: 354925
This commit is contained in:
Alexey Bataev 2019-02-26 20:52:16 +00:00
parent f9ef9f868c
commit ddc181d256
2 changed files with 45 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
clang/lib/Sema/SemaExpr.cpp
View File

@ -13909,6 +13909,11 @@ ExprResult Sema::BuildVAArgExpr(SourceLocation BuiltinLoc,
}
}
// 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.)

40
clang/test/OpenMP/nvptx_va_arg_delayed_diags.c Normal file
View File

@ -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;
}