Skip to content

Commit 160431e

Browse files
Update table gen and use one test case in _alias.s
1 parent c120faa commit 160431e

File tree

2 files changed

+8
-11
lines changed

2 files changed

+8
-11
lines changed

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -2726,8 +2726,8 @@ class IntDSBVHStackRtn<LLVMType vdst, LLVMType data1> :
27262726
[ImmArg<ArgIndex<3>>, IntrWillReturn, IntrNoCallback, IntrNoFree]
27272727
>;
27282728

2729-
def int_amdgcn_ds_bvh_stack_rtn : IntDSBVHStackRtn</* vdst = */ llvm_i32_ty,
2730-
/* data1 = */ llvm_v4i32_ty>;
2729+
def int_amdgcn_ds_bvh_stack_rtn : IntDSBVHStackRtn<vdst = llvm_i32_ty,
2730+
data1 = llvm_v4i32_ty>;
27312731

27322732
def int_amdgcn_s_wait_event_export_ready :
27332733
ClangBuiltin<"__builtin_amdgcn_s_wait_event_export_ready">,
@@ -2804,14 +2804,14 @@ def int_amdgcn_wmma_bf16_16x16x16_bf16 : AMDGPUWmmaIntrinsicOPSEL<llvm_anyint_ty
28042804
// GFX12 Intrinsics
28052805
//===----------------------------------------------------------------------===//
28062806

2807-
def int_amdgcn_ds_bvh_stack_push4_pop1_rtn : IntDSBVHStackRtn</* vdst = */ llvm_i32_ty,
2808-
/* data1 = */ llvm_v4i32_ty>;
2807+
def int_amdgcn_ds_bvh_stack_push4_pop1_rtn : IntDSBVHStackRtn<vdst = llvm_i32_ty,
2808+
data1 = llvm_v4i32_ty>;
28092809

2810-
def int_amdgcn_ds_bvh_stack_push8_pop1_rtn : IntDSBVHStackRtn</* vdst = */ llvm_i32_ty,
2811-
/* data1 = */ llvm_v8i32_ty>;
2810+
def int_amdgcn_ds_bvh_stack_push8_pop1_rtn : IntDSBVHStackRtn<vdst = llvm_i32_ty,
2811+
data1 = llvm_v8i32_ty>;
28122812

2813-
def int_amdgcn_ds_bvh_stack_push8_pop2_rtn : IntDSBVHStackRtn</* vdst = */ llvm_i64_ty,
2814-
/* data1 = */ llvm_v8i32_ty>;
2813+
def int_amdgcn_ds_bvh_stack_push8_pop2_rtn : IntDSBVHStackRtn<vdst = llvm_i64_ty,
2814+
data1 = llvm_v8i32_ty>;
28152815

28162816
// llvm.amdgcn.permlane16.var <old> <src0> <src1> <fi> <bound_control>
28172817
def int_amdgcn_permlane16_var : ClangBuiltin<"__builtin_amdgcn_permlane16_var">,

llvm/test/MC/AMDGPU/gfx12_asm_ds_alias.s

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

3939
ds_bvh_stack_rtn_b32 v1, v0, v1, v[2:5]
4040
// GFX12: ds_bvh_stack_push4_pop1_rtn_b32 v1, v0, v1, v[2:5] ; encoding: [0x00,0x00,0x80,0xdb,0x00,0x01,0x02,0x01]
41-
42-
ds_bvh_stack_rtn_b32 v1, v0, v1, v[2:5] offset:1
43-
// GFX12: ds_bvh_stack_push4_pop1_rtn_b32 v1, v0, v1, v[2:5] offset:1 ; encoding: [0x01,0x00,0x80,0xdb,0x00,0x01,0x02,0x01]

0 commit comments

Comments
 (0)