[PATCH] D77329: [AMDGPU] Allow AGPR in inline asm

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: rampitec, arsenm.
Herald added subscribers: kerbowa, t-tye, tpr, dstuttard, nhaehnle, wdng, jvesely, kzhuravl.
rampitec accepted this revision.
rampitec added a comment.
This revision is now accepted and ready to land.

Thanks.

Could you also update AMDGPUTargetInfo::GCCRegNames[] (in a separate change)? It is used in clobber constraints.
JBTW, it does not support register tuples even for V and S now.



================
Comment at: clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl:11-14
+  float acc_c;
+  float reg_a;
+  float reg_b;
+  float reg_c;
----------------
These mostly aren't the right types?


https://reviews.llvm.org/D77329

Files:
  clang/lib/Basic/Targets/AMDGPU.h
  clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl


Index: clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
===================================================================
--- clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
+++ clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
@@ -1,8 +1,21 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn %s | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm -O0 -o - -triple amdgcn %s | FileCheck %s
 
 kernel void test_long(int arg0) {
   long v15_16;
-  // CHECK: tail call i64 asm sideeffect "v_lshlrev_b64 v[15:16], 0, $0", "={v[15:16]},v"(i32 %arg0)
+  // CHECK: call i64 asm sideeffect "v_lshlrev_b64 v[15:16], 0, $0", "={v[15:16]},v"
   __asm volatile("v_lshlrev_b64 v[15:16], 0, %0" : "={v[15:16]}"(v15_16) : "v"(arg0));
 }
+
+kernel void test_agpr() {
+  float acc_c;
+  float reg_a;
+  float reg_b;
+  float reg_c;
+  // CHECK: call float asm "v_mfma_f32_32x32x1f32 $0, $1, $2, $3", "=a,v,v,a"
+  __asm ("v_mfma_f32_32x32x1f32 %0, %1, %2, %3" : "=a"(acc_c) : "v"(reg_a), "v"(reg_b), "a"(acc_c));
+  // CHECK: call float asm sideeffect "v_mfma_f32_32x32x1f32 a[0:31], $0, $1, a[0:31]", "={a[0:31]},v,v,{a[0:31]}"
+  __asm volatile("v_mfma_f32_32x32x1f32 a[0:31], %0, %1, a[0:31]" : "={a[0:31]}"(acc_c) : "v"(reg_a),"v"(reg_b), "{a[0:31]}"(acc_c));
+  // CHECK: call float asm "v_accvgpr_read_b32 $0, $1", "={a1},{a1}"
+  __asm ("v_accvgpr_read_b32 %0, %1" : "={a1}"(reg_c) : "{a1}"(acc_c));
+}
\ No newline at end of file
Index: clang/lib/Basic/Targets/AMDGPU.h
===================================================================
--- clang/lib/Basic/Targets/AMDGPU.h
+++ clang/lib/Basic/Targets/AMDGPU.h
@@ -114,11 +114,14 @@
   /// Accepted register names: (n, m is unsigned integer, n < m)
   /// v
   /// s
+  /// a
   /// {vn}, {v[n]}
   /// {sn}, {s[n]}
+  /// {an}, {a[n]}
   /// {S} , where S is a special register name
   ////{v[n:m]}
   /// {s[n:m]}
+  /// {a[n:m]}
   bool validateAsmConstraint(const char *&Name,
                              TargetInfo::ConstraintInfo &Info) const override {
     static const ::llvm::StringSet<> SpecialRegs({
@@ -135,7 +138,7 @@
     }
     if (S.empty())
       return false;
-    if (S.front() != 'v' && S.front() != 's') {
+    if (S.front() != 'v' && S.front() != 's' && S.front() != 'a') {
       if (!HasLeftParen)
         return false;
       auto E = S.find('}');
@@ -153,7 +156,7 @@
     if (!HasLeftParen) {
       if (!S.empty())
         return false;
-      // Found s or v.
+      // Found s, v or a.
       Info.setAllowsRegister();
       Name = S.data() - 1;
       return true;
@@ -184,7 +187,8 @@
     S = S.drop_front();
     if (!S.empty())
       return false;
-    // Found {vn}, {sn}, {v[n]}, {s[n]}, {v[n:m]}, or {s[n:m]}.
+    // Found {vn}, {sn}, {an}, {v[n]}, {s[n]}, {a[n]}, {v[n:m]}, {s[n:m]}
+    // or {a[n:m]}.
     Info.setAllowsRegister();
     Name = S.data() - 1;
     return true;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D77329.254587.patch
Type: text/x-patch
Size: 2896 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200402/0c625550/attachment-0001.bin>