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