Skip to content

Commit 3a29dfe

Browse files
Acim-MaravicAcim-Maravic
authored
[LLVM][AMDGPU] Add Intrinsic and Builtin for ds_bpermute_fi_b32 (#124616)
1 parent 978e083 commit 3a29dfe

File tree

8 files changed

+194
-4
lines changed

8 files changed

+194
-4
lines changed

‎clang/include/clang/Basic/BuiltinsAMDGPU.def

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -504,6 +504,8 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4i16, "V4sV4s*1", "nc", "gf
504504
TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4f16, "V4hV4h*1", "nc", "gfx12-insts,wavefrontsize64")
505505
TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4bf16, "V4yV4y*1", "nc", "gfx12-insts,wavefrontsize64")
506506

507+
TARGET_BUILTIN(__builtin_amdgcn_ds_bpermute_fi_b32, "iii", "nc", "gfx12-insts")
508+
507509
//===----------------------------------------------------------------------===//
508510
// WMMA builtins.
509511
// Postfix w32 indicates the builtin requires wavefront size of 32.

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22

33
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -verify -emit-llvm -o - %s
44

5-
void test_s_sleep_var(int d)
6-
{
7-
__builtin_amdgcn_s_sleep_var(d); // expected-error {{'__builtin_amdgcn_s_sleep_var' needs target feature gfx12-insts}}
5+
void builtin_test_unsupported(int a, int b) {
6+
__builtin_amdgcn_s_sleep_var(a); // expected-error {{'__builtin_amdgcn_s_sleep_var' needs target feature gfx12-insts}}
7+
b = __builtin_amdgcn_ds_bpermute_fi_b32(a, b); // expected-error {{'__builtin_amdgcn_ds_bpermute_fi_b32' needs target feature gfx12-insts}}
88
}

‎clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -296,3 +296,26 @@ void test_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t rsrc, unsigned int len)
296296
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 128, len);
297297
__builtin_amdgcn_s_buffer_prefetch_data(rsrc, 0, 31);
298298
}
299+
300+
// CHECK-LABEL: @test_ds_bpermute_fi_b32(
301+
// CHECK-NEXT: entry:
302+
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8, addrspace(5)
303+
// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
304+
// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
305+
// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
306+
// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
307+
// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
308+
// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]], align 8
309+
// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
310+
// CHECK-NEXT: store i32 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 4
311+
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
312+
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[B_ADDR_ASCAST]], align 4
313+
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 [[TMP0]], i32 [[TMP1]])
314+
// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(1), ptr [[OUT_ADDR_ASCAST]], align 8
315+
// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[TMP3]], align 4
316+
// CHECK-NEXT: ret void
317+
//
318+
void test_ds_bpermute_fi_b32(global int* out, int a, int b)
319+
{
320+
*out = __builtin_amdgcn_ds_bpermute_fi_b32(a, b);
321+
}

‎llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2923,6 +2923,12 @@ def int_amdgcn_s_prefetch_data :
29232923
"", [SDNPMemOperand]
29242924
>;
29252925

