[NVPTX] [TableGen] Use new features of TableGen to simplify and clarify

[Paul-C-Anagnostopoulos]

Summary

This patch updates one of the NVPTX TableGen files to use the new features of TableGen.

[Paul-C-Anagnostopoulos]

Oh, let me get rid of the old //// lines and update this patch.

[Paul-C-Anagnostopoulos]

Cleaned out the spurious lines with ////.

[tra]

LGTM.
I'd add NFC to the subject of this patch.

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

60

I'd probably keep this as 0/1. IMO it's more suitable for a bitmask of threads.

[Paul-C-Anagnostopoulos]

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

60

Yes, I'm trying to make that distinction carefully.

[Paul-C-Anagnostopoulos]

I see you accepted this revision. Shall I clean up the other .td files as part of this revision, or do them separately?

[tra]

You can add more cleanup changes here and ping me when they are ready for review.

[Paul-C-Anagnostopoulos]

Cleaned up the other TableGen files.

Not much to do: very clean files already.

[tra]

LGTM.

llvm/lib/Target/NVPTX/NVPTXInstrFormats.td

class NVPTXInst<dag outs, dag ins, string asmstr, list<dag> pattern>
   let Namespace = "NVPTX";
   dag OutOperandList = outs;
   dag InOperandList = ins;
   let AsmString = asmstr;
   let Pattern = pattern;
 
   // TSFlagFields
   bits<4> VecInstType = VecNOP.Value;
-  bit IsSimpleMove = 0;
-  bit IsLoad = 0;
-  bit IsStore = 0;
-
-  bit IsTex = 0;
-  bit IsSust = 0;
-  bit IsSurfTexQuery = 0;
-  bit IsTexModeUnified = 0;
+  bit IsSimpleMove = false;
+  bit IsLoad = false;
+  bit IsStore = false;
+
+  bit IsTex = false;
+  bit IsSust = false;
+  bit IsSurfTexQuery = false;
+  bit IsTexModeUnified = false;
 
   // The following field is encoded as log2 of the vector size minus one,
   // with 0 meaning the operation is not a surface instruction.  For example,
   // if IsSuld == 2, then the instruction is a suld instruction with vector size
   // 2**(2-1) = 2.
   bits<2> IsSuld = 0;
 
