Index: llvm/lib/Target/NVPTX/NVPTXIntrinsics.td =================================================================== --- llvm/lib/Target/NVPTX/NVPTXIntrinsics.td +++ llvm/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -51,19 +51,19 @@ // Generates list of n sequential register names. // E.g. RegNames<3,"r">.ret -> ["r0", "r1", "r2" ] class RegSeq { - list ret = !if(n, !listconcat(RegSeq.ret, - [prefix # !add(n, -1)]), + list ret = !if(n, !listconcat(RegSeq.ret, + [prefix # !sub(n, 1)]), []); } class THREADMASK_INFO { - list ret = !if(sync, [0,1], [0]); + list ret = !if(sync, [false,true], [false]); } //----------------------------------- // Synchronization and shuffle functions //----------------------------------- -let isConvergent = 1 in { +let isConvergent = true in { def INT_BARRIER0 : NVPTXInst<(outs), (ins), "bar.sync \t0;", [(int_nvvm_barrier0)]>; @@ -173,12 +173,12 @@ )]; } -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.ret in { def : SHFL_INSTR, @@ -274,7 +274,7 @@ defm MATCH_ALLP_SYNC_64 : MATCH_ALLP_SYNC; -} // isConvergent = 1 +} // isConvergent = true //----------------------------------- // Explicit Memory Fence Functions @@ -1548,7 +1548,8 @@ !cast( "int_nvvm_atomic_" # OpStr # "_" # SpaceStr # "_" # IntTypeStr - # !if(!eq(ScopeStr,""), "", "_" # ScopeStr)), + # !if(!empty(ScopeStr), "", "_" # ScopeStr)), +//// # !if(!eq(ScopeStr,""), "", "_" # ScopeStr)), regclass, ImmType, Imm, ImmTy, Preds>; } multiclass ATOM3N_impl( "int_nvvm_atomic_" # OpStr # "_" # SpaceStr # "_" # IntTypeStr - # !if(!eq(ScopeStr,""), "", "_" # ScopeStr)), + # !if(!empty(ScopeStr), "", "_" # ScopeStr)), +//// # !if(!eq(ScopeStr,""), "", "_" # ScopeStr)), regclass, ImmType, Imm, ImmTy, Preds>; } @@ -2131,7 +2133,7 @@ (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", @@ -2147,7 +2149,7 @@ []> ; } -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}};", []> ; @@ -2159,7 +2161,7 @@ // 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), @@ -2242,7 +2244,7 @@ // 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, @@ -2925,7 +2927,7 @@ // 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, @@ -3610,7 +3612,7 @@ //=== 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), @@ -3922,7 +3924,7 @@ // .trap variant -let IsSuld = 1 in { +let IsSuld = true in { def SULD_1D_I8_TRAP : NVPTXInst<(outs Int16Regs:$r), (ins Int64Regs:$s, Int32Regs:$x), @@ -4233,7 +4235,7 @@ } // .zero variant -let IsSuld = 1 in { +let IsSuld = true in { def SULD_1D_I8_ZERO : NVPTXInst<(outs Int16Regs:$r), (ins Int64Regs:$s, Int32Regs:$x), @@ -4547,7 +4549,7 @@ // 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];", @@ -4604,7 +4606,7 @@ // 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];", @@ -4663,7 +4665,7 @@ //===- Surface Stores -----------------------------------------------------===// -let IsSust = 1 in { +let IsSust = true in { // Unformatted // .clamp variant def SUST_B_1D_B8_CLAMP @@ -7361,7 +7363,8 @@ !eq(ptx_elt_type, "b1") : Int32Regs); // Instruction input/output arguments for the fragment. - list ptx_regs = !foreach(tmp, regs, regclass); + list ptx_regs = !listsplat(regclass, !size(regs)); +//// list ptx_regs = !foreach(tmp, regs, regclass); // List of register names for the fragment -- ["ra0", "ra1",...] list reg_names = RegSeq.ret; @@ -7450,12 +7453,14 @@ // 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; + true: AS_match.generic)>; // Build AS-constrained pattern. let IntrinsicPattern = BuildPatternPF.ret; @@ -7490,14 +7495,15 @@ // 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), +//// !dag(ops, !foreach(tmp, Frag.regs, node), Frag.reg_names), !if(WithStride, (ops node:$ldm), (ops))); // Build PatFrag that only matches particular address space. PatFrag IntrFrag = PatFrag; + true: AS_match.generic)>; // Build AS-constrained pattern. let IntrinsicPattern = BuildPatternPF.ret; @@ -7518,7 +7524,7 @@ // Create all load/store variants defset list 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