1 // REQUIRES: x86-registered-target
2 // REQUIRES: nvptx-registered-target
3 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \
4 // RUN: -verify -DEXPECT_VA_ARG_ERR -DEXPECT_VARARG_ERR %s
5 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -fsyntax-only \
6 // RUN: -fcuda-allow-variadic-functions -verify -DEXPECT_VA_ARG_ERR %s
7 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify \
8 // RUN: -DEXPECT_VARARG_ERR %s
9
10 #include <stdarg.h>
11 #include "Inputs/cuda.h"
12
13 __device__ void foo() {
14 va_list list;
15 va_arg(list, int);
16 #ifdef EXPECT_VA_ARG_ERR
17 // expected-error@-2 {{CUDA device code does not support va_arg}}
18 #endif
19 }
20
21 void bar() {
22 va_list list;
23 va_arg(list, int); // OK: host-only
24 }
25
26 __device__ void baz() {
27 #if !defined(__CUDA_ARCH__)
28 va_list list;
29 va_arg(list, int); // OK: only seen when compiling for host
30 #endif
31 }
32
33 __device__ void vararg(const char* x, ...) {}
34 #ifdef EXPECT_VARARG_ERR
35 // expected-error@-2 {{CUDA device code does not support variadic functions}}
36 #endif
37
38 template <typename T>
39 __device__ void vararg(T t, ...) {}
40 #ifdef EXPECT_VARARG_ERR
41 // expected-error@-2 {{CUDA device code does not support variadic functions}}
42 #endif
43
44 extern "C" __device__ int printf(const char* fmt, ...); // OK, special case.
45
46 // Definition of printf not allowed.
47 extern "C" __device__ int printf(const char* fmt, ...) { return 0; }
test_shm_mq_setup(int64 queue_size,int32 nworkers,dsm_segment ** segp,shm_mq_handle ** output,shm_mq_handle ** input)48 #ifdef EXPECT_VARARG_ERR
49 // expected-error@-2 {{CUDA device code does not support variadic functions}}
50 #endif
51
52 namespace ns {
53 __device__ int printf(const char* fmt, ...);
54 #ifdef EXPECT_VARARG_ERR
55 // expected-error@-2 {{CUDA device code does not support variadic functions}}
56 #endif
57 }
58