-  let TSFlags{3-0}   = VecInstType;
-  let TSFlags{4-4}   = IsSimpleMove;
-  let TSFlags{5-5}   = IsLoad;
-  let TSFlags{6-6}   = IsStore;
-  let TSFlags{7}     = IsTex;
-  let TSFlags{9-8}   = IsSuld;
-  let TSFlags{10}    = IsSust;
-  let TSFlags{11}    = IsSurfTexQuery;
-  let TSFlags{12}    = IsTexModeUnified;
+  let TSFlags{3...0}   = VecInstType;
+  let TSFlags{4...4}   = IsSimpleMove;
+  let TSFlags{5...5}   = IsLoad;
+  let TSFlags{6...6}   = IsStore;
+  let TSFlags{7}       = IsTex;
+  let TSFlags{9...8}   = IsSuld;
+  let TSFlags{10}      = IsSust;
+  let TSFlags{11}      = IsSurfTexQuery;
+  let TSFlags{12}      = IsTexModeUnified;
 }

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

 //===- NVPTXInstrInfo.td - NVPTX Instruction defs -------------*- tblgen-*-===//
 //
 // 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 describes the PTX instructions in TableGen format.
 //
 //===----------------------------------------------------------------------===//
 
 include "NVPTXInstrFormats.td"
 
 // A NOP instruction
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   def NOP : NVPTXInst<(outs), (ins), "", []>;
 }
 
 let OperandType = "OPERAND_IMMEDIATE" in {
   def f16imm : Operand<f16>;
 }
 
 // List of vector specific properties

 //===----------------------------------------------------------------------===//
 // NVPTX Instructions.
 //===----------------------------------------------------------------------===//
 
 //-----------------------------------
 // Type Conversion
 //-----------------------------------
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   // Generate a cvt to the given type from all possible types.  Each instance
   // takes a CvtMode immediate that defines the conversion mode to use.  It can
   // be CvtNONE to omit a conversion mode.
   multiclass CVT_FROM_ALL<string FromName, RegisterClass RC> {
     def _s8 :
       NVPTXInst<(outs RC:$dst),
                 (ins Int16Regs:$src, CvtMode:$mode),
                 !strconcat("cvt${mode:base}${mode:ftz}${mode:sat}.",
: NVPTXInst<(outs RC:$d),
                 (ins RC:$a, Int32Regs:$b, i32imm:$c),
                 !strconcat("bfe.", TyStr, " \t$d, $a, $b, $c;"), []>;
   def rii
     : NVPTXInst<(outs RC:$d),
                 (ins RC:$a, i32imm:$b, i32imm:$c),
                 !strconcat("bfe.", TyStr, " \t$d, $a, $b, $c;"), []>;
 }
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   defm BFE_S32 : BFE<"s32", Int32Regs>;
   defm BFE_U32 : BFE<"u32", Int32Regs>;
   defm BFE_S64 : BFE<"s64", Int64Regs>;
   defm BFE_U64 : BFE<"u64", Int64Regs>;
 }
 
 //-----------------------------------
 // Comparison instructions (setp, set)
 //-----------------------------------
 
 // FIXME: This doesn't cover versions of set and setp that combine with a
 // boolean predicate, e.g. setp.eq.and.b16.
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   multiclass SETP<string TypeStr, RegisterClass RC, Operand ImmCls> {
     def rr :
       NVPTXInst<(outs Int1Regs:$dst), (ins RC:$a, RC:$b, CmpMode:$cmp),
                 !strconcat("setp${cmp:base}${cmp:ftz}.", TypeStr,
                            " \t$dst, $a, $b;"), []>;
     def ri :
       NVPTXInst<(outs Int1Regs:$dst), (ins RC:$a, ImmCls:$b, CmpMode:$cmp),
                 !strconcat("setp${cmp:base}${cmp:ftz}.", TypeStr,
def SETP_f16x2rr :
                 []>,
                 Requires<[useFP16Math]>;
 
 
 // FIXME: This doesn't appear to be correct.  The "set" mnemonic has the form
 // "set.CmpOp{.ftz}.dtype.stype", where dtype is the type of the destination
 // reg, either u32, s32, or f32.  Anyway these aren't used at the moment.
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   multiclass SET<string TypeStr, RegisterClass RC, Operand ImmCls> {
     def rr : NVPTXInst<(outs Int32Regs:$dst),
                        (ins RC:$a, RC:$b, CmpMode:$cmp),
                        !strconcat("set$cmp.", TypeStr, " \t$dst, $a, $b;"), []>;
     def ri : NVPTXInst<(outs Int32Regs:$dst),
                        (ins RC:$a, ImmCls:$b, CmpMode:$cmp),
                        !strconcat("set$cmp.", TypeStr, " \t$dst, $a, $b;"), []>;
     def ir : NVPTXInst<(outs Int32Regs:$dst),

 //-----------------------------------
 // Selection instructions (selp)
 //-----------------------------------
 
 // FIXME: Missing slct
 
 // selp instructions that don't have any pattern matches; we explicitly use
 // them within this file.
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   multiclass SELP<string TypeStr, RegisterClass RC, Operand ImmCls> {
     def rr : NVPTXInst<(outs RC:$dst),
                        (ins RC:$a, RC:$b, Int1Regs:$p),
                        !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"), []>;
     def ri : NVPTXInst<(outs RC:$dst),
                        (ins RC:$a, ImmCls:$b, Int1Regs:$p),
                        !strconcat("selp.", TypeStr, " \t$dst, $a, $b, $p;"), []>;
     def ir : NVPTXInst<(outs RC:$dst),

 def MOV_ADDR : NVPTXInst<(outs Int32Regs:$dst), (ins imem:$a),
                          "mov.u32 \t$dst, $a;",
                          [(set Int32Regs:$dst, (Wrapper tglobaladdr:$a))]>;
 def MOV_ADDR64 : NVPTXInst<(outs Int64Regs:$dst), (ins imem:$a),
                            "mov.u64 \t$dst, $a;",
                            [(set Int64Regs:$dst, (Wrapper tglobaladdr:$a))]>;
 
 // Get pointer to local stack.
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   def MOV_DEPOT_ADDR :    NVPTXInst<(outs Int32Regs:$d), (ins i32imm:$num),
                                      "mov.u32 \t$d, __local_depot$num;", []>;
   def MOV_DEPOT_ADDR_64 : NVPTXInst<(outs Int64Regs:$d), (ins i32imm:$num),
                                     "mov.u64 \t$d, __local_depot$num;", []>;
 }
 
 
 // copyPhysreg is hard-coded in NVPTXInstrInfo.cpp
SDNode<"NVPTXISD::PseudoUseParam", SDTPseudoUseParamProfile,
          [SDNPHasChain, SDNPOutGlue, SDNPInGlue, SDNPSideEffect]>;
 def RETURNNode :
   SDNode<"NVPTXISD::RETURN", SDTCallArgMarkProfile,
          [SDNPHasChain, SDNPSideEffect]>;
 def ProxyReg :
   SDNode<"NVPTXISD::ProxyReg", SDTProxyRegProfile,
          [SDNPHasChain, SDNPOutGlue, SDNPInGlue, SDNPSideEffect]>;
 
-let mayLoad = 1 in {
+let mayLoad = true in {
   class LoadParamMemInst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs regclass:$dst), (ins i32imm:$b),
                   !strconcat("ld.param", opstr, " \t$dst, [retval0+$b];"),
                   []>;
 
   class LoadParamV2MemInst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs regclass:$dst, regclass:$dst2), (ins i32imm:$b),
                   !strconcat("ld.param.v2", opstr,
                              " \t{{$dst, $dst2}}, [retval0+$b];"), []>;
 
   class LoadParamV4MemInst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs regclass:$dst, regclass:$dst2, regclass:$dst3,
                         regclass:$dst4),
                   (ins i32imm:$b),
                   !strconcat("ld.param.v4", opstr,
                              " \t{{$dst, $dst2, $dst3, $dst4}}, [retval0+$b];"),
                   []>;
 }
 
 class LoadParamRegInst<NVPTXRegClass regclass, string opstr> :
       NVPTXInst<(outs regclass:$dst), (ins i32imm:$b),
                 !strconcat("mov", opstr, " \t$dst, retval$b;"),
                 [(set regclass:$dst, (LoadParam (i32 0), (i32 imm:$b)))]>;
 
-let mayStore = 1 in {
+let mayStore = true in {
   class StoreParamInst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs), (ins regclass:$val, i32imm:$a, i32imm:$b),
                   !strconcat("st.param", opstr, " \t[param$a+$b], $val;"),
                   []>;
 
   class StoreParamV2Inst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs), (ins regclass:$val, regclass:$val2,
                                i32imm:$a, i32imm:$b),

 def : Pat<(select Int32Regs:$pred, Float32Regs:$a, Float32Regs:$b),
           (SELP_f32rr Float32Regs:$a, Float32Regs:$b,
           (SETP_b32ri (ANDb32ri Int32Regs:$pred, 1), 1, CmpEQ))>;
 def : Pat<(select Int32Regs:$pred, Float64Regs:$a, Float64Regs:$b),
           (SELP_f64rr Float64Regs:$a, Float64Regs:$b,
           (SETP_b32ri (ANDb32ri Int32Regs:$pred, 1), 1, CmpEQ))>;
 
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   // pack a set of smaller int registers to a larger int register
   def V4I16toI64 : NVPTXInst<(outs Int64Regs:$d),
                              (ins Int16Regs:$s1, Int16Regs:$s2,
                                   Int16Regs:$s3, Int16Regs:$s4),
                              "mov.b64 \t$d, {{$s1, $s2, $s3, $s4}};", []>;
   def V2I16toI32 : NVPTXInst<(outs Int32Regs:$d),
                              (ins Int16Regs:$s1, Int16Regs:$s2),
                              "mov.b32 \t$d, {{$s1, $s2}};", []>;
def I64toV2I32 : NVPTXInst<(outs Int32Regs:$d1, Int32Regs:$d2),
                              (ins Int64Regs:$s),
                              "mov.b64 \t{{$d1, $d2}}, $s;", []>;
   def F64toV2F32 : NVPTXInst<(outs Float32Regs:$d1, Float32Regs:$d2),
                              (ins Float64Regs:$s),
                              "mov.b64 \t{{$d1, $d2}}, $s;", []>;
 
 }
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   // Extract element of f16x2 register. PTX does not provide any way
   // to access elements of f16x2 vector directly, so we need to
   // extract it using a temporary register.
   def F16x2toF16_0 : NVPTXInst<(outs Float16Regs:$dst),
                                (ins Float16x2Regs:$src),
                                "{{ .reg .b16 \t%tmp_hi;\n\t"
                                "  mov.b32 \t{$dst, %tmp_hi}, $src; }}",
                                [(set Float16Regs:$dst,
let hasSideEffects = false in {
   // Split an i32 into two f16
   def SplitI32toF16x2  : NVPTXInst<(outs Float16Regs:$lo, Float16Regs:$hi),
                                    (ins Int32Regs:$src),
                                    "mov.b32 \t{{$lo, $hi}}, $src;",
                                    []>;
 }
 
 // Count leading zeros
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   def CLZr32 : NVPTXInst<(outs Int32Regs:$d), (ins Int32Regs:$a),
                          "clz.b32 \t$d, $a;", []>;
   def CLZr64 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
                          "clz.b64 \t$d, $a;", []>;
 }
 
 // 32-bit has a direct PTX instruction
 def : Pat<(ctlz Int32Regs:$a), (CLZr32 Int32Regs:$a)>;

 // "mov b32reg, {b16imm, b16reg}", so we don't do this optimization.
 def : Pat<(i16 (ctlz Int16Regs:$a)),
           (SUBi16ri (CVT_u16_u32
            (CLZr32 (CVT_u32_u16 Int16Regs:$a, CvtNONE)), CvtNONE), 16)>;
 def : Pat<(i32 (zext (i16 (ctlz Int16Regs:$a)))),
           (SUBi32ri (CLZr32 (CVT_u32_u16 Int16Regs:$a, CvtNONE)), 16)>;
 
 // Population count
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   def POPCr32 : NVPTXInst<(outs Int32Regs:$d), (ins Int32Regs:$a),
                           "popc.b32 \t$d, $a;", []>;
   def POPCr64 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
                           "popc.b64 \t$d, $a;", []>;
 }
 
 // 32-bit has a direct PTX instruction
 def : Pat<(ctpop Int32Regs:$a), (POPCr32 Int32Regs:$a)>;

llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

class PTX {
   // (i32 0) will be XForm'ed to the currently used PTX version.
   dag version = (PTXVerXform (i32 0));
 }
 def ptx : PTX;
 
 // Generates list of n sequential register names.
 // E.g. RegNames<3,"r">.ret -> ["r0", "r1", "r2" ]
 class RegSeq<int n, string prefix> {
-  list<string> ret = !if(n, !listconcat(RegSeq<!add(n,-1), prefix>.ret,
-                                        [prefix # !add(n, -1)]),
+  list<string> ret = !if(n, !listconcat(RegSeq<!sub(n, 1), prefix>.ret,
+                                        [prefix # !sub(n, 1)]),
                             []);
 }
 
 class THREADMASK_INFO<bit sync> {
-  list<bit> ret = !if(sync, [0,1], [0]);
+  list<bit> ret = !if(sync, [0, 1], [0]);
 }
 
 //-----------------------------------
 // Synchronization and shuffle functions
 //-----------------------------------
-let isConvergent = 1 in {
+let isConvergent = true in {
 def INT_BARRIER0 : NVPTXInst<(outs), (ins),
                   "bar.sync \t0;",
       [(int_nvvm_barrier0)]>;
 def INT_BARRIERN : NVPTXInst<(outs), (ins Int32Regs:$src1),
                   "bar.sync \t$src1;",
       [(int_nvvm_barrier_n Int32Regs:$src1)]>;
 def INT_BARRIER : NVPTXInst<(outs), (ins Int32Regs:$src1, Int32Regs:$src2),
                   "bar.sync \t$src1, $src2;",
let Pattern = [!con(
              !subst(outs, set,
              !subst(i32imm, imm, tmp))),
       (set !foreach(tmp, InOperandList,
              !subst(ins, Intr,
              !subst(i32imm, imm, tmp))))
   )];
 }
 
-foreach sync = [0, 1] in {
+foreach sync = [false, true] in {
   foreach mode = ["up", "down", "bfly", "idx"] in {
     foreach regclass = ["i32", "f32"] in {
-      foreach return_pred = [0, 1] in {
-        foreach offset_imm = [0, 1] in {
-          foreach mask_imm = [0, 1] in {
+      foreach return_pred = [false, true] in {
+        foreach offset_imm = [false, true] in {
+          foreach mask_imm = [false, true] in {
             foreach threadmask_imm = THREADMASK_INFO<sync>.ret in {
               def : SHFL_INSTR<sync, mode, regclass, return_pred,
                                offset_imm, mask_imm, threadmask_imm>,
                     Requires<!if(sync, [hasSM30], [hasSM30, hasSHFL])>;
             }
           }
         }
       }
def rr : NVPTXInst<(outs regclass:$dest, Int1Regs:$pred),
               [(set regclass:$dest, Int1Regs:$pred, (IntOp Int32Regs:$mask, regclass:$value))]>,
            Requires<[hasPTX60, hasSM70]>;
 }
 defm MATCH_ALLP_SYNC_32 : MATCH_ALLP_SYNC<Int32Regs, "b32", int_nvvm_match_all_sync_i32p,
                                          i32imm>;
 defm MATCH_ALLP_SYNC_64 : MATCH_ALLP_SYNC<Int64Regs, "b64", int_nvvm_match_all_sync_i64p,
                                          i64imm>;
 
-} // isConvergent = 1
+} // isConvergent = true
 
 //-----------------------------------
 // Explicit Memory Fence Functions
 //-----------------------------------
 class MEMBAR<string StrOp, Intrinsic IntOP> :
               NVPTXInst<(outs), (ins),
             StrOp, [(IntOP)]>;
 
multiclass ATOM2N_impl<string OpStr, string IntTypeStr, string TypeStr,
                        ValueType ImmTy, list<Predicate> Preds> {
   defm : ATOM2P_impl<"atom" # !if(!eq(SpaceStr, "gen"), "", "." # SpaceStr)
                             # !if(!eq(ScopeStr, "gpu"), "", "." # ScopeStr)
                             # "." # OpStr # "." # TypeStr
                             # " \t$result, [$src], $b;",
                      !cast<Intrinsic>(
                             "int_nvvm_atomic_" # OpStr
                             # "_" # SpaceStr # "_" # IntTypeStr
-                            # !if(!eq(ScopeStr,""), "", "_" # ScopeStr)),
+                            # !if(!empty(ScopeStr), "", "_" # ScopeStr)),
                      regclass, ImmType, Imm, ImmTy, Preds>;
 }
 multiclass ATOM3N_impl<string OpStr, string IntTypeStr, string TypeStr,
                        string ScopeStr, string SpaceStr,
                        NVPTXRegClass regclass, Operand ImmType, SDNode Imm,
                        ValueType ImmTy, list<Predicate> Preds> {
   defm : ATOM3P_impl<"atom" # !if(!eq(SpaceStr, "gen"), "", "." # SpaceStr)
                             # !if(!eq(ScopeStr, "gpu"), "", "." # ScopeStr)
                             # "." # OpStr # "." # TypeStr
                             # " \t$result, [$src], $b, $c;",
                      !cast<Intrinsic>(
                             "int_nvvm_atomic_" # OpStr
                             # "_" # SpaceStr # "_" # IntTypeStr
-                            # !if(!eq(ScopeStr,""), "", "_" # ScopeStr)),
+                            # !if(!empty(ScopeStr), "", "_" # ScopeStr)),
                      regclass, ImmType, Imm, ImmTy, Preds>;
 }
 
 // Constructs variants for different address spaces.
 // For now we only need variants for generic space pointers.
 multiclass ATOM2A_impl<string OpStr, string IntTypeStr, string TypeStr,
                        string ScopeStr, NVPTXRegClass regclass, Operand ImmType,
                        SDNode Imm, ValueType ImmTy, list<Predicate> Preds> {

 def : Pat<(int_nvvm_rotate_b32 Int32Regs:$src, (i32 imm:$amt)),
           (ROT32imm_sw Int32Regs:$src, imm:$amt, (SUB_FRM_32 node:$amt))>,
       Requires<[noHWROT32]> ;
 
 def : Pat<(int_nvvm_rotate_b32 Int32Regs:$src, Int32Regs:$amt),
           (ROTL32reg_sw Int32Regs:$src, Int32Regs:$amt)>,
       Requires<[noHWROT32]> ;
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   def GET_LO_INT64 : NVPTXInst<(outs Int32Regs:$dst), (ins Int64Regs:$src),
     !strconcat("{{\n\t",
                ".reg .b32 %dummy;\n\t",
                "mov.b64 \t{$dst,%dummy}, $src;\n\t",
                "}}"),
           []> ;
 
   def GET_HI_INT64 : NVPTXInst<(outs Int32Regs:$dst), (ins Int64Regs:$src),
     !strconcat("{{\n\t",
                ".reg .b32 %dummy;\n\t",
                "mov.b64 \t{%dummy,$dst}, $src;\n\t",
                "}}"),
           []> ;
 }
 
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   def PACK_TWO_INT32
     : NVPTXInst<(outs Int64Regs:$dst), (ins Int32Regs:$lo, Int32Regs:$hi),
                 "mov.b64 \t$dst, {{$lo, $hi}};", []> ;
 }
 
 def : Pat<(int_nvvm_swap_lo_hi_b64 Int64Regs:$src),
           (PACK_TWO_INT32 (GET_HI_INT64 Int64Regs:$src),
                           (GET_LO_INT64 Int64Regs:$src))> ;
 
 // Funnel shift, requires >= sm_32.  Does not trap if amt is out of range, so
 // no side effects.
-let hasSideEffects = 0 in {
+let hasSideEffects = false in {
   def SHF_L_WRAP_B32_IMM
     : NVPTXInst<(outs Int32Regs:$dst),
                 (ins  Int32Regs:$lo, Int32Regs:$hi, i32imm:$amt),
                 "shf.l.wrap.b32 \t$dst, $lo, $hi, $amt;",[]>,
       Requires<[hasHWROT32]>;
 
   def SHF_L_WRAP_B32_REG
     : NVPTXInst<(outs Int32Regs:$dst),

 //-----------------------------------
 // Texture Intrinsics
 //-----------------------------------
 
 // NOTE: For Fermi support, any new texture/surface/sampler intrinsics must be
 // also defined in NVPTXReplaceImageHandles.cpp
 
 // texmode_independent
-let IsTex = 1, IsTexModeUnified = 0 in {
+let IsTex = true, IsTexModeUnified = false in {
 // Texture fetch instructions using handles
 def TEX_1D_F32_S32
   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                     Float32Regs:$b, Float32Regs:$a),
               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
               "tex.1d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
               []>;
 def TEX_1D_F32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
               "tld4.a.2d.v4.u32.f32 \t\\{$v0, $v1, $v2, $v3\\}, "
               "[$t, $s, \\{$x, $y\\}];",
               []>;
 }
 
 
 // texmode_unified
-let IsTex = 1, IsTexModeUnified = 1 in {
+let IsTex = true, IsTexModeUnified = true in {
 // Texture fetch instructions using handles
 def TEX_UNIFIED_1D_F32_S32
   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                     Float32Regs:$b, Float32Regs:$a),
               (ins Int64Regs:$t, Int32Regs:$x),
               "tex.1d.v4.f32.s32 \t\\{$r, $g, $b, $a\\}, [$t, \\{$x\\}];",
               []>;
 def TEX_UNIFIED_1D_F32_F32
: NVPTXInst<(outs Int32Regs:$v0, Int32Regs:$v1,
               "[$t, \\{$x, $y\\}];",
               []>;
 }
 
 
 
 //=== Surface load instructions
 // .clamp variant
-let IsSuld = 1 in {
+let IsSuld = true in {
 def SULD_1D_I8_CLAMP
   : NVPTXInst<(outs Int16Regs:$r),
               (ins Int64Regs:$s, Int32Regs:$x),
               "suld.b.1d.b8.clamp \\{$r\\}, [$s, \\{$x\\}];",
               []>;
 def SULD_1D_I16_CLAMP
   : NVPTXInst<(outs Int16Regs:$r),
               (ins Int64Regs:$s, Int32Regs:$x),
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
               "suld.b.3d.v4.b32.clamp \\{$r, $g, $b, $a\\}, "
               "[$s, \\{$x, $y, $z, $z\\}];",
               []>;
 }
 
 
 // .trap variant
-let IsSuld = 1 in {
+let IsSuld = true in {
 def SULD_1D_I8_TRAP
   : NVPTXInst<(outs Int16Regs:$r),
               (ins Int64Regs:$s, Int32Regs:$x),
               "suld.b.1d.b8.trap \\{$r\\}, [$s, \\{$x\\}];",
               []>;
 def SULD_1D_I16_TRAP
   : NVPTXInst<(outs Int16Regs:$r),
               (ins Int64Regs:$s, Int32Regs:$x),
def SULD_3D_V4I32_TRAP
   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
               "suld.b.3d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
               "[$s, \\{$x, $y, $z, $z\\}];",
               []>;
 }
 
 // .zero variant
-let IsSuld = 1 in {
+let IsSuld = true in {
 def SULD_1D_I8_ZERO
   : NVPTXInst<(outs Int16Regs:$r),
               (ins Int64Regs:$s, Int32Regs:$x),
               "suld.b.1d.b8.zero \\{$r\\}, [$s, \\{$x\\}];",
               []>;
 def SULD_1D_I16_ZERO
   : NVPTXInst<(outs Int16Regs:$r),
               (ins Int64Regs:$s, Int32Regs:$x),
: NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
               "[$s, \\{$x, $y, $z, $z\\}];",
               []>;
 }
 
 //-----------------------------------
 // Texture Query Intrinsics
 //-----------------------------------
 
-let IsSurfTexQuery = 1 in {
+let IsSurfTexQuery = true in {
 def TXQ_CHANNEL_ORDER
   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
               "txq.channel_order.b32 \t$d, [$a];",
               []>;
 def TXQ_CHANNEL_DATA_TYPE
   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
               "txq.channel_data_type.b32 \t$d, [$a];",
               []>;

 def : Pat<(int_nvvm_txq_num_mipmap_levels Int64Regs:$a),
           (TXQ_NUM_MIPMAP_LEVELS Int64Regs:$a)>;
 
 
 //-----------------------------------
 // Surface Query Intrinsics
 //-----------------------------------
 
-let IsSurfTexQuery = 1 in {
+let IsSurfTexQuery = true in {
 def SUQ_CHANNEL_ORDER
   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
               "suq.channel_order.b32 \t$d, [$a];",
               []>;
 def SUQ_CHANNEL_DATA_TYPE
   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
               "suq.channel_data_type.b32 \t$d, [$a];",
               []>;
: NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
               [(set Int1Regs:$d, (int_nvvm_istypep_surface Int64Regs:$a))]>;
 def ISTYPEP_TEXTURE
   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
               "istypep.texref \t$d, $a;",
               [(set Int1Regs:$d, (int_nvvm_istypep_texture Int64Regs:$a))]>;
 
 //===- Surface Stores -----------------------------------------------------===//
 
-let IsSust = 1 in {
+let IsSust = true in {
 // Unformatted
 // .clamp variant
 def SUST_B_1D_B8_CLAMP
   : NVPTXInst<(outs),
               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
               "sust.b.1d.b8.clamp \t[$s, \\{$x\\}], \\{$r\\};",
               []>;
 def SUST_B_1D_B16_CLAMP
NVPTXRegClass regclass = !cond(
     !eq(ptx_elt_type, "s32") : Int32Regs,
     !eq(ptx_elt_type, "s8") : Int32Regs,
     !eq(ptx_elt_type, "u8") : Int32Regs,
     !eq(ptx_elt_type, "s4") : Int32Regs,
     !eq(ptx_elt_type, "u4") : Int32Regs,
     !eq(ptx_elt_type, "b1") : Int32Regs);
 
   // Instruction input/output arguments for the fragment.
-  list<NVPTXRegClass> ptx_regs = !foreach(tmp, regs, regclass);
+  list<NVPTXRegClass> ptx_regs = !listsplat(regclass, !size(regs));
 
   // List of register names for the fragment -- ["ra0", "ra1",...]
   list<string> reg_names = RegSeq<!size(ptx_regs), "r"#frag>.ret;
 
   // Generates "{{$r0, $r1,.... $rN-1}}" for use in asm string construction.
   string regstring = "{{$" # !interleave(reg_names, ", $") # "}}";
 
   // Predicates for particular fragment variant. Technically those are
class WMMA_LOAD<WMMA_REGINFO Frag, string Layout, string Space, bit WithStride,
   : WMMA_INSTR<WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.record,
                               [!con((ins SrcOp:$src),
                                     !if(WithStride, (ins Int32Regs:$ldm), (ins)))]>,
     Requires<Frag.Predicates> {
   // Load/store intrinsics are overloaded on pointer's address space.
   // To match the right intrinsic, we need to build AS-constrained PatFrag.
   // Operands is a dag equivalent in shape to Args, but using (ops node:$name, .....).
   dag PFOperands = !if(WithStride, (ops node:$src, node:$ldm), (ops node:$src));
+  dag PFOperandsIntr = !if(WithStride, (Intr node:$src, node:$ldm), (Intr node:$src));
   // Build PatFrag that only matches particular address space.
   PatFrag IntrFrag = PatFrag<PFOperands,
-                             !foreach(tmp, PFOperands, !subst(ops, Intr, tmp)),
+                             PFOperandsIntr,
                              !cond(!eq(Space, ".shared"): AS_match.shared,
                                    !eq(Space, ".global"): AS_match.global,
-                                   1: AS_match.generic)>;
+                                   true: AS_match.generic)>;
   // Build AS-constrained pattern.
   let IntrinsicPattern = BuildPatternPF<IntrFrag, Args>.ret;
 
   let OutOperandList = Frag.Outs;
   let InOperandList = !con(Args, (ins MmaCode:$ptx));
   let AsmString = "wmma.load."
                   # Frag.frag
                   # ".sync"
: WMMA_INSTR<WMMA_NAME_LDST<"store", Frag, Layout, WithStride>.record,
                      Frag.Ins,
                      !if(WithStride, (ins Int32Regs:$ldm), (ins)))]>,
     Requires<Frag.Predicates> {
 
   // Load/store intrinsics are overloaded on pointer's address space.
   // To match the right intrinsic, we need to build AS-constrained PatFrag.
   // Operands is a dag equivalent in shape to Args, but using (ops node:$name, .....).
   dag PFOperands = !con((ops node:$dst),
-                        !dag(ops, !foreach(tmp, Frag.regs, node), Frag.reg_names),
+                        !dag(ops, !listsplat(node, !size(Frag.regs)), Frag.reg_names),
                         !if(WithStride, (ops node:$ldm), (ops)));
   // Build PatFrag that only matches particular address space.
   PatFrag IntrFrag = PatFrag<PFOperands,
                              !foreach(tmp, PFOperands, !subst(ops, Intr, tmp)),
                              !cond(!eq(Space, ".shared"): AS_match.shared,
                                    !eq(Space, ".global"): AS_match.global,
-                                   1: AS_match.generic)>;
+                                   true: AS_match.generic)>;
   // Build AS-constrained pattern.
   let IntrinsicPattern = BuildPatternPF<IntrFrag, Args>.ret;
 
   let InOperandList  = !con(Args, (ins MmaCode:$ptx));
   let OutOperandList = (outs);
   let AsmString = "wmma.store.d.sync"
                   # "${ptx:aligned}"
                   # "." # Layout
                   # "." # Frag.geom
                   # Space
                   # "." # Frag.ptx_elt_type
                   # " \t[$dst],"
                   # Frag.regstring
                   # !if(WithStride, ", $ldm", "")
                   # ";";
 }
 
 // Create all load/store variants
 defset list<WMMA_INSTR> MMA_LDSTs  = {
   foreach layout = ["row", "col"] in {
-    foreach stride = [0, 1] in {
+    foreach stride = [false, true] in {
       foreach space = [".global", ".shared", ""] in {
         foreach addr = [imem, Int32Regs, Int64Regs, MEMri, MEMri64] in {
           foreach frag = NVVM_MMA_OPS.all_ld_ops in
             foreach _ = NVVM_MMA_SUPPORTED<[frag], layout>.ret in
               def : WMMA_LOAD<WMMA_REGINFO<frag>, layout, space, stride, addr>;
           foreach frag = NVVM_MMA_OPS.all_st_ops in
             foreach _ = NVVM_MMA_SUPPORTED<[frag], layout>.ret in
               def : WMMA_STORE_D<WMMA_REGINFO<frag>, layout, space, stride, addr>;

llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td

 def VRFrame         : NVPTXReg<"%SP">;
 def VRFrameLocal    : NVPTXReg<"%SPL">;
 
 // Special Registers used as the stack
 def VRDepot  : NVPTXReg<"%Depot">;
 
 // We use virtual registers, but define a few physical registers here to keep
 // SDAG and the MachineInstr layers happy.
-foreach i = 0-4 in {
+foreach i = 0...4 in {
   def P#i  : NVPTXReg<"%p"#i>;  // Predicate
   def RS#i : NVPTXReg<"%rs"#i>; // 16-bit
   def R#i  : NVPTXReg<"%r"#i>;  // 32-bit
   def RL#i : NVPTXReg<"%rd"#i>; // 64-bit
   def H#i  : NVPTXReg<"%h"#i>;  // 16-bit float
   def HH#i : NVPTXReg<"%hh"#i>; // 2x16-bit float
   def F#i  : NVPTXReg<"%f"#i>;  // 32-bit float
   def FL#i : NVPTXReg<"%fd"#i>; // 64-bit float
 
   // Arguments
   def ia#i : NVPTXReg<"%ia"#i>;
   def la#i : NVPTXReg<"%la"#i>;
   def fa#i : NVPTXReg<"%fa"#i>;
   def da#i : NVPTXReg<"%da"#i>;
 }
 
-foreach i = 0-31 in {
+foreach i = 0...31 in {
   def ENVREG#i : NVPTXReg<"%envreg"#i>;
 }
 
 //===----------------------------------------------------------------------===//
 //  Register classes
 //===----------------------------------------------------------------------===//
 def Int1Regs : NVPTXRegClass<[i1], 8, (add (sequence "P%u", 0, 4))>;
 def Int16Regs : NVPTXRegClass<[i16], 16, (add (sequence "RS%u", 0, 4))>;