[clang] 3c6c2ec - [AMDGPU] Added 'A' constraint for inline assembler

[Dmitry Preobrazhensky via cfe-commits]

Author: Dmitry Preobrazhensky
Date: 2020-05-25T17:47:06+03:00
New Revision: 3c6c2ecd6efa393e7a8422d88e5d4ada0970e47e

URL: https://github.com/llvm/llvm-project/commit/3c6c2ecd6efa393e7a8422d88e5d4ada0970e47e
DIFF: https://github.com/llvm/llvm-project/commit/3c6c2ecd6efa393e7a8422d88e5d4ada0970e47e.diff

LOG: [AMDGPU] Added 'A' constraint for inline assembler

Summary: 'A' constraint requires an immediate int or fp constant that can be inlined in an instruction encoding.
This is the second part of the change. The llvm part has been committed as b087b91c9170.
See https://reviews.llvm.org/D78494

Reviewers: arsenm, rampitec

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

Added: 
    

Modified: 
    clang/lib/Basic/Targets/AMDGPU.h
    clang/test/Sema/inline-asm-validate-amdgpu.cl

Removed: 
    


################################################################################
diff  --git a/clang/lib/Basic/Targets/AMDGPU.h b/clang/lib/Basic/Targets/AMDGPU.h
index d0e88e223e95..6c9060aa3f7b 100644
--- a/clang/lib/Basic/Targets/AMDGPU.h
+++ b/clang/lib/Basic/Targets/AMDGPU.h
@@ -131,6 +131,11 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
     });
 
     StringRef S(Name);
+    if (S == "A") {
+      Info.setRequiresImmediate();
+      return true;
+    }
+
     bool HasLeftParen = false;
     if (S.front() == '{') {
       HasLeftParen = true;

diff  --git a/clang/test/Sema/inline-asm-validate-amdgpu.cl b/clang/test/Sema/inline-asm-validate-amdgpu.cl
index 51009ecb3f1e..3d6488227ef2 100644
--- a/clang/test/Sema/inline-asm-validate-amdgpu.cl
+++ b/clang/test/Sema/inline-asm-validate-amdgpu.cl
@@ -17,6 +17,10 @@ kernel void test () {
 
   // 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