[AMDGPU] Added 'A' constraint for inline assembler

[dp]

Summary

Added 'A' constraint which requires an integer or a floating-point inline constant.

[arsenm]

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

23

Missing validation for acceptable values?

[dp]

This comment was removed by dp.

[dp]

Sorry, I updated a wrong diff.

clang/lib/Basic/Targets/AMDGPU.h

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);
+    if (S == "A") {
+      Info.setRequiresImmediate();
+      return true;
+    }
+
     bool HasLeftParen = false;
     if (S.front() == '{') {
       HasLeftParen = true;
       S = S.drop_front();
     }
     if (S.empty())
       return false;
     if (S.front() != 'v' && S.front() != 's' && S.front() != 'a') {

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

kernel void test () {
 
   __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
+  __asm__ ("s_mov_b32 %0, %1" : "=s" (sgpr) : "A" (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;