[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;



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;

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits