[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 b087b91c91.
See https://reviews.llvm.org/D78494

Reviewers: arsenm, rampitec

Differential Revision: https://reviews.llvm.org/D79493
This commit is contained in:
Dmitry Preobrazhensky 2020-05-25 17:45:18 +03:00
parent 7b1dc0015a
commit 3c6c2ecd6e
2 changed files with 9 additions and 0 deletions

View File

@ -131,6 +131,11 @@ public:
});
StringRef S(Name);
if (S == "A") {
Info.setRequiresImmediate();
return true;
}
bool HasLeftParen = false;
if (S.front() == '{') {
HasLeftParen = true;

View File

@ -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