[AMDGPU] Allow AGPR in inline asm

[arsenm]

clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl

13–16

These mostly aren't the right types?

[rampitec]

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.

[yaxunl]

added agprs to GCCRegNames and fixed types in test

[rampitec]

clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl

16

reg_c is also float32.

36

New line.

[yaxunl]

fix test

clang/lib/Basic/Targets/AMDGPU.h

public:
 
   ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {
     return None;
   }
 
   /// 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({
         "exec", "vcc", "flat_scratch", "m0", "scc", "tba", "tma",
         "flat_scratch_lo", "flat_scratch_hi", "vcc_lo", "vcc_hi", "exec_lo",
         "exec_hi", "tma_lo", "tma_hi", "tba_lo", "tba_hi",
     });
 
     StringRef S(Name);
     bool HasLeftParen = false;
     if (S.front() == '{') {
       HasLeftParen = true;
       S = S.drop_front();
     }
     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('}');
       if (!SpecialRegs.count(S.substr(0, E)))
         return false;
       S = S.drop_front(E + 1);
       if (!S.empty())
         return false;
       // Found {S} where S is a special register.
       Info.setAllowsRegister();
       Name = S.data() - 1;
       return true;
     }
     S = S.drop_front();
     if (!HasLeftParen) {
       if (!S.empty())
         return false;
-      // Found s or v.
+      // Found s, v or a.
       Info.setAllowsRegister();
       Name = S.data() - 1;
       return true;
     }
     bool HasLeftBracket = false;
     if (!S.empty() && S.front() == '[') {
       HasLeftBracket = true;
       S = S.drop_front();
if (HasLeftBracket) {
         return false;
       S = S.drop_front();
     }
     if (S.empty() || S.front() != '}')
       return false;
     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;
   }
 
   // \p Constraint will be left pointing at the last character of
   // the constraint.  In practice, it won't be changed unless the
   // constraint is longer than one character.

clang/lib/Basic/Targets/AMDGPU.cpp

const char *const AMDGPUTargetInfo::GCCRegNames[] = {
   "s68", "s69", "s70", "s71", "s72", "s73", "s74", "s75", "s76",
   "s77", "s78", "s79", "s80", "s81", "s82", "s83", "s84", "s85",
   "s86", "s87", "s88", "s89", "s90", "s91", "s92", "s93", "s94",
   "s95", "s96", "s97", "s98", "s99", "s100", "s101", "s102", "s103",
   "s104", "s105", "s106", "s107", "s108", "s109", "s110", "s111", "s112",
   "s113", "s114", "s115", "s116", "s117", "s118", "s119", "s120", "s121",
   "s122", "s123", "s124", "s125", "s126", "s127", "exec", "vcc", "scc",
   "m0", "flat_scratch", "exec_lo", "exec_hi", "vcc_lo", "vcc_hi",
-  "flat_scratch_lo", "flat_scratch_hi"
+  "flat_scratch_lo", "flat_scratch_hi",
+  "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", "a32", "a33", "a34", "a35",
+  "a36", "a37", "a38", "a39", "a40", "a41", "a42", "a43", "a44",
+  "a45", "a46", "a47", "a48", "a49", "a50", "a51", "a52", "a53",
+  "a54", "a55", "a56", "a57", "a58", "a59", "a60", "a61", "a62",
+  "a63", "a64", "a65", "a66", "a67", "a68", "a69", "a70", "a71",
+  "a72", "a73", "a74", "a75", "a76", "a77", "a78", "a79", "a80",
+  "a81", "a82", "a83", "a84", "a85", "a86", "a87", "a88", "a89",
+  "a90", "a91", "a92", "a93", "a94", "a95", "a96", "a97", "a98",
+  "a99", "a100", "a101", "a102", "a103", "a104", "a105", "a106", "a107",
+  "a108", "a109", "a110", "a111", "a112", "a113", "a114", "a115", "a116",
+  "a117", "a118", "a119", "a120", "a121", "a122", "a123", "a124", "a125",
+  "a126", "a127", "a128", "a129", "a130", "a131", "a132", "a133", "a134",
+  "a135", "a136", "a137", "a138", "a139", "a140", "a141", "a142", "a143",
+  "a144", "a145", "a146", "a147", "a148", "a149", "a150", "a151", "a152",
+  "a153", "a154", "a155", "a156", "a157", "a158", "a159", "a160", "a161",
+  "a162", "a163", "a164", "a165", "a166", "a167", "a168", "a169", "a170",
+  "a171", "a172", "a173", "a174", "a175", "a176", "a177", "a178", "a179",
+  "a180", "a181", "a182", "a183", "a184", "a185", "a186", "a187", "a188",
+  "a189", "a190", "a191", "a192", "a193", "a194", "a195", "a196", "a197",
+  "a198", "a199", "a200", "a201", "a202", "a203", "a204", "a205", "a206",
+  "a207", "a208", "a209", "a210", "a211", "a212", "a213", "a214", "a215",
+  "a216", "a217", "a218", "a219", "a220", "a221", "a222", "a223", "a224",
+  "a225", "a226", "a227", "a228", "a229", "a230", "a231", "a232", "a233",
+  "a234", "a235", "a236", "a237", "a238", "a239", "a240", "a241", "a242",
+  "a243", "a244", "a245", "a246", "a247", "a248", "a249", "a250", "a251",
+  "a252", "a253", "a254", "a255"
 };
 
 ArrayRef<const char *> AMDGPUTargetInfo::getGCCRegNames() const {
   return llvm::makeArrayRef(GCCRegNames);
 }
 
 bool AMDGPUTargetInfo::initFeatureMap(
     llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags, StringRef CPU,

clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl

 // 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
+
+typedef float float32 __attribute__((ext_vector_type(32)));
 
 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() {
+  float32 acc_c;
+  float reg_a;
+  float reg_b;
+  float32 reg_c;
+  // 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}"
+  __asm ("v_mfma_f32_32x32x1f32 %0, %1, %2, %3"
+         : "=a"(acc_c)
+         : "v"(reg_a), "v"(reg_b), "a"(reg_c)
+         : "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");
+
+   // 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]}"
+  __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]}"(reg_c));
+
+  // CHECK: call float asm "v_accvgpr_read_b32 $0, $1", "={a1},{a1}"
+  __asm ("v_accvgpr_read_b32 %0, %1"
+         : "={a1}"(reg_a)
+         : "{a1}"(reg_b));
+}