1 // RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -triple x86_64-unknown-linux-gnu -Wuninitialized
2 // RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -DHOST_USED -triple x86_64-unknown-linux-gnu -Wuninitialized
3 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE_NOT_USED -triple nvptx-unknown-cuda -Wuninitialized
4 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -triple nvptx-unknown-cuda -Wuninitialized
5 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -DDEVICE_USED -triple nvptx-unknown-cuda -Wuninitialized
6
7 // REQUIRES: x86-registered-target
8 // REQUIRES: nvptx-registered-target
9
10 #if (defined(HOST) && !defined(HOST_USED)) || defined(DEVICE_NOT_USED)
11 // expected-no-diagnostics
12 #endif
13
14 #include "Inputs/cuda.h"
15
t1(int r)16 static __device__ __host__ void t1(int r) {
17 __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
18 : [ r ] "+r"(r)
19 : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
20 }
21
t2(signed char input)22 static __device__ __host__ unsigned t2(signed char input) {
23 unsigned output;
24 __asm__("xyz"
25 : "=a"(output)
26 : "0"(input));
27 return output;
28 }
29
t3(double x)30 static __device__ __host__ double t3(double x) {
31 register long double result;
32 __asm __volatile("frndint"
33 : "=t"(result)
34 : "0"(x));
35 return result;
36 }
37
t4(unsigned char a,unsigned char b)38 static __device__ __host__ unsigned char t4(unsigned char a, unsigned char b) {
39 unsigned int la = a;
40 unsigned int lb = b;
41 unsigned int bigres;
42 unsigned char res;
43 __asm__("0:\n1:\n"
44 : [ bigres ] "=la"(bigres)
45 : [ la ] "0"(la), [ lb ] "c"(lb)
46 : "edx", "cc");
47 res = bigres;
48 return res;
49 }
50
t5(void)51 static __device__ __host__ void t5(void) {
52 __asm__ __volatile__(
53 "finit"
54 :
55 :
56 : "st", "st(1)", "st(2)", "st(3)",
57 "st(4)", "st(5)", "st(6)", "st(7)",
58 "fpsr", "fpcr");
59 }
60
61 typedef long long __m256i __attribute__((__vector_size__(32)));
t6(__m256i * p)62 static __device__ __host__ void t6(__m256i *p) {
63 __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p)
64 : "ymm0");
65 }
66
t7(__m256i * p)67 static __device__ __host__ void t7(__m256i *p) {
68 __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p)
69 : "r0");
70 }
71
72 #ifdef DEVICE
m()73 __device__ int m() {
74 t1(0);
75 t2(0);
76 t3(0);
77 t4(0, 0);
78 t5();
79 t6(0);
80 #ifdef DEVICE_USED
81 t7(0);
82 #endif // DEVICE_USED
83 return 0;
84 }
85 #endif // DEVICE
86
87 #ifdef HOST
main()88 __host__ int main() {
89 t1(0);
90 t2(0);
91 t3(0);
92 t4(0, 0);
93 t5();
94 t6(0);
95 #ifdef HOST_USED
96 t7(0);
97 #endif // HOST_USED
98 return 0;
99 }
100 #endif // HOST
101
102 #if defined(HOST_USED)
103 // expected-error@69 {{unknown register name 'r0' in asm}}
104 // expected-note@96 {{called by 'main'}}
105 #elif defined(DEVICE)
106 // expected-error@19 {{invalid input constraint 'mx' in asm}}
107 // expected-error@25 {{invalid output constraint '=a' in asm}}
108 // expected-error@33 {{invalid output constraint '=t' in asm}}
109 // expected-error@44 {{invalid output constraint '=la' in asm}}
110 // expected-error@56 {{unknown register name 'st' in asm}}
111 // expected-error@64 {{unknown register name 'ymm0' in asm}}
112 // expected-note@74 {{called by 'm'}}
113 // expected-note@75 {{called by 'm'}}
114 // expected-note@76 {{called by 'm'}}
115 // expected-note@77 {{called by 'm'}}
116 // expected-note@78 {{called by 'm'}}
117 // expected-note@79 {{called by 'm'}}
118 #endif
119