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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits