https://github.com/Acim-Maravic updated
https://github.com/llvm/llvm-project/pull/124616
>From c9c461bf868e5f874dc0881d007e88dae7a83d43 Mon Sep 17 00:00:00 2001
From: Acim Maravic <acim.mara...@amd.com>
Date: Tue, 28 Jan 2025 20:42:28 +0100
Subject: [PATCH] [LLVM][AMDGPU] Add Intrinsic and Builtin for
ds_bpermute_fi_b32
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 +
.../builtins-amdgcn-gfx11-err.cl | 6 +-
.../CodeGenOpenCL/builtins-amdgcn-gfx12.cl | 23 +++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 +
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 1 +
llvm/lib/Target/AMDGPU/DSInstructions.td | 3 +-
.../AMDGPU/llvm.amdgcn.ds.bpermute.fi.b32.ll | 154 ++++++++++++++++++
llvm/test/MC/AMDGPU/gfx11_unsupported.s | 3 +
8 files changed, 194 insertions(+), 4 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.fi.b32.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 1b29a8e359c205..39e295aced96b2 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -504,6 +504,8 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4i16,
"V4sV4s*1", "nc", "gf
TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4f16, "V4hV4h*1", "nc",
"gfx12-insts,wavefrontsize64")
TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4bf16, "V4yV4y*1", "nc",
"gfx12-insts,wavefrontsize64")
+TARGET_BUILTIN(__builtin_amdgcn_ds_bpermute_fi_b32, "iii", "nc", "gfx12-insts")
+
//===----------------------------------------------------------------------===//
// WMMA builtins.
// Postfix w32 indicates the builtin requires wavefront size of 32.
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
index 08f70a25276f17..d518fe3a11a81c 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
@@ -2,7 +2,7 @@
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -verify
-emit-llvm -o - %s
-void test_s_sleep_var(int d)
-{
- __builtin_amdgcn_s_sleep_var(d); // expected-error
{{'__builtin_amdgcn_s_sleep_var' needs target feature gfx12-insts}}
+void builtin_test_unsupported(int a, int b) {
+ __builtin_amdgcn_s_sleep_var(a); // expected-error
{{'__builtin_amdgcn_s_sleep_var' needs target feature gfx12-insts}}
+ b = __builtin_amdgcn_ds_bpermute_fi_b32(a, b); // expected-error
{{'__builtin_amdgcn_ds_bpermute_fi_b32' needs target feature gfx12-insts}}
}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
index 5b5ae419f0a4a9..234ad4fd8cde61 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
@@ -296,3 +296,26 @@ void test_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t
rsrc, unsigned int len)
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 128, len);
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 0, 31);
}
+
+// CHECK-LABEL: @test_ds_bpermute_fi_b32(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[OUT_ADDR]] to ptr
+// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[A_ADDR]] to ptr
+// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[B_ADDR]] to ptr
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]],
align 8
+// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32
[[TMP0]], i32 [[TMP1]])
+// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
+// CHECK-NEXT: ret void
+//
+void test_ds_bpermute_fi_b32(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_ds_bpermute_fi_b32(a, b);
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index cc3584833202bf..f721d5267cd2a0 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2923,6 +2923,12 @@ def int_amdgcn_s_prefetch_data :
"", [SDNPMemOperand]
>;
+// llvm.amdgcn.ds.bpermute.fi.b32 <index> <src>
+def int_amdgcn_ds_bpermute_fi_b32 :
+ ClangBuiltin<"__builtin_amdgcn_ds_bpermute_fi_b32">,
+ Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
+ [IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
+
//===----------------------------------------------------------------------===//
// Deep learning intrinsics.
//===----------------------------------------------------------------------===//
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index 224c368cff4a1f..2e5f42c3bdc405 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4675,6 +4675,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const
MachineInstr &MI) const {
case Intrinsic::amdgcn_set_inactive:
case Intrinsic::amdgcn_set_inactive_chain_arg:
case Intrinsic::amdgcn_permlane64:
+ case Intrinsic::amdgcn_ds_bpermute_fi_b32:
return getDefaultMappingAllVGPR(MI);
case Intrinsic::amdgcn_cvt_pkrtz:
if (Subtarget.hasSALUFloatInsts() && isSALUMapping(MI))
diff --git a/llvm/lib/Target/AMDGPU/DSInstructions.td
b/llvm/lib/Target/AMDGPU/DSInstructions.td
index bc217e10e0fbd7..9ca853befba73c 100644
--- a/llvm/lib/Target/AMDGPU/DSInstructions.td
+++ b/llvm/lib/Target/AMDGPU/DSInstructions.td
@@ -699,7 +699,6 @@ def DS_PERMUTE_B32 : DS_1A1D_PERMUTE <"ds_permute_b32",
int_amdgcn_ds_permute>;
def DS_BPERMUTE_B32 : DS_1A1D_PERMUTE <"ds_bpermute_b32",
int_amdgcn_ds_bpermute>;
-def DS_BPERMUTE_FI_B32 : DS_1A1D_PERMUTE <"ds_bpermute_fi_b32">;
}
} // let SubtargetPredicate = isGFX8Plus
@@ -737,6 +736,8 @@ defm DS_COND_SUB_U32 :
DS_1A1D_NORET_mc<"ds_cond_sub_u32">;
defm DS_COND_SUB_RTN_U32 : DS_1A1D_RET_mc<"ds_cond_sub_rtn_u32", VGPR_32>;
defm DS_SUB_CLAMP_U32 : DS_1A1D_NORET_mc<"ds_sub_clamp_u32">;
defm DS_SUB_CLAMP_RTN_U32 : DS_1A1D_RET_mc<"ds_sub_clamp_rtn_u32", VGPR_32>;
+def DS_BPERMUTE_FI_B32 : DS_1A1D_PERMUTE <"ds_bpermute_fi_b32",
+ int_amdgcn_ds_bpermute_fi_b32>;
multiclass DSAtomicRetNoRetPatIntrinsic_mc<DS_Pseudo inst, DS_Pseudo noRetInst,
ValueType vt, string frag> {
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.fi.b32.ll
b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.fi.b32.ll
new file mode 100644
index 00000000000000..ea85055b65d127
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.fi.b32.ll
@@ -0,0 +1,154 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
UTC_ARGS: --version 5
+; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 -verify-machineinstrs
< %s | FileCheck -check-prefixes=GFX12,GFX12-SDAG %s
+; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 -verify-machineinstrs
< %s | FileCheck -check-prefixes=GFX12,GFX12-GISEL %s
+
+declare i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32, i32) #0
+
+define amdgpu_kernel void @ds_bpermute_fi_b32(ptr addrspace(1) %out, i32
%index, i32 %src) nounwind {
+; GFX12-LABEL: ds_bpermute_fi_b32:
+; GFX12: ; %bb.0:
+; GFX12-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
+; GFX12-NEXT: s_wait_kmcnt 0x0
+; GFX12-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3
+; GFX12-NEXT: ds_bpermute_fi_b32 v0, v0, v1
+; GFX12-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-NEXT: s_wait_dscnt 0x0
+; GFX12-NEXT: global_store_b32 v1, v0, s[0:1]
+; GFX12-NEXT: s_endpgm
+ %bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %index, i32 %src) #0
+ store i32 %bpermute, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_kernel void @ds_bpermute_fi_b32_imm_offset(ptr addrspace(1)
%out, i32 %base_index, i32 %src) nounwind {
+; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_imm_offset:
+; GFX12-SDAG: ; %bb.0:
+; GFX12-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
+; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-SDAG-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3
+; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v0, v0, v1 offset:4
+; GFX12-SDAG-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
+; GFX12-SDAG-NEXT: global_store_b32 v1, v0, s[0:1]
+; GFX12-SDAG-NEXT: s_endpgm
+;
+; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_imm_offset:
+; GFX12-GISEL: ; %bb.0:
+; GFX12-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
+; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-GISEL-NEXT: s_add_co_i32 s2, s2, 4
+; GFX12-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
+; GFX12-GISEL-NEXT: v_dual_mov_b32 v1, s3 :: v_dual_mov_b32 v0, s2
+; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v0, v0, v1
+; GFX12-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
+; GFX12-GISEL-NEXT: global_store_b32 v1, v0, s[0:1]
+; GFX12-GISEL-NEXT: s_endpgm
+ %index = add i32 %base_index, 4
+ %bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %index, i32 %src) #0
+ store i32 %bpermute, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define amdgpu_kernel void @ds_bpermute_fi_b32_imm_index(ptr addrspace(1) %out,
i32 %base_index, i32 %src) nounwind {
+; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_imm_index:
+; GFX12-SDAG: ; %bb.0:
+; GFX12-SDAG-NEXT: s_load_b32 s0, s[4:5], 0x30
+; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
+; GFX12-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24
+; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v1, v0, v1 offset:64
+; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
+; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-SDAG-NEXT: global_store_b32 v0, v1, s[0:1]
+; GFX12-SDAG-NEXT: s_endpgm
+;
+; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_imm_index:
+; GFX12-GISEL: ; %bb.0:
+; GFX12-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x30
+; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-GISEL-NEXT: v_dual_mov_b32 v0, 64 :: v_dual_mov_b32 v1, s0
+; GFX12-GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x24
+; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v0, v0, v1
+; GFX12-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
+; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-GISEL-NEXT: global_store_b32 v1, v0, s[0:1]
+; GFX12-GISEL-NEXT: s_endpgm
+ %bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 64, i32 %src) #0
+ store i32 %bpermute, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define void @ds_bpermute_fi_b32_add_shl(ptr addrspace(1) %out, i32
%base_index, i32 %src) nounwind {
+; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_add_shl:
+; GFX12-SDAG: ; %bb.0:
+; GFX12-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX12-SDAG-NEXT: s_wait_expcnt 0x0
+; GFX12-SDAG-NEXT: s_wait_samplecnt 0x0
+; GFX12-SDAG-NEXT: s_wait_bvhcnt 0x0
+; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-SDAG-NEXT: v_lshlrev_b32_e32 v2, 2, v2
+; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v2, v2, v3 offset:4
+; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
+; GFX12-SDAG-NEXT: global_store_b32 v[0:1], v2, off
+; GFX12-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_add_shl:
+; GFX12-GISEL: ; %bb.0:
+; GFX12-GISEL-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX12-GISEL-NEXT: s_wait_expcnt 0x0
+; GFX12-GISEL-NEXT: s_wait_samplecnt 0x0
+; GFX12-GISEL-NEXT: s_wait_bvhcnt 0x0
+; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-GISEL-NEXT: v_lshl_add_u32 v2, v2, 2, 4
+; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v2, v2, v3
+; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
+; GFX12-GISEL-NEXT: global_store_b32 v[0:1], v2, off
+; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %index = add i32 %base_index, 1
+ %byte_index = shl i32 %index, 2
+ %bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %byte_index, i32
%src) #0
+ store i32 %bpermute, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+define void @ds_bpermute_fi_b32_or_shl(ptr addrspace(1) %out, i32 %base_index,
i32 %src) nounwind {
+; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_or_shl:
+; GFX12-SDAG: ; %bb.0:
+; GFX12-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX12-SDAG-NEXT: s_wait_expcnt 0x0
+; GFX12-SDAG-NEXT: s_wait_samplecnt 0x0
+; GFX12-SDAG-NEXT: s_wait_bvhcnt 0x0
+; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
+; GFX12-SDAG-NEXT: v_and_b32_e32 v2, 62, v2
+; GFX12-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX12-SDAG-NEXT: v_lshlrev_b32_e32 v2, 2, v2
+; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v2, v2, v3 offset:4
+; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
+; GFX12-SDAG-NEXT: global_store_b32 v[0:1], v2, off
+; GFX12-SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_or_shl:
+; GFX12-GISEL: ; %bb.0:
+; GFX12-GISEL-NEXT: s_wait_loadcnt_dscnt 0x0
+; GFX12-GISEL-NEXT: s_wait_expcnt 0x0
+; GFX12-GISEL-NEXT: s_wait_samplecnt 0x0
+; GFX12-GISEL-NEXT: s_wait_bvhcnt 0x0
+; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
+; GFX12-GISEL-NEXT: v_and_b32_e32 v2, 62, v2
+; GFX12-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX12-GISEL-NEXT: v_lshl_or_b32 v2, v2, 2, 4
+; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v2, v2, v3
+; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
+; GFX12-GISEL-NEXT: global_store_b32 v[0:1], v2, off
+; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31]
+ %masked = and i32 %base_index, 62
+ %index = or i32 %masked, 1
+ %byte_index = shl i32 %index, 2
+ %bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %byte_index, i32
%src) #0
+ store i32 %bpermute, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+attributes #0 = { nounwind readnone convergent }
diff --git a/llvm/test/MC/AMDGPU/gfx11_unsupported.s
b/llvm/test/MC/AMDGPU/gfx11_unsupported.s
index 640ffb65a76bf5..d5d10ce548c4d0 100644
--- a/llvm/test/MC/AMDGPU/gfx11_unsupported.s
+++ b/llvm/test/MC/AMDGPU/gfx11_unsupported.s
@@ -2067,3 +2067,6 @@ ds_subrev_rtn_u64 v[5:6], v1, v[2:3]
v_dot2c_f32_bf16 v5, v1, v2
// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
+
+ds_bpermute_fi_b32 v5, v1, v2
+// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
