[llvm] [LLVM][NVPTX]: Add intrinsic for setmaxnreg (PR #77289)

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Tue Jan 9 11:05:45 PST 2024


================
@@ -6031,6 +6032,16 @@ void Verifier::visitIntrinsicCall(Intrinsic::ID ID, CallBase &Call) {
           "Value for inactive lanes must be a VGPR function argument", &Call);
     break;
   }
+  case Intrinsic::nvvm_setmaxnreg_inc_sync_aligned_u32:
+  case Intrinsic::nvvm_setmaxnreg_dec_sync_aligned_u32: {
+    Value *V = Call.getArgOperand(0);
+    unsigned RegCount = cast<ConstantInt>(V)->getZExtValue();
+    Check(RegCount % 8 == 0,
+          "reg_count argument to nvvm.setmaxnreg must be divisible by 8");
----------------
Artem-B wrote:

Nit. "must be in multiples of 8" may work a bit better.

After all all integers are divisible by 8, it's just the remainder for some is non-zero. :-)

https://github.com/llvm/llvm-project/pull/77289


More information about the llvm-commits mailing list