2531 lines · plain
1//===- NVPTXInstrInfo.td - NVPTX Instruction defs -------------*- tblgen-*-===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// This file describes the PTX instructions in TableGen format.10//11//===----------------------------------------------------------------------===//12 13include "NVPTXInstrFormats.td"14 15let OperandType = "OPERAND_IMMEDIATE" in {16 def f16imm : Operand<f16>;17 def bf16imm : Operand<bf16>;18}19 20//===----------------------------------------------------------------------===//21// NVPTX Operand Definitions.22//===----------------------------------------------------------------------===//23 24def brtarget : Operand<OtherVT>;25 26// CVT conversion modes27// These must match the enum in NVPTX.h28def CvtNONE : PatLeaf<(i32 0x0)>;29def CvtRNI : PatLeaf<(i32 0x1)>;30def CvtRZI : PatLeaf<(i32 0x2)>;31def CvtRMI : PatLeaf<(i32 0x3)>;32def CvtRPI : PatLeaf<(i32 0x4)>;33def CvtRN : PatLeaf<(i32 0x5)>;34def CvtRZ : PatLeaf<(i32 0x6)>;35def CvtRM : PatLeaf<(i32 0x7)>;36def CvtRP : PatLeaf<(i32 0x8)>;37def CvtRNA : PatLeaf<(i32 0x9)>;38def CvtRS : PatLeaf<(i32 0xA)>;39 40def CvtNONE_FTZ : PatLeaf<(i32 0x10)>;41def CvtRNI_FTZ : PatLeaf<(i32 0x11)>;42def CvtRZI_FTZ : PatLeaf<(i32 0x12)>;43def CvtRMI_FTZ : PatLeaf<(i32 0x13)>;44def CvtRPI_FTZ : PatLeaf<(i32 0x14)>;45def CvtRN_FTZ : PatLeaf<(i32 0x15)>;46def CvtRZ_FTZ : PatLeaf<(i32 0x16)>;47def CvtRM_FTZ : PatLeaf<(i32 0x17)>;48def CvtRP_FTZ : PatLeaf<(i32 0x18)>;49 50def CvtSAT : PatLeaf<(i32 0x20)>;51def CvtSAT_FTZ : PatLeaf<(i32 0x30)>;52 53def CvtNONE_RELU : PatLeaf<(i32 0x40)>;54def CvtRN_RELU : PatLeaf<(i32 0x45)>;55def CvtRZ_RELU : PatLeaf<(i32 0x46)>;56def CvtRS_RELU : PatLeaf<(i32 0x4A)>;57 58def CvtMode : Operand<i32> {59 let PrintMethod = "printCvtMode";60}61 62// FTZ flag63 64def FTZ : PatLeaf<(i1 1)>;65def NoFTZ : PatLeaf<(i1 0)>;66 67def getFTZFlag : SDNodeXForm<imm, [{68 (void)N;69 return CurDAG->getTargetConstant(useF32FTZ() ? 1 : 0, SDLoc(), MVT::i1);70}]>;71 72def FTZFlag : OperandWithDefaultOps<i1, (ops (getFTZFlag (i1 0)))> {73 let PrintMethod = "printFTZFlag";74}75 76// Compare modes77// These must match the enum in NVPTX.h78def CmpEQ : PatLeaf<(i32 0)>;79def CmpNE : PatLeaf<(i32 1)>;80 81def CmpMode : Operand<i32> {82 let PrintMethod = "printCmpMode";83}84 85// PRMT modes86// These must match the enum in NVPTX.h87def PrmtNONE : PatLeaf<(i32 0x0)>;88def PrmtF4E : PatLeaf<(i32 0x1)>;89def PrmtB4E : PatLeaf<(i32 0x2)>;90def PrmtRC8 : PatLeaf<(i32 0x3)>;91def PrmtECL : PatLeaf<(i32 0x4)>;92def PrmtECR : PatLeaf<(i32 0x5)>;93def PrmtRC16 : PatLeaf<(i32 0x6)>;94 95def PrmtMode : Operand<i32> {96 let PrintMethod = "printPrmtMode";97}98 99 100//===----------------------------------------------------------------------===//101// NVPTX Instruction Predicate Definitions102//===----------------------------------------------------------------------===//103 104// Checks PTX version and family-specific and architecture-specific SM versions.105// For example, sm_100{f/a} and any future variants in the same family will match106// for any PTX version greater than or equal to `PTXVersion`.107class PTXWithFamilySMs<int PTXVersion, list<int> SMVersions> :108 Predicate<"Subtarget->hasPTXWithFamilySMs(" # PTXVersion # ", {" #109 !interleave(SMVersions, ", ") # "})">;110 111// Checks PTX version and architecture-specific SM versions.112// For example, sm_100{a} will match for any PTX version113// greater than or equal to `PTXVersion`.114class PTXWithAccelSMs<int PTXVersion, list<int> SMVersions> :115 Predicate<"Subtarget->hasPTXWithAccelSMs(" # PTXVersion # ", {" #116 !interleave(SMVersions, ", ") # "})">;117 118// Helper predicate to call a subtarget method.119class callSubtarget<string SubtargetMethod> : Predicate<"Subtarget->" # SubtargetMethod # "()">;120 121def hasAtomAddF64 : Predicate<"Subtarget->hasAtomAddF64()">;122def hasAtomScope : Predicate<"Subtarget->hasAtomScope()">;123def hasAtomBitwise64 : Predicate<"Subtarget->hasAtomBitwise64()">;124def hasAtomMinMax64 : Predicate<"Subtarget->hasAtomMinMax64()">;125def hasAtomSwap128 : Predicate<"Subtarget->hasAtomSwap128()">;126def hasClusters : Predicate<"Subtarget->hasClusters()">;127def hasPTXASUnreachableBug : Predicate<"Subtarget->hasPTXASUnreachableBug()">;128def noPTXASUnreachableBug : Predicate<"!Subtarget->hasPTXASUnreachableBug()">;129def hasOptEnabled : Predicate<"TM.getOptLevel() != CodeGenOptLevel::None">;130def hasArchAccelFeatures : Predicate<"Subtarget->hasArchAccelFeatures()">;131 132def doF32FTZ : Predicate<"useF32FTZ()">;133def doNoF32FTZ : Predicate<"!useF32FTZ()">;134def doRsqrtOpt : Predicate<"doRsqrtOpt()">;135def doMADWideOpt : Predicate<"doMADWideOpt()">;136 137def hasHWROT32 : Predicate<"Subtarget->hasHWROT32()">;138def noHWROT32 : Predicate<"!Subtarget->hasHWROT32()">;139def hasDotInstructions : Predicate<"Subtarget->hasDotInstructions()">;140def hasTcgen05Instructions : Predicate<"Subtarget->hasTcgen05Instructions()">;141def hasTcgen05MMAScaleInputDImm : Predicate<"Subtarget->hasTcgen05MMAScaleInputDImm()">;142def hasF32x2Instructions : Predicate<"Subtarget->hasF32x2Instructions()">;143 144class hasPTX<int version>: Predicate<"Subtarget->getPTXVersion() >= " # version>;145class hasSM<int version>: Predicate<"Subtarget->getSmVersion() >= " # version>;146 147// Explicit records for arch-accelerated SM versions148def hasSM90a : Predicate<"Subtarget->getSmVersion() == 90 && Subtarget->hasArchAccelFeatures()">;149def hasSM100a : Predicate<"Subtarget->getSmVersion() == 100 && Subtarget->hasArchAccelFeatures()">;150def hasSM101a : Predicate<"Subtarget->getSmVersion() == 101 && Subtarget->hasArchAccelFeatures()">;151def hasSM120a : Predicate<"Subtarget->getSmVersion() == 120 && Subtarget->hasArchAccelFeatures()">;152 153def hasSM100aOrSM103a :154 Predicate<"(Subtarget->getSmVersion() == 100 || " # 155 "Subtarget->getSmVersion() == 103) " #156 "&& Subtarget->hasArchAccelFeatures()">;157 158// non-sync shfl instructions are not available on sm_70+ in PTX6.4+159def hasSHFL : Predicate<"!(Subtarget->getSmVersion() >= 70"160 "&& Subtarget->getPTXVersion() >= 64)">;161 162def useFP16Math: Predicate<"Subtarget->allowFP16Math()">;163def hasBF16Math: Predicate<"Subtarget->hasBF16Math()">;164 165 166//===----------------------------------------------------------------------===//167// Some Common Instruction Class Templates168//===----------------------------------------------------------------------===//169 170class OneUse1<SDPatternOperator operator>171 : PatFrag<(ops node:$A), (operator node:$A), [{ return N->hasOneUse(); }]>;172class OneUse2<SDPatternOperator operator>173 : PatFrag<(ops node:$A, node:$B), (operator node:$A, node:$B), [{ return N->hasOneUse(); }]>;174 175 176class zeroinitializer<ValueType vt> : 177 PatLeaf<(vt (bitconvert (!cast<ValueType>("i" # vt.Size) 0)))>;178 179 180def fpimm_pos_inf : FPImmLeaf<fAny, [{ return Imm.isPosInfinity(); }]>;181def fpimm_0 : FPImmLeaf<fAny, [{ return Imm.isZero(); }]>;182def fpimm_1 : FPImmLeaf<fAny, [{ return Imm.isExactlyValue(1.0); }]>;183def fpimm_neg_1 : FPImmLeaf<fAny, [{ return Imm.isExactlyValue(-1.0); }]>;184 185 186// Operands which can hold a Register or an Immediate.187//188// Unfortunately, since most register classes can hold multiple types, we must189// use the 'Any' type for these.190 191def RI1 : Operand<i1>;192def RI16 : Operand<Any>;193def RI32 : Operand<Any>;194def RI64 : Operand<Any>;195 196// Utility class to wrap up information about a register and DAG type for more197// convenient iteration and parameterization198class RegTyInfo<ValueType ty, NVPTXRegClass rc, string ptx_type, Operand imm, SDNode imm_node,199 bit supports_imm = 1> {200 ValueType Ty = ty;201 NVPTXRegClass RC = rc;202 Operand Imm = imm;203 SDNode ImmNode = imm_node;204 bit SupportsImm = supports_imm;205 int Size = ty.Size;206 string PtxType = ptx_type;207}208 209def I1RT : RegTyInfo<i1, B1, "pred", i1imm, imm>;210def I16RT : RegTyInfo<i16, B16, "b16", i16imm, imm>;211def I32RT : RegTyInfo<i32, B32, "b32", i32imm, imm>;212def I64RT : RegTyInfo<i64, B64, "b64", i64imm, imm>;213 214def F32RT : RegTyInfo<f32, B32, "f32", f32imm, fpimm>;215def F64RT : RegTyInfo<f64, B64, "f64", f64imm, fpimm>;216def F16RT : RegTyInfo<f16, B16, "f16", f16imm, fpimm, supports_imm = 0>;217def BF16RT : RegTyInfo<bf16, B16, "bf16", bf16imm, fpimm, supports_imm = 0>;218 219def F16X2RT : RegTyInfo<v2f16, B32, "f16x2", ?, ?, supports_imm = 0>;220def BF16X2RT : RegTyInfo<v2bf16, B32, "bf16x2", ?, ?, supports_imm = 0>;221def F32X2RT : RegTyInfo<v2f32, B64, "f32x2", ?, ?, supports_imm = 0>;222 223 224// This class provides a basic wrapper around an NVPTXInst that abstracts the225// specific syntax of most PTX instructions. It automatically handles the226// construction of the asm string based on the provided dag arguments.227// For example, the following asm-strings would be computed:228//229// * BasicFlagsNVPTXInst<(outs B32:$dst),230// (ins B32:$a, B32:$b), (ins),231// "add.s32">;232// ---> "add.s32 \t$dst, $a, $b;"233//234// * BasicFlagsNVPTXInst<(outs B32:$d),235// (ins B32:$a, B32:$b, Hexu32imm:$c),236// (ins PrmtMode:$mode),237// "prmt.b32${mode}">;238// ---> "prmt.b32${mode} \t$d, $a, $b, $c;"239//240// * BasicFlagsNVPTXInst<(outs B64:$state),241// (ins ADDR:$addr),242// "mbarrier.arrive.b64">;243// ---> "mbarrier.arrive.b64 \t$state, [$addr];"244//245class BasicFlagsNVPTXInst<dag outs_dag, dag ins_dag, dag flags_dag, string asmstr,246 list<dag> pattern = []>247 : NVPTXInst<248 outs_dag,249 !con(ins_dag, flags_dag),250 !strconcat(251 asmstr,252 !if(!and(!empty(ins_dag), !empty(outs_dag)), "",253 !strconcat(254 " \t",255 !interleave(256 !foreach(i, !range(!size(outs_dag)),257 "$" # !getdagname(outs_dag, i)),258 "|"),259 !if(!or(!empty(ins_dag), !empty(outs_dag)), "", ", "),260 !interleave(261 !foreach(i, !range(!size(ins_dag)),262 !if(!eq(!cast<string>(!getdagarg<DAGOperand>(ins_dag, i)), "ADDR"),263 "[$" # !getdagname(ins_dag, i) # "]",264 "$" # !getdagname(ins_dag, i)265 )266 ),267 ", "))),268 ";"),269 pattern>;270 271class BasicNVPTXInst<dag outs, dag insv, string asmstr, list<dag> pattern = []>272 : BasicFlagsNVPTXInst<outs, insv, (ins), asmstr, pattern>;273 274 275multiclass I3Inst<string op_str, SDPatternOperator op_node, RegTyInfo t,276 bit commutative, list<Predicate> requires = []> {277 def rr :278 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a, t.RC:$b),279 op_str,280 [(set t.Ty:$dst, (op_node t.Ty:$a, t.Ty:$b))]>,281 Requires<requires>;282 def ri :283 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a, t.Imm:$b),284 op_str,285 [(set t.Ty:$dst, (op_node t.Ty:$a, (t.Ty imm:$b)))]>,286 Requires<requires>;287 if !not(commutative) then288 def ir :289 BasicNVPTXInst<(outs t.RC:$dst), (ins t.Imm:$a, t.RC:$b),290 op_str,291 [(set t.Ty:$dst, (op_node (t.Ty imm:$a), t.Ty:$b))]>,292 Requires<requires>;293}294 295// Template for instructions which take three int64, int32, or int16 args.296// The instructions are named "<OpcStr><Width>" (e.g. "add.s64").297multiclass I3<string op_str, SDPatternOperator op_node, bit commutative> {298 foreach t = [I16RT, I32RT, I64RT] in299 defm t.Size# : I3Inst<op_str # t.Size, op_node, t, commutative>;300}301 302class I16x2<string OpcStr, SDNode OpNode> :303 BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),304 OpcStr # "16x2",305 [(set v2i16:$dst, (OpNode v2i16:$a, v2i16:$b))]>,306 Requires<[hasPTX<80>, hasSM<90>]>;307 308// Template for instructions which take 3 int args. The instructions are309// named "<OpcStr>.s32" (e.g. "addc.cc.s32").310multiclass ADD_SUB_INT_CARRY<string op_str, SDNode op_node, bit commutative> {311 let hasSideEffects = 1 in {312 defm i32 : I3Inst<op_str # ".s32", op_node, I32RT, commutative>;313 defm i64 : I3Inst<op_str # ".s64", op_node, I64RT, commutative,314 requires = [hasPTX<43>]>;315 }316}317 318// Template for minimum/maximum instructions.319//320// Also defines ftz (flush subnormal inputs and results to sign-preserving321// zero) variants for fp32 functions.322multiclass FMINIMUMMAXIMUM<string OpcStr, bit NaN, SDPatternOperator OpNode> {323 defvar nan_str = !if(NaN, ".NaN", "");324 if !not(NaN) then {325 def _f64_rr :326 BasicNVPTXInst<(outs B64:$dst),327 (ins B64:$a, B64:$b),328 OpcStr # ".f64",329 [(set f64:$dst, (OpNode f64:$a, f64:$b))]>;330 def _f64_ri :331 BasicNVPTXInst<(outs B64:$dst),332 (ins B64:$a, f64imm:$b),333 OpcStr # ".f64",334 [(set f64:$dst, (OpNode f64:$a, fpimm:$b))]>;335 }336 def _f32_rr :337 BasicFlagsNVPTXInst<(outs B32:$dst),338 (ins B32:$a, B32:$b),339 (ins FTZFlag:$ftz),340 OpcStr # "$ftz" # nan_str # ".f32",341 [(set f32:$dst, (OpNode f32:$a, f32:$b))]>;342 def _f32_ri :343 BasicFlagsNVPTXInst<(outs B32:$dst),344 (ins B32:$a, f32imm:$b),345 (ins FTZFlag:$ftz),346 OpcStr # "$ftz" # nan_str # ".f32",347 [(set f32:$dst, (OpNode f32:$a, fpimm:$b))]>;348 349 def _f16_rr :350 BasicFlagsNVPTXInst<(outs B16:$dst),351 (ins B16:$a, B16:$b),352 (ins FTZFlag:$ftz),353 OpcStr # "$ftz" # nan_str # ".f16",354 [(set f16:$dst, (OpNode f16:$a, f16:$b))]>,355 Requires<[useFP16Math]>;356 357 def _f16x2_rr :358 BasicFlagsNVPTXInst<(outs B32:$dst),359 (ins B32:$a, B32:$b),360 (ins FTZFlag:$ftz),361 OpcStr # "$ftz" # nan_str # ".f16x2",362 [(set v2f16:$dst, (OpNode v2f16:$a, v2f16:$b))]>,363 Requires<[useFP16Math, hasSM<80>, hasPTX<70>]>;364 def _bf16_rr :365 BasicNVPTXInst<(outs B16:$dst),366 (ins B16:$a, B16:$b),367 OpcStr # nan_str # ".bf16",368 [(set bf16:$dst, (OpNode bf16:$a, bf16:$b))]>,369 Requires<[hasBF16Math, hasSM<80>, hasPTX<70>]>;370 def _bf16x2_rr :371 BasicNVPTXInst<(outs B32:$dst),372 (ins B32:$a, B32:$b),373 OpcStr # nan_str # ".bf16x2",374 [(set v2bf16:$dst, (OpNode v2bf16:$a, v2bf16:$b))]>,375 Requires<[hasBF16Math, hasSM<80>, hasPTX<70>]>;376}377 378// Template for 3-input minimum/maximum instructions379// (sm_100+/PTX 8.8 and f32 only)380//381// Also defines ftz (flush subnormal inputs and results to sign-preserving382// zero) variants for fp32 functions.383multiclass FMINIMUMMAXIMUM3<string OpcStr, bit NaN, SDNode OpNode> {384 defvar nan_str = !if(NaN, ".NaN", "");385 def f32rrr :386 BasicFlagsNVPTXInst<(outs B32:$dst),387 (ins B32:$a, B32:$b, B32:$c),388 (ins FTZFlag:$ftz),389 OpcStr # "$ftz" # nan_str # ".f32",390 [(set f32:$dst, (OpNode f32:$a, f32:$b, f32:$c))]>,391 Requires<[hasPTX<88>, hasSM<100>]>;392 def f32rri :393 BasicFlagsNVPTXInst<(outs B32:$dst),394 (ins B32:$a, B32:$b, f32imm:$c),395 (ins FTZFlag:$ftz),396 OpcStr # "$ftz" # nan_str # ".f32",397 [(set f32:$dst, (OpNode f32:$a, f32:$b, fpimm:$c))]>,398 Requires<[hasPTX<88>, hasSM<100>]>;399 def f32rii :400 BasicFlagsNVPTXInst<(outs B32:$dst),401 (ins B32:$a, f32imm:$b, f32imm:$c),402 (ins FTZFlag:$ftz),403 OpcStr # "$ftz" # nan_str # ".f32",404 [(set f32:$dst, (OpNode f32:$a, fpimm:$b, fpimm:$c))]>,405 Requires<[hasPTX<88>, hasSM<100>]>;406}407 408// Template for instructions which take three FP args. The409// instructions are named "<OpcStr>.f<Width>" (e.g. "add.f64").410//411// Also defines ftz (flush subnormal inputs and results to sign-preserving412// zero) variants for fp32/fp16 functions.413//414// This multiclass should be used for nodes that can be folded to make fma ops.415// In this case, we use the ".rn" variant when FMA is disabled, as this behaves416// just like the non ".rn" op, but prevents ptxas from creating FMAs.417multiclass F3<string op_str, SDPatternOperator op_pat> {418 def f64rr :419 BasicNVPTXInst<(outs B64:$dst),420 (ins B64:$a, B64:$b),421 op_str # ".f64",422 [(set f64:$dst, (op_pat f64:$a, f64:$b))]>;423 def f64ri :424 BasicNVPTXInst<(outs B64:$dst),425 (ins B64:$a, f64imm:$b),426 op_str # ".f64",427 [(set f64:$dst, (op_pat f64:$a, fpimm:$b))]>;428 def f32rr :429 BasicFlagsNVPTXInst<(outs B32:$dst),430 (ins B32:$a, B32:$b),431 (ins FTZFlag:$ftz),432 op_str # "$ftz.f32",433 [(set f32:$dst, (op_pat f32:$a, f32:$b))]>;434 def f32ri :435 BasicFlagsNVPTXInst<(outs B32:$dst),436 (ins B32:$a, f32imm:$b),437 (ins FTZFlag:$ftz),438 op_str # "$ftz.f32",439 [(set f32:$dst, (op_pat f32:$a, fpimm:$b))]>;440 441 def f16rr :442 BasicFlagsNVPTXInst<(outs B16:$dst),443 (ins B16:$a, B16:$b),444 (ins FTZFlag:$ftz),445 op_str # "$ftz.f16",446 [(set f16:$dst, (op_pat f16:$a, f16:$b))]>,447 Requires<[useFP16Math]>;448 def f32x2rr :449 BasicFlagsNVPTXInst<(outs B64:$dst),450 (ins B64:$a, B64:$b),451 (ins FTZFlag:$ftz),452 op_str # "$ftz.f32x2",453 [(set v2f32:$dst, (op_pat v2f32:$a, v2f32:$b))]>,454 Requires<[hasF32x2Instructions]>;455 def f16x2rr :456 BasicFlagsNVPTXInst<(outs B32:$dst),457 (ins B32:$a, B32:$b),458 (ins FTZFlag:$ftz),459 op_str # "$ftz.f16x2",460 [(set v2f16:$dst, (op_pat v2f16:$a, v2f16:$b))]>,461 Requires<[useFP16Math]>;462 def bf16rr :463 BasicNVPTXInst<(outs B16:$dst),464 (ins B16:$a, B16:$b),465 op_str # ".bf16",466 [(set bf16:$dst, (op_pat bf16:$a, bf16:$b))]>,467 Requires<[hasBF16Math]>;468 469 def bf16x2rr :470 BasicNVPTXInst<(outs B32:$dst),471 (ins B32:$a, B32:$b),472 op_str # ".bf16x2",473 [(set v2bf16:$dst, (op_pat v2bf16:$a, v2bf16:$b))]>,474 Requires<[hasBF16Math]>;475}476 477class BinOpAllowsFMA<SDPatternOperator operator>478 : PatFrag<(ops node:$A, node:$B),479 (operator node:$A, node:$B), [{480 return allowFMA() || N->getFlags().hasAllowContract();481}]>;482 483multiclass F3_fma_component<string op_str, SDNode op_node> {484 defm "" : F3<op_str, BinOpAllowsFMA<op_node>>;485 defm _rn : F3<op_str # ".rn", op_node>;486}487 488// Template for operations which take two f32 or f64 operands. Provides three489// instructions: <OpcStr>.f64, <OpcStr>.f32, and <OpcStr>.ftz.f32 (flush490// subnormal inputs and results to zero).491multiclass F2<string OpcStr, SDNode OpNode> {492 def f64 : BasicNVPTXInst<(outs B64:$dst), (ins B64:$a),493 OpcStr # ".f64",494 [(set f64:$dst, (OpNode f64:$a))]>;495 def f32 : BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$a),496 (ins FTZFlag:$ftz),497 OpcStr # "$ftz.f32",498 [(set f32:$dst, (OpNode f32:$a))]>;499}500 501multiclass F2_Support_Half<string OpcStr, SDNode OpNode> {502 def bf16 : BasicNVPTXInst<(outs B16:$dst), (ins B16:$a),503 OpcStr # ".bf16",504 [(set bf16:$dst, (OpNode bf16:$a))]>,505 Requires<[hasSM<80>, hasPTX<70>]>;506 def bf16x2 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a),507 OpcStr # ".bf16x2",508 [(set v2bf16:$dst, (OpNode v2bf16:$a))]>,509 Requires<[hasSM<80>, hasPTX<70>]>;510 def f16 : BasicFlagsNVPTXInst<(outs B16:$dst), (ins B16:$a),511 (ins FTZFlag:$ftz),512 OpcStr # "$ftz.f16",513 [(set f16:$dst, (OpNode f16:$a))]>,514 Requires<[hasSM<53>, hasPTX<65>]>;515 def f16x2 : BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$a),516 (ins FTZFlag:$ftz),517 OpcStr # "$ftz.f16x2",518 [(set v2f16:$dst, (OpNode v2f16:$a))]>,519 Requires<[hasSM<53>, hasPTX<65>]>;520 521}522 523//===----------------------------------------------------------------------===//524// NVPTX Instructions.525//===----------------------------------------------------------------------===//526 527//-----------------------------------528// Type Conversion529//-----------------------------------530 531let hasSideEffects = false in {532 // Generate a cvt to the given type from all possible types. Each instance533 // takes a CvtMode immediate that defines the conversion mode to use. It can534 // be CvtNONE to omit a conversion mode.535 multiclass CVT_FROM_ALL<string ToType, RegisterClass RC, list<Predicate> Preds = []> {536 foreach sign = ["s", "u"] in {537 def _ # sign # "8" :538 BasicFlagsNVPTXInst<(outs RC:$dst),539 (ins B16:$src), (ins CvtMode:$mode),540 "cvt${mode:base}${mode:ftz}${mode:sat}." # ToType # "." # sign # "8">,541 Requires<Preds>;542 def _ # sign # "16" :543 BasicFlagsNVPTXInst<(outs RC:$dst),544 (ins B16:$src), (ins CvtMode:$mode),545 "cvt${mode:base}${mode:ftz}${mode:sat}." # ToType # "." # sign # "16">,546 Requires<Preds>;547 def _ # sign # "32" :548 BasicFlagsNVPTXInst<(outs RC:$dst),549 (ins B32:$src), (ins CvtMode:$mode),550 "cvt${mode:base}${mode:ftz}${mode:sat}." # ToType # "." # sign # "32">,551 Requires<Preds>;552 def _ # sign # "64" :553 BasicFlagsNVPTXInst<(outs RC:$dst),554 (ins B64:$src), (ins CvtMode:$mode),555 "cvt${mode:base}${mode:ftz}${mode:sat}." # ToType # "." # sign # "64">,556 Requires<Preds>;557 }558 def _f16 :559 BasicFlagsNVPTXInst<(outs RC:$dst),560 (ins B16:$src), (ins CvtMode:$mode),561 "cvt${mode:base}${mode:ftz}${mode:sat}." # ToType # ".f16">,562 Requires<Preds>;563 def _bf16 :564 BasicFlagsNVPTXInst<(outs RC:$dst),565 (ins B16:$src), (ins CvtMode:$mode),566 "cvt${mode:base}${mode:ftz}${mode:relu}${mode:sat}." # ToType # ".bf16">,567 Requires<!if(!eq(ToType, "f32"),568 // bf16->f32 was introduced early.569 [hasPTX<71>, hasSM<80>],570 // bf16->everything else needs sm90/ptx78571 [hasPTX<78>, hasSM<90>])>;572 def _f32 :573 BasicFlagsNVPTXInst<(outs RC:$dst),574 (ins B32:$src), (ins CvtMode:$mode),575 "cvt${mode:base}${mode:ftz}${mode:relu}${mode:sat}." # ToType # ".f32">,576 Requires<!if(!eq(ToType, "bf16"),577 // f32->bf16 was introduced early.578 [hasPTX<70>, hasSM<80>],579 Preds)>;580 def _f64 :581 BasicFlagsNVPTXInst<(outs RC:$dst),582 (ins B64:$src), (ins CvtMode:$mode),583 "cvt${mode:base}${mode:ftz}${mode:sat}." # ToType # ".f64">,584 Requires<Preds>;585 }586 587 // Generate cvts from all types to all types.588 foreach sign = ["s", "u"] in {589 defm CVT_ # sign # "8" : CVT_FROM_ALL<sign # "8", B16>;590 defm CVT_ # sign # "16" : CVT_FROM_ALL<sign # "16", B16>;591 defm CVT_ # sign # "32" : CVT_FROM_ALL<sign # "32", B32>;592 defm CVT_ # sign # "64" : CVT_FROM_ALL<sign # "64", B64>;593 }594 defm CVT_f16 : CVT_FROM_ALL<"f16", B16>;595 defm CVT_bf16 : CVT_FROM_ALL<"bf16", B16, [hasPTX<78>, hasSM<90>]>;596 defm CVT_f32 : CVT_FROM_ALL<"f32", B32>;597 defm CVT_f64 : CVT_FROM_ALL<"f64", B64>;598 599 multiclass CVT_FROM_FLOAT_SATFINITE<string ToName, RegisterClass RC> {600 def _f32_sf :601 BasicFlagsNVPTXInst<(outs RC:$dst),602 (ins B32:$src), (ins CvtMode:$mode),603 "cvt${mode:base}${mode:relu}.satfinite." # ToName # ".f32">;604 }605 defm CVT_bf16 : CVT_FROM_FLOAT_SATFINITE<"bf16", B16>;606 defm CVT_f16 : CVT_FROM_FLOAT_SATFINITE<"f16", B16>;607 608 // These cvts are different from those above: The source and dest registers609 // are of the same type.610 def CVT_INREG_s16_s8 : BasicNVPTXInst<(outs B16:$dst), (ins B16:$src), "cvt.s16.s8">;611 def CVT_INREG_s32_s8 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$src), "cvt.s32.s8">;612 def CVT_INREG_s32_s16 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$src), "cvt.s32.s16">;613 def CVT_INREG_s64_s8 : BasicNVPTXInst<(outs B64:$dst), (ins B64:$src), "cvt.s64.s8">;614 def CVT_INREG_s64_s16 : BasicNVPTXInst<(outs B64:$dst), (ins B64:$src), "cvt.s64.s16">;615 def CVT_INREG_s64_s32 : BasicNVPTXInst<(outs B64:$dst), (ins B64:$src), "cvt.s64.s32">;616 617 multiclass CVT_FROM_FLOAT_V2_SM80<string FromName, RegisterClass RC> {618 def _f32 :619 BasicFlagsNVPTXInst<(outs RC:$dst),620 (ins B32:$src1, B32:$src2), (ins CvtMode:$mode),621 "cvt${mode:base}${mode:relu}." # FromName # ".f32">,622 Requires<[hasPTX<70>, hasSM<80>]>;623 624 def _f32_sf :625 BasicFlagsNVPTXInst<(outs RC:$dst),626 (ins B32:$src1, B32:$src2), (ins CvtMode:$mode),627 "cvt${mode:base}${mode:relu}.satfinite." # FromName # ".f32">;628 }629 630 defm CVT_f16x2 : CVT_FROM_FLOAT_V2_SM80<"f16x2", B32>;631 defm CVT_bf16x2 : CVT_FROM_FLOAT_V2_SM80<"bf16x2", B32>;632 633 multiclass CVT_FROM_FLOAT_V2_RS<string FromName, RegisterClass RC> {634 def _f32_rs :635 BasicFlagsNVPTXInst<(outs RC:$dst),636 (ins B32:$src1, B32:$src2, B32:$src3),637 (ins CvtMode:$mode),638 "cvt${mode:base}${mode:relu}." # FromName # ".f32">;639 640 def _f32_rs_sf :641 BasicFlagsNVPTXInst<(outs RC:$dst),642 (ins B32:$src1, B32:$src2, B32:$src3),643 (ins CvtMode:$mode),644 "cvt${mode:base}${mode:relu}.satfinite." # FromName # ".f32">;645 }646 647 defm CVT_f16x2 : CVT_FROM_FLOAT_V2_RS<"f16x2", B32>;648 defm CVT_bf16x2 : CVT_FROM_FLOAT_V2_RS<"bf16x2", B32>;649 650 // FP8 conversions.651 multiclass CVT_TO_F8X2<string F8Name> {652 def _f32 :653 BasicFlagsNVPTXInst<(outs B16:$dst),654 (ins B32:$src1, B32:$src2), (ins CvtMode:$mode),655 "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.f32">,656 Requires<[hasPTX<81>, hasSM<89>]>;657 def _f16x2 :658 BasicFlagsNVPTXInst<(outs B16:$dst),659 (ins B32:$src), (ins CvtMode:$mode),660 "cvt${mode:base}.satfinite${mode:relu}." # F8Name # "x2.f16x2">,661 Requires<[hasPTX<81>, hasSM<89>]>;662 }663 664 defm CVT_e4m3x2 : CVT_TO_F8X2<"e4m3">;665 defm CVT_e5m2x2 : CVT_TO_F8X2<"e5m2">;666 667 class CVT_f16x2_fp8<string F8Name> :668 BasicFlagsNVPTXInst<(outs B32:$dst),669 (ins B16:$src), (ins CvtMode:$mode),670 "cvt${mode:base}${mode:relu}.f16x2." # F8Name # "x2">,671 Requires<[hasPTX<81>, hasSM<89>]>;672 673 def CVT_f16x2_e4m3x2 : CVT_f16x2_fp8<"e4m3">;674 def CVT_f16x2_e5m2x2 : CVT_f16x2_fp8<"e5m2">;675 676 class CVT_TO_FP8X4<string F8Name> :677 NVPTXInst<(outs B32:$dst),678 (ins B32:$src1, B32:$src2, B32:$src3, B32:$src4, B32:$src5, CvtMode:$mode),679 "cvt${mode:base}${mode:relu}.satfinite." # F8Name # 680 "x4.f32 \t$dst, {{$src1, $src2, $src3, $src4}}, $src5;">;681 682 def CVT_e4m3x4_f32x4_rs_sf : CVT_TO_FP8X4<"e4m3">;683 def CVT_e5m2x4_f32x4_rs_sf : CVT_TO_FP8X4<"e5m2">;684 685 // Float to TF32 conversions686 multiclass CVT_TO_TF32<string Modifier, list<Predicate> Preds = [hasPTX<78>, hasSM<90>]> {687 defvar Intr = !cast<Intrinsic>("int_nvvm_f2tf32_" # !subst(".", "_", Modifier));688 689 def NAME : BasicNVPTXInst<(outs B32:$dst), (ins B32:$src),690 "cvt." # Modifier # ".tf32.f32",691 [(set i32:$dst, (Intr f32:$src))]>,692 Requires<Preds>;693 }694 695 defm CVT_to_tf32_rn : CVT_TO_TF32<"rn">;696 defm CVT_to_tf32_rz : CVT_TO_TF32<"rz">;697 defm CVT_to_tf32_rn_relu : CVT_TO_TF32<"rn.relu">;698 defm CVT_to_tf32_rz_relu : CVT_TO_TF32<"rz.relu">;699 defm CVT_to_tf32_rna : CVT_TO_TF32<"rna", [hasPTX<70>, hasSM<80>]>;700 defm CVT_to_tf32_rna_satf : CVT_TO_TF32<"rna.satfinite", [hasPTX<81>, hasSM<80>]>;701 702 defm CVT_to_tf32_rn_satf : CVT_TO_TF32<"rn.satfinite", [hasPTX<86>, hasSM<100>]>;703 defm CVT_to_tf32_rz_satf : CVT_TO_TF32<"rz.satfinite", [hasPTX<86>, hasSM<100>]>;704 defm CVT_to_tf32_rn_relu_satf : CVT_TO_TF32<"rn.relu.satfinite", [hasPTX<86>, hasSM<100>]>;705 defm CVT_to_tf32_rz_relu_satf : CVT_TO_TF32<"rz.relu.satfinite", [hasPTX<86>, hasSM<100>]>;706 707 // FP6 conversions.708 foreach type = ["e2m3x2", "e3m2x2"] in {709 def CVT_ # type # _f32_sf : BasicFlagsNVPTXInst<(outs B16:$dst),710 (ins B32:$src1, B32:$src2), (ins CvtMode:$mode),711 "cvt${mode:base}.satfinite${mode:relu}." # type # ".f32">;712 def CVT_f16x2_ # type : BasicFlagsNVPTXInst<(outs B32:$dst),713 (ins B16:$src), (ins CvtMode:$mode),714 "cvt${mode:base}${mode:relu}.f16x2." # type>;715 }716 717 class CVT_TO_FP6X4<string F6Name> :718 NVPTXInst<(outs B32:$dst),719 (ins B32:$src1, B32:$src2, B32:$src3, B32:$src4, B32:$src5, CvtMode:$mode),720 "cvt${mode:base}${mode:relu}.satfinite." # F6Name #721 "x4.f32 \t$dst, {{$src1, $src2, $src3, $src4}}, $src5;">;722 723 def CVT_e2m3x4_f32x4_rs_sf : CVT_TO_FP6X4<"e2m3">;724 def CVT_e3m2x4_f32x4_rs_sf : CVT_TO_FP6X4<"e3m2">;725 726 // FP4 conversions.727 def CVT_e2m1x2_f32_sf : NVPTXInst<(outs B16:$dst),728 (ins B32:$src1, B32:$src2, CvtMode:$mode),729 !strconcat("{{ \n\t",730 ".reg .b8 \t%e2m1x2_out; \n\t",731 "cvt${mode:base}.satfinite${mode:relu}.e2m1x2.f32 \t%e2m1x2_out, $src1, $src2; \n\t",732 "cvt.u16.u8 \t$dst, %e2m1x2_out; \n\t",733 "}}"), []>;734 735 def CVT_f16x2_e2m1x2 : NVPTXInst<(outs B32:$dst),736 (ins B16:$src, CvtMode:$mode),737 !strconcat("{{ \n\t",738 ".reg .b8 \t%e2m1x2_in; \n\t",739 "cvt.u8.u16 \t%e2m1x2_in, $src; \n\t",740 "cvt${mode:base}${mode:relu}.f16x2.e2m1x2 \t$dst, %e2m1x2_in; \n\t",741 "}}"), []>;742 743 def CVT_e2m1x4_f32x4_rs_sf :744 NVPTXInst<(outs B16:$dst),745 (ins B32:$src1, B32:$src2, B32:$src3, B32:$src4, B32:$src5, CvtMode:$mode),746 "cvt${mode:base}${mode:relu}.satfinite.e2m1x4.f32 \t" # 747 "$dst, {{$src1, $src2, $src3, $src4}}, $src5;">;748 749 // UE8M0x2 conversions.750 class CVT_f32_to_ue8m0x2<string sat = ""> :751 BasicFlagsNVPTXInst<(outs B16:$dst),752 (ins B32:$src1, B32:$src2), (ins CvtMode:$mode),753 "cvt${mode:base}" # sat # ".ue8m0x2.f32">;754 755 class CVT_bf16x2_to_ue8m0x2<string sat = ""> :756 BasicFlagsNVPTXInst<(outs B16:$dst),757 (ins B32:$src), (ins CvtMode:$mode),758 "cvt${mode:base}" # sat # ".ue8m0x2.bf16x2">;759 760 def CVT_ue8m0x2_f32 : CVT_f32_to_ue8m0x2;761 def CVT_ue8m0x2_f32_sf : CVT_f32_to_ue8m0x2<".satfinite">;762 def CVT_ue8m0x2_bf16x2 : CVT_bf16x2_to_ue8m0x2;763 def CVT_ue8m0x2_bf16x2_sf : CVT_bf16x2_to_ue8m0x2<".satfinite">;764 765 def CVT_bf16x2_ue8m0x2 :766 BasicNVPTXInst<(outs B32:$dst),767 (ins B16:$src),768 "cvt.rn.bf16x2.ue8m0x2">;769 770}771 772def fpround_oneuse : OneUse1<fpround>;773def : Pat<(v2bf16 (build_vector (bf16 (fpround_oneuse f32:$lo)),774 (bf16 (fpround_oneuse f32:$hi)))),775 (CVT_bf16x2_f32 $hi, $lo, CvtRN)>,776 Requires<[hasPTX<70>, hasSM<80>, hasBF16Math]>;777 778def : Pat<(v2f16 (build_vector (f16 (fpround_oneuse f32:$lo)),779 (f16 (fpround_oneuse f32:$hi)))),780 (CVT_f16x2_f32 $hi, $lo, CvtRN)>,781 Requires<[hasPTX<70>, hasSM<80>, useFP16Math]>;782 783//-----------------------------------784// Selection instructions (selp)785//-----------------------------------786 787// TODO: Missing slct788 789// selp instructions that don't have any pattern matches; we explicitly use790// them within this file.791let hasSideEffects = false in {792 multiclass SELP_PATTERN<string TypeStr, RegTyInfo t> {793 defvar asm_str = "selp." # TypeStr;794 def rr :795 BasicNVPTXInst<(outs t.RC:$dst),796 (ins t.RC:$a, t.RC:$b, B1:$p),797 asm_str,798 [(set t.Ty:$dst, (select i1:$p, t.Ty:$a, t.Ty:$b))]>;799 def ri :800 BasicNVPTXInst<(outs t.RC:$dst),801 (ins t.RC:$a, t.Imm:$b, B1:$p),802 asm_str,803 [(set t.Ty:$dst, (select i1:$p, t.Ty:$a, t.ImmNode:$b))]>;804 def ir :805 BasicNVPTXInst<(outs t.RC:$dst),806 (ins t.Imm:$a, t.RC:$b, B1:$p),807 asm_str,808 [(set t.Ty:$dst, (select i1:$p, t.ImmNode:$a, t.Ty:$b))]>;809 def ii :810 BasicNVPTXInst<(outs t.RC:$dst),811 (ins t.Imm:$a, t.Imm:$b, B1:$p),812 asm_str,813 [(set t.Ty:$dst, (select i1:$p, t.ImmNode:$a, t.ImmNode:$b))]>;814 }815}816 817// Don't pattern match on selp.{s,u}{16,32,64} -- selp.b{16,32,64} is just as818// good.819defm SELP_b16 : SELP_PATTERN<"b16", I16RT>;820defm SELP_b32 : SELP_PATTERN<"b32", I32RT>;821defm SELP_b64 : SELP_PATTERN<"b64", I64RT>;822defm SELP_f16 : SELP_PATTERN<"b16", F16RT>;823defm SELP_bf16 : SELP_PATTERN<"b16", BF16RT>;824defm SELP_f32 : SELP_PATTERN<"f32", F32RT>;825defm SELP_f64 : SELP_PATTERN<"f64", F64RT>;826 827// This does not work as tablegen fails to infer the type of 'imm'.828// def v2f16imm : Operand<v2f16>;829// defm SELP_f16x2 : SELP_PATTERN<"b32", v2f16, B32, v2f16imm, imm>;830 831foreach vt = [v2f16, v2bf16, v2i16, v4i8] in {832def : Pat<(vt (select i1:$p, vt:$a, vt:$b)),833 (SELP_b32rr $a, $b, $p)>;834}835 836foreach vt = [v2f32, v2i32] in {837def : Pat<(vt (select i1:$p, vt:$a, vt:$b)),838 (SELP_b64rr $a, $b, $p)>;839}840 841//-----------------------------------842// Test Instructions843//-----------------------------------844 845def fabs_oneuse : OneUse1<fabs>;846 847def TESTINF_f32r : BasicNVPTXInst<(outs B1:$p), (ins B32:$a),848 "testp.infinite.f32",849 [(set i1:$p, (seteq (fabs_oneuse f32:$a), fpimm_pos_inf))]>;850def TESTINF_f64r : BasicNVPTXInst<(outs B1:$p), (ins B64:$a),851 "testp.infinite.f64",852 [(set i1:$p, (seteq (fabs_oneuse f64:$a), fpimm_pos_inf))]>;853 854//-----------------------------------855// Integer Arithmetic856//-----------------------------------857 858// int16, int32, and int64 signed addition. Since nvptx is 2's complement, we859// also use these for unsigned arithmetic.860defm ADD : I3<"add.s", add, commutative = true>;861defm SUB : I3<"sub.s", sub, commutative = false>;862 863def ADD16x2 : I16x2<"add.s", add>;864 865// int32 and int64 addition and subtraction with carry-out.866defm ADDCC : ADD_SUB_INT_CARRY<"add.cc", addc, commutative = true>;867defm SUBCC : ADD_SUB_INT_CARRY<"sub.cc", subc, commutative = false>;868 869// int32 and int64 addition and subtraction with carry-in and carry-out.870defm ADDCCC : ADD_SUB_INT_CARRY<"addc.cc", adde, commutative = true>;871defm SUBCCC : ADD_SUB_INT_CARRY<"subc.cc", sube, commutative = false>;872 873defm MULT : I3<"mul.lo.s", mul, commutative = true>;874 875defm MUL_HI_S : I3<"mul.hi.s", mulhs, commutative = true>;876defm MUL_HI_U : I3<"mul.hi.u", mulhu, commutative = true>;877 878defm SDIV : I3<"div.s", sdiv, commutative = false>;879defm UDIV : I3<"div.u", udiv, commutative = false>;880 881// The ri versions of rem.s and rem.u won't be selected; DAGCombiner::visitSREM882// will lower it.883defm SREM : I3<"rem.s", srem, commutative = false>;884defm UREM : I3<"rem.u", urem, commutative = false>;885 886foreach t = [I16RT, I32RT, I64RT] in {887 def ABS_S # t.Size :888 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a),889 "abs.s" # t.Size,890 [(set t.Ty:$dst, (abs t.Ty:$a))]>;891 892 def NEG_S # t.Size :893 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src),894 "neg.s" # t.Size,895 [(set t.Ty:$dst, (ineg t.Ty:$src))]>;896}897 898// Integer min/max.899defm SMAX : I3<"max.s", smax, commutative = true>;900defm UMAX : I3<"max.u", umax, commutative = true>;901defm SMIN : I3<"min.s", smin, commutative = true>;902defm UMIN : I3<"min.u", umin, commutative = true>;903 904def SMAX16x2 : I16x2<"max.s", smax>;905def UMAX16x2 : I16x2<"max.u", umax>;906def SMIN16x2 : I16x2<"min.s", smin>;907def UMIN16x2 : I16x2<"min.u", umin>;908 909let Predicates = [hasPTX<80>, hasSM<90>] in {910 911 def MIN_RELU_S32 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),912 "min.relu.s32",913 [(set i32:$dst, (smax (smin i32:$a, i32:$b), 0))]>;914 def MAX_RELU_S32 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),915 "max.relu.s32",916 [(set i32:$dst, (smax (smax i32:$a, i32:$b), 0))]>;917 def MIN_RELU_S16x2 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),918 "min.relu.s16x2",919 [(set v2i16:$dst, (smax (smin v2i16:$a, v2i16:$b),920 zeroinitializer<v2i16>))]>;921 def MAX_RELU_S16x2 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),922 "max.relu.s16x2",923 [(set v2i16:$dst, (smax (smax v2i16:$a, v2i16:$b),924 zeroinitializer<v2i16>))]>;925}926 927//928// Wide multiplication929//930 931def SDTMulWide : SDTypeProfile<1, 2, [SDTCisInt<0>, SDTCisInt<1>, SDTCisSameAs<1, 2>]>;932def smul_wide : SDNode<"NVPTXISD::MUL_WIDE_SIGNED", SDTMulWide, [SDNPCommutative]>;933def umul_wide : SDNode<"NVPTXISD::MUL_WIDE_UNSIGNED", SDTMulWide, [SDNPCommutative]>;934 935 936multiclass MULWIDEInst<string suffix, SDPatternOperator op, RegTyInfo big_t, RegTyInfo small_t> {937 def suffix # _rr :938 BasicNVPTXInst<(outs big_t.RC:$dst), (ins small_t.RC:$a, small_t.RC:$b), 939 "mul.wide." # suffix,940 [(set big_t.Ty:$dst, (op small_t.Ty:$a, small_t.Ty:$b))]>;941 def suffix # _ri :942 BasicNVPTXInst<(outs big_t.RC:$dst), (ins small_t.RC:$a, small_t.Imm:$b), 943 "mul.wide." # suffix,944 [(set big_t.Ty:$dst, (op small_t.Ty:$a, imm:$b))]>;945}946 947defm MUL_WIDE : MULWIDEInst<"s32", smul_wide, I64RT, I32RT>;948defm MUL_WIDE : MULWIDEInst<"u32", umul_wide, I64RT, I32RT>;949defm MUL_WIDE : MULWIDEInst<"s16", smul_wide, I32RT, I16RT>;950defm MUL_WIDE : MULWIDEInst<"u16", umul_wide, I32RT, I16RT>;951 952//953// Integer multiply-add954//955multiclass MADInst<string suffix, SDPatternOperator op, RegTyInfo big_t, RegTyInfo small_t> {956 def rrr:957 BasicNVPTXInst<(outs big_t.RC:$dst),958 (ins small_t.RC:$a, small_t.RC:$b, big_t.RC:$c),959 "mad." # suffix,960 [(set big_t.Ty:$dst, (add (OneUse2<op> small_t.Ty:$a, small_t.Ty:$b), big_t.Ty:$c))]>;961 def rri:962 BasicNVPTXInst<(outs big_t.RC:$dst),963 (ins small_t.RC:$a, small_t.RC:$b, big_t.Imm:$c),964 "mad." # suffix,965 [(set big_t.Ty:$dst, (add (OneUse2<op> small_t.Ty:$a, small_t.Ty:$b), imm:$c))]>;966 def rir:967 BasicNVPTXInst<(outs big_t.RC:$dst),968 (ins small_t.RC:$a, small_t.Imm:$b, big_t.RC:$c),969 "mad." # suffix,970 [(set big_t.Ty:$dst, (add (OneUse2<op> small_t.Ty:$a, imm:$b), big_t.Ty:$c))]>;971 def rii:972 BasicNVPTXInst<(outs big_t.RC:$dst),973 (ins small_t.RC:$a, small_t.Imm:$b, big_t.Imm:$c),974 "mad." # suffix,975 [(set big_t.Ty:$dst, (add (OneUse2<op> small_t.Ty:$a, imm:$b), imm:$c))]>;976}977 978let Predicates = [hasOptEnabled] in {979 defm MAD_LO_S16 : MADInst<"lo.s16", mul, I16RT, I16RT>;980 defm MAD_LO_S32 : MADInst<"lo.s32", mul, I32RT, I32RT>;981 defm MAD_LO_S64 : MADInst<"lo.s64", mul, I64RT, I64RT>;982 983 // Generating mad.wide causes a regression in some cases: 984 // https://github.com/llvm/llvm-project/pull/150477#issuecomment-3191367837985 // Only do so when the user requests it.986 let Predicates = [doMADWideOpt] in {987 defm MAD_WIDE_U16 : MADInst<"wide.u16", umul_wide, I32RT, I16RT>;988 defm MAD_WIDE_S16 : MADInst<"wide.s16", smul_wide, I32RT, I16RT>;989 defm MAD_WIDE_U32 : MADInst<"wide.u32", umul_wide, I64RT, I32RT>;990 defm MAD_WIDE_S32 : MADInst<"wide.s32", smul_wide, I64RT, I32RT>;991 }992}993 994//-----------------------------------995// Floating Point Arithmetic996//-----------------------------------997 998defm FADD : F3_fma_component<"add", fadd>;999defm FSUB : F3_fma_component<"sub", fsub>;1000defm FMUL : F3_fma_component<"mul", fmul>;1001 1002def fminnum_or_fminimumnum : PatFrags<(ops node:$a, node:$b),1003 [(fminnum node:$a, node:$b),1004 (fminimumnum node:$a, node:$b)]>;1005def fmaxnum_or_fmaximumnum : PatFrags<(ops node:$a, node:$b),1006 [(fmaxnum node:$a, node:$b),1007 (fmaximumnum node:$a, node:$b)]>;1008 1009defm MIN : FMINIMUMMAXIMUM<"min", /* NaN */ false, fminnum_or_fminimumnum>;1010defm MAX : FMINIMUMMAXIMUM<"max", /* NaN */ false, fmaxnum_or_fmaximumnum>;1011defm MIN_NAN : FMINIMUMMAXIMUM<"min", /* NaN */ true, fminimum>;1012defm MAX_NAN : FMINIMUMMAXIMUM<"max", /* NaN */ true, fmaximum>;1013 1014def nvptx_fminnum3 : SDNode<"NVPTXISD::FMINNUM3", SDTFPTernaryOp,1015 [SDNPCommutative]>;1016def nvptx_fmaxnum3 : SDNode<"NVPTXISD::FMAXNUM3", SDTFPTernaryOp,1017 [SDNPCommutative]>;1018def nvptx_fminimum3 : SDNode<"NVPTXISD::FMINIMUM3", SDTFPTernaryOp,1019 [SDNPCommutative]>;1020def nvptx_fmaximum3 : SDNode<"NVPTXISD::FMAXIMUM3", SDTFPTernaryOp,1021 [SDNPCommutative]>;1022 1023defm FMIN3 : FMINIMUMMAXIMUM3<"min", /* NaN */ false, nvptx_fminnum3>;1024defm FMAX3 : FMINIMUMMAXIMUM3<"max", /* NaN */ false, nvptx_fmaxnum3>;1025defm FMINNAN3 : FMINIMUMMAXIMUM3<"min", /* NaN */ true, nvptx_fminimum3>;1026defm FMAXNAN3 : FMINIMUMMAXIMUM3<"max", /* NaN */ true, nvptx_fmaximum3>;1027 1028defm FABS : F2<"abs", fabs>;1029defm FNEG : F2<"neg", fneg>;1030defm FABS_H: F2_Support_Half<"abs", fabs>;1031defm FNEG_H: F2_Support_Half<"neg", fneg>;1032 1033defm FSQRT : F2<"sqrt.rn", fsqrt>;1034 1035//1036// F16 NEG1037//1038class FNEG16<RegTyInfo t> :1039 BasicFlagsNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src), (ins FTZFlag:$ftz),1040 "neg$ftz." # t.PtxType,1041 [(set t.Ty:$dst, (fneg t.Ty:$src))]>;1042 1043let Predicates = [useFP16Math, hasPTX<60>, hasSM<53>] in {1044 def NEG_F16 : FNEG16<F16RT>;1045 def NEG_F16x2 : FNEG16<F16X2RT>;1046}1047let Predicates = [hasBF16Math, hasPTX<70>, hasSM<80>] in {1048 def NEG_BF16 : FNEG16<BF16RT>;1049 def NEG_BF16x2 : FNEG16<BF16X2RT>;1050}1051 1052//1053// EX21054//1055 1056class FEXP2Inst<RegTyInfo t, dag flags, string flag_str> :1057 BasicFlagsNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src),1058 flags, "ex2.approx" # flag_str # "." # t.PtxType,1059 [(set t.Ty:$dst, (fexp2 t.Ty:$src))]>;1060 1061def EX2_APPROX_f32 : FEXP2Inst<F32RT, (ins FTZFlag:$ftz), "$ftz">;1062 1063let Predicates = [useFP16Math, hasPTX<70>, hasSM<75>] in {1064 def EX2_APPROX_f16 : FEXP2Inst<F16RT, (ins), "">;1065 def EX2_APPROX_f16x2 : FEXP2Inst<F16X2RT, (ins), "">;1066}1067let Predicates = [hasPTX<78>, hasSM<90>] in {1068 def EX2_APPROX_bf16 : FEXP2Inst<BF16RT, (ins), ".ftz">;1069 def EX2_APPROX_bf16x2 : FEXP2Inst<BF16X2RT, (ins), ".ftz">;1070}1071 1072// F64 division1073//1074def FRCP64r :1075 BasicNVPTXInst<(outs B64:$dst),1076 (ins B64:$b),1077 "rcp.rn.f64",1078 [(set f64:$dst, (fdiv fpimm_1, f64:$b))]>;1079def FDIV64rr :1080 BasicNVPTXInst<(outs B64:$dst),1081 (ins B64:$a, B64:$b),1082 "div.rn.f64",1083 [(set f64:$dst, (fdiv f64:$a, f64:$b))]>;1084def FDIV64ri :1085 BasicNVPTXInst<(outs B64:$dst),1086 (ins B64:$a, f64imm:$b),1087 "div.rn.f64",1088 [(set f64:$dst, (fdiv f64:$a, fpimm:$b))]>;1089 1090// fdiv will be converted to rcp1091// fneg (fdiv 1.0, X) => fneg (rcp.rn X)1092def : Pat<(fdiv fpimm_neg_1, f64:$b),1093 (FNEGf64 (FRCP64r $b))>;1094 1095//1096// F32 Approximate reciprocal1097//1098 1099def fdiv_approx : PatFrag<(ops node:$a, node:$b),1100 (fdiv node:$a, node:$b), [{1101 return getDivF32Level(N) == NVPTX::DivPrecisionLevel::Approx;1102}]>;1103 1104 1105def RCP_APPROX_F32_r :1106 BasicFlagsNVPTXInst<(outs B32:$dst),1107 (ins B32:$b), (ins FTZFlag:$ftz),1108 "rcp.approx$ftz.f32",1109 [(set f32:$dst, (fdiv_approx fpimm_1, f32:$b))]>;1110 1111//1112// F32 Approximate division1113//1114def DIV_APPROX_F32_rr :1115 BasicFlagsNVPTXInst<(outs B32:$dst),1116 (ins B32:$a, B32:$b), (ins FTZFlag:$ftz),1117 "div.approx$ftz.f32",1118 [(set f32:$dst, (fdiv_approx f32:$a, f32:$b))]>;1119def DIV_APPROX_F32_ri :1120 BasicFlagsNVPTXInst<(outs B32:$dst),1121 (ins B32:$a, f32imm:$b), (ins FTZFlag:$ftz),1122 "div.approx$ftz.f32",1123 [(set f32:$dst, (fdiv_approx f32:$a, fpimm:$b))]>;1124//1125// F32 Semi-accurate reciprocal1126//1127// rcp.approx gives the same result as div.full(1.0f, a) and is faster.1128//1129 1130def fdiv_full : PatFrag<(ops node:$a, node:$b),1131 (fdiv node:$a, node:$b), [{1132 return getDivF32Level(N) == NVPTX::DivPrecisionLevel::Full;1133}]>;1134 1135 1136def : Pat<(fdiv_full fpimm_1, f32:$b),1137 (RCP_APPROX_F32_r $b)>;1138 1139//1140// F32 Semi-accurate division1141//1142def FDIV32rr :1143 BasicFlagsNVPTXInst<(outs B32:$dst),1144 (ins B32:$a, B32:$b), (ins FTZFlag:$ftz),1145 "div.full$ftz.f32",1146 [(set f32:$dst, (fdiv_full f32:$a, f32:$b))]>;1147def FDIV32ri :1148 BasicFlagsNVPTXInst<(outs B32:$dst),1149 (ins B32:$a, f32imm:$b), (ins FTZFlag:$ftz),1150 "div.full$ftz.f32",1151 [(set f32:$dst, (fdiv_full f32:$a, fpimm:$b))]>;1152//1153// F32 Accurate reciprocal1154//1155 1156def fdiv_ftz : PatFrag<(ops node:$a, node:$b),1157 (fdiv node:$a, node:$b), [{1158 return getDivF32Level(N) == NVPTX::DivPrecisionLevel::IEEE754;1159}]>;1160 1161def FRCP32r_prec :1162 BasicFlagsNVPTXInst<(outs B32:$dst),1163 (ins B32:$b), (ins FTZFlag:$ftz),1164 "rcp.rn$ftz.f32",1165 [(set f32:$dst, (fdiv_ftz fpimm_1, f32:$b))]>;1166//1167// F32 Accurate division1168//1169def FDIV32rr_prec :1170 BasicFlagsNVPTXInst<(outs B32:$dst),1171 (ins B32:$a, B32:$b), (ins FTZFlag:$ftz),1172 "div.rn$ftz.f32",1173 [(set f32:$dst, (fdiv_ftz f32:$a, f32:$b))]>;1174def FDIV32ri_prec :1175 BasicFlagsNVPTXInst<(outs B32:$dst),1176 (ins B32:$a, f32imm:$b), (ins FTZFlag:$ftz),1177 "div.rn$ftz.f32",1178 [(set f32:$dst, (fdiv_ftz f32:$a, fpimm:$b))]>;1179 1180def : Pat<(fdiv fpimm_1, f32:$b), (FRCP32r_prec $b, NoFTZ)>;1181def : Pat<(fdiv f32:$a, f32:$b), (FDIV32rr_prec $a, $b, NoFTZ)>;1182def : Pat<(fdiv f32:$a, fpimm:$b), (FDIV32ri_prec $a, fpimm:$b, NoFTZ)>;1183 1184//1185// FMA1186//1187 1188multiclass FMA<RegTyInfo t, bit allow_ftz = true, list<Predicate> preds = []> {1189 defvar flag_str = !if(allow_ftz, "$ftz", "");1190 defvar flag_ops = !if(allow_ftz, (ins FTZFlag:$ftz), (ins));1191 defvar op_str = "fma.rn" # flag_str # "." # t.PtxType;1192 1193 let Predicates = preds in {1194 def rrr : BasicFlagsNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a, t.RC:$b, t.RC:$c),1195 flag_ops, op_str,1196 [(set t.Ty:$dst, (fma t.Ty:$a, t.Ty:$b, t.Ty:$c))]>;1197 1198 if t.SupportsImm then {1199 def rri : BasicFlagsNVPTXInst<(outs t.RC:$dst),1200 (ins t.RC:$a, t.RC:$b, t.Imm:$c),1201 flag_ops, op_str,1202 [(set t.Ty:$dst, (fma t.Ty:$a, t.Ty:$b, fpimm:$c))]>;1203 def rir : BasicFlagsNVPTXInst<(outs t.RC:$dst),1204 (ins t.RC:$a, t.Imm:$b, t.RC:$c),1205 flag_ops, op_str,1206 [(set t.Ty:$dst, (fma t.Ty:$a, fpimm:$b, t.Ty:$c))]>;1207 def rii : BasicFlagsNVPTXInst<(outs t.RC:$dst),1208 (ins t.RC:$a, t.Imm:$b, t.Imm:$c),1209 flag_ops, op_str,1210 [(set t.Ty:$dst, (fma t.Ty:$a, fpimm:$b, fpimm:$c))]>;1211 def iir : BasicFlagsNVPTXInst<(outs t.RC:$dst),1212 (ins t.Imm:$a, t.Imm:$b, t.RC:$c),1213 flag_ops, op_str,1214 [(set t.Ty:$dst, (fma fpimm:$a, fpimm:$b, t.Ty:$c))]>;1215 }1216 }1217}1218 1219defm FMA_F16 : FMA<F16RT, allow_ftz = true, preds = [useFP16Math]>;1220defm FMA_F16x2 : FMA<F16X2RT, allow_ftz = true, preds = [useFP16Math]>;1221defm FMA_BF16 : FMA<BF16RT, allow_ftz = false, preds = [hasBF16Math]>;1222defm FMA_BF16x2 : FMA<BF16X2RT, allow_ftz = false, preds = [hasBF16Math]>;1223defm FMA_F32 : FMA<F32RT, allow_ftz = true>;1224defm FMA_F32x2 : FMA<F32X2RT, allow_ftz = true, preds = [hasF32x2Instructions]>;1225defm FMA_F64 : FMA<F64RT, allow_ftz = false>;1226 1227// sin/cos/tanh1228 1229class UnaryOpAllowsApproxFn<SDPatternOperator operator>1230 : PatFrag<(ops node:$A), (operator node:$A), [{1231 return N->getFlags().hasApproximateFuncs();1232}]>;1233 1234def SIN_APPROX_f32 :1235 BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$src), (ins FTZFlag:$ftz),1236 "sin.approx$ftz.f32",1237 [(set f32:$dst, (UnaryOpAllowsApproxFn<fsin> f32:$src))]>;1238def COS_APPROX_f32 :1239 BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$src), (ins FTZFlag:$ftz),1240 "cos.approx$ftz.f32",1241 [(set f32:$dst, (UnaryOpAllowsApproxFn<fcos> f32:$src))]>;1242def TANH_APPROX_f32 :1243 BasicNVPTXInst<(outs B32:$dst), (ins B32:$src), "tanh.approx.f32",1244 [(set f32:$dst, (UnaryOpAllowsApproxFn<ftanh> f32:$src))]>,1245 Requires<[hasPTX<70>, hasSM<75>]>;1246 1247//-----------------------------------1248// Bitwise operations1249//-----------------------------------1250 1251// Template for three-arg bitwise operations. Takes three args, Creates .b16,1252// .b32, .b64, and .pred (predicate registers -- i.e., i1) versions of OpcStr.1253multiclass BITWISE<string OpcStr, SDNode OpNode> {1254 foreach t = [I1RT, I16RT, I32RT, I64RT] in1255 defm _ # t.PtxType : I3Inst<OpcStr # "." # t.PtxType, OpNode, t, commutative = true>;1256}1257 1258defm OR : BITWISE<"or", or>;1259defm AND : BITWISE<"and", and>;1260defm XOR : BITWISE<"xor", xor>;1261 1262// PTX does not support mul on predicates, convert to and instructions1263def : Pat<(mul i1:$a, i1:$b), (AND_predrr $a, $b)>;1264def : Pat<(mul i1:$a, imm:$b), (AND_predri $a, imm:$b)>;1265 1266foreach op = [add, sub] in {1267 def : Pat<(op i1:$a, i1:$b), (XOR_predrr $a, $b)>;1268 def : Pat<(op i1:$a, imm:$b), (XOR_predri $a, imm:$b)>;1269}1270 1271// These transformations were once reliably performed by instcombine, but thanks1272// to poison semantics they are no longer safe for LLVM IR, perform them here1273// instead.1274def : Pat<(select i1:$a, i1:$b, 0), (AND_predrr $a, $b)>;1275def : Pat<(select i1:$a, 1, i1:$b), (OR_predrr $a, $b)>;1276 1277// Lower logical v2i16/v4i8 ops as bitwise ops on b32.1278foreach vt = [v2i16, v4i8] in {1279 def : Pat<(or vt:$a, vt:$b), (OR_b32rr $a, $b)>;1280 def : Pat<(xor vt:$a, vt:$b), (XOR_b32rr $a, $b)>;1281 def : Pat<(and vt:$a, vt:$b), (AND_b32rr $a, $b)>;1282 1283 // The constants get legalized into a bitcast from i32, so that's what we need1284 // to match here.1285 def: Pat<(or vt:$a, (vt (bitconvert (i32 imm:$b)))),1286 (OR_b32ri $a, imm:$b)>;1287 def: Pat<(xor vt:$a, (vt (bitconvert (i32 imm:$b)))),1288 (XOR_b32ri $a, imm:$b)>;1289 def: Pat<(and vt:$a, (vt (bitconvert (i32 imm:$b)))),1290 (AND_b32ri $a, imm:$b)>;1291}1292 1293foreach t = [I1RT, I16RT, I32RT, I64RT] in1294 def NOT_ # t.PtxType : BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src),1295 "not." # t.PtxType,1296 [(set t.Ty:$dst, (not t.Ty:$src))]>;1297 1298// Template for left/right shifts. Takes three operands,1299// [dest (reg), src (reg), shift (reg or imm)].1300// dest and src may be int64, int32, or int16, but shift is always int32.1301//1302// This template also defines a 32-bit shift (imm, imm) instruction.1303multiclass SHIFT<string OpcStr, SDNode OpNode> {1304 let hasSideEffects = false in {1305 foreach t = [I64RT, I32RT, I16RT] in {1306 def t.Size # _rr :1307 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a, B32:$b),1308 OpcStr # t.Size,1309 [(set t.Ty:$dst, (OpNode t.Ty:$a, i32:$b))]>;1310 def t.Size # _ri :1311 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a, i32imm:$b),1312 OpcStr # t.Size,1313 [(set t.Ty:$dst, (OpNode t.Ty:$a, (i32 imm:$b)))]>;1314 def t.Size # _ii :1315 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a, i32imm:$b),1316 OpcStr # t.Size,1317 [(set t.Ty:$dst, (OpNode (t.Ty imm:$a), (i32 imm:$b)))]>;1318 }1319 }1320}1321 1322defm SHL : SHIFT<"shl.b", shl>;1323defm SRA : SHIFT<"shr.s", sra>;1324defm SRL : SHIFT<"shr.u", srl>;1325 1326// Bit-reverse1327foreach t = [I64RT, I32RT] in1328 def BREV_ # t.PtxType :1329 BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a),1330 "brev." # t.PtxType,1331 [(set t.Ty:$dst, (bitreverse t.Ty:$a))]>;1332 1333 1334//1335// BFE - bit-field extract1336//1337 1338// Template for BFE/BFI instructions.1339// Args: [dest (reg), src (reg), start (reg or imm), end (reg or imm)].1340// Start may be an imm only if end is also an imm. FIXME: Is this a1341// restriction in PTX?1342//1343// dest and src may be int32 or int64, but start and end are always int32.1344def SDTBFI :1345 SDTypeProfile<1, 4, [SDTCisInt<0>, SDTCisSameAs<0, 1>, SDTCisSameAs<0, 2>, 1346 SDTCisVT<3, i32>, SDTCisVT<4, i32>]>;1347def bfi : SDNode<"NVPTXISD::BFI", SDTBFI>;1348 1349def SDTPRMT :1350 SDTypeProfile<1, 4, [SDTCisVT<0, i32>, SDTCisVT<1, i32>,1351 SDTCisVT<2, i32>, SDTCisVT<3, i32>, SDTCisVT<4, i32>]>;1352def prmt : SDNode<"NVPTXISD::PRMT", SDTPRMT>;1353 1354multiclass BFE<string Instr, RegisterClass RC> {1355 def rrr1356 : BasicNVPTXInst<(outs RC:$d), (ins RC:$a, B32:$b, B32:$c), Instr>;1357 def rri1358 : BasicNVPTXInst<(outs RC:$d), (ins RC:$a, B32:$b, i32imm:$c), Instr>;1359 def rii1360 : BasicNVPTXInst<(outs RC:$d), (ins RC:$a, i32imm:$b, i32imm:$c), Instr>;1361}1362 1363multiclass BFI<string Instr, ValueType T, RegisterClass RC, Operand ImmCls> {1364 def rrrr1365 : BasicNVPTXInst<(outs RC:$f),1366 (ins RC:$a, RC:$b, B32:$c, B32:$d),1367 Instr,1368 [(set T:$f, (bfi T:$a, T:$b, i32:$c, i32:$d))]>;1369 def rrri1370 : BasicNVPTXInst<(outs RC:$f),1371 (ins RC:$a, RC:$b, B32:$c, i32imm:$d),1372 Instr,1373 [(set T:$f, (bfi T:$a, T:$b, i32:$c, imm:$d))]>;1374 def rrii1375 : BasicNVPTXInst<(outs RC:$f),1376 (ins RC:$a, RC:$b, i32imm:$c, i32imm:$d),1377 Instr,1378 [(set T:$f, (bfi T:$a, T:$b, imm:$c, imm:$d))]>;1379 def irrr1380 : BasicNVPTXInst<(outs RC:$f),1381 (ins ImmCls:$a, RC:$b, B32:$c, B32:$d),1382 Instr,1383 [(set T:$f, (bfi (T imm:$a), T:$b, i32:$c, i32:$d))]>;1384 def irri1385 : BasicNVPTXInst<(outs RC:$f),1386 (ins ImmCls:$a, RC:$b, B32:$c, i32imm:$d),1387 Instr,1388 [(set T:$f, (bfi (T imm:$a), T:$b, i32:$c, imm:$d))]>;1389 def irii1390 : BasicNVPTXInst<(outs RC:$f),1391 (ins ImmCls:$a, RC:$b, i32imm:$c, i32imm:$d),1392 Instr,1393 [(set T:$f, (bfi (T imm:$a), T:$b, imm:$c, imm:$d))]>;1394}1395 1396def Hexu32imm : Operand<i32> {1397 let PrintMethod = "printHexu32imm";1398}1399 1400let hasSideEffects = false in {1401 // order is somewhat important here. signed/unsigned variants match1402 // the same patterns, so the first one wins. Having unsigned byte extraction1403 // has the benefit of always having zero in unused bits, which makes some1404 // optimizations easier (e.g. no need to mask them).1405 defm BFE_U32 : BFE<"bfe.u32", B32>;1406 defm BFE_S32 : BFE<"bfe.s32", B32>;1407 defm BFE_U64 : BFE<"bfe.u64", B64>;1408 defm BFE_S64 : BFE<"bfe.s64", B64>;1409 1410 defm BFI_B32 : BFI<"bfi.b32", i32, B32, i32imm>;1411 defm BFI_B64 : BFI<"bfi.b64", i64, B64, i64imm>;1412 1413 def PRMT_B32rrr1414 : BasicFlagsNVPTXInst<(outs B32:$d),1415 (ins B32:$a, B32:$b, B32:$c),1416 (ins PrmtMode:$mode),1417 "prmt.b32$mode",1418 [(set i32:$d, (prmt i32:$a, i32:$b, i32:$c, imm:$mode))]>;1419 def PRMT_B32rri1420 : BasicFlagsNVPTXInst<(outs B32:$d),1421 (ins B32:$a, B32:$b, Hexu32imm:$c),1422 (ins PrmtMode:$mode),1423 "prmt.b32$mode",1424 [(set i32:$d, (prmt i32:$a, i32:$b, imm:$c, imm:$mode))]>;1425 def PRMT_B32rir1426 : BasicFlagsNVPTXInst<(outs B32:$d),1427 (ins B32:$a, i32imm:$b, B32:$c),1428 (ins PrmtMode:$mode),1429 "prmt.b32$mode",1430 [(set i32:$d, (prmt i32:$a, imm:$b, i32:$c, imm:$mode))]>;1431 def PRMT_B32rii1432 : BasicFlagsNVPTXInst<(outs B32:$d),1433 (ins B32:$a, i32imm:$b, Hexu32imm:$c),1434 (ins PrmtMode:$mode),1435 "prmt.b32$mode",1436 [(set i32:$d, (prmt i32:$a, imm:$b, imm:$c, imm:$mode))]>;1437 def PRMT_B32irr1438 : BasicFlagsNVPTXInst<(outs B32:$d),1439 (ins i32imm:$a, B32:$b, B32:$c), (ins PrmtMode:$mode),1440 "prmt.b32$mode",1441 [(set i32:$d, (prmt imm:$a, i32:$b, i32:$c, imm:$mode))]>;1442 def PRMT_B32iri1443 : BasicFlagsNVPTXInst<(outs B32:$d),1444 (ins i32imm:$a, B32:$b, Hexu32imm:$c), (ins PrmtMode:$mode),1445 "prmt.b32$mode",1446 [(set i32:$d, (prmt imm:$a, i32:$b, imm:$c, imm:$mode))]>;1447 def PRMT_B32iir1448 : BasicFlagsNVPTXInst<(outs B32:$d),1449 (ins i32imm:$a, i32imm:$b, B32:$c), (ins PrmtMode:$mode),1450 "prmt.b32$mode",1451 [(set i32:$d, (prmt imm:$a, imm:$b, i32:$c, imm:$mode))]>;1452 1453}1454 1455// PRMT folding patterns1456def : Pat<(fshr i32:$hi, i32:$lo, (shl i32:$amt, (i32 3))),1457 (PRMT_B32rrr $lo, $hi, $amt, PrmtF4E)>;1458 1459 1460def byte_extract_prmt : ImmLeaf<i32, [{1461 return (Imm == 0x7770) || (Imm == 0x7771) || (Imm == 0x7772) || (Imm == 0x7773);1462}]>;1463 1464def to_sign_extend_selector : SDNodeXForm<imm, [{1465 const APInt &V = N->getAPIntValue();1466 const APInt B = V.trunc(4);1467 const APInt BSext = B | 8;1468 const APInt R = BSext.concat(BSext).concat(BSext).concat(B).zext(32);1469 return CurDAG->getTargetConstant(R, SDLoc(N), MVT::i32);1470}]>;1471 1472 1473// byte extraction + signed/unsigned extension to i32.1474def : Pat<(i32 (sext_inreg (prmt i32:$s, 0, byte_extract_prmt:$sel, PrmtNONE), i8)),1475 (PRMT_B32rii $s, 0, (to_sign_extend_selector $sel), PrmtNONE)>;1476 1477// byte extraction + signed extension to i161478def : Pat<(i16 (sext_inreg (trunc (prmt i32:$s, 0, byte_extract_prmt:$sel, PrmtNONE)), i8)),1479 (CVT_u16_u32 (PRMT_B32rii $s, 0, (to_sign_extend_selector $sel), PrmtNONE), CvtNONE)>;1480 1481 1482// Byte extraction via shift/trunc/sext1483def : Pat<(i16 (sext_inreg (trunc i32:$s), i8)), (CVT_s8_s32 $s, CvtNONE)>;1484def : Pat<(i16 (sext_inreg (trunc i64:$s), i8)), (CVT_s8_s64 $s, CvtNONE)>;1485 1486def : Pat<(sext_inreg (srl i32:$s, (i32 imm:$o)), i8), (BFE_S32rii $s, imm:$o, 8)>;1487def : Pat<(sext_inreg (srl i64:$s, (i32 imm:$o)), i8), (BFE_S64rii $s, imm:$o, 8)>;1488 1489def : Pat<(i16 (sext_inreg (trunc (srl i32:$s, (i32 imm:$o))), i8)),1490 (CVT_s8_s32 (BFE_S32rii $s, imm:$o, 8), CvtNONE)>;1491def : Pat<(i16 (sext_inreg (trunc (srl i64:$s, (i32 imm:$o))), i8)),1492 (CVT_s8_s64 (BFE_S64rii $s, imm:$o, 8), CvtNONE)>;1493 1494def : Pat<(i16 (sra (i16 (trunc i32:$s)), (i32 8))),1495 (CVT_s8_s32 (BFE_S32rii $s, 8, 8), CvtNONE)>;1496 1497//-----------------------------------1498// Comparison instructions (setp, set)1499//-----------------------------------1500 1501// FIXME: This doesn't cover versions of set and setp that combine with a1502// boolean predicate, e.g. setp.eq.and.b16.1503def cond2cc : SDNodeXForm<cond, [{1504 return getPTXCmpMode(*N);1505}]>;1506 1507multiclass FSETP<RegTyInfo t, bit allow_ftz = true> {1508 defvar ftz_str = !if(allow_ftz, "$ftz", "");1509 defvar op_str = "setp.${cmp:FCmp}" # ftz_str # "." # t.PtxType;1510 defvar flags = !con((ins CmpMode:$cmp), !if(allow_ftz, (ins FTZFlag:$ftz), (ins)));1511 let hasSideEffects = false in {1512 def rr :1513 BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.RC:$a, t.RC:$b),1514 flags, op_str>;1515 1516 if t.SupportsImm then {1517 def ri :1518 BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.RC:$a, t.Imm:$b),1519 flags, op_str>;1520 def ir :1521 BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.Imm:$a, t.RC:$b),1522 flags, op_str>;1523 }1524 }1525 def : Pat<(i1 (setcc t.Ty:$a, t.Ty:$b, cond:$cc)),1526 (!cast<NVPTXInst>(NAME # "rr") $a, $b, (cond2cc $cc))>;1527 if t.SupportsImm then {1528 def : Pat<(i1 (setcc t.Ty:$a, fpimm:$b, cond:$cc)),1529 (!cast<NVPTXInst>(NAME # "ri") $a, fpimm:$b, (cond2cc $cc))>;1530 def : Pat<(i1 (setcc fpimm:$a, t.Ty:$b, cond:$cc)),1531 (!cast<NVPTXInst>(NAME # "ir") fpimm:$a, $b, (cond2cc $cc))>;1532 }1533}1534 1535multiclass ISETP<RegTyInfo t> {1536 defvar op_str = "setp.${cmp:ICmp}.${cmp:IType}" # t.Size;1537 let hasSideEffects = false in {1538 def rr :1539 BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.RC:$a, t.RC:$b),1540 (ins CmpMode:$cmp), op_str>;1541 def ri :1542 BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.RC:$a, t.Imm:$b),1543 (ins CmpMode:$cmp), op_str>;1544 def ir :1545 BasicFlagsNVPTXInst<(outs B1:$dst), (ins t.Imm:$a, t.RC:$b),1546 (ins CmpMode:$cmp), op_str>;1547 }1548 def : Pat<(i1 (setcc t.Ty:$a, t.Ty:$b, cond:$cc)),1549 (!cast<NVPTXInst>(NAME # "rr") $a, $b, (cond2cc $cc))>;1550 def : Pat<(i1 (setcc t.Ty:$a, imm:$b, cond:$cc)),1551 (!cast<NVPTXInst>(NAME # "ri") $a, imm:$b, (cond2cc $cc))>;1552 def : Pat<(i1 (setcc imm:$a, t.Ty:$b, cond:$cc)),1553 (!cast<NVPTXInst>(NAME # "ir") imm:$a, $b, (cond2cc $cc))>;1554}1555 1556defm SETP_i16 : ISETP<I16RT>;1557defm SETP_i32 : ISETP<I32RT>;1558defm SETP_i64 : ISETP<I64RT>;1559 1560defm SETP_f32 : FSETP<F32RT>;1561defm SETP_f64 : FSETP<F64RT, allow_ftz = false>;1562let Predicates = [useFP16Math] in1563 defm SETP_f16 : FSETP<F16RT>;1564let Predicates = [hasBF16Math, hasPTX<78>, hasSM<90>] in1565 defm SETP_bf16 : FSETP<BF16RT>;1566 1567def SETP_f16x2rr :1568 BasicFlagsNVPTXInst<(outs B1:$p, B1:$q),1569 (ins B32:$a, B32:$b), (ins CmpMode:$cmp, FTZFlag:$ftz),1570 "setp.${cmp:FCmp}$ftz.f16x2">,1571 Requires<[useFP16Math]>;1572 1573def SETP_bf16x2rr :1574 BasicFlagsNVPTXInst<(outs B1:$p, B1:$q),1575 (ins B32:$a, B32:$b), (ins CmpMode:$cmp, FTZFlag:$ftz),1576 "setp.${cmp:FCmp}$ftz.bf16x2">,1577 Requires<[hasBF16Math, hasPTX<78>, hasSM<90>]>;1578 1579//-----------------------------------1580// Data Movement (Load / Store, Move)1581//-----------------------------------1582 1583def addr : ComplexPattern<pAny, 2, "SelectADDR">;1584 1585def ADDR_base : Operand<pAny>;1586def ADDR : Operand<pAny> {1587 let PrintMethod = "printMemOperand";1588 let MIOperandInfo = (ops ADDR_base, i32imm);1589}1590 1591def UsedBytesMask : Operand<i32> {1592 let PrintMethod = "printUsedBytesMaskPragma";1593}1594 1595def RegOrSink : Operand<Any> {1596 let PrintMethod = "printRegisterOrSinkSymbol";1597}1598 1599def AtomicCode : Operand<i32> {1600 let PrintMethod = "printAtomicCode";1601}1602 1603def MmaCode : Operand<i32> {1604 let PrintMethod = "printMmaCode";1605}1606 1607// Get pointer to local stack.1608let hasSideEffects = false in {1609 def MOV_DEPOT_ADDR : NVPTXInst<(outs B32:$d), (ins i32imm:$num),1610 "mov.b32 \t$d, __local_depot$num;">;1611 def MOV_DEPOT_ADDR_64 : NVPTXInst<(outs B64:$d), (ins i32imm:$num),1612 "mov.b64 \t$d, __local_depot$num;">;1613}1614 1615let hasSideEffects = false in {1616 let isMoveReg = true, isAsCheapAsAMove = true in1617 class MOVr<RegisterClass RC, string OpStr> :1618 BasicNVPTXInst<(outs RC:$dst), (ins RC:$src), "mov." # OpStr>;1619 1620 let isMoveImm = true, isAsCheapAsAMove = true in1621 class MOVi<RegTyInfo t, string suffix> :1622 BasicNVPTXInst<(outs t.RC:$dst), (ins t.Imm:$src),1623 "mov." # suffix,1624 [(set t.Ty:$dst, t.ImmNode:$src)]>;1625 1626 // We don't want to set isAsCheapAsAMove to true for these instructions as1627 // this would prevent CSE and resulted in regressions (see discussion after1628 // PR-145581 in llvm-project).1629 class MovSymInst<RegTyInfo t> :1630 BasicNVPTXInst<(outs t.RC:$dst), (ins Operand<t.Ty>:$src),1631 "mov.b" # t.Size>;1632}1633 1634def MOV_B1_r : MOVr<B1, "pred">;1635def MOV_B16_r : MOVr<B16, "b16">;1636def MOV_B32_r : MOVr<B32, "b32">;1637def MOV_B64_r : MOVr<B64, "b64">;1638def MOV_B128_r : MOVr<B128, "b128">;1639 1640def MOV_B1_i : MOVi<I1RT, "pred">;1641def MOV_B16_i : MOVi<I16RT, "b16">;1642def MOV_B32_i : MOVi<I32RT, "b32">;1643def MOV_B64_i : MOVi<I64RT, "b64">;1644def MOV_F16_i : MOVi<F16RT, "b16">;1645def MOV_BF16_i : MOVi<BF16RT, "b16">;1646def MOV_F32_i : MOVi<F32RT, "b32">;1647def MOV_F64_i : MOVi<F64RT, "b64">;1648 1649def MOV_B32_sym : MovSymInst<I32RT>;1650def MOV_B64_sym : MovSymInst<I64RT>;1651 1652 1653def to_tglobaladdr : SDNodeXForm<globaladdr, [{1654 return CurDAG->getTargetGlobalAddress(N->getGlobal(), SDLoc(N),1655 N->getValueType(0), N->getOffset(),1656 N->getTargetFlags());1657}]>;1658 1659def to_texternsym : SDNodeXForm<externalsym, [{1660 return CurDAG->getTargetExternalSymbol(N->getSymbol(), N->getValueType(0),1661 N->getTargetFlags());1662}]>;1663 1664def to_tframeindex : SDNodeXForm<frameindex, [{1665 return CurDAG->getTargetFrameIndex(N->getIndex(), N->getValueType(0));1666}]>;1667 1668def : Pat<(i32 globaladdr:$dst), (MOV_B32_sym (to_tglobaladdr $dst))>;1669def : Pat<(i64 globaladdr:$dst), (MOV_B64_sym (to_tglobaladdr $dst))>;1670 1671def : Pat<(i32 externalsym:$dst), (MOV_B32_sym (to_texternsym $dst))>;1672def : Pat<(i64 externalsym:$dst), (MOV_B64_sym (to_texternsym $dst))>;1673 1674//---- Copy Frame Index ----1675def LEA_ADDRi : NVPTXInst<(outs B32:$dst), (ins ADDR:$addr),1676 "add.u32 \t$dst, ${addr:add};">;1677def LEA_ADDRi64 : NVPTXInst<(outs B64:$dst), (ins ADDR:$addr),1678 "add.u64 \t$dst, ${addr:add};">;1679 1680def : Pat<(i32 frameindex:$fi), (LEA_ADDRi (to_tframeindex $fi), 0)>;1681def : Pat<(i64 frameindex:$fi), (LEA_ADDRi64 (to_tframeindex $fi), 0)>;1682 1683//-----------------------------------1684// Comparison and Selection1685//-----------------------------------1686// TODO: These patterns seem very specific and brittle. We should try to find1687// a more general solution.1688 1689def cond_signed : PatLeaf<(cond), [{1690 return isSignedIntSetCC(N->get());1691}]>;1692 1693// A 16-bit signed comparison of sign-extended byte extracts can be converted1694// to 32-bit comparison if we change the PRMT to sign-extend the extracted1695// bytes.1696def : Pat<(setcc (i16 (sext_inreg (trunc (prmt i32:$a, 0, byte_extract_prmt:$sel_a, PrmtNONE)), i8)),1697 (i16 (sext_inreg (trunc (prmt i32:$b, 0, byte_extract_prmt:$sel_b, PrmtNONE)), i8)),1698 cond_signed:$cc),1699 (SETP_i32rr (PRMT_B32rii i32:$a, 0, (to_sign_extend_selector $sel_a), PrmtNONE),1700 (PRMT_B32rii i32:$b, 0, (to_sign_extend_selector $sel_b), PrmtNONE),1701 (cond2cc $cc))>;1702 1703def SDTDeclareArrayParam :1704 SDTypeProfile<0, 3, [SDTCisVT<0, i32>, SDTCisVT<1, i32>, SDTCisVT<2, i32>]>;1705def SDTDeclareScalarParam :1706 SDTypeProfile<0, 2, [SDTCisVT<0, i32>, SDTCisVT<1, i32>]>;1707def SDTMoveParamProfile : SDTypeProfile<1, 1, [SDTCisInt<0>, SDTCisSameAs<0, 1>]>;1708 1709def SDTProxyReg : SDTypeProfile<1, 1, [SDTCisSameAs<0, 1>]>;1710 1711// These nodes represent a parameter declaration. In PTX this will look like:1712// .param .align 16 .b8 param0[1024];1713// .param .b32 retval0;1714//1715// DeclareArrayParam(Chain, Externalsym, Align, Size, Glue)1716// DeclareScalarParam(Chain, Externalsym, Size, Glue)1717def declare_array_param :1718 SDNode<"NVPTXISD::DeclareArrayParam", SDTDeclareArrayParam,1719 [SDNPHasChain, SDNPOutGlue, SDNPInGlue, SDNPSideEffect]>;1720def declare_scalar_param :1721 SDNode<"NVPTXISD::DeclareScalarParam", SDTDeclareScalarParam,1722 [SDNPHasChain, SDNPOutGlue, SDNPInGlue, SDNPSideEffect]>;1723 1724def MoveParam :1725 SDNode<"NVPTXISD::MoveParam", SDTMoveParamProfile, []>;1726def proxy_reg :1727 SDNode<"NVPTXISD::ProxyReg", SDTProxyReg, [SDNPHasChain]>;1728 1729 /// CALL(Chain, IsConvergent, IsIndirectCall/IsUniform, NumReturns,1730 /// NumParams, Callee, Proto)1731def SDTCallProfile : SDTypeProfile<0, 6,1732 [SDTCisVT<0, i32>, SDTCisVT<1, i32>, SDTCisVT<2, i32>,1733 SDTCisVT<3, i32>, SDTCisVT<5, i32>]>;1734def call : SDNode<"NVPTXISD::CALL", SDTCallProfile, [SDNPHasChain, SDNPSideEffect]>;1735 1736/// CALL(Chain, IsConvergent, IsIndirectCall/IsUniform, NumReturns,1737/// NumParams, Callee, Proto)1738 1739def CallOperand : Operand<i32> { let PrintMethod = "printCallOperand"; }1740 1741foreach is_convergent = [0, 1] in {1742 defvar convergent_suffix = !if(is_convergent, "_conv", "");1743 1744 let isCall = 1, isConvergent = is_convergent in {1745 def CALL # convergent_suffix :1746 NVPTXInst<(outs),1747 (ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params, 1748 i32imm:$proto),1749 "call${rets:RetList} $addr, (${params:ParamList}), prototype_$proto;">;1750 1751 def CALL_UNI # convergent_suffix :1752 NVPTXInst<(outs),1753 (ins ADDR_base:$addr, CallOperand:$rets, CallOperand:$params),1754 "call.uni${rets:RetList} $addr, (${params:ParamList});">;1755 }1756 1757 defvar call_inst = !cast<NVPTXInst>("CALL" # convergent_suffix);1758 def : Pat<(call is_convergent, 1, imm:$rets, imm:$params, i32:$addr, imm:$proto),1759 (call_inst $addr, imm:$rets, imm:$params, imm:$proto)>;1760 def : Pat<(call is_convergent, 1, imm:$rets, imm:$params, i64:$addr, imm:$proto),1761 (call_inst $addr, imm:$rets, imm:$params, imm:$proto)>;1762 1763 defvar call_uni_inst = !cast<NVPTXInst>("CALL_UNI" # convergent_suffix);1764 def : Pat<(call is_convergent, 0, imm:$rets, imm:$params, globaladdr:$addr, 0),1765 (call_uni_inst (to_tglobaladdr $addr), imm:$rets, imm:$params)>;1766}1767 1768def DECLARE_PARAM_array :1769 NVPTXInst<(outs), (ins i32imm:$a, i32imm:$align, i32imm:$size),1770 ".param .align $align .b8 \t$a[$size];">;1771def DECLARE_PARAM_scalar :1772 NVPTXInst<(outs), (ins i32imm:$a, i32imm:$size),1773 ".param .b$size \t$a;">;1774 1775def : Pat<(declare_array_param externalsym:$a, imm:$align, imm:$size),1776 (DECLARE_PARAM_array (to_texternsym $a), imm:$align, imm:$size)>;1777def : Pat<(declare_scalar_param externalsym:$a, imm:$size),1778 (DECLARE_PARAM_scalar (to_texternsym $a), imm:$size)>;1779 1780// Call prototype wrapper, this is a dummy instruction that just prints it's1781// operand which is string defining the prototype.1782def SDTCallPrototype : SDTypeProfile<0, 1, [SDTCisInt<0>]>;1783def CallPrototype :1784 SDNode<"NVPTXISD::CallPrototype", SDTCallPrototype,1785 [SDNPHasChain, SDNPSideEffect]>;1786def ProtoIdent : Operand<i32> { let PrintMethod = "printProtoIdent"; }1787def CALL_PROTOTYPE :1788 NVPTXInst<(outs), (ins ProtoIdent:$ident),1789 "$ident", [(CallPrototype (i32 texternalsym:$ident))]>;1790 1791 1792foreach t = [I32RT, I64RT] in {1793 defvar inst_name = "MOV" # t.Size # "_PARAM";1794 def inst_name : BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src), "mov.b" # t.Size>;1795 def : Pat<(MoveParam (t.Ty externalsym:$src)),1796 (!cast<NVPTXInst>(inst_name) (t.Ty (to_texternsym $src)))>;1797}1798 1799multiclass ProxyRegInst<string SzStr, NVPTXRegClass rc> {1800 def NAME : BasicNVPTXInst<(outs rc:$dst), (ins rc:$src),1801 "mov." # SzStr>;1802 foreach vt = rc.RegTypes in1803 def : Pat<(vt (proxy_reg vt:$src)), (!cast<NVPTXInst>(NAME) $src)>;1804}1805 1806defm ProxyRegB1 : ProxyRegInst<"pred", B1>;1807defm ProxyRegB16 : ProxyRegInst<"b16", B16>;1808defm ProxyRegB32 : ProxyRegInst<"b32", B32>;1809defm ProxyRegB64 : ProxyRegInst<"b64", B64>;1810 1811 1812// Callseq start and end1813 1814// Note: these nodes are marked as SDNPMayStore and SDNPMayLoad because1815// they define the scope in which the declared params may be used. Therefore1816// we add these flags to ensure ld.param and st.param are not sunk or hoisted1817// out of that scope.1818 1819def callseq_start : SDNode<"ISD::CALLSEQ_START",1820 SDCallSeqStart<[SDTCisVT<0, i32>, SDTCisVT<1, i32>]>,1821 [SDNPHasChain, SDNPOutGlue,1822 SDNPSideEffect, SDNPMayStore, SDNPMayLoad]>;1823def callseq_end : SDNode<"ISD::CALLSEQ_END",1824 SDCallSeqEnd<[SDTCisVT<0, i32>, SDTCisVT<1, i32>]>,1825 [SDNPHasChain, SDNPOptInGlue, SDNPOutGlue,1826 SDNPSideEffect, SDNPMayStore, SDNPMayLoad]>;1827 1828def Callseq_Start :1829 NVPTXInst<(outs), (ins i32imm:$amt1, i32imm:$amt2),1830 "\\{ // callseq $amt1, $amt2",1831 [(callseq_start timm:$amt1, timm:$amt2)]>;1832def Callseq_End :1833 NVPTXInst<(outs), (ins i32imm:$amt1, i32imm:$amt2),1834 "\\} // callseq $amt1",1835 [(callseq_end timm:$amt1, timm:$amt2)]>;1836 1837//1838// Load / Store Handling1839//1840class LD<NVPTXRegClass regclass>1841 : NVPTXInst<1842 (outs regclass:$dst),1843 (ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,1844 AtomicCode:$Sign, i32imm:$fromWidth, UsedBytesMask:$usedBytes,1845 ADDR:$addr),1846 "${usedBytes}"1847 "ld${sem:sem}${scope:scope}${addsp:addsp}.${Sign:sign}$fromWidth "1848 "\t$dst, [$addr];">;1849 1850let mayLoad=1, hasSideEffects=0 in {1851 def LD_i16 : LD<B16>;1852 def LD_i32 : LD<B32>;1853 def LD_i64 : LD<B64>;1854}1855 1856class ST<DAGOperand O>1857 : NVPTXInst<1858 (outs),1859 (ins O:$src,1860 AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$toWidth,1861 ADDR:$addr),1862 "st${sem:sem}${scope:scope}${addsp:addsp}.b$toWidth"1863 " \t[$addr], $src;">;1864 1865let mayStore=1, hasSideEffects=0 in {1866 def ST_i16 : ST<RI16>;1867 def ST_i32 : ST<RI32>;1868 def ST_i64 : ST<RI64>;1869}1870 1871// The following is used only in and after vector elementizations. Vector1872// elementization happens at the machine instruction level, so the following1873// instructions never appear in the DAG.1874multiclass LD_VEC<NVPTXRegClass regclass, bit support_v8 = false> {1875 def _v2 : NVPTXInst<1876 (outs regclass:$dst1, regclass:$dst2),1877 (ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,1878 AtomicCode:$Sign, i32imm:$fromWidth, UsedBytesMask:$usedBytes,1879 ADDR:$addr),1880 "${usedBytes}"1881 "ld${sem:sem}${scope:scope}${addsp:addsp}.v2.${Sign:sign}$fromWidth "1882 "\t{{$dst1, $dst2}}, [$addr];">;1883 def _v4 : NVPTXInst<1884 (outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4),1885 (ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,1886 AtomicCode:$Sign, i32imm:$fromWidth, UsedBytesMask:$usedBytes,1887 ADDR:$addr),1888 "${usedBytes}"1889 "ld${sem:sem}${scope:scope}${addsp:addsp}.v4.${Sign:sign}$fromWidth "1890 "\t{{$dst1, $dst2, $dst3, $dst4}}, [$addr];">;1891 if support_v8 then1892 def _v8 : NVPTXInst<1893 (outs regclass:$dst1, regclass:$dst2, regclass:$dst3, regclass:$dst4,1894 regclass:$dst5, regclass:$dst6, regclass:$dst7, regclass:$dst8),1895 (ins AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp,1896 AtomicCode:$Sign, i32imm:$fromWidth, UsedBytesMask:$usedBytes,1897 ADDR:$addr),1898 "${usedBytes}"1899 "ld${sem:sem}${scope:scope}${addsp:addsp}.v8.${Sign:sign}$fromWidth "1900 "\t{{$dst1, $dst2, $dst3, $dst4, $dst5, $dst6, $dst7, $dst8}}, "1901 "[$addr];">;1902}1903let mayLoad=1, hasSideEffects=0 in {1904 defm LDV_i16 : LD_VEC<B16>;1905 defm LDV_i32 : LD_VEC<B32, support_v8 = true>;1906 defm LDV_i64 : LD_VEC<B64>;1907}1908 1909multiclass ST_VEC<DAGOperand O, bit support_v8 = false> {1910 def _v2 : NVPTXInst<1911 (outs),1912 (ins O:$src1, O:$src2,1913 AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,1914 ADDR:$addr),1915 "st${sem:sem}${scope:scope}${addsp:addsp}.v2.b$fromWidth "1916 "\t[$addr], {{$src1, $src2}};">;1917 def _v4 : NVPTXInst<1918 (outs),1919 (ins RegOrSink:$src1, RegOrSink:$src2, RegOrSink:$src3, RegOrSink:$src4,1920 AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,1921 ADDR:$addr),1922 "st${sem:sem}${scope:scope}${addsp:addsp}.v4.b$fromWidth "1923 "\t[$addr], {{$src1, $src2, $src3, $src4}};">;1924 if support_v8 then1925 def _v8 : NVPTXInst<1926 (outs),1927 (ins RegOrSink:$src1, RegOrSink:$src2, RegOrSink:$src3, RegOrSink:$src4,1928 RegOrSink:$src5, RegOrSink:$src6, RegOrSink:$src7, RegOrSink:$src8,1929 AtomicCode:$sem, AtomicCode:$scope, AtomicCode:$addsp, i32imm:$fromWidth,1930 ADDR:$addr),1931 "st${sem:sem}${scope:scope}${addsp:addsp}.v8.b$fromWidth "1932 "\t[$addr], "1933 "{{$src1, $src2, $src3, $src4, $src5, $src6, $src7, $src8}};">;1934}1935 1936let mayStore=1, hasSideEffects=0 in {1937 defm STV_i16 : ST_VEC<RI16>;1938 defm STV_i32 : ST_VEC<RI32, support_v8 = true>;1939 defm STV_i64 : ST_VEC<RI64>;1940}1941 1942//---- Conversion ----1943 1944foreach rc = [B16, B32, B64] in1945 foreach ta = rc.RegTypes in1946 foreach tb = rc.RegTypes in1947 if !ne(ta, tb) then1948 def : Pat<(ta (bitconvert tb:$a)),1949 (ta rc:$a)>;1950 1951// NOTE: pred->fp are currently sub-optimal due to an issue in TableGen where1952// we cannot specify floating-point literals in isel patterns. Therefore, we1953// use an integer selp to select either 1 (or -1 in case of signed) or 01954// and then cvt to floating-point.1955 1956// sint -> f161957def : Pat<(f16 (sint_to_fp i1:$a)), (CVT_f16_s32 (SELP_b32ii -1, 0, $a), CvtRN)>;1958def : Pat<(f16 (sint_to_fp i16:$a)), (CVT_f16_s16 $a, CvtRN)>;1959def : Pat<(f16 (sint_to_fp i32:$a)), (CVT_f16_s32 $a, CvtRN)>;1960def : Pat<(f16 (sint_to_fp i64:$a)), (CVT_f16_s64 $a, CvtRN)>;1961 1962// uint -> f161963def : Pat<(f16 (uint_to_fp i1:$a)), (CVT_f16_u32 (SELP_b32ii 1, 0, $a), CvtRN)>;1964def : Pat<(f16 (uint_to_fp i16:$a)), (CVT_f16_u16 $a, CvtRN)>;1965def : Pat<(f16 (uint_to_fp i32:$a)), (CVT_f16_u32 $a, CvtRN)>;1966def : Pat<(f16 (uint_to_fp i64:$a)), (CVT_f16_u64 $a, CvtRN)>;1967 1968// sint -> bf161969let Predicates = [hasPTX<78>, hasSM<90>] in {1970 def : Pat<(bf16 (sint_to_fp i1:$a)), (CVT_bf16_s32 (SELP_b32ii 1, 0, $a), CvtRN)>;1971 def : Pat<(bf16 (sint_to_fp i16:$a)), (CVT_bf16_s16 $a, CvtRN)>;1972 def : Pat<(bf16 (sint_to_fp i32:$a)), (CVT_bf16_s32 $a, CvtRN)>;1973 def : Pat<(bf16 (sint_to_fp i64:$a)), (CVT_bf16_s64 $a, CvtRN)>;1974}1975 1976// uint -> bf161977let Predicates = [hasPTX<78>, hasSM<90>] in {1978 def : Pat<(bf16 (uint_to_fp i1:$a)), (CVT_bf16_u32 (SELP_b32ii 1, 0, $a), CvtRN)>;1979 def : Pat<(bf16 (uint_to_fp i16:$a)), (CVT_bf16_u16 $a, CvtRN)>;1980 def : Pat<(bf16 (uint_to_fp i32:$a)), (CVT_bf16_u32 $a, CvtRN)>;1981 def : Pat<(bf16 (uint_to_fp i64:$a)), (CVT_bf16_u64 $a, CvtRN)>;1982}1983 1984// sint -> f321985def : Pat<(f32 (sint_to_fp i1:$a)), (CVT_f32_s32 (SELP_b32ii -1, 0, $a), CvtRN)>;1986def : Pat<(f32 (sint_to_fp i16:$a)), (CVT_f32_s16 $a, CvtRN)>;1987def : Pat<(f32 (sint_to_fp i32:$a)), (CVT_f32_s32 $a, CvtRN)>;1988def : Pat<(f32 (sint_to_fp i64:$a)), (CVT_f32_s64 $a, CvtRN)>;1989 1990// uint -> f321991def : Pat<(f32 (uint_to_fp i1:$a)), (CVT_f32_u32 (SELP_b32ii 1, 0, $a), CvtRN)>;1992def : Pat<(f32 (uint_to_fp i16:$a)), (CVT_f32_u16 $a, CvtRN)>;1993def : Pat<(f32 (uint_to_fp i32:$a)), (CVT_f32_u32 $a, CvtRN)>;1994def : Pat<(f32 (uint_to_fp i64:$a)), (CVT_f32_u64 $a, CvtRN)>;1995 1996// sint -> f641997def : Pat<(f64 (sint_to_fp i1:$a)), (CVT_f64_s32 (SELP_b32ii -1, 0, $a), CvtRN)>;1998def : Pat<(f64 (sint_to_fp i16:$a)), (CVT_f64_s16 $a, CvtRN)>;1999def : Pat<(f64 (sint_to_fp i32:$a)), (CVT_f64_s32 $a, CvtRN)>;2000def : Pat<(f64 (sint_to_fp i64:$a)), (CVT_f64_s64 $a, CvtRN)>;2001 2002// uint -> f642003def : Pat<(f64 (uint_to_fp i1:$a)), (CVT_f64_u32 (SELP_b32ii 1, 0, $a), CvtRN)>;2004def : Pat<(f64 (uint_to_fp i16:$a)), (CVT_f64_u16 $a, CvtRN)>;2005def : Pat<(f64 (uint_to_fp i32:$a)), (CVT_f64_u32 $a, CvtRN)>;2006def : Pat<(f64 (uint_to_fp i64:$a)), (CVT_f64_u64 $a, CvtRN)>;2007 2008 2009// f16 -> sint2010def : Pat<(i1 (fp_to_sint f16:$a)), (SETP_i16ri $a, 0, CmpEQ)>;2011def : Pat<(i16 (fp_to_sint f16:$a)), (CVT_s16_f16 $a, CvtRZI)>;2012def : Pat<(i32 (fp_to_sint f16:$a)), (CVT_s32_f16 $a, CvtRZI)>;2013def : Pat<(i64 (fp_to_sint f16:$a)), (CVT_s64_f16 $a, CvtRZI)>;2014 2015// f16 -> uint2016def : Pat<(i1 (fp_to_uint f16:$a)), (SETP_i16ri $a, 0, CmpEQ)>;2017def : Pat<(i16 (fp_to_uint f16:$a)), (CVT_u16_f16 $a, CvtRZI)>;2018def : Pat<(i32 (fp_to_uint f16:$a)), (CVT_u32_f16 $a, CvtRZI)>;2019def : Pat<(i64 (fp_to_uint f16:$a)), (CVT_u64_f16 $a, CvtRZI)>;2020 2021// bf16 -> sint2022def : Pat<(i1 (fp_to_sint bf16:$a)), (SETP_i16ri $a, 0, CmpEQ)>;2023def : Pat<(i16 (fp_to_sint bf16:$a)), (CVT_s16_bf16 $a, CvtRZI)>;2024def : Pat<(i32 (fp_to_sint bf16:$a)), (CVT_s32_bf16 $a, CvtRZI)>;2025def : Pat<(i64 (fp_to_sint bf16:$a)), (CVT_s64_bf16 $a, CvtRZI)>;2026 2027// bf16 -> uint2028def : Pat<(i1 (fp_to_uint bf16:$a)), (SETP_i16ri $a, 0, CmpEQ)>;2029def : Pat<(i16 (fp_to_uint bf16:$a)), (CVT_u16_bf16 $a, CvtRZI)>;2030def : Pat<(i32 (fp_to_uint bf16:$a)), (CVT_u32_bf16 $a, CvtRZI)>;2031def : Pat<(i64 (fp_to_uint bf16:$a)), (CVT_u64_bf16 $a, CvtRZI)>;2032// f32 -> sint2033let Predicates = [doF32FTZ] in {2034 def : Pat<(i16 (fp_to_sint f32:$a)), (CVT_s16_f32 $a, CvtRZI_FTZ)>;2035 def : Pat<(i32 (fp_to_sint f32:$a)), (CVT_s32_f32 $a, CvtRZI_FTZ)>;2036 def : Pat<(i64 (fp_to_sint f32:$a)), (CVT_s64_f32 $a, CvtRZI_FTZ)>;2037}2038def : Pat<(i1 (fp_to_sint f32:$a)), (SETP_i32ri $a, 0, CmpEQ)>;2039def : Pat<(i16 (fp_to_sint f32:$a)), (CVT_s16_f32 $a, CvtRZI)>;2040def : Pat<(i32 (fp_to_sint f32:$a)), (CVT_s32_f32 $a, CvtRZI)>;2041def : Pat<(i64 (fp_to_sint f32:$a)), (CVT_s64_f32 $a, CvtRZI)>;2042 2043// f32 -> uint2044let Predicates = [doF32FTZ] in {2045 def : Pat<(i16 (fp_to_uint f32:$a)), (CVT_u16_f32 $a, CvtRZI_FTZ)>;2046 def : Pat<(i32 (fp_to_uint f32:$a)), (CVT_u32_f32 $a, CvtRZI_FTZ)>;2047 def : Pat<(i64 (fp_to_uint f32:$a)), (CVT_u64_f32 $a, CvtRZI_FTZ)>;2048}2049def : Pat<(i1 (fp_to_uint f32:$a)), (SETP_i32ri $a, 0, CmpEQ)>;2050def : Pat<(i16 (fp_to_uint f32:$a)), (CVT_u16_f32 $a, CvtRZI)>;2051def : Pat<(i32 (fp_to_uint f32:$a)), (CVT_u32_f32 $a, CvtRZI)>;2052def : Pat<(i64 (fp_to_uint f32:$a)), (CVT_u64_f32 $a, CvtRZI)>;2053 2054// f64 -> sint2055def : Pat<(i1 (fp_to_sint f64:$a)), (SETP_i64ri $a, 0, CmpEQ)>;2056def : Pat<(i16 (fp_to_sint f64:$a)), (CVT_s16_f64 $a, CvtRZI)>;2057def : Pat<(i32 (fp_to_sint f64:$a)), (CVT_s32_f64 $a, CvtRZI)>;2058def : Pat<(i64 (fp_to_sint f64:$a)), (CVT_s64_f64 $a, CvtRZI)>;2059 2060// f64 -> uint2061def : Pat<(i1 (fp_to_uint f64:$a)), (SETP_i64ri $a, 0, CmpEQ)>;2062def : Pat<(i16 (fp_to_uint f64:$a)), (CVT_u16_f64 $a, CvtRZI)>;2063def : Pat<(i32 (fp_to_uint f64:$a)), (CVT_u32_f64 $a, CvtRZI)>;2064def : Pat<(i64 (fp_to_uint f64:$a)), (CVT_u64_f64 $a, CvtRZI)>;2065 2066// sext i12067def : Pat<(i16 (sext i1:$a)), (SELP_b16ii -1, 0, $a)>;2068def : Pat<(i32 (sext i1:$a)), (SELP_b32ii -1, 0, $a)>;2069def : Pat<(i64 (sext i1:$a)), (SELP_b64ii -1, 0, $a)>;2070 2071// zext i12072def : Pat<(i16 (zext i1:$a)), (SELP_b16ii 1, 0, $a)>;2073def : Pat<(i32 (zext i1:$a)), (SELP_b32ii 1, 0, $a)>;2074def : Pat<(i64 (zext i1:$a)), (SELP_b64ii 1, 0, $a)>;2075 2076// anyext i12077def : Pat<(i16 (anyext i1:$a)), (SELP_b16ii -1, 0, $a)>;2078def : Pat<(i32 (anyext i1:$a)), (SELP_b32ii -1, 0, $a)>;2079def : Pat<(i64 (anyext i1:$a)), (SELP_b64ii -1, 0, $a)>;2080 2081// sext i162082def : Pat<(i32 (sext i16:$a)), (CVT_s32_s16 $a, CvtNONE)>;2083def : Pat<(i64 (sext i16:$a)), (CVT_s64_s16 $a, CvtNONE)>;2084 2085// zext i162086def : Pat<(i32 (zext i16:$a)), (CVT_u32_u16 $a, CvtNONE)>;2087def : Pat<(i64 (zext i16:$a)), (CVT_u64_u16 $a, CvtNONE)>;2088 2089// anyext i162090def : Pat<(i32 (anyext i16:$a)), (CVT_u32_u16 $a, CvtNONE)>;2091def : Pat<(i64 (anyext i16:$a)), (CVT_u64_u16 $a, CvtNONE)>;2092 2093// sext i322094def : Pat<(i64 (sext i32:$a)), (CVT_s64_s32 $a, CvtNONE)>;2095 2096// zext i322097def : Pat<(i64 (zext i32:$a)), (CVT_u64_u32 $a, CvtNONE)>;2098 2099// anyext i322100def : Pat<(i64 (anyext i32:$a)), (CVT_u64_u32 $a, CvtNONE)>;2101 2102 2103// truncate i642104def : Pat<(i32 (trunc i64:$a)), (CVT_u32_u64 $a, CvtNONE)>;2105def : Pat<(i16 (trunc i64:$a)), (CVT_u16_u64 $a, CvtNONE)>;2106def : Pat<(i1 (trunc i64:$a)), (SETP_i64ri (AND_b64ri $a, 1), 0, CmpNE)>;2107 2108// truncate i322109def : Pat<(i16 (trunc i32:$a)), (CVT_u16_u32 $a, CvtNONE)>;2110def : Pat<(i1 (trunc i32:$a)), (SETP_i32ri (AND_b32ri $a, 1), 0, CmpNE)>;2111 2112// truncate i162113def : Pat<(i1 (trunc i16:$a)), (SETP_i16ri (AND_b16ri $a, 1), 0, CmpNE)>;2114 2115// sext_inreg2116def : Pat<(sext_inreg i16:$a, i8), (CVT_INREG_s16_s8 $a)>;2117def : Pat<(sext_inreg i32:$a, i8), (CVT_INREG_s32_s8 $a)>;2118def : Pat<(sext_inreg i32:$a, i16), (CVT_INREG_s32_s16 $a)>;2119def : Pat<(sext_inreg i64:$a, i8), (CVT_INREG_s64_s8 $a)>;2120def : Pat<(sext_inreg i64:$a, i16), (CVT_INREG_s64_s16 $a)>;2121def : Pat<(sext_inreg i64:$a, i32), (CVT_INREG_s64_s32 $a)>;2122 2123let hasSideEffects = false in {2124 // pack a set of smaller int registers to a larger int register2125 def V4I16toI64 : NVPTXInst<(outs B64:$d),2126 (ins B16:$s1, B16:$s2,2127 B16:$s3, B16:$s4),2128 "mov.b64 \t$d, {{$s1, $s2, $s3, $s4}};">;2129 def V2I16toI32 : NVPTXInst<(outs B32:$d),2130 (ins B16:$s1, B16:$s2),2131 "mov.b32 \t$d, {{$s1, $s2}};">;2132 def V2I32toI64 : NVPTXInst<(outs B64:$d),2133 (ins B32:$s1, B32:$s2),2134 "mov.b64 \t$d, {{$s1, $s2}};">;2135 def V2I64toI128 : NVPTXInst<(outs B128:$d),2136 (ins B64:$s1, B64:$s2),2137 "mov.b128 \t$d, {{$s1, $s2}};">;2138 2139 // unpack a larger int register to a set of smaller int registers2140 def I64toV4I16 : NVPTXInst<(outs B16:$d1, B16:$d2,2141 B16:$d3, B16:$d4),2142 (ins B64:$s),2143 "mov.b64 \t{{$d1, $d2, $d3, $d4}}, $s;">;2144 def I32toV2I16 : NVPTXInst<(outs B16:$d1, B16:$d2),2145 (ins B32:$s),2146 "mov.b32 \t{{$d1, $d2}}, $s;">;2147 def I64toV2I32 : NVPTXInst<(outs B32:$d1, B32:$d2),2148 (ins B64:$s),2149 "mov.b64 \t{{$d1, $d2}}, $s;">;2150 def I128toV2I64: NVPTXInst<(outs B64:$d1, B64:$d2),2151 (ins B128:$s),2152 "mov.b128 \t{{$d1, $d2}}, $s;">;2153 2154 def I32toI16H : NVPTXInst<(outs B16:$high), (ins B32:$s),2155 "{{ .reg .b16 tmp; mov.b32 {tmp, $high}, $s; }}">;2156 def I32toI16L : NVPTXInst<(outs B16:$low), (ins B32:$s),2157 "{{ .reg .b16 tmp; mov.b32 {$low, tmp}, $s; }}">;2158 def I64toI32H : NVPTXInst<(outs B32:$high), (ins B64:$s),2159 "{{ .reg .b32 tmp; mov.b64 {tmp, $high}, $s; }}">;2160 def I64toI32L : NVPTXInst<(outs B32:$low), (ins B64:$s),2161 "{{ .reg .b32 tmp; mov.b64 {$low, tmp}, $s; }}">;2162 2163 // PTX 7.1 lets you avoid a temp register and just use _ as a "sink" for the2164 // unused high/low part.2165 let Predicates = [hasPTX<71>] in {2166 def I32toI16H_Sink : NVPTXInst<(outs B16:$high), (ins B32:$s),2167 "mov.b32 \t{{_, $high}}, $s;">;2168 def I32toI16L_Sink : NVPTXInst<(outs B16:$low), (ins B32:$s),2169 "mov.b32 \t{{$low, _}}, $s;">;2170 def I64toI32H_Sink : NVPTXInst<(outs B32:$high), (ins B64:$s),2171 "mov.b64 \t{{_, $high}}, $s;">;2172 def I64toI32L_Sink : NVPTXInst<(outs B32:$low), (ins B64:$s),2173 "mov.b64 \t{{$low, _}}, $s;">;2174 }2175}2176 2177let Predicates = [hasPTX<71>] in {2178 def : Pat<(i16 (trunc (srl i32:$s, (i32 16)))), (I32toI16H_Sink i32:$s)>;2179 def : Pat<(i16 (trunc (sra i32:$s, (i32 16)))), (I32toI16H_Sink i32:$s)>;2180 def : Pat<(i32 (trunc (srl i64:$s, (i32 32)))), (I64toI32H_Sink i64:$s)>;2181 def : Pat<(i32 (trunc (sra i64:$s, (i32 32)))), (I64toI32H_Sink i64:$s)>;2182}2183 2184// Fall back to the old way if we don't have PTX 7.1.2185def : Pat<(i16 (trunc (srl i32:$s, (i32 16)))), (I32toI16H $s)>;2186def : Pat<(i16 (trunc (sra i32:$s, (i32 16)))), (I32toI16H $s)>;2187def : Pat<(i32 (trunc (srl i64:$s, (i32 32)))), (I64toI32H $s)>;2188def : Pat<(i32 (trunc (sra i64:$s, (i32 32)))), (I64toI32H $s)>;2189 2190def: Pat<(i32 (sext (extractelt v2i16:$src, 0))),2191 (CVT_INREG_s32_s16 $src)>;2192 2193// Handle extracting one element from the pair (32-bit types)2194foreach vt = [v2f16, v2bf16, v2i16] in {2195 def : Pat<(extractelt vt:$src, 0), (I32toI16L_Sink $src)>, Requires<[hasPTX<71>]>;2196 def : Pat<(extractelt vt:$src, 1), (I32toI16H_Sink $src)>, Requires<[hasPTX<71>]>;2197 2198 def : Pat<(extractelt vt:$src, 0), (I32toI16L $src)>;2199 def : Pat<(extractelt vt:$src, 1), (I32toI16H $src)>;2200 2201 def : Pat<(vt (build_vector vt.ElementType:$a, vt.ElementType:$b)), 2202 (V2I16toI32 $a, $b)>;2203}2204 2205// Handle extracting one element from the pair (64-bit types)2206foreach vt = [v2f32, v2i32] in {2207 def : Pat<(extractelt vt:$src, 0), (I64toI32L_Sink $src)>, Requires<[hasPTX<71>]>;2208 def : Pat<(extractelt vt:$src, 1), (I64toI32H_Sink $src)>, Requires<[hasPTX<71>]>;2209 2210 def : Pat<(extractelt vt:$src, 0), (I64toI32L $src)>;2211 def : Pat<(extractelt vt:$src, 1), (I64toI32H $src)>;2212 2213 def : Pat<(vt (build_vector vt.ElementType:$a, vt.ElementType:$b)), 2214 (V2I32toI64 $a, $b)>;2215}2216 2217def: Pat<(v2i16 (scalar_to_vector i16:$a)),2218 (CVT_u32_u16 $a, CvtNONE)>;2219 2220// This node is similar to ISD::BUILD_VECTOR except that the output may be2221// implicitly bitcast to a scalar. This allows for the representation of2222// packing move instructions for vector types which are not legal i.e. v2i322223def nvptx_build_vector : SDNode<"NVPTXISD::BUILD_VECTOR", SDTypeProfile<1, 2, []>, []>;2224 2225def : Pat<(i64 (nvptx_build_vector i32:$a, i32:$b)),2226 (V2I32toI64 $a, $b)>;2227 2228//2229// Funnel-Shift2230//2231 2232// Create SDNodes so they can be used in the DAG code, e.g.2233// NVPTXISelLowering (LowerShiftLeftParts and LowerShiftRightParts)2234def fshl_clamp : SDNode<"NVPTXISD::FSHL_CLAMP", SDTIntShiftDOp, []>;2235def fshr_clamp : SDNode<"NVPTXISD::FSHR_CLAMP", SDTIntShiftDOp, []>;2236 2237// Funnel shift, requires >= sm_32. Does not trap if amt is out of range, so2238// no side effects.2239let hasSideEffects = false in {2240 multiclass ShfInst<string mode, SDNode op> {2241 def _i2242 : BasicNVPTXInst<(outs B32:$dst),2243 (ins B32:$lo, B32:$hi, i32imm:$amt),2244 "shf." # mode # ".b32",2245 [(set i32:$dst,2246 (op i32:$hi, i32:$lo, (i32 imm:$amt)))]>,2247 Requires<[hasHWROT32]>;2248 2249 def _r2250 : BasicNVPTXInst<(outs B32:$dst),2251 (ins B32:$lo, B32:$hi, B32:$amt),2252 "shf." # mode # ".b32",2253 [(set i32:$dst,2254 (op i32:$hi, i32:$lo, i32:$amt))]>,2255 Requires<[hasHWROT32]>;2256 }2257 2258 defm SHF_L_CLAMP : ShfInst<"l.clamp", fshl_clamp>;2259 defm SHF_R_CLAMP : ShfInst<"r.clamp", fshr_clamp>;2260 defm SHF_L_WRAP : ShfInst<"l.wrap", fshl>;2261 defm SHF_R_WRAP : ShfInst<"r.wrap", fshr>;2262}2263 2264def : Pat<(i32 (int_nvvm_fshl_clamp i32:$hi, i32:$lo, i32:$amt)),2265 (SHF_L_CLAMP_r $lo, $hi, $amt)>;2266def : Pat<(i32 (int_nvvm_fshl_clamp i32:$hi, i32:$lo, (i32 imm:$amt))),2267 (SHF_L_CLAMP_i $lo, $hi, imm:$amt)>;2268def : Pat<(i32 (int_nvvm_fshr_clamp i32:$hi, i32:$lo, i32:$amt)),2269 (SHF_R_CLAMP_r $lo, $hi, $amt)>;2270def : Pat<(i32 (int_nvvm_fshr_clamp i32:$hi, i32:$lo, (i32 imm:$amt))),2271 (SHF_R_CLAMP_i $lo, $hi, imm:$amt)>;2272 2273let hasSideEffects = false in {2274 foreach RT = [I32RT, I64RT] in {2275 // Count leading zeros2276 def CLZr # RT.Size : BasicNVPTXInst<(outs B32:$d), (ins RT.RC:$a),2277 "clz.b" # RT.Size,2278 [(set i32:$d, (ctlz RT.Ty:$a))]>;2279 2280 // Population count2281 def POPCr # RT.Size : BasicNVPTXInst<(outs B32:$d), (ins RT.RC:$a),2282 "popc.b" # RT.Size,2283 [(set i32:$d, (ctpop RT.Ty:$a))]>;2284 }2285}2286 2287// fpround f32 -> f162288def : Pat<(f16 (fpround f32:$a)), (CVT_f16_f32 $a, CvtRN)>;2289 2290// fpround f32 -> bf162291def : Pat<(bf16 (fpround f32:$a)), (CVT_bf16_f32 $a, CvtRN)>, 2292 Requires<[hasPTX<70>, hasSM<80>]>;2293 2294// fpround f64 -> f162295def : Pat<(f16 (fpround f64:$a)), (CVT_f16_f64 $a, CvtRN)>;2296 2297// fpround f64 -> bf162298def : Pat<(bf16 (fpround f64:$a)), (CVT_bf16_f64 $a, CvtRN)>, 2299 Requires<[hasPTX<78>, hasSM<90>]>;2300 2301// fpround f64 -> f322302def : Pat<(f32 (fpround f64:$a)), (CVT_f32_f64 $a, CvtRN_FTZ)>, Requires<[doF32FTZ]>;2303def : Pat<(f32 (fpround f64:$a)), (CVT_f32_f64 $a, CvtRN)>;2304 2305// fpextend f16 -> f322306def : Pat<(f32 (fpextend f16:$a)), (CVT_f32_f16 $a, CvtNONE_FTZ)>, Requires<[doF32FTZ]>;2307def : Pat<(f32 (fpextend f16:$a)), (CVT_f32_f16 $a, CvtNONE)>;2308// fpextend bf16 -> f322309def : Pat<(f32 (fpextend bf16:$a)), (CVT_f32_bf16 $a, CvtNONE_FTZ)>, Requires<[doF32FTZ, hasPTX<78>, hasSM<90>]>;2310def : Pat<(f32 (fpextend bf16:$a)), (CVT_f32_bf16 $a, CvtNONE)>, Requires<[hasPTX<71>, hasSM<80>]>;2311 2312// fpextend f16 -> f642313def : Pat<(f64 (fpextend f16:$a)), (CVT_f64_f16 $a, CvtNONE)>;2314 2315// fpextend bf16 -> f642316def : Pat<(f64 (fpextend bf16:$a)), (CVT_f64_bf16 $a, CvtNONE)>, Requires<[hasPTX<78>, hasSM<90>]>;2317 2318// fpextend f32 -> f642319def : Pat<(f64 (fpextend f32:$a)), (CVT_f64_f32 $a, CvtNONE_FTZ)>, Requires<[doF32FTZ]>;2320def : Pat<(f64 (fpextend f32:$a)), (CVT_f64_f32 $a, CvtNONE)>;2321 2322def retglue : SDNode<"NVPTXISD::RET_GLUE", SDTNone,2323 [SDNPHasChain, SDNPOptInGlue]>;2324 2325// fceil, ffloor, froundeven, ftrunc.2326 2327multiclass CVT_ROUND<SDNode OpNode, PatLeaf Mode, PatLeaf ModeFTZ> {2328 def : Pat<(OpNode f16:$a), (CVT_f16_f16 $a, Mode)>;2329 def : Pat<(OpNode bf16:$a), (CVT_bf16_bf16 $a, Mode)>;2330 def : Pat<(OpNode f32:$a), (CVT_f32_f32 $a, ModeFTZ)>, Requires<[doF32FTZ]>;2331 def : Pat<(OpNode f32:$a), (CVT_f32_f32 $a, Mode)>, Requires<[doNoF32FTZ]>;2332 def : Pat<(OpNode f64:$a), (CVT_f64_f64 $a, Mode)>;2333}2334 2335defm : CVT_ROUND<fceil, CvtRPI, CvtRPI_FTZ>;2336defm : CVT_ROUND<ffloor, CvtRMI, CvtRMI_FTZ>;2337defm : CVT_ROUND<froundeven, CvtRNI, CvtRNI_FTZ>;2338defm : CVT_ROUND<ftrunc, CvtRZI, CvtRZI_FTZ>;2339 2340// nearbyint and rint are implemented as rounding to nearest even. This isn't2341// strictly correct, because it causes us to ignore the rounding mode. But it2342// matches what CUDA's "libm" does.2343 2344defm : CVT_ROUND<fnearbyint, CvtRNI, CvtRNI_FTZ>;2345defm : CVT_ROUND<frint, CvtRNI, CvtRNI_FTZ>;2346 2347//-----------------------------------2348// Control-flow2349//-----------------------------------2350 2351let isTerminator=1 in {2352 let isReturn=1, isBarrier=1 in2353 def Return : BasicNVPTXInst<(outs), (ins), "ret", [(retglue)]>;2354 2355 let isBranch=1 in {2356 def CBranch : NVPTXInst<(outs), (ins B1:$a, brtarget:$target),2357 "@$a bra \t$target;",2358 [(brcond i1:$a, bb:$target)]>;2359 2360 let isBarrier=1 in2361 def GOTO : BasicNVPTXInst<(outs), (ins brtarget:$target),2362 "bra.uni", [(br bb:$target)]>;2363 }2364}2365 2366 2367// trap instruction2368def trapinst : BasicNVPTXInst<(outs), (ins), "trap", [(trap)]>, Requires<[noPTXASUnreachableBug]>;2369// Emit an `exit` as well to convey to ptxas that `trap` exits the CFG.2370// This won't be necessary in a future version of ptxas.2371def trapexitinst : NVPTXInst<(outs), (ins), "trap; exit;", [(trap)]>, Requires<[hasPTXASUnreachableBug]>;2372// brkpt instruction2373def debugtrapinst : BasicNVPTXInst<(outs), (ins), "brkpt", [(debugtrap)]>;2374 2375def SDTDynAllocaOp :2376 SDTypeProfile<1, 2, [SDTCisSameAs<0, 1>, SDTCisInt<1>, SDTCisVT<2, i32>]>;2377 2378def dyn_alloca :2379 SDNode<"NVPTXISD::DYNAMIC_STACKALLOC", SDTDynAllocaOp,2380 [SDNPHasChain, SDNPSideEffect]>;2381 2382foreach t = [I32RT, I64RT] in {2383 def DYNAMIC_STACKALLOC # t.Size :2384 BasicNVPTXInst<(outs t.RC:$ptr),2385 (ins t.RC:$size, i32imm:$align),2386 "alloca.u" # t.Size,2387 [(set t.Ty:$ptr, (dyn_alloca t.Ty:$size, timm:$align))]>,2388 Requires<[hasPTX<73>, hasSM<52>]>;2389}2390 2391//2392// BRX2393//2394 2395def SDTBrxStartProfile : SDTypeProfile<0, 1, [SDTCisInt<0>]>;2396def SDTBrxItemProfile : SDTypeProfile<0, 1, [SDTCisVT<0, OtherVT>]>;2397def SDTBrxEndProfile : SDTypeProfile<0, 3, [SDTCisVT<0, OtherVT>, SDTCisInt<1>, SDTCisInt<2>]>;2398 2399def brx_start :2400 SDNode<"NVPTXISD::BrxStart", SDTBrxStartProfile,2401 [SDNPHasChain, SDNPOutGlue, SDNPSideEffect]>;2402def brx_item :2403 SDNode<"NVPTXISD::BrxItem", SDTBrxItemProfile,2404 [SDNPHasChain, SDNPOutGlue, SDNPInGlue, SDNPSideEffect]>;2405def brx_end :2406 SDNode<"NVPTXISD::BrxEnd", SDTBrxEndProfile,2407 [SDNPHasChain, SDNPInGlue, SDNPSideEffect]>;2408 2409let isTerminator = 1, isBranch = 1, isIndirectBranch = 1, isNotDuplicable = 1 in {2410 2411 def BRX_START :2412 NVPTXInst<(outs), (ins i32imm:$id),2413 "$$L_brx_$id: .branchtargets",2414 [(brx_start (i32 imm:$id))]>;2415 2416 def BRX_ITEM :2417 NVPTXInst<(outs), (ins brtarget:$target),2418 "\t$target,",2419 [(brx_item bb:$target)]>;2420 2421 def BRX_END :2422 NVPTXInst<(outs), (ins brtarget:$target, B32:$val, i32imm:$id),2423 "\t$target;\n\tbrx.idx \t$val, $$L_brx_$id;",2424 [(brx_end bb:$target, i32:$val, (i32 imm:$id))]> {2425 let isBarrier = 1;2426 }2427}2428 2429 2430foreach a_type = ["s", "u"] in {2431 foreach b_type = ["s", "u"] in {2432 2433 def DOT4_ # a_type # b_type :2434 BasicNVPTXInst<(outs B32:$dst),2435 (ins B32:$a, B32:$b, B32:$c),2436 "dp4a." # a_type # "32." # b_type # "32",2437 [(set i32:$dst,2438 (!cast<Intrinsic>("int_nvvm_idp4a_" # a_type # "_" # b_type)2439 i32:$a, i32:$b, i32:$c))]>,2440 Requires<[hasDotInstructions]>;2441 2442 foreach is_hi = [0, -1] in {2443 defvar lohi_suffix = !if(is_hi, "hi", "lo");2444 2445 def DOT2_ # lohi_suffix # _ # a_type # b_type :2446 BasicNVPTXInst<(outs B32:$dst),2447 (ins B32:$a, B32:$b, B32:$c),2448 "dp2a." # lohi_suffix # "." # a_type # "32." # b_type # "32",2449 [(set i32:$dst,2450 (!cast<Intrinsic>("int_nvvm_idp2a_" # a_type # "_" # b_type)2451 i32:$a, i32:$b, is_hi, i32:$c))]>,2452 Requires<[hasDotInstructions]>;2453 }2454 }2455}2456 2457//2458// Stack Manipulation2459//2460 2461def SDTStackRestore : SDTypeProfile<0, 1, [SDTCisInt<0>]>;2462 2463def stackrestore :2464 SDNode<"NVPTXISD::STACKRESTORE", SDTStackRestore,2465 [SDNPHasChain, SDNPSideEffect]>;2466 2467def stacksave :2468 SDNode<"NVPTXISD::STACKSAVE", SDTIntLeaf,2469 [SDNPHasChain, SDNPSideEffect]>;2470 2471let Predicates = [hasPTX<73>, hasSM<52>] in {2472 foreach t = [I32RT, I64RT] in {2473 def STACKRESTORE_ # t.Size :2474 BasicNVPTXInst<(outs), (ins t.RC:$ptr),2475 "stackrestore.u" # t.Size,2476 [(stackrestore t.Ty:$ptr)]>;2477 2478 def STACKSAVE_ # t.Size :2479 BasicNVPTXInst<(outs t.RC:$dst), (ins),2480 "stacksave.u" # t.Size,2481 [(set t.Ty:$dst, (t.Ty stacksave))]>;2482 }2483}2484 2485include "NVPTXIntrinsics.td"2486 2487////////////////////////////////////////////////////////////////////////////////2488// PTX Fence instructions2489////////////////////////////////////////////////////////////////////////////////2490 2491class NVPTXFenceInst<string scope, string sem, Predicate ptx>:2492 BasicNVPTXInst<(outs), (ins), "fence."#sem#"."#scope>,2493 Requires<[ptx, hasSM<70>]>;2494 2495foreach scope = ["sys", "gpu", "cluster", "cta"] in {2496 def atomic_thread_fence_seq_cst_#scope: NVPTXFenceInst<scope, "sc", hasPTX<60>>;2497 def atomic_thread_fence_acq_rel_#scope: NVPTXFenceInst<scope, "acq_rel", hasPTX<60>>;2498 def atomic_thread_fence_acquire_#scope: NVPTXFenceInst<scope, "acquire", hasPTX<87>>;2499 def atomic_thread_fence_release_#scope: NVPTXFenceInst<scope, "release", hasPTX<87>>;2500}2501 2502// Perform substitution if fma only has one use, and also if instruction has2503// nnan instruction flag or if the TM has NoNaNsFPMath2504def NVPTX_fma_oneuse_and_nnan : PatFrag<(ops node:$a, node:$b, node:$c),2505 (fma node:$a, node:$b, node:$c), [{2506 return N->hasOneUse() &&2507 (N->getFlags().hasNoNaNs() || TM.Options.NoNaNsFPMath);2508}]>;2509// fmaxnum will differentiate between signed and unsigned zeros soon, so this2510// PatFrag is for a fmaxnum node with nsz2511def NVPTX_fmaxnum_nsz : PatFrag<(ops node:$a, node:$b),2512 (fmaxnum node:$a, node:$b), [{2513 return N->getFlags().hasNoSignedZeros() || TM.Options.NoSignedZerosFPMath;2514}]>;2515 2516class FMARELUInst<RegTyInfo t, bit allow_ftz, PatFrag zero_pat>2517 : BasicFlagsNVPTXInst<(outs t.RC:$dst), (ins t.RC:$a, t.RC:$b, t.RC:$c),2518 !if(allow_ftz, (ins FTZFlag:$ftz), (ins)),2519 "fma.rn" # !if(allow_ftz, "$ftz", "") # ".relu." # t.PtxType,2520 [(set t.Ty:$dst, (NVPTX_fmaxnum_nsz (NVPTX_fma_oneuse_and_nnan t.Ty:$a, t.Ty:$b, t.Ty:$c), zero_pat))]>;2521 2522let Predicates = [useFP16Math, hasPTX<70>, hasSM<80>] in {2523 def FMARELU_F16 : FMARELUInst<F16RT, true, fpimm_0>;2524 def FMARELU_F16X2 : FMARELUInst<F16X2RT, true, zeroinitializer<v2f16>>;2525}2526 2527let Predicates = [hasBF16Math, hasPTX<70>, hasSM<80>] in {2528 def FMARELU_BF16 : FMARELUInst<BF16RT, false, fpimm_0>;2529 def FMARELU_BF16X2 : FMARELUInst<BF16X2RT, false, zeroinitializer<v2bf16>>;2530}2531