[AMDGPU] Allow flexible register names in inline asm constraints

[yaxunl]

Summary

Currently AMDGPU inline asm only allow "v" and "s" as register names in constraints.

This patch allows the following register names in constraints: (n, m is unsigned integer, n < m)

v

s

{vn} or {v[n]}

{sn} or {s[n]}

{S} , where S is a special register name

{v[n:m]}

{s[n:m]}

[arsenm]

test/Sema/inline-asm-validate-amdgpu.cl

26

I don't understand. The backend parsed syntax is v[3:4]. Why should v3:4 be valid? Also in this example you are using a 64-bit input to a 32-bit operand

[arsenm]

test/Sema/inline-asm-validate-amdgpu.cl

26

i.e. this won't parse in the backend and is invalid. The brackets are required

[yaxunl]

test/Sema/inline-asm-validate-amdgpu.cl

26

Will fix the syntax about [].

The test did not consider validity of the inline assembly for the backend since FE only check format. I agree it is better to make it valid for backend and will try to fix that.

[yaxunl]

Revised by Matt's comments.

[b-sumner]

The assembler accepts v[N] in addition to vN. I'm not sure if that is needed here.

[yaxunl]

In D37568#863735, @b-sumner wrote:

The assembler accepts v[N] in addition to vN. I'm not sure if that is needed here.

Then we'd better also allow that in constraints to avoid confusion of users.

[yaxunl]

Allow {v[n]} and {s[n]}. Add more tests.

[arsenm]

lib/Basic/Targets/AMDGPU.h

69

Typo wheere

91

I'm not sure I understand these data() - 1s.

test/Sema/inline-asm-validate-amdgpu.cl

55

While you're here can we add some tests for the immediate constraints? There was a bug report recently when using s_trap with the i constraints for the constant operand.

[yaxunl]

lib/Basic/Targets/AMDGPU.h

91

The caller of this function expects Name is on the last parsed char and will increase it before check the next char.

test/Sema/inline-asm-validate-amdgpu.cl

55

what's the syntax of the immediate constraints? And some examples? Thanks.

[yaxunl]

Fix typo.

[yaxunl]

test/Sema/inline-asm-validate-amdgpu.cl

55

Ping. Can we leave this for another patch? Since there are people waiting for this feature.

[yaxunl]

Ping.

Brian, Stas,

Can you review this since Matt is on vacation? Thanks.

[b-sumner]

LGTM. I think we can leave immediates to another patch.

lib/Basic/Targets/AMDGPU.h

 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_CLANG_LIB_BASIC_TARGETS_AMDGPU_H
 #define LLVM_CLANG_LIB_BASIC_TARGETS_AMDGPU_H
 
 #include "clang/AST/Type.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/TargetOptions.h"
