//===- IntrinsicsNVVM.td - Defines NVVM intrinsics ---------*- tablegen -*-===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
// This file defines all of the NVVM-specific intrinsics for use with NVPTX.
//
//===----------------------------------------------------------------------===//

//===----------------------------------------------------------------------===//
// Guidelines on NVPTX Intrinsic design
//===----------------------------------------------------------------------===//
//
// The NVPTX intrinsics are used to model instructions in the PTX ISA.
// While simpler intrinsics can represent certain features effectively,
// more complex instructions like TMA and MMA are not as straightforward
// to model. A single variant of these complex instructions can expand
// into hundreds of intrinsics. Additionally, any expansion in the
// corresponding ISA can exponentially increase these numbers, making it
// difficult to manage them in the IR and backend passes. Therefore,
// a careful design of intrinsic interfaces can ease maintenance and
// contribute to a sustainable, long-term solution.
//
// The default approach is to have a 1:1 match between the intrinsic and
// the instruction where the instruction suffixes map to the intrinsic name
// and the instruction arguments map to the intrinsic arguments or return
// value.
//
// However, when there are too many instruction/intrinsic variants like
// the TMA/MMA family, it is desirable to encode some variants as a
// constant argument, referred to as 'flags'.
// TODO: Add a guideline to quantify the metric on 'how many intrinsics' here.
//
// Below are a set of guidelines that may help in choosing
// an appropriate design for the complex intrinsics:
//
// 1. Each flag argument represents one set of instruction modifiers.
//    These flags are compile-time integer constants.
//
// 2. When an intrinsic uses flags, document it with details of the
//    flag usage in the ``NVPTXUsage.rst`` file.
// 3. Annotate all flag arguments with ImmArg<ArgIdx<>>.
// 4. Place the flag arguments at the end of the (actual)argument list.
//
// 5. Use `i1` for boolean flags and `i8` for others. Usually,
//    the `i8` types represent an `enum` encoding the family of
//    modifiers.
// 6. Note that, the specific variant for non-boolean flags may not be
//    obvious in the IR. So, maintain consistency between the enum value
//    definitions and their usage in the backend.
//    * Provide a meaningful default value in the enums wherever applicable.
//    * TODO: Investigate auto-upgrade capability for intrinsics
//      when only flag value mappings change.
//
// 7. Identify the key features of an intrinsic and distinguish between
//    first-order and supplementary information. Typically, encoding the
//    first-order information in the intrinsic name while using flags
//    for supplementary details improves readability.
//    For example:
//
//    i. For MMA intrinsics, 'dense' vs. 'sparse' is a fundamental feature,
//    whereas an optional scaling applied to matrices is relatively secondary.
//
//    ii. For TMAs, the mode of copy (e.g., 'Tile' or 'Im2col') is a first-order
//    information, while features like an optional cache hint tend to be
//    secondary.
//
// 8. If there are invalid combinations within a set of modifiers, avoid
//    encoding them as flags, as much as possible. This helps reduce the
//    need for error handling of unsupported cases in the backend.
//    For example, some 'cvt' intrinsics support only a subset of the
//    possible rounding modes; so it is preferable not to encode the
//    rounding modes as flags.
// 9. Similarly, when there are invalid combinations across a set of
//    modifiers, avoid encoding them as flags to prevent additional
//    complexity in error handling.
//
// 10. Maintain a consistent design within an intrinsic family, including
//     argument ordering as well as the usage and ordering of flags.
// 11. When designing an intrinsic corresponding to an instruction or its variant,
//     consider the entire instruction family. This may reveal common features
//     that can be modelled consistently across the family.
//
// In summary, strive to balance the aspects mentioned above, to achieve
// a scalable design with maximum readability.
//===----------------------------------------------------------------------===//

// The following intrinsics were once defined here, but are now auto-upgraded
// to target-generic LLVM intrinsics.
//
//   * llvm.nvvm.brev32              --> llvm.bitreverse.i32
//   * llvm.nvvm.brev64              --> llvm.bitreverse.i64
//   * llvm.nvvm.clz.i               --> llvm.ctlz.i32
//   * llvm.nvvm.clz.ll              --> trunc i64 llvm.ctlz.i64(x) to i32
//   * llvm.nvvm.popc.i              --> llvm.ctpop.i32
//   * llvm.nvvm.popc.ll             --> trunc i64 llvm.ctpop.i64 to i32
//   * llvm.nvvm.abs.i               --> llvm.abs.i32(x, true)
//   * llvm.nvvm.abs.ll              --> llvm.abs.i64(x, true)
//   * llvm.nvvm.max.i               --> select(x sge y, x, y)
//   * llvm.nvvm.max.ll              --> ibid.
//   * llvm.nvvm.max.ui              --> select(x uge y, x, y)
//   * llvm.nvvm.max.ull             --> ibid.
//   * llvm.nvvm.max.i               --> select(x sle y, x, y)
//   * llvm.nvvm.max.ll              --> ibid.
//   * llvm.nvvm.max.ui              --> select(x ule y, x, y)
//   * llvm.nvvm.max.ull             --> ibid.
//   * llvm.nvvm.h2f                 --> llvm.convert.to.fp16.f32
//   * llvm.nvvm.bitcast.f2i         --> bitcast
//   * llvm.nvvm.bitcast.i2f         --> ibid.
//   * llvm.nvvm.bitcast.d2ll        --> ibid.
//   * llvm.nvvm.bitcast.ll2d        --> ibid.
//   * llvm.nvvm.ptr.gen.to.global   --> addrspacecast
//   * llvm.nvvm.ptr.gen.to.shared   --> ibid.
//   * llvm.nvvm.ptr.gen.to.constant --> ibid.
//   * llvm.nvvm.ptr.gen.to.local    --> ibid.
//   * llvm.nvvm.ptr.gen.to.param    --> ibid.
//   * llvm.nvvm.ptr.global.to.gen   --> ibid.
//   * llvm.nvvm.ptr.shared.to.gen   --> ibid.
//   * llvm.nvvm.ptr.constant.to.gen --> ibid.
//   * llvm.nvvm.ptr.local.to.gen    --> ibid.
//   * llvm.nvvm.ptr.param.to.gen    --> ibid.
//   * llvm.nvvm.ldg.global.i        --> load addrspace(1) !load.invariant
//   * llvm.nvvm.ldg.global.f        --> ibid.
//   * llvm.nvvm.ldg.global.p        --> ibid.
//   * llvm.nvvm.swap.lo.hi.b64      --> llvm.fshl(x, x, 32)
//   * llvm.nvvm.atomic.load.inc.32  --> atomicrmw uinc_wrap
//   * llvm.nvvm.atomic.load.dec.32  --> atomicrmw udec_wrap
//   * llvm.nvvm.barrier0            --> llvm.nvvm.barrier.cta.sync.aligned.all(0)
//   * llvm.nvvm.barrier.n           --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
//   * llvm.nvvm.bar.sync            --> llvm.nvvm.barrier.cta.sync.aligned.all(x)
//   * llvm.nvvm.barrier             --> llvm.nvvm.barrier.cta.sync.aligned(x, y)
//   * llvm.nvvm.barrier.sync        --> llvm.nvvm.barrier.cta.sync.all(x)
//   * llvm.nvvm.barrier.sync.cnt    --> llvm.nvvm.barrier.cta.sync(x, y)
//   * llvm.nvvm.barrier0.popc       --> llvm.nvvm.barrier.cta.red.popc.aligned.all(0, c)
//   * llvm.nvvm.barrier0.and        --> llvm.nvvm.barrier.cta.red.and.aligned.all(0, z)
//   * llvm.nvvm.barrier0.or         --> llvm.nvvm.barrier.cta.red.or.aligned.all(0, z)

def llvm_global_ptr_ty  : LLVMQualPointerType<1>;         // (global)ptr
def llvm_shared_ptr_ty  : LLVMQualPointerType<3>;         // (shared)ptr
def llvm_constant_ptr_ty: LLVMQualPointerType<4>;         // (const)ptr
def llvm_local_ptr_ty   : LLVMQualPointerType<5>;         // (local)ptr
def llvm_tmem_ptr_ty    : LLVMQualPointerType<6>;         // (tensor memory)ptr
def llvm_shared_cluster_ptr_ty : LLVMQualPointerType<7>;  // (shared_cluster)ptr

//
// MISC
//

defvar WARP_SIZE = 32;

// Note: the maximum grid size in the x-dimension is the lower value of 65535
// on sm_20. We conservatively use the larger value here as it required for
// sm_30+ and also correct for sm_20.
defvar MAX_GRID_SIZE_X = 0x7fffffff;
defvar MAX_GRID_SIZE_Y = 0xffff;
defvar MAX_GRID_SIZE_Z = 0xffff;

defvar MAX_BLOCK_SIZE_X = 1024;
defvar MAX_BLOCK_SIZE_Y = 1024;
defvar MAX_BLOCK_SIZE_Z = 64;

// Helper class that concatenates list elements with
// a given separator 'sep' and returns the result.
// Handles empty strings.
class StrJoin<string sep, list<string> str_list> {
  string ret = !foldl("", str_list, a, b,
               !if(!eq(a, ""), b, !if(!eq(b, ""), a, !strconcat(a, sep, b))));
}

// Helper class that represents a 'fragment' of an NVPTX *MMA instruction.
// Geom: m<M>n<N>k<K>. E.g. m8n32k16
// Frag: [a|b|c|d] ([x1|x2|x4] for ldmatrix)
// PtxEltType: PTX type for the element.
class WMMA_REGS<string Geom, string Frag, string PtxEltType, bit IsSparse = false> {
  string geom = Geom;
  string frag = Frag;
  string ptx_elt_type = PtxEltType;
  string gft = Geom#":"#Frag#":"#ptx_elt_type;
  string gf = Geom#":"#Frag;
  string ft = frag#":"#ptx_elt_type;
  bit isSparse = IsSparse;
  list<LLVMType> regs = !if(!eq(isSparse, true),
    !cond(
      // mma sparse ops use other fragments for some arguments
      !eq(gft,"m16n8k16:a:bf16") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k16:a:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k16:b:bf16") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k16:b:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k16:c:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k16:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k16:d:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k16:d:f32") : !listsplat(llvm_float_ty, 4),

      !eq(gft,"m16n8k32:a:bf16") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:a:f16") : !listsplat(llvm_v2f16_ty, 4),
      !eq(gft,"m16n8k32:b:bf16") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:b:f16") : !listsplat(llvm_v2f16_ty, 4),
      !eq(gft,"m16n8k32:c:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k32:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k32:d:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k32:d:f32") : !listsplat(llvm_float_ty, 4),

      !eq(gft,"m16n8k16:a:tf32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k16:b:tf32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k16:c:tf32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k16:d:tf32") : !listsplat(llvm_float_ty, 4),

      !eq(gft,"m16n8k8:a:tf32") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k8:b:tf32") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k8:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k8:d:f32") : !listsplat(llvm_float_ty, 4),

      !eq(gft,"m16n8k32:a:u8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:a:s8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:u8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:s8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:d:s32") : !listsplat(llvm_i32_ty, 4),

      !eq(gft,"m16n8k64:a:u8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:a:s8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:a:e4m3") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:a:e5m2") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:a:e3m2") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:a:e2m3") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:a:e2m1") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:u8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:s8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:e4m3") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:e5m2") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:e3m2") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:e2m3") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:e2m1") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:c:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k64:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k64:d:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k64:d:f32") : !listsplat(llvm_float_ty, 4),

      !eq(gft,"m16n8k64:a:u4") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k64:a:s4") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k64:b:u4") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k64:b:s4") : !listsplat(llvm_i32_ty, 2),

      !eq(gft,"m16n8k64:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:d:s32") : !listsplat(llvm_i32_ty, 4),

      !eq(gft,"m16n8k128:a:u4") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:a:s4") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:a:e2m1") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:b:u4") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:b:s4") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:b:e2m1") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k128:d:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:d:f32") : !listsplat(llvm_float_ty, 4),
    ),
    !cond(
      // mma fp ops use smaller fragments than wmma fp ops
      !eq(gft,"m8n8k4:a:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m8n8k4:b:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k8:a:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k8:b:f16") : [llvm_v2f16_ty],
      !eq(gft,"m16n8k8:c:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k8:d:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k8:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k8:d:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k16:a:f16") : !listsplat(llvm_v2f16_ty, 4),
      !eq(gft,"m16n8k16:b:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k16:c:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k16:d:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k16:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k16:d:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k4:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k4:d:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k32:c:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k32:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k32:d:f16") : !listsplat(llvm_v2f16_ty, 2),
      !eq(gft,"m16n8k32:d:f32") : !listsplat(llvm_float_ty, 4),

      // mma.block_scale e2m1 (mxf4, mxf4nvf4) -> f32 @ m16n8k64
      !eq(gft,"m16n8k64:c:f32") : !listsplat(llvm_float_ty, 4),
      !eq(gft,"m16n8k64:d:f32") : !listsplat(llvm_float_ty, 4),

      // wmma fp16 -> fp16/fp32 @  m16n16k16/m8n32k16/m32n8k16
      // All other supported geometries use the same fragment format for f32 and
      // f16, so we only need to consider {fragment, type}.
      !eq(ft,"a:f16") : !listsplat(llvm_v2f16_ty, 8),
      !eq(ft,"b:f16") : !listsplat(llvm_v2f16_ty, 8),
      !eq(ft,"c:f16") : !listsplat(llvm_v2f16_ty, 4),
      !eq(ft,"d:f16") : !listsplat(llvm_v2f16_ty, 4),
      !eq(ft,"c:f32") : !listsplat(llvm_float_ty, 8),
      !eq(ft,"d:f32") : !listsplat(llvm_float_ty, 8),

      // wmma tf32 -> s32 @ m16n16k8
      !eq(gft,"m16n16k8:a:tf32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n16k8:b:tf32") : !listsplat(llvm_i32_ty, 4),

      // mma tf32 -> s32 @ m16n16k8/m16n8k8
      !eq(gft,"m16n8k4:a:tf32") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k4:b:tf32") : [llvm_i32_ty],
      !eq(gft,"m16n8k8:a:tf32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k8:b:tf32") : !listsplat(llvm_i32_ty, 2),

      !eq(gft,"m8n8k4:a:f64") : [llvm_double_ty],
      !eq(gft,"m8n8k4:b:f64") : [llvm_double_ty],
      !eq(gft,"m8n8k4:c:f64") : !listsplat(llvm_double_ty, 2),
      !eq(gft,"m8n8k4:d:f64") : !listsplat(llvm_double_ty, 2),

      !eq(gft,"m16n8k4:a:f64") : !listsplat(llvm_double_ty, 2),
      !eq(gft,"m16n8k4:b:f64") : [llvm_double_ty],
      !eq(gft,"m16n8k4:c:f64") : !listsplat(llvm_double_ty, 4),
      !eq(gft,"m16n8k4:d:f64") : !listsplat(llvm_double_ty, 4),

      !eq(gft,"m16n8k8:a:f64") : !listsplat(llvm_double_ty, 4),
      !eq(gft,"m16n8k8:b:f64") : !listsplat(llvm_double_ty, 2),
      !eq(gft,"m16n8k8:c:f64") : !listsplat(llvm_double_ty, 4),
      !eq(gft,"m16n8k8:d:f64") : !listsplat(llvm_double_ty, 4),

      !eq(gft,"m16n8k16:a:f64") : !listsplat(llvm_double_ty, 8),
      !eq(gft,"m16n8k16:b:f64") : !listsplat(llvm_double_ty, 4),
      !eq(gft,"m16n8k16:c:f64") : !listsplat(llvm_double_ty, 4),
      !eq(gft,"m16n8k16:d:f64") : !listsplat(llvm_double_ty, 4),

      // wmma bf16 -> s32 @ m16n16k16/m8n32k16/m32n8k16
      !eq(gft,"m16n16k16:a:bf16") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n16k16:b:bf16") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m8n32k16:a:bf16") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m8n32k16:b:bf16") : !listsplat(llvm_i32_ty, 8),
      !eq(gft,"m32n8k16:a:bf16") : !listsplat(llvm_i32_ty, 8),
      !eq(gft,"m32n8k16:b:bf16") : !listsplat(llvm_i32_ty, 2),

      // mma bf16 -> s32 @ m16n8k16/m16n8k8
      !eq(gft,"m16n8k16:a:bf16") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k16:b:bf16") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k8:a:bf16") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k8:b:bf16") : [llvm_i32_ty],

      // wmma u8/s8 -> s32 @ m16n16k16/m8n32k16/m32n8k16
      !eq(gft,"m16n16k16:a:u8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n16k16:a:s8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n16k16:b:u8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n16k16:b:s8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n16k16:c:s32") : !listsplat(llvm_i32_ty, 8),
      !eq(gft,"m16n16k16:d:s32") : !listsplat(llvm_i32_ty, 8),

      !eq(gft,"m8n32k16:a:u8") : [llvm_i32_ty],
      !eq(gft,"m8n32k16:a:s8") : [llvm_i32_ty],
      !eq(gft,"m8n32k16:b:u8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m8n32k16:b:s8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m8n32k16:c:s32") : !listsplat(llvm_i32_ty, 8),
      !eq(gft,"m8n32k16:d:s32") : !listsplat(llvm_i32_ty, 8),

      !eq(gft,"m32n8k16:a:u8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m32n8k16:a:s8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m32n8k16:b:u8") : [llvm_i32_ty],
      !eq(gft,"m32n8k16:b:s8") : [llvm_i32_ty],
      !eq(gft,"m32n8k16:c:s32") : !listsplat(llvm_i32_ty, 8),
      !eq(gft,"m32n8k16:d:s32") : !listsplat(llvm_i32_ty, 8),

      // mma u8/s8 -> s32 @ m8n8k16/m16n8k16/m16n8k32
      !eq(gft,"m8n8k16:a:u8") : [llvm_i32_ty],
      !eq(gft,"m8n8k16:a:s8") : [llvm_i32_ty],
      !eq(gft,"m8n8k16:b:u8") : [llvm_i32_ty],
      !eq(gft,"m8n8k16:b:s8") : [llvm_i32_ty],
      !eq(gft,"m8n8k16:c:s32") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m8n8k16:d:s32") : !listsplat(llvm_i32_ty, 2),

      !eq(gft,"m16n8k16:a:u8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k16:a:s8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k16:b:u8") : [llvm_i32_ty],
      !eq(gft,"m16n8k16:b:s8") : [llvm_i32_ty],
      !eq(gft,"m16n8k16:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k16:d:s32") : !listsplat(llvm_i32_ty, 4),

      !eq(gft,"m16n8k32:a:u8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:a:s8") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:b:u8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:s8") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:d:s32") : !listsplat(llvm_i32_ty, 4),

      // wmma/mma u4/s4 -> s32 @ m8n8k32 (u4/s4)
      !eq(gft,"m8n8k32:a:u4") : [llvm_i32_ty],
      !eq(gft,"m8n8k32:a:s4") : [llvm_i32_ty],
      !eq(gft,"m8n8k32:b:u4") : [llvm_i32_ty],
      !eq(gft,"m8n8k32:b:s4") : [llvm_i32_ty],
      !eq(gft,"m8n8k32:c:s32") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m8n8k32:d:s32") : !listsplat(llvm_i32_ty, 2),

      !eq(gft,"m16n8k32:a:u4") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:a:s4") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:u4") : [llvm_i32_ty],
      !eq(gft,"m16n8k32:b:s4") : [llvm_i32_ty],
      !eq(gft,"m16n8k32:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:d:s32") : !listsplat(llvm_i32_ty, 4),

      !eq(gft,"m16n8k64:a:u4") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:a:s4") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:u4") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k64:b:s4") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k64:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:d:s32") : !listsplat(llvm_i32_ty, 4),

      // mma e4m3/e5m2 -> f16/f32 @ m16n8k16
      !eq(gft,"m16n8k16:a:e4m3") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k16:a:e5m2") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k16:b:e4m3") : [llvm_i32_ty],
      !eq(gft,"m16n8k16:b:e5m2") : [llvm_i32_ty],
      // mma e4m3/e5m2/e3m2/e2m3/e2m1 -> f32 @ m16n8k32
      !eq(gft,"m16n8k32:a:e4m3") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:a:e5m2") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:a:e3m2") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:a:e2m3") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:a:e2m1") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k32:b:e4m3") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:e5m2") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:e3m2") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:e2m3") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k32:b:e2m1") : !listsplat(llvm_i32_ty, 2),
      // mma e2m1 -> f32 @m16n8k64
      !eq(gft,"m16n8k64:a:e2m1") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k64:b:e2m1") : !listsplat(llvm_i32_ty, 2),

      // wmma/mma b1 -> s32 @ m8n8k128(b1)
      !eq(gft,"m8n8k128:a:b1") : [llvm_i32_ty],
      !eq(gft,"m8n8k128:b:b1") : [llvm_i32_ty],
      !eq(gft,"m8n8k128:c:s32") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m8n8k128:d:s32") : !listsplat(llvm_i32_ty, 2),

      !eq(gft,"m16n8k128:a:b1") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k128:b:b1") : [llvm_i32_ty],
      !eq(gft,"m16n8k128:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k128:d:s32") : !listsplat(llvm_i32_ty, 4),

      !eq(gft,"m16n8k256:a:b1") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k256:b:b1") : !listsplat(llvm_i32_ty, 2),
      !eq(gft,"m16n8k256:c:s32") : !listsplat(llvm_i32_ty, 4),
      !eq(gft,"m16n8k256:d:s32") : !listsplat(llvm_i32_ty, 4),

      // ldmatrix b16 -> s32 @ m8n8
      !eq(gf,"m8n8:x1") : !listsplat(llvm_i32_ty, 1),
      !eq(gf,"m8n8:x2") : !listsplat(llvm_i32_ty, 2),
      !eq(gf,"m8n8:x4") : !listsplat(llvm_i32_ty, 4),

      // ldmatrix b8, b8x16.b6x16_p32, b8x16.b4x16_p64 -> s32 @ m16n16
      !eq(gf,"m16n16:x1") : !listsplat(llvm_i32_ty, 2),
      !eq(gf,"m16n16:x2") : !listsplat(llvm_i32_ty, 4),

      // ldmatrix b8x16.b6x16_p32, b8x16.b4x16_p64 -> s32 @ m8n16
      !eq(gf,"m8n16:x1") : !listsplat(llvm_i32_ty, 1),
      !eq(gf,"m8n16:x2") : !listsplat(llvm_i32_ty, 2),
      !eq(gf,"m8n16:x4") : !listsplat(llvm_i32_ty, 4),

      // stmatrix b8 -> s32 @ m16n8
      !eq(gf,"m16n8:x1") : !listsplat(llvm_i32_ty, 1),
      !eq(gf,"m16n8:x2") : !listsplat(llvm_i32_ty, 2),
      !eq(gf,"m16n8:x4") : !listsplat(llvm_i32_ty, 4),
    )
  );
}

