clang/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
  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
  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
  30 static __device__ __host__ double t3(double x) {
  31   long double result;
  32   __asm __volatile("frndint"
  33                    : "=t"(result)
  34                    : "0"(x));
  35   return result;
  36 }
  37
  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
  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)));
  62 static __device__ __host__ void t6(__m256i *p) {
  63   __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
  64                    : "ymm0");
  65 }
  66
  67 static __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