[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