2926+
// llvm.amdgcn.ds.bpermute.fi.b32 <index> <src>
2927+
def int_amdgcn_ds_bpermute_fi_b32 :
2928+
ClangBuiltin<"__builtin_amdgcn_ds_bpermute_fi_b32">,
2929+
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
2930+
[IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
2931+
29262932
//===----------------------------------------------------------------------===//
29272933
// Deep learning intrinsics.
29282934
//===----------------------------------------------------------------------===//

‎llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4675,6 +4675,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const {
46754675
case Intrinsic::amdgcn_set_inactive:
46764676
case Intrinsic::amdgcn_set_inactive_chain_arg:
46774677
case Intrinsic::amdgcn_permlane64:
4678+
case Intrinsic::amdgcn_ds_bpermute_fi_b32:
46784679
return getDefaultMappingAllVGPR(MI);
46794680
case Intrinsic::amdgcn_cvt_pkrtz:
46804681
if (Subtarget.hasSALUFloatInsts() && isSALUMapping(MI))

‎llvm/lib/Target/AMDGPU/DSInstructions.td

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -699,7 +699,6 @@ def DS_PERMUTE_B32 : DS_1A1D_PERMUTE <"ds_permute_b32",
699699
int_amdgcn_ds_permute>;
700700
def DS_BPERMUTE_B32 : DS_1A1D_PERMUTE <"ds_bpermute_b32",
701701
int_amdgcn_ds_bpermute>;
702-
def DS_BPERMUTE_FI_B32 : DS_1A1D_PERMUTE <"ds_bpermute_fi_b32">;
703702
}
704703

705704
} // let SubtargetPredicate = isGFX8Plus
@@ -737,6 +736,8 @@ defm DS_COND_SUB_U32 : DS_1A1D_NORET_mc<"ds_cond_sub_u32">;
737736
defm DS_COND_SUB_RTN_U32 : DS_1A1D_RET_mc<"ds_cond_sub_rtn_u32", VGPR_32>;
738737
defm DS_SUB_CLAMP_U32 : DS_1A1D_NORET_mc<"ds_sub_clamp_u32">;
739738
defm DS_SUB_CLAMP_RTN_U32 : DS_1A1D_RET_mc<"ds_sub_clamp_rtn_u32", VGPR_32>;
739+
def DS_BPERMUTE_FI_B32 : DS_1A1D_PERMUTE <"ds_bpermute_fi_b32",
740+
int_amdgcn_ds_bpermute_fi_b32>;
740741