class WMMA_NAME_LDST<string Op, WMMA_REGS Frag, string Layout, int WithStride> {
  string intr_name = "llvm.nvvm.wmma."
                # Frag.geom
                # "." # Op
                # "." # Frag.frag
                # "." # Layout
                # !if(WithStride, ".stride", "")
                # "." # Frag.ptx_elt_type
                ;
  // TODO(tra): record name should ideally use the same field order as the intrinsic.
  // E.g. string record = !subst("llvm", "int",
  //                      !subst(".", "_", llvm));
  string record_name = "int_nvvm_wmma_"
                # Frag.geom
                # "_" # Op
                # "_" # Frag.frag
                # "_" # Frag.ptx_elt_type
                # "_" # Layout
                # !if(WithStride, "_stride", "");
}

class MMA_SIGNATURE<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
  list<WMMA_REGS> id_frags = !cond(
     // FP8/F8F6F4 ops are identified by A,B inputs & accomulator & result type.
     !or(!eq(A.ptx_elt_type, "e4m3"),
         !eq(A.ptx_elt_type, "e5m2"),
         !eq(A.ptx_elt_type, "e3m2"),
         !eq(A.ptx_elt_type, "e2m3"),
         !eq(A.ptx_elt_type, "e2m1")): [D, A, B, C],
     // FP16 ops are identified by accumulator & result type.
     !eq(A.ptx_elt_type, "f16") : [D, C],
     // other ops are identified by input types.
     !ne(A.ptx_elt_type, B.ptx_elt_type): [A, B],
     true: [A]
     );
   string ret = !foldl("", id_frags, a, b, !strconcat(a, "_", b.ptx_elt_type));
}

class WMMA_NAME<string ALayout, string BLayout, int Satfinite, string Rnd, string b1op,
                WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
  string signature = MMA_SIGNATURE<A, B, C, D>.ret;
  string record_name = "int_nvvm_wmma_"
                  # A.geom
                  # "_mma"
                  # !subst(".", "_", b1op)
                  # "_" # ALayout
                  # "_" # BLayout
                  # !if(!ne(Rnd, ""), !strconcat("_", Rnd), "")
                  # signature
                  # !if(Satfinite, "_satfinite", "");
}

class MMA_NAME<string ALayout, string BLayout, int Satfinite, string b1op, string Kind,
               WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
  string signature = MMA_SIGNATURE<A, B, C, D>.ret;
  string record_name = "int_nvvm_mma"
                  # !subst(".", "_", b1op)
                  # "_" # A.geom
                  # "_" # ALayout
                  # "_" # BLayout
                  # !if(!ne(Kind, ""), !strconcat("_", !subst("::", "_", Kind)), "")
                  # !if(Satfinite, "_satfinite", "")
                  # signature;
}

class MMA_BLOCK_SCALE_NAME<string Kind, string SType, string ScaleVecSize,
                           WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> {
  string signature = MMA_SIGNATURE<A, B, C, D>.ret;
  string record_name = "int_nvvm_mma_block_scale"
                  # "_" # A.geom
                  # "_row_col"
                  # "_" # Kind
                  # !subst(".", "_", ScaleVecSize)
                  # signature
                  # "_" # SType;
}

class MMA_SP_NAME<string Metadata, string Kind, int Satfinite,
                  WMMA_REGS A, WMMA_REGS B,
                  WMMA_REGS C, WMMA_REGS D> {
  string signature = MMA_SIGNATURE<A, B, C, D>.ret;
  string record_name = "int_nvvm_mma"
                  # "_" # !subst("::", "_", Metadata)
                  # "_" # A.geom
                  # "_row_col"
                  # !if(!ne(Kind, ""), !strconcat("_", !subst("::", "_", Kind)), "")
                  # !if(Satfinite, "_satfinite", "")
                  # signature;
}

class MMA_SP_BLOCK_SCALE_NAME<string Kind, string SType, string ScaleVecSize,
                              WMMA_REGS A, WMMA_REGS B,
                              WMMA_REGS C, WMMA_REGS D> {
  string signature = MMA_SIGNATURE<A, B, C, D>.ret;
  string record_name = "int_nvvm_mma_sp_ordered_metadata_block_scale"
                  # "_" # A.geom
                  # "_row_col"
                  # "_" # Kind
                  # !subst(".", "_", ScaleVecSize)
                  # signature
                  # "_" # SType;
}

// Helper class that takes an intrinsic name and construct a record name.
// Additionally, sets `intr_name` to be non-empty if the default name assigned
// to this intrinsic will not match the name given.
class IntrinsicName<string name> {
  string record_name = !subst(".", "_",
                  !subst("llvm.", "int_", name));
  // Use explicit intrinsic name if it has an _ in it, else rely on LLVM
  // assigned default name.
  string intr_name = !if(!ne(!find(name, "_"), -1), name, "");
}

class LDMATRIX_NAME<WMMA_REGS Frag, int Trans> {
  defvar name = "llvm.nvvm.ldmatrix.sync.aligned"
                # "." # Frag.geom
                # "." # Frag.frag
                # !if(Trans, ".trans", "")
                # "." # Frag.ptx_elt_type
                ;
  string intr_name = IntrinsicName<name>.intr_name;
  string record_name = IntrinsicName<name>.record_name;
}

class STMATRIX_NAME<WMMA_REGS Frag, int Trans> {
  defvar name = "llvm.nvvm.stmatrix.sync.aligned"
                # "." # Frag.geom
                # "." # Frag.frag
                # !if(Trans, ".trans", "")
                # "." # Frag.ptx_elt_type
                ;
  string intr_name = IntrinsicName<name>.intr_name;
  string record_name = IntrinsicName<name>.record_name;
}

// Generates list of 4-tuples of WMMA_REGS representing a valid MMA op.
//   Geom: list of supported geometries.
//   TypeN: PTX type of the corresponding fragment's element.
//   TypeB and TypeD may be empty if it must match that of TypeA or TypeC.
class MMA_OPS<list<string> Geom, list<string> TypeA, list<string> TypeB,
            list<string> TypeC, list<string> TypeD, bit IsSparse = false> {
  list<list<WMMA_REGS>> ret =
     !foldl([]<list<WMMA_REGS>>, Geom, t1, geom, !listconcat(t1,
     !foldl([]<list<WMMA_REGS>>, TypeA, t2, type_a, !listconcat(t2,
     !foldl([]<list<WMMA_REGS>>, !if(!size(TypeB), TypeB, [type_a]), t3, type_b, !listconcat(t3,
     !foldl([]<list<WMMA_REGS>>, TypeC, t4, type_c, !listconcat(t4,
     !foldl([]<list<WMMA_REGS>>, !if(!size(TypeD), TypeD, [type_c]), t5, type_d, !listconcat(t5,
            [[WMMA_REGS<geom, "a", type_a, IsSparse>,
              WMMA_REGS<geom, "b", type_b, IsSparse>,
              WMMA_REGS<geom, "c", type_c, IsSparse>,
              WMMA_REGS<geom, "d", type_d, IsSparse>]]))))))))));
   // Debugging aid for readable representation of the list above.
   list<list<string>> ops = !foreach(x, ret, [x[0].gft, x[1].gft, x[2].gft, x[3].gft]);
}


class MMA_LDST_OPS<list<string> Geom, list<string> Frags, list<string> Types> {
  list<WMMA_REGS> ret =
     !foldl([]<WMMA_REGS>, Geom, t1, geom, !listconcat(t1,
     !foldl([]<WMMA_REGS>, Frags, t2, frag, !listconcat(t2,
     !foldl([]<WMMA_REGS>, Types, t3, type, !listconcat(t3,
            [WMMA_REGS<geom, frag, type>]))))));
   // Debugging aid for readable representation of the list above.
   list<string> ops = !foreach(x, ret, x.gft);
}

class LDMATRIX_OPS<list<string> Geom, list<string> Frags, list<string> Types> {
  list<WMMA_REGS> ret =
     !foldl([]<WMMA_REGS>, Geom, t1, geom, !listconcat(t1,
     !foldl([]<WMMA_REGS>, Frags, t2, frag, !listconcat(t2,
     !foldl([]<WMMA_REGS>, Types, t3, type, !listconcat(t3,
            [WMMA_REGS<geom, frag, type>]))))));
   // Debugging aid for readable representation of the list above.
   list<string> ops = !foreach(x, ret, x.gft);
}

class STMATRIX_OPS<list<string> Geom, list<string> Frags, list<string> Types> {
  list<WMMA_REGS> ret =
     !foldl([]<WMMA_REGS>, Geom, t1, geom, !listconcat(t1,
     !foldl([]<WMMA_REGS>, Frags, t2, frag, !listconcat(t2,
     !foldl([]<WMMA_REGS>, Types, t3, type, !listconcat(t3,
            [WMMA_REGS<geom, frag, type>]))))));
   // Debugging aid for readable representation of the list above.
   list<string> ops = !foreach(x, ret, x.gft);
}

// Creates list of valid combinations of fragments. This is the main list that
// drives generation of corresponding intrinsics and instructions.
class NVVM_MMA_OPS {
  list<list<WMMA_REGS>> tf32_wmma_ops = MMA_OPS<
            ["m16n16k8"],
            ["tf32"], [], ["f32"], []>.ret;
  list<list<WMMA_REGS>> bf16_wmma_ops = MMA_OPS<
            ["m16n16k16", "m32n8k16", "m8n32k16"],
            ["bf16"], [], ["f32"], []>.ret;
  list<list<WMMA_REGS>> f64_wmma_ops = MMA_OPS<
            ["m8n8k4"],
            ["f64"], [], ["f64"], []>.ret;
  list<list<WMMA_REGS>> fp_wmma_ops = MMA_OPS<
            ["m16n16k16", "m32n8k16", "m8n32k16"],
            ["f16"], [], ["f16", "f32"], ["f16", "f32"]>.ret;
  list<list<WMMA_REGS>> int_wmma_ops = MMA_OPS<
            ["m16n16k16", "m32n8k16", "m8n32k16"],
            ["s8", "u8"], [], ["s32"], []>.ret;
  list<list<WMMA_REGS>> subint_wmma_ops = MMA_OPS<
            ["m8n8k32"],
            ["s4", "u4"], [], ["s32"], []>.ret;
  list<list<WMMA_REGS>> bit_wmma_ops = MMA_OPS<
            ["m8n8k128"],
            ["b1"], [], ["s32"], []>.ret;
  list<list<WMMA_REGS>> all_wmma_ops = !listconcat(
            tf32_wmma_ops, bf16_wmma_ops, f64_wmma_ops,
            fp_wmma_ops, int_wmma_ops, subint_wmma_ops, bit_wmma_ops);

  list<list<WMMA_REGS>> tf32_mma_ops = MMA_OPS<
            ["m16n8k4", "m16n8k8"],
            ["tf32"], [], ["f32"], []>.ret;
  list<list<WMMA_REGS>> bf16_mma_ops = MMA_OPS<
            ["m16n8k16", "m16n8k8"],
            ["bf16"], [], ["f32"], []>.ret;
  list<list<WMMA_REGS>> f64_mma_ops = MMA_OPS<
            ["m8n8k4", "m16n8k4", "m16n8k8", "m16n8k16"],
            ["f64"], [], ["f64"], []>.ret;
  list<list<WMMA_REGS>> fp_mma_ops = MMA_OPS<
            ["m8n8k4", "m16n8k8", "m16n8k16"],
            ["f16"], [], ["f16", "f32"], ["f16", "f32"]>.ret;
  list<list<WMMA_REGS>> int_mma_ops = MMA_OPS<
            ["m8n8k16", "m16n8k16", "m16n8k32"],
            ["s8", "u8"], ["s8", "u8"], ["s32"], []>.ret;
  // m16n8k32 fp8 variants are intersected with f8f6f4 variants
  // and processed there
  list<list<WMMA_REGS>> fp8_mma_ops = MMA_OPS<
            ["m16n8k16"],
            ["e4m3", "e5m2"], ["e4m3", "e5m2"],
            ["f16", "f32"], ["f16", "f32"]>.ret;
  // it also contains e4m3/e5m2 from fp8 variants
  list<list<WMMA_REGS>> f8f6f4_mma_ops = MMA_OPS<
            ["m16n8k32"],
            ["e4m3", "e5m2", "e3m2", "e2m3", "e2m1"],
            ["e4m3", "e5m2", "e3m2", "e2m3", "e2m1"],
            ["f16", "f32"], ["f16", "f32"]>.ret;
  list<list<WMMA_REGS>> subint_mma_ops = MMA_OPS<
            ["m8n8k32", "m16n8k32", "m16n8k64"],
            ["s4", "u4"], ["s4", "u4"], ["s32"], []>.ret;
  list<list<WMMA_REGS>> bit_mma_ops = MMA_OPS<
            ["m8n8k128", "m16n8k128", "m16n8k256"],
            ["b1"], [], ["s32"], []>.ret;
  list<list<WMMA_REGS>> all_mma_ops = !listconcat(
            tf32_mma_ops, bf16_mma_ops, f64_mma_ops,
            fp_mma_ops, fp8_mma_ops, f8f6f4_mma_ops,
            int_mma_ops, subint_mma_ops, bit_mma_ops);

  list<list<WMMA_REGS>> mxf4_mma_ops = MMA_OPS<
    ["m16n8k64"], ["e2m1"], ["e2m1"], ["f32"], ["f32"]
  >.ret;

  list<list<WMMA_REGS>> mxf8f6f4_mma_ops = MMA_OPS<
    ["m16n8k32"], ["e2m1", "e2m3", "e3m2", "e5m2", "e4m3"],
                  ["e2m1", "e2m3", "e3m2", "e5m2", "e4m3"], ["f32"], ["f32"]
  >.ret;

  list<list<WMMA_REGS>> all_mma_block_scale_ops = !listconcat(
            mxf4_mma_ops, mxf8f6f4_mma_ops);

  list<list<WMMA_REGS>> bf16_mma_sp_ops = MMA_OPS<
            ["m16n8k16", "m16n8k32"],
            ["bf16"], [], ["f32"], [], true>.ret;
  list<list<WMMA_REGS>> tf32_mma_sp_ops = MMA_OPS<
            ["m16n8k8", "m16n8k16"],
            ["tf32"], [], ["f32"], [], true>.ret;
  list<list<WMMA_REGS>> fp_mma_sp_ops = MMA_OPS<
            ["m16n8k16", "m16n8k32"],
            ["f16"], [], ["f16", "f32"], ["f16", "f32"], true>.ret;
  list<list<WMMA_REGS>> fp8_mma_sp_ops = MMA_OPS<
            ["m16n8k64"],
            ["e4m3", "e5m2", "e3m2", "e2m3", "e2m1"],
            ["e4m3", "e5m2", "e3m2", "e2m3", "e2m1"],
            ["f16", "f32"], ["f16", "f32"], true>.ret;
  list<list<WMMA_REGS>> subint_mma_sp_ops = MMA_OPS<
            ["m16n8k64", "m16n8k128"],
            ["s4", "u4"], ["s4", "u4"], ["s32"], [], true>.ret;
  list<list<WMMA_REGS>> int_mma_sp_ops = MMA_OPS<
            ["m16n8k32", "m16n8k64"],
            ["s8", "u8"], ["s8", "u8"], ["s32"], [], true>.ret;
  list<list<WMMA_REGS>> all_mma_sp_ops = !listconcat(
            bf16_mma_sp_ops, tf32_mma_sp_ops, fp_mma_sp_ops, fp8_mma_sp_ops,
            subint_mma_sp_ops, int_mma_sp_ops);

  // combines available geoms and types for mxf4 and mxf4nvf4 kinds
  list<list<WMMA_REGS>> mxf4xx_mma_sp_ops = MMA_OPS<
            ["m16n8k128"],
            ["e2m1"], ["e2m1"], ["f32"], [], true>.ret;
  list<list<WMMA_REGS>> mxf8f6f4_mma_sp_ops = MMA_OPS<
            ["m16n8k64"],
            ["e4m3", "e5m2", "e3m2", "e2m3", "e2m1"],
            ["e4m3", "e5m2", "e3m2", "e2m3", "e2m1"],
            ["f32"], [], true>.ret;
  list<list<WMMA_REGS>> all_mma_sp_block_scale_ops = !listconcat(
            mxf4xx_mma_sp_ops, mxf8f6f4_mma_sp_ops);

  list<WMMA_REGS> ldst_ab_ops = MMA_LDST_OPS<
            ["m16n16k16", "m32n8k16", "m8n32k16"],
            ["a", "b"], ["f16", "u8", "s8", "bf16"]>.ret;
  list<WMMA_REGS> ldst_cd_ops = MMA_LDST_OPS<
            ["m16n16k16", "m32n8k16", "m8n32k16"],
            ["c", "d"], ["f16", "f32", "s32"]>.ret;
  list<WMMA_REGS> ldst_tf32_ab_ops = MMA_LDST_OPS<
            ["m16n16k8"],
            ["a", "b"], ["tf32"]>.ret;
  list<WMMA_REGS> ldst_tf32_cd_ops = MMA_LDST_OPS<
            ["m16n16k8"],
            ["c", "d"], ["f32"]>.ret;
  list<WMMA_REGS> ldst_f64_abcd_ops = MMA_LDST_OPS<
            ["m8n8k4"],
            ["a", "b", "c", "d"], ["f64"]>.ret;
  list<WMMA_REGS> ldst_subint_ab_ops = MMA_LDST_OPS<
            ["m8n8k32"], ["a", "b"], ["s4","u4"]>.ret;
  list<WMMA_REGS> ldst_bit_ab_ops = MMA_LDST_OPS<
            ["m8n8k128"], ["a", "b"], ["b1"]>.ret;
  list<WMMA_REGS> ldst_subint_cd_ops = MMA_LDST_OPS<
            ["m8n8k32", "m8n8k128"],  ["c", "d"], ["s32"]>.ret;
  list<WMMA_REGS> all_ldst_ops = !listconcat(ldst_ab_ops, ldst_cd_ops,
                                             ldst_tf32_ab_ops,
                                             ldst_tf32_cd_ops,
                                             ldst_f64_abcd_ops,
                                             ldst_subint_ab_ops,
                                             ldst_bit_ab_ops,
                                             ldst_subint_cd_ops);
  // Separate A/B/C fragments (loads) from D (stores).
  list<WMMA_REGS> all_ld_ops = !filter(op, all_ldst_ops, !ne(op.frag, "d"));
  list<WMMA_REGS> all_st_ops = !filter(op, all_ldst_ops, !eq(op.frag, "d"));

  list<WMMA_REGS> ldmatrix_b16_ops = LDMATRIX_OPS<
    ["m8n8"], ["x1", "x2", "x4"], ["b16"]>.ret;

  list<WMMA_REGS> ldmatrix_geom_m16n16_ops = LDMATRIX_OPS<
    ["m16n16"], ["x1", "x2"], ["b8", "b8x16.b6x16_p32", "b8x16.b4x16_p64"]>.ret;

  list<WMMA_REGS> ldmatrix_geom_m8n16_ops = LDMATRIX_OPS<
    ["m8n16"], ["x1", "x2", "x4"], ["b8x16.b6x16_p32", "b8x16.b4x16_p64"]>.ret;

  list<WMMA_REGS> stmatrix_b16_ops = STMATRIX_OPS<
    ["m8n8"], ["x1", "x2", "x4"], ["b16"]>.ret;

  list<WMMA_REGS> stmatrix_b8_ops = STMATRIX_OPS<
    ["m16n8"], ["x1", "x2", "x4"], ["b8"]>.ret;

  list<WMMA_REGS> all_ldmatrix_ops = !listconcat(ldmatrix_b16_ops,
                                                 ldmatrix_geom_m16n16_ops,
                                                 ldmatrix_geom_m8n16_ops);

  list<WMMA_REGS> all_stmatrix_ops = !listconcat(stmatrix_b16_ops,
                                                 stmatrix_b8_ops);
}

def NVVM_MMA_OPS : NVVM_MMA_OPS;

// Returns true if this combination of fragment and layout for WMMA load/store
// ops is supported; false otherwise.
// E.g.
// if NVVM_WMMA_LDST_SUPPORTED<...>.ret then
//   def : FOO<>; // The record will only be defined for supported ops.
//
class NVVM_WMMA_LDST_SUPPORTED<WMMA_REGS frag, string layout> {
  string f = frag.frag;
  string t = frag.ptx_elt_type;

  bit ret = !cond(
    // Sub-int load and store requires A fragment to be of row layout and B
    // fragments to be of column layout.
    !and(!or(!eq(t, "b1"),
             !eq(t, "u4"),
             !eq(t, "s4")),
         !or(!and(!eq(f, "a"),
                  !ne(layout, "row")),
             !and(!eq(f, "b"),
                  !ne(layout, "col")))) : false,
    true: true
  );
}

// Returns true if this combination of layout/satf/rnd for WMMA ops is
// supported; false otherwise.
// E.g.
// if NVVM_WMMA_SUPPORTED<...>.ret then
//   def : FOO<>; // The record will only be defined for supported ops.
//
class NVVM_WMMA_SUPPORTED<list<WMMA_REGS> frags, string layout_a, string layout_b, int satf, string rnd> {
  // WMMA ops check both layouts.
  string layout = layout_a # ":" # layout_b;
  string t = frags[0].ptx_elt_type;

  bit ret = !cond(
    // only f64 wmma functions support rnd options
    // any non f64 type that uses a rnd value is invalid
    !and(!ne(t, "f64"), !ne(rnd, "")) : false,

    // satf is only valid for select types
    !and(!eq(satf, 1),
         !ne(t, "s8"),
         !ne(t, "u8"),
         !ne(t, "s4"),
         !ne(t, "u4"),
         !ne(t, "f16")): false,

    // Sub-int wmma requires row/column layout
    !and(!or(!eq(t, "s4"),
             !eq(t, "u4"),
             !eq(t, "b1")),
         !ne(layout, "row:col")) : false,
    true: true
  );
}

class NVVM_MMA_B1OPS<list<WMMA_REGS> frags> {
  list<string> ret = !cond(
    !eq(frags[0].ptx_elt_type, "b1") : [".xor.popc", ".and.popc"],
    true: [""]
  );
}

// Returns true if this combination of layout/satf for MMA ops is supported;
// false otherwise.
// E.g.
// if NVVM_MMA_SUPPORTED<...>.ret then
//   def : FOO<>; // The record will only be defined for supported ops.
//
class NVVM_MMA_SUPPORTED<list<WMMA_REGS> frags, string layout_a, string layout_b,
                         string kind, int satf> {
  // MMA ops check both layouts.
  string layout = layout_a # ":" # layout_b;
  string a_type = frags[0].ptx_elt_type;
  string b_type = frags[1].ptx_elt_type;
  string c_type = frags[2].ptx_elt_type;
  string d_type = frags[3].ptx_elt_type;
  string geom = frags[0].geom;

  // gcd is a shortcut used to identify instructions that depend on
  // geom+frag_c+frag_d.
  string gcd = geom # ":" # c_type # d_type;
  bit ret = !cond(

    // Limit satf to valid types
    !and(!eq(satf, 1),
         !ne(a_type, "s8"),
         !ne(a_type, "u8"),
         !ne(a_type, "s4"),
         !ne(a_type, "u4")): false,

    // m8n8k4 has no C=f32 D=f16 variant.
    !eq(gcd, "m8n8k4:f32f16"): false,

    // only m8n8k4 for f16 does not require row:col layout
    !and(!ne(layout, "row:col"),
         !or(!ne(geom, "m8n8k4"),
             !ne(a_type, "f16"))) : false,

    // m16n8k8 requires A and B to be the same type and C and D to be the same
    // type.
    !and(!eq(geom, "m16n8k8"),
         !or(!ne(a_type, b_type),
             !ne(c_type, d_type))): false,

    // m16n8k16/m16n8k32 requires C and D to be the same type
    !and(!or(!eq(geom, "m16n8k16"),
             !eq(geom, "m16n8k32")),
         !ne(c_type, d_type)): false,

    // Limit kind to valid types and geometries
    !and(!ne(kind, ""),
         !or(!ne(geom, "m16n8k32"),
             !and(!ne(a_type, "e4m3"),
                  !ne(a_type, "e5m2"),
                  !ne(a_type, "e3m2"),
                  !ne(a_type, "e2m3"),
                  !ne(a_type, "e2m1")))): false,

    // Limit m16n8k16/m16n8k32 with no kind to valid types
    !and(!eq(kind, ""),
         !or(!eq(geom, "m16n8k16"),
             !eq(geom, "m16n8k32")),
             !or(!eq(a_type, "e3m2"),
                 !eq(a_type, "e2m3"),
                 !eq(a_type, "e2m1"),
                 !eq(b_type, "e3m2"),
                 !eq(b_type, "e2m3"),
                 !eq(b_type, "e2m1"))): false,

    // All other are OK.
    true: true
  );
}

class NVVM_MMA_BLOCK_SCALE_SUPPORTED<list<WMMA_REGS> frags, string kind, string stype, string scale_vec_size> {
  string geom = frags[0].geom;

  bit ret = !cond(
    !and(!eq(geom, "m16n8k64"),
         !eq(kind, "mxf4"),
         !or(!eq(scale_vec_size, ""),
             !eq(scale_vec_size, ".scale_2x")),
         !eq(stype, "ue8m0")) : true,
    !and(!eq(geom, "m16n8k64"),
         !eq(kind, "mxf4nvf4"),
         !eq(scale_vec_size, ".scale_2x"),
         !eq(stype, "ue8m0")) : true,
    !and(!eq(geom, "m16n8k64"),
         !eq(kind, "mxf4nvf4"),
         !eq(scale_vec_size, ".scale_4x"),
         !or(!eq(stype, "ue4m3"),
             !eq(stype, "ue8m0"))) : true,
    !and(!eq(geom, "m16n8k32"),
         !eq(kind, "mxf8f6f4"),
         !or(!eq(scale_vec_size, ""),
             !eq(scale_vec_size, ".scale_1x")),
         !eq(stype, "ue8m0")) : true,
    true: false
  );
}

// Returns true if the fragment is valid for ldmatrix ops is supported;
// false otherwise.
// E.g.
// if NVVM_LDMATRIX_SUPPORTED<...>.ret then
//   def : FOO<>; // The record will only be defined for supported ops.
//
class NVVM_LDMATRIX_SUPPORTED<WMMA_REGS frag, bit trans> {
  string g = frag.geom;
  string t = frag.ptx_elt_type;

  bit ret = !cond(
    !and(!eq(g, "m8n8"), !eq(t, "b16")): true,
    !and(!eq(g, "m16n16"), !eq(t, "b8"), !eq(trans, 1)): true,
    !and(!eq(g, "m16n16"), !eq(t, "b8x16.b6x16_p32"), !eq(trans, 1)): true,
    !and(!eq(g, "m16n16"), !eq(t, "b8x16.b4x16_p64"), !eq(trans, 1)): true,
    !and(!eq(g, "m8n16"), !eq(t, "b8"), !eq(trans, 0)): true,
    !and(!eq(g, "m8n16"), !eq(t, "b8x16.b6x16_p32"), !eq(trans, 0)): true,
    !and(!eq(g, "m8n16"), !eq(t, "b8x16.b4x16_p64"), !eq(trans, 0)): true,
    true: false
  );
}

// Returns true if the fragment is valid for stmatrix ops is supported;
// false otherwise.
class NVVM_STMATRIX_SUPPORTED<WMMA_REGS frag, bit trans> {
  string g = frag.geom;
  string t = frag.ptx_elt_type;

  bit ret = !cond(
    !and(!eq(g, "m8n8"), !eq(t, "b16")): true,
    !and(!eq(g, "m16n8"), !eq(t, "b8"), !eq(trans, 1)): true,
    true: false
  );
}


// Returns true if this combination of layout/kind/satf for MMA.SP ops is supported;
// false otherwise.
// E.g.
// if NVVM_MMA_SP_SUPPORTED<...>.ret then
//   def : FOO<>; // The record will only be defined for supported ops.
//
class NVVM_MMA_SP_SUPPORTED<list<WMMA_REGS> frags, string metadata,
                            string kind, int satf> {
  // MMA.SP ops check both layouts.
  string a_type = frags[0].ptx_elt_type;
  string b_type = frags[1].ptx_elt_type;
  string c_type = frags[2].ptx_elt_type;
  string d_type = frags[3].ptx_elt_type;
  string geom = frags[0].geom;

  bit is_int = !or(!eq(a_type, "s8"),
                   !eq(a_type, "u8"),
                   !eq(a_type, "s4"),
                   !eq(a_type, "u4"));

  bit ret = !cond(

    // Limit satf to valid types
    !and(!eq(satf, 1),
         !eq(is_int, 0)): false,

    // f16/bf16/tf32 requires A and B to be the same type.
    !and(!or(!eq(a_type, "f16"),
             !eq(a_type, "bf16"),
             !eq(a_type, "tf32")),
         !ne(a_type, b_type)): false,

    // m16n8k16, m16n8k32 and m16n8k64 requires C and D to be the same type.
    !and(!or(!eq(geom, "m16n8k16"),
             !eq(geom, "m16n8k32"),
             !eq(geom, "m16n8k64")),
         !ne(c_type, d_type)): false,

    !and(!eq(kind, ""),
         !or(!eq(a_type, "e3m2"),
             !eq(a_type, "e2m3"),
             !eq(a_type, "e2m1"),
             !eq(b_type, "e3m2"),
             !eq(b_type, "e2m3"),
             !eq(b_type, "e2m1"))): false,

    !and(!eq(kind, ""),
         !eq(geom, "m16n8k64"),
         !or(!eq(c_type, "f16"),
             !eq(d_type, "f16"))): false,

    !and(!ne(kind, ""),
         !or(!eq(metadata, "sp"),
             !ne(geom, "m16n8k64"),
             !eq(is_int, 1))): false,

    // All other are OK.
    true: true
  );
}


// Returns true if this combination of kind/scale_vec_size/stype
// for MMA.SP ops is supported;
// false otherwise.
// E.g.
// if NVVM_MMA_SP_BLOCK_SCALE_SUPPORTED<...>.ret then
//   def : FOO<>; // The record will only be defined for supported ops.
//
class NVVM_MMA_SP_BLOCK_SCALE_SUPPORTED<list<WMMA_REGS> frags, string kind,
                                        string stype, string scale_vec_size> {
  // MMA.SP ops check both layouts.
  string a_type = frags[0].ptx_elt_type;
  string b_type = frags[1].ptx_elt_type;
  string c_type = frags[2].ptx_elt_type;
  string d_type = frags[3].ptx_elt_type;
  string geom = frags[0].geom;

  bit ret = !cond(
    !and(!eq(geom, "m16n8k128"),
         !eq(kind, "mxf4"),
         !eq(stype, "ue8m0"),
         !or(!eq(scale_vec_size, ""),
             !eq(scale_vec_size, ".scale_2x"))): true,

    !and(!eq(geom, "m16n8k128"),
         !eq(kind, "mxf4nvf4"),
         !eq(stype, "ue8m0"),
         !eq(scale_vec_size, ".scale_2x")): true,

    !and(!eq(geom, "m16n8k128"),
         !eq(kind, "mxf4nvf4"),
         !or(!eq(stype, "ue4m3"),
             !eq(stype, "ue8m0")),
         !eq(scale_vec_size, ".scale_4x")): true,

    !and(!eq(geom, "m16n8k64"),
         !eq(kind, "mxf8f6f4"),
         !eq(stype, "ue8m0"),
         !or(!eq(scale_vec_size, ""),
             !eq(scale_vec_size, ".scale_1x"))): true,

    // All other are NOT OK.
    true: false
  );
}


class SHFL_INFO<bit sync, string mode, string type, bit return_pred> {
  string Suffix = !if(sync, "sync_", "")
                  # mode # "_"
                  # type
                  # !if(return_pred, "p", "");

  string Name = "int_nvvm_shfl_" # Suffix;
  bit withGccBuiltin = !not(return_pred);
  LLVMType OpType = !cond(
    !eq(type,"i32"): llvm_i32_ty,
    !eq(type,"f32"): llvm_float_ty);
  list<LLVMType> RetTy = !if(return_pred, [OpType, llvm_i1_ty], [OpType]);
  list<LLVMType> ArgsTy = !if(sync,
    [llvm_i32_ty, OpType, llvm_i32_ty, llvm_i32_ty],
    [OpType, llvm_i32_ty, llvm_i32_ty]);
}

class NVVM_TCGEN05_LDST_ACCESS_SIZE<string Shape, int Num, string ElemType = "i32"> {
  int shift = !cond(!eq(Shape, "16x128b"): 1,
                    !eq(Shape, "16x256b"): 2,
                    true : 0);

  int veclen = !shl(1, !add(Num, shift));

  int valid = !le(veclen, 128);
  LLVMType type = !cond(!eq(veclen,   1): LLVMType<!cast<ValueType>(ElemType)>,
                        !eq(veclen,   2): LLVMType<!cast<ValueType>("v"#2#ElemType)>,
                        !eq(veclen,   4): LLVMType<!cast<ValueType>("v"#4#ElemType)>,
                        !eq(veclen,   8): LLVMType<!cast<ValueType>("v"#8#ElemType)>,
                        !eq(veclen,  16): LLVMType<!cast<ValueType>("v"#16#ElemType)>,
                        !eq(veclen,  32): LLVMType<!cast<ValueType>("v"#32#ElemType)>,
                        !eq(veclen,  64): LLVMType<!cast<ValueType>("v"#64#ElemType)>,
                        !eq(veclen, 128): LLVMType<!cast<ValueType>("v"#128#ElemType)>,
                        true : llvm_void_ty);
}

class NVVM_TCGEN05_MMA_BASE<string Space, bit IsSparse, string Kind = ""> {
  LLVMType a_operand_type = !if(!eq(Space, "tensor"),
                                  llvm_tmem_ptr_ty, llvm_i64_ty);
  list<LLVMType> common_args = !listconcat(
                                  [llvm_tmem_ptr_ty, // d
                                   a_operand_type,   // a
                                   llvm_i64_ty,      // b
                                   llvm_i32_ty,      // idesc
                                   llvm_i1_ty],      // enable_input_d
                                  !if(!eq(IsSparse, 1), [llvm_tmem_ptr_ty], [])); // spmetadata
  list<IntrinsicProperty> common_intr_props = !listconcat(
                            [IntrArgMemOnly, WriteOnly<ArgIndex<0>>],
                            !if(!eq(Space, "tensor"), [ReadOnly<ArgIndex<1>>], [])
                          );

  string Prefix = "llvm.nvvm.tcgen05.mma";
  string SpSpaceKindStr = !if(IsSparse, ".sp", "")
                          # "." # Space
                          # !if(!ne(Kind, ""), "." # Kind, "");
}

class NVVM_TCGEN05_MMA<bit IsSparse, string Space,
                       bit IsAShift, bit IsScaleInputD>:
        NVVM_TCGEN05_MMA_BASE<Space, IsSparse> {
  string name = Prefix
                # SpSpaceKindStr
                # !if(IsScaleInputD, ".scale_d", "")
                # !if(IsAShift, ".ashift", "");
  string intr_name = IntrinsicName<name>.intr_name;
  string record_name = IntrinsicName<name>.record_name;
}

class NVVM_TCGEN05_MMA_BLOCKSCALE<bit IsSparse, string Space,
                                  string Kind, string ScaleVecSize>:
        NVVM_TCGEN05_MMA_BASE<Space, IsSparse, Kind> {
  string name = Prefix
                # SpSpaceKindStr
                # ".block_scale" # ScaleVecSize;
  string intr_name = IntrinsicName<name>.intr_name;
  string record_name = IntrinsicName<name>.record_name;
}

class NVVM_TCGEN05_MMA_WS<bit IsSparse, string Space, bit IsZeroColMask>:
        NVVM_TCGEN05_MMA_BASE<Space, IsSparse> {
  string name = Prefix # ".ws"
                # SpSpaceKindStr
                # !if(IsZeroColMask, ".zero_col_mask", "");
  string intr_name = IntrinsicName<name>.intr_name;
  string record_name = IntrinsicName<name>.record_name;
}

class NVVM_TCGEN05_MMA_DISABLE_OUTPUT_LANE<bit IsSparse, string Space,
                                           int CtaGroup, bit IsAShift,
                                           bit IsScaleInputD>:
        NVVM_TCGEN05_MMA_BASE<Space, IsSparse> {
  string name = Prefix
                # SpSpaceKindStr
                # !if(IsScaleInputD, ".scale_d", "")
                # ".disable_output_lane.cg" # CtaGroup
                # !if(IsAShift, ".ashift", "");
  string intr_name = IntrinsicName<name>.intr_name;
  string record_name = IntrinsicName<name>.record_name;
}

class NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<string Kind, string ScaleVecSize> {
  bit ret = !cond(
    !and(!eq(Kind, "mxf8f6f4"), !eq(ScaleVecSize, ""))         : true,
    !and(!eq(Kind, "mxf4"),     !eq(ScaleVecSize, ""))         : true,
    !and(!eq(Kind, "mxf4nvf4"), !eq(ScaleVecSize, ".block16")) : true,
    !and(!eq(Kind, "mxf4"),     !eq(ScaleVecSize, ".block32")) : true,
    !and(!eq(Kind, "mxf4nvf4"), !eq(ScaleVecSize, ".block32")) : true,
    !and(!eq(Kind, "mxf8f6f4"), !eq(ScaleVecSize, ".block32")) : true,
    true: false
  );
}

class TexVector<string name, list<LLVMType> types> {
  string Name = name;
  list<LLVMType> Types = types;
}

def TV_I8 : TexVector<"i8", [llvm_i16_ty]>;
def TV_I16 : TexVector<"i16", [llvm_i16_ty]>;
def TV_I32 : TexVector<"i32", [llvm_i32_ty]>;
def TV_I64 : TexVector<"i64", [llvm_i64_ty]>;
def TV_V2I8 : TexVector<"v2i8", !listsplat(llvm_i16_ty, 2)>;
def TV_V2I16 : TexVector<"v2i16", !listsplat(llvm_i16_ty, 2)>;
def TV_V2I32 : TexVector<"v2i32", !listsplat(llvm_i32_ty, 2)>;
def TV_V2I64 : TexVector<"v2i64", !listsplat(llvm_i64_ty, 2)>;
def TV_V4I8 : TexVector<"v4i8", !listsplat(llvm_i16_ty, 4)>;
def TV_V4I16 : TexVector<"v4i16", !listsplat(llvm_i16_ty, 4)>;
def TV_V4I32 : TexVector<"v4i32", !listsplat(llvm_i32_ty, 4)>;


def V4F32 : TexVector<"v4f32", !listsplat(llvm_float_ty, 4)>;
def V4S32 : TexVector<"v4s32", !listsplat(llvm_i32_ty, 4)>;
def V4U32 : TexVector<"v4u32", !listsplat(llvm_i32_ty, 4)>;

class NVVMBuiltin :
  ClangBuiltin<!strconcat("__", !substr(NAME, !size("int_")))> {
    assert !eq(!substr(NAME, 0, !size("int_nvvm_")), "int_nvvm_"),
           "NVVMBuiltin must be a NVVM intrinsic starting with 'int_nvvm_'";
}

class NVVMPureIntrinsic<list<LLVMType> ret_types,
                    list<LLVMType> param_types = [],
                    list<IntrinsicProperty> intr_properties = [],
                    string name = ""> :
  DefaultAttrsIntrinsic<ret_types, param_types,
                        intr_properties # [IntrNoMem, IntrSpeculatable], name> {}

let TargetPrefix = "nvvm" in {

  //
  // PRMT - permute
  //
  def int_nvvm_prmt : NVVMBuiltin,
    PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;

  foreach mode = ["f4e", "b4e"] in
    def int_nvvm_prmt_ # mode :
        PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;

  // Note: these variants also have 2 source operands but only one will ever
  // be used so we eliminate the other operand in the IR (0 is used as the
  // placeholder in the backend).
  foreach mode = ["rc8", "ecl", "ecr", "rc16"] in
    def int_nvvm_prmt_ # mode :
        PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;

  //
  // Nanosleep
  //
  def int_nvvm_nanosleep : NVVMBuiltin,
      DefaultAttrsIntrinsic<[], [llvm_i32_ty],
                            [IntrConvergent, IntrNoMem, IntrHasSideEffects]>;

  //
  // Performance Monitor Events (pm events) intrinsics
  //
  def int_nvvm_pm_event_mask : NVVMBuiltin,
      DefaultAttrsIntrinsic<[], [llvm_i16_ty],
                [IntrConvergent, IntrNoMem, IntrHasSideEffects,
                 ImmArg<ArgIndex<0>>]>;
  //
  // Min Max
  //
  let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
                        IntrNoCreateUndefOrPoison] in {
    foreach operation = ["min", "max"] in {
      def int_nvvm_f # operation # _d : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;

      foreach variant = ["", "_xorsign_abs"] in {
        foreach nan = ["", "_nan"] in {
          foreach ftz = ["", "_ftz"] in {
            def int_nvvm_f # operation # ftz # nan # variant # _f : NVVMBuiltin,
              DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;

            def int_nvvm_f # operation # ftz # nan # variant # _f16 :
              DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty]>;

            def int_nvvm_f # operation # ftz # nan # variant # _f16x2 :
              DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty]>;

            def int_nvvm_f # operation # ftz # nan # variant # _bf16 : NVVMBuiltin,
              DefaultAttrsIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty, llvm_bfloat_ty]>;

            def int_nvvm_f # operation # ftz # nan # variant # _bf16x2 : NVVMBuiltin,
              DefaultAttrsIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty, llvm_v2bf16_ty]>;
          } // ftz
        } // nan
      } // variant
    } // operation
  }

  //
  // Multiplication
  //
  let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative, 
                        IntrNoCreateUndefOrPoison] in {
    foreach sign = ["", "u"] in {
      def int_nvvm_mulhi_ # sign # s : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty]>;

      def int_nvvm_mulhi_ # sign # i : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;

      def int_nvvm_mulhi_ # sign # ll : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty]>;

      def int_nvvm_mul24_ # sign # i : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
    }

    foreach rnd = ["rn", "rz", "rm", "rp"] in {
      foreach ftz = ["", "_ftz"] in
        def int_nvvm_mul_ # rnd # ftz # _f : NVVMBuiltin,
            DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;

      def int_nvvm_mul_ # rnd # _d : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
    }
    
    foreach ftz = ["", "_ftz"] in {
      def int_nvvm_mul_rn # ftz # _sat_f16 : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty]>;

      def int_nvvm_mul_rn # ftz # _sat_v2f16 : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty]>;
    } // ftz
  }

  //
  // Div
  //
  let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
    foreach ftz = ["", "_ftz"] in {
      def int_nvvm_div_approx # ftz # _f : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;

      def int_nvvm_div_full # ftz : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;
    }

    foreach rnd = ["rn", "rz", "rm", "rp"] in {
      foreach ftz = ["", "_ftz"] in
        def int_nvvm_div_ # rnd # ftz # _f : NVVMBuiltin,
            DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;

      def int_nvvm_div_ # rnd # _d : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
    }
  }

  //
  // Sad - Sum of Absolute Differences
  //
  foreach sign = ["", "u"] in {
    def int_nvvm_sad_ # sign # s : NVVMBuiltin,
        PureIntrinsic<[llvm_i16_ty], [llvm_i16_ty, llvm_i16_ty, llvm_i16_ty]>;

    def int_nvvm_sad_ # sign # i : NVVMBuiltin,
        PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;

    def int_nvvm_sad_ # sign # ll : NVVMBuiltin,
        PureIntrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i64_ty]>;
  }

  //
  // Floor  Ceil
  //
  foreach op = ["floor", "ceil"] in {
    foreach ftz = ["", "_ftz"] in
      def int_nvvm_ # op # ftz # _f : NVVMBuiltin,
          PureIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
    def int_nvvm_ # op # _d : NVVMBuiltin,
        PureIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
  }

  //
  // Abs
  //
  foreach ftz = ["", "_ftz"] in
    def int_nvvm_fabs # ftz :
      PureIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;

  //
  // Neg bf16, bf16x2
  //
  def int_nvvm_neg_bf16 : NVVMBuiltin,
    PureIntrinsic<[llvm_bfloat_ty], [llvm_bfloat_ty]>;
  def int_nvvm_neg_bf16x2 : NVVMBuiltin,
    PureIntrinsic<[llvm_v2bf16_ty], [llvm_v2bf16_ty]>;

  //
  // Round
  //
  foreach ftz = ["", "_ftz"] in
    def int_nvvm_round # ftz # _f : NVVMBuiltin,
        PureIntrinsic<[llvm_float_ty], [llvm_float_ty]>;

  def int_nvvm_round_d : NVVMBuiltin,
      PureIntrinsic<[llvm_double_ty], [llvm_double_ty]>;

  //
  // Trunc
  //
  foreach ftz = ["", "_ftz"] in
    def int_nvvm_trunc # ftz # _f : NVVMBuiltin,
        PureIntrinsic<[llvm_float_ty], [llvm_float_ty]>;

  def int_nvvm_trunc_d : NVVMBuiltin,
      PureIntrinsic<[llvm_double_ty], [llvm_double_ty]>;

  //
  // Saturate
  //
  foreach ftz = ["", "_ftz"] in
    def int_nvvm_saturate # ftz # _f : NVVMBuiltin,
        PureIntrinsic<[llvm_float_ty], [llvm_float_ty]>;

  def int_nvvm_saturate_d : NVVMBuiltin,
      PureIntrinsic<[llvm_double_ty], [llvm_double_ty]>;

  //
  // Exp2  Log2
  //
  let IntrProperties = [IntrNoMem] in {
    foreach ftz = ["", "_ftz"] in
      def int_nvvm_ex2_approx # ftz :
          PureIntrinsic<[llvm_anyfloat_ty], [LLVMMatchType<0>]>;

    foreach ftz = ["", "_ftz"] in
      def int_nvvm_lg2_approx # ftz # _f : NVVMBuiltin,
          PureIntrinsic<[llvm_float_ty], [llvm_float_ty]>;

    def int_nvvm_lg2_approx_d : NVVMBuiltin,
        PureIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
  }

  //
  // Sin  Cos
  //
  foreach op = ["sin", "cos"] in
    foreach ftz = ["", "_ftz"] in
      def int_nvvm_ # op # _approx # ftz # _f : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty],
                                [IntrNoMem, IntrNoCreateUndefOrPoison]>;

  //
  // Fma
  //
  foreach variant = ["", "_sat", "_relu"] in {
    foreach ftz = ["", "_ftz"] in {
      def int_nvvm_fma_rn # ftz # variant # _f16 :
        PureIntrinsic<[llvm_half_ty],
          [llvm_half_ty, llvm_half_ty, llvm_half_ty]>;

      def int_nvvm_fma_rn # ftz # variant # _f16x2 :
        PureIntrinsic<[llvm_v2f16_ty],
          [llvm_v2f16_ty, llvm_v2f16_ty, llvm_v2f16_ty]>;
    } // ftz
  } // variant
  
  foreach relu = ["", "_relu"] in { 
    def int_nvvm_fma_rn # relu # _bf16 : NVVMBuiltin,
      PureIntrinsic<[llvm_bfloat_ty],
        [llvm_bfloat_ty, llvm_bfloat_ty, llvm_bfloat_ty]>;

    def int_nvvm_fma_rn # relu # _bf16x2 : NVVMBuiltin,
      PureIntrinsic<[llvm_v2bf16_ty],
        [llvm_v2bf16_ty, llvm_v2bf16_ty, llvm_v2bf16_ty]>;
    
    // oob (out-of-bounds) - clamps the result to 0 if either of the operand is 
    // an OOB NaN value.
    def int_nvvm_fma_rn_oob # relu : PureIntrinsic<[llvm_anyfloat_ty],
        [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>]>;
  } // relu

  foreach rnd = ["_rn", "_rz", "_rm", "_rp"] in {
    foreach ftz = ["", "_ftz"] in {
      foreach sat = ["", "_sat"] in {
        def int_nvvm_fma # rnd # ftz # sat # _f : NVVMBuiltin,
          PureIntrinsic<[llvm_float_ty],
            [llvm_float_ty, llvm_float_ty, llvm_float_ty]>;
      } // sat
    } // ftz
    def int_nvvm_fma # rnd # _d : NVVMBuiltin,
      PureIntrinsic<[llvm_double_ty],
        [llvm_double_ty, llvm_double_ty, llvm_double_ty]>;
  } // rnd

  //
  // Rcp
  //
  let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
    foreach rnd = ["rn", "rz", "rm", "rp"] in {
      foreach ftz = ["", "_ftz"] in
        def int_nvvm_rcp_ # rnd # ftz # _f : NVVMBuiltin,
            DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;

      def int_nvvm_rcp_ # rnd # _d : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
    }

    def int_nvvm_rcp_approx_ftz_f : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
    def int_nvvm_rcp_approx_ftz_d : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
  }

  //
  // Sqrt
  //
  let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
    foreach rnd = ["rn", "rz", "rm", "rp"] in {
      foreach ftz = ["", "_ftz"] in
        def int_nvvm_sqrt_ # rnd # ftz # _f : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;

      def int_nvvm_sqrt_ # rnd # _d : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
    }

    def int_nvvm_sqrt_f : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;

    foreach ftz = ["", "_ftz"] in
      def int_nvvm_sqrt_approx # ftz # _f : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
  }

  //
  // Rsqrt
  //
  let IntrProperties = [IntrNoMem, IntrNoCreateUndefOrPoison] in {
    foreach ftz = ["", "_ftz"] in {
      def int_nvvm_rsqrt_approx # ftz # _f : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
      def int_nvvm_rsqrt_approx # ftz # _d : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
    }
  }

  //
  // Add
  //

  let IntrProperties = [IntrNoMem, IntrSpeculatable, Commutative,
                        IntrNoCreateUndefOrPoison] in {
    foreach rnd = ["_rn", "_rz", "_rm", "_rp"] in {
      foreach ftz = ["", "_ftz"] in {
        foreach sat = ["", "_sat"] in {
          def int_nvvm_add # rnd # ftz # sat # _f : NVVMBuiltin,
            DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]>;
        } // sat
      } // ftz
      def int_nvvm_add # rnd # _d : NVVMBuiltin,
          DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty]>;
    }
    
    foreach ftz = ["", "_ftz"] in {
      def int_nvvm_add_rn # ftz # _sat_f16 : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_half_ty], [llvm_half_ty, llvm_half_ty]>;

      def int_nvvm_add_rn # ftz # _sat_v2f16 : NVVMBuiltin,
        DefaultAttrsIntrinsic<[llvm_v2f16_ty], [llvm_v2f16_ty, llvm_v2f16_ty]>;
        
    } // ftz
  }

  //
  // Dot Product
  //
  foreach a_type = ["s", "u"] in {
    foreach b_type = ["s", "u"] in {
      def int_nvvm_idp4a_ # a_type # _ # b_type :
          PureIntrinsic<[llvm_i32_ty],
              [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
      def int_nvvm_idp2a_ # a_type # _ # b_type :
          PureIntrinsic<[llvm_i32_ty],
            [llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
            [ImmArg<ArgIndex<2>>]>;
    }
  }

  //
  // Funnel-shift
  //
  foreach direction = ["l", "r"] in
    def int_nvvm_fsh # direction # _clamp :
      PureIntrinsic<[llvm_anyint_ty],
                    [LLVMMatchType<0>, LLVMMatchType<0>, LLVMMatchType<0>]>;

  //
  // FLO - Find Leading One
  //
  foreach sign = ["s", "u"] in
    def int_nvvm_flo_ # sign :
      PureIntrinsic<[llvm_i32_ty], [llvm_anyint_ty, llvm_i1_ty],
                    [ImmArg<ArgIndex<1>>]>;

  //
  // szext
  //
  foreach ext = ["sext", "zext"] in
    foreach mode = ["wrap", "clamp"] in
      def int_nvvm_ # ext # _ # mode :
        PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;

  //
  // BMSK - bit mask
  //
  foreach mode = ["wrap", "clamp"] in
    def int_nvvm_bmsk_ # mode :
      PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;

  //
  // FNS - Find the n-th set bit
  //
  def int_nvvm_fns : NVVMBuiltin,
      PureIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;

  //
  // Convert
  //
  // TODO: All these intrinsics are defined as NVVMPureIntrinsic, this attaches the
  //       IntrSpeculatable property to them. Consider if some of these should
  //       have this attribute removed as they may be too expensive.
  //
  def int_nvvm_lohi_i2d : NVVMBuiltin,
      PureIntrinsic<[llvm_double_ty], [llvm_i32_ty, llvm_i32_ty]>;

  def int_nvvm_d2i_lo : NVVMBuiltin,
      PureIntrinsic<[llvm_i32_ty], [llvm_double_ty]>;
  def int_nvvm_d2i_hi : NVVMBuiltin,
      PureIntrinsic<[llvm_i32_ty], [llvm_double_ty]>;

  foreach rnd = ["rn", "rz", "rm", "rp"] in {
    foreach ftz = ["", "_ftz"] in
      def int_nvvm_d2f_ # rnd # ftz : NVVMBuiltin,
          PureIntrinsic<[llvm_float_ty], [llvm_double_ty]>;

    foreach sign = ["", "u"] in {

      def int_nvvm_d2 # sign # i_ # rnd : NVVMBuiltin,
          PureIntrinsic<[llvm_i32_ty], [llvm_double_ty]>;

      def int_nvvm_ # sign # i2d_ # rnd : NVVMBuiltin,
        PureIntrinsic<[llvm_double_ty], [llvm_i32_ty]>;

      foreach ftz = ["", "_ftz"] in
        def int_nvvm_f2 # sign # i_ # rnd # ftz : NVVMBuiltin,
            PureIntrinsic<[llvm_i32_ty], [llvm_float_ty]>;

      def int_nvvm_ # sign # i2f_ # rnd : NVVMBuiltin,
          PureIntrinsic<[llvm_float_ty], [llvm_i32_ty]>;

      foreach ftz = ["", "_ftz"] in
        def int_nvvm_f2 # sign # ll_ # rnd # ftz : NVVMBuiltin,
            PureIntrinsic<[llvm_i64_ty], [llvm_float_ty]>;

      def int_nvvm_d2 # sign # ll_ # rnd : NVVMBuiltin,
        PureIntrinsic<[llvm_i64_ty], [llvm_double_ty]>;

      def int_nvvm_ # sign # ll2f_ # rnd : NVVMBuiltin,
          PureIntrinsic<[llvm_float_ty], [llvm_i64_ty]>;

      def int_nvvm_ # sign # ll2d_ # rnd : NVVMBuiltin,
          PureIntrinsic<[llvm_double_ty], [llvm_i64_ty]>;

    } // sign
  } // rnd

  foreach ftz = ["", "_ftz"] in {
    def int_nvvm_f2h_rn # ftz : NVVMBuiltin,
        PureIntrinsic<[llvm_i16_ty], [llvm_float_ty]>;

    def int_nvvm_bf2h_rn # ftz : NVVMBuiltin,
        PureIntrinsic<[llvm_i16_ty], [llvm_bfloat_ty]>;
  }

  foreach rnd = ["rn", "rz"] in {
    foreach relu = ["", "_relu"] in {
      foreach satfinite = ["", "_satfinite"] in {
        def int_nvvm_ff2bf16x2_ # rnd # relu # satfinite : NVVMBuiltin,
            PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty]>;

        def int_nvvm_ff2f16x2_ # rnd # relu # satfinite : NVVMBuiltin,
            PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty]>;

        def int_nvvm_f2bf16_ # rnd # relu # satfinite : NVVMBuiltin,
            PureIntrinsic<[llvm_bfloat_ty], [llvm_float_ty]>;

        def int_nvvm_f2f16_ # rnd # relu # satfinite : NVVMBuiltin,
            PureIntrinsic<[llvm_half_ty], [llvm_float_ty]>;
      }
    }
  }

  // RS rounding mode (Stochastic Rounding) conversions for f16x2, bf16x2 types
  // The last i32 operand provides the random bits for the conversion
  foreach relu = ["", "_relu"] in {
    foreach satfinite = ["", "_satfinite"] in {
      def int_nvvm_ff2f16x2_rs # relu # satfinite : NVVMBuiltin,
          PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>;

      def int_nvvm_ff2bf16x2_rs # relu # satfinite : NVVMBuiltin,
          PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>;
    }
  }

  foreach satfinite = ["", "_satfinite"] in {
    def int_nvvm_f2tf32_rna # satfinite : NVVMBuiltin,
        PureIntrinsic<[llvm_i32_ty], [llvm_float_ty]>;

    foreach rnd = ["rn", "rz"] in
      foreach relu = ["", "_relu"] in
        def int_nvvm_f2tf32_ # rnd # relu # satfinite : NVVMBuiltin,
            PureIntrinsic<[llvm_i32_ty], [llvm_float_ty]>;
  }

  foreach relu = ["", "_relu"] in {
    def int_nvvm_ff_to_s2f6x2_rn # relu # _satfinite_scale_n2_ue8m0 :
        PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty, llvm_i16_ty]>;

    def int_nvvm_bf16x2_to_s2f6x2_rn # relu # _satfinite_scale_n2_ue8m0 :
        PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty, llvm_i16_ty]>;

    def int_nvvm_s2f6x2_to_bf16x2_rn # relu # _satfinite_scale_n2_ue8m0 :
        PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;

    // No satfinite variants
    def int_nvvm_s2f6x2_to_bf16x2_rn # relu # _scale_n2_ue8m0 :
        PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
  }

  foreach type = ["e4m3x2", "e5m2x2"] in {
    foreach relu = ["", "_relu"] in {
      def int_nvvm_ff_to_ # type # _rn # relu : NVVMBuiltin,
          PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>;

      def int_nvvm_f16x2_to_ # type # _rn # relu : NVVMBuiltin,
          PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>;

      def int_nvvm_ # type # _to_f16x2_rn # relu : NVVMBuiltin,
          PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>;
          
      def int_nvvm_bf16x2_to_ # type # _rn # relu # _satfinite
        : PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
      
      foreach satfinite = ["", "_satfinite"] in {
        def int_nvvm_ # type # _to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0
            : PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
      }
    }
  }

  // RS rounding mode (Stochastic Rounding) conversions for f8x4 types
  // The last i32 operand provides the random bits for the conversion
  foreach type = ["e4m3x4", "e5m2x4"] in {
    foreach relu = ["", "_relu"] in {
      def int_nvvm_f32x4_to_ # type # _rs # relu # _satfinite : NVVMBuiltin,
          PureIntrinsic<[llvm_v4i8_ty], [llvm_v4f32_ty, llvm_i32_ty]>;
    }
  }

  // FP4 conversions.
  foreach relu = ["", "_relu"] in {
    def int_nvvm_ff_to_e2m1x2_rn # relu # _satfinite : NVVMBuiltin,
        PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>;

    def int_nvvm_e2m1x2_to_f16x2_rn # relu : NVVMBuiltin,
        PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>;
        
    def int_nvvm_f16x2_to_e2m1x2_rn # relu # _satfinite
      : PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>;

    def int_nvvm_bf16x2_to_e2m1x2_rn # relu # _satfinite
      : PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
    
    foreach satfinite = ["", "_satfinite"] in {
      def int_nvvm_e2m1x2_to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0
        : PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
    }
  }

  // RS rounding mode (Stochastic Rounding) conversions for f4x4 type
  // The last i32 operand provides the random bits for the conversion
  foreach relu = ["", "_relu"] in {
    def int_nvvm_f32x4_to_e2m1x4_rs # relu # _satfinite : NVVMBuiltin,
        PureIntrinsic<[llvm_i16_ty], [llvm_v4f32_ty, llvm_i32_ty]>;
  }

  // FP6 conversions.
  foreach type = ["e2m3x2", "e3m2x2"] in {
    foreach relu = ["", "_relu"] in {
      def int_nvvm_ff_to_ # type # _rn # relu # _satfinite : NVVMBuiltin,
          PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>;

      def int_nvvm_ # type # _to_f16x2_rn # relu : NVVMBuiltin,
          PureIntrinsic<[llvm_v2f16_ty], [llvm_i16_ty]>;
          
      def int_nvvm_f16x2_to_ # type # _rn # relu # _satfinite
        : PureIntrinsic<[llvm_i16_ty], [llvm_v2f16_ty]>;

      def int_nvvm_bf16x2_to_ # type # _rn # relu # _satfinite
        : PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;
      
      foreach satfinite = ["", "_satfinite"] in {
        def int_nvvm_ # type # _to_bf16x2_rn # relu # satfinite # _scale_n2_ue8m0
            : PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty, llvm_i16_ty]>;
      }
    }
  }

  // RS rounding mode (Stochastic Rounding) conversions for f6x4 types
  // The last i32 operand provides the random bits for the conversion
  foreach type = ["e2m3x4", "e3m2x4"] in {
    foreach relu = ["", "_relu"] in {
      def int_nvvm_f32x4_to_ # type # _rs # relu # _satfinite : NVVMBuiltin,
          PureIntrinsic<[llvm_v4i8_ty], [llvm_v4f32_ty, llvm_i32_ty]>;
    }
  }

  // UE8M0x2 conversions.
  foreach rmode = ["_rz", "_rp"] in {
    foreach satmode = ["", "_satfinite"] in {
      defvar suffix = rmode # satmode;
      def int_nvvm_ff_to_ue8m0x2 # suffix : NVVMBuiltin,
          PureIntrinsic<[llvm_i16_ty], [llvm_float_ty, llvm_float_ty]>;

      def int_nvvm_bf16x2_to_ue8m0x2 # suffix : NVVMBuiltin,
          PureIntrinsic<[llvm_i16_ty], [llvm_v2bf16_ty]>;

    }
  }

  def int_nvvm_ue8m0x2_to_bf16x2 : NVVMBuiltin,
      PureIntrinsic<[llvm_v2bf16_ty], [llvm_i16_ty]>;

  //
  // Bar.Sync
  //
  def int_nvvm_bar_warp_sync : NVVMBuiltin,
      Intrinsic<[], [llvm_i32_ty], [IntrConvergent, IntrNoCallback]>;

  multiclass BarrierCTAIntrinsic<list<LLVMType> ret_types = [],
                                 list<LLVMType> extra_param_types = [],
                                 bit has_all_variant = true> {
    let IntrProperties = [IntrConvergent, IntrNoCallback] in {
      foreach align = ["", "_aligned"] in {
        def align # _count :
          Intrinsic<ret_types, [llvm_i32_ty, llvm_i32_ty] # extra_param_types>;
        if has_all_variant then
          def align # _all :
            Intrinsic<ret_types, [llvm_i32_ty] # extra_param_types>;
      }
    }
  }

  // barrier{.cta}.sync{.aligned}      a{, b};
  // barrier{.cta}.arrive{.aligned}    a, b;
  defm int_nvvm_barrier_cta_sync : BarrierCTAIntrinsic<>;
  defm int_nvvm_barrier_cta_arrive : BarrierCTAIntrinsic<has_all_variant = false>;

  // barrier{.cta}.red.popc{.aligned}.u32  d, a{, b}, {!}c;
  // barrier{.cta}.red.op{.aligned}.pred   p, a{, b}, {!}c;
  // .op = { .and, .or };
  defm int_nvvm_barrier_cta_red_popc : BarrierCTAIntrinsic<[llvm_i32_ty], [llvm_i1_ty]>;
  defm int_nvvm_barrier_cta_red_and : BarrierCTAIntrinsic<[llvm_i1_ty], [llvm_i1_ty]>;
  defm int_nvvm_barrier_cta_red_or : BarrierCTAIntrinsic<[llvm_i1_ty], [llvm_i1_ty]>;

  let IntrProperties = [IntrConvergent, IntrNoCallback] in {
    // barrier.cluster.[wait, arrive, arrive.relaxed]
    def int_nvvm_barrier_cluster_arrive : Intrinsic<[]>;
    def int_nvvm_barrier_cluster_arrive_relaxed : Intrinsic<[]>;
    def int_nvvm_barrier_cluster_wait : Intrinsic<[]>;

    // 'aligned' versions of the above barrier.cluster.* intrinsics
    def int_nvvm_barrier_cluster_arrive_aligned : Intrinsic<[]>;
    def int_nvvm_barrier_cluster_arrive_relaxed_aligned : Intrinsic<[]>;
    def int_nvvm_barrier_cluster_wait_aligned : Intrinsic<[]>;
  }

  //
  // Membar / Fence
  //
  let IntrProperties = [IntrNoCallback] in {
    def int_nvvm_membar_cta : NVVMBuiltin, Intrinsic<[]>;
    def int_nvvm_membar_gl : NVVMBuiltin, Intrinsic<[]>;
    def int_nvvm_membar_sys : NVVMBuiltin, Intrinsic<[]>;
    def int_nvvm_fence_sc_cluster : Intrinsic<[]>;

    // Operation fence
    def int_nvvm_fence_mbarrier_init_release_cluster: Intrinsic<[], [], [],
            "llvm.nvvm.fence.mbarrier_init.release.cluster">;

    // Thread fence
    def int_nvvm_fence_acquire_sync_restrict_space_cluster_scope_cluster :
          Intrinsic<[], [], [],
            "llvm.nvvm.fence.acquire.sync_restrict.space.cluster.scope.cluster">;

    def int_nvvm_fence_release_sync_restrict_space_cta_scope_cluster :
          Intrinsic<[], [], [],
            "llvm.nvvm.fence.release.sync_restrict.space.cta.scope.cluster">;

  //
  // Proxy fence (uni-directional)
  //

  def int_nvvm_fence_proxy_async_generic_acquire_sync_restrict_space_cluster_scope_cluster :
        Intrinsic<[], [], [],
          "llvm.nvvm.fence.proxy.async_generic.acquire.sync_restrict.space.cluster.scope.cluster">;

  def int_nvvm_fence_proxy_async_generic_release_sync_restrict_space_cta_scope_cluster :
        Intrinsic<[], [], [],
          "llvm.nvvm.fence.proxy.async_generic.release.sync_restrict.space.cta.scope.cluster">;

  foreach scope = ["cta", "cluster", "gpu", "sys"] in {

    def int_nvvm_fence_proxy_tensormap_generic_release_ # scope :
          Intrinsic<[], [], [],
          "llvm.nvvm.fence.proxy.tensormap_generic.release." # scope>;

    // The imm-arg 'size' can only be 128.
    def int_nvvm_fence_proxy_tensormap_generic_acquire_ # scope :
          Intrinsic<[], [llvm_ptr_ty, llvm_i32_ty], [],
                    "llvm.nvvm.fence.proxy.tensormap_generic.acquire." # scope> {
      let IntrProperties = [IntrNoCallback, IntrArgMemOnly,
                            ImmArg<ArgIndex<1>>, Range<ArgIndex<1>, 128, 129>];
    }
  }

  //
  // Proxy fence (bi-directional)
  //
    foreach proxykind = ["alias", "async", "async.global", "async.shared_cta",
                          "async.shared_cluster"] in {
      defvar Intr = IntrinsicName<"llvm.nvvm.fence.proxy." # proxykind>;
      def Intr.record_name: Intrinsic<[], [], [], Intr.intr_name>;
    }
  }

