|
150
|
1 // REQUIRES: amdgpu-registered-target
|
|
173
|
2 // RUN: %clang_cc1 -emit-llvm -O0 -o - -triple amdgcn %s | FileCheck %s
|
|
|
3
|
|
|
4 typedef float float32 __attribute__((ext_vector_type(32)));
|
|
150
|
5
|
|
|
6 kernel void test_long(int arg0) {
|
|
|
7 long v15_16;
|
|
173
|
8 // CHECK: call i64 asm sideeffect "v_lshlrev_b64 v[15:16], 0, $0", "={v[15:16]},v"
|
|
150
|
9 __asm volatile("v_lshlrev_b64 v[15:16], 0, %0" : "={v[15:16]}"(v15_16) : "v"(arg0));
|
|
|
10 }
|
|
173
|
11
|
|
|
12 kernel void test_agpr() {
|
|
|
13 float32 acc_c;
|
|
|
14 float reg_a;
|
|
|
15 float reg_b;
|
|
|
16 float32 reg_c;
|
|
|
17 // CHECK: call <32 x float> asm "v_mfma_f32_32x32x1f32 $0, $1, $2, $3", "=a,v,v,a,~{a0},~{a1},~{a2},~{a3},~{a4},~{a5},~{a6},~{a7},~{a8},~{a9},~{a10},~{a11},~{a12},~{a13},~{a14},~{a15},~{a16},~{a17},~{a18},~{a19},~{a20},~{a21},~{a22},~{a23},~{a24},~{a25},~{a26},~{a27},~{a28},~{a29},~{a30},~{a31}"
|
|
|
18 __asm ("v_mfma_f32_32x32x1f32 %0, %1, %2, %3"
|
|
|
19 : "=a"(acc_c)
|
|
|
20 : "v"(reg_a), "v"(reg_b), "a"(reg_c)
|
|
|
21 : "a0", "a1", "a2", "a3", "a4", "a5", "a6", "a7",
|
|
|
22 "a8", "a9", "a10", "a11", "a12", "a13", "a14", "a15",
|
|
|
23 "a16", "a17", "a18", "a19", "a20", "a21", "a22", "a23",
|
|
|
24 "a24", "a25", "a26", "a27", "a28", "a29", "a30", "a31");
|
|
|
25
|
|
|
26 // CHECK: call <32 x float> asm sideeffect "v_mfma_f32_32x32x1f32 a[0:31], $0, $1, a[0:31]", "={a[0:31]},v,v,{a[0:31]}"
|
|
|
27 __asm volatile("v_mfma_f32_32x32x1f32 a[0:31], %0, %1, a[0:31]"
|
|
|
28 : "={a[0:31]}"(acc_c)
|
|
|
29 : "v"(reg_a),"v"(reg_b), "{a[0:31]}"(reg_c));
|
|
|
30
|
|
|
31 // CHECK: call float asm "v_accvgpr_read_b32 $0, $1", "={a1},{a1}"
|
|
|
32 __asm ("v_accvgpr_read_b32 %0, %1"
|
|
|
33 : "={a1}"(reg_a)
|
|
|
34 : "{a1}"(reg_b));
|
|
|
35 }
|
|
221
|
36
|
|
|
37 kernel void test_constraint_DA() {
|
|
|
38 const long x = 0x200000001;
|
|
|
39 int res;
|
|
|
40 // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DA"(i64 8589934593)
|
|
|
41 __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DA"(x));
|
|
|
42 }
|
|
|
43
|
|
|
44 kernel void test_constraint_DB() {
|
|
|
45 const long x = 0x200000001;
|
|
|
46 int res;
|
|
|
47 // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DB"(i64 8589934593)
|
|
|
48 __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DB"(x));
|
|
|
49 }
|
