[AMDGPU] Added support of new inline assembler constraints

Added support for constraints 'I', 'J', 'L', 'B', 'C', 'Kf', 'DA', 'DB'.

See https://gcc.gnu.org/onlinedocs/gcc/Machine-Constraints.html#Machine-Constraints.

Reviewers: arsenm, rampitec

Differential Revision: https://reviews.llvm.org/D81657

diff --git a/clang/lib/Basic/Targets/AMDGPU.h b/clang/lib/Basic/Targets/AMDGPU.h
index 387b91a..d039449 100644 (file)
--- a/clang/lib/Basic/Targets/AMDGPU.h
+++ b/clang/lib/Basic/Targets/AMDGPU.h
@@ -130,8 +130,26 @@ public:
         "exec_hi", "tma_lo", "tma_hi", "tba_lo", "tba_hi",
     });
 
+    switch (*Name) {
+    case 'I':
+      Info.setRequiresImmediate(-16, 64);
+      return true;
+    case 'J':
+      Info.setRequiresImmediate(-32768, 32767);
+      return true;
+    case 'A':
+    case 'B':
+    case 'C':
+      Info.setRequiresImmediate();
+      return true;
+    default:
+      break;
+    }
+
     StringRef S(Name);
-    if (S == "A") {
+
+    if (S == "DA" || S == "DB") {
+      Name++;
       Info.setRequiresImmediate();
       return true;
     }
@@ -203,6 +221,12 @@ public:
   // the constraint.  In practice, it won't be changed unless the
   // constraint is longer than one character.
   std::string convertConstraint(const char *&Constraint) const override {
+
+    StringRef S(Constraint);
+    if (S == "DA" || S == "DB") {
+      return std::string("^") + std::string(Constraint++, 2);
+    }
+
     const char *Begin = Constraint;
     TargetInfo::ConstraintInfo Info("", "");
     if (validateAsmConstraint(Constraint, Info))

diff --git a/clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl b/clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
index 3709077..259c123 100644 (file)
--- a/clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl
@@ -33,3 +33,17 @@ kernel void test_agpr() {
          : "={a1}"(reg_a)
          : "{a1}"(reg_b));
 }
+
+kernel void test_constraint_DA() {
+  const long x = 0x200000001;
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DA"(i64 8589934593)
+  __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DA"(x));
+}
+
+kernel void test_constraint_DB() {
+  const long x = 0x200000001;
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1 & 0xFFFFFFFF", "=v,^DB"(i64 8589934593)
+  __asm volatile("v_mov_b32 %0, %1 & 0xFFFFFFFF" : "=v"(res) : "DB"(x));
+}

diff --git a/clang/test/Sema/inline-asm-validate-amdgpu.cl b/clang/test/Sema/inline-asm-validate-amdgpu.cl
index 3d64882..418952c 100644 (file)
--- a/clang/test/Sema/inline-asm-validate-amdgpu.cl
+++ b/clang/test/Sema/inline-asm-validate-amdgpu.cl
@@ -18,9 +18,35 @@ kernel void test () {
   // vgpr constraints
   __asm__ ("v_mov_b32 %0, %1" : "=v" (vgpr) : "v" (imm) : );
 
-  // 'A' constraint
+  // 'I' constraint (an immediate integer in the range -16 to 64)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (imm) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (-16) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (64) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (-17) : ); // expected-error {{value '-17' out of range for constraint 'I'}}
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "I" (65) : ); // expected-error {{value '65' out of range for constraint 'I'}}
+
+  // 'J' constraint (an immediate 16-bit signed integer)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (imm) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (-32768) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (32767) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (-32769) : ); // expected-error {{value '-32769' out of range for constraint 'J'}}
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "J" (32768) : ); // expected-error {{value '32768' out of range for constraint 'J'}}
+
+  // 'A' constraint (an immediate constant that can be inlined)
   __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "A" (imm) : );
 
+  // 'B' constraint (an immediate 32-bit signed integer)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "B" (imm) : );
+
+  // 'C' constraint (an immediate 32-bit unsigned integer or 'A' constraint)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "C" (imm) : );
+
+  // 'DA' constraint (an immediate 64-bit constant that can be split into two 'A' constants)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "DA" (imm) : );
+
+  // 'DB' constraint (an immediate 64-bit constant that can be split into two 'B' constants)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "DB" (imm) : );
+
 }
 
 __kernel void