//
// Async Copy
//
let IntrProperties = [IntrConvergent, IntrNoCallback] in {
  def int_nvvm_cp_async_mbarrier_arrive : NVVMBuiltin,
      Intrinsic<[], [llvm_ptr_ty]>;
  def int_nvvm_cp_async_mbarrier_arrive_shared : NVVMBuiltin,
      Intrinsic<[], [llvm_shared_ptr_ty]>;
  def int_nvvm_cp_async_mbarrier_arrive_noinc : NVVMBuiltin,
      Intrinsic<[], [llvm_ptr_ty]>;
  def int_nvvm_cp_async_mbarrier_arrive_noinc_shared : NVVMBuiltin,
      Intrinsic<[], [llvm_shared_ptr_ty]>;
}

multiclass CP_ASYNC_SHARED_GLOBAL {
  let IntrProperties = [IntrArgMemOnly, IntrNoCallback, NoAlias<ArgIndex<0>>,
                        NoAlias<ArgIndex<1>>, WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>] in {
    def NAME : Intrinsic<[], [llvm_shared_ptr_ty, llvm_global_ptr_ty]>;
    def _s : Intrinsic<[], [llvm_shared_ptr_ty, llvm_global_ptr_ty, llvm_i32_ty]>;
  }
}

defm int_nvvm_cp_async_ca_shared_global_4 : CP_ASYNC_SHARED_GLOBAL;
defm int_nvvm_cp_async_ca_shared_global_8 : CP_ASYNC_SHARED_GLOBAL;
defm int_nvvm_cp_async_ca_shared_global_16 : CP_ASYNC_SHARED_GLOBAL;
defm int_nvvm_cp_async_cg_shared_global_16 : CP_ASYNC_SHARED_GLOBAL;

def int_nvvm_cp_async_commit_group : NVVMBuiltin, Intrinsic<[]>;

def int_nvvm_cp_async_wait_group : NVVMBuiltin,
    Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;

def int_nvvm_cp_async_wait_all : NVVMBuiltin, Intrinsic<[]>;

// cp.async.bulk variants of the commit/wait group
def int_nvvm_cp_async_bulk_commit_group : Intrinsic<[]>;

def int_nvvm_cp_async_bulk_wait_group :
    Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;

def int_nvvm_cp_async_bulk_wait_group_read :
    Intrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>]>;

// mbarrier
foreach is_shared = [true, false] in {
  defvar mbarrier_ptr_ty = !if(is_shared, llvm_shared_ptr_ty, llvm_ptr_ty);
  defvar shared = !if(is_shared, "_shared", "");

  def int_nvvm_mbarrier_init # shared : NVVMBuiltin,
      Intrinsic<[], [mbarrier_ptr_ty, llvm_i32_ty],
                [IntrConvergent, IntrNoCallback]>;

  def int_nvvm_mbarrier_inval # shared : NVVMBuiltin,
      Intrinsic<[], [mbarrier_ptr_ty],
                [IntrConvergent, IntrWriteMem, IntrArgMemOnly, IntrNoCallback,
                 WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>]>;

  let IntrProperties = [IntrConvergent, IntrNoCallback] in {
    foreach drop = ["", "_drop"] in {
      def int_nvvm_mbarrier_arrive # drop # shared : NVVMBuiltin,
          Intrinsic<[llvm_i64_ty], [mbarrier_ptr_ty]>;
      def int_nvvm_mbarrier_arrive # drop # _noComplete # shared : NVVMBuiltin,
          Intrinsic<[llvm_i64_ty], [mbarrier_ptr_ty, llvm_i32_ty]>;
    }

    def int_nvvm_mbarrier_test_wait # shared : NVVMBuiltin,
        Intrinsic<[llvm_i1_ty], [mbarrier_ptr_ty, llvm_i64_ty]>;
  }
}

def int_nvvm_mbarrier_pending_count : NVVMBuiltin,
    NVVMPureIntrinsic<[llvm_i32_ty], [llvm_i64_ty]>;

// mbarrier.{expect_tx/complete_tx}
foreach op = ["expect_tx", "complete_tx"] in {
  foreach scope = ["scope_cta", "scope_cluster"] in {
    foreach space = ["space_cta", "space_cluster"] in {
      defvar suffix = StrJoin<"_", [op, scope, space]>.ret;
      defvar mbar_addr_ty = !if(!eq(space, "space_cta"),
                            llvm_shared_ptr_ty, llvm_shared_cluster_ptr_ty);

      def int_nvvm_mbarrier_ # suffix :
          Intrinsic<[], [mbar_addr_ty, llvm_i32_ty],
          [IntrConvergent, IntrArgMemOnly, IntrNoCallback]>;
    } // space
  } // scope
} // op

// mbarrier.arrive and mbarrier.arrive.expect_tx
// mbarrier.arrive_drop and mbarrier.arrive_drop.expect_tx
foreach op = ["arrive", "arrive_expect_tx",
              "arrive_drop", "arrive_drop_expect_tx"] in {
  foreach scope = ["scope_cta", "scope_cluster"] in {
    foreach space = ["space_cta", "space_cluster"] in {
      defvar suffix = StrJoin<"_", [scope, space]>.ret;
      defvar mbar_addr_ty = !if(!eq(space, "space_cta"),
                            llvm_shared_ptr_ty, llvm_shared_cluster_ptr_ty);
      defvar args_ty = [mbar_addr_ty, // mbar_address_ptr
                        llvm_i32_ty]; // tx-count

      // mbarriers in shared_cluster space cannot return any value.
      defvar mbar_ret_ty = !if(!eq(space, "space_cta"),
                              [llvm_i64_ty], []<LLVMType>);

      def int_nvvm_mbarrier_ # op # "_" # suffix:
          Intrinsic<mbar_ret_ty, args_ty,
          [IntrConvergent, IntrNoCallback]>;
      def int_nvvm_mbarrier_ # op # "_relaxed_" # suffix :
          Intrinsic<mbar_ret_ty, args_ty,
          [IntrConvergent, IntrArgMemOnly, IntrNoCallback]>;
    } // space
  } // scope
} // op

// mbarrier.{test_wait and try_wait}
foreach op = ["test_wait", "try_wait"] in {
  foreach scope = ["scope_cta", "scope_cluster"] in {
    foreach parity = [true, false] in {
      foreach time_limit = !if(!eq(op, "try_wait"), [true, false], [false]) in {
        defvar base_args = [llvm_shared_ptr_ty]; // mbar_ptr
        defvar parity_args = !if(parity, [llvm_i32_ty], [llvm_i64_ty]);
        defvar tl_args = !if(time_limit, [llvm_i32_ty], []<LLVMType>);
        defvar args = !listconcat(base_args, parity_args, tl_args);
        defvar tmp_op = StrJoin<"_", [op,
                           !if(parity, "parity", ""),
                           !if(time_limit, "tl", "")]>.ret;
        defvar suffix = StrJoin<"_", [scope, "space_cta"]>.ret;

        def int_nvvm_mbarrier_ # tmp_op # "_" # suffix :
            Intrinsic<[llvm_i1_ty], args,
            [IntrConvergent, NoCapture<ArgIndex<0>>, IntrNoCallback]>;
        def int_nvvm_mbarrier_ # tmp_op # "_relaxed_" # suffix :
            Intrinsic<[llvm_i1_ty], args,
            [IntrConvergent, NoCapture<ArgIndex<0>>, IntrNoCallback,
             IntrArgMemOnly, IntrReadMem]>;
      } // tl
    } // parity
  } // scope
} // op

// Generated within nvvm. Use for ldu on sm_20 or later.  Second arg is the
// pointer's alignment.
let IntrProperties = [IntrReadMem, IntrArgMemOnly, IntrNoCallback, IntrWillReturn, NoCapture<ArgIndex<0>>] in {
  def int_nvvm_ldu_global_i : Intrinsic<[llvm_anyint_ty], [llvm_anyptr_ty, llvm_i32_ty]>;
  def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty], [llvm_anyptr_ty, llvm_i32_ty]>;
  def int_nvvm_ldu_global_p : Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty, llvm_i32_ty]>;
}

// Represents an explicit hole in the LLVM IR type system. It may be inserted by
// the compiler in cases where a pointer is of the wrong type. In the backend
// this intrinsic will be folded away and not equate to any instruction. It
// should not be used by any frontend and should only be considered well defined
// when added in the following cases:
//
//  - NVPTXLowerArgs: When wrapping a byval pointer argument to a kernel
//    function to convert the address space from generic (0) to param (101).
//    This accounts for the fact that the parameter symbols will occupy this
//    space when lowered during ISel.
//
def int_nvvm_internal_addrspace_wrap :
  NVVMPureIntrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty],
                [NoUndef<ArgIndex<0>>, NoUndef<RetIndex>]>;

// Move intrinsics, used in nvvm internally

let IntrProperties = [IntrNoMem] in {
  def int_nvvm_move_i16 : DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i16_ty]>;
  def int_nvvm_move_i32 : DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty]>;
  def int_nvvm_move_i64 : DefaultAttrsIntrinsic<[llvm_i64_ty], [llvm_i64_ty]>;
  def int_nvvm_move_float : DefaultAttrsIntrinsic<[llvm_float_ty], [llvm_float_ty]>;
  def int_nvvm_move_double : DefaultAttrsIntrinsic<[llvm_double_ty], [llvm_double_ty]>;
  def int_nvvm_move_ptr : DefaultAttrsIntrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty]>;
}

// For getting the handle from a texture or surface variable
def int_nvvm_texsurf_handle
  : NVVMPureIntrinsic<[llvm_i64_ty], [llvm_metadata_ty, llvm_anyptr_ty]>;
def int_nvvm_texsurf_handle_internal
  : NVVMPureIntrinsic<[llvm_i64_ty], [llvm_anyptr_ty]>;

/// Error / Warn
def int_nvvm_compiler_error : Intrinsic<[], [llvm_anyptr_ty]>;
def int_nvvm_compiler_warn : Intrinsic<[], [llvm_anyptr_ty]>;

def int_nvvm_reflect : NVVMBuiltin, NVVMPureIntrinsic<[llvm_i32_ty], [llvm_ptr_ty]>;

// isspacep.{const, global, local, shared}
foreach space = ["const", "global", "local", "shared", "shared_cluster"] in
  def int_nvvm_isspacep_ # space : NVVMBuiltin,
    NVVMPureIntrinsic<[llvm_i1_ty], [llvm_ptr_ty], [NoCapture<ArgIndex<0>>]>;

//
// Texture Fetch
//
let IntrProperties = [IntrReadMem, IntrNoCallback, IntrNoFree, IntrWillReturn] in {
  foreach is_unified = [true, false] in {
    defvar mode = !if(is_unified, "_unified", "");
    defvar addr_args = !if(is_unified, [llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty]);

    foreach vec = [V4F32, V4S32, V4U32] in {
      foreach is_array = [true, false] in {
        defvar array = !if(is_array, "_array", "");
        defvar array_args = !if(is_array, [llvm_i32_ty], []<LLVMType>);
        defvar base_args = !listconcat(addr_args, array_args);

        def int_nvvm_tex # mode # _1d # array # _ # vec.Name # _s32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_i32_ty, 1)>;
        def int_nvvm_tex # mode # _1d # array # _ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 1)>;
        def int_nvvm_tex # mode # _1d # array # _level_ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 2)>;
        def int_nvvm_tex # mode # _1d # array # _grad_ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 3)>;

        def int_nvvm_tex # mode # _2d # array # _ # vec.Name # _s32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_i32_ty, 2)>;
        def int_nvvm_tex # mode # _2d # array # _ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 2)>;
        def int_nvvm_tex # mode # _2d # array # _level_ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 3)>;
        def int_nvvm_tex # mode # _2d # array # _grad_ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 6)>;

        if !not(is_array) then {
          def int_nvvm_tex # mode # _3d_ # vec.Name # _s32
            : Intrinsic<vec.Types, base_args # !listsplat(llvm_i32_ty, 3)>;
          def int_nvvm_tex # mode # _3d_ # vec.Name # _f32
            : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 3)>;
          def int_nvvm_tex # mode # _3d_level_ # vec.Name # _f32
            : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 4)>;
          def int_nvvm_tex # mode # _3d_grad_ # vec.Name # _f32
            : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 9)>;
        }

        def int_nvvm_tex # mode # _cube # array # _ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 3)>;
        def int_nvvm_tex # mode # _cube # array # _level_ # vec.Name # _f32
          : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 4)>;

        if is_unified then
          def int_nvvm_tex # mode # _cube # array # _grad_ # vec.Name # _f32
            : Intrinsic<vec.Types, base_args # !listsplat(llvm_float_ty, 9)>;
      } // is_array

      foreach comp = ["r", "g", "b", "a"] in {
        def int_nvvm_tld4 # mode # _ # comp # _2d_ # vec.Name # _f32
          : Intrinsic<vec.Types, addr_args # !listsplat(llvm_float_ty, 2)>;
      } // comp
    } // vec
  } // is_unified
} // IntrProperties = [IntrReadMem]

