| 1 | // REQUIRES: amdgpu-registered-target |
|---|
| 2 | // RUN: %clang_cc1 -triple amdgcn -fsyntax-only -verify %s |
|---|
| 3 | |
|---|
| 4 | #pragma OPENCL EXTENSION cl_khr_fp64 : enable |
|---|
| 5 | |
|---|
| 6 | kernel void test () { |
|---|
| 7 | |
|---|
| 8 | int sgpr = 0, vgpr = 0, imm = 0; |
|---|
| 9 | |
|---|
| 10 | // sgpr constraints |
|---|
| 11 | __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "s" (imm) : ); |
|---|
| 12 | |
|---|
| 13 | __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec}" (imm) : ); |
|---|
| 14 | __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exe" (imm) : ); // expected-error {{invalid input constraint '{exe' in asm}} |
|---|
| 15 | __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec" (imm) : ); // expected-error {{invalid input constraint '{exec' in asm}} |
|---|
| 16 | __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec}a" (imm) : ); // expected-error {{invalid input constraint '{exec}a' in asm}} |
|---|
| 17 | |
|---|
| 18 | // vgpr constraints |
|---|
| 19 | __asm__ ("v_mov_b32 %0, %1" : "=v" (vgpr) : "v" (imm) : ); |
|---|
| 20 | } |
|---|
| 21 | |
|---|
| 22 | __kernel void |
|---|
| 23 | test_float(const __global float *a, const __global float *b, __global float *c, unsigned i) |
|---|
| 24 | { |
|---|
| 25 | float ai = a[i]; |
|---|
| 26 | float bi = b[i]; |
|---|
| 27 | float ci; |
|---|
| 28 | |
|---|
| 29 | __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); |
|---|
| 30 | __asm("v_add_f32_e32 v1, v2, v3" : ""(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '' in asm}} |
|---|
| 31 | __asm("v_add_f32_e32 v1, v2, v3" : "="(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '=' in asm}} |
|---|
| 32 | __asm("v_add_f32_e32 v1, v2, v3" : "={a}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={a}' in asm}} |
|---|
| 33 | __asm("v_add_f32_e32 v1, v2, v3" : "={"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={' in asm}} |
|---|
| 34 | __asm("v_add_f32_e32 v1, v2, v3" : "={}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={}' in asm}} |
|---|
| 35 | __asm("v_add_f32_e32 v1, v2, v3" : "={v"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v' in asm}} |
|---|
| 36 | __asm("v_add_f32_e32 v1, v2, v3" : "={v1a}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1a}' in asm}} |
|---|
| 37 | __asm("v_add_f32_e32 v1, v2, v3" : "={va}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={va}' in asm}} |
|---|
| 38 | __asm("v_add_f32_e32 v1, v2, v3" : "={v1}a"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1}a' in asm}} |
|---|
| 39 | __asm("v_add_f32_e32 v1, v2, v3" : "={v1"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1' in asm}} |
|---|
| 40 | __asm("v_add_f32_e32 v1, v2, v3" : "=v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '=v1}' in asm}} |
|---|
| 41 | |
|---|
| 42 | __asm("v_add_f32_e32 v1, v2, v3" : "={v[1]}"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : ); |
|---|
| 43 | __asm("v_add_f32_e32 v1, v2, v3" : "={v[1}"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : ); // expected-error {{invalid output constraint '={v[1}' in asm}} |
|---|
| 44 | __asm("v_add_f32_e32 v1, v2, v3" : "={v[1]"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : ); // expected-error {{invalid output constraint '={v[1]' in asm}} |
|---|
| 45 | __asm("v_add_f32_e32 v1, v2, v3" : "={v[a]}"(ci) : "{v[2]}"(ai), "{v[3]}"(bi) : ); // expected-error {{invalid output constraint '={v[a]}' in asm}} |
|---|
| 46 | |
|---|
| 47 | __asm("v_add_f32_e32 v1, v2, v3" : "=v"(ci) : "v"(ai), "v"(bi) : ); |
|---|
| 48 | __asm("v_add_f32_e32 v1, v2, v3" : "=v1"(ci) : "v2"(ai), "v3"(bi) : ); /// expected-error {{invalid output constraint '=v1' in asm}} |
|---|
| 49 | |
|---|
| 50 | __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{a}"(ai), "{v3}"(bi) : ); // expected-error {{invalid input constraint '{a}' in asm}} |
|---|
| 51 | __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{a}"(bi) : ); // expected-error {{invalid input constraint '{a}' in asm}} |
|---|
| 52 | c[i] = ci; |
|---|
| 53 | } |
|---|
| 54 | |
|---|
| 55 | __kernel void |
|---|
| 56 | test_double(const __global double *a, const __global double *b, __global double *c, unsigned i) |
|---|
| 57 | { |
|---|
| 58 | double ai = a[i]; |
|---|
| 59 | double bi = b[i]; |
|---|
| 60 | double ci; |
|---|
| 61 | |
|---|
| 62 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); |
|---|
| 63 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "=v{[1:2]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '=v{[1:2]}' in asm}} |
|---|
| 64 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]a}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2]a}' in asm}} |
|---|
| 65 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]}a"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2]}a' in asm}} |
|---|
| 66 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:' in asm}} |
|---|
| 67 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:]}' in asm}} |
|---|
| 68 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[:2]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[:2]}' in asm}} |
|---|
| 69 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2]"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2]' in asm}} |
|---|
| 70 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[1:2}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[1:2}' in asm}} |
|---|
| 71 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "={v[2:1]}"(ci) : "{v[3:4]}"(ai), "{v[5:6]}"(bi) : ); //expected-error {{invalid output constraint '={v[2:1]}' in asm}} |
|---|
| 72 | |
|---|
| 73 | __asm("v_add_f64_e64 v[1:2], v[3:4], v[5:6]" : "=v[1:2]"(ci) : "v[3:4]"(ai), "v[5:6]"(bi) : ); //expected-error {{invalid output constraint '=v[1:2]' in asm}} |
|---|
| 74 | |
|---|
| 75 | c[i] = ci; |
|---|
| 76 | } |
|---|
| 77 | |
|---|
| 78 | void test_long(int arg0) { |
|---|
| 79 | long v15_16; |
|---|
| 80 | __asm volatile("v_lshlrev_b64 v[15:16], 0, %0" : "={v[15:16]}"(v15_16) : "v"(arg0)); |
|---|
| 81 | } |
|---|
| 82 | |
|---|