[PATCH] D79493: [AMDGPU] Added 'A' constraint for inline assembler

[Dmitry Preobrazhensky via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rG3c6c2ecd6efa: [AMDGPU] Added 'A' constraint for inline assembler (authored by dp).
Herald added a project: clang.
Herald added a subscriber: cfe-commits.

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D79493/new/

https://reviews.llvm.org/D79493

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


Index: clang/test/Sema/inline-asm-validate-amdgpu.cl
===================================================================
--- clang/test/Sema/inline-asm-validate-amdgpu.cl
+++ clang/test/Sema/inline-asm-validate-amdgpu.cl
@@ -17,6 +17,10 @@
 
   // 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
Index: clang/lib/Basic/Targets/AMDGPU.h
===================================================================
--- clang/lib/Basic/Targets/AMDGPU.h
+++ clang/lib/Basic/Targets/AMDGPU.h
@@ -131,6 +131,11 @@
     });
 
     StringRef S(Name);
+    if (S == "A") {
+      Info.setRequiresImmediate();
+      return true;
+    }
+
     bool HasLeftParen = false;
     if (S.front() == '{') {
       HasLeftParen = true;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D79493.266023.patch
Type: text/x-patch
Size: 847 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200525/ecdac269/attachment-0001.bin>