//=== Surface Load
foreach clamp = ["clamp", "trap", "zero"] in {
  foreach vec = [TV_I8, TV_I16, TV_I32, TV_I64,
                TV_V2I8, TV_V2I16, TV_V2I32, TV_V2I64,
                TV_V4I8, TV_V4I16, TV_V4I32] in {

    let IntrProperties = [IntrNoCallback, IntrNoFree, IntrReadMem]
                         # !if(!ne(clamp, "trap"), [IntrWillReturn], []<IntrinsicProperty>) in {

      def int_nvvm_suld_1d_ # vec.Name # _ # clamp
        : Intrinsic<vec.Types,
                    [llvm_i64_ty, llvm_i32_ty]>;

      def int_nvvm_suld_1d_array_ # vec.Name # _ # clamp
        : Intrinsic<vec.Types,
                    [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty]>;

      def int_nvvm_suld_2d_ # vec.Name # _ # clamp
        : Intrinsic<vec.Types,
                    [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty]>;

      def int_nvvm_suld_2d_array_ # vec.Name # _ # clamp
        : Intrinsic<vec.Types,
                    [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;

      def int_nvvm_suld_3d_ # vec.Name # _ # clamp
        : Intrinsic<vec.Types,
                    [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty]>;
    }
  } // vec
} // clamp

//===- Texture Query ------------------------------------------------------===//

foreach query = ["channel_order", "channel_data_type", "width", "height",
                 "depth", "array_size", "num_samples", "num_mipmap_levels"] in
  def int_nvvm_txq_ # query : NVVMBuiltin,
    DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem]>;

//===- Surface Query ------------------------------------------------------===//

foreach query = ["channel_order", "channel_data_type", "width", "height",
                 "depth", "array_size"] in
  def int_nvvm_suq_ # query : NVVMBuiltin,
    DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem]>;

//===- Handle Query -------------------------------------------------------===//

foreach type = ["sampler", "surface", "texture"] in
  def int_nvvm_istypep_ # type : NVVMBuiltin,
    DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem]>;

//===- Surface Stores -----------------------------------------------------===//

multiclass SurfaceStoreIntrinsics<string clamp, TexVector vec> {
  let IntrProperties = [IntrNoCallback, IntrNoFree, IntrWriteMem] #
                       !if(!ne(clamp, "trap"), [IntrWillReturn], []<IntrinsicProperty>) in {
    def _1d_ # vec.Name # _ # clamp : NVVMBuiltin,
        Intrinsic<[], [llvm_i64_ty, llvm_i32_ty] # vec.Types>;

    def _1d_array_ # vec.Name # _ # clamp : NVVMBuiltin,
        Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty] # vec.Types>;

    def _2d_ # vec.Name # _ # clamp : NVVMBuiltin,
        Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty] # vec.Types>;

    def _2d_array_ # vec.Name # _ # clamp : NVVMBuiltin,
        Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty] # vec.Types>;

    def _3d_ # vec.Name # _ # clamp : NVVMBuiltin,
        Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty] # vec.Types>;
  }
}

// Unformatted
foreach clamp = ["clamp", "trap", "zero"] in
  foreach vec = [TV_I8, TV_I16, TV_I32, TV_I64,
                 TV_V2I8, TV_V2I16, TV_V2I32, TV_V2I64,
                 TV_V4I8, TV_V4I16, TV_V4I32] in
    defm int_nvvm_sust_b : SurfaceStoreIntrinsics<clamp, vec>;

// Formatted
foreach vec = [TV_I8, TV_I16, TV_I32,
               TV_V2I8, TV_V2I16, TV_V2I32,
               TV_V4I8, TV_V4I16, TV_V4I32] in
  defm int_nvvm_sust_p : SurfaceStoreIntrinsics<"trap", vec>;

//
// Accessing special registers.
//
class PTXReadSRegIntrinsicNB_r32<list<IntrinsicProperty> properties = [], string name = "">
  : NVVMPureIntrinsic<[llvm_i32_ty], [], [NoUndef<RetIndex>] # properties,
                      !if(!eq(name, ""), "", !strconcat("llvm.nvvm.read.ptx.sreg.", name))>;

class PTXReadSRegIntrinsic_r32<list<IntrinsicProperty> properties = []>
  : PTXReadSRegIntrinsicNB_r32<properties>, NVVMBuiltin;

multiclass PTXReadSRegIntrinsic_v4i32<list<list<IntrinsicProperty>> properties = [[], [], [], []]> {
  assert !eq(!size(properties), 4), "properties must be a list of 4 lists";
  defvar suffixes = ["_x", "_y", "_z", "_w"];
  foreach i = !range(suffixes) in
    def suffixes[i] : PTXReadSRegIntrinsic_r32<properties[i]>;
}

// Same, but without automatic clang builtins. It will be used for
// registers that require particular GPU or PTX version.
multiclass PTXReadSRegIntrinsicNB_v4i32<list<list<IntrinsicProperty>> properties = [[], [], [], []]> {
  assert !eq(!size(properties), 4), "properties must be a list of 4 lists";
  defvar suffixes = ["_x", "_y", "_z", "_w"];
  foreach i = !range(suffixes) in
    def suffixes[i] : PTXReadSRegIntrinsicNB_r32<properties[i]>;
}

// Intrinsics to read registers with non-constant values. E.g. the values that
// do change over the kernel lifetime. Such reads should not be CSE'd.
class PTXReadNCSRegIntrinsic<LLVMType ty>
  : Intrinsic<[ty], [], [IntrInaccessibleMemOnly, IntrNoCallback,
                         IntrNoFree, IntrWillReturn, NoUndef<RetIndex>]>,
    NVVMBuiltin;

defvar MAX_BLOCK_ID_RANGE = [[Range<RetIndex, 0, MAX_BLOCK_SIZE_X>],
                             [Range<RetIndex, 0, MAX_BLOCK_SIZE_Y>],
                             [Range<RetIndex, 0, MAX_BLOCK_SIZE_Z>],
                             [Range<RetIndex, 0, 1>]];

defvar MAX_BLOCK_NID_RANGE = [[Range<RetIndex, 1, !add(MAX_BLOCK_SIZE_X, 1)>],
                              [Range<RetIndex, 1, !add(MAX_BLOCK_SIZE_Y, 1)>],
                              [Range<RetIndex, 1, !add(MAX_BLOCK_SIZE_Z, 1)>],
                              [Range<RetIndex, 0, 1>]];

defvar MAX_GRID_ID_RANGE = [[Range<RetIndex, 0, MAX_GRID_SIZE_X>],
                            [Range<RetIndex, 0, MAX_GRID_SIZE_Y>],
                            [Range<RetIndex, 0, MAX_GRID_SIZE_Z>],
                            [Range<RetIndex, 0, 1>]];

defvar MAX_GRID_NID_RANGE = [[Range<RetIndex, 1, !add(MAX_GRID_SIZE_X, 1)>],
                             [Range<RetIndex, 1, !add(MAX_GRID_SIZE_Y, 1)>],
                             [Range<RetIndex, 1, !add(MAX_GRID_SIZE_Z, 1)>],
                             [Range<RetIndex, 0, 1>]];

defm int_nvvm_read_ptx_sreg_tid : PTXReadSRegIntrinsic_v4i32<MAX_BLOCK_ID_RANGE>;
defm int_nvvm_read_ptx_sreg_ntid : PTXReadSRegIntrinsic_v4i32<MAX_BLOCK_NID_RANGE>;

def int_nvvm_read_ptx_sreg_laneid
  : PTXReadSRegIntrinsic_r32<[Range<RetIndex, 0, WARP_SIZE>]>;

def int_nvvm_read_ptx_sreg_warpid : PTXReadSRegIntrinsic_r32;
def int_nvvm_read_ptx_sreg_nwarpid : PTXReadSRegIntrinsic_r32;

defm int_nvvm_read_ptx_sreg_ctaid : PTXReadSRegIntrinsic_v4i32<MAX_GRID_ID_RANGE>;
defm int_nvvm_read_ptx_sreg_nctaid : PTXReadSRegIntrinsic_v4i32<MAX_GRID_NID_RANGE>;

def int_nvvm_read_ptx_sreg_smid : PTXReadSRegIntrinsic_r32;
def int_nvvm_read_ptx_sreg_nsmid : PTXReadSRegIntrinsic_r32;
def int_nvvm_read_ptx_sreg_gridid : PTXReadSRegIntrinsic_r32;

def int_nvvm_read_ptx_sreg_lanemask_eq : PTXReadSRegIntrinsic_r32;
def int_nvvm_read_ptx_sreg_lanemask_le : PTXReadSRegIntrinsic_r32;
def int_nvvm_read_ptx_sreg_lanemask_lt : PTXReadSRegIntrinsic_r32;
def int_nvvm_read_ptx_sreg_lanemask_ge : PTXReadSRegIntrinsic_r32;
def int_nvvm_read_ptx_sreg_lanemask_gt : PTXReadSRegIntrinsic_r32;

def int_nvvm_read_ptx_sreg_clock : PTXReadNCSRegIntrinsic<llvm_i32_ty>;
def int_nvvm_read_ptx_sreg_clock64 : PTXReadNCSRegIntrinsic<llvm_i64_ty>;

def int_nvvm_read_ptx_sreg_globaltimer : PTXReadNCSRegIntrinsic<llvm_i64_ty>;
def int_nvvm_read_ptx_sreg_globaltimer_lo : PTXReadNCSRegIntrinsic<llvm_i32_ty>;

foreach i = 0...4 in
  def int_nvvm_read_ptx_sreg_pm # i : PTXReadNCSRegIntrinsic<llvm_i32_ty>;

def int_nvvm_read_ptx_sreg_warpsize
  : PTXReadSRegIntrinsic_r32<[Range<RetIndex, WARP_SIZE, !add(WARP_SIZE, 1)>]>;

foreach i = 0...31 in
  def int_nvvm_read_ptx_sreg_envreg # i : PTXReadSRegIntrinsic_r32;

// sm90+, PTX7.8+

// Note: Since clusters are subdivisions of the grid, we conservatively use the
// maximum grid size as an upper bound for the clusterid and cluster_ctaid. In
// practice, the clusterid will likely be much smaller. The CUDA programming
// guide recommends 8 as a maximum portable value and H100s support 16.

defm int_nvvm_read_ptx_sreg_clusterid : PTXReadSRegIntrinsicNB_v4i32<MAX_GRID_ID_RANGE>;
defm int_nvvm_read_ptx_sreg_nclusterid : PTXReadSRegIntrinsicNB_v4i32<MAX_GRID_NID_RANGE>;
defm int_nvvm_read_ptx_sreg_cluster_ctaid : PTXReadSRegIntrinsicNB_v4i32<MAX_GRID_ID_RANGE>;
defm int_nvvm_read_ptx_sreg_cluster_nctaid : PTXReadSRegIntrinsicNB_v4i32<MAX_GRID_NID_RANGE>;

def int_nvvm_read_ptx_sreg_cluster_ctarank : PTXReadSRegIntrinsicNB_r32;
def int_nvvm_read_ptx_sreg_cluster_nctarank : PTXReadSRegIntrinsicNB_r32;

//
// Reserved Shared Memory Intrinsics
//
def int_nvvm_read_ptx_sreg_reserved_smem_offset_begin
  : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_begin">;
def int_nvvm_read_ptx_sreg_reserved_smem_offset_end
  : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_end">;
def int_nvvm_read_ptx_sreg_reserved_smem_offset_cap
  : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_cap">;
def int_nvvm_read_ptx_sreg_reserved_smem_offset_0
  : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_0">;
def int_nvvm_read_ptx_sreg_reserved_smem_offset_1
  : PTXReadSRegIntrinsicNB_r32<name = "reserved_smem_offset_1">;

def int_nvvm_read_ptx_sreg_total_smem_size :
    PTXReadSRegIntrinsicNB_r32<name = "total_smem_size">;
def int_nvvm_read_ptx_sreg_aggr_smem_size :
    PTXReadSRegIntrinsicNB_r32<name = "aggr_smem_size">;
def int_nvvm_read_ptx_sreg_dynamic_smem_size :
    PTXReadSRegIntrinsicNB_r32<name = "dynamic_smem_size">;

//
// SHUFFLE
//
// Non-sync SHFL variants are pure functions of their SSA operands and the active
// thread mask. Memory properties (IntrReadMem + IntrInaccessibleMemOnly + IntrWillReturn) 
// model reading the implicit mask state to ensure correct ordering with other convergent
// operations, preventing unsafe reordering while allowing DCE of truly dead code.
//
// Sync shfl variants have synchronization side effects modeled as
// inaccessible memory accesses to prevent DCE and reordering.
foreach is_sync = [false, true] in {
  defvar CommonIntrinsicProps = [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback];
  defvar IntrinsicProps = !if(is_sync, CommonIntrinsicProps,
                              !listconcat(CommonIntrinsicProps, [IntrReadMem, IntrWillReturn]));
  let IntrProperties = IntrinsicProps in {
    foreach mode = ["up", "down", "bfly", "idx"] in {
      foreach type = ["i32", "f32"] in {
        foreach return_pred = [false, true] in {
          defvar i = SHFL_INFO<is_sync, mode, type, return_pred>;
          if i.withGccBuiltin then
            def i.Name : NVVMBuiltin, Intrinsic<i.RetTy, i.ArgsTy>;
          else
            def i.Name : Intrinsic<i.RetTy, i.ArgsTy>;
        }
      }
    }
  }
}

//
// VOTE
//
let IntrProperties = [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback] in {
  def int_nvvm_vote_all : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i1_ty]>;
  def int_nvvm_vote_any : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i1_ty]>;
  def int_nvvm_vote_uni : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i1_ty]>;
  def int_nvvm_vote_ballot : NVVMBuiltin, Intrinsic<[llvm_i32_ty], [llvm_i1_ty]>;
}

//
// VOTE.SYNC
//
let IntrProperties = [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback] in {
  def int_nvvm_vote_all_sync : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i32_ty, llvm_i1_ty]>;
  def int_nvvm_vote_any_sync : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i32_ty, llvm_i1_ty]>;
  def int_nvvm_vote_uni_sync : NVVMBuiltin, Intrinsic<[llvm_i1_ty], [llvm_i32_ty, llvm_i1_ty]>;
  def int_nvvm_vote_ballot_sync : NVVMBuiltin, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i1_ty]>;
}

//
// ACTIVEMASK
//
def int_nvvm_activemask : NVVMBuiltin,
  Intrinsic<[llvm_i32_ty], [],
            [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback, IntrHasSideEffects]>;

//
// MATCH.SYNC
//
let IntrProperties = [IntrInaccessibleMemOnly, IntrConvergent, IntrNoCallback] in {
  // match.any.sync.b32 mask, value
  def int_nvvm_match_any_sync_i32 : NVVMBuiltin,
    Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;
  // match.any.sync.b64 mask, value
  def int_nvvm_match_any_sync_i64 : NVVMBuiltin,
    Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i64_ty]>;

  // match.all instruction have two variants -- one returns a single value, another
  // returns a pair {value, predicate}. We currently only implement the latter as
  // that's the variant exposed by CUDA API.

  // match.all.sync.b32p mask, value
  def int_nvvm_match_all_sync_i32p :
    Intrinsic<[llvm_i32_ty, llvm_i1_ty], [llvm_i32_ty, llvm_i32_ty]>;
  // match.all.sync.b64p mask, value
  def int_nvvm_match_all_sync_i64p :
    Intrinsic<[llvm_i32_ty, llvm_i1_ty], [llvm_i32_ty, llvm_i64_ty]>;
}
//
// ELECT.SYNC
//
// elect.sync dst|pred, membermask
def int_nvvm_elect_sync :
  DefaultAttrsIntrinsic<[llvm_i32_ty, llvm_i1_ty], [llvm_i32_ty],
                        [IntrInaccessibleMemOnly, IntrConvergent]>;

//
// REDUX.SYNC
//
// redux.sync.op.u32 dst, src, membermask;

let IntrProperties = [IntrConvergent, IntrInaccessibleMemOnly, IntrNoCallback] in {
  foreach op = ["umin", "umax", "add", "min", "max", "and", "xor", "or"] in
    def int_nvvm_redux_sync_ # op : NVVMBuiltin,
      Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty]>;

  // redux.sync.op.{abs}.{NaN}.f32 dst, src, membermask;
  foreach binOp = ["min", "max"] in
    foreach abs = ["", "_abs"] in
      foreach NaN = ["", "_NaN"] in
        def int_nvvm_redux_sync_f # binOp # abs # NaN : NVVMBuiltin,
          Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty]>;
}

//
// WGMMA fence instructions
//
// wgmma.fence.sync.aligned;
def int_nvvm_wgmma_fence_sync_aligned : Intrinsic<[], [], [IntrConvergent]>;

// wgmma.commit_group.sync.aligned;
def int_nvvm_wgmma_commit_group_sync_aligned
  : Intrinsic<[], [], [IntrConvergent], "llvm.nvvm.wgmma.commit_group.sync.aligned">;

// wgmma.wait_group.sync.aligned N;
def int_nvvm_wgmma_wait_group_sync_aligned
  : Intrinsic<[], [llvm_i64_ty], [IntrConvergent, ImmArg<ArgIndex<0>>], "llvm.nvvm.wgmma.wait_group.sync.aligned">;

//
// WMMA instructions
//
// WMMA.LOAD
class NVVM_WMMA_LD<WMMA_REGS Frag, string Layout, int WithStride>
  : Intrinsic<Frag.regs,
              !if(WithStride, [llvm_anyptr_ty, llvm_i32_ty], [llvm_anyptr_ty]),
              [IntrWillReturn, IntrReadMem, IntrArgMemOnly, IntrNoCallback, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>],
              WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.intr_name>;

// WMMA.STORE.D
class NVVM_WMMA_ST<WMMA_REGS Frag, string Layout, int WithStride>
  : Intrinsic<[],
              !listconcat(
                [llvm_anyptr_ty],
                Frag.regs,
                !if(WithStride, [llvm_i32_ty], [])),
              [IntrWriteMem, IntrArgMemOnly, IntrNoCallback, WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>],
              WMMA_NAME_LDST<"store", Frag, Layout, WithStride>.intr_name>;

// Create all load/store variants
foreach layout = ["row", "col"] in {
  foreach stride = [0, 1] in {
    foreach frag = NVVM_MMA_OPS.all_ld_ops in
      if NVVM_WMMA_LDST_SUPPORTED<frag, layout>.ret then
        def WMMA_NAME_LDST<"load", frag, layout, stride>.record_name
             : NVVM_WMMA_LD<frag, layout, stride>;
    foreach frag = NVVM_MMA_OPS.all_st_ops in
      if NVVM_WMMA_LDST_SUPPORTED<frag, layout>.ret then
        def WMMA_NAME_LDST<"store", frag, layout, stride>.record_name
             : NVVM_WMMA_ST<frag, layout, stride>;
  }
}

// WMMA.MMA
class NVVM_MMA<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D>
  : Intrinsic<D.regs,
              !listconcat(A.regs, B.regs, C.regs),
              [IntrNoMem, IntrNoCallback]>;

foreach layout_a = ["row", "col"] in {
  foreach layout_b = ["row", "col"] in {
    foreach satf = [0, 1] in {
      foreach rnd = ["", "rn", "rz", "rm", "rp"] in {
        foreach op = NVVM_MMA_OPS.all_wmma_ops in {
          foreach b1op = NVVM_MMA_B1OPS<op>.ret in {
            if NVVM_WMMA_SUPPORTED<op, layout_a, layout_b, satf, rnd>.ret then {
              def WMMA_NAME<layout_a, layout_b, satf, rnd, b1op,
                                op[0], op[1], op[2], op[3]>.record_name
                : NVVM_MMA<op[0], op[1], op[2], op[3]>;
            }
          } // b1op
        } // op
      } // rnd
    } // satf
  } // layout_b
} // layout_a

foreach layout_a = ["row", "col"] in {
  foreach layout_b = ["row", "col"] in {
    foreach satf = [0, 1] in {
      foreach op = NVVM_MMA_OPS.all_mma_ops in {
        foreach b1op = NVVM_MMA_B1OPS<op>.ret in {
          foreach kind = ["", "kind::f8f6f4"] in {
            if NVVM_MMA_SUPPORTED<op, layout_a, layout_b, kind, satf>.ret then {
                def MMA_NAME<layout_a, layout_b, satf, b1op, kind, op[0], op[1], op[2], op[3]>.record_name
                : NVVM_MMA<op[0], op[1], op[2], op[3]>;
            }
          } // kind
        } // b1op
      } // op
    } // satf
  } // layout_b
} // layout_a

