Skip to content

AMDGPU: Add noundef to mbcnt intrinsic returns #136304

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 2 commits into from
Closed

AMDGPU: Add noundef to mbcnt intrinsic returns #136304

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
4022bab
AMDGPU: Add range to wavefrontsize intrinsic declaration
arsenm Apr 18, 2025
050b743
AMDGPU: Add noundef to mbcnt intrinsic returns
arsenm Apr 18, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 13 additions & 11 deletions llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -234,9 +234,11 @@ def int_amdgcn_init_exec_from_input : Intrinsic<[],
def int_amdgcn_init_whole_wave : Intrinsic<[llvm_i1_ty], [], [
IntrHasSideEffects, IntrNoMem, IntrConvergent]>;

def int_amdgcn_wavefrontsize :
ClangBuiltin<"__builtin_amdgcn_wavefrontsize">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>, IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_wavefrontsize
: ClangBuiltin<"__builtin_amdgcn_wavefrontsize">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [],
[NoUndef<RetIndex>, Range<RetIndex, 32, 65>,
IntrNoMem, IntrSpeculatable]>;

// Represent a relocation constant.
def int_amdgcn_reloc_constant : DefaultAttrsIntrinsic<
Expand Down Expand Up @@ -2213,15 +2215,15 @@ def int_amdgcn_live_mask : DefaultAttrsIntrinsic <[llvm_i1_ty],
[], [NoUndef<RetIndex>, IntrReadMem, IntrInaccessibleMemOnly]
>;

def int_amdgcn_mbcnt_lo :
ClangBuiltin<"__builtin_amdgcn_mbcnt_lo">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem]>;
def int_amdgcn_mbcnt_lo
: ClangBuiltin<"__builtin_amdgcn_mbcnt_lo">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[NoUndef<RetIndex>, IntrNoMem]>;

def int_amdgcn_mbcnt_hi :
ClangBuiltin<"__builtin_amdgcn_mbcnt_hi">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[IntrNoMem]>;
def int_amdgcn_mbcnt_hi
: ClangBuiltin<"__builtin_amdgcn_mbcnt_hi">,
DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
[NoUndef<RetIndex>, IntrNoMem]>;

// llvm.amdgcn.ds.swizzle src offset
def int_amdgcn_ds_swizzle :
Expand Down
21 changes: 21 additions & 0 deletions llvm/test/Assembler/amdgcn-intrinsic-attributes.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,4 +18,25 @@ define i32 @ds_consume(ptr addrspace(3) %ptr) {
ret i32 %ret
}

; CHECK: declare noundef i32 @llvm.amdgcn.mbcnt.hi(i32, i32) #1
define i32 @mbcnt_hi(i32 %a, i32 %b) {
%ret = call i32 @llvm.amdgcn.mbcnt.hi(i32 %a, i32 %b)
ret i32 %ret
}

; CHECK: declare noundef i32 @llvm.amdgcn.mbcnt.lo(i32, i32) #1
define i32 @mbcnt_lo(i32 %a, i32 %b) {
%ret = call i32 @llvm.amdgcn.mbcnt.lo(i32 %a, i32 %b)
ret i32 %ret
}

; Test assumed range
; CHECK: declare noundef range(i32 32, 65) i32 @llvm.amdgcn.wavefrontsize() #2
define i32 @wavefrontsize() {
%ret = call i32 @llvm.amdgcn.wavefrontsize()
ret i32 %ret
}

; CHECK: attributes #0 = { convergent nocallback nofree nounwind willreturn memory(argmem: readwrite) }
; CHECK: attributes #1 = { nocallback nofree nosync nounwind willreturn memory(none) }
; CHECK: attributes #2 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
4 changes: 2 additions & 2 deletions llvm/test/Transforms/InstCombine/AMDGPU/llvm.amdgcn.wavefrontsize.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ define amdgpu_kernel void @fold_and_optimize_wavefrontsize(ptr addrspace(1) noca
; OPT-SAME: ptr addrspace(1) captures(none) [[ARG:%.*]]) {
; OPT-NEXT: [[BB:.*:]]
; OPT-NEXT: [[TMP:%.*]] = tail call i32 @llvm.amdgcn.wavefrontsize() #[[ATTR1]]
; OPT-NEXT: [[TMP1:%.*]] = icmp ugt i32 [[TMP]], 32
; OPT-NEXT: [[TMP1:%.*]] = icmp samesign ugt i32 [[TMP]], 32
; OPT-NEXT: [[TMP2:%.*]] = select i1 [[TMP1]], i32 2, i32 1
; OPT-NEXT: store i32 [[TMP2]], ptr addrspace(1) [[ARG]], align 4
; OPT-NEXT: ret void
Expand Down Expand Up @@ -69,7 +69,7 @@ define amdgpu_kernel void @fold_and_optimize_if_wavefrontsize(ptr addrspace(1) n
; OPT-SAME: ptr addrspace(1) captures(none) [[ARG:%.*]]) {
; OPT-NEXT: [[BB:.*:]]
; OPT-NEXT: [[TMP:%.*]] = tail call i32 @llvm.amdgcn.wavefrontsize() #[[ATTR1]]
; OPT-NEXT: [[TMP1:%.*]] = icmp ugt i32 [[TMP]], 32
; OPT-NEXT: [[TMP1:%.*]] = icmp samesign ugt i32 [[TMP]], 32
; OPT-NEXT: br i1 [[TMP1]], label %[[BB2:.*]], label %[[BB3:.*]]
; OPT: [[BB2]]:
; OPT-NEXT: store i32 1, ptr addrspace(1) [[ARG]], align 4
Expand Down
Loading