741742
multiclass DSAtomicRetNoRetPatIntrinsic_mc<DS_Pseudo inst, DS_Pseudo noRetInst,
742743
ValueType vt, string frag> {

‎llvm/test/CodeGen/AMDGPU/llvm.amdgcn.ds.bpermute.fi.b32.ll

Lines changed: 154 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,154 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc -global-isel=0 -mtriple=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX12,GFX12-SDAG %s
3+
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1200 -verify-machineinstrs < %s | FileCheck -check-prefixes=GFX12,GFX12-GISEL %s
4+
5+
declare i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32, i32) #0
6+
7+
define amdgpu_kernel void @ds_bpermute_fi_b32(ptr addrspace(1) %out, i32 %index, i32 %src) nounwind {
8+
; GFX12-LABEL: ds_bpermute_fi_b32:
9+
; GFX12: ; %bb.0:
10+
; GFX12-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
11+
; GFX12-NEXT: s_wait_kmcnt 0x0
12+
; GFX12-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3
13+
; GFX12-NEXT: ds_bpermute_fi_b32 v0, v0, v1
14+
; GFX12-NEXT: v_mov_b32_e32 v1, 0
15+
; GFX12-NEXT: s_wait_dscnt 0x0
16+
; GFX12-NEXT: global_store_b32 v1, v0, s[0:1]
17+
; GFX12-NEXT: s_endpgm
18+
%bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %index, i32 %src) #0
19+
store i32 %bpermute, ptr addrspace(1) %out, align 4
20+
ret void
21+
}
22+
23+
define amdgpu_kernel void @ds_bpermute_fi_b32_imm_offset(ptr addrspace(1) %out, i32 %base_index, i32 %src) nounwind {
24+
; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_imm_offset:
25+
; GFX12-SDAG: ; %bb.0:
26+
; GFX12-SDAG-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
27+
; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
28+
; GFX12-SDAG-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3
29+
; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v0, v0, v1 offset:4
30+
; GFX12-SDAG-NEXT: v_mov_b32_e32 v1, 0
31+
; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
32+
; GFX12-SDAG-NEXT: global_store_b32 v1, v0, s[0:1]
33+
; GFX12-SDAG-NEXT: s_endpgm
34+
;
35+
; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_imm_offset:
36+
; GFX12-GISEL: ; %bb.0:
37+
; GFX12-GISEL-NEXT: s_load_b128 s[0:3], s[4:5], 0x24
38+
; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
39+
; GFX12-GISEL-NEXT: s_add_co_i32 s2, s2, 4
40+
; GFX12-GISEL-NEXT: s_delay_alu instid0(SALU_CYCLE_1)
41+
; GFX12-GISEL-NEXT: v_dual_mov_b32 v1, s3 :: v_dual_mov_b32 v0, s2
42+
; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v0, v0, v1
43+
; GFX12-GISEL-NEXT: v_mov_b32_e32 v1, 0
44+
; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
45+
; GFX12-GISEL-NEXT: global_store_b32 v1, v0, s[0:1]
46+
; GFX12-GISEL-NEXT: s_endpgm
47+
%index = add i32 %base_index, 4
48+
%bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %index, i32 %src) #0
49+
store i32 %bpermute, ptr addrspace(1) %out, align 4
50+
ret void
51+
}
52+
53+
define amdgpu_kernel void @ds_bpermute_fi_b32_imm_index(ptr addrspace(1) %out, i32 %base_index, i32 %src) nounwind {
54+
; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_imm_index:
55+
; GFX12-SDAG: ; %bb.0:
56+
; GFX12-SDAG-NEXT: s_load_b32 s0, s[4:5], 0x30
57+
; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
58+
; GFX12-SDAG-NEXT: v_dual_mov_b32 v0, 0 :: v_dual_mov_b32 v1, s0
59+
; GFX12-SDAG-NEXT: s_load_b64 s[0:1], s[4:5], 0x24
60+
; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v1, v0, v1 offset:64
61+
; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
62+
; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
63+
; GFX12-SDAG-NEXT: global_store_b32 v0, v1, s[0:1]
64+
; GFX12-SDAG-NEXT: s_endpgm
65+
;
66+
; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_imm_index:
67+
; GFX12-GISEL: ; %bb.0:
68+
; GFX12-GISEL-NEXT: s_load_b32 s0, s[4:5], 0x30
69+
; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
70+
; GFX12-GISEL-NEXT: v_dual_mov_b32 v0, 64 :: v_dual_mov_b32 v1, s0
71+
; GFX12-GISEL-NEXT: s_load_b64 s[0:1], s[4:5], 0x24
72+
; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v0, v0, v1
73+
; GFX12-GISEL-NEXT: v_mov_b32_e32 v1, 0
74+
; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
75+
; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
76+
; GFX12-GISEL-NEXT: global_store_b32 v1, v0, s[0:1]
77+
; GFX12-GISEL-NEXT: s_endpgm
78+
%bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 64, i32 %src) #0
79+
store i32 %bpermute, ptr addrspace(1) %out, align 4
80+
ret void
81+
}
82+
83+
define void @ds_bpermute_fi_b32_add_shl(ptr addrspace(1) %out, i32 %base_index, i32 %src) nounwind {
84+
; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_add_shl:
85+
; GFX12-SDAG: ; %bb.0:
86+
; GFX12-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0
87+
; GFX12-SDAG-NEXT: s_wait_expcnt 0x0
88+
; GFX12-SDAG-NEXT: s_wait_samplecnt 0x0
89+
; GFX12-SDAG-NEXT: s_wait_bvhcnt 0x0
90+
; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
91+
; GFX12-SDAG-NEXT: v_lshlrev_b32_e32 v2, 2, v2
92+
; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v2, v2, v3 offset:4
93+
; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
94+
; GFX12-SDAG-NEXT: global_store_b32 v[0:1], v2, off
95+
; GFX12-SDAG-NEXT: s_setpc_b64 s[30:31]
96+
;
97+
; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_add_shl:
98+
; GFX12-GISEL: ; %bb.0:
99+
; GFX12-GISEL-NEXT: s_wait_loadcnt_dscnt 0x0
100+
; GFX12-GISEL-NEXT: s_wait_expcnt 0x0
101+
; GFX12-GISEL-NEXT: s_wait_samplecnt 0x0
102+
; GFX12-GISEL-NEXT: s_wait_bvhcnt 0x0
103+
; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
104+
; GFX12-GISEL-NEXT: v_lshl_add_u32 v2, v2, 2, 4
105+
; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v2, v2, v3
106+
; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
107+
; GFX12-GISEL-NEXT: global_store_b32 v[0:1], v2, off
108+
; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31]
109+
%index = add i32 %base_index, 1
110+
%byte_index = shl i32 %index, 2
111+
%bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %byte_index, i32 %src) #0
112+
store i32 %bpermute, ptr addrspace(1) %out, align 4
113+
ret void
114+
}
115+
116+
define void @ds_bpermute_fi_b32_or_shl(ptr addrspace(1) %out, i32 %base_index, i32 %src) nounwind {
117+
; GFX12-SDAG-LABEL: ds_bpermute_fi_b32_or_shl:
118+
; GFX12-SDAG: ; %bb.0:
119+
; GFX12-SDAG-NEXT: s_wait_loadcnt_dscnt 0x0
120+
; GFX12-SDAG-NEXT: s_wait_expcnt 0x0
121+
; GFX12-SDAG-NEXT: s_wait_samplecnt 0x0
122+
; GFX12-SDAG-NEXT: s_wait_bvhcnt 0x0
123+
; GFX12-SDAG-NEXT: s_wait_kmcnt 0x0
124+
; GFX12-SDAG-NEXT: v_and_b32_e32 v2, 62, v2
125+
; GFX12-SDAG-NEXT: s_delay_alu instid0(VALU_DEP_1)
126+
; GFX12-SDAG-NEXT: v_lshlrev_b32_e32 v2, 2, v2
127+
; GFX12-SDAG-NEXT: ds_bpermute_fi_b32 v2, v2, v3 offset:4
128+
; GFX12-SDAG-NEXT: s_wait_dscnt 0x0
129+
; GFX12-SDAG-NEXT: global_store_b32 v[0:1], v2, off
130+
; GFX12-SDAG-NEXT: s_setpc_b64 s[30:31]
131+
;
132+
; GFX12-GISEL-LABEL: ds_bpermute_fi_b32_or_shl:
133+
; GFX12-GISEL: ; %bb.0:
134+
; GFX12-GISEL-NEXT: s_wait_loadcnt_dscnt 0x0
135+
; GFX12-GISEL-NEXT: s_wait_expcnt 0x0
136+
; GFX12-GISEL-NEXT: s_wait_samplecnt 0x0
137+
; GFX12-GISEL-NEXT: s_wait_bvhcnt 0x0
138+
; GFX12-GISEL-NEXT: s_wait_kmcnt 0x0
139+
; GFX12-GISEL-NEXT: v_and_b32_e32 v2, 62, v2
140+
; GFX12-GISEL-NEXT: s_delay_alu instid0(VALU_DEP_1)
141+
; GFX12-GISEL-NEXT: v_lshl_or_b32 v2, v2, 2, 4
142+
; GFX12-GISEL-NEXT: ds_bpermute_fi_b32 v2, v2, v3
143+
; GFX12-GISEL-NEXT: s_wait_dscnt 0x0
144+
; GFX12-GISEL-NEXT: global_store_b32 v[0:1], v2, off
145+
; GFX12-GISEL-NEXT: s_setpc_b64 s[30:31]
146+
%masked = and i32 %base_index, 62
147+
%index = or i32 %masked, 1
148+
%byte_index = shl i32 %index, 2
149+
%bpermute = call i32 @llvm.amdgcn.ds.bpermute.fi.b32(i32 %byte_index, i32 %src) #0
150+
store i32 %bpermute, ptr addrspace(1) %out, align 4
151+
ret void
152+
}
153+
154+
attributes #0 = { nounwind readnone convergent }

‎llvm/test/MC/AMDGPU/gfx11_unsupported.s

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2067,3 +2067,6 @@ ds_subrev_rtn_u64 v[5:6], v1, v[2:3]
20672067

20682068
v_dot2c_f32_bf16 v5, v1, v2
20692069
// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU
2070+
2071+
ds_bpermute_fi_b32 v5, v1, v2
2072+
// CHECK: :[[@LINE-1]]:{{[0-9]+}}: error: instruction not supported on this GPU

0 commit comments

Comments
 (0)