[llvm] a2e2038 - [AMDGPU] Add missing attributes to llvm.amdgcn.init.whole.wave (#222493)

[via llvm-commits]

Author: Shilei Tian
Date: 2026-09-10T03:05:07Z
New Revision: a2e2038fc932039ee5e101432b4a549c92fdf9c8

URL: https://github.com/llvm/llvm-project/commit/a2e2038fc932039ee5e101432b4a549c92fdf9c8
DIFF: https://github.com/llvm/llvm-project/commit/a2e2038fc932039ee5e101432b4a549c92fdf9c8.diff

LOG: [AMDGPU] Add missing attributes to llvm.amdgcn.init.whole.wave (#222493)

`llvm.amdgcn.init.whole.wave` only sets the EXEC mask and reports
whether the lane was active on entry. It never calls into the module,
never frees memory, and always returns, so it can carry `nocallback`,
`nofree`, and `willreturn`.

The neighboring `llvm.amdgcn.init.exec` and
`llvm.amdgcn.init.exec.from.input` already have all three; this brings
`init.whole.wave` in line and lets the attributor propagate them to
callers.

Added: 
    

Modified: 
    llvm/include/llvm/IR/IntrinsicsAMDGPU.td

Removed: 
    


################################################################################
diff  --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 9aa4eca59b96a..6b29b8612828b 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -247,7 +247,7 @@ def int_amdgcn_init_exec_from_input : Intrinsic<[],
 // (mostly related to WWM CSR handling) that 
diff erentiate it from using
 // a plain `amdgcn.init.exec -1`.
 def int_amdgcn_init_whole_wave : Intrinsic<[llvm_i1_ty], [], [
-    IntrHasSideEffects, IntrNoMem, IntrConvergent]>;
+    IntrHasSideEffects, IntrNoMem, IntrConvergent, IntrNoCallback, IntrNoFree, IntrWillReturn]>;
 
 def int_amdgcn_wavefrontsize
     : ClangBuiltin<"__builtin_amdgcn_wavefrontsize">,