[AMDGPU] Added 'A' constraint for inline assembler

[dp]

Summary

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

[dp]

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

22

arsenm wrote: Missing validation for acceptable values?

I'm not sure it is possible to validate input values at this point because we do not know input value size. For example, 0x3118 is an inlineable value for f16, but not for f32.

[arsenm]

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

22

I would expect us to know the type from the C type, but I think we are accepting int for 16-bit asm currently

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;