yaxunl created this revision. yaxunl added reviewers: rampitec, arsenm. Herald added subscribers: kerbowa, t-tye, tpr, dstuttard, nhaehnle, wdng, jvesely, kzhuravl. rampitec accepted this revision. rampitec added a comment. This revision is now accepted and ready to land.
Thanks. Could you also update AMDGPUTargetInfo::GCCRegNames[] (in a separate change)? It is used in clobber constraints. JBTW, it does not support register tuples even for V and S now. ================ Comment at: clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl:11-14 + float acc_c; + float reg_a; + float reg_b; + float reg_c; ---------------- These mostly aren't the right types? https://reviews.llvm.org/D77329 Files: clang/lib/Basic/Targets/AMDGPU.h clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl Index: clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl =================================================================== --- clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl +++ clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl @@ -1,8 +1,21 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -O0 -o - -triple amdgcn %s | FileCheck %s kernel void test_long(int arg0) { long v15_16; - // CHECK: tail call i64 asm sideeffect "v_lshlrev_b64 v[15:16], 0, $0", "={v[15:16]},v"(i32 %arg0) + // CHECK: call i64 asm sideeffect "v_lshlrev_b64 v[15:16], 0, $0", "={v[15:16]},v" __asm volatile("v_lshlrev_b64 v[15:16], 0, %0" : "={v[15:16]}"(v15_16) : "v"(arg0)); } + +kernel void test_agpr() { + float acc_c; + float reg_a; + float reg_b; + float reg_c; + // CHECK: call float asm "v_mfma_f32_32x32x1f32 $0, $1, $2, $3", "=a,v,v,a" + __asm ("v_mfma_f32_32x32x1f32 %0, %1, %2, %3" : "=a"(acc_c) : "v"(reg_a), "v"(reg_b), "a"(acc_c)); + // CHECK: call float asm sideeffect "v_mfma_f32_32x32x1f32 a[0:31], $0, $1, a[0:31]", "={a[0:31]},v,v,{a[0:31]}" + __asm volatile("v_mfma_f32_32x32x1f32 a[0:31], %0, %1, a[0:31]" : "={a[0:31]}"(acc_c) : "v"(reg_a),"v"(reg_b), "{a[0:31]}"(acc_c)); + // CHECK: call float asm "v_accvgpr_read_b32 $0, $1", "={a1},{a1}" + __asm ("v_accvgpr_read_b32 %0, %1" : "={a1}"(reg_c) : "{a1}"(acc_c)); +} \ No newline at end of file Index: clang/lib/Basic/Targets/AMDGPU.h =================================================================== --- clang/lib/Basic/Targets/AMDGPU.h +++ clang/lib/Basic/Targets/AMDGPU.h @@ -114,11 +114,14 @@ /// Accepted register names: (n, m is unsigned integer, n < m) /// v /// s + /// a /// {vn}, {v[n]} /// {sn}, {s[n]} + /// {an}, {a[n]} /// {S} , where S is a special register name ////{v[n:m]} /// {s[n:m]} + /// {a[n:m]} bool validateAsmConstraint(const char *&Name, TargetInfo::ConstraintInfo &Info) const override { static const ::llvm::StringSet<> SpecialRegs({ @@ -135,7 +138,7 @@ } if (S.empty()) return false; - if (S.front() != 'v' && S.front() != 's') { + if (S.front() != 'v' && S.front() != 's' && S.front() != 'a') { if (!HasLeftParen) return false; auto E = S.find('}'); @@ -153,7 +156,7 @@ if (!HasLeftParen) { if (!S.empty()) return false; - // Found s or v. + // Found s, v or a. Info.setAllowsRegister(); Name = S.data() - 1; return true; @@ -184,7 +187,8 @@ S = S.drop_front(); if (!S.empty()) return false; - // Found {vn}, {sn}, {v[n]}, {s[n]}, {v[n:m]}, or {s[n:m]}. + // Found {vn}, {sn}, {an}, {v[n]}, {s[n]}, {a[n]}, {v[n:m]}, {s[n:m]} + // or {a[n:m]}. Info.setAllowsRegister(); Name = S.data() - 1; return true;
Index: clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl =================================================================== --- clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl +++ clang/test/CodeGenOpenCL/inline-asm-amdgcn.cl @@ -1,8 +1,21 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn %s | FileCheck %s +// RUN: %clang_cc1 -emit-llvm -O0 -o - -triple amdgcn %s | FileCheck %s kernel void test_long(int arg0) { long v15_16; - // CHECK: tail call i64 asm sideeffect "v_lshlrev_b64 v[15:16], 0, $0", "={v[15:16]},v"(i32 %arg0) + // CHECK: call i64 asm sideeffect "v_lshlrev_b64 v[15:16], 0, $0", "={v[15:16]},v" __asm volatile("v_lshlrev_b64 v[15:16], 0, %0" : "={v[15:16]}"(v15_16) : "v"(arg0)); } + +kernel void test_agpr() { + float acc_c; + float reg_a; + float reg_b; + float reg_c; + // CHECK: call float asm "v_mfma_f32_32x32x1f32 $0, $1, $2, $3", "=a,v,v,a" + __asm ("v_mfma_f32_32x32x1f32 %0, %1, %2, %3" : "=a"(acc_c) : "v"(reg_a), "v"(reg_b), "a"(acc_c)); + // CHECK: call float asm sideeffect "v_mfma_f32_32x32x1f32 a[0:31], $0, $1, a[0:31]", "={a[0:31]},v,v,{a[0:31]}" + __asm volatile("v_mfma_f32_32x32x1f32 a[0:31], %0, %1, a[0:31]" : "={a[0:31]}"(acc_c) : "v"(reg_a),"v"(reg_b), "{a[0:31]}"(acc_c)); + // CHECK: call float asm "v_accvgpr_read_b32 $0, $1", "={a1},{a1}" + __asm ("v_accvgpr_read_b32 %0, %1" : "={a1}"(reg_c) : "{a1}"(acc_c)); +} \ No newline at end of file Index: clang/lib/Basic/Targets/AMDGPU.h =================================================================== --- clang/lib/Basic/Targets/AMDGPU.h +++ clang/lib/Basic/Targets/AMDGPU.h @@ -114,11 +114,14 @@ /// Accepted register names: (n, m is unsigned integer, n < m) /// v /// s + /// a /// {vn}, {v[n]} /// {sn}, {s[n]} + /// {an}, {a[n]} /// {S} , where S is a special register name ////{v[n:m]} /// {s[n:m]} + /// {a[n:m]} bool validateAsmConstraint(const char *&Name, TargetInfo::ConstraintInfo &Info) const override { static const ::llvm::StringSet<> SpecialRegs({ @@ -135,7 +138,7 @@ } if (S.empty()) return false; - if (S.front() != 'v' && S.front() != 's') { + if (S.front() != 'v' && S.front() != 's' && S.front() != 'a') { if (!HasLeftParen) return false; auto E = S.find('}'); @@ -153,7 +156,7 @@ if (!HasLeftParen) { if (!S.empty()) return false; - // Found s or v. + // Found s, v or a. Info.setAllowsRegister(); Name = S.data() - 1; return true; @@ -184,7 +187,8 @@ S = S.drop_front(); if (!S.empty()) return false; - // Found {vn}, {sn}, {v[n]}, {s[n]}, {v[n:m]}, or {s[n:m]}. + // Found {vn}, {sn}, {an}, {v[n]}, {s[n]}, {a[n]}, {v[n:m]}, {s[n:m]} + // or {a[n:m]}. Info.setAllowsRegister(); Name = S.data() - 1; return true;
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits