119 lines · plain
1// RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -triple x86_64-unknown-linux-gnu -Wuninitialized2// RUN: %clang_cc1 -fsyntax-only -verify %s -DHOST -DHOST_USED -triple x86_64-unknown-linux-gnu -Wuninitialized3// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE_NOT_USED -triple nvptx-unknown-cuda -Wuninitialized4// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -triple nvptx-unknown-cuda -Wuninitialized5// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DDEVICE -DDEVICE_USED -triple nvptx-unknown-cuda -Wuninitialized6 7// REQUIRES: x86-registered-target8// REQUIRES: nvptx-registered-target9 10#if (defined(HOST) && !defined(HOST_USED)) || defined(DEVICE_NOT_USED)11// expected-no-diagnostics12#endif13 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 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 DEVICE73__device__ int m() {74 t1(0);75 t2(0);76 t3(0);77 t4(0, 0);78 t5();79 t6(0);80#ifdef DEVICE_USED81 t7(0);82#endif // DEVICE_USED83 return 0;84}85#endif // DEVICE86 87#ifdef HOST88__host__ int main() {89 t1(0);90 t2(0);91 t3(0);92 t4(0, 0);93 t5();94 t6(0);95#ifdef HOST_USED96 t7(0);97#endif // HOST_USED98 return 0;99}100#endif // HOST101 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#endif119