LLVM  16.0.0git
NVPTXISelLowering.cpp
Go to the documentation of this file.
1 //===-- NVPTXISelLowering.cpp - NVPTX DAG Lowering Implementation ---------===//
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-exception
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This file defines the interfaces that NVPTX uses to lower LLVM code into a
10 // selection DAG.
11 //
12 //===----------------------------------------------------------------------===//
13 
14 #include "NVPTXISelLowering.h"
16 #include "NVPTX.h"
17 #include "NVPTXSubtarget.h"
18 #include "NVPTXTargetMachine.h"
19 #include "NVPTXTargetObjectFile.h"
20 #include "NVPTXUtilities.h"
21 #include "llvm/ADT/APInt.h"
22 #include "llvm/ADT/STLExtras.h"
23 #include "llvm/ADT/SmallVector.h"
24 #include "llvm/ADT/StringRef.h"
25 #include "llvm/CodeGen/Analysis.h"
33 #include "llvm/IR/Argument.h"
34 #include "llvm/IR/Attributes.h"
35 #include "llvm/IR/Constants.h"
36 #include "llvm/IR/DataLayout.h"
37 #include "llvm/IR/DerivedTypes.h"
38 #include "llvm/IR/FPEnv.h"
39 #include "llvm/IR/Function.h"
40 #include "llvm/IR/GlobalValue.h"
41 #include "llvm/IR/Instruction.h"
42 #include "llvm/IR/Instructions.h"
43 #include "llvm/IR/IntrinsicsNVPTX.h"
44 #include "llvm/IR/Module.h"
45 #include "llvm/IR/Type.h"
46 #include "llvm/IR/Value.h"
47 #include "llvm/Support/Casting.h"
48 #include "llvm/Support/CodeGen.h"
55 #include <algorithm>
56 #include <cassert>
57 #include <cmath>
58 #include <cstdint>
59 #include <iterator>
60 #include <sstream>
61 #include <string>
62 #include <utility>
63 #include <vector>
64 
65 #define DEBUG_TYPE "nvptx-lower"
66 
67 using namespace llvm;
68 
69 static std::atomic<unsigned> GlobalUniqueCallSite;
70 
72  "nvptx-sched4reg",
73  cl::desc("NVPTX Specific: schedule for register pressue"), cl::init(false));
74 
76  "nvptx-fma-level", cl::Hidden,
77  cl::desc("NVPTX Specific: FMA contraction (0: don't do it"
78  " 1: do it 2: do it aggressively"),
79  cl::init(2));
80 
82  "nvptx-prec-divf32", cl::Hidden,
83  cl::desc("NVPTX Specifies: 0 use div.approx, 1 use div.full, 2 use"
84  " IEEE Compliant F32 div.rnd if available."),
85  cl::init(2));
86 
88  "nvptx-prec-sqrtf32", cl::Hidden,
89  cl::desc("NVPTX Specific: 0 use sqrt.approx, 1 use sqrt.rn."),
90  cl::init(true));
91 
93  if (UsePrecDivF32.getNumOccurrences() > 0) {
94  // If nvptx-prec-div32=N is used on the command-line, always honor it
95  return UsePrecDivF32;
96  } else {
97  // Otherwise, use div.approx if fast math is enabled
98  if (getTargetMachine().Options.UnsafeFPMath)
99  return 0;
100  else
101  return 2;
102  }
103 }
104 
106  if (UsePrecSqrtF32.getNumOccurrences() > 0) {
107  // If nvptx-prec-sqrtf32 is used on the command-line, always honor it
108  return UsePrecSqrtF32;
109  } else {
110  // Otherwise, use sqrt.approx if fast math is enabled
112  }
113 }
114 
118 }
119 
120 static bool IsPTXVectorType(MVT VT) {
121  switch (VT.SimpleTy) {
122  default:
123  return false;
124  case MVT::v2i1:
125  case MVT::v4i1:
126  case MVT::v2i8:
127  case MVT::v4i8:
128  case MVT::v2i16:
129  case MVT::v4i16:
130  case MVT::v2i32:
131  case MVT::v4i32:
132  case MVT::v2i64:
133  case MVT::v2f16:
134  case MVT::v4f16:
135  case MVT::v8f16: // <4 x f16x2>
136  case MVT::v2bf16:
137  case MVT::v4bf16:
138  case MVT::v8bf16: // <4 x bf16x2>
139  case MVT::v2f32:
140  case MVT::v4f32:
141  case MVT::v2f64:
142  return true;
143  }
144 }
145 
146 /// ComputePTXValueVTs - For the given Type \p Ty, returns the set of primitive
147 /// EVTs that compose it. Unlike ComputeValueVTs, this will break apart vectors
148 /// into their primitive components.
149 /// NOTE: This is a band-aid for code that expects ComputeValueVTs to return the
150 /// same number of types as the Ins/Outs arrays in LowerFormalArguments,
151 /// LowerCall, and LowerReturn.
152 static void ComputePTXValueVTs(const TargetLowering &TLI, const DataLayout &DL,
153  Type *Ty, SmallVectorImpl<EVT> &ValueVTs,
155  uint64_t StartingOffset = 0) {
156  SmallVector<EVT, 16> TempVTs;
157  SmallVector<uint64_t, 16> TempOffsets;
158 
159  // Special case for i128 - decompose to (i64, i64)
160  if (Ty->isIntegerTy(128)) {
161  ValueVTs.push_back(EVT(MVT::i64));
162  ValueVTs.push_back(EVT(MVT::i64));
163 
164  if (Offsets) {
165  Offsets->push_back(StartingOffset + 0);
166  Offsets->push_back(StartingOffset + 8);
167  }
168 
169  return;
170  }
171 
172  // Given a struct type, recursively traverse the elements with custom ComputePTXValueVTs.
173  if (StructType *STy = dyn_cast<StructType>(Ty)) {
174  auto const *SL = DL.getStructLayout(STy);
175  auto ElementNum = 0;
176  for(auto *EI : STy->elements()) {
177  ComputePTXValueVTs(TLI, DL, EI, ValueVTs, Offsets,
178  StartingOffset + SL->getElementOffset(ElementNum));
179  ++ElementNum;
180  }
181  return;
182  }
183 
184  ComputeValueVTs(TLI, DL, Ty, TempVTs, &TempOffsets, StartingOffset);
185  for (unsigned i = 0, e = TempVTs.size(); i != e; ++i) {
186  EVT VT = TempVTs[i];
187  uint64_t Off = TempOffsets[i];
188  // Split vectors into individual elements, except for v2f16, which
189  // we will pass as a single scalar.
190  if (VT.isVector()) {
191  unsigned NumElts = VT.getVectorNumElements();
192  EVT EltVT = VT.getVectorElementType();
193  // Vectors with an even number of f16 elements will be passed to
194  // us as an array of v2f16 elements. We must match this so we
195  // stay in sync with Ins/Outs.
196  if ((EltVT == MVT::f16 || EltVT == MVT::f16) && NumElts % 2 == 0) {
197  EltVT = EltVT == MVT::f16 ? MVT::v2f16 : MVT::v2bf16;
198  NumElts /= 2;
199  }
200  for (unsigned j = 0; j != NumElts; ++j) {
201  ValueVTs.push_back(EltVT);
202  if (Offsets)
203  Offsets->push_back(Off + j * EltVT.getStoreSize());
204  }
205  } else {
206  ValueVTs.push_back(VT);
207  if (Offsets)
208  Offsets->push_back(Off);
209  }
210  }
211 }
212 
213 /// PromoteScalarIntegerPTX
214 /// Used to make sure the arguments/returns are suitable for passing
215 /// and promote them to a larger size if they're not.
216 ///
217 /// The promoted type is placed in \p PromoteVT if the function returns true.
218 static bool PromoteScalarIntegerPTX(const EVT &VT, MVT *PromotedVT) {
219  if (VT.isScalarInteger()) {
220  switch (PowerOf2Ceil(VT.getFixedSizeInBits())) {
221  default:
223  "Promotion is not suitable for scalars of size larger than 64-bits");
224  case 1:
225  *PromotedVT = MVT::i1;
226  break;
227  case 2:
228  case 4:
229  case 8:
230  *PromotedVT = MVT::i8;
231  break;
232  case 16:
233  *PromotedVT = MVT::i16;
234  break;
235  case 32:
236  *PromotedVT = MVT::i32;
237  break;
238  case 64:
239  *PromotedVT = MVT::i64;
240  break;
241  }
242  return EVT(*PromotedVT) != VT;
243  }
244  return false;
245 }
246 
247 // Check whether we can merge loads/stores of some of the pieces of a
248 // flattened function parameter or return value into a single vector
249 // load/store.
250 //
251 // The flattened parameter is represented as a list of EVTs and
252 // offsets, and the whole structure is aligned to ParamAlignment. This
253 // function determines whether we can load/store pieces of the
254 // parameter starting at index Idx using a single vectorized op of
255 // size AccessSize. If so, it returns the number of param pieces
256 // covered by the vector op. Otherwise, it returns 1.
258  unsigned Idx, uint32_t AccessSize, const SmallVectorImpl<EVT> &ValueVTs,
259  const SmallVectorImpl<uint64_t> &Offsets, Align ParamAlignment) {
260 
261  // Can't vectorize if param alignment is not sufficient.
262  if (ParamAlignment < AccessSize)
263  return 1;
264  // Can't vectorize if offset is not aligned.
265  if (Offsets[Idx] & (AccessSize - 1))
266  return 1;
267 
268  EVT EltVT = ValueVTs[Idx];
269  unsigned EltSize = EltVT.getStoreSize();
270 
271  // Element is too large to vectorize.
272  if (EltSize >= AccessSize)
273  return 1;
274 
275  unsigned NumElts = AccessSize / EltSize;
276  // Can't vectorize if AccessBytes if not a multiple of EltSize.
277  if (AccessSize != EltSize * NumElts)
278  return 1;
279 
280  // We don't have enough elements to vectorize.
281  if (Idx + NumElts > ValueVTs.size())
282  return 1;
283 
284  // PTX ISA can only deal with 2- and 4-element vector ops.
285  if (NumElts != 4 && NumElts != 2)
286  return 1;
287 
288  for (unsigned j = Idx + 1; j < Idx + NumElts; ++j) {
289  // Types do not match.
290  if (ValueVTs[j] != EltVT)
291  return 1;
292 
293  // Elements are not contiguous.
294  if (Offsets[j] - Offsets[j - 1] != EltSize)
295  return 1;
296  }
297  // OK. We can vectorize ValueVTs[i..i+NumElts)
298  return NumElts;
299 }
300 
301 // Flags for tracking per-element vectorization state of loads/stores
302 // of a flattened function parameter or return value.
304  PVF_INNER = 0x0, // Middle elements of a vector.
305  PVF_FIRST = 0x1, // First element of the vector.
306  PVF_LAST = 0x2, // Last element of the vector.
307  // Scalar is effectively a 1-element vector.
309 };
310 
311 // Computes whether and how we can vectorize the loads/stores of a
312 // flattened function parameter or return value.
313 //
314 // The flattened parameter is represented as the list of ValueVTs and
315 // Offsets, and is aligned to ParamAlignment bytes. We return a vector
316 // of the same size as ValueVTs indicating how each piece should be
317 // loaded/stored (i.e. as a scalar, or as part of a vector
318 // load/store).
322  Align ParamAlignment) {
323  // Set vector size to match ValueVTs and mark all elements as
324  // scalars by default.
326  VectorInfo.assign(ValueVTs.size(), PVF_SCALAR);
327 
328  // Check what we can vectorize using 128/64/32-bit accesses.
329  for (int I = 0, E = ValueVTs.size(); I != E; ++I) {
330  // Skip elements we've already processed.
331  assert(VectorInfo[I] == PVF_SCALAR && "Unexpected vector info state.");
332  for (unsigned AccessSize : {16, 8, 4, 2}) {
333  unsigned NumElts = CanMergeParamLoadStoresStartingAt(
334  I, AccessSize, ValueVTs, Offsets, ParamAlignment);
335  // Mark vectorized elements.
336  switch (NumElts) {
337  default:
338  llvm_unreachable("Unexpected return value");
339  case 1:
340  // Can't vectorize using this size, try next smaller size.
341  continue;
342  case 2:
343  assert(I + 1 < E && "Not enough elements.");
344  VectorInfo[I] = PVF_FIRST;
345  VectorInfo[I + 1] = PVF_LAST;
346  I += 1;
347  break;
348  case 4:
349  assert(I + 3 < E && "Not enough elements.");
350  VectorInfo[I] = PVF_FIRST;
351  VectorInfo[I + 1] = PVF_INNER;
352  VectorInfo[I + 2] = PVF_INNER;
353  VectorInfo[I + 3] = PVF_LAST;
354  I += 3;
355  break;
356  }
357  // Break out of the inner loop because we've already succeeded
358  // using largest possible AccessSize.
359  break;
360  }
361  }
362  return VectorInfo;
363 }
364 
365 // NVPTXTargetLowering Constructor.
367  const NVPTXSubtarget &STI)
368  : TargetLowering(TM), nvTM(&TM), STI(STI) {
369  // always lower memset, memcpy, and memmove intrinsics to load/store
370  // instructions, rather
371  // then generating calls to memset, mempcy or memmove.
372  MaxStoresPerMemset = (unsigned) 0xFFFFFFFF;
373  MaxStoresPerMemcpy = (unsigned) 0xFFFFFFFF;
374  MaxStoresPerMemmove = (unsigned) 0xFFFFFFFF;
375 
378 
379  // Jump is Expensive. Don't create extra control flow for 'and', 'or'
380  // condition branches.
381  setJumpIsExpensive(true);
382 
383  // Wide divides are _very_ slow. Try to reduce the width of the divide if
384  // possible.
385  addBypassSlowDiv(64, 32);
386 
387  // By default, use the Source scheduling
388  if (sched4reg)
390  else
392 
393  auto setFP16OperationAction = [&](unsigned Op, MVT VT, LegalizeAction Action,
394  LegalizeAction NoF16Action) {
395  setOperationAction(Op, VT, STI.allowFP16Math() ? Action : NoF16Action);
396  };
397 
398  addRegisterClass(MVT::i1, &NVPTX::Int1RegsRegClass);
399  addRegisterClass(MVT::i16, &NVPTX::Int16RegsRegClass);
400  addRegisterClass(MVT::i32, &NVPTX::Int32RegsRegClass);
401  addRegisterClass(MVT::i64, &NVPTX::Int64RegsRegClass);
402  addRegisterClass(MVT::f32, &NVPTX::Float32RegsRegClass);
403  addRegisterClass(MVT::f64, &NVPTX::Float64RegsRegClass);
404  addRegisterClass(MVT::f16, &NVPTX::Float16RegsRegClass);
405  addRegisterClass(MVT::v2f16, &NVPTX::Float16x2RegsRegClass);
406  addRegisterClass(MVT::bf16, &NVPTX::Float16RegsRegClass);
407  addRegisterClass(MVT::v2bf16, &NVPTX::Float16x2RegsRegClass);
408 
409  // Conversion to/from FP16/FP16x2 is always legal.
416 
417  setFP16OperationAction(ISD::SETCC, MVT::f16, Legal, Promote);
418  setFP16OperationAction(ISD::SETCC, MVT::v2f16, Legal, Expand);
419 
420  // Operations not directly supported by NVPTX.
425  }
426 
427  // Some SIGN_EXTEND_INREG can be done using cvt instruction.
428  // For others we will expand to a SHL/SRA pair.
434 
441 
444 
445  // TODO: we may consider expanding ROTL/ROTR on older GPUs. Currently on GPUs
446  // that don't have h/w rotation we lower them to multi-instruction assembly.
447  // See ROT*_sw in NVPTXIntrInfo.td
452 
460 
461  // Indirect branch is not supported.
462  // This also disables Jump Table creation.
465 
468 
469  // We want to legalize constant related memmove and memcopy
470  // intrinsics.
472 
473  // Turn FP extload into load/fpextend
483  // Turn FP truncstore into trunc + store.
484  // FIXME: vector types should also be expanded
488 
489  // PTX does not support load / store predicate registers
492 
493  for (MVT VT : MVT::integer_valuetypes()) {
497  }
498 
499  // This is legal in NVPTX
504 
505  // TRAP can be lowered to PTX trap
507 
508  // Register custom handling for vector loads/stores
509  for (MVT VT : MVT::fixedlen_vector_valuetypes()) {
510  if (IsPTXVectorType(VT)) {
514  }
515  }
516 
517  // Custom handling for i8 intrinsics
519 
520  for (const auto& Ty : {MVT::i16, MVT::i32, MVT::i64}) {
526 
529  }
530 
535  if (STI.getPTXVersion() >= 43) {
540  }
541 
545 
546  // PTX does not directly support SELP of i1, so promote to i32 first
548 
549  // PTX cannot multiply two i64s in a single instruction.
552 
553  // We have some custom DAG combine patterns for these nodes
555  ISD::SREM, ISD::UREM});
556 
557  // setcc for f16x2 needs special handling to prevent legalizer's
558  // attempt to scalarize it due to v2i1 not being legal.
559  if (STI.allowFP16Math())
561 
562  // Promote fp16 arithmetic if fp16 hardware isn't available or the
563  // user passed --nvptx-no-fp16-math. The flag is useful because,
564  // although sm_53+ GPUs have some sort of FP16 support in
565  // hardware, only sm_53 and sm_60 have full implementation. Others
566  // only have token amount of hardware and are likely to run faster
567  // by using fp32 units instead.
568  for (const auto &Op : {ISD::FADD, ISD::FMUL, ISD::FSUB, ISD::FMA}) {
569  setFP16OperationAction(Op, MVT::f16, Legal, Promote);
570  setFP16OperationAction(Op, MVT::v2f16, Legal, Expand);
571  }
572 
573  // f16/f16x2 neg was introduced in PTX 60, SM_53.
574  const bool IsFP16FP16x2NegAvailable = STI.getSmVersion() >= 53 &&
575  STI.getPTXVersion() >= 60 &&
576  STI.allowFP16Math();
577  for (const auto &VT : {MVT::f16, MVT::v2f16})
579  IsFP16FP16x2NegAvailable ? Legal : Expand);
580 
581  // (would be) Library functions.
582 
583  // These map to conversion instructions for scalar FP types.
584  for (const auto &Op : {ISD::FCEIL, ISD::FFLOOR, ISD::FNEARBYINT, ISD::FRINT,
590  }
591 
596 
597 
598  // 'Expand' implements FCOPYSIGN without calling an external library.
603 
604  // These map to corresponding instructions for f32/f64. f16 must be
605  // promoted to f32. v2f16 is expanded to f16, which is then promoted
606  // to f32.
607  for (const auto &Op :
613  }
614  // max.f16, max.f16x2 and max.NaN are supported on sm_80+.
615  auto GetMinMaxAction = [&](LegalizeAction NotSm80Action) {
616  bool IsAtLeastSm80 = STI.getSmVersion() >= 80 && STI.getPTXVersion() >= 70;
617  return IsAtLeastSm80 ? Legal : NotSm80Action;
618  };
619  for (const auto &Op : {ISD::FMINNUM, ISD::FMAXNUM}) {
620  setFP16OperationAction(Op, MVT::f16, GetMinMaxAction(Promote), Promote);
623  setFP16OperationAction(Op, MVT::v2f16, GetMinMaxAction(Expand), Expand);
624  }
625  for (const auto &Op : {ISD::FMINIMUM, ISD::FMAXIMUM}) {
626  setFP16OperationAction(Op, MVT::f16, GetMinMaxAction(Expand), Expand);
627  setOperationAction(Op, MVT::f32, GetMinMaxAction(Expand));
628  setFP16OperationAction(Op, MVT::v2f16, GetMinMaxAction(Expand), Expand);
629  }
630 
631  // No FEXP2, FLOG2. The PTX ex2 and log2 functions are always approximate.
632  // No FPOW or FREM in PTX.
633 
634  // Now deduce the information based on the above mentioned
635  // actions
637 
639 }
640 
641 const char *NVPTXTargetLowering::getTargetNodeName(unsigned Opcode) const {
642  switch ((NVPTXISD::NodeType)Opcode) {
644  break;
645  case NVPTXISD::CALL:
646  return "NVPTXISD::CALL";
647  case NVPTXISD::RET_FLAG:
648  return "NVPTXISD::RET_FLAG";
650  return "NVPTXISD::LOAD_PARAM";
651  case NVPTXISD::Wrapper:
652  return "NVPTXISD::Wrapper";
654  return "NVPTXISD::DeclareParam";
656  return "NVPTXISD::DeclareScalarParam";
658  return "NVPTXISD::DeclareRet";
660  return "NVPTXISD::DeclareScalarRet";
662  return "NVPTXISD::DeclareRetParam";
663  case NVPTXISD::PrintCall:
664  return "NVPTXISD::PrintCall";
666  return "NVPTXISD::PrintConvergentCall";
668  return "NVPTXISD::PrintCallUni";
670  return "NVPTXISD::PrintConvergentCallUni";
671  case NVPTXISD::LoadParam:
672  return "NVPTXISD::LoadParam";
674  return "NVPTXISD::LoadParamV2";
676  return "NVPTXISD::LoadParamV4";
678  return "NVPTXISD::StoreParam";
680  return "NVPTXISD::StoreParamV2";
682  return "NVPTXISD::StoreParamV4";
684  return "NVPTXISD::StoreParamS32";
686  return "NVPTXISD::StoreParamU32";
688  return "NVPTXISD::CallArgBegin";
689  case NVPTXISD::CallArg:
690  return "NVPTXISD::CallArg";
692  return "NVPTXISD::LastCallArg";
694  return "NVPTXISD::CallArgEnd";
695  case NVPTXISD::CallVoid:
696  return "NVPTXISD::CallVoid";
697  case NVPTXISD::CallVal:
698  return "NVPTXISD::CallVal";
700  return "NVPTXISD::CallSymbol";
701  case NVPTXISD::Prototype:
702  return "NVPTXISD::Prototype";
703  case NVPTXISD::MoveParam:
704  return "NVPTXISD::MoveParam";
706  return "NVPTXISD::StoreRetval";
708  return "NVPTXISD::StoreRetvalV2";
710  return "NVPTXISD::StoreRetvalV4";
712  return "NVPTXISD::PseudoUseParam";
713  case NVPTXISD::RETURN:
714  return "NVPTXISD::RETURN";
716  return "NVPTXISD::CallSeqBegin";
718  return "NVPTXISD::CallSeqEnd";
720  return "NVPTXISD::CallPrototype";
721  case NVPTXISD::ProxyReg:
722  return "NVPTXISD::ProxyReg";
723  case NVPTXISD::LoadV2:
724  return "NVPTXISD::LoadV2";
725  case NVPTXISD::LoadV4:
726  return "NVPTXISD::LoadV4";
727  case NVPTXISD::LDGV2:
728  return "NVPTXISD::LDGV2";
729  case NVPTXISD::LDGV4:
730  return "NVPTXISD::LDGV4";
731  case NVPTXISD::LDUV2:
732  return "NVPTXISD::LDUV2";
733  case NVPTXISD::LDUV4:
734  return "NVPTXISD::LDUV4";
735  case NVPTXISD::StoreV2:
736  return "NVPTXISD::StoreV2";
737  case NVPTXISD::StoreV4:
738  return "NVPTXISD::StoreV4";
740  return "NVPTXISD::FUN_SHFL_CLAMP";
742  return "NVPTXISD::FUN_SHFR_CLAMP";
743  case NVPTXISD::IMAD:
744  return "NVPTXISD::IMAD";
746  return "NVPTXISD::SETP_F16X2";
747  case NVPTXISD::Dummy:
748  return "NVPTXISD::Dummy";
750  return "NVPTXISD::MUL_WIDE_SIGNED";
752  return "NVPTXISD::MUL_WIDE_UNSIGNED";
753  case NVPTXISD::Tex1DFloatS32: return "NVPTXISD::Tex1DFloatS32";
754  case NVPTXISD::Tex1DFloatFloat: return "NVPTXISD::Tex1DFloatFloat";
756  return "NVPTXISD::Tex1DFloatFloatLevel";
758  return "NVPTXISD::Tex1DFloatFloatGrad";
759  case NVPTXISD::Tex1DS32S32: return "NVPTXISD::Tex1DS32S32";
760  case NVPTXISD::Tex1DS32Float: return "NVPTXISD::Tex1DS32Float";
762  return "NVPTXISD::Tex1DS32FloatLevel";
764  return "NVPTXISD::Tex1DS32FloatGrad";
765  case NVPTXISD::Tex1DU32S32: return "NVPTXISD::Tex1DU32S32";
766  case NVPTXISD::Tex1DU32Float: return "NVPTXISD::Tex1DU32Float";
768  return "NVPTXISD::Tex1DU32FloatLevel";
770  return "NVPTXISD::Tex1DU32FloatGrad";
771  case NVPTXISD::Tex1DArrayFloatS32: return "NVPTXISD::Tex1DArrayFloatS32";
772  case NVPTXISD::Tex1DArrayFloatFloat: return "NVPTXISD::Tex1DArrayFloatFloat";
774  return "NVPTXISD::Tex1DArrayFloatFloatLevel";
776  return "NVPTXISD::Tex1DArrayFloatFloatGrad";
777  case NVPTXISD::Tex1DArrayS32S32: return "NVPTXISD::Tex1DArrayS32S32";
778  case NVPTXISD::Tex1DArrayS32Float: return "NVPTXISD::Tex1DArrayS32Float";
780  return "NVPTXISD::Tex1DArrayS32FloatLevel";
782  return "NVPTXISD::Tex1DArrayS32FloatGrad";
783  case NVPTXISD::Tex1DArrayU32S32: return "NVPTXISD::Tex1DArrayU32S32";
784  case NVPTXISD::Tex1DArrayU32Float: return "NVPTXISD::Tex1DArrayU32Float";
786  return "NVPTXISD::Tex1DArrayU32FloatLevel";
788  return "NVPTXISD::Tex1DArrayU32FloatGrad";
789  case NVPTXISD::Tex2DFloatS32: return "NVPTXISD::Tex2DFloatS32";
790  case NVPTXISD::Tex2DFloatFloat: return "NVPTXISD::Tex2DFloatFloat";
792  return "NVPTXISD::Tex2DFloatFloatLevel";
794  return "NVPTXISD::Tex2DFloatFloatGrad";
795  case NVPTXISD::Tex2DS32S32: return "NVPTXISD::Tex2DS32S32";
796  case NVPTXISD::Tex2DS32Float: return "NVPTXISD::Tex2DS32Float";
798  return "NVPTXISD::Tex2DS32FloatLevel";
800  return "NVPTXISD::Tex2DS32FloatGrad";
801  case NVPTXISD::Tex2DU32S32: return "NVPTXISD::Tex2DU32S32";
802  case NVPTXISD::Tex2DU32Float: return "NVPTXISD::Tex2DU32Float";
804  return "NVPTXISD::Tex2DU32FloatLevel";
806  return "NVPTXISD::Tex2DU32FloatGrad";
807  case NVPTXISD::Tex2DArrayFloatS32: return "NVPTXISD::Tex2DArrayFloatS32";
808  case NVPTXISD::Tex2DArrayFloatFloat: return "NVPTXISD::Tex2DArrayFloatFloat";
810  return "NVPTXISD::Tex2DArrayFloatFloatLevel";
812  return "NVPTXISD::Tex2DArrayFloatFloatGrad";
813  case NVPTXISD::Tex2DArrayS32S32: return "NVPTXISD::Tex2DArrayS32S32";
814  case NVPTXISD::Tex2DArrayS32Float: return "NVPTXISD::Tex2DArrayS32Float";
816  return "NVPTXISD::Tex2DArrayS32FloatLevel";
818  return "NVPTXISD::Tex2DArrayS32FloatGrad";
819  case NVPTXISD::Tex2DArrayU32S32: return "NVPTXISD::Tex2DArrayU32S32";
820  case NVPTXISD::Tex2DArrayU32Float: return "NVPTXISD::Tex2DArrayU32Float";
822  return "NVPTXISD::Tex2DArrayU32FloatLevel";
824  return "NVPTXISD::Tex2DArrayU32FloatGrad";
825  case NVPTXISD::Tex3DFloatS32: return "NVPTXISD::Tex3DFloatS32";
826  case NVPTXISD::Tex3DFloatFloat: return "NVPTXISD::Tex3DFloatFloat";
828  return "NVPTXISD::Tex3DFloatFloatLevel";
830  return "NVPTXISD::Tex3DFloatFloatGrad";
831  case NVPTXISD::Tex3DS32S32: return "NVPTXISD::Tex3DS32S32";
832  case NVPTXISD::Tex3DS32Float: return "NVPTXISD::Tex3DS32Float";
834  return "NVPTXISD::Tex3DS32FloatLevel";
836  return "NVPTXISD::Tex3DS32FloatGrad";
837  case NVPTXISD::Tex3DU32S32: return "NVPTXISD::Tex3DU32S32";
838  case NVPTXISD::Tex3DU32Float: return "NVPTXISD::Tex3DU32Float";
840  return "NVPTXISD::Tex3DU32FloatLevel";
842  return "NVPTXISD::Tex3DU32FloatGrad";
843  case NVPTXISD::TexCubeFloatFloat: return "NVPTXISD::TexCubeFloatFloat";
845  return "NVPTXISD::TexCubeFloatFloatLevel";
846  case NVPTXISD::TexCubeS32Float: return "NVPTXISD::TexCubeS32Float";
848  return "NVPTXISD::TexCubeS32FloatLevel";
849  case NVPTXISD::TexCubeU32Float: return "NVPTXISD::TexCubeU32Float";
851  return "NVPTXISD::TexCubeU32FloatLevel";
853  return "NVPTXISD::TexCubeArrayFloatFloat";
855  return "NVPTXISD::TexCubeArrayFloatFloatLevel";
857  return "NVPTXISD::TexCubeArrayS32Float";
859  return "NVPTXISD::TexCubeArrayS32FloatLevel";
861  return "NVPTXISD::TexCubeArrayU32Float";
863  return "NVPTXISD::TexCubeArrayU32FloatLevel";
865  return "NVPTXISD::Tld4R2DFloatFloat";
867  return "NVPTXISD::Tld4G2DFloatFloat";
869  return "NVPTXISD::Tld4B2DFloatFloat";
871  return "NVPTXISD::Tld4A2DFloatFloat";
873  return "NVPTXISD::Tld4R2DS64Float";
875  return "NVPTXISD::Tld4G2DS64Float";
877  return "NVPTXISD::Tld4B2DS64Float";
879  return "NVPTXISD::Tld4A2DS64Float";
881  return "NVPTXISD::Tld4R2DU64Float";
883  return "NVPTXISD::Tld4G2DU64Float";
885  return "NVPTXISD::Tld4B2DU64Float";
887  return "NVPTXISD::Tld4A2DU64Float";
888 
890  return "NVPTXISD::TexUnified1DFloatS32";
892  return "NVPTXISD::TexUnified1DFloatFloat";
894  return "NVPTXISD::TexUnified1DFloatFloatLevel";
896  return "NVPTXISD::TexUnified1DFloatFloatGrad";
898  return "NVPTXISD::TexUnified1DS32S32";
900  return "NVPTXISD::TexUnified1DS32Float";
902  return "NVPTXISD::TexUnified1DS32FloatLevel";
904  return "NVPTXISD::TexUnified1DS32FloatGrad";
906  return "NVPTXISD::TexUnified1DU32S32";
908  return "NVPTXISD::TexUnified1DU32Float";
910  return "NVPTXISD::TexUnified1DU32FloatLevel";
912  return "NVPTXISD::TexUnified1DU32FloatGrad";
914  return "NVPTXISD::TexUnified1DArrayFloatS32";
916  return "NVPTXISD::TexUnified1DArrayFloatFloat";
918  return "NVPTXISD::TexUnified1DArrayFloatFloatLevel";
920  return "NVPTXISD::TexUnified1DArrayFloatFloatGrad";
922  return "NVPTXISD::TexUnified1DArrayS32S32";
924  return "NVPTXISD::TexUnified1DArrayS32Float";
926  return "NVPTXISD::TexUnified1DArrayS32FloatLevel";
928  return "NVPTXISD::TexUnified1DArrayS32FloatGrad";
930  return "NVPTXISD::TexUnified1DArrayU32S32";
932  return "NVPTXISD::TexUnified1DArrayU32Float";
934  return "NVPTXISD::TexUnified1DArrayU32FloatLevel";
936  return "NVPTXISD::TexUnified1DArrayU32FloatGrad";
938  return "NVPTXISD::TexUnified2DFloatS32";
940  return "NVPTXISD::TexUnified2DFloatFloat";
942  return "NVPTXISD::TexUnified2DFloatFloatLevel";
944  return "NVPTXISD::TexUnified2DFloatFloatGrad";
946  return "NVPTXISD::TexUnified2DS32S32";
948  return "NVPTXISD::TexUnified2DS32Float";
950  return "NVPTXISD::TexUnified2DS32FloatLevel";
952  return "NVPTXISD::TexUnified2DS32FloatGrad";
954  return "NVPTXISD::TexUnified2DU32S32";
956  return "NVPTXISD::TexUnified2DU32Float";
958  return "NVPTXISD::TexUnified2DU32FloatLevel";
960  return "NVPTXISD::TexUnified2DU32FloatGrad";
962  return "NVPTXISD::TexUnified2DArrayFloatS32";
964  return "NVPTXISD::TexUnified2DArrayFloatFloat";
966  return "NVPTXISD::TexUnified2DArrayFloatFloatLevel";
968  return "NVPTXISD::TexUnified2DArrayFloatFloatGrad";
970  return "NVPTXISD::TexUnified2DArrayS32S32";
972  return "NVPTXISD::TexUnified2DArrayS32Float";
974  return "NVPTXISD::TexUnified2DArrayS32FloatLevel";
976  return "NVPTXISD::TexUnified2DArrayS32FloatGrad";
978  return "NVPTXISD::TexUnified2DArrayU32S32";
980  return "NVPTXISD::TexUnified2DArrayU32Float";
982  return "NVPTXISD::TexUnified2DArrayU32FloatLevel";
984  return "NVPTXISD::TexUnified2DArrayU32FloatGrad";
986  return "NVPTXISD::TexUnified3DFloatS32";
988  return "NVPTXISD::TexUnified3DFloatFloat";
990  return "NVPTXISD::TexUnified3DFloatFloatLevel";
992  return "NVPTXISD::TexUnified3DFloatFloatGrad";
994  return "NVPTXISD::TexUnified3DS32S32";
996  return "NVPTXISD::TexUnified3DS32Float";
998  return "NVPTXISD::TexUnified3DS32FloatLevel";
1000  return "NVPTXISD::TexUnified3DS32FloatGrad";
1002  return "NVPTXISD::TexUnified3DU32S32";
1004  return "NVPTXISD::TexUnified3DU32Float";
1006  return "NVPTXISD::TexUnified3DU32FloatLevel";
1008  return "NVPTXISD::TexUnified3DU32FloatGrad";
1010  return "NVPTXISD::TexUnifiedCubeFloatFloat";
1012  return "NVPTXISD::TexUnifiedCubeFloatFloatLevel";
1014  return "NVPTXISD::TexUnifiedCubeS32Float";
1016  return "NVPTXISD::TexUnifiedCubeS32FloatLevel";
1018  return "NVPTXISD::TexUnifiedCubeU32Float";
1020  return "NVPTXISD::TexUnifiedCubeU32FloatLevel";
1022  return "NVPTXISD::TexUnifiedCubeArrayFloatFloat";
1024  return "NVPTXISD::TexUnifiedCubeArrayFloatFloatLevel";
1026  return "NVPTXISD::TexUnifiedCubeArrayS32Float";
1028  return "NVPTXISD::TexUnifiedCubeArrayS32FloatLevel";
1030  return "NVPTXISD::TexUnifiedCubeArrayU32Float";
1032  return "NVPTXISD::TexUnifiedCubeArrayU32FloatLevel";
1034  return "NVPTXISD::Tld4UnifiedR2DFloatFloat";
1036  return "NVPTXISD::Tld4UnifiedG2DFloatFloat";
1038  return "NVPTXISD::Tld4UnifiedB2DFloatFloat";
1040  return "NVPTXISD::Tld4UnifiedA2DFloatFloat";
1042  return "NVPTXISD::Tld4UnifiedR2DS64Float";
1044  return "NVPTXISD::Tld4UnifiedG2DS64Float";
1046  return "NVPTXISD::Tld4UnifiedB2DS64Float";
1048  return "NVPTXISD::Tld4UnifiedA2DS64Float";
1050  return "NVPTXISD::Tld4UnifiedR2DU64Float";
1052  return "NVPTXISD::Tld4UnifiedG2DU64Float";
1054  return "NVPTXISD::Tld4UnifiedB2DU64Float";
1056  return "NVPTXISD::Tld4UnifiedA2DU64Float";
1057 
1058  case NVPTXISD::Suld1DI8Clamp: return "NVPTXISD::Suld1DI8Clamp";
1059  case NVPTXISD::Suld1DI16Clamp: return "NVPTXISD::Suld1DI16Clamp";
1060  case NVPTXISD::Suld1DI32Clamp: return "NVPTXISD::Suld1DI32Clamp";
1061  case NVPTXISD::Suld1DI64Clamp: return "NVPTXISD::Suld1DI64Clamp";
1062  case NVPTXISD::Suld1DV2I8Clamp: return "NVPTXISD::Suld1DV2I8Clamp";
1063  case NVPTXISD::Suld1DV2I16Clamp: return "NVPTXISD::Suld1DV2I16Clamp";
1064  case NVPTXISD::Suld1DV2I32Clamp: return "NVPTXISD::Suld1DV2I32Clamp";
1065  case NVPTXISD::Suld1DV2I64Clamp: return "NVPTXISD::Suld1DV2I64Clamp";
1066  case NVPTXISD::Suld1DV4I8Clamp: return "NVPTXISD::Suld1DV4I8Clamp";
1067  case NVPTXISD::Suld1DV4I16Clamp: return "NVPTXISD::Suld1DV4I16Clamp";
1068  case NVPTXISD::Suld1DV4I32Clamp: return "NVPTXISD::Suld1DV4I32Clamp";
1069 
1070  case NVPTXISD::Suld1DArrayI8Clamp: return "NVPTXISD::Suld1DArrayI8Clamp";
1071  case NVPTXISD::Suld1DArrayI16Clamp: return "NVPTXISD::Suld1DArrayI16Clamp";
1072  case NVPTXISD::Suld1DArrayI32Clamp: return "NVPTXISD::Suld1DArrayI32Clamp";
1073  case NVPTXISD::Suld1DArrayI64Clamp: return "NVPTXISD::Suld1DArrayI64Clamp";
1074  case NVPTXISD::Suld1DArrayV2I8Clamp: return "NVPTXISD::Suld1DArrayV2I8Clamp";
1075  case NVPTXISD::Suld1DArrayV2I16Clamp:return "NVPTXISD::Suld1DArrayV2I16Clamp";
1076  case NVPTXISD::Suld1DArrayV2I32Clamp:return "NVPTXISD::Suld1DArrayV2I32Clamp";
1077  case NVPTXISD::Suld1DArrayV2I64Clamp:return "NVPTXISD::Suld1DArrayV2I64Clamp";
1078  case NVPTXISD::Suld1DArrayV4I8Clamp: return "NVPTXISD::Suld1DArrayV4I8Clamp";
1079  case NVPTXISD::Suld1DArrayV4I16Clamp:return "NVPTXISD::Suld1DArrayV4I16Clamp";
1080  case NVPTXISD::Suld1DArrayV4I32Clamp:return "NVPTXISD::Suld1DArrayV4I32Clamp";
1081 
1082  case NVPTXISD::Suld2DI8Clamp: return "NVPTXISD::Suld2DI8Clamp";
1083  case NVPTXISD::Suld2DI16Clamp: return "NVPTXISD::Suld2DI16Clamp";
1084  case NVPTXISD::Suld2DI32Clamp: return "NVPTXISD::Suld2DI32Clamp";
1085  case NVPTXISD::Suld2DI64Clamp: return "NVPTXISD::Suld2DI64Clamp";
1086  case NVPTXISD::Suld2DV2I8Clamp: return "NVPTXISD::Suld2DV2I8Clamp";
1087  case NVPTXISD::Suld2DV2I16Clamp: return "NVPTXISD::Suld2DV2I16Clamp";
1088  case NVPTXISD::Suld2DV2I32Clamp: return "NVPTXISD::Suld2DV2I32Clamp";
1089  case NVPTXISD::Suld2DV2I64Clamp: return "NVPTXISD::Suld2DV2I64Clamp";
1090  case NVPTXISD::Suld2DV4I8Clamp: return "NVPTXISD::Suld2DV4I8Clamp";
1091  case NVPTXISD::Suld2DV4I16Clamp: return "NVPTXISD::Suld2DV4I16Clamp";
1092  case NVPTXISD::Suld2DV4I32Clamp: return "NVPTXISD::Suld2DV4I32Clamp";
1093 
1094  case NVPTXISD::Suld2DArrayI8Clamp: return "NVPTXISD::Suld2DArrayI8Clamp";
1095  case NVPTXISD::Suld2DArrayI16Clamp: return "NVPTXISD::Suld2DArrayI16Clamp";
1096  case NVPTXISD::Suld2DArrayI32Clamp: return "NVPTXISD::Suld2DArrayI32Clamp";
1097  case NVPTXISD::Suld2DArrayI64Clamp: return "NVPTXISD::Suld2DArrayI64Clamp";
1098  case NVPTXISD::Suld2DArrayV2I8Clamp: return "NVPTXISD::Suld2DArrayV2I8Clamp";
1099  case NVPTXISD::Suld2DArrayV2I16Clamp:return "NVPTXISD::Suld2DArrayV2I16Clamp";
1100  case NVPTXISD::Suld2DArrayV2I32Clamp:return "NVPTXISD::Suld2DArrayV2I32Clamp";
1101  case NVPTXISD::Suld2DArrayV2I64Clamp:return "NVPTXISD::Suld2DArrayV2I64Clamp";
1102  case NVPTXISD::Suld2DArrayV4I8Clamp: return "NVPTXISD::Suld2DArrayV4I8Clamp";
1103  case NVPTXISD::Suld2DArrayV4I16Clamp:return "NVPTXISD::Suld2DArrayV4I16Clamp";
1104  case NVPTXISD::Suld2DArrayV4I32Clamp:return "NVPTXISD::Suld2DArrayV4I32Clamp";
1105 
1106  case NVPTXISD::Suld3DI8Clamp: return "NVPTXISD::Suld3DI8Clamp";
1107  case NVPTXISD::Suld3DI16Clamp: return "NVPTXISD::Suld3DI16Clamp";
1108  case NVPTXISD::Suld3DI32Clamp: return "NVPTXISD::Suld3DI32Clamp";
1109  case NVPTXISD::Suld3DI64Clamp: return "NVPTXISD::Suld3DI64Clamp";
1110  case NVPTXISD::Suld3DV2I8Clamp: return "NVPTXISD::Suld3DV2I8Clamp";
1111  case NVPTXISD::Suld3DV2I16Clamp: return "NVPTXISD::Suld3DV2I16Clamp";
1112  case NVPTXISD::Suld3DV2I32Clamp: return "NVPTXISD::Suld3DV2I32Clamp";
1113  case NVPTXISD::Suld3DV2I64Clamp: return "NVPTXISD::Suld3DV2I64Clamp";
1114  case NVPTXISD::Suld3DV4I8Clamp: return "NVPTXISD::Suld3DV4I8Clamp";
1115  case NVPTXISD::Suld3DV4I16Clamp: return "NVPTXISD::Suld3DV4I16Clamp";
1116  case NVPTXISD::Suld3DV4I32Clamp: return "NVPTXISD::Suld3DV4I32Clamp";
1117 
1118  case NVPTXISD::Suld1DI8Trap: return "NVPTXISD::Suld1DI8Trap";
1119  case NVPTXISD::Suld1DI16Trap: return "NVPTXISD::Suld1DI16Trap";
1120  case NVPTXISD::Suld1DI32Trap: return "NVPTXISD::Suld1DI32Trap";
1121  case NVPTXISD::Suld1DI64Trap: return "NVPTXISD::Suld1DI64Trap";
1122  case NVPTXISD::Suld1DV2I8Trap: return "NVPTXISD::Suld1DV2I8Trap";
1123  case NVPTXISD::Suld1DV2I16Trap: return "NVPTXISD::Suld1DV2I16Trap";
1124  case NVPTXISD::Suld1DV2I32Trap: return "NVPTXISD::Suld1DV2I32Trap";
1125  case NVPTXISD::Suld1DV2I64Trap: return "NVPTXISD::Suld1DV2I64Trap";
1126  case NVPTXISD::Suld1DV4I8Trap: return "NVPTXISD::Suld1DV4I8Trap";
1127  case NVPTXISD::Suld1DV4I16Trap: return "NVPTXISD::Suld1DV4I16Trap";
1128  case NVPTXISD::Suld1DV4I32Trap: return "NVPTXISD::Suld1DV4I32Trap";
1129 
1130  case NVPTXISD::Suld1DArrayI8Trap: return "NVPTXISD::Suld1DArrayI8Trap";
1131  case NVPTXISD::Suld1DArrayI16Trap: return "NVPTXISD::Suld1DArrayI16Trap";
1132  case NVPTXISD::Suld1DArrayI32Trap: return "NVPTXISD::Suld1DArrayI32Trap";
1133  case NVPTXISD::Suld1DArrayI64Trap: return "NVPTXISD::Suld1DArrayI64Trap";
1134  case NVPTXISD::Suld1DArrayV2I8Trap: return "NVPTXISD::Suld1DArrayV2I8Trap";
1135  case NVPTXISD::Suld1DArrayV2I16Trap: return "NVPTXISD::Suld1DArrayV2I16Trap";
1136  case NVPTXISD::Suld1DArrayV2I32Trap: return "NVPTXISD::Suld1DArrayV2I32Trap";
1137  case NVPTXISD::Suld1DArrayV2I64Trap: return "NVPTXISD::Suld1DArrayV2I64Trap";
1138  case NVPTXISD::Suld1DArrayV4I8Trap: return "NVPTXISD::Suld1DArrayV4I8Trap";
1139  case NVPTXISD::Suld1DArrayV4I16Trap: return "NVPTXISD::Suld1DArrayV4I16Trap";
1140  case NVPTXISD::Suld1DArrayV4I32Trap: return "NVPTXISD::Suld1DArrayV4I32Trap";
1141 
1142  case NVPTXISD::Suld2DI8Trap: return "NVPTXISD::Suld2DI8Trap";
1143  case NVPTXISD::Suld2DI16Trap: return "NVPTXISD::Suld2DI16Trap";
1144  case NVPTXISD::Suld2DI32Trap: return "NVPTXISD::Suld2DI32Trap";
1145  case NVPTXISD::Suld2DI64Trap: return "NVPTXISD::Suld2DI64Trap";
1146  case NVPTXISD::Suld2DV2I8Trap: return "NVPTXISD::Suld2DV2I8Trap";
1147  case NVPTXISD::Suld2DV2I16Trap: return "NVPTXISD::Suld2DV2I16Trap";
1148  case NVPTXISD::Suld2DV2I32Trap: return "NVPTXISD::Suld2DV2I32Trap";
1149  case NVPTXISD::Suld2DV2I64Trap: return "NVPTXISD::Suld2DV2I64Trap";
1150  case NVPTXISD::Suld2DV4I8Trap: return "NVPTXISD::Suld2DV4I8Trap";
1151  case NVPTXISD::Suld2DV4I16Trap: return "NVPTXISD::Suld2DV4I16Trap";
1152  case NVPTXISD::Suld2DV4I32Trap: return "NVPTXISD::Suld2DV4I32Trap";
1153 
1154  case NVPTXISD::Suld2DArrayI8Trap: return "NVPTXISD::Suld2DArrayI8Trap";
1155  case NVPTXISD::Suld2DArrayI16Trap: return "NVPTXISD::Suld2DArrayI16Trap";
1156  case NVPTXISD::Suld2DArrayI32Trap: return "NVPTXISD::Suld2DArrayI32Trap";
1157  case NVPTXISD::Suld2DArrayI64Trap: return "NVPTXISD::Suld2DArrayI64Trap";
1158  case NVPTXISD::Suld2DArrayV2I8Trap: return "NVPTXISD::Suld2DArrayV2I8Trap";
1159  case NVPTXISD::Suld2DArrayV2I16Trap: return "NVPTXISD::Suld2DArrayV2I16Trap";
1160  case NVPTXISD::Suld2DArrayV2I32Trap: return "NVPTXISD::Suld2DArrayV2I32Trap";
1161  case NVPTXISD::Suld2DArrayV2I64Trap: return "NVPTXISD::Suld2DArrayV2I64Trap";
1162  case NVPTXISD::Suld2DArrayV4I8Trap: return "NVPTXISD::Suld2DArrayV4I8Trap";
1163  case NVPTXISD::Suld2DArrayV4I16Trap: return "NVPTXISD::Suld2DArrayV4I16Trap";
1164  case NVPTXISD::Suld2DArrayV4I32Trap: return "NVPTXISD::Suld2DArrayV4I32Trap";
1165 
1166  case NVPTXISD::Suld3DI8Trap: return "NVPTXISD::Suld3DI8Trap";
1167  case NVPTXISD::Suld3DI16Trap: return "NVPTXISD::Suld3DI16Trap";
1168  case NVPTXISD::Suld3DI32Trap: return "NVPTXISD::Suld3DI32Trap";
1169  case NVPTXISD::Suld3DI64Trap: return "NVPTXISD::Suld3DI64Trap";
1170  case NVPTXISD::Suld3DV2I8Trap: return "NVPTXISD::Suld3DV2I8Trap";
1171  case NVPTXISD::Suld3DV2I16Trap: return "NVPTXISD::Suld3DV2I16Trap";
1172  case NVPTXISD::Suld3DV2I32Trap: return "NVPTXISD::Suld3DV2I32Trap";
1173  case NVPTXISD::Suld3DV2I64Trap: return "NVPTXISD::Suld3DV2I64Trap";
1174  case NVPTXISD::Suld3DV4I8Trap: return "NVPTXISD::Suld3DV4I8Trap";
1175  case NVPTXISD::Suld3DV4I16Trap: return "NVPTXISD::Suld3DV4I16Trap";
1176  case NVPTXISD::Suld3DV4I32Trap: return "NVPTXISD::Suld3DV4I32Trap";
1177 
1178  case NVPTXISD::Suld1DI8Zero: return "NVPTXISD::Suld1DI8Zero";
1179  case NVPTXISD::Suld1DI16Zero: return "NVPTXISD::Suld1DI16Zero";
1180  case NVPTXISD::Suld1DI32Zero: return "NVPTXISD::Suld1DI32Zero";
1181  case NVPTXISD::Suld1DI64Zero: return "NVPTXISD::Suld1DI64Zero";
1182  case NVPTXISD::Suld1DV2I8Zero: return "NVPTXISD::Suld1DV2I8Zero";
1183  case NVPTXISD::Suld1DV2I16Zero: return "NVPTXISD::Suld1DV2I16Zero";
1184  case NVPTXISD::Suld1DV2I32Zero: return "NVPTXISD::Suld1DV2I32Zero";
1185  case NVPTXISD::Suld1DV2I64Zero: return "NVPTXISD::Suld1DV2I64Zero";
1186  case NVPTXISD::Suld1DV4I8Zero: return "NVPTXISD::Suld1DV4I8Zero";
1187  case NVPTXISD::Suld1DV4I16Zero: return "NVPTXISD::Suld1DV4I16Zero";
1188  case NVPTXISD::Suld1DV4I32Zero: return "NVPTXISD::Suld1DV4I32Zero";
1189 
1190  case NVPTXISD::Suld1DArrayI8Zero: return "NVPTXISD::Suld1DArrayI8Zero";
1191  case NVPTXISD::Suld1DArrayI16Zero: return "NVPTXISD::Suld1DArrayI16Zero";
1192  case NVPTXISD::Suld1DArrayI32Zero: return "NVPTXISD::Suld1DArrayI32Zero";
1193  case NVPTXISD::Suld1DArrayI64Zero: return "NVPTXISD::Suld1DArrayI64Zero";
1194  case NVPTXISD::Suld1DArrayV2I8Zero: return "NVPTXISD::Suld1DArrayV2I8Zero";
1195  case NVPTXISD::Suld1DArrayV2I16Zero: return "NVPTXISD::Suld1DArrayV2I16Zero";
1196  case NVPTXISD::Suld1DArrayV2I32Zero: return "NVPTXISD::Suld1DArrayV2I32Zero";
1197  case NVPTXISD::Suld1DArrayV2I64Zero: return "NVPTXISD::Suld1DArrayV2I64Zero";
1198  case NVPTXISD::Suld1DArrayV4I8Zero: return "NVPTXISD::Suld1DArrayV4I8Zero";
1199  case NVPTXISD::Suld1DArrayV4I16Zero: return "NVPTXISD::Suld1DArrayV4I16Zero";
1200  case NVPTXISD::Suld1DArrayV4I32Zero: return "NVPTXISD::Suld1DArrayV4I32Zero";
1201 
1202  case NVPTXISD::Suld2DI8Zero: return "NVPTXISD::Suld2DI8Zero";
1203  case NVPTXISD::Suld2DI16Zero: return "NVPTXISD::Suld2DI16Zero";
1204  case NVPTXISD::Suld2DI32Zero: return "NVPTXISD::Suld2DI32Zero";
1205  case NVPTXISD::Suld2DI64Zero: return "NVPTXISD::Suld2DI64Zero";
1206  case NVPTXISD::Suld2DV2I8Zero: return "NVPTXISD::Suld2DV2I8Zero";
1207  case NVPTXISD::Suld2DV2I16Zero: return "NVPTXISD::Suld2DV2I16Zero";
1208  case NVPTXISD::Suld2DV2I32Zero: return "NVPTXISD::Suld2DV2I32Zero";
1209  case NVPTXISD::Suld2DV2I64Zero: return "NVPTXISD::Suld2DV2I64Zero";
1210  case NVPTXISD::Suld2DV4I8Zero: return "NVPTXISD::Suld2DV4I8Zero";
1211  case NVPTXISD::Suld2DV4I16Zero: return "NVPTXISD::Suld2DV4I16Zero";
1212  case NVPTXISD::Suld2DV4I32Zero: return "NVPTXISD::Suld2DV4I32Zero";
1213 
1214  case NVPTXISD::Suld2DArrayI8Zero: return "NVPTXISD::Suld2DArrayI8Zero";
1215  case NVPTXISD::Suld2DArrayI16Zero: return "NVPTXISD::Suld2DArrayI16Zero";
1216  case NVPTXISD::Suld2DArrayI32Zero: return "NVPTXISD::Suld2DArrayI32Zero";
1217  case NVPTXISD::Suld2DArrayI64Zero: return "NVPTXISD::Suld2DArrayI64Zero";
1218  case NVPTXISD::Suld2DArrayV2I8Zero: return "NVPTXISD::Suld2DArrayV2I8Zero";
1219  case NVPTXISD::Suld2DArrayV2I16Zero: return "NVPTXISD::Suld2DArrayV2I16Zero";
1220  case NVPTXISD::Suld2DArrayV2I32Zero: return "NVPTXISD::Suld2DArrayV2I32Zero";
1221  case NVPTXISD::Suld2DArrayV2I64Zero: return "NVPTXISD::Suld2DArrayV2I64Zero";
1222  case NVPTXISD::Suld2DArrayV4I8Zero: return "NVPTXISD::Suld2DArrayV4I8Zero";
1223  case NVPTXISD::Suld2DArrayV4I16Zero: return "NVPTXISD::Suld2DArrayV4I16Zero";
1224  case NVPTXISD::Suld2DArrayV4I32Zero: return "NVPTXISD::Suld2DArrayV4I32Zero";
1225 
1226  case NVPTXISD::Suld3DI8Zero: return "NVPTXISD::Suld3DI8Zero";
1227  case NVPTXISD::Suld3DI16Zero: return "NVPTXISD::Suld3DI16Zero";
1228  case NVPTXISD::Suld3DI32Zero: return "NVPTXISD::Suld3DI32Zero";
1229  case NVPTXISD::Suld3DI64Zero: return "NVPTXISD::Suld3DI64Zero";
1230  case NVPTXISD::Suld3DV2I8Zero: return "NVPTXISD::Suld3DV2I8Zero";
1231  case NVPTXISD::Suld3DV2I16Zero: return "NVPTXISD::Suld3DV2I16Zero";
1232  case NVPTXISD::Suld3DV2I32Zero: return "NVPTXISD::Suld3DV2I32Zero";
1233  case NVPTXISD::Suld3DV2I64Zero: return "NVPTXISD::Suld3DV2I64Zero";
1234  case NVPTXISD::Suld3DV4I8Zero: return "NVPTXISD::Suld3DV4I8Zero";
1235  case NVPTXISD::Suld3DV4I16Zero: return "NVPTXISD::Suld3DV4I16Zero";
1236  case NVPTXISD::Suld3DV4I32Zero: return "NVPTXISD::Suld3DV4I32Zero";
1237  }
1238  return nullptr;
1239 }
1240 
1243  if (!VT.isScalableVector() && VT.getVectorNumElements() != 1 &&
1244  VT.getScalarType() == MVT::i1)
1245  return TypeSplitVector;
1246  if (VT == MVT::v2f16)
1247  return TypeLegal;
1249 }
1250 
1252  int Enabled, int &ExtraSteps,
1253  bool &UseOneConst,
1254  bool Reciprocal) const {
1256  (Enabled == ReciprocalEstimate::Unspecified && !usePrecSqrtF32())))
1257  return SDValue();
1258 
1259  if (ExtraSteps == ReciprocalEstimate::Unspecified)
1260  ExtraSteps = 0;
1261 
1262  SDLoc DL(Operand);
1263  EVT VT = Operand.getValueType();
1264  bool Ftz = useF32FTZ(DAG.getMachineFunction());
1265 
1266  auto MakeIntrinsicCall = [&](Intrinsic::ID IID) {
1267  return DAG.getNode(ISD::INTRINSIC_WO_CHAIN, DL, VT,
1268  DAG.getConstant(IID, DL, MVT::i32), Operand);
1269  };
1270 
1271  // The sqrt and rsqrt refinement processes assume we always start out with an
1272  // approximation of the rsqrt. Therefore, if we're going to do any refinement
1273  // (i.e. ExtraSteps > 0), we must return an rsqrt. But if we're *not* doing
1274  // any refinement, we must return a regular sqrt.
1275  if (Reciprocal || ExtraSteps > 0) {
1276  if (VT == MVT::f32)
1277  return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_rsqrt_approx_ftz_f
1278  : Intrinsic::nvvm_rsqrt_approx_f);
1279  else if (VT == MVT::f64)
1280  return MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d);
1281  else
1282  return SDValue();
1283  } else {
1284  if (VT == MVT::f32)
1285  return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_sqrt_approx_ftz_f
1286  : Intrinsic::nvvm_sqrt_approx_f);
1287  else {
1288  // There's no sqrt.approx.f64 instruction, so we emit
1289  // reciprocal(rsqrt(x)). This is faster than
1290  // select(x == 0, 0, x * rsqrt(x)). (In fact, it's faster than plain
1291  // x * rsqrt(x).)
1292  return DAG.getNode(
1294  DAG.getConstant(Intrinsic::nvvm_rcp_approx_ftz_d, DL, MVT::i32),
1295  MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d));
1296  }
1297  }
1298 }
1299 
1300 SDValue
1302  SDLoc dl(Op);
1303  const GlobalAddressSDNode *GAN = cast<GlobalAddressSDNode>(Op);
1304  auto PtrVT = getPointerTy(DAG.getDataLayout(), GAN->getAddressSpace());
1305  Op = DAG.getTargetGlobalAddress(GAN->getGlobal(), dl, PtrVT);
1306  return DAG.getNode(NVPTXISD::Wrapper, dl, PtrVT, Op);
1307 }
1308 
1310  const DataLayout &DL, Type *retTy, const ArgListTy &Args,
1311  const SmallVectorImpl<ISD::OutputArg> &Outs, MaybeAlign retAlignment,
1312  const CallBase &CB, unsigned UniqueCallSite) const {
1313  auto PtrVT = getPointerTy(DL);
1314 
1315  bool isABI = (STI.getSmVersion() >= 20);
1316  assert(isABI && "Non-ABI compilation is not supported");
1317  if (!isABI)
1318  return "";
1319 
1320  std::stringstream O;
1321  O << "prototype_" << UniqueCallSite << " : .callprototype ";
1322 
1323  if (retTy->getTypeID() == Type::VoidTyID) {
1324  O << "()";
1325  } else {
1326  O << "(";
1327  if (retTy->isFloatingPointTy() || (retTy->isIntegerTy() && !retTy->isIntegerTy(128))) {
1328  unsigned size = 0;
1329  if (auto *ITy = dyn_cast<IntegerType>(retTy)) {
1330  size = ITy->getBitWidth();
1331  } else {
1332  assert(retTy->isFloatingPointTy() &&
1333  "Floating point type expected here");
1334  size = retTy->getPrimitiveSizeInBits();
1335  }
1336  // PTX ABI requires all scalar return values to be at least 32
1337  // bits in size. fp16 normally uses .b16 as its storage type in
1338  // PTX, so its size must be adjusted here, too.
1340 
1341  O << ".param .b" << size << " _";
1342  } else if (isa<PointerType>(retTy)) {
1343  O << ".param .b" << PtrVT.getSizeInBits() << " _";
1344  } else if (retTy->isAggregateType() || retTy->isVectorTy() ||
1345  retTy->isIntegerTy(128)) {
1346  O << ".param .align " << (retAlignment ? retAlignment->value() : 0)
1347  << " .b8 _[" << DL.getTypeAllocSize(retTy) << "]";
1348  } else {
1349  llvm_unreachable("Unknown return type");
1350  }
1351  O << ") ";
1352  }
1353  O << "_ (";
1354 
1355  bool first = true;
1356 
1357  const Function *F = CB.getFunction();
1358  for (unsigned i = 0, e = Args.size(), OIdx = 0; i != e; ++i, ++OIdx) {
1359  Type *Ty = Args[i].Ty;
1360  if (!first) {
1361  O << ", ";
1362  }
1363  first = false;
1364 
1365  if (!Outs[OIdx].Flags.isByVal()) {
1366  if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
1367  unsigned ParamAlign = 0;
1368  const CallInst *CallI = cast<CallInst>(&CB);
1369  // +1 because index 0 is reserved for return type alignment
1370  if (!getAlign(*CallI, i + 1, ParamAlign))
1371  ParamAlign = getFunctionParamOptimizedAlign(F, Ty, DL).value();
1372  O << ".param .align " << ParamAlign << " .b8 ";
1373  O << "_";
1374  O << "[" << DL.getTypeAllocSize(Ty) << "]";
1375  // update the index for Outs
1376  SmallVector<EVT, 16> vtparts;
1377  ComputeValueVTs(*this, DL, Ty, vtparts);
1378  if (unsigned len = vtparts.size())
1379  OIdx += len - 1;
1380  continue;
1381  }
1382  // i8 types in IR will be i16 types in SDAG
1383  assert((getValueType(DL, Ty) == Outs[OIdx].VT ||
1384  (getValueType(DL, Ty) == MVT::i8 && Outs[OIdx].VT == MVT::i16)) &&
1385  "type mismatch between callee prototype and arguments");
1386  // scalar type
1387  unsigned sz = 0;
1388  if (isa<IntegerType>(Ty)) {
1389  sz = cast<IntegerType>(Ty)->getBitWidth();
1390  sz = promoteScalarArgumentSize(sz);
1391  } else if (isa<PointerType>(Ty)) {
1392  sz = PtrVT.getSizeInBits();
1393  } else if (Ty->isHalfTy())
1394  // PTX ABI requires all scalar parameters to be at least 32
1395  // bits in size. fp16 normally uses .b16 as its storage type
1396  // in PTX, so its size must be adjusted here, too.
1397  sz = 32;
1398  else
1399  sz = Ty->getPrimitiveSizeInBits();
1400  O << ".param .b" << sz << " ";
1401  O << "_";
1402  continue;
1403  }
1404 
1405  Align ParamByValAlign = Outs[OIdx].Flags.getNonZeroByValAlign();
1406 
1407  // Try to increase alignment. This code matches logic in LowerCall when
1408  // alignment increase is performed to increase vectorization options.
1409  Type *ETy = Args[i].IndirectType;
1410  Align AlignCandidate = getFunctionParamOptimizedAlign(F, ETy, DL);
1411  ParamByValAlign = std::max(ParamByValAlign, AlignCandidate);
1412 
1413  O << ".param .align " << ParamByValAlign.value() << " .b8 ";
1414  O << "_";
1415  O << "[" << Outs[OIdx].Flags.getByValSize() << "]";
1416  }
1417  O << ");";
1418  return O.str();
1419 }
1420 
1421 Align NVPTXTargetLowering::getArgumentAlignment(SDValue Callee,
1422  const CallBase *CB, Type *Ty,
1423  unsigned Idx,
1424  const DataLayout &DL) const {
1425  if (!CB) {
1426  // CallSite is zero, fallback to ABI type alignment
1427  return DL.getABITypeAlign(Ty);
1428  }
1429 
1430  unsigned Alignment = 0;
1431  const Function *DirectCallee = CB->getCalledFunction();
1432 
1433  if (!DirectCallee) {
1434  // We don't have a direct function symbol, but that may be because of
1435  // constant cast instructions in the call.
1436 
1437  // With bitcast'd call targets, the instruction will be the call
1438  if (const auto *CI = dyn_cast<CallInst>(CB)) {
1439  // Check if we have call alignment metadata
1440  if (getAlign(*CI, Idx, Alignment))
1441  return Align(Alignment);
1442  }
1443  DirectCallee = getMaybeBitcastedCallee(CB);
1444  }
1445 
1446  // Check for function alignment information if we found that the
1447  // ultimate target is a Function
1448  if (DirectCallee) {
1449  if (getAlign(*DirectCallee, Idx, Alignment))
1450  return Align(Alignment);
1451  // If alignment information is not available, fall back to the
1452  // default function param optimized type alignment
1453  return getFunctionParamOptimizedAlign(DirectCallee, Ty, DL);
1454  }
1455 
1456  // Call is indirect, fall back to the ABI type alignment
1457  return DL.getABITypeAlign(Ty);
1458 }
1459 
1461  SmallVectorImpl<SDValue> &InVals) const {
1462  SelectionDAG &DAG = CLI.DAG;
1463  SDLoc dl = CLI.DL;
1465  SmallVectorImpl<SDValue> &OutVals = CLI.OutVals;
1467  SDValue Chain = CLI.Chain;
1468  SDValue Callee = CLI.Callee;
1469  bool &isTailCall = CLI.IsTailCall;
1470  ArgListTy &Args = CLI.getArgs();
1471  Type *RetTy = CLI.RetTy;
1472  const CallBase *CB = CLI.CB;
1473  const DataLayout &DL = DAG.getDataLayout();
1474 
1475  bool isABI = (STI.getSmVersion() >= 20);
1476  assert(isABI && "Non-ABI compilation is not supported");
1477  if (!isABI)
1478  return Chain;
1479 
1480  unsigned UniqueCallSite = GlobalUniqueCallSite.fetch_add(1);
1481  SDValue TempChain = Chain;
1482  Chain = DAG.getCALLSEQ_START(Chain, UniqueCallSite, 0, dl);
1483  SDValue InFlag = Chain.getValue(1);
1484 
1485  unsigned ParamCount = 0;
1486  // Args.size() and Outs.size() need not match.
1487  // Outs.size() will be larger
1488  // * if there is an aggregate argument with multiple fields (each field
1489  // showing up separately in Outs)
1490  // * if there is a vector argument with more than typical vector-length
1491  // elements (generally if more than 4) where each vector element is
1492  // individually present in Outs.
1493  // So a different index should be used for indexing into Outs/OutVals.
1494  // See similar issue in LowerFormalArguments.
1495  unsigned OIdx = 0;
1496  // Declare the .params or .reg need to pass values
1497  // to the function
1498  for (unsigned i = 0, e = Args.size(); i != e; ++i, ++OIdx) {
1499  EVT VT = Outs[OIdx].VT;
1500  Type *Ty = Args[i].Ty;
1501  bool IsByVal = Outs[OIdx].Flags.isByVal();
1502 
1505 
1506  assert((!IsByVal || Args[i].IndirectType) &&
1507  "byval arg must have indirect type");
1508  Type *ETy = (IsByVal ? Args[i].IndirectType : Ty);
1509  ComputePTXValueVTs(*this, DL, ETy, VTs, &Offsets);
1510 
1511  Align ArgAlign;
1512  if (IsByVal) {
1513  // The ByValAlign in the Outs[OIdx].Flags is always set at this point,
1514  // so we don't need to worry whether it's naturally aligned or not.
1515  // See TargetLowering::LowerCallTo().
1516  ArgAlign = Outs[OIdx].Flags.getNonZeroByValAlign();
1517 
1518  // Try to increase alignment to enhance vectorization options.
1519  ArgAlign = std::max(ArgAlign, getFunctionParamOptimizedAlign(
1520  getMaybeBitcastedCallee(CB), ETy, DL));
1521 
1522  // Enforce minumum alignment of 4 to work around ptxas miscompile
1523  // for sm_50+. See corresponding alignment adjustment in
1524  // emitFunctionParamList() for details.
1525  ArgAlign = std::max(ArgAlign, Align(4));
1526  } else {
1527  ArgAlign = getArgumentAlignment(Callee, CB, Ty, ParamCount + 1, DL);
1528  }
1529 
1530  unsigned TypeSize =
1531  (IsByVal ? Outs[OIdx].Flags.getByValSize() : DL.getTypeAllocSize(Ty));
1532  SDVTList DeclareParamVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1533 
1534  bool NeedAlign; // Does argument declaration specify alignment?
1535  if (IsByVal ||
1536  (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128))) {
1537  // declare .param .align <align> .b8 .param<n>[<size>];
1538  SDValue DeclareParamOps[] = {
1539  Chain, DAG.getConstant(ArgAlign.value(), dl, MVT::i32),
1540  DAG.getConstant(ParamCount, dl, MVT::i32),
1541  DAG.getConstant(TypeSize, dl, MVT::i32), InFlag};
1542  Chain = DAG.getNode(NVPTXISD::DeclareParam, dl, DeclareParamVTs,
1543  DeclareParamOps);
1544  NeedAlign = true;
1545  } else {
1546  // declare .param .b<size> .param<n>;
1547  if (VT.isInteger() || VT.isFloatingPoint()) {
1548  // PTX ABI requires integral types to be at least 32 bits in
1549  // size. FP16 is loaded/stored using i16, so it's handled
1550  // here as well.
1552  }
1553  SDValue DeclareScalarParamOps[] = {
1554  Chain, DAG.getConstant(ParamCount, dl, MVT::i32),
1555  DAG.getConstant(TypeSize * 8, dl, MVT::i32),
1556  DAG.getConstant(0, dl, MVT::i32), InFlag};
1557  Chain = DAG.getNode(NVPTXISD::DeclareScalarParam, dl, DeclareParamVTs,
1558  DeclareScalarParamOps);
1559  NeedAlign = false;
1560  }
1561  InFlag = Chain.getValue(1);
1562 
1563  // PTX Interoperability Guide 3.3(A): [Integer] Values shorter
1564  // than 32-bits are sign extended or zero extended, depending on
1565  // whether they are signed or unsigned types. This case applies
1566  // only to scalar parameters and not to aggregate values.
1567  bool ExtendIntegerParam =
1568  Ty->isIntegerTy() && DL.getTypeAllocSizeInBits(Ty) < 32;
1569 
1570  auto VectorInfo = VectorizePTXValueVTs(VTs, Offsets, ArgAlign);
1571  SmallVector<SDValue, 6> StoreOperands;
1572  for (unsigned j = 0, je = VTs.size(); j != je; ++j) {
1573  EVT EltVT = VTs[j];
1574  int CurOffset = Offsets[j];
1575  MaybeAlign PartAlign;
1576  if (NeedAlign)
1577  PartAlign = commonAlignment(ArgAlign, CurOffset);
1578 
1579  // New store.
1580  if (VectorInfo[j] & PVF_FIRST) {
1581  assert(StoreOperands.empty() && "Unfinished preceding store.");
1582  StoreOperands.push_back(Chain);
1583  StoreOperands.push_back(DAG.getConstant(ParamCount, dl, MVT::i32));
1584  StoreOperands.push_back(DAG.getConstant(CurOffset, dl, MVT::i32));
1585  }
1586 
1587  SDValue StVal = OutVals[OIdx];
1588 
1589  MVT PromotedVT;
1590  if (PromoteScalarIntegerPTX(EltVT, &PromotedVT)) {
1591  EltVT = EVT(PromotedVT);
1592  }
1593  if (PromoteScalarIntegerPTX(StVal.getValueType(), &PromotedVT)) {
1595  Outs[OIdx].Flags.isSExt() ? ISD::SIGN_EXTEND : ISD::ZERO_EXTEND;
1596  StVal = DAG.getNode(Ext, dl, PromotedVT, StVal);
1597  }
1598 
1599  if (IsByVal) {
1600  auto PtrVT = getPointerTy(DL);
1601  SDValue srcAddr = DAG.getNode(ISD::ADD, dl, PtrVT, StVal,
1602  DAG.getConstant(CurOffset, dl, PtrVT));
1603  StVal = DAG.getLoad(EltVT, dl, TempChain, srcAddr, MachinePointerInfo(),
1604  PartAlign);
1605  } else if (ExtendIntegerParam) {
1606  assert(VTs.size() == 1 && "Scalar can't have multiple parts.");
1607  // zext/sext to i32
1608  StVal = DAG.getNode(Outs[OIdx].Flags.isSExt() ? ISD::SIGN_EXTEND
1609  : ISD::ZERO_EXTEND,
1610  dl, MVT::i32, StVal);
1611  }
1612 
1613  if (!ExtendIntegerParam && EltVT.getSizeInBits() < 16) {
1614  // Use 16-bit registers for small stores as it's the
1615  // smallest general purpose register size supported by NVPTX.
1616  StVal = DAG.getNode(ISD::ANY_EXTEND, dl, MVT::i16, StVal);
1617  }
1618 
1619  // Record the value to store.
1620  StoreOperands.push_back(StVal);
1621 
1622  if (VectorInfo[j] & PVF_LAST) {
1623  unsigned NumElts = StoreOperands.size() - 3;
1625  switch (NumElts) {
1626  case 1:
1628  break;
1629  case 2:
1631  break;
1632  case 4:
1634  break;
1635  default:
1636  llvm_unreachable("Invalid vector info.");
1637  }
1638 
1639  StoreOperands.push_back(InFlag);
1640 
1641  // Adjust type of the store op if we've extended the scalar
1642  // return value.
1643  EVT TheStoreType = ExtendIntegerParam ? MVT::i32 : EltVT;
1644 
1645  Chain = DAG.getMemIntrinsicNode(
1646  Op, dl, DAG.getVTList(MVT::Other, MVT::Glue), StoreOperands,
1647  TheStoreType, MachinePointerInfo(), PartAlign,
1649  InFlag = Chain.getValue(1);
1650 
1651  // Cleanup.
1652  StoreOperands.clear();
1653  }
1654  if (!IsByVal)
1655  ++OIdx;
1656  }
1657  assert(StoreOperands.empty() && "Unfinished parameter store.");
1658  if (!IsByVal && VTs.size() > 0)
1659  --OIdx;
1660  ++ParamCount;
1661  }
1662 
1663  GlobalAddressSDNode *Func = dyn_cast<GlobalAddressSDNode>(Callee.getNode());
1664  MaybeAlign retAlignment = None;
1665 
1666  // Handle Result
1667  if (Ins.size() > 0) {
1668  SmallVector<EVT, 16> resvtparts;
1669  ComputeValueVTs(*this, DL, RetTy, resvtparts);
1670 
1671  // Declare
1672  // .param .align 16 .b8 retval0[<size-in-bytes>], or
1673  // .param .b<size-in-bits> retval0
1674  unsigned resultsz = DL.getTypeAllocSizeInBits(RetTy);
1675  // Emit ".param .b<size-in-bits> retval0" instead of byte arrays only for
1676  // these three types to match the logic in
1677  // NVPTXAsmPrinter::printReturnValStr and NVPTXTargetLowering::getPrototype.
1678  // Plus, this behavior is consistent with nvcc's.
1679  if (RetTy->isFloatingPointTy() || RetTy->isPointerTy() ||
1680  (RetTy->isIntegerTy() && !RetTy->isIntegerTy(128))) {
1681  resultsz = promoteScalarArgumentSize(resultsz);
1682  SDVTList DeclareRetVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1683  SDValue DeclareRetOps[] = { Chain, DAG.getConstant(1, dl, MVT::i32),
1684  DAG.getConstant(resultsz, dl, MVT::i32),
1685  DAG.getConstant(0, dl, MVT::i32), InFlag };
1686  Chain = DAG.getNode(NVPTXISD::DeclareRet, dl, DeclareRetVTs,
1687  DeclareRetOps);
1688  InFlag = Chain.getValue(1);
1689  } else {
1690  retAlignment = getArgumentAlignment(Callee, CB, RetTy, 0, DL);
1691  assert(retAlignment && "retAlignment is guaranteed to be set");
1692  SDVTList DeclareRetVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1693  SDValue DeclareRetOps[] = {
1694  Chain, DAG.getConstant(retAlignment->value(), dl, MVT::i32),
1695  DAG.getConstant(resultsz / 8, dl, MVT::i32),
1696  DAG.getConstant(0, dl, MVT::i32), InFlag};
1697  Chain = DAG.getNode(NVPTXISD::DeclareRetParam, dl, DeclareRetVTs,
1698  DeclareRetOps);
1699  InFlag = Chain.getValue(1);
1700  }
1701  }
1702 
1703  // Both indirect calls and libcalls have nullptr Func. In order to distinguish
1704  // between them we must rely on the call site value which is valid for
1705  // indirect calls but is always null for libcalls.
1706  bool isIndirectCall = !Func && CB;
1707 
1708  if (isa<ExternalSymbolSDNode>(Callee)) {
1709  Function* CalleeFunc = nullptr;
1710 
1711  // Try to find the callee in the current module.
1712  Callee = DAG.getSymbolFunctionGlobalAddress(Callee, &CalleeFunc);
1713  assert(CalleeFunc != nullptr && "Libcall callee must be set.");
1714 
1715  // Set the "libcall callee" attribute to indicate that the function
1716  // must always have a declaration.
1717  CalleeFunc->addFnAttr("nvptx-libcall-callee", "true");
1718  }
1719 
1720  if (isIndirectCall) {
1721  // This is indirect function call case : PTX requires a prototype of the
1722  // form
1723  // proto_0 : .callprototype(.param .b32 _) _ (.param .b32 _);
1724  // to be emitted, and the label has to used as the last arg of call
1725  // instruction.
1726  // The prototype is embedded in a string and put as the operand for a
1727  // CallPrototype SDNode which will print out to the value of the string.
1728  SDVTList ProtoVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1729  std::string Proto =
1730  getPrototype(DL, RetTy, Args, Outs, retAlignment, *CB, UniqueCallSite);
1731  const char *ProtoStr =
1732  nvTM->getManagedStrPool()->getManagedString(Proto.c_str())->c_str();
1733  SDValue ProtoOps[] = {
1734  Chain, DAG.getTargetExternalSymbol(ProtoStr, MVT::i32), InFlag,
1735  };
1736  Chain = DAG.getNode(NVPTXISD::CallPrototype, dl, ProtoVTs, ProtoOps);
1737  InFlag = Chain.getValue(1);
1738  }
1739  // Op to just print "call"
1740  SDVTList PrintCallVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1741  SDValue PrintCallOps[] = {
1742  Chain, DAG.getConstant((Ins.size() == 0) ? 0 : 1, dl, MVT::i32), InFlag
1743  };
1744  // We model convergent calls as separate opcodes.
1746  if (CLI.IsConvergent)
1749  Chain = DAG.getNode(Opcode, dl, PrintCallVTs, PrintCallOps);
1750  InFlag = Chain.getValue(1);
1751 
1752  // Ops to print out the function name
1753  SDVTList CallVoidVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1754  SDValue CallVoidOps[] = { Chain, Callee, InFlag };
1755  Chain = DAG.getNode(NVPTXISD::CallVoid, dl, CallVoidVTs, CallVoidOps);
1756  InFlag = Chain.getValue(1);
1757 
1758  // Ops to print out the param list
1759  SDVTList CallArgBeginVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1760  SDValue CallArgBeginOps[] = { Chain, InFlag };
1761  Chain = DAG.getNode(NVPTXISD::CallArgBegin, dl, CallArgBeginVTs,
1762  CallArgBeginOps);
1763  InFlag = Chain.getValue(1);
1764 
1765  for (unsigned i = 0, e = ParamCount; i != e; ++i) {
1766  unsigned opcode;
1767  if (i == (e - 1))
1768  opcode = NVPTXISD::LastCallArg;
1769  else
1770  opcode = NVPTXISD::CallArg;
1771  SDVTList CallArgVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1772  SDValue CallArgOps[] = { Chain, DAG.getConstant(1, dl, MVT::i32),
1773  DAG.getConstant(i, dl, MVT::i32), InFlag };
1774  Chain = DAG.getNode(opcode, dl, CallArgVTs, CallArgOps);
1775  InFlag = Chain.getValue(1);
1776  }
1777  SDVTList CallArgEndVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1778  SDValue CallArgEndOps[] = { Chain,
1779  DAG.getConstant(isIndirectCall ? 0 : 1, dl, MVT::i32),
1780  InFlag };
1781  Chain = DAG.getNode(NVPTXISD::CallArgEnd, dl, CallArgEndVTs, CallArgEndOps);
1782  InFlag = Chain.getValue(1);
1783 
1784  if (isIndirectCall) {
1785  SDVTList PrototypeVTs = DAG.getVTList(MVT::Other, MVT::Glue);
1786  SDValue PrototypeOps[] = {
1787  Chain, DAG.getConstant(UniqueCallSite, dl, MVT::i32), InFlag};
1788  Chain = DAG.getNode(NVPTXISD::Prototype, dl, PrototypeVTs, PrototypeOps);
1789  InFlag = Chain.getValue(1);
1790  }
1791 
1792  SmallVector<SDValue, 16> ProxyRegOps;
1793  SmallVector<Optional<MVT>, 16> ProxyRegTruncates;
1794 
1795  // Generate loads from param memory/moves from registers for result
1796  if (Ins.size() > 0) {
1799  ComputePTXValueVTs(*this, DL, RetTy, VTs, &Offsets, 0);
1800  assert(VTs.size() == Ins.size() && "Bad value decomposition");
1801 
1802  Align RetAlign = getArgumentAlignment(Callee, CB, RetTy, 0, DL);
1803  auto VectorInfo = VectorizePTXValueVTs(VTs, Offsets, RetAlign);
1804 
1805  SmallVector<EVT, 6> LoadVTs;
1806  int VecIdx = -1; // Index of the first element of the vector.
1807 
1808  // PTX Interoperability Guide 3.3(A): [Integer] Values shorter than
1809  // 32-bits are sign extended or zero extended, depending on whether
1810  // they are signed or unsigned types.
1811  bool ExtendIntegerRetVal =
1812  RetTy->isIntegerTy() && DL.getTypeAllocSizeInBits(RetTy) < 32;
1813 
1814  for (unsigned i = 0, e = VTs.size(); i != e; ++i) {
1815  bool needTruncate = false;
1816  EVT TheLoadType = VTs[i];
1817  EVT EltType = Ins[i].VT;
1818  Align EltAlign = commonAlignment(RetAlign, Offsets[i]);
1819  MVT PromotedVT;
1820 
1821  if (PromoteScalarIntegerPTX(TheLoadType, &PromotedVT)) {
1822  TheLoadType = EVT(PromotedVT);
1823  EltType = EVT(PromotedVT);
1824  needTruncate = true;
1825  }
1826 
1827  if (ExtendIntegerRetVal) {
1828  TheLoadType = MVT::i32;
1829  EltType = MVT::i32;
1830  needTruncate = true;
1831  } else if (TheLoadType.getSizeInBits() < 16) {
1832  if (VTs[i].isInteger())
1833  needTruncate = true;
1834  EltType = MVT::i16;
1835  }
1836 
1837  // Record index of the very first element of the vector.
1838  if (VectorInfo[i] & PVF_FIRST) {
1839  assert(VecIdx == -1 && LoadVTs.empty() && "Orphaned operand list.");
1840  VecIdx = i;
1841  }
1842 
1843  LoadVTs.push_back(EltType);
1844 
1845  if (VectorInfo[i] & PVF_LAST) {
1846  unsigned NumElts = LoadVTs.size();
1847  LoadVTs.push_back(MVT::Other);
1848  LoadVTs.push_back(MVT::Glue);
1850  switch (NumElts) {
1851  case 1:
1853  break;
1854  case 2:
1856  break;
1857  case 4:
1859  break;
1860  default:
1861  llvm_unreachable("Invalid vector info.");
1862  }
1863 
1864  SDValue LoadOperands[] = {
1865  Chain, DAG.getConstant(1, dl, MVT::i32),
1866  DAG.getConstant(Offsets[VecIdx], dl, MVT::i32), InFlag};
1867  SDValue RetVal = DAG.getMemIntrinsicNode(
1868  Op, dl, DAG.getVTList(LoadVTs), LoadOperands, TheLoadType,
1869  MachinePointerInfo(), EltAlign,
1871 
1872  for (unsigned j = 0; j < NumElts; ++j) {
1873  ProxyRegOps.push_back(RetVal.getValue(j));
1874 
1875  if (needTruncate)
1876  ProxyRegTruncates.push_back(Optional<MVT>(Ins[VecIdx + j].VT));
1877  else
1878  ProxyRegTruncates.push_back(Optional<MVT>());
1879  }
1880 
1881  Chain = RetVal.getValue(NumElts);
1882  InFlag = RetVal.getValue(NumElts + 1);
1883 
1884  // Cleanup
1885  VecIdx = -1;
1886  LoadVTs.clear();
1887  }
1888  }
1889  }
1890 
1891  Chain =
1892  DAG.getCALLSEQ_END(Chain, UniqueCallSite, UniqueCallSite + 1, InFlag, dl);
1893  InFlag = Chain.getValue(1);
1894 
1895  // Append ProxyReg instructions to the chain to make sure that `callseq_end`
1896  // will not get lost. Otherwise, during libcalls expansion, the nodes can become
1897  // dangling.
1898  for (unsigned i = 0; i < ProxyRegOps.size(); ++i) {
1899  SDValue Ret = DAG.getNode(
1900  NVPTXISD::ProxyReg, dl,
1901  DAG.getVTList(ProxyRegOps[i].getSimpleValueType(), MVT::Other, MVT::Glue),
1902  { Chain, ProxyRegOps[i], InFlag }
1903  );
1904 
1905  Chain = Ret.getValue(1);
1906  InFlag = Ret.getValue(2);
1907 
1908  if (ProxyRegTruncates[i]) {
1909  Ret = DAG.getNode(ISD::TRUNCATE, dl, ProxyRegTruncates[i].value(), Ret);
1910  }
1911 
1912  InVals.push_back(Ret);
1913  }
1914 
1915  // set isTailCall to false for now, until we figure out how to express
1916  // tail call optimization in PTX
1917  isTailCall = false;
1918  return Chain;
1919 }
1920 
1921 // By default CONCAT_VECTORS is lowered by ExpandVectorBuildThroughStack()
1922 // (see LegalizeDAG.cpp). This is slow and uses local memory.
1923 // We use extract/insert/build vector just as what LegalizeOp() does in llvm 2.5
1924 SDValue
1925 NVPTXTargetLowering::LowerCONCAT_VECTORS(SDValue Op, SelectionDAG &DAG) const {
1926  SDNode *Node = Op.getNode();
1927  SDLoc dl(Node);
1929  unsigned NumOperands = Node->getNumOperands();
1930  for (unsigned i = 0; i < NumOperands; ++i) {
1931  SDValue SubOp = Node->getOperand(i);
1932  EVT VVT = SubOp.getNode()->getValueType(0);
1933  EVT EltVT = VVT.getVectorElementType();
1934  unsigned NumSubElem = VVT.getVectorNumElements();
1935  for (unsigned j = 0; j < NumSubElem; ++j) {
1936  Ops.push_back(DAG.getNode(ISD::EXTRACT_VECTOR_ELT, dl, EltVT, SubOp,
1937  DAG.getIntPtrConstant(j, dl)));
1938  }
1939  }
1940  return DAG.getBuildVector(Node->getValueType(0), dl, Ops);
1941 }
1942 
1943 // We can init constant f16x2 with a single .b32 move. Normally it
1944 // would get lowered as two constant loads and vector-packing move.
1945 // mov.b16 %h1, 0x4000;
1946 // mov.b16 %h2, 0x3C00;
1947 // mov.b32 %hh2, {%h2, %h1};
1948 // Instead we want just a constant move:
1949 // mov.b32 %hh2, 0x40003C00
1950 //
1951 // This results in better SASS code with CUDA 7.x. Ptxas in CUDA 8.0
1952 // generates good SASS in both cases.
1953 SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op,
1954  SelectionDAG &DAG) const {
1955  //return Op;
1956  if (!(Op->getValueType(0) == MVT::v2f16 &&
1957  isa<ConstantFPSDNode>(Op->getOperand(0)) &&
1958  isa<ConstantFPSDNode>(Op->getOperand(1))))
1959  return Op;
1960 
1961  APInt E0 =
1962  cast<ConstantFPSDNode>(Op->getOperand(0))->getValueAPF().bitcastToAPInt();
1963  APInt E1 =
1964  cast<ConstantFPSDNode>(Op->getOperand(1))->getValueAPF().bitcastToAPInt();
1965  SDValue Const =
1966  DAG.getConstant(E1.zext(32).shl(16) | E0.zext(32), SDLoc(Op), MVT::i32);
1967  return DAG.getNode(ISD::BITCAST, SDLoc(Op), MVT::v2f16, Const);
1968 }
1969 
1970 SDValue NVPTXTargetLowering::LowerEXTRACT_VECTOR_ELT(SDValue Op,
1971  SelectionDAG &DAG) const {
1972  SDValue Index = Op->getOperand(1);
1973  // Constant index will be matched by tablegen.
1974  if (isa<ConstantSDNode>(Index.getNode()))
1975  return Op;
1976 
1977  // Extract individual elements and select one of them.
1978  SDValue Vector = Op->getOperand(0);
1979  EVT VectorVT = Vector.getValueType();
1980  assert(VectorVT == MVT::v2f16 && "Unexpected vector type.");
1981  EVT EltVT = VectorVT.getVectorElementType();
1982 
1983  SDLoc dl(Op.getNode());
1984  SDValue E0 = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, dl, EltVT, Vector,
1985  DAG.getIntPtrConstant(0, dl));
1986  SDValue E1 = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, dl, EltVT, Vector,
1987  DAG.getIntPtrConstant(1, dl));
1988  return DAG.getSelectCC(dl, Index, DAG.getIntPtrConstant(0, dl), E0, E1,
1990 }
1991 
1992 /// LowerShiftRightParts - Lower SRL_PARTS, SRA_PARTS, which
1993 /// 1) returns two i32 values and take a 2 x i32 value to shift plus a shift
1994 /// amount, or
1995 /// 2) returns two i64 values and take a 2 x i64 value to shift plus a shift
1996 /// amount.
1997 SDValue NVPTXTargetLowering::LowerShiftRightParts(SDValue Op,
1998  SelectionDAG &DAG) const {
1999  assert(Op.getNumOperands() == 3 && "Not a double-shift!");
2000  assert(Op.getOpcode() == ISD::SRA_PARTS || Op.getOpcode() == ISD::SRL_PARTS);
2001 
2002  EVT VT = Op.getValueType();
2003  unsigned VTBits = VT.getSizeInBits();
2004  SDLoc dl(Op);
2005  SDValue ShOpLo = Op.getOperand(0);
2006  SDValue ShOpHi = Op.getOperand(1);
2007  SDValue ShAmt = Op.getOperand(2);
2008  unsigned Opc = (Op.getOpcode() == ISD::SRA_PARTS) ? ISD::SRA : ISD::SRL;
2009 
2010  if (VTBits == 32 && STI.getSmVersion() >= 35) {
2011  // For 32bit and sm35, we can use the funnel shift 'shf' instruction.
2012  // {dHi, dLo} = {aHi, aLo} >> Amt
2013  // dHi = aHi >> Amt
2014  // dLo = shf.r.clamp aLo, aHi, Amt
2015 
2016  SDValue Hi = DAG.getNode(Opc, dl, VT, ShOpHi, ShAmt);
2017  SDValue Lo = DAG.getNode(NVPTXISD::FUN_SHFR_CLAMP, dl, VT, ShOpLo, ShOpHi,
2018  ShAmt);
2019 
2020  SDValue Ops[2] = { Lo, Hi };
2021  return DAG.getMergeValues(Ops, dl);
2022  }
2023  else {
2024  // {dHi, dLo} = {aHi, aLo} >> Amt
2025  // - if (Amt>=size) then
2026  // dLo = aHi >> (Amt-size)
2027  // dHi = aHi >> Amt (this is either all 0 or all 1)
2028  // else
2029  // dLo = (aLo >>logic Amt) | (aHi << (size-Amt))
2030  // dHi = aHi >> Amt
2031 
2032  SDValue RevShAmt = DAG.getNode(ISD::SUB, dl, MVT::i32,
2033  DAG.getConstant(VTBits, dl, MVT::i32),
2034  ShAmt);
2035  SDValue Tmp1 = DAG.getNode(ISD::SRL, dl, VT, ShOpLo, ShAmt);
2036  SDValue ExtraShAmt = DAG.getNode(ISD::SUB, dl, MVT::i32, ShAmt,
2037  DAG.getConstant(VTBits, dl, MVT::i32));
2038  SDValue Tmp2 = DAG.getNode(ISD::SHL, dl, VT, ShOpHi, RevShAmt);
2039  SDValue FalseVal = DAG.getNode(ISD::OR, dl, VT, Tmp1, Tmp2);
2040  SDValue TrueVal = DAG.getNode(Opc, dl, VT, ShOpHi, ExtraShAmt);
2041 
2042  SDValue Cmp = DAG.getSetCC(dl, MVT::i1, ShAmt,
2043  DAG.getConstant(VTBits, dl, MVT::i32),
2044  ISD::SETGE);
2045  SDValue Hi = DAG.getNode(Opc, dl, VT, ShOpHi, ShAmt);
2046  SDValue Lo = DAG.getNode(ISD::SELECT, dl, VT, Cmp, TrueVal, FalseVal);
2047 
2048  SDValue Ops[2] = { Lo, Hi };
2049  return DAG.getMergeValues(Ops, dl);
2050  }
2051 }
2052 
2053 /// LowerShiftLeftParts - Lower SHL_PARTS, which
2054 /// 1) returns two i32 values and take a 2 x i32 value to shift plus a shift
2055 /// amount, or
2056 /// 2) returns two i64 values and take a 2 x i64 value to shift plus a shift
2057 /// amount.
2058 SDValue NVPTXTargetLowering::LowerShiftLeftParts(SDValue Op,
2059  SelectionDAG &DAG) const {
2060  assert(Op.getNumOperands() == 3 && "Not a double-shift!");
2061  assert(Op.getOpcode() == ISD::SHL_PARTS);
2062 
2063  EVT VT = Op.getValueType();
2064  unsigned VTBits = VT.getSizeInBits();
2065  SDLoc dl(Op);
2066  SDValue ShOpLo = Op.getOperand(0);
2067  SDValue ShOpHi = Op.getOperand(1);
2068  SDValue ShAmt = Op.getOperand(2);
2069 
2070  if (VTBits == 32 && STI.getSmVersion() >= 35) {
2071  // For 32bit and sm35, we can use the funnel shift 'shf' instruction.
2072  // {dHi, dLo} = {aHi, aLo} << Amt
2073  // dHi = shf.l.clamp aLo, aHi, Amt
2074  // dLo = aLo << Amt
2075 
2076  SDValue Hi = DAG.getNode(NVPTXISD::FUN_SHFL_CLAMP, dl, VT, ShOpLo, ShOpHi,
2077  ShAmt);
2078  SDValue Lo = DAG.getNode(ISD::SHL, dl, VT, ShOpLo, ShAmt);
2079 
2080  SDValue Ops[2] = { Lo, Hi };
2081  return DAG.getMergeValues(Ops, dl);
2082  }
2083  else {
2084  // {dHi, dLo} = {aHi, aLo} << Amt
2085  // - if (Amt>=size) then
2086  // dLo = aLo << Amt (all 0)
2087  // dLo = aLo << (Amt-size)
2088  // else
2089  // dLo = aLo << Amt
2090  // dHi = (aHi << Amt) | (aLo >> (size-Amt))
2091 
2092  SDValue RevShAmt = DAG.getNode(ISD::SUB, dl, MVT::i32,
2093  DAG.getConstant(VTBits, dl, MVT::i32),
2094  ShAmt);
2095  SDValue Tmp1 = DAG.getNode(ISD::SHL, dl, VT, ShOpHi, ShAmt);
2096  SDValue ExtraShAmt = DAG.getNode(ISD::SUB, dl, MVT::i32, ShAmt,
2097  DAG.getConstant(VTBits, dl, MVT::i32));
2098  SDValue Tmp2 = DAG.getNode(ISD::SRL, dl, VT, ShOpLo, RevShAmt);
2099  SDValue FalseVal = DAG.getNode(ISD::OR, dl, VT, Tmp1, Tmp2);
2100  SDValue TrueVal = DAG.getNode(ISD::SHL, dl, VT, ShOpLo, ExtraShAmt);
2101 
2102  SDValue Cmp = DAG.getSetCC(dl, MVT::i1, ShAmt,
2103  DAG.getConstant(VTBits, dl, MVT::i32),
2104  ISD::SETGE);
2105  SDValue Lo = DAG.getNode(ISD::SHL, dl, VT, ShOpLo, ShAmt);
2106  SDValue Hi = DAG.getNode(ISD::SELECT, dl, VT, Cmp, TrueVal, FalseVal);
2107 
2108  SDValue Ops[2] = { Lo, Hi };
2109  return DAG.getMergeValues(Ops, dl);
2110  }
2111 }
2112 
2113 SDValue NVPTXTargetLowering::LowerFROUND(SDValue Op, SelectionDAG &DAG) const {
2114  EVT VT = Op.getValueType();
2115 
2116  if (VT == MVT::f32)
2117  return LowerFROUND32(Op, DAG);
2118 
2119  if (VT == MVT::f64)
2120  return LowerFROUND64(Op, DAG);
2121 
2122  llvm_unreachable("unhandled type");
2123 }
2124 
2125 // This is the the rounding method used in CUDA libdevice in C like code:
2126 // float roundf(float A)
2127 // {
2128 // float RoundedA = (float) (int) ( A > 0 ? (A + 0.5f) : (A - 0.5f));
2129 // RoundedA = abs(A) > 0x1.0p23 ? A : RoundedA;
2130 // return abs(A) < 0.5 ? (float)(int)A : RoundedA;
2131 // }
2132 SDValue NVPTXTargetLowering::LowerFROUND32(SDValue Op,
2133  SelectionDAG &DAG) const {
2134  SDLoc SL(Op);
2135  SDValue A = Op.getOperand(0);
2136  EVT VT = Op.getValueType();
2137 
2138  SDValue AbsA = DAG.getNode(ISD::FABS, SL, VT, A);
2139 
2140  // RoundedA = (float) (int) ( A > 0 ? (A + 0.5f) : (A - 0.5f))
2141  SDValue Bitcast = DAG.getNode(ISD::BITCAST, SL, MVT::i32, A);
2142  const int SignBitMask = 0x80000000;
2143  SDValue Sign = DAG.getNode(ISD::AND, SL, MVT::i32, Bitcast,
2144  DAG.getConstant(SignBitMask, SL, MVT::i32));
2145  const int PointFiveInBits = 0x3F000000;
2146  SDValue PointFiveWithSignRaw =
2147  DAG.getNode(ISD::OR, SL, MVT::i32, Sign,
2148  DAG.getConstant(PointFiveInBits, SL, MVT::i32));
2149  SDValue PointFiveWithSign =
2150  DAG.getNode(ISD::BITCAST, SL, VT, PointFiveWithSignRaw);
2151  SDValue AdjustedA = DAG.getNode(ISD::FADD, SL, VT, A, PointFiveWithSign);
2152  SDValue RoundedA = DAG.getNode(ISD::FTRUNC, SL, VT, AdjustedA);
2153 
2154  // RoundedA = abs(A) > 0x1.0p23 ? A : RoundedA;
2155  EVT SetCCVT = getSetCCResultType(DAG.getDataLayout(), *DAG.getContext(), VT);
2156  SDValue IsLarge =
2157  DAG.getSetCC(SL, SetCCVT, AbsA, DAG.getConstantFP(pow(2.0, 23.0), SL, VT),
2158  ISD::SETOGT);
2159  RoundedA = DAG.getNode(ISD::SELECT, SL, VT, IsLarge, A, RoundedA);
2160 
2161  // return abs(A) < 0.5 ? (float)(int)A : RoundedA;
2162  SDValue IsSmall =DAG.getSetCC(SL, SetCCVT, AbsA,
2163  DAG.getConstantFP(0.5, SL, VT), ISD::SETOLT);
2164  SDValue RoundedAForSmallA = DAG.getNode(ISD::FTRUNC, SL, VT, A);
2165  return DAG.getNode(ISD::SELECT, SL, VT, IsSmall, RoundedAForSmallA, RoundedA);
2166 }
2167 
2168 // The implementation of round(double) is similar to that of round(float) in
2169 // that they both separate the value range into three regions and use a method
2170 // specific to the region to round the values. However, round(double) first
2171 // calculates the round of the absolute value and then adds the sign back while
2172 // round(float) directly rounds the value with sign.
2173 SDValue NVPTXTargetLowering::LowerFROUND64(SDValue Op,
2174  SelectionDAG &DAG) const {
2175  SDLoc SL(Op);
2176  SDValue A = Op.getOperand(0);
2177  EVT VT = Op.getValueType();
2178 
2179  SDValue AbsA = DAG.getNode(ISD::FABS, SL, VT, A);
2180 
2181  // double RoundedA = (double) (int) (abs(A) + 0.5f);
2182  SDValue AdjustedA = DAG.getNode(ISD::FADD, SL, VT, AbsA,
2183  DAG.getConstantFP(0.5, SL, VT));
2184  SDValue RoundedA = DAG.getNode(ISD::FTRUNC, SL, VT, AdjustedA);
2185 
2186  // RoundedA = abs(A) < 0.5 ? (double)0 : RoundedA;
2187  EVT SetCCVT = getSetCCResultType(DAG.getDataLayout(), *DAG.getContext(), VT);
2188  SDValue IsSmall =DAG.getSetCC(SL, SetCCVT, AbsA,
2189  DAG.getConstantFP(0.5, SL, VT), ISD::SETOLT);
2190  RoundedA = DAG.getNode(ISD::SELECT, SL, VT, IsSmall,
2191  DAG.getConstantFP(0, SL, VT),
2192  RoundedA);
2193 
2194  // Add sign to rounded_A
2195  RoundedA = DAG.getNode(ISD::FCOPYSIGN, SL, VT, RoundedA, A);
2196  DAG.getNode(ISD::FTRUNC, SL, VT, A);
2197 
2198  // RoundedA = abs(A) > 0x1.0p52 ? A : RoundedA;
2199  SDValue IsLarge =
2200  DAG.getSetCC(SL, SetCCVT, AbsA, DAG.getConstantFP(pow(2.0, 52.0), SL, VT),
2201  ISD::SETOGT);
2202  return DAG.getNode(ISD::SELECT, SL, VT, IsLarge, A, RoundedA);
2203 }
2204 
2205 
2206 
2207 SDValue
2209  switch (Op.getOpcode()) {
2210  case ISD::RETURNADDR:
2211  return SDValue();
2212  case ISD::FRAMEADDR:
2213  return SDValue();
2214  case ISD::GlobalAddress:
2215  return LowerGlobalAddress(Op, DAG);
2217  return Op;
2218  case ISD::BUILD_VECTOR:
2219  return LowerBUILD_VECTOR(Op, DAG);
2221  return Op;
2223  return LowerEXTRACT_VECTOR_ELT(Op, DAG);
2224  case ISD::CONCAT_VECTORS:
2225  return LowerCONCAT_VECTORS(Op, DAG);
2226  case ISD::STORE:
2227  return LowerSTORE(Op, DAG);
2228  case ISD::LOAD:
2229  return LowerLOAD(Op, DAG);
2230  case ISD::SHL_PARTS:
2231  return LowerShiftLeftParts(Op, DAG);
2232  case ISD::SRA_PARTS:
2233  case ISD::SRL_PARTS:
2234  return LowerShiftRightParts(Op, DAG);
2235  case ISD::SELECT:
2236  return LowerSelect(Op, DAG);
2237  case ISD::FROUND:
2238  return LowerFROUND(Op, DAG);
2239  default:
2240  llvm_unreachable("Custom lowering not defined for operation");
2241  }
2242 }
2243 
2244 SDValue NVPTXTargetLowering::LowerSelect(SDValue Op, SelectionDAG &DAG) const {
2245  SDValue Op0 = Op->getOperand(0);
2246  SDValue Op1 = Op->getOperand(1);
2247  SDValue Op2 = Op->getOperand(2);
2248  SDLoc DL(Op.getNode());
2249 
2250  assert(Op.getValueType() == MVT::i1 && "Custom lowering enabled only for i1");
2251 
2252  Op1 = DAG.getNode(ISD::ANY_EXTEND, DL, MVT::i32, Op1);
2253  Op2 = DAG.getNode(ISD::ANY_EXTEND, DL, MVT::i32, Op2);
2254  SDValue Select = DAG.getNode(ISD::SELECT, DL, MVT::i32, Op0, Op1, Op2);
2255  SDValue Trunc = DAG.getNode(ISD::TRUNCATE, DL, MVT::i1, Select);
2256 
2257  return Trunc;
2258 }
2259 
2260 SDValue NVPTXTargetLowering::LowerLOAD(SDValue Op, SelectionDAG &DAG) const {
2261  if (Op.getValueType() == MVT::i1)
2262  return LowerLOADi1(Op, DAG);
2263 
2264  // v2f16 is legal, so we can't rely on legalizer to handle unaligned
2265  // loads and have to handle it here.
2266  if (Op.getValueType() == MVT::v2f16) {
2267  LoadSDNode *Load = cast<LoadSDNode>(Op);
2268  EVT MemVT = Load->getMemoryVT();
2270  MemVT, *Load->getMemOperand())) {
2271  SDValue Ops[2];
2272  std::tie(Ops[0], Ops[1]) = expandUnalignedLoad(Load, DAG);
2273  return DAG.getMergeValues(Ops, SDLoc(Op));
2274  }
2275  }
2276 
2277  return SDValue();
2278 }
2279 
2280 // v = ld i1* addr
2281 // =>
2282 // v1 = ld i8* addr (-> i16)
2283 // v = trunc i16 to i1
2284 SDValue NVPTXTargetLowering::LowerLOADi1(SDValue Op, SelectionDAG &DAG) const {
2285  SDNode *Node = Op.getNode();
2286  LoadSDNode *LD = cast<LoadSDNode>(Node);
2287  SDLoc dl(Node);
2288  assert(LD->getExtensionType() == ISD::NON_EXTLOAD);
2289  assert(Node->getValueType(0) == MVT::i1 &&
2290  "Custom lowering for i1 load only");
2291  SDValue newLD = DAG.getLoad(MVT::i16, dl, LD->getChain(), LD->getBasePtr(),
2292  LD->getPointerInfo(), LD->getAlign(),
2293  LD->getMemOperand()->getFlags());
2294  SDValue result = DAG.getNode(ISD::TRUNCATE, dl, MVT::i1, newLD);
2295  // The legalizer (the caller) is expecting two values from the legalized
2296  // load, so we build a MergeValues node for it. See ExpandUnalignedLoad()
2297  // in LegalizeDAG.cpp which also uses MergeValues.
2298  SDValue Ops[] = { result, LD->getChain() };
2299  return DAG.getMergeValues(Ops, dl);
2300 }
2301 
2302 SDValue NVPTXTargetLowering::LowerSTORE(SDValue Op, SelectionDAG &DAG) const {
2303  StoreSDNode *Store = cast<StoreSDNode>(Op);
2304  EVT VT = Store->getMemoryVT();
2305 
2306  if (VT == MVT::i1)
2307  return LowerSTOREi1(Op, DAG);
2308 
2309  // v2f16 is legal, so we can't rely on legalizer to handle unaligned
2310  // stores and have to handle it here.
2311  if (VT == MVT::v2f16 &&
2313  VT, *Store->getMemOperand()))
2314  return expandUnalignedStore(Store, DAG);
2315 
2316  if (VT.isVector())
2317  return LowerSTOREVector(Op, DAG);
2318 
2319  return SDValue();
2320 }
2321 
2322 SDValue
2323 NVPTXTargetLowering::LowerSTOREVector(SDValue Op, SelectionDAG &DAG) const {
2324  SDNode *N = Op.getNode();
2325  SDValue Val = N->getOperand(1);
2326  SDLoc DL(N);
2327  EVT ValVT = Val.getValueType();
2328 
2329  if (ValVT.isVector()) {
2330  // We only handle "native" vector sizes for now, e.g. <4 x double> is not
2331  // legal. We can (and should) split that into 2 stores of <2 x double> here
2332  // but I'm leaving that as a TODO for now.
2333  if (!ValVT.isSimple())
2334  return SDValue();
2335  switch (ValVT.getSimpleVT().SimpleTy) {
2336  default:
2337  return SDValue();
2338  case MVT::v2i8:
2339  case MVT::v2i16:
2340  case MVT::v2i32:
2341  case MVT::v2i64:
2342  case MVT::v2f16:
2343  case MVT::v2bf16:
2344  case MVT::v2f32:
2345  case MVT::v2f64:
2346  case MVT::v4i8:
2347  case MVT::v4i16:
2348  case MVT::v4i32:
2349  case MVT::v4f16:
2350  case MVT::v4bf16:
2351  case MVT::v4f32:
2352  case MVT::v8f16: // <4 x f16x2>
2353  case MVT::v8bf16: // <4 x bf16x2>
2354  // This is a "native" vector type
2355  break;
2356  }
2357 
2358  MemSDNode *MemSD = cast<MemSDNode>(N);
2359  const DataLayout &TD = DAG.getDataLayout();
2360 
2361  Align Alignment = MemSD->getAlign();
2362  Align PrefAlign =
2363  TD.getPrefTypeAlign(ValVT.getTypeForEVT(*DAG.getContext()));
2364  if (Alignment < PrefAlign) {
2365  // This store is not sufficiently aligned, so bail out and let this vector
2366  // store be scalarized. Note that we may still be able to emit smaller
2367  // vector stores. For example, if we are storing a <4 x float> with an
2368  // alignment of 8, this check will fail but the legalizer will try again
2369  // with 2 x <2 x float>, which will succeed with an alignment of 8.
2370  return SDValue();
2371  }
2372 
2373  unsigned Opcode = 0;
2374  EVT EltVT = ValVT.getVectorElementType();
2375  unsigned NumElts = ValVT.getVectorNumElements();
2376 
2377  // Since StoreV2 is a target node, we cannot rely on DAG type legalization.
2378  // Therefore, we must ensure the type is legal. For i1 and i8, we set the
2379  // stored type to i16 and propagate the "real" type as the memory type.
2380  bool NeedExt = false;
2381  if (EltVT.getSizeInBits() < 16)
2382  NeedExt = true;
2383 
2384  bool StoreF16x2 = false;
2385  switch (NumElts) {
2386  default:
2387  return SDValue();
2388  case 2:
2389  Opcode = NVPTXISD::StoreV2;
2390  break;
2391  case 4:
2392  Opcode = NVPTXISD::StoreV4;
2393  break;
2394  case 8:
2395  // v8f16 is a special case. PTX doesn't have st.v8.f16
2396  // instruction. Instead, we split the vector into v2f16 chunks and
2397  // store them with st.v4.b32.
2398  assert((EltVT == MVT::f16 || EltVT == MVT::bf16) &&
2399  "Wrong type for the vector.");
2400  Opcode = NVPTXISD::StoreV4;
2401  StoreF16x2 = true;
2402  break;
2403  }
2404 
2406 
2407  // First is the chain
2408  Ops.push_back(N->getOperand(0));
2409 
2410  if (StoreF16x2) {
2411  // Combine f16,f16 -> v2f16
2412  NumElts /= 2;
2413  for (unsigned i = 0; i < NumElts; ++i) {
2415  DAG.getIntPtrConstant(i * 2, DL));
2417  DAG.getIntPtrConstant(i * 2 + 1, DL));
2418  SDValue V2 = DAG.getNode(ISD::BUILD_VECTOR, DL, MVT::v2f16, E0, E1);
2419  Ops.push_back(V2);
2420  }
2421  } else {
2422  // Then the split values
2423  for (unsigned i = 0; i < NumElts; ++i) {
2424  SDValue ExtVal = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, EltVT, Val,
2425  DAG.getIntPtrConstant(i, DL));
2426  if (NeedExt)
2427  ExtVal = DAG.getNode(ISD::ANY_EXTEND, DL, MVT::i16, ExtVal);
2428  Ops.push_back(ExtVal);
2429  }
2430  }
2431 
2432  // Then any remaining arguments
2433  Ops.append(N->op_begin() + 2, N->op_end());
2434 
2435  SDValue NewSt =
2436  DAG.getMemIntrinsicNode(Opcode, DL, DAG.getVTList(MVT::Other), Ops,
2437  MemSD->getMemoryVT(), MemSD->getMemOperand());
2438 
2439  // return DCI.CombineTo(N, NewSt, true);
2440  return NewSt;
2441  }
2442 
2443  return SDValue();
2444 }
2445 
2446 // st i1 v, addr
2447 // =>
2448 // v1 = zxt v to i16
2449 // st.u8 i16, addr
2450 SDValue NVPTXTargetLowering::LowerSTOREi1(SDValue Op, SelectionDAG &DAG) const {
2451  SDNode *Node = Op.getNode();
2452  SDLoc dl(Node);
2453  StoreSDNode *ST = cast<StoreSDNode>(Node);
2454  SDValue Tmp1 = ST->getChain();
2455  SDValue Tmp2 = ST->getBasePtr();
2456  SDValue Tmp3 = ST->getValue();
2457  assert(Tmp3.getValueType() == MVT::i1 && "Custom lowering for i1 store only");
2458  Tmp3 = DAG.getNode(ISD::ZERO_EXTEND, dl, MVT::i16, Tmp3);
2459  SDValue Result =
2460  DAG.getTruncStore(Tmp1, dl, Tmp3, Tmp2, ST->getPointerInfo(), MVT::i8,
2461  ST->getAlign(), ST->getMemOperand()->getFlags());
2462  return Result;
2463 }
2464 
2465 SDValue
2466 NVPTXTargetLowering::getParamSymbol(SelectionDAG &DAG, int idx, EVT v) const {
2467  std::string ParamSym;
2468  raw_string_ostream ParamStr(ParamSym);
2469 
2470  ParamStr << DAG.getMachineFunction().getName() << "_param_" << idx;
2471  ParamStr.flush();
2472 
2473  std::string *SavedStr =
2474  nvTM->getManagedStrPool()->getManagedString(ParamSym.c_str());
2475  return DAG.getTargetExternalSymbol(SavedStr->c_str(), v);
2476 }
2477 
2479  SDValue Chain, CallingConv::ID CallConv, bool isVarArg,
2480  const SmallVectorImpl<ISD::InputArg> &Ins, const SDLoc &dl,
2481  SelectionDAG &DAG, SmallVectorImpl<SDValue> &InVals) const {
2482  MachineFunction &MF = DAG.getMachineFunction();
2483  const DataLayout &DL = DAG.getDataLayout();
2484  auto PtrVT = getPointerTy(DAG.getDataLayout());
2485 
2486  const Function *F = &MF.getFunction();
2487  const AttributeList &PAL = F->getAttributes();
2488  const TargetLowering *TLI = STI.getTargetLowering();
2489 
2490  SDValue Root = DAG.getRoot();
2491  std::vector<SDValue> OutChains;
2492 
2493  bool isABI = (STI.getSmVersion() >= 20);
2494  assert(isABI && "Non-ABI compilation is not supported");
2495  if (!isABI)
2496  return Chain;
2497 
2498  std::vector<Type *> argTypes;
2499  std::vector<const Argument *> theArgs;
2500  for (const Argument &I : F->args()) {
2501  theArgs.push_back(&I);
2502  argTypes.push_back(I.getType());
2503  }
2504  // argTypes.size() (or theArgs.size()) and Ins.size() need not match.
2505  // Ins.size() will be larger
2506  // * if there is an aggregate argument with multiple fields (each field
2507  // showing up separately in Ins)
2508  // * if there is a vector argument with more than typical vector-length
2509  // elements (generally if more than 4) where each vector element is
2510  // individually present in Ins.
2511  // So a different index should be used for indexing into Ins.
2512  // See similar issue in LowerCall.
2513  unsigned InsIdx = 0;
2514 
2515  int idx = 0;
2516  for (unsigned i = 0, e = theArgs.size(); i != e; ++i, ++idx, ++InsIdx) {
2517  Type *Ty = argTypes[i];
2518 
2519  if (theArgs[i]->use_empty()) {
2520  // argument is dead
2521  if (Ty->isAggregateType() || Ty->isIntegerTy(128)) {
2522  SmallVector<EVT, 16> vtparts;
2523 
2524  ComputePTXValueVTs(*this, DAG.getDataLayout(), Ty, vtparts);
2525  assert(vtparts.size() > 0 && "empty aggregate type not expected");
2526  for (unsigned parti = 0, parte = vtparts.size(); parti != parte;
2527  ++parti) {
2528  InVals.push_back(DAG.getNode(ISD::UNDEF, dl, Ins[InsIdx].VT));
2529  ++InsIdx;
2530  }
2531  if (vtparts.size() > 0)
2532  --InsIdx;
2533  continue;
2534  }
2535  if (Ty->isVectorTy()) {
2536  EVT ObjectVT = getValueType(DL, Ty);
2537  unsigned NumRegs = TLI->getNumRegisters(F->getContext(), ObjectVT);
2538  for (unsigned parti = 0; parti < NumRegs; ++parti) {
2539  InVals.push_back(DAG.getNode(ISD::UNDEF, dl, Ins[InsIdx].VT));
2540  ++InsIdx;
2541  }
2542  if (NumRegs > 0)
2543  --InsIdx;
2544  continue;
2545  }
2546  InVals.push_back(DAG.getNode(ISD::UNDEF, dl, Ins[InsIdx].VT));
2547  continue;
2548  }
2549 
2550  // In the following cases, assign a node order of "idx+1"
2551  // to newly created nodes. The SDNodes for params have to
2552  // appear in the same order as their order of appearance
2553  // in the original function. "idx+1" holds that order.
2554  if (!PAL.hasParamAttr(i, Attribute::ByVal)) {
2555  bool aggregateIsPacked = false;
2556  if (StructType *STy = dyn_cast<StructType>(Ty))
2557  aggregateIsPacked = STy->isPacked();
2558 
2561  ComputePTXValueVTs(*this, DL, Ty, VTs, &Offsets, 0);
2562  assert(VTs.size() > 0 && "Unexpected empty type.");
2563  auto VectorInfo =
2564  VectorizePTXValueVTs(VTs, Offsets, DL.getABITypeAlign(Ty));
2565 
2566  SDValue Arg = getParamSymbol(DAG, idx, PtrVT);
2567  int VecIdx = -1; // Index of the first element of the current vector.
2568  for (unsigned parti = 0, parte = VTs.size(); parti != parte; ++parti) {
2569  if (VectorInfo[parti] & PVF_FIRST) {
2570  assert(VecIdx == -1 && "Orphaned vector.");
2571  VecIdx = parti;
2572  }
2573 
2574  // That's the last element of this store op.
2575  if (VectorInfo[parti] & PVF_LAST) {
2576  unsigned NumElts = parti - VecIdx + 1;
2577  EVT EltVT = VTs[parti];
2578  // i1 is loaded/stored as i8.
2579  EVT LoadVT = EltVT;
2580  if (EltVT == MVT::i1)
2581  LoadVT = MVT::i8;
2582  else if (EltVT == MVT::v2f16)
2583  // getLoad needs a vector type, but it can't handle
2584  // vectors which contain v2f16 elements. So we must load
2585  // using i32 here and then bitcast back.
2586  LoadVT = MVT::i32;
2587 
2588  EVT VecVT = EVT::getVectorVT(F->getContext(), LoadVT, NumElts);
2589  SDValue VecAddr =
2590  DAG.getNode(ISD::ADD, dl, PtrVT, Arg,
2591  DAG.getConstant(Offsets[VecIdx], dl, PtrVT));
2593  EltVT.getTypeForEVT(F->getContext()), ADDRESS_SPACE_PARAM));
2594  SDValue P =
2595  DAG.getLoad(VecVT, dl, Root, VecAddr,
2596  MachinePointerInfo(srcValue), aggregateIsPacked,
2599  if (P.getNode())
2600  P.getNode()->setIROrder(idx + 1);
2601  for (unsigned j = 0; j < NumElts; ++j) {
2602  SDValue Elt = DAG.getNode(ISD::EXTRACT_VECTOR_ELT, dl, LoadVT, P,
2603  DAG.getIntPtrConstant(j, dl));
2604  // We've loaded i1 as an i8 and now must truncate it back to i1
2605  if (EltVT == MVT::i1)
2606  Elt = DAG.getNode(ISD::TRUNCATE, dl, MVT::i1, Elt);
2607  // v2f16 was loaded as an i32. Now we must bitcast it back.
2608  else if (EltVT == MVT::v2f16)
2609  Elt = DAG.getNode(ISD::BITCAST, dl, MVT::v2f16, Elt);
2610 
2611  // If a promoted integer type is used, truncate down to the original
2612  MVT PromotedVT;
2613  if (PromoteScalarIntegerPTX(EltVT, &PromotedVT)) {
2614  Elt = DAG.getNode(ISD::TRUNCATE, dl, EltVT, Elt);
2615  }
2616 
2617  // Extend the element if necessary (e.g. an i8 is loaded
2618  // into an i16 register)
2619  if (Ins[InsIdx].VT.isInteger() &&
2620  Ins[InsIdx].VT.getFixedSizeInBits() >
2621  LoadVT.getFixedSizeInBits()) {
2622  unsigned Extend = Ins[InsIdx].Flags.isSExt() ? ISD::SIGN_EXTEND
2623  : ISD::ZERO_EXTEND;
2624  Elt = DAG.getNode(Extend, dl, Ins[InsIdx].VT, Elt);
2625  }
2626  InVals.push_back(Elt);
2627  }
2628 
2629  // Reset vector tracking state.
2630  VecIdx = -1;
2631  }
2632  ++InsIdx;
2633  }
2634  if (VTs.size() > 0)
2635  --InsIdx;
2636  continue;
2637  }
2638 
2639  // Param has ByVal attribute
2640  // Return MoveParam(param symbol).
2641  // Ideally, the param symbol can be returned directly,
2642  // but when SDNode builder decides to use it in a CopyToReg(),
2643  // machine instruction fails because TargetExternalSymbol
2644  // (not lowered) is target dependent, and CopyToReg assumes
2645  // the source is lowered.
2646  EVT ObjectVT = getValueType(DL, Ty);
2647  assert(ObjectVT == Ins[InsIdx].VT &&
2648  "Ins type did not match function type");
2649  SDValue Arg = getParamSymbol(DAG, idx, PtrVT);
2650  SDValue p = DAG.getNode(NVPTXISD::MoveParam, dl, ObjectVT, Arg);
2651  if (p.getNode())
2652  p.getNode()->setIROrder(idx + 1);
2653  InVals.push_back(p);
2654  }
2655 
2656  // Clang will check explicit VarArg and issue error if any. However, Clang
2657  // will let code with
2658  // implicit var arg like f() pass. See bug 617733.
2659  // We treat this case as if the arg list is empty.
2660  // if (F.isVarArg()) {
2661  // assert(0 && "VarArg not supported yet!");
2662  //}
2663 
2664  if (!OutChains.empty())
2665  DAG.setRoot(DAG.getNode(ISD::TokenFactor, dl, MVT::Other, OutChains));
2666 
2667  return Chain;
2668 }
2669 
2670 SDValue
2672  bool isVarArg,
2673  const SmallVectorImpl<ISD::OutputArg> &Outs,
2674  const SmallVectorImpl<SDValue> &OutVals,
2675  const SDLoc &dl, SelectionDAG &DAG) const {
2676  const MachineFunction &MF = DAG.getMachineFunction();
2677  const Function &F = MF.getFunction();
2678  Type *RetTy = MF.getFunction().getReturnType();
2679 
2680  bool isABI = (STI.getSmVersion() >= 20);
2681  assert(isABI && "Non-ABI compilation is not supported");
2682  if (!isABI)
2683  return Chain;
2684 
2685  const DataLayout &DL = DAG.getDataLayout();
2686  SmallVector<SDValue, 16> PromotedOutVals;
2689  ComputePTXValueVTs(*this, DL, RetTy, VTs, &Offsets);
2690  assert(VTs.size() == OutVals.size() && "Bad return value decomposition");
2691 
2692  for (unsigned i = 0, e = VTs.size(); i != e; ++i) {
2693  SDValue PromotedOutVal = OutVals[i];
2694  MVT PromotedVT;
2695  if (PromoteScalarIntegerPTX(VTs[i], &PromotedVT)) {
2696  VTs[i] = EVT(PromotedVT);
2697  }
2698  if (PromoteScalarIntegerPTX(PromotedOutVal.getValueType(), &PromotedVT)) {
2700  Outs[i].Flags.isSExt() ? ISD::SIGN_EXTEND : ISD::ZERO_EXTEND;
2701  PromotedOutVal = DAG.getNode(Ext, dl, PromotedVT, PromotedOutVal);
2702  }
2703  PromotedOutVals.push_back(PromotedOutVal);
2704  }
2705 
2706  auto VectorInfo = VectorizePTXValueVTs(
2707  VTs, Offsets,
2708  RetTy->isSized() ? getFunctionParamOptimizedAlign(&F, RetTy, DL)
2709  : Align(1));
2710 
2711  // PTX Interoperability Guide 3.3(A): [Integer] Values shorter than
2712  // 32-bits are sign extended or zero extended, depending on whether
2713  // they are signed or unsigned types.
2714  bool ExtendIntegerRetVal =
2715  RetTy->isIntegerTy() && DL.getTypeAllocSizeInBits(RetTy) < 32;
2716 
2717  SmallVector<SDValue, 6> StoreOperands;
2718  for (unsigned i = 0, e = VTs.size(); i != e; ++i) {
2719  // New load/store. Record chain and offset operands.
2720  if (VectorInfo[i] & PVF_FIRST) {
2721  assert(StoreOperands.empty() && "Orphaned operand list.");
2722  StoreOperands.push_back(Chain);
2723  StoreOperands.push_back(DAG.getConstant(Offsets[i], dl, MVT::i32));
2724  }
2725 
2726  SDValue OutVal = OutVals[i];
2727  SDValue RetVal = PromotedOutVals[i];
2728 
2729  if (ExtendIntegerRetVal) {
2730  RetVal = DAG.getNode(Outs[i].Flags.isSExt() ? ISD::SIGN_EXTEND
2731  : ISD::ZERO_EXTEND,
2732  dl, MVT::i32, RetVal);
2733  } else if (OutVal.getValueSizeInBits() < 16) {
2734  // Use 16-bit registers for small load-stores as it's the
2735  // smallest general purpose register size supported by NVPTX.
2736  RetVal = DAG.getNode(ISD::ANY_EXTEND, dl, MVT::i16, RetVal);
2737  }
2738 
2739  // Record the value to return.
2740  StoreOperands.push_back(RetVal);
2741 
2742  // That's the last element of this store op.
2743  if (VectorInfo[i] & PVF_LAST) {
2745  unsigned NumElts = StoreOperands.size() - 2;
2746  switch (NumElts) {
2747  case 1:
2749  break;
2750  case 2:
2752  break;
2753  case 4:
2755  break;
2756  default:
2757  llvm_unreachable("Invalid vector info.");
2758  }
2759 
2760  // Adjust type of load/store op if we've extended the scalar
2761  // return value.
2762  EVT TheStoreType = ExtendIntegerRetVal ? MVT::i32 : VTs[i];
2763  Chain = DAG.getMemIntrinsicNode(
2764  Op, dl, DAG.getVTList(MVT::Other), StoreOperands, TheStoreType,
2766  // Cleanup vector state.
2767  StoreOperands.clear();
2768  }
2769  }
2770 
2771  return DAG.getNode(NVPTXISD::RET_FLAG, dl, MVT::Other, Chain);
2772 }
2773 
2775  SDValue Op, std::string &Constraint, std::vector<SDValue> &Ops,
2776  SelectionDAG &DAG) const {
2777  if (Constraint.length() > 1)
2778  return;
2779  else
2780  TargetLowering::LowerAsmOperandForConstraint(Op, Constraint, Ops, DAG);
2781 }
2782 
2783 static unsigned getOpcForTextureInstr(unsigned Intrinsic) {
2784  switch (Intrinsic) {
2785  default:
2786  return 0;
2787 
2788  case Intrinsic::nvvm_tex_1d_v4f32_s32:
2789  return NVPTXISD::Tex1DFloatS32;
2790  case Intrinsic::nvvm_tex_1d_v4f32_f32:
2792  case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
2794  case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
2796  case Intrinsic::nvvm_tex_1d_v4s32_s32:
2797  return NVPTXISD::Tex1DS32S32;
2798  case Intrinsic::nvvm_tex_1d_v4s32_f32:
2799  return NVPTXISD::Tex1DS32Float;
2800  case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
2802  case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
2804  case Intrinsic::nvvm_tex_1d_v4u32_s32:
2805  return NVPTXISD::Tex1DU32S32;
2806  case Intrinsic::nvvm_tex_1d_v4u32_f32:
2807  return NVPTXISD::Tex1DU32Float;
2808  case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
2810  case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
2812 
2813  case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
2815  case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
2817  case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
2819  case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
2821  case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
2823  case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
2825  case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
2827  case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
2829  case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
2831  case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
2833  case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
2835  case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
2837 
2838  case Intrinsic::nvvm_tex_2d_v4f32_s32:
2839  return NVPTXISD::Tex2DFloatS32;
2840  case Intrinsic::nvvm_tex_2d_v4f32_f32:
2842  case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
2844  case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
2846  case Intrinsic::nvvm_tex_2d_v4s32_s32:
2847  return NVPTXISD::Tex2DS32S32;
2848  case Intrinsic::nvvm_tex_2d_v4s32_f32:
2849  return NVPTXISD::Tex2DS32Float;
2850  case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
2852  case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
2854  case Intrinsic::nvvm_tex_2d_v4u32_s32:
2855  return NVPTXISD::Tex2DU32S32;
2856  case Intrinsic::nvvm_tex_2d_v4u32_f32:
2857  return NVPTXISD::Tex2DU32Float;
2858  case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
2860  case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
2862 
2863  case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
2865  case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
2867  case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
2869  case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
2871  case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
2873  case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
2875  case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
2877  case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
2879  case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
2881  case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
2883  case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
2885  case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
2887 
2888  case Intrinsic::nvvm_tex_3d_v4f32_s32:
2889  return NVPTXISD::Tex3DFloatS32;
2890  case Intrinsic::nvvm_tex_3d_v4f32_f32:
2892  case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
2894  case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
2896  case Intrinsic::nvvm_tex_3d_v4s32_s32:
2897  return NVPTXISD::Tex3DS32S32;
2898  case Intrinsic::nvvm_tex_3d_v4s32_f32:
2899  return NVPTXISD::Tex3DS32Float;
2900  case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
2902  case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
2904  case Intrinsic::nvvm_tex_3d_v4u32_s32:
2905  return NVPTXISD::Tex3DU32S32;
2906  case Intrinsic::nvvm_tex_3d_v4u32_f32:
2907  return NVPTXISD::Tex3DU32Float;
2908  case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
2910  case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
2912 
2913  case Intrinsic::nvvm_tex_cube_v4f32_f32:
2915  case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
2917  case Intrinsic::nvvm_tex_cube_v4s32_f32:
2919  case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
2921  case Intrinsic::nvvm_tex_cube_v4u32_f32:
2923  case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
2925 
2926  case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
2928  case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
2930  case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
2932  case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
2934  case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
2936  case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
2938 
2939  case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
2941  case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
2943  case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
2945  case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
2947  case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
2949  case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
2951  case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
2953  case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
2955  case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
2957  case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
2959  case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
2961  case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
2963 
2964  case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
2966  case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
2968  case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
2970  case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
2972  case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
2974  case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
2976  case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
2978  case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
2980  case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
2982  case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
2984  case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
2986  case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
2988 
2989  case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
2991  case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
2993  case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
2995  case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
2997  case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
2999  case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
3001  case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
3003  case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
3005  case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
3007  case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
3009  case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
3011  case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
3013 
3014  case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
3016  case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
3018  case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
3020  case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
3022  case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
3024  case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
3026  case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
3028  case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
3030  case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
3032  case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
3034  case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
3036  case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
3038 
3039  case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
3041  case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
3043  case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
3045  case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
3047  case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
3049  case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
3051  case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
3053  case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
3055  case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
3057  case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
3059  case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
3061  case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
3063 
3064  case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
3066  case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
3068  case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
3070  case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
3072  case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
3074  case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
3076  case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
3078  case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
3080  case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
3082  case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
3084  case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
3086  case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
3088 
3089  case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
3091  case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
3093  case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
3095  case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
3097  case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
3099  case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
3101 
3102  case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
3104  case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
3106  case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
3108  case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
3110  case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
3112  case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
3114 
3115  case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
3117  case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
3119  case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
3121  case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
3123  case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
3125  case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
3127  case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
3129  case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
3131  case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
3133  case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
3135  case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
3137  case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
3139  }
3140 }
3141 
3142 static unsigned getOpcForSurfaceInstr(unsigned Intrinsic) {
3143  switch (Intrinsic) {
3144  default:
3145  return 0;
3146  case Intrinsic::nvvm_suld_1d_i8_clamp:
3147  return NVPTXISD::Suld1DI8Clamp;
3148  case Intrinsic::nvvm_suld_1d_i16_clamp:
3149  return NVPTXISD::Suld1DI16Clamp;
3150  case Intrinsic::nvvm_suld_1d_i32_clamp:
3151  return NVPTXISD::Suld1DI32Clamp;
3152  case Intrinsic::nvvm_suld_1d_i64_clamp:
3153  return NVPTXISD::Suld1DI64Clamp;
3154  case Intrinsic::nvvm_suld_1d_v2i8_clamp:
3156  case Intrinsic::nvvm_suld_1d_v2i16_clamp:
3158  case Intrinsic::nvvm_suld_1d_v2i32_clamp:
3160  case Intrinsic::nvvm_suld_1d_v2i64_clamp:
3162  case Intrinsic::nvvm_suld_1d_v4i8_clamp:
3164  case Intrinsic::nvvm_suld_1d_v4i16_clamp:
3166  case Intrinsic::nvvm_suld_1d_v4i32_clamp:
3168  case Intrinsic::nvvm_suld_1d_array_i8_clamp:
3170  case Intrinsic::nvvm_suld_1d_array_i16_clamp:
3172  case Intrinsic::nvvm_suld_1d_array_i32_clamp:
3174  case Intrinsic::nvvm_suld_1d_array_i64_clamp:
3176  case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
3178  case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
3180  case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
3182  case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
3184  case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
3186  case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
3188  case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
3190  case Intrinsic::nvvm_suld_2d_i8_clamp:
3191  return NVPTXISD::Suld2DI8Clamp;
3192  case Intrinsic::nvvm_suld_2d_i16_clamp:
3193  return NVPTXISD::Suld2DI16Clamp;
3194  case Intrinsic::nvvm_suld_2d_i32_clamp:
3195  return NVPTXISD::Suld2DI32Clamp;
3196  case Intrinsic::nvvm_suld_2d_i64_clamp:
3197  return NVPTXISD::Suld2DI64Clamp;
3198  case Intrinsic::nvvm_suld_2d_v2i8_clamp:
3200  case Intrinsic::nvvm_suld_2d_v2i16_clamp:
3202  case Intrinsic::nvvm_suld_2d_v2i32_clamp:
3204  case Intrinsic::nvvm_suld_2d_v2i64_clamp:
3206  case Intrinsic::nvvm_suld_2d_v4i8_clamp:
3208  case Intrinsic::nvvm_suld_2d_v4i16_clamp:
3210  case Intrinsic::nvvm_suld_2d_v4i32_clamp:
3212  case Intrinsic::nvvm_suld_2d_array_i8_clamp:
3214  case Intrinsic::nvvm_suld_2d_array_i16_clamp:
3216  case Intrinsic::nvvm_suld_2d_array_i32_clamp:
3218  case Intrinsic::nvvm_suld_2d_array_i64_clamp:
3220  case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
3222  case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
3224  case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
3226  case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
3228  case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
3230  case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
3232  case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
3234  case Intrinsic::nvvm_suld_3d_i8_clamp:
3235  return NVPTXISD::Suld3DI8Clamp;
3236  case Intrinsic::nvvm_suld_3d_i16_clamp:
3237  return NVPTXISD::Suld3DI16Clamp;
3238  case Intrinsic::nvvm_suld_3d_i32_clamp:
3239  return NVPTXISD::Suld3DI32Clamp;
3240  case Intrinsic::nvvm_suld_3d_i64_clamp:
3241  return NVPTXISD::Suld3DI64Clamp;
3242  case Intrinsic::nvvm_suld_3d_v2i8_clamp:
3244  case Intrinsic::nvvm_suld_3d_v2i16_clamp:
3246  case Intrinsic::nvvm_suld_3d_v2i32_clamp:
3248  case Intrinsic::nvvm_suld_3d_v2i64_clamp:
3250  case Intrinsic::nvvm_suld_3d_v4i8_clamp:
3252  case Intrinsic::nvvm_suld_3d_v4i16_clamp:
3254  case Intrinsic::nvvm_suld_3d_v4i32_clamp:
3256  case Intrinsic::nvvm_suld_1d_i8_trap:
3257  return NVPTXISD::Suld1DI8Trap;
3258  case Intrinsic::nvvm_suld_1d_i16_trap:
3259  return NVPTXISD::Suld1DI16Trap;
3260  case Intrinsic::nvvm_suld_1d_i32_trap:
3261  return NVPTXISD::Suld1DI32Trap;
3262  case Intrinsic::nvvm_suld_1d_i64_trap:
3263  return NVPTXISD::Suld1DI64Trap;
3264  case Intrinsic::nvvm_suld_1d_v2i8_trap:
3265  return NVPTXISD::Suld1DV2I8Trap;
3266  case Intrinsic::nvvm_suld_1d_v2i16_trap:
3268  case Intrinsic::nvvm_suld_1d_v2i32_trap:
3270  case Intrinsic::nvvm_suld_1d_v2i64_trap:
3272  case Intrinsic::nvvm_suld_1d_v4i8_trap:
3273  return NVPTXISD::Suld1DV4I8Trap;
3274  case Intrinsic::nvvm_suld_1d_v4i16_trap:
3276  case Intrinsic::nvvm_suld_1d_v4i32_trap:
3278  case Intrinsic::nvvm_suld_1d_array_i8_trap:
3280  case Intrinsic::nvvm_suld_1d_array_i16_trap:
3282  case Intrinsic::nvvm_suld_1d_array_i32_trap:
3284  case Intrinsic::nvvm_suld_1d_array_i64_trap:
3286  case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
3288  case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
3290  case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
3292  case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
3294  case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
3296  case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
3298  case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
3300  case Intrinsic::nvvm_suld_2d_i8_trap:
3301  return NVPTXISD::Suld2DI8Trap;
3302  case Intrinsic::nvvm_suld_2d_i16_trap:
3303  return NVPTXISD::Suld2DI16Trap;
3304  case Intrinsic::nvvm_suld_2d_i32_trap:
3305  return NVPTXISD::Suld2DI32Trap;
3306  case Intrinsic::nvvm_suld_2d_i64_trap:
3307  return NVPTXISD::Suld2DI64Trap;
3308  case Intrinsic::nvvm_suld_2d_v2i8_trap:
3309  return NVPTXISD::Suld2DV2I8Trap;
3310  case Intrinsic::nvvm_suld_2d_v2i16_trap:
3312  case Intrinsic::nvvm_suld_2d_v2i32_trap:
3314  case Intrinsic::nvvm_suld_2d_v2i64_trap:
3316  case Intrinsic::nvvm_suld_2d_v4i8_trap:
3317  return NVPTXISD::Suld2DV4I8Trap;
3318  case Intrinsic::nvvm_suld_2d_v4i16_trap:
3320  case Intrinsic::nvvm_suld_2d_v4i32_trap:
3322  case Intrinsic::nvvm_suld_2d_array_i8_trap:
3324  case Intrinsic::nvvm_suld_2d_array_i16_trap:
3326  case Intrinsic::nvvm_suld_2d_array_i32_trap:
3328  case Intrinsic::nvvm_suld_2d_array_i64_trap:
3330  case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
3332  case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
3334  case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
3336  case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
3338  case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
3340  case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
3342  case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
3344  case Intrinsic::nvvm_suld_3d_i8_trap:
3345  return NVPTXISD::Suld3DI8Trap;
3346  case Intrinsic::nvvm_suld_3d_i16_trap:
3347  return NVPTXISD::Suld3DI16Trap;
3348  case Intrinsic::nvvm_suld_3d_i32_trap:
3349  return NVPTXISD::Suld3DI32Trap;
3350  case Intrinsic::nvvm_suld_3d_i64_trap:
3351  return NVPTXISD::Suld3DI64Trap;
3352  case Intrinsic::nvvm_suld_3d_v2i8_trap:
3353  return NVPTXISD::Suld3DV2I8Trap;
3354  case Intrinsic::nvvm_suld_3d_v2i16_trap:
3356  case Intrinsic::nvvm_suld_3d_v2i32_trap:
3358  case Intrinsic::nvvm_suld_3d_v2i64_trap:
3360  case Intrinsic::nvvm_suld_3d_v4i8_trap:
3361  return NVPTXISD::Suld3DV4I8Trap;
3362  case Intrinsic::nvvm_suld_3d_v4i16_trap:
3364  case Intrinsic::nvvm_suld_3d_v4i32_trap:
3366  case Intrinsic::nvvm_suld_1d_i8_zero:
3367  return NVPTXISD::Suld1DI8Zero;
3368  case Intrinsic::nvvm_suld_1d_i16_zero:
3369  return NVPTXISD::Suld1DI16Zero;
3370  case Intrinsic::nvvm_suld_1d_i32_zero:
3371  return NVPTXISD::Suld1DI32Zero;
3372  case Intrinsic::nvvm_suld_1d_i64_zero:
3373  return NVPTXISD::Suld1DI64Zero;
3374  case Intrinsic::nvvm_suld_1d_v2i8_zero:
3375  return NVPTXISD::Suld1DV2I8Zero;
3376  case Intrinsic::nvvm_suld_1d_v2i16_zero:
3378  case Intrinsic::nvvm_suld_1d_v2i32_zero:
3380  case Intrinsic::nvvm_suld_1d_v2i64_zero:
3382  case Intrinsic::nvvm_suld_1d_v4i8_zero:
3383  return NVPTXISD::Suld1DV4I8Zero;
3384  case Intrinsic::nvvm_suld_1d_v4i16_zero:
3386  case Intrinsic::nvvm_suld_1d_v4i32_zero:
3388  case Intrinsic::nvvm_suld_1d_array_i8_zero:
3390  case Intrinsic::nvvm_suld_1d_array_i16_zero:
3392  case Intrinsic::nvvm_suld_1d_array_i32_zero:
3394  case Intrinsic::nvvm_suld_1d_array_i64_zero:
3396  case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
3398  case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
3400  case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
3402  case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
3404  case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
3406  case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
3408  case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
3410  case Intrinsic::nvvm_suld_2d_i8_zero:
3411  return NVPTXISD::Suld2DI8Zero;
3412  case Intrinsic::nvvm_suld_2d_i16_zero:
3413  return NVPTXISD::Suld2DI16Zero;
3414  case Intrinsic::nvvm_suld_2d_i32_zero:
3415  return NVPTXISD::Suld2DI32Zero;
3416  case Intrinsic::nvvm_suld_2d_i64_zero:
3417  return NVPTXISD::Suld2DI64Zero;
3418  case Intrinsic::nvvm_suld_2d_v2i8_zero:
3419  return NVPTXISD::Suld2DV2I8Zero;
3420  case Intrinsic::nvvm_suld_2d_v2i16_zero:
3422  case Intrinsic::nvvm_suld_2d_v2i32_zero:
3424  case Intrinsic::nvvm_suld_2d_v2i64_zero:
3426  case Intrinsic::nvvm_suld_2d_v4i8_zero:
3427  return NVPTXISD::Suld2DV4I8Zero;
3428  case Intrinsic::nvvm_suld_2d_v4i16_zero:
3430  case Intrinsic::nvvm_suld_2d_v4i32_zero:
3432  case Intrinsic::nvvm_suld_2d_array_i8_zero:
3434  case Intrinsic::nvvm_suld_2d_array_i16_zero:
3436  case Intrinsic::nvvm_suld_2d_array_i32_zero:
3438  case Intrinsic::nvvm_suld_2d_array_i64_zero:
3440  case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
3442  case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
3444  case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
3446  case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
3448  case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
3450  case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
3452  case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
3454  case Intrinsic::nvvm_suld_3d_i8_zero:
3455  return NVPTXISD::Suld3DI8Zero;
3456  case Intrinsic::nvvm_suld_3d_i16_zero:
3457  return NVPTXISD::Suld3DI16Zero;
3458  case Intrinsic::nvvm_suld_3d_i32_zero:
3459  return NVPTXISD::Suld3DI32Zero;
3460  case Intrinsic::nvvm_suld_3d_i64_zero:
3461  return NVPTXISD::Suld3DI64Zero;
3462  case Intrinsic::nvvm_suld_3d_v2i8_zero:
3463  return NVPTXISD::Suld3DV2I8Zero;
3464  case Intrinsic::nvvm_suld_3d_v2i16_zero:
3466  case Intrinsic::nvvm_suld_3d_v2i32_zero:
3468  case Intrinsic::nvvm_suld_3d_v2i64_zero:
3470  case Intrinsic::nvvm_suld_3d_v4i8_zero:
3471  return NVPTXISD::Suld3DV4I8Zero;
3472  case Intrinsic::nvvm_suld_3d_v4i16_zero:
3474  case Intrinsic::nvvm_suld_3d_v4i32_zero:
3476  }
3477 }
3478 
3479 // llvm.ptx.memcpy.const and llvm.ptx.memmove.const need to be modeled as
3480 // TgtMemIntrinsic
3481 // because we need the information that is only available in the "Value" type
3482 // of destination
3483 // pointer. In particular, the address space information.
3485  IntrinsicInfo &Info, const CallInst &I,
3486  MachineFunction &MF, unsigned Intrinsic) const {
3487  switch (Intrinsic) {
3488  default:
3489  return false;
3490  case Intrinsic::nvvm_match_all_sync_i32p:
3491  case Intrinsic::nvvm_match_all_sync_i64p:
3493  // memVT is bogus. These intrinsics have IntrInaccessibleMemOnly attribute
3494  // in order to model data exchange with other threads, but perform no real
3495  // memory accesses.
3496  Info.memVT = MVT::i1;
3497 
3498  // Our result depends on both our and other thread's arguments.
3500  return true;
3501  case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col:
3502  case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row:
3503  case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col_stride:
3504  case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row_stride:
3505  case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col:
3506  case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row:
3507  case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col_stride:
3508  case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row_stride:
3509  case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col:
3510  case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row:
3511  case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col_stride:
3512  case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row_stride:
3513  case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col:
3514  case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row:
3515  case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col_stride:
3516  case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row_stride:
3517  case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col:
3518  case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row:
3519  case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col_stride:
3520  case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row_stride:
3521  case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col:
3522  case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row:
3523  case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col_stride:
3524  case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row_stride: {
3526  Info.memVT = MVT::v8f16;
3527  Info.ptrVal = I.getArgOperand(0);
3528  Info.offset = 0;
3530  Info.align = Align(16);
3531  return true;
3532  }
3533  case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col:
3534  case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col_stride:
3535  case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col_stride:
3536  case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col:
3537  case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row:
3538  case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row_stride:
3539  case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row_stride:
3540  case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row:
3541  case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col:
3542  case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col_stride:
3543  case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row:
3544  case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row_stride:
3545  case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col:
3546  case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col_stride:
3547  case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col_stride:
3548  case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col:
3549  case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row:
3550  case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row_stride:
3551  case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row_stride:
3552  case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row:
3553  case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col:
3554  case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col_stride:
3555  case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row:
3556  case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row_stride: {
3558  Info.memVT = MVT::v2i32;
3559  Info.ptrVal = I.getArgOperand(0);
3560  Info.offset = 0;
3562  Info.align = Align(8);
3563  return true;
3564  }
3565 
3566  case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col:
3567  case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col_stride:
3568  case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col_stride:
3569  case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col:
3570  case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row:
3571  case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row_stride:
3572  case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row_stride:
3573  case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row:
3574  case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col:
3575  case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col_stride:
3576  case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row:
3577  case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row_stride:
3578  case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col:
3579  case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col_stride:
3580  case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row:
3581  case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row_stride:
3582 
3583  case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col:
3584  case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col_stride:
3585  case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col_stride:
3586  case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col:
3587  case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row:
3588  case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row_stride:
3589  case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row_stride:
3590  case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row:
3591  case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col:
3592  case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col_stride:
3593  case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row:
3594  case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row_stride:
3595  case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col:
3596  case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col_stride:
3597  case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row:
3598  case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row_stride:
3599  case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16:
3600  case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_trans_b16: {
3602  Info.memVT = MVT::v4i32;
3603  Info.ptrVal = I.getArgOperand(0);
3604  Info.offset = 0;
3606  Info.align = Align(16);
3607  return true;
3608  }
3609 
3610  case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col:
3611  case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col_stride:
3612  case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col_stride:
3613  case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col:
3614  case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row:
3615  case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row_stride:
3616  case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row_stride:
3617  case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row:
3618 
3619  case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col:
3620  case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col_stride:
3621  case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col_stride:
3622  case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col:
3623  case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row:
3624  case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row_stride:
3625  case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row_stride:
3626  case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row:
3627  case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row:
3628  case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row_stride:
3629  case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col:
3630  case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col_stride:
3631  case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row:
3632  case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row_stride:
3633  case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row_stride:
3634  case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row:
3635  case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col:
3636  case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col_stride:
3637  case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col_stride:
3638  case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col:
3639  case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16:
3640  case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_trans_b16: {
3642  Info.memVT = MVT::i32;
3643  Info.ptrVal = I.getArgOperand(0);
3644  Info.offset = 0;
3646  Info.align = Align(4);
3647  return true;
3648  }
3649 
3650  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col:
3651  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row:
3652  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col_stride:
3653  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row_stride:
3654  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col:
3655  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row:
3656  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col_stride:
3657  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row_stride:
3658  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col:
3659  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row:
3660  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col_stride:
3661  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row_stride: {
3663  Info.memVT = MVT::v4f16;
3664  Info.ptrVal = I.getArgOperand(0);
3665  Info.offset = 0;
3667  Info.align = Align(16);
3668  return true;
3669  }
3670 
3671  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col:
3672  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row:
3673  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col_stride:
3674  case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row_stride:
3675  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col:
3676  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row:
3677  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col_stride:
3678  case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row_stride:
3679  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col:
3680  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row:
3681  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col_stride:
3682  case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row_stride:
3683  case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col:
3684  case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row:
3685  case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col_stride:
3686  case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row_stride: {
3688  Info.memVT = MVT::v8f32;
3689  Info.ptrVal = I.getArgOperand(0);
3690  Info.offset = 0;
3692  Info.align = Align(16);
3693  return true;
3694  }
3695 
3696  case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col:
3697  case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col_stride:
3698  case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row:
3699  case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row_stride:
3700 
3701  case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col:
3702  case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col_stride:
3703  case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row:
3704  case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row_stride:
3705 
3706  case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col:
3707  case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col_stride:
3708  case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row:
3709  case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row_stride:
3710  case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col:
3711  case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col_stride:
3712  case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row:
3713  case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row_stride:
3714  case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col:
3715  case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col_stride:
3716  case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row:
3717  case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row_stride: {
3719  Info.memVT = MVT::v8i32;
3720  Info.ptrVal = I.getArgOperand(0);
3721  Info.offset = 0;
3723  Info.align = Align(16);
3724  return true;
3725  }
3726 
3727  case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col:
3728  case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col_stride:
3729  case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row:
3730  case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row_stride:
3731  case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col:
3732  case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col_stride:
3733  case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row:
3734  case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row_stride:
3735  case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16:
3736  case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_trans_b16: {
3738  Info.memVT = MVT::v2i32;
3739  Info.ptrVal = I.getArgOperand(0);
3740  Info.offset = 0;
3742  Info.align = Align(8);
3743  return true;
3744  }
3745 
3746  case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col:
3747  case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col_stride:
3748  case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row:
3749  case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row_stride:
3750 
3751  case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col:
3752  case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col_stride:
3753  case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row:
3754  case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row_stride: {
3756  Info.memVT = MVT::f64;
3757  Info.ptrVal = I.getArgOperand(0);
3758  Info.offset = 0;
3760  Info.align = Align(8);
3761  return true;
3762  }
3763 
3764  case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col:
3765  case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col_stride:
3766  case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row:
3767  case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row_stride: {
3769  Info.memVT = MVT::v2f64;
3770  Info.ptrVal = I.getArgOperand(0);
3771  Info.offset = 0;
3773  Info.align = Align(16);
3774  return true;
3775  }
3776 
3777  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col:
3778  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row:
3779  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col_stride:
3780  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row_stride:
3781  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col:
3782  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row:
3783  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col_stride:
3784  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row_stride:
3785  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col:
3786  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row:
3787  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col_stride:
3788  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row_stride: {
3789  Info.opc = ISD::INTRINSIC_VOID;
3790  Info.memVT = MVT::v4f16;
3791  Info.ptrVal = I.getArgOperand(0);
3792  Info.offset = 0;
3794  Info.align = Align(16);
3795  return true;
3796  }
3797 
3798  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col:
3799  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row:
3800  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col_stride:
3801  case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row_stride:
3802  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col:
3803  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row:
3804  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col_stride:
3805  case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row_stride:
3806  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col:
3807  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row:
3808  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col_stride:
3809  case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row_stride:
3810  case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col:
3811  case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row:
3812  case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col_stride:
3813  case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row_stride: {
3814  Info.opc = ISD::INTRINSIC_VOID;
3815  Info.memVT = MVT::v8f32;
3816  Info.ptrVal = I.getArgOperand(0);
3817  Info.offset = 0;
3819  Info.align = Align(16);
3820  return true;
3821  }
3822 
3823  case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col:
3824  case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col_stride:
3825  case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row:
3826  case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row_stride:
3827  case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col:
3828  case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col_stride:
3829  case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row:
3830  case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row_stride:
3831  case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col:
3832  case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col_stride:
3833  case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row:
3834  case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row_stride: {
3835  Info.opc = ISD::INTRINSIC_VOID;
3836  Info.memVT = MVT::v8i32;
3837  Info.ptrVal = I.getArgOperand(0);
3838  Info.offset = 0;
3840  Info.align = Align(16);
3841  return true;
3842  }
3843 
3844  case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col:
3845  case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col_stride:
3846  case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row:
3847  case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row_stride:
3848  case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col:
3849  case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col_stride:
3850  case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row:
3851  case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row_stride: {
3852  Info.opc = ISD::INTRINSIC_VOID;
3853  Info.memVT = MVT::v2i32;
3854  Info.ptrVal = I.getArgOperand(0);
3855  Info.offset = 0;
3857  Info.align = Align(8);
3858  return true;
3859  }
3860 
3861  case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col:
3862  case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col_stride:
3863  case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row:
3864  case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row_stride: {
3865  Info.opc = ISD::INTRINSIC_VOID;
3866  Info.memVT = MVT::v2f64;
3867  Info.ptrVal = I.getArgOperand(0);
3868  Info.offset = 0;
3870  Info.align = Align(16);
3871  return true;
3872  }
3873 
3874  case Intrinsic::nvvm_atomic_load_inc_32:
3875  case Intrinsic::nvvm_atomic_load_dec_32:
3876 
3877  case Intrinsic::nvvm_atomic_add_gen_f_cta:
3878  case Intrinsic::nvvm_atomic_add_gen_f_sys:
3879  case Intrinsic::nvvm_atomic_add_gen_i_cta:
3880  case Intrinsic::nvvm_atomic_add_gen_i_sys:
3881  case Intrinsic::nvvm_atomic_and_gen_i_cta:
3882  case Intrinsic::nvvm_atomic_and_gen_i_sys:
3883  case Intrinsic::nvvm_atomic_cas_gen_i_cta:
3884  case Intrinsic::nvvm_atomic_cas_gen_i_sys:
3885  case Intrinsic::nvvm_atomic_dec_gen_i_cta:
3886  case Intrinsic::nvvm_atomic_dec_gen_i_sys:
3887  case Intrinsic::nvvm_atomic_inc_gen_i_cta:
3888  case Intrinsic::nvvm_atomic_inc_gen_i_sys:
3889  case Intrinsic::nvvm_atomic_max_gen_i_cta:
3890  case Intrinsic::nvvm_atomic_max_gen_i_sys:
3891  case Intrinsic::nvvm_atomic_min_gen_i_cta:
3892  case Intrinsic::nvvm_atomic_min_gen_i_sys:
3893  case Intrinsic::nvvm_atomic_or_gen_i_cta:
3894  case Intrinsic::nvvm_atomic_or_gen_i_sys:
3895  case Intrinsic::nvvm_atomic_exch_gen_i_cta:
3896  case Intrinsic::nvvm_atomic_exch_gen_i_sys:
3897  case Intrinsic::nvvm_atomic_xor_gen_i_cta:
3898  case Intrinsic::nvvm_atomic_xor_gen_i_sys: {
3899  auto &DL = I.getModule()->getDataLayout();
3901  Info.memVT = getValueType(DL, I.getType());
3902  Info.ptrVal = I.getArgOperand(0);
3903  Info.offset = 0;
3905  Info.align.reset();
3906  return true;
3907  }
3908 
3909  case Intrinsic::nvvm_ldu_global_i:
3910  case Intrinsic::nvvm_ldu_global_f:
3911  case Intrinsic::nvvm_ldu_global_p: {
3912  auto &DL = I.getModule()->getDataLayout();
3914  if (Intrinsic == Intrinsic::nvvm_ldu_global_i)
3915  Info.memVT = getValueType(DL, I.getType());
3916  else if(Intrinsic == Intrinsic::nvvm_ldu_global_p)
3917  Info.memVT = getPointerTy(DL);
3918  else
3919  Info.memVT = getValueType(DL, I.getType());
3920  Info.ptrVal = I.getArgOperand(0);
3921  Info.offset = 0;
3923  Info.align = cast<ConstantInt>(I.getArgOperand(1))->getMaybeAlignValue();
3924 
3925  return true;
3926  }
3927  case Intrinsic::nvvm_ldg_global_i:
3928  case Intrinsic::nvvm_ldg_global_f:
3929  case Intrinsic::nvvm_ldg_global_p: {
3930  auto &DL = I.getModule()->getDataLayout();
3931 
3933  if (Intrinsic == Intrinsic::nvvm_ldg_global_i)
3934  Info.memVT = getValueType(DL, I.getType());
3935  else if(Intrinsic == Intrinsic::nvvm_ldg_global_p)
3936  Info.memVT = getPointerTy(DL);
3937  else
3938  Info.memVT = getValueType(DL, I.getType());
3939  Info.ptrVal = I.getArgOperand(0);
3940  Info.offset = 0;
3942  Info.align = cast<ConstantInt>(I.getArgOperand(1))->getMaybeAlignValue();
3943 
3944  return true;
3945  }
3946 
3947  case Intrinsic::nvvm_tex_1d_v4f32_s32:
3948  case Intrinsic::nvvm_tex_1d_v4f32_f32:
3949  case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
3950  case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
3951  case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
3952  case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
3953  case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
3954  case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
3955  case Intrinsic::nvvm_tex_2d_v4f32_s32:
3956  case Intrinsic::nvvm_tex_2d_v4f32_f32:
3957  case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
3958  case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
3959  case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
3960  case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
3961  case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
3962  case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
3963  case Intrinsic::nvvm_tex_3d_v4f32_s32:
3964  case Intrinsic::nvvm_tex_3d_v4f32_f32:
3965  case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
3966  case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
3967  case Intrinsic::nvvm_tex_cube_v4f32_f32:
3968  case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
3969  case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
3970  case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
3971  case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
3972  case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
3973  case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
3974  case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
3975  case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
3976  case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
3977  case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
3978  case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
3979  case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
3980  case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
3981  case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
3982  case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
3983  case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
3984  case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
3985  case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
3986  case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
3987  case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
3988  case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
3989  case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
3990  case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
3991  case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
3992  case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
3993  case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
3994  case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
3995  case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
3996  case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
3997  case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
3998  case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
3999  case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
4000  case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
4001  case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
4002  case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
4003  Info.opc = getOpcForTextureInstr(Intrinsic);
4004  Info.memVT = MVT::v4f32;
4005  Info.ptrVal = nullptr;
4006  Info.offset = 0;
4008  Info.align = Align(16);
4009  return true;
4010 
4011  case Intrinsic::nvvm_tex_1d_v4s32_s32:
4012  case Intrinsic::nvvm_tex_1d_v4s32_f32:
4013  case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
4014  case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
4015  case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
4016  case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
4017  case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
4018  case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
4019  case Intrinsic::nvvm_tex_2d_v4s32_s32:
4020  case Intrinsic::nvvm_tex_2d_v4s32_f32:
4021  case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
4022  case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
4023  case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
4024  case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
4025  case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
4026  case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
4027  case Intrinsic::nvvm_tex_3d_v4s32_s32:
4028  case Intrinsic::nvvm_tex_3d_v4s32_f32:
4029  case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
4030  case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
4031  case Intrinsic::nvvm_tex_cube_v4s32_f32:
4032  case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
4033  case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
4034  case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
4035  case Intrinsic::nvvm_tex_cube_v4u32_f32:
4036  case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
4037  case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
4038  case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
4039  case Intrinsic::nvvm_tex_1d_v4u32_s32:
4040  case Intrinsic::nvvm_tex_1d_v4u32_f32:
4041  case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
4042  case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
4043  case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
4044  case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
4045  case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
4046  case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
4047  case Intrinsic::nvvm_tex_2d_v4u32_s32:
4048  case Intrinsic::nvvm_tex_2d_v4u32_f32:
4049  case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
4050  case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
4051  case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
4052  case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
4053  case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
4054  case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
4055  case Intrinsic::nvvm_tex_3d_v4u32_s32:
4056  case Intrinsic::nvvm_tex_3d_v4u32_f32:
4057  case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
4058  case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
4059  case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
4060  case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
4061  case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
4062  case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
4063  case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
4064  case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
4065  case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
4066  case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
4067  case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
4068  case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
4069  case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
4070  case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
4071  case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
4072  case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
4073  case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
4074  case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
4075  case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
4076  case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
4077  case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
4078  case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
4079  case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
4080  case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
4081  case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
4082  case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
4083  case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
4084  case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
4085  case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
4086  case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
4087  case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
4088  case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
4089  case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
4090  case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
4091  case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
4092  case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
4093  case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
4094  case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
4095  case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
4096  case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
4097  case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
4098  case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
4099  case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
4100  case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
4101  case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
4102  case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
4103  case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
4104  case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
4105  case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
4106  case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
4107  case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
4108  case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
4109  case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
4110  case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
4111  case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
4112  case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
4113  case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
4114  case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
4115  case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
4116  case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
4117  case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
4118  case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
4119  case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
4120  case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
4121  case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
4122  case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
4123  Info.opc = getOpcForTextureInstr(Intrinsic);
4124  Info.memVT = MVT::v4i32;
4125  Info.ptrVal = nullptr;
4126  Info.offset = 0;
4128  Info.align = Align(16);
4129  return true;
4130 
4131  case Intrinsic::nvvm_suld_1d_i8_clamp:
4132  case Intrinsic::nvvm_suld_1d_v2i8_clamp:
4133  case Intrinsic::nvvm_suld_1d_v4i8_clamp:
4134  case Intrinsic::nvvm_suld_1d_array_i8_clamp:
4135  case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
4136  case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
4137  case Intrinsic::nvvm_suld_2d_i8_clamp:
4138  case Intrinsic::nvvm_suld_2d_v2i8_clamp:
4139  case Intrinsic::nvvm_suld_2d_v4i8_clamp:
4140  case Intrinsic::nvvm_suld_2d_array_i8_clamp:
4141  case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
4142  case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
4143  case Intrinsic::nvvm_suld_3d_i8_clamp:
4144  case Intrinsic::nvvm_suld_3d_v2i8_clamp:
4145  case Intrinsic::nvvm_suld_3d_v4i8_clamp:
4146  case Intrinsic::nvvm_suld_1d_i8_trap:
4147  case Intrinsic::nvvm_suld_1d_v2i8_trap:
4148  case Intrinsic::nvvm_suld_1d_v4i8_trap:
4149  case Intrinsic::nvvm_suld_1d_array_i8_trap:
4150  case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
4151  case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
4152  case Intrinsic::nvvm_suld_2d_i8_trap:
4153  case Intrinsic::nvvm_suld_2d_v2i8_trap:
4154  case Intrinsic::nvvm_suld_2d_v4i8_trap:
4155  case Intrinsic::nvvm_suld_2d_array_i8_trap:
4156  case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
4157  case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
4158  case Intrinsic::nvvm_suld_3d_i8_trap:
4159  case Intrinsic::nvvm_suld_3d_v2i8_trap:
4160  case Intrinsic::nvvm_suld_3d_v4i8_trap:
4161  case Intrinsic::nvvm_suld_1d_i8_zero:
4162  case Intrinsic::nvvm_suld_1d_v2i8_zero:
4163  case Intrinsic::nvvm_suld_1d_v4i8_zero:
4164  case Intrinsic::nvvm_suld_1d_array_i8_zero:
4165  case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
4166  case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
4167  case Intrinsic::nvvm_suld_2d_i8_zero:
4168  case Intrinsic::nvvm_suld_2d_v2i8_zero:
4169  case Intrinsic::nvvm_suld_2d_v4i8_zero:
4170  case Intrinsic::nvvm_suld_2d_array_i8_zero:
4171  case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
4172  case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
4173  case Intrinsic::nvvm_suld_3d_i8_zero:
4174  case Intrinsic::nvvm_suld_3d_v2i8_zero:
4175  case Intrinsic::nvvm_suld_3d_v4i8_zero:
4176  Info.opc = getOpcForSurfaceInstr(Intrinsic);
4177  Info.memVT = MVT::i8;
4178  Info.ptrVal = nullptr;
4179  Info.offset = 0;
4181  Info.align = Align(16);
4182  return true;
4183 
4184  case Intrinsic::nvvm_suld_1d_i16_clamp:
4185  case Intrinsic::nvvm_suld_1d_v2i16_clamp:
4186  case Intrinsic::nvvm_suld_1d_v4i16_clamp:
4187  case Intrinsic::nvvm_suld_1d_array_i16_clamp:
4188  case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
4189  case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
4190  case Intrinsic::nvvm_suld_2d_i16_clamp:
4191  case Intrinsic::nvvm_suld_2d_v2i16_clamp:
4192  case Intrinsic::nvvm_suld_2d_v4i16_clamp:
4193  case Intrinsic::nvvm_suld_2d_array_i16_clamp:
4194  case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
4195  case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
4196  case Intrinsic::nvvm_suld_3d_i16_clamp:
4197  case Intrinsic::nvvm_suld_3d_v2i16_clamp:
4198  case Intrinsic::nvvm_suld_3d_v4i16_clamp:
4199  case Intrinsic::nvvm_suld_1d_i16_trap:
4200  case Intrinsic::nvvm_suld_1d_v2i16_trap:
4201  case Intrinsic::nvvm_suld_1d_v4i16_trap:
4202  case Intrinsic::nvvm_suld_1d_array_i16_trap:
4203  case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
4204  case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
4205  case Intrinsic::nvvm_suld_2d_i16_trap:
4206  case Intrinsic::nvvm_suld_2d_v2i16_trap:
4207  case Intrinsic::nvvm_suld_2d_v4i16_trap:
4208  case Intrinsic::nvvm_suld_2d_array_i16_trap:
4209  case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
4210  case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
4211  case Intrinsic::nvvm_suld_3d_i16_trap:
4212  case Intrinsic::nvvm_suld_3d_v2i16_trap:
4213  case Intrinsic::nvvm_suld_3d_v4i16_trap:
4214  case Intrinsic::nvvm_suld_1d_i16_zero:
4215  case Intrinsic::nvvm_suld_1d_v2i16_zero:
4216  case Intrinsic::nvvm_suld_1d_v4i16_zero:
4217  case Intrinsic::nvvm_suld_1d_array_i16_zero:
4218  case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
4219  case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
4220  case Intrinsic::nvvm_suld_2d_i16_zero:
4221  case Intrinsic::nvvm_suld_2d_v2i16_zero:
4222  case Intrinsic::nvvm_suld_2d_v4i16_zero:
4223  case Intrinsic::nvvm_suld_2d_array_i16_zero:
4224  case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
4225  case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
4226  case Intrinsic::nvvm_suld_3d_i16_zero:
4227  case Intrinsic::nvvm_suld_3d_v2i16_zero:
4228  case Intrinsic::nvvm_suld_3d_v4i16_zero:
4229  Info.opc = getOpcForSurfaceInstr(Intrinsic);
4230  Info.memVT = MVT::i16;
4231  Info.ptrVal = nullptr;
4232  Info.offset = 0;
4234  Info.align = Align(16);
4235  return true;
4236 
4237  case Intrinsic::nvvm_suld_1d_i32_clamp:
4238  case Intrinsic::nvvm_suld_1d_v2i32_clamp:
4239  case Intrinsic::nvvm_suld_1d_v4i32_clamp:
4240  case Intrinsic::nvvm_suld_1d_array_i32_clamp:
4241  case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
4242  case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
4243  case Intrinsic::nvvm_suld_2d_i32_clamp:
4244  case Intrinsic::nvvm_suld_2d_v2i32_clamp:
4245  case Intrinsic::nvvm_suld_2d_v4i32_clamp:
4246  case Intrinsic::nvvm_suld_2d_array_i32_clamp:
4247  case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
4248  case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
4249  case Intrinsic::nvvm_suld_3d_i32_clamp:
4250  case Intrinsic::nvvm_suld_3d_v2i32_clamp:
4251  case Intrinsic::nvvm_suld_3d_v4i32_clamp:
4252  case Intrinsic::nvvm_suld_1d_i32_trap:
4253  case Intrinsic::nvvm_suld_1d_v2i32_trap:
4254  case Intrinsic::nvvm_suld_1d_v4i32_trap:
4255  case Intrinsic::nvvm_suld_1d_array_i32_trap:
4256  case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
4257  case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
4258  case Intrinsic::nvvm_suld_2d_i32_trap:
4259  case Intrinsic::nvvm_suld_2d_v2i32_trap:
4260  case Intrinsic::nvvm_suld_2d_v4i32_trap:
4261  case Intrinsic::nvvm_suld_2d_array_i32_trap:
4262  case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
4263  case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
4264  case Intrinsic::nvvm_suld_3d_i32_trap:
4265  case Intrinsic::nvvm_suld_3d_v2i32_trap:
4266  case Intrinsic::nvvm_suld_3d_v4i32_trap:
4267  case Intrinsic::nvvm_suld_1d_i32_zero:
4268  case Intrinsic::nvvm_suld_1d_v2i32_zero:
4269  case Intrinsic::nvvm_suld_1d_v4i32_zero:
4270  case Intrinsic::nvvm_suld_1d_array_i32_zero:
4271  case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
4272  case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
4273  case Intrinsic::nvvm_suld_2d_i32_zero:
4274  case Intrinsic::nvvm_suld_2d_v2i32_zero:
4275  case Intrinsic::nvvm_suld_2d_v4i32_zero:
4276  case Intrinsic::nvvm_suld_2d_array_i32_zero:
4277  case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
4278  case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
4279  case Intrinsic::nvvm_suld_3d_i32_zero:
4280  case Intrinsic::nvvm_suld_3d_v2i32_zero:
4281  case Intrinsic::nvvm_suld_3d_v4i32_zero:
4282  Info.opc = getOpcForSurfaceInstr(Intrinsic);
4283  Info.memVT = MVT::i32;
4284  Info.ptrVal = nullptr;
4285  Info.offset = 0;
4287  Info.align = Align(16);
4288  return true;
4289 
4290  case Intrinsic::nvvm_suld_1d_i64_clamp:
4291  case Intrinsic::nvvm_suld_1d_v2i64_clamp:
4292  case Intrinsic::nvvm_suld_1d_array_i64_clamp:
4293  case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
4294  case Intrinsic::nvvm_suld_2d_i64_clamp:
4295  case Intrinsic::nvvm_suld_2d_v2i64_clamp:
4296  case Intrinsic::nvvm_suld_2d_array_i64_clamp:
4297  case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
4298  case Intrinsic::nvvm_suld_3d_i64_clamp:
4299  case Intrinsic::nvvm_suld_3d_v2i64_clamp:
4300  case Intrinsic::nvvm_suld_1d_i64_trap:
4301  case Intrinsic::nvvm_suld_1d_v2i64_trap:
4302  case Intrinsic::nvvm_suld_1d_array_i64_trap:
4303  case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
4304  case Intrinsic::nvvm_suld_2d_i64_trap:
4305  case Intrinsic::nvvm_suld_2d_v2i64_trap:
4306  case Intrinsic::nvvm_suld_2d_array_i64_trap:
4307  case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
4308  case Intrinsic::nvvm_suld_3d_i64_trap:
4309  case Intrinsic::nvvm_suld_3d_v2i64_trap:
4310  case Intrinsic::nvvm_suld_1d_i64_zero:
4311  case Intrinsic::nvvm_suld_1d_v2i64_zero:
4312  case Intrinsic::nvvm_suld_1d_array_i64_zero:
4313  case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
4314  case Intrinsic::nvvm_suld_2d_i64_zero:
4315  case Intrinsic::nvvm_suld_2d_v2i64_zero:
4316  case Intrinsic::nvvm_suld_2d_array_i64_zero:
4317  case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
4318  case Intrinsic::nvvm_suld_3d_i64_zero:
4319  case Intrinsic::nvvm_suld_3d_v2i64_zero:
4320  Info.opc = getOpcForSurfaceInstr(Intrinsic);
4321  Info.memVT = MVT::i64;
4322  Info.ptrVal = nullptr;
4323  Info.offset = 0;
4325  Info.align = Align(16);
4326  return true;
4327  }
4328  return false;
4329 }
4330 
4331 /// getFunctionParamOptimizedAlign - since function arguments are passed via
4332 /// .param space, we may want to increase their alignment in a way that
4333 /// ensures that we can effectively vectorize their loads & stores. We can
4334 /// increase alignment only if the function has internal or has private
4335 /// linkage as for other linkage types callers may already rely on default
4336 /// alignment. To allow using 128-bit vectorized loads/stores, this function
4337 /// ensures that alignment is 16 or greater.
4339  const Function *F, Type *ArgTy, const DataLayout &DL) const {
4340  const uint64_t ABITypeAlign = DL.getABITypeAlign(ArgTy).value();
4341 
4342  // If a function has linkage different from internal or private, we
4343  // must use default ABI alignment as external users rely on it. Same
4344  // for a function that may be called from a function pointer.
4345  if (!F || !F->hasLocalLinkage() ||
4346  F->hasAddressTaken(/*Users=*/nullptr,
4347  /*IgnoreCallbackUses=*/false,
4348  /*IgnoreAssumeLikeCalls=*/true,
4349  /*IgnoreLLVMUsed=*/true))
4350  return Align(ABITypeAlign);
4351 
4352  assert(!isKernelFunction(*F) && "Expect kernels to have non-local linkage");
4353  return Align(std::max(uint64_t(16), ABITypeAlign));
4354 }
4355 
4356 /// isLegalAddressingMode - Return true if the addressing mode represented
4357 /// by AM is legal for this target, for a load/store of the specified type.
4358 /// Used to guide target specific optimizations, like loop strength reduction
4359 /// (LoopStrengthReduce.cpp) and memory optimization for address mode
4360 /// (CodeGenPrepare.cpp)
4362  const AddrMode &AM, Type *Ty,
4363  unsigned AS, Instruction *I) const {
4364  // AddrMode - This represents an addressing mode of:
4365  // BaseGV + BaseOffs + BaseReg + Scale*ScaleReg
4366  //
4367  // The legal address modes are
4368  // - [avar]
4369  // - [areg]
4370  // - [areg+immoff]
4371  // - [immAddr]
4372 
4373  if (AM.BaseGV) {
4374  return !AM.BaseOffs && !AM.HasBaseReg && !AM.Scale;
4375  }
4376 
4377  switch (AM.Scale) {
4378  case 0: // "r", "r+i" or "i" is allowed
4379  break;
4380  case 1:
4381  if (AM.HasBaseReg) // "r+r+i" or "r+r" is not allowed.
4382  return false;
4383  // Otherwise we have r+i.
4384  break;
4385  default:
4386  // No scale > 1 is allowed
4387  return false;
4388  }
4389  return true;
4390 }
4391 
4392 //===----------------------------------------------------------------------===//
4393 // NVPTX Inline Assembly Support
4394 //===----------------------------------------------------------------------===//
4395 
4396 /// getConstraintType - Given a constraint letter, return the type of
4397 /// constraint it is for this target.
4400  if (Constraint.size() == 1) {
4401  switch (Constraint[0]) {
4402  default:
4403  break;
4404  case 'b':
4405  case 'r':
4406  case 'h':
4407  case 'c':
4408  case 'l':
4409  case 'f':
4410  case 'd':
4411  case '0':
4412  case 'N':
4413  return C_RegisterClass;
4414  }
4415  }
4416  return TargetLowering::getConstraintType(Constraint);
4417 }
4418 
4419 std::pair<unsigned, const TargetRegisterClass *>
4421  StringRef Constraint,
4422  MVT VT) const {
4423  if (Constraint.size() == 1) {
4424  switch (Constraint[0]) {
4425  case 'b':
4426  return std::make_pair(0U, &NVPTX::Int1RegsRegClass);
4427  case 'c':
4428  return std::make_pair(0U, &NVPTX::Int16RegsRegClass);
4429  case 'h':
4430  return std::make_pair(0U, &NVPTX::Int16RegsRegClass);
4431  case 'r':
4432  return std::make_pair(0U, &NVPTX::Int32RegsRegClass);
4433  case 'l':
4434  case 'N':
4435  return std::make_pair(0U, &NVPTX::Int64RegsRegClass);
4436  case 'f':
4437  return std::make_pair(0U, &NVPTX::Float32RegsRegClass);
4438  case 'd':
4439  return std::make_pair(0U, &NVPTX::Float64RegsRegClass);
4440  }
4441  }
4442  return TargetLowering::getRegForInlineAsmConstraint(TRI, Constraint, VT);
4443 }
4444 
4445 //===----------------------------------------------------------------------===//
4446 // NVPTX DAG Combining
4447 //===----------------------------------------------------------------------===//
4448 
4450  CodeGenOpt::Level OptLevel) const {
4451  // Always honor command-line argument
4452  if (FMAContractLevelOpt.getNumOccurrences() > 0)
4453  return FMAContractLevelOpt > 0;
4454 
4455  // Do not contract if we're not optimizing the code.
4456  if (OptLevel == 0)
4457  return false;
4458 
4459  // Honor TargetOptions flags that explicitly say fusion is okay.
4461  return true;
4462 
4463  return allowUnsafeFPMath(MF);
4464 }
4465 
4467  // Honor TargetOptions flags that explicitly say unsafe math is okay.
4468  if (MF.getTarget().Options.UnsafeFPMath)
4469  return true;
4470 
4471  // Allow unsafe math if unsafe-fp-math attribute explicitly says so.
4472  const Function &F = MF.getFunction();
4473  return F.getFnAttribute("unsafe-fp-math").getValueAsBool();
4474 }
4475 
4476 /// PerformADDCombineWithOperands - Try DAG combinations for an ADD with
4477 /// operands N0 and N1. This is a helper for PerformADDCombine that is
4478 /// called with the default ope