+#include "llvm/ADT/StringSet.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/Support/Compiler.h"
 
 namespace clang {
 namespace targets {
 
 class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
 
enum GPUKind {
     GK_CAYMAN,
     GK_GFX6,
     GK_GFX7,
     GK_GFX8,
     GK_GFX9
   } GPU;
 
   bool hasFP64 : 1;
   bool hasFMAF : 1;
   bool hasLDEXPF : 1;
   const AddrSpace AS;
 
   static bool hasFullSpeedFMAF32(StringRef GPUName) {
     return parseAMDGCNName(GPUName) >= GK_GFX9;
   }
 
   static bool isAMDGCN(const llvm::Triple &TT) {
     return TT.getArch() == llvm::Triple::amdgcn;
   }
 
   static bool isGenericZero(const llvm::Triple &TT) {
     return TT.getEnvironmentName() == "amdgiz" ||
            TT.getEnvironmentName() == "amdgizcl";
   }
 
 public:
   AMDGPUTargetInfo(const llvm::Triple &Triple, const TargetOptions &Opts);
 
   void setAddressSpaceMap(bool DefaultIsPrivate);
 
   void adjust(LangOptions &Opts) override;
 
   uint64_t getPointerWidthV(unsigned AddrSpace) const override {
     if (GPU <= GK_CAYMAN)
       return 32;
 
     if (AddrSpace == AS.Private || AddrSpace == AS.Local) {
       return 32;
     }
public:
   const char *getClobbers() const override { return ""; }
 
   ArrayRef<const char *> getGCCRegNames() const override;
 
   ArrayRef<TargetInfo::GCCRegAlias> getGCCRegAliases() const override {
     return None;
   }
 
+  /// Accepted register names: (n, m is unsigned integer, n < m)
+  /// v
+  /// s
+  /// {vn}
+  /// {sn}
+  /// {S} , wheere S is a special register name
+  ////{v[n:m]}
+  /// {s[n:m]}
   bool validateAsmConstraint(const char *&Name,
                              TargetInfo::ConstraintInfo &Info) const override {
-    switch (*Name) {
-    default:
-      break;
-    case 'v': // vgpr
-    case 's': // sgpr
+    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.front() != 'v' && S.front() != 's') {
+      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.
       Info.setAllowsRegister();
+      Name = S.data() - 1;
       return true;
     }
+    bool HasLeftBracket = false;
+    if (S.front() == '[') {
+      HasLeftBracket = true;
+      S = S.drop_front();
+    }
+    unsigned long long N;
+    if (consumeUnsignedInteger(S, 10, N))
+      return false;
+    if (S.front() == ':') {
+      if (!HasLeftBracket)
+        return false;
+      S = S.drop_front();
+      unsigned long long M;
+      if (consumeUnsignedInteger(S, 10, M) || N >= M)
         return false;
+      if (S.front() != ']')
+        return false;
+      S = S.drop_front();
+    }
+    if (S.front() != '}')
+      return false;
+    S = S.drop_front();
+    if (!S.empty())
+      return false;
+    // Found {vn}, {sn}, {v[n:m]}, or {s[n:m]}.
+    Info.setAllowsRegister();
+    Name = S.data() - 1;
+    return true;
   }
 
   bool
   initFeatureMap(llvm::StringMap<bool> &Features, DiagnosticsEngine &Diags,
                  StringRef CPU,
                  const std::vector<std::string> &FeatureVec) const override;
 
   void adjustTargetOptions(const CodeGenOptions &CGOpts,

test/CodeGenOpenCL/amdgcn-inline-asm.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn -O0 -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @ker
+__kernel void
+ker(const __global float *a, const __global float *b, __global float *c, unsigned i)
+{
+    float ai = a[i];
+    float bi = b[i];
+    float ci;
+    // CHECK: call float asm "v_add_f32_e32 v1, v2, v3", "={v1},{v2},{v3}"(float %{{.*}}, float %{{.*}})
+    __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : );
+    // CHECK: call float asm "v_add_f32_e32 $0, $1, $2", "={v1},{v2},{v3}"(float %{{.*}}, float %{{.*}})
+    __asm("v_add_f32_e32 %0, %1, %2" : "={v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : );
+    c[i] = ci;
+}

test/Sema/inline-asm-validate-amdgpu.cl

 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -x cl -triple amdgcn -fsyntax-only  %s
-// expected-no-diagnostics
+// RUN: %clang_cc1 -triple amdgcn -fsyntax-only -verify %s
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
 
 kernel void test () {
 
   int sgpr = 0, vgpr = 0, imm = 0;
 
   // sgpr constraints
   __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "s" (imm) : );
 
   // vgpr constraints
   __asm__ ("v_mov_b32 %0, %1" : "=v" (vgpr) : "v" (imm) : );
 }
+
+__kernel void
+test_float(const __global float *a, const __global float *b, __global float *c, unsigned i)
+{
+    float ai = a[i];
+    float bi = b[i];
+    float ci;
+
+    __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{v3}"(bi) : );
+    __asm("v_add_f32_e32 v1, v2, v3" : ""(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '' in asm}}
+    __asm("v_add_f32_e32 v1, v2, v3" : "="(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '=' in asm}}
+    __asm("v_add_f32_e32 v1, v2, v3" : "={a}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={a}' in asm}}
+    __asm("v_add_f32_e32 v1, v2, v3" : "={v1a}"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1a}' in asm}}
+    __asm("v_add_f32_e32 v1, v2, v3" : "={v1}a"(ci) : "{v2}"(ai), "{v3}"(bi) : ); // expected-error {{invalid output constraint '={v1}a' in asm}}
+
+    __asm("v_add_f32_e32 v1, v2, exec" : "={v1}"(ci) : "{v2}"(ai), "{exec}"(bi) : );
+    __asm("v_add_f32_e32 v1, v2, exec" : "={v1}"(ci) : "{v2}"(ai), "{exec}a"(bi) : ); // expected-error {{invalid input constraint '{exec}a' in asm}}
+
+    __asm("v_add_f32_e32 v1, v2, v3" : "=v"(ci) : "v"(ai), "v"(bi) : );
+    __asm("v_add_f32_e32 v1, v2, v3" : "=v1"(ci) : "v2"(ai), "v3"(bi) : ); /// expected-error {{invalid output constraint '=v1' in asm}}
+
+    __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{a}"(ai), "{v3}"(bi) : ); // expected-error {{invalid input constraint '{a}' in asm}}
+    __asm("v_add_f32_e32 v1, v2, v3" : "={v1}"(ci) : "{v2}"(ai), "{a}"(bi) : ); // expected-error {{invalid input constraint '{a}' in asm}}
+    c[i] = ci;
+}
+
+__kernel void
+test_double(const __global double *a, const __global double *b, __global double *c, unsigned i)
+{
+    double ai = a[i];
+    double bi = b[i];
+    double ci;
+
+    __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) : );
+    __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}}
+    __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}}
+    __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}}
+
+    __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}}
+
+    c[i] = ci;
+}