Clang Project

clang_source_code/test/SemaCUDA/asm_delayed_diags.cu
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
16static __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
22static __device__ __host__ unsigned t2(signed char input) {
23  unsigned output;
24  __asm__("xyz"
25          : "=a"(output)
26          : "0"(input));
27  return output;
28}
29
30static __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
38static __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
51static __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
61typedef long long __m256i __attribute__((__vector_size__(32)));
62static __device__ __host__ void t6(__m256i *p) {
63  __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
64                   : "ymm0");
65}
66
67static __device__ __host__ void t7(__m256i *p) {
68  __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
69                   : "r0");
70}
71
72#ifdef DEVICE
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
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
CodeThinker, Enjoy your reading code