1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
| // 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
void foo(int r, ...) {
#ifdef IMMEDIATE
// expected-error@+4 {{CUDA device code does not support va_arg}}
#endif // IMMEDIATE
__builtin_va_list list;
__builtin_va_start(list, r);
(void)__builtin_va_arg(list, int);
__builtin_va_end(list);
}
#ifdef IMMEDIATE
#pragma omp declare target to(foo)
#endif //IMMEDIATE
#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;
}
|