[llvm] [LLVM][NVPTX]: Add intrinsic for setmaxnreg (PR #77289)
Guray Ozen via llvm-commits
llvm-commits at lists.llvm.org
Mon Jan 8 09:09:01 PST 2024
================
@@ -4710,4 +4710,10 @@ def int_nvvm_is_explicit_cluster
[IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.is_explicit_cluster">;
+// Setmaxnreg intrinsic
+def int_nvvm_setmaxnreg_sync_aligned_u32
+ : DefaultAttrsIntrinsic<[], [llvm_i32_ty, llvm_i32_ty],
+ [IntrConvergent, IntrNoMem, IntrHasSideEffects, ImmArg<ArgIndex<1>>],
----------------
grypp wrote:
Does "IntrConvergent" means convergence within a warpgroup? This instruction helps establish smaller register sizes for first warpgroup and larger register sizes for second warpgroup. At least this is how we use for warp specilized GEMM kernels
https://github.com/llvm/llvm-project/pull/77289
More information about the llvm-commits
mailing list