class NVVM_MMA_BLOCK_SCALE<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D>
  : Intrinsic<D.regs,
              !listconcat(A.regs, B.regs, C.regs,
              [
                llvm_i32_ty,              // scale-a-data
                llvm_i16_ty, llvm_i16_ty, // byte-id-a, thread-id-a
                llvm_i32_ty,              // scale-b-data,
                llvm_i16_ty, llvm_i16_ty, // byte-id-b, thread-id-b
              ]),
              [IntrNoMem, IntrNoCallback]>;

foreach kind = ["mxf4", "mxf4nvf4", "mxf8f6f4"] in {
  foreach scale_vec_size = ["", ".scale_1x", ".scale_2x", ".scale_4x"] in {
    foreach stype = ["ue8m0", "ue4m3"] in {
      foreach op = NVVM_MMA_OPS.all_mma_block_scale_ops in {
        if NVVM_MMA_BLOCK_SCALE_SUPPORTED<op, kind, stype, scale_vec_size>.ret then {
          def MMA_BLOCK_SCALE_NAME<kind, stype, scale_vec_size,
                                   op[0], op[1], op[2], op[3]>.record_name
            : NVVM_MMA_BLOCK_SCALE<op[0], op[1], op[2], op[3]>;
        }
      } // op
    } // stype
  } // scale_vec_size
} // kind

// MMA.SP
class NVVM_MMA_SP<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D>
  : Intrinsic<D.regs,
              !listconcat(A.regs, B.regs, C.regs, [llvm_i32_ty], [llvm_i32_ty])> {
    int pos = !size(!listconcat(A.regs, B.regs, C.regs, [llvm_i32_ty]));

    // The range [0;num_threads) is for the sparsity selector that indicates the threads
    // which contribute metadata.
    int num_threads = !if(!or(!and(!eq(A.geom, "m16n8k32"), !eq(A.ptx_elt_type, "bf16")),
                              !and(!eq(A.geom, "m16n8k32"), !eq(A.ptx_elt_type, "f16")),
                              !and(!eq(A.geom, "m16n8k16"), !eq(A.ptx_elt_type, "tf32")),
                              !and(!eq(A.geom, "m16n8k32"), !eq(A.ptx_elt_type, "u8")),
                              !and(!eq(A.geom, "m16n8k32"), !eq(A.ptx_elt_type, "s8")),
                              !and(!eq(A.geom, "m16n8k64"), !eq(A.ptx_elt_type, "u4")),
                              !and(!eq(A.geom, "m16n8k64"), !eq(A.ptx_elt_type, "s4"))),
                          2,
                          !if(!and(!eq(A.geom, "m16n8k64"),
                                  !or(!eq(A.ptx_elt_type, "u8"),
                                      !eq(A.ptx_elt_type, "s8"),
                                      !eq(A.ptx_elt_type, "e4m3"),
                                      !eq(A.ptx_elt_type, "e5m2"),
                                      !eq(A.ptx_elt_type, "e3m2"),
                                      !eq(A.ptx_elt_type, "e2m3"),
                                      !eq(A.ptx_elt_type, "e2m1"))),
                              1,
                              !if(!and(!eq(A.geom, "m16n8k128"),
                                       !or(!eq(A.ptx_elt_type, "s4"),
                                           !eq(A.ptx_elt_type, "u4"))),
                                  1, 4)));
    let IntrProperties = [IntrNoMem, IntrNoCallback, ImmArg<ArgIndex<pos>>,
                          Range<ArgIndex<pos>, 0, num_threads>];
}

foreach metadata = ["sp", "sp::ordered_metadata"] in {
  foreach kind = ["", "kind::f8f6f4"] in {
    foreach satf = [0, 1] in {
      foreach op = NVVM_MMA_OPS.all_mma_sp_ops in {
        if NVVM_MMA_SP_SUPPORTED<op, metadata, kind, satf>.ret then {
          def MMA_SP_NAME<metadata, kind, satf,
                          op[0], op[1], op[2], op[3]>.record_name
            : NVVM_MMA_SP<op[0], op[1], op[2], op[3]>;
        }
      } // op
    } // satf
  } // kind
} // metadata

// MMA.SP BLOCK SCALE
class NVVM_MMA_SP_BLOCK_SCALE<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D>
  : Intrinsic<D.regs,
              !listconcat(A.regs, B.regs, C.regs,
              [
                llvm_i32_ty, // metadata
                llvm_i32_ty, // sparsity selector
                llvm_i32_ty, // scale-a-data
                llvm_i16_ty, llvm_i16_ty, // byte-id-a, thread-id-a
                llvm_i32_ty, // scale-b-data
                llvm_i16_ty, llvm_i16_ty, // byte-id-b, thread-id-b
              ])> {
    int pos = !size(!listconcat(A.regs, B.regs, C.regs, [llvm_i32_ty]));

    // The range [0;num_threads) is for the sparsity selector that indicates the threads
    // which contribute metadata.
    // According to PTX ISA 9.0, the sparsity selector is always 0
    // for sparse MMA block scale instructions
    int num_threads = 1;
    let IntrProperties = [IntrNoMem, IntrNoCallback, ImmArg<ArgIndex<pos>>,
                          Range<ArgIndex<pos>, 0, num_threads>];
}

// According to PTX ISA 9.0
// a_layout = ["row"], b_layout = ["col"], spvariant = ["sp::ordered_metadata"]
foreach kind = ["mxf4", "mxf4nvf4", "mxf8f6f4"] in {
  foreach scale_vec_size = ["", ".scale_1x", ".scale_2x", ".scale_4x"] in {
    foreach stype = ["ue8m0", "ue4m3"] in {
      foreach op = NVVM_MMA_OPS.all_mma_sp_block_scale_ops in {
        if NVVM_MMA_SP_BLOCK_SCALE_SUPPORTED<op, kind, stype, scale_vec_size>.ret then {
          def MMA_SP_BLOCK_SCALE_NAME<kind, stype, scale_vec_size,
                                      op[0], op[1], op[2], op[3]>.record_name
            : NVVM_MMA_SP_BLOCK_SCALE<op[0], op[1], op[2], op[3]>;
        }
      } // op
    } // stype
  } // scale_vec_size
} // kind

// LDMATRIX
class NVVM_LDMATRIX<WMMA_REGS Frag, int Transposed>
  : Intrinsic<Frag.regs, [llvm_anyptr_ty],
              [IntrReadMem, IntrArgMemOnly, IntrNoCallback, IntrConvergent,
               ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>],
              LDMATRIX_NAME<Frag, Transposed>.intr_name>;

foreach transposed = [0, 1] in {
  foreach frag = NVVM_MMA_OPS.all_ldmatrix_ops in {
    if NVVM_LDMATRIX_SUPPORTED<frag, transposed>.ret then {
      def LDMATRIX_NAME<frag, transposed>.record_name
        : NVVM_LDMATRIX<frag, transposed>;
    }
  }
}

// STMATRIX
class NVVM_STMATRIX<WMMA_REGS Frag, int Transposed>
  : Intrinsic<[],
          !listconcat([llvm_anyptr_ty], Frag.regs),
          [IntrWriteMem, IntrArgMemOnly, IntrNoCallback, IntrConvergent,
           WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>],
          STMATRIX_NAME<Frag, Transposed>.intr_name>;

foreach transposed = [0, 1] in {
  foreach frag = NVVM_MMA_OPS.all_stmatrix_ops in {
    if NVVM_STMATRIX_SUPPORTED<frag, transposed>.ret then {
      def STMATRIX_NAME<frag, transposed>.record_name
        : NVVM_STMATRIX<frag, transposed>;
    }
  }
}

// movmatrix
def int_nvvm_movmatrix_sync_aligned_m8n8_trans_b16
  : Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
              [IntrInaccessibleMemOnly, IntrConvergent,
               IntrNoCallback, IntrNoFree]>;

// MAPA
let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>] in {
  def int_nvvm_mapa
    : DefaultAttrsIntrinsic<[llvm_ptr_ty], [llvm_ptr_ty, llvm_i32_ty]>;
  def int_nvvm_mapa_shared_cluster
    : DefaultAttrsIntrinsic<[llvm_shared_cluster_ptr_ty], [llvm_shared_ptr_ty, llvm_i32_ty]>;
}

// GETCTARANK
let IntrProperties = [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>] in {
  def int_nvvm_getctarank
    : DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_ptr_ty]>;
  def int_nvvm_getctarank_shared_cluster
    : DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_shared_ptr_ty]>;
}

def int_nvvm_is_explicit_cluster
  : NVVMPureIntrinsic<[llvm_i1_ty], [], [NoUndef<RetIndex>],
              "llvm.nvvm.is_explicit_cluster">;

// Setmaxnreg inc/dec intrinsics
// The imm-arg should be in the range: 24 <= val <= 256
foreach op = ["dec", "inc"] in
  def int_nvvm_setmaxnreg_ # op # _sync_aligned_u32
    : DefaultAttrsIntrinsic<[], [llvm_i32_ty],
              [IntrConvergent, IntrNoMem, IntrHasSideEffects,
               ImmArg<ArgIndex<0>>, Range<ArgIndex<0>, 24, 257>]>;

// Exit
def int_nvvm_exit : NVVMBuiltin,
    Intrinsic<[], [], [IntrConvergent, IntrInaccessibleMemOnly, IntrNoReturn]>;

class DefaultAttrsIntrinsicFlags<list<LLVMType> ret_types,
                list<LLVMType> param_types,
                list<LLVMType> flags,
                list<IntrinsicProperty> intr_properties,
                string name = "">
  : DefaultAttrsIntrinsic<
        ret_types,
        !listconcat(param_types, flags),
        !listconcat(intr_properties,
                    !foreach(i, !range(flags),
                        ImmArg<ArgIndex<!add(i, !size(param_types))>>)),
        name>;

// TMA Tensor Copy Intrinsics: S2G -> From Shared to Global memory variants
foreach dim = 1...5 in {
  defvar tensor_dim_args = !listsplat(llvm_i32_ty, dim);
  foreach mode = !if(!ge(dim, 3), ["tile", "im2col"], ["tile"]) in {
    def int_nvvm_cp_async_bulk_tensor_s2g_ # mode # _ # dim # d :
      DefaultAttrsIntrinsicFlags<[],
          !listconcat([llvm_shared_ptr_ty,  // src_smem_ptr
                       llvm_ptr_ty],        // tensormap_ptr
                      tensor_dim_args,      // actual tensor dims
                      [llvm_i64_ty]),       // cache_hint
          [llvm_i1_ty],                     // Flag for cache_hint
          [IntrConvergent, ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>]>;

    // Intrinsics for TMA Copy with reduction
    foreach red_op = ["add", "min", "max", "inc", "dec", "and", "or", "xor"] in
      def int_nvvm_cp_async_bulk_tensor_reduce_ # red_op # _ # mode # _ # dim # d :
        DefaultAttrsIntrinsicFlags<[],
            !listconcat([llvm_shared_ptr_ty,  // src_smem_ptr
                         llvm_ptr_ty],        // tensormap_ptr
                         tensor_dim_args,     // actual tensor dims
                        [llvm_i64_ty]),       // cache_hint
          [llvm_i1_ty],                       // Flag for cache_hint
          [IntrConvergent, ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>]>;
  }
}

// TMA S2G tile::scatter4
def int_nvvm_cp_async_bulk_tensor_s2g_tile_scatter4_2d
  : DefaultAttrsIntrinsicFlags<[],
      !listconcat([llvm_shared_ptr_ty,        // src_smem_ptr
                   llvm_ptr_ty],              // tensormap_ptr
                  !listsplat(llvm_i32_ty, 5), // dims
                  [llvm_i64_ty]),             // cache_hint
      [llvm_i1_ty],                           // Flag for cache_hint
      [IntrConvergent, ReadOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>]>;

// TMA Tensor Copy Intrinsics: G2S -> From Global to Shared memory variants
foreach dim = 1...5 in {
  defvar tensor_dim_args = !listsplat(llvm_i32_ty, dim);

  foreach mode = !if(!ge(dim, 3), ["tile", "im2col", "im2col_w", "im2col_w_128"], ["tile"]) in {
    defvar is_im2col = !eq(mode, "im2col");
    defvar is_im2colw = !or(!eq(mode, "im2col_w"), !eq(mode, "im2col_w_128"));

    // For im2col_w/w128 modes, the num_offsets is always 2.
    // For im2col mode, the num_offsets is (dim - 2).
    defvar num_im2col_offsets = !if(is_im2colw, 2, !if(is_im2col, !add(dim, -2), 0));
    defvar im2col_offsets_args = !listsplat(llvm_i16_ty, num_im2col_offsets);

    defvar g2s_params = !listconcat(
                          [llvm_shared_cluster_ptr_ty, // dst_ptr
                           llvm_shared_ptr_ty,  // mbarrier_ptr
                           llvm_ptr_ty],        // tensormap_ptr
                          tensor_dim_args,      // actual tensor dims
                          im2col_offsets_args,  // im2col offsets
                          [llvm_i16_ty,         // cta_mask
                           llvm_i64_ty]);       // cache_hint
    defvar g2s_flags = [llvm_i1_ty,             // Flag for cta_mask
                        llvm_i1_ty,             // Flag for cache_hint
                        llvm_i32_ty];           // Flag for cta_group
    defvar cta_group_idx = !add(
                             !size(g2s_params),
                             !sub(!size(g2s_flags), 1));
    defvar g2s_props = [IntrConvergent,
                        WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>,
                        // Allowed values for cta_group are {0,1,2} i.e [0, 3).
                        Range<ArgIndex<cta_group_idx>, 0, 3>];
    def int_nvvm_cp_async_bulk_tensor_g2s_ # mode # _ # dim # d :
      DefaultAttrsIntrinsicFlags<[], g2s_params, g2s_flags, g2s_props>;

    def int_nvvm_cp_async_bulk_tensor_prefetch_ # mode # _ # dim # d :
      DefaultAttrsIntrinsicFlags<[],
          !listconcat([llvm_ptr_ty],        // tensormap_ptr
                       tensor_dim_args,     // actual tensor dims
                       im2col_offsets_args, // im2col offsets
                      [llvm_i64_ty]),       // cache_hint
          [llvm_i1_ty],                     // Flag for cache_hint
          [IntrConvergent, ReadOnly<ArgIndex<0>>]>;

    def int_nvvm_cp_async_bulk_tensor_g2s_cta_ # mode # _ # dim # d :
      DefaultAttrsIntrinsicFlags<[],
          !listconcat([llvm_shared_ptr_ty,  // dst_ptr
                       llvm_shared_ptr_ty,  // mbarrier_ptr
                       llvm_ptr_ty],        // tensormap_ptr
                       tensor_dim_args,     // actual tensor dims
                       im2col_offsets_args, // im2col offsets
                       [llvm_i64_ty]),      // cache_hint
          [llvm_i1_ty],                     // Flag for cache_hint
          [IntrConvergent, WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>]>;
  }
}

// TMA copy for tile::gather4
def int_nvvm_cp_async_bulk_tensor_g2s_tile_gather4_2d
  : DefaultAttrsIntrinsicFlags<[],
      !listconcat(
      [llvm_shared_cluster_ptr_ty, // dst_shared_cluster_ptr
       llvm_shared_ptr_ty,         // mbarrier_ptr
       llvm_ptr_ty],               // tensormap_ptr
       !listsplat(llvm_i32_ty, 5), // co-ordinates
      [llvm_i16_ty,                // cta_mask
       llvm_i64_ty]),              // cache_hint
      [llvm_i1_ty,                 // Flag for cta_mask
       llvm_i1_ty,                 // Flag for cache_hint
       llvm_i32_ty],               // Flag for cta_group
      [IntrConvergent,
       WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>,
       // Allowed values for cta_group are {0,1,2} i.e [0, 3).
       Range<ArgIndex<12>, 0, 3>]>;

def int_nvvm_cp_async_bulk_tensor_g2s_cta_tile_gather4_2d
  : DefaultAttrsIntrinsicFlags<[],
      !listconcat(
      [llvm_shared_ptr_ty,         // dst_shared_ptr
       llvm_shared_ptr_ty,         // mbarrier_ptr
       llvm_ptr_ty],               // tensormap_ptr
       !listsplat(llvm_i32_ty, 5), // co-ordinates
      [llvm_i64_ty]),              // cache_hint
      [llvm_i1_ty],                // Flag for cache_hint
      [IntrConvergent,
       WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>]>;

// TMA prefetch for tile::gather4
def int_nvvm_cp_async_bulk_tensor_prefetch_tile_gather4_2d
  : DefaultAttrsIntrinsicFlags<[],
      !listconcat([llvm_ptr_ty],               // tensormap_ptr
                  !listsplat(llvm_i32_ty, 5),  // co-ordinates
                  [llvm_i64_ty]),              // cache_hint
      [llvm_i1_ty],                            // Flag for cache_hint
      [IntrConvergent, ReadOnly<ArgIndex<0>>]>;

// Intrinsics for Prefetch and Prefetchu
let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>] in {
  foreach level = ["L1", "L2"] in {
    def int_nvvm_prefetch_ # level : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
    def int_nvvm_prefetch_global_ # level : DefaultAttrsIntrinsic<[], [llvm_global_ptr_ty]>;
    def int_nvvm_prefetch_local_ # level : DefaultAttrsIntrinsic<[], [llvm_local_ptr_ty]>;
  }

  def int_nvvm_prefetch_tensormap : DefaultAttrsIntrinsic<[], [llvm_anyptr_ty]>;

  foreach eviction_priority = ["evict_normal", "evict_last"] in
    def int_nvvm_prefetch_global_L2_ # eviction_priority : DefaultAttrsIntrinsic<[], [llvm_global_ptr_ty]>;

  def int_nvvm_prefetchu_L1 : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
}

// applypriority
let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>,
                      ImmArg<ArgIndex<1>>] in {
  def int_nvvm_applypriority_global_L2_evict_normal
    : DefaultAttrsIntrinsic<[], [llvm_global_ptr_ty, llvm_i64_ty]>;

  def int_nvvm_applypriority_L2_evict_normal
    : DefaultAttrsIntrinsic<[], [llvm_ptr_ty, llvm_i64_ty]>;
}

// discard
let IntrProperties = [NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<1>>, IntrHasSideEffects] in {
  def int_nvvm_discard_global_L2 : DefaultAttrsIntrinsic<[], [llvm_global_ptr_ty, llvm_i64_ty]>;
  def int_nvvm_discard_L2 : DefaultAttrsIntrinsic<[], [llvm_ptr_ty, llvm_i64_ty]>;
}

// Intrinsics for Bulk Copy using TMA (non-tensor)
// From Global to Shared Cluster
def int_nvvm_cp_async_bulk_global_to_shared_cluster
  : DefaultAttrsIntrinsicFlags<[],
      [llvm_shared_cluster_ptr_ty, // dst_shared_cluster_ptr
       llvm_shared_ptr_ty,         // mbarrier_ptr
       llvm_global_ptr_ty,         // src_gmem_ptr
       llvm_i32_ty,                // copy_size
       llvm_i16_ty,                // cta_mask
       llvm_i64_ty],               // cache_hint
      [llvm_i1_ty,                 // Flag for cta_mask
       llvm_i1_ty],                // Flag for cache_hint
      [IntrConvergent, IntrArgMemOnly,
       WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>]>;

// From Global to Shared CTA
def int_nvvm_cp_async_bulk_global_to_shared_cta
  : DefaultAttrsIntrinsicFlags<[],
      [llvm_shared_ptr_ty, // dst_shared_cta_ptr
       llvm_shared_ptr_ty, // mbarrier_ptr
       llvm_global_ptr_ty, // src_gmem_ptr
       llvm_i32_ty,        // copy_size
       llvm_i64_ty],       // cache_hint
      [llvm_i1_ty],        // Flag for cache_hint
      [IntrConvergent, IntrArgMemOnly,
       WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>]>;

// From Shared CTA to Shared Cluster
def int_nvvm_cp_async_bulk_shared_cta_to_cluster
  : DefaultAttrsIntrinsic<[],
      [llvm_shared_cluster_ptr_ty, // dst_shared_cluster_ptr
       llvm_shared_ptr_ty,         // mbarrier_ptr
       llvm_shared_ptr_ty,         // src_smem_ptr
       llvm_i32_ty],               // copy_size
      [IntrConvergent, IntrArgMemOnly,
       WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<2>>]>;

// From Shared CTA to Global memory
def int_nvvm_cp_async_bulk_shared_cta_to_global
  : DefaultAttrsIntrinsicFlags<[],
      [llvm_global_ptr_ty, // dst_gmem_ptr
       llvm_shared_ptr_ty, // src_smem_ptr
       llvm_i32_ty,        // copy_size
       llvm_i64_ty],       // cache_hint
      [llvm_i1_ty],        // Flag for cache_hint
      [IntrConvergent, IntrArgMemOnly,
       WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>]>;

// From Shared CTA to Global memory with bytemask
def int_nvvm_cp_async_bulk_shared_cta_to_global_bytemask
  : DefaultAttrsIntrinsic<[],
      [llvm_global_ptr_ty, // dst_gmem_ptr
       llvm_shared_ptr_ty, // src_smem_ptr
       llvm_i32_ty,        // copy_size
       llvm_i64_ty,        // cache_hint
       llvm_i1_ty,         // Flag for cache_hint
       llvm_i16_ty],       // byte_mask
      [IntrConvergent, IntrArgMemOnly,
       WriteOnly<ArgIndex<0>>, ReadOnly<ArgIndex<1>>,
       ImmArg<ArgIndex<4>>]>;

// Intrinsics for Bulk Copy Prefetch L2
def int_nvvm_cp_async_bulk_prefetch_L2
  : DefaultAttrsIntrinsicFlags<[],
      [llvm_global_ptr_ty, // src_gmem_ptr
       llvm_i32_ty,        // copy_size
       llvm_i64_ty],       // cache_hint
      [llvm_i1_ty],        // Flag for cache_hint
      [IntrConvergent, IntrArgMemOnly,
       NoCapture<ArgIndex<0>>, ReadOnly<ArgIndex<0>>]>;

def int_nvvm_griddepcontrol_launch_dependents : Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>;
def int_nvvm_griddepcontrol_wait : Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>;

//
// Tcgen05 family of Intrinsics
//

// Tcgen05 alloc/dealloc related intrinsics

foreach cta_group = ["cg1", "cg2"] in {
  def int_nvvm_tcgen05_alloc_ # cta_group : Intrinsic<[],
    [llvm_ptr_ty,        // dst_ptr
     llvm_i32_ty] ,      // num_columns
    [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
     WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>]>;

  def int_nvvm_tcgen05_alloc_shared_ # cta_group : Intrinsic<[],
    [llvm_shared_ptr_ty, // dst_ptr
     llvm_i32_ty],       // num_columns
    [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
     WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>]>;

  def int_nvvm_tcgen05_dealloc_ # cta_group : Intrinsic<[],
    [llvm_tmem_ptr_ty,   // tmem_addr
     llvm_i32_ty],       // num_columns
    [IntrConvergent, IntrArgMemOnly,
     NoCapture<ArgIndex<0>>]>;

  def int_nvvm_tcgen05_relinq_alloc_permit_ # cta_group : Intrinsic<[], [],
    [IntrConvergent, IntrInaccessibleMemOnly]>;

  def int_nvvm_tcgen05_commit_ # cta_group : Intrinsic<[],
    [llvm_ptr_ty],        // mbar_ptr
    [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
     NoCapture<ArgIndex<0>>]>;

  def int_nvvm_tcgen05_commit_shared_ # cta_group : Intrinsic<[],
    [llvm_shared_ptr_ty], // mbar_ptr
    [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
     NoCapture<ArgIndex<0>>]>;

  def int_nvvm_tcgen05_commit_mc_ # cta_group : Intrinsic<[],
    [llvm_ptr_ty, llvm_i16_ty], // mbar_ptr, cta_mask
    [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
     NoCapture<ArgIndex<0>>]>;

  def int_nvvm_tcgen05_commit_mc_shared_ # cta_group : Intrinsic<[],
    [llvm_shared_ptr_ty, llvm_i16_ty], // mbar_ptr, cta_mask
    [IntrConvergent, IntrInaccessibleMemOrArgMemOnly,
     NoCapture<ArgIndex<0>>]>;

  def int_nvvm_tcgen05_shift_down_ # cta_group : Intrinsic<[],
    [llvm_tmem_ptr_ty],   // tmem_addr
    [IntrConvergent, IntrArgMemOnly,
     NoCapture<ArgIndex<0>>]>;
}

// Tcgen05 wait_ld/st intrinsics
def int_nvvm_tcgen05_wait_ld : Intrinsic<[], [],
  [IntrConvergent, IntrInaccessibleMemOnly]>;
def int_nvvm_tcgen05_wait_st : Intrinsic<[], [],
  [IntrConvergent, IntrInaccessibleMemOnly]>;

// Tcgen05 Fence intrinsics
def int_nvvm_tcgen05_fence_before_thread_sync : Intrinsic<[], [],
  [IntrNoMem, IntrHasSideEffects]>;
def int_nvvm_tcgen05_fence_after_thread_sync : Intrinsic<[], [],
  [IntrNoMem, IntrHasSideEffects]>;

// Tcgen05 cp intrinsics
foreach cta_group = ["cg1", "cg2"] in {
  foreach src_fmt = ["", "b6x16_p32", "b4x16_p64"] in {
    foreach shape = ["128x256b", "4x256b", "128x128b",
                     "64x128b_warpx2_02_13",
                     "64x128b_warpx2_01_23",
                     "32x128b_warpx4"] in {
      defvar name = "llvm.nvvm.tcgen05.cp."  #
                    StrJoin<".", [shape, src_fmt, cta_group]>.ret;

      defvar intrinsic_name = IntrinsicName<name>;
      def intrinsic_name.record_name : Intrinsic<[],
        [llvm_tmem_ptr_ty,   // tmem_addr
         llvm_i64_ty],       // smem descriptor
        [IntrConvergent, IntrInaccessibleMemOrArgMemOnly, NoCapture<ArgIndex<0>>],
        intrinsic_name.intr_name>;
    }
  }
}

// Tcgen05 ld intrinsics
class NVVM_TCGEN05_LD<string Shape, int Num> :
        Intrinsic<[NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num>.type],
                  !listconcat([llvm_tmem_ptr_ty],
                              !if(!eq(Shape, "16x32bx2"), [llvm_i64_ty], []),
                              [llvm_i1_ty]),
                  !listconcat([IntrConvergent, IntrArgMemOnly, NoCapture<ArgIndex<0>>],
                              !if(!eq(Shape, "16x32bx2"),
                                [ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<2>>],
                                [ImmArg<ArgIndex<1>>]))>;

// Tcgen05 ld.red intrinsics
class NVVM_TCGEN05_LD_RED<string Shape, int Num, string RedValTy> :
        DefaultAttrsIntrinsicFlags<
          [NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num, RedValTy>.type, LLVMType<!cast<ValueType>(RedValTy)>],
                  [llvm_tmem_ptr_ty],
                  !listconcat(!if(!eq(Shape, "16x32bx2"), [llvm_i64_ty], []),
                              [llvm_i32_ty], // RedOp
                              !if(!eq(RedValTy, "f32"), [llvm_i1_ty, llvm_i1_ty], [])), // abs, nan
                  [IntrConvergent, IntrArgMemOnly, NoCapture<ArgIndex<0>>]>;

// Tcgen05 st intrinsics
class NVVM_TCGEN05_ST<string Shape, int Num> :
        Intrinsic<[],
                  !listconcat([llvm_tmem_ptr_ty],
                              !if(!eq(Shape, "16x32bx2"), [llvm_i64_ty], []),
                              [NVVM_TCGEN05_LDST_ACCESS_SIZE<Shape, Num>.type],
                              [llvm_i1_ty]),
                  !listconcat([IntrConvergent, IntrArgMemOnly, NoCapture<ArgIndex<0>>],
                              !if(!eq(Shape, "16x32bx2"),
                                [ImmArg<ArgIndex<1>>, ImmArg<ArgIndex<3>>],
                                [ImmArg<ArgIndex<2>>]))>;

// tcgen05 ld/st intrinsics
foreach shape = ["16x64b", "16x128b", "16x256b", "32x32b", "16x32bx2"] in {
  foreach num = 0...8 in {
    if NVVM_TCGEN05_LDST_ACCESS_SIZE<shape, num>.valid then {
      def int_nvvm_tcgen05_ld_ # shape # _x # !shl(1, num) :
            NVVM_TCGEN05_LD<shape, num>;
      def int_nvvm_tcgen05_st_ # shape # _x # !shl(1, num) :
            NVVM_TCGEN05_ST<shape, num>;
    }
  }
}

// tcgen05.ld.red intrinsics
foreach shape = ["32x32b", "16x32bx2"] in {
  // num starts with x2 (1 << 1) as tcgen05.ld.red does not support x1
  foreach num = 1...8 in {
    foreach ty = ["f32", "i32"] in {
      if NVVM_TCGEN05_LDST_ACCESS_SIZE<shape, num>.valid then {

        def IntrinsicName<"llvm.nvvm.tcgen05.ld.red." # shape
                          # ".x" # !shl(1, num) # "." # ty>.record_name:
              NVVM_TCGEN05_LD_RED<shape, num, ty>;
      } // valid
    } // ty
  } // num
} // shape

//
// Bulk store intrinsics
//
let IntrProperties = [IntrArgMemOnly, IntrWriteMem, WriteOnly<ArgIndex<0>>,
                      NoCapture<ArgIndex<0>>, ImmArg<ArgIndex<2>>] in {
  def int_nvvm_st_bulk :
      DefaultAttrsIntrinsic<[], [llvm_ptr_ty, llvm_i64_ty, llvm_i64_ty]>;

  def int_nvvm_st_bulk_shared_cta :
      DefaultAttrsIntrinsic<[], [llvm_shared_ptr_ty, llvm_i64_ty, llvm_i64_ty]>;
}

//
// clusterlaunchcontorl Intrinsics
//

// clusterlaunchcontrol.try_cancel

def int_nvvm_clusterlaunchcontrol_try_cancel_async_shared
    : DefaultAttrsIntrinsic<[], [llvm_shared_ptr_ty, llvm_shared_ptr_ty],
                [IntrHasSideEffects, IntrArgMemOnly],
                "llvm.nvvm.clusterlaunchcontrol.try_cancel.async.shared">;

def int_nvvm_clusterlaunchcontrol_try_cancel_async_multicast_shared
    : DefaultAttrsIntrinsic<[], [llvm_shared_ptr_ty, llvm_shared_ptr_ty],
                [IntrHasSideEffects, IntrArgMemOnly],
                "llvm.nvvm.clusterlaunchcontrol.try_cancel.async.multicast.shared">;

// clusterlaunchcontrol.query_cancel.is_canceled

def int_nvvm_clusterlaunchcontrol_query_cancel_is_canceled
    : NVVMPureIntrinsic<[llvm_i1_ty], [llvm_i128_ty], [],
                    "llvm.nvvm.clusterlaunchcontrol.query_cancel.is_canceled">;

foreach dim = ["x", "y", "z"] in
  def int_nvvm_clusterlaunchcontrol_query_cancel_get_first_ctaid_ # dim
    : NVVMPureIntrinsic<[llvm_i32_ty], [llvm_i128_ty], [],
                    "llvm.nvvm.clusterlaunchcontrol.query_cancel.get_first_ctaid." # dim>;

//
// tcgen05.mma intrinsics
//

foreach sparse = [0, 1] in {
  foreach space = ["tensor", "shared"] in {
    foreach scale_d = [0, 1] in {
      foreach ashift = !if(!eq(space, "tensor"), [0, 1], [0]) in {
        defvar mma = NVVM_TCGEN05_MMA<sparse, space, ashift, scale_d>;
        defvar args = !listconcat(
          mma.common_args,
          !if(!eq(scale_d, 1), [llvm_i64_ty], []) // scale_d_imm
        );
        defvar flags = [llvm_i32_ty,   // kind
                        llvm_i32_ty,   // cta_group
                        llvm_i32_ty];  // collector_usage_a
        defvar nargs = !size(args);
        defvar scale_d_imm = ArgIndex<!sub(nargs, 1)>;
        defvar scale_d_imm_range = [ImmArg<scale_d_imm>, Range<scale_d_imm, 0, 16>];

        // Check if this is the specific llvm.nvvm.tcgen05.mma.tensor intrinsic.
        defvar is_target_intrinsic = !and(!eq(sparse, 0),
                                          !eq(space, "tensor"),
                                          !eq(scale_d, 0),
                                          !eq(ashift, 0));

        defvar base_properties = !listconcat(
          mma.common_intr_props,
          !if(!eq(scale_d, 1), scale_d_imm_range, []),
          [Range<ArgIndex<nargs>, 0, !if(!eq(scale_d, 1), 2, 4)>, // kind
           Range<ArgIndex<!add(nargs, 1)>, 1, 3>,                 // cta_group
           Range<ArgIndex<!add(nargs, 2)>, 0,
                          !if(!eq(ashift, 1), 2, 4)>              // collector_usage
          ]
        );

        defvar intrinsic_properties = !if(is_target_intrinsic,
          !listconcat(base_properties,
            [ArgInfo<ArgIndex<nargs>, [ArgName<"kind">, ImmArgPrinter<"printTcgen05MMAKind">]>,
             ArgInfo<ArgIndex<!add(nargs, 1)>, [ArgName<"cta_group">]>,
             ArgInfo<ArgIndex<!add(nargs, 2)>, [ArgName<"collector">, ImmArgPrinter<"printTcgen05CollectorUsageOp">]>]),
          base_properties);

        def mma.record_name:
              DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
                mma.intr_name>;
      } // ashift
    } // scale_d
  } // space
} // sparse

//
// tcgen05.mma disable_output_lane intrinsics
//
foreach sparse = [0, 1] in {
  foreach space = ["tensor", "shared"] in {
    foreach cta_group = [1, 2] in {
      foreach scale_d = [0, 1] in {
        foreach ashift = !if(!eq(space, "tensor"), [0, 1], [0]) in {
          defvar mma = NVVM_TCGEN05_MMA_DISABLE_OUTPUT_LANE<
                         sparse, space, cta_group, ashift, scale_d>;
          defvar disable_output_lane_type =
                  !if(!eq(cta_group, 1), llvm_v4i32_ty, llvm_v8i32_ty);
          defvar args = !listconcat(
                          mma.common_args,
                          !if(!eq(scale_d, 1), [llvm_i64_ty], []),
                          [disable_output_lane_type]
                        );
          defvar flags = [llvm_i32_ty,  // kind_flag
                          llvm_i32_ty]; // collector_usage_a_flag
          defvar nargs = !size(args);
          defvar scale_d_imm = ArgIndex<!sub(nargs, 2)>;
          defvar scale_d_imm_range = [ImmArg<scale_d_imm>, Range<scale_d_imm, 0, 16>];
          defvar intrinsic_properties = !listconcat(
                    mma.common_intr_props,
                    !if(!eq(scale_d, 1), scale_d_imm_range, []),
                    [Range<ArgIndex<nargs>, 0, !if(!eq(scale_d, 1), 2, 4)>,
                     Range<ArgIndex<!add(nargs, 1)>, 0, !if(!eq(ashift, 1), 2, 4)>]
                  );

          def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
                                  intrinsic_properties, mma.intr_name>;
        } // ashift
      } // scale_d
    } // cta_group
  } // space
} // sparse

//
// tcgen05.mma block_scale intrinsics
//
foreach sparse = [0, 1] in {
  foreach space = ["tensor", "shared"] in {
    foreach kind = ["mxf8f6f4", "mxf4", "mxf4nvf4"] in {
      foreach scale_vec_size = ["", ".block16", ".block32"] in {
        defvar mma = NVVM_TCGEN05_MMA_BLOCKSCALE<sparse, space, kind, scale_vec_size>;
        defvar args = !listconcat(mma.common_args,
                                  [llvm_tmem_ptr_ty,   // scale_a
                                   llvm_tmem_ptr_ty]); // scale_b
        defvar flags = [llvm_i32_ty,  // cta_group
                        llvm_i32_ty]; // collector_usage_a
        defvar nargs = !size(args);
        defvar cta_group = ArgIndex<nargs>;
        defvar collector_usage = ArgIndex<!add(nargs, 1)>;

        if NVVM_TCGEN05_MMA_BLOCKSCALE_SUPPORTED<kind, scale_vec_size>.ret then {
          def mma.record_name : DefaultAttrsIntrinsicFlags<[], args, flags,
            !listconcat(mma.common_intr_props,
                        [Range<cta_group, 1, 3>,
                         Range<collector_usage, 0, 4>]),
            mma.intr_name>;
        }
      } // scale_vec_size
    } // kind
  } // space
} // sparse

//
// tcgen05.mma ws intrinsics
//
foreach sparse = [0, 1] in {
  foreach space = ["tensor", "shared"] in {
    foreach zero_col_mask = [0, 1] in {
      defvar mma = NVVM_TCGEN05_MMA_WS<sparse, space, zero_col_mask>;
      defvar args = !listconcat(
        mma.common_args,
        !if(!eq(zero_col_mask, 1), [llvm_i64_ty], [])
      );
      defvar flags = [llvm_i32_ty,        // kind
                      llvm_i32_ty,        // collector_buffer_b
                      llvm_i32_ty];       // collector_usage_b_op
      defvar nargs = !size(args);
      defvar intrinsic_properties = !listconcat(
        mma.common_intr_props,
        [Range<ArgIndex<nargs>, 0, 4>,
         Range<ArgIndex<!add(nargs, 1)>, 0, 4>,
         Range<ArgIndex<!add(nargs, 2)>, 0, 4>]
      );

      def mma.record_name:
            DefaultAttrsIntrinsicFlags<[], args, flags, intrinsic_properties,
              mma.intr_name>;
    } // zero_col_mask
  } // space
} // sparse

//
// tensormap.replace intrinsics
//

let IntrProperties = [IntrArgMemOnly, IntrWriteMem, NoCapture<ArgIndex<0>>] in {
  def int_nvvm_tensormap_replace_global_address :
    DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i64_ty], []>;
  def int_nvvm_tensormap_replace_rank :
    DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty], []>;
}

let IntrProperties = [IntrArgMemOnly, ImmArg<ArgIndex<1>>, IntrWriteMem,
                      NoCapture<ArgIndex<0>>] in {
  def int_nvvm_tensormap_replace_global_stride :
    DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty, llvm_i64_ty], []>;
  foreach tmap_field = ["box_dim", "global_dim", "element_stride"] in {
    def int_nvvm_tensormap_replace_ # tmap_field :
      DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty, llvm_i32_ty], []>;
  }
}

def int_nvvm_tensormap_replace_elemtype : 
  DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty], 
    [IntrArgMemOnly, IntrWriteMem, ImmArg<ArgIndex<1>>, NoCapture<ArgIndex<0>>,
     Range<ArgIndex<1>, 0, 16>,
     ArgInfo<ArgIndex<1>, [ArgName<"elemtype">, 
                           ImmArgPrinter<"printTensormapElemType">]>]>;
def int_nvvm_tensormap_replace_interleave_layout :
  DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty], 
    [IntrArgMemOnly, IntrWriteMem, ImmArg<ArgIndex<1>>, NoCapture<ArgIndex<0>>,
     Range<ArgIndex<1>, 0, 3>,
     ArgInfo<ArgIndex<1>, [ArgName<"interleave_layout">, 
                           ImmArgPrinter<"printTensormapInterleaveLayout">]>]>;
def int_nvvm_tensormap_replace_swizzle_mode :
  DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty], 
    [IntrArgMemOnly, IntrWriteMem, ImmArg<ArgIndex<1>>, NoCapture<ArgIndex<0>>,
     Range<ArgIndex<1>, 0, 5>,
     ArgInfo<ArgIndex<1>, [ArgName<"swizzle_mode">, 
                           ImmArgPrinter<"printTensormapSwizzleMode">]>]>;
def int_nvvm_tensormap_replace_swizzle_atomicity :
  DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty], 
    [IntrArgMemOnly, IntrWriteMem, ImmArg<ArgIndex<1>>, NoCapture<ArgIndex<0>>,
     Range<ArgIndex<1>, 0, 4>,
     ArgInfo<ArgIndex<1>, [ArgName<"swizzle_atomicity">, 
                           ImmArgPrinter<"printTensormapSwizzleAtomicity">]>]>;
def int_nvvm_tensormap_replace_fill_mode :
  DefaultAttrsIntrinsic<[], [llvm_anyptr_ty, llvm_i32_ty], 
    [IntrArgMemOnly, IntrWriteMem, ImmArg<ArgIndex<1>>, NoCapture<ArgIndex<0>>,
     Range<ArgIndex<1>, 0, 2>,
     ArgInfo<ArgIndex<1>, [ArgName<"fill_mode">, 
                           ImmArgPrinter<"printTensormapFillMode">]>]>;

} // let TargetPrefix = "nvvm"