[all-commits] [llvm/llvm-project] 340cc1: [LLVM][NVPTX]: Add intrinsic for setmaxnreg (#77289)

[Durgadoss R via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 340cc1702e21128b62799c5dfbf2875c3c2c96a1
      https://github.com/llvm/llvm-project/commit/340cc1702e21128b62799c5dfbf2875c3c2c96a1
  Author: Durgadoss R <durgadossr at nvidia.com>
  Date:   2024-01-09 (Tue, 09 Jan 2024)

  Changed paths:
    M llvm/include/llvm/IR/IntrinsicsNVVM.td
    M llvm/lib/IR/Verifier.cpp
    M llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
    M llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
    A llvm/test/CodeGen/NVPTX/setmaxnreg.ll
    A llvm/test/Verifier/NVPTX/lit.local.cfg
    A llvm/test/Verifier/NVPTX/setmaxnreg.ll

  Log Message:
  -----------
  [LLVM][NVPTX]: Add intrinsic for setmaxnreg (#77289)

This patch adds an intrinsic for setmaxnreg PTX instruction.
* PTX Doc link for this instruction:
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#miscellaneous-instructions-setmaxnreg

* The i32 argument, an immediate value, specifies the actual
  absolute register count for the instruction.
* The `setmaxnreg` instruction is available in SM90a.
  So, this patch adds 'hasSM90a' predicate to use in
  the NVPTX backend.
* lit tests are added to verify the lowering of the intrinsic.
* Verifier logic (and tests) are added to test the register
  count range and divisibility-by-8 requirements.

Signed-off-by: Durgadoss R <durgadossr at nvidia.com>