[AMDGPU] Added support of new inline assembler constraints

[dp]

Summary

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

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

[arsenm]

clang/lib/Basic/Targets/AMDGPU.h

156

Not sure what this means? Only accepts -1 as a value? Why would you ever need this special case?

[dp]

Corrected a typo in tests.

[dp]

clang/lib/Basic/Targets/AMDGPU.h

156

I have no idea why GCC guys added this constraint - this is an interesting question. Anyway we adding this for the sake of compatibility with GCC.

[arsenm]

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

41–42

Can you add tests for this with i64/i16 and the packed vectors?

[dp]

Added more tests for 'Kf' constraint.

clang/lib/Basic/Targets/AMDGPU.h

public:
   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",
     });
 
+    switch (*Name) {
+    case 'I':
+      Info.setRequiresImmediate(-16, 64);
+      return true;
+    case 'J':
+      Info.setRequiresImmediate(-32768, 32767);
+      return true;
+    case 'L':
+      Info.setRequiresImmediate(0, 65535);
+      return true;
+    case 'A':
+    case 'B':
+    case 'C':
+      Info.setRequiresImmediate();
+      return true;
+    default:
+      break;
+    }
+
     StringRef S(Name);
-    if (S == "A") {
+
+    if (S == "Kf") {
+      Name++;
+      Info.setRequiresImmediate(-1, -1);
+      return true;
+    }
+    if (S == "DA" || S == "DB") {
+      Name++;
       Info.setRequiresImmediate();
       return true;
     }
 
     bool HasLeftParen = false;
     if (S.front() == '{') {
       HasLeftParen = true;
       S = S.drop_front();
bool validateAsmConstraint(const char *&Name,
     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.
   std::string convertConstraint(const char *&Constraint) const override {
+
+    StringRef S(Constraint);
+    if (S == "Kf" || S == "DA" || S == "DB") {
+      return std::string("^") + std::string(Constraint++, 2);
+    }
+
     const char *Begin = Constraint;
     TargetInfo::ConstraintInfo Info("", "");
     if (validateAsmConstraint(Constraint, Info))
       return std::string(Begin).substr(0, Constraint - Begin + 1);
 
     Constraint = Begin;
     return std::string(1, *Constraint);
   }

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

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -emit-llvm -O0 -o - -triple amdgcn %s | FileCheck %s
 
 typedef float float32 __attribute__((ext_vector_type(32)));
+typedef short __attribute__((ext_vector_type(2))) short2;
 
 kernel void test_long(int arg0) {
   long v15_16;
   // 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() {
__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));
 }
+
+kernel void test_constraint_Kf_i32() {
+  const int m1 = -1;
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(i32 -1)
+  __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1));
+}
+
+kernel void test_constraint_Kf_i64() {
+  const long m1 = -1;
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(i64 -1)
+  __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1));
+}
+
+kernel void test_constraint_Kf_i16() {
+  const short m1 = -1;
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(i16 -1)
+  __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1));
+}
+
+kernel void test_constraint_Kf_i16x2() {
+  const short2 m1 = (short2)(-1, -1);
+  int res;
+  // CHECK: call i32 asm sideeffect "v_mov_b32 $0, $1", "=v,^Kf"(<2 x i16> <i16 -1, i16 -1>)
+  __asm volatile("v_mov_b32 %0, %1" : "=v"(res) : "Kf"(m1));
+}
+
+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));
+}

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

 // REQUIRES: amdgpu-registered-target
 // 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;
+  int sgpr = 0, vgpr = 0, imm = 0, m1 = -1;
 
   // sgpr constraints
   __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "s" (imm) : );
 
   __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec}" (imm) : );
   __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exe" (imm) : ); // expected-error {{invalid input constraint '{exe' in asm}}
   __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec" (imm) : ); // expected-error {{invalid input constraint '{exec' in asm}}
   __asm__ ("s_mov_b32 %0, %1" : "={s1}" (sgpr) : "{exec}a" (imm) : ); // expected-error {{invalid input constraint '{exec}a' in asm}}
 
   // 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'}}
+
+  // 'L' constraint (an immediate 15-bit unsigned integer)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (imm) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (0) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (65535) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (-1) : ); // expected-error {{value '-1' out of range for constraint 'L'}}
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "L" (65536) : ); // expected-error {{value '65536' out of range for constraint 'L'}}
+
+  // '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) : );
+
+  // 'Kf' constraint (-1)
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (m1) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (-1) : );
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (-2) : ); // expected-error {{value '-2' out of range for constraint 'Kf'}}
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "Kf" (0) : ); // expected-error {{value '0' out of range for constraint 'Kf'}}
+
+  // '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
 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;