CoCalc -- NVPTXISelDAGToDAG.cpp

GitHub Repository: freebsd/freebsd-src
Path: blob/main/contrib/llvm-project/llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
³⁵²⁷¹ views
1
//===-- NVPTXISelDAGToDAG.cpp - A dag to dag inst selector for NVPTX ------===//
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 an instruction selector for the NVPTX target.
10
//
11
//===----------------------------------------------------------------------===//
12

13
#include "NVPTXISelDAGToDAG.h"
14
#include "MCTargetDesc/NVPTXBaseInfo.h"
15
#include "NVPTXUtilities.h"
16
#include "llvm/Analysis/ValueTracking.h"
17
#include "llvm/CodeGen/ISDOpcodes.h"
18
#include "llvm/IR/GlobalValue.h"
19
#include "llvm/IR/Instructions.h"
20
#include "llvm/IR/IntrinsicsNVPTX.h"
21
#include "llvm/Support/AtomicOrdering.h"
22
#include "llvm/Support/CommandLine.h"
23
#include "llvm/Support/Debug.h"
24
#include "llvm/Support/ErrorHandling.h"
25
#include "llvm/Support/raw_ostream.h"
26
#include "llvm/Target/TargetIntrinsicInfo.h"
27

28
using namespace llvm;
29

30
#define DEBUG_TYPE "nvptx-isel"
31
#define PASS_NAME "NVPTX DAG->DAG Pattern Instruction Selection"
32

33
static cl::opt<bool>
34
    EnableRsqrtOpt("nvptx-rsqrt-approx-opt", cl::init(true), cl::Hidden,
35
                   cl::desc("Enable reciprocal sqrt optimization"));
36

37
/// createNVPTXISelDag - This pass converts a legalized DAG into a
38
/// NVPTX-specific DAG, ready for instruction scheduling.
39
FunctionPass *llvm::createNVPTXISelDag(NVPTXTargetMachine &TM,
40
                                       llvm::CodeGenOptLevel OptLevel) {
41
  return new NVPTXDAGToDAGISelLegacy(TM, OptLevel);
42
}
43

44
NVPTXDAGToDAGISelLegacy::NVPTXDAGToDAGISelLegacy(NVPTXTargetMachine &tm,
45
                                                 CodeGenOptLevel OptLevel)
46
    : SelectionDAGISelLegacy(
47
          ID, std::make_unique<NVPTXDAGToDAGISel>(tm, OptLevel)) {}
48

49
char NVPTXDAGToDAGISelLegacy::ID = 0;
50

51
INITIALIZE_PASS(NVPTXDAGToDAGISelLegacy, DEBUG_TYPE, PASS_NAME, false, false)
52

53
NVPTXDAGToDAGISel::NVPTXDAGToDAGISel(NVPTXTargetMachine &tm,
54
                                     CodeGenOptLevel OptLevel)
55
    : SelectionDAGISel(tm, OptLevel), TM(tm) {
56
  doMulWide = (OptLevel > CodeGenOptLevel::None);
57
}
58

59
bool NVPTXDAGToDAGISel::runOnMachineFunction(MachineFunction &MF) {
60
  Subtarget = &MF.getSubtarget<NVPTXSubtarget>();
61
  return SelectionDAGISel::runOnMachineFunction(MF);
62
}
63

64
int NVPTXDAGToDAGISel::getDivF32Level() const {
65
  return Subtarget->getTargetLowering()->getDivF32Level();
66
}
67

68
bool NVPTXDAGToDAGISel::usePrecSqrtF32() const {
69
  return Subtarget->getTargetLowering()->usePrecSqrtF32();
70
}
71

72
bool NVPTXDAGToDAGISel::useF32FTZ() const {
73
  return Subtarget->getTargetLowering()->useF32FTZ(*MF);
74
}
75

76
bool NVPTXDAGToDAGISel::allowFMA() const {
77
  const NVPTXTargetLowering *TL = Subtarget->getTargetLowering();
78
  return TL->allowFMA(*MF, OptLevel);
79
}
80

81
bool NVPTXDAGToDAGISel::allowUnsafeFPMath() const {
82
  const NVPTXTargetLowering *TL = Subtarget->getTargetLowering();
83
  return TL->allowUnsafeFPMath(*MF);
84
}
85

86
bool NVPTXDAGToDAGISel::doRsqrtOpt() const { return EnableRsqrtOpt; }
87

88
/// Select - Select instructions not customized! Used for
89
/// expanded, promoted and normal instructions.
90
void NVPTXDAGToDAGISel::Select(SDNode *N) {
91

92
  if (N->isMachineOpcode()) {
93
    N->setNodeId(-1);
94
    return; // Already selected.
95
  }
96

97
  switch (N->getOpcode()) {
98
  case ISD::LOAD:
99
  case ISD::ATOMIC_LOAD:
100
    if (tryLoad(N))
101
      return;
102
    break;
103
  case ISD::STORE:
104
  case ISD::ATOMIC_STORE:
105
    if (tryStore(N))
106
      return;
107
    break;
108
  case ISD::EXTRACT_VECTOR_ELT:
109
    if (tryEXTRACT_VECTOR_ELEMENT(N))
110
      return;
111
    break;
112
  case NVPTXISD::SETP_F16X2:
113
    SelectSETP_F16X2(N);
114
    return;
115
  case NVPTXISD::SETP_BF16X2:
116
    SelectSETP_BF16X2(N);
117
    return;
118
  case NVPTXISD::LoadV2:
119
  case NVPTXISD::LoadV4:
120
    if (tryLoadVector(N))
121
      return;
122
    break;
123
  case NVPTXISD::LDGV2:
124
  case NVPTXISD::LDGV4:
125
  case NVPTXISD::LDUV2:
126
  case NVPTXISD::LDUV4:
127
    if (tryLDGLDU(N))
128
      return;
129
    break;
130
  case NVPTXISD::StoreV2:
131
  case NVPTXISD::StoreV4:
132
    if (tryStoreVector(N))
133
      return;
134
    break;
135
  case NVPTXISD::LoadParam:
136
  case NVPTXISD::LoadParamV2:
137
  case NVPTXISD::LoadParamV4:
138
    if (tryLoadParam(N))
139
      return;
140
    break;
141
  case NVPTXISD::StoreRetval:
142
  case NVPTXISD::StoreRetvalV2:
143
  case NVPTXISD::StoreRetvalV4:
144
    if (tryStoreRetval(N))
145
      return;
146
    break;
147
  case NVPTXISD::StoreParam:
148
  case NVPTXISD::StoreParamV2:
149
  case NVPTXISD::StoreParamV4:
150
  case NVPTXISD::StoreParamS32:
151
  case NVPTXISD::StoreParamU32:
152
    if (tryStoreParam(N))
153
      return;
154
    break;
155
  case ISD::INTRINSIC_WO_CHAIN:
156
    if (tryIntrinsicNoChain(N))
157
      return;
158
    break;
159
  case ISD::INTRINSIC_W_CHAIN:
160
    if (tryIntrinsicChain(N))
161
      return;
162
    break;
163
  case NVPTXISD::Tex1DFloatS32:
164
  case NVPTXISD::Tex1DFloatFloat:
165
  case NVPTXISD::Tex1DFloatFloatLevel:
166
  case NVPTXISD::Tex1DFloatFloatGrad:
167
  case NVPTXISD::Tex1DS32S32:
168
  case NVPTXISD::Tex1DS32Float:
169
  case NVPTXISD::Tex1DS32FloatLevel:
170
  case NVPTXISD::Tex1DS32FloatGrad:
171
  case NVPTXISD::Tex1DU32S32:
172
  case NVPTXISD::Tex1DU32Float:
173
  case NVPTXISD::Tex1DU32FloatLevel:
174
  case NVPTXISD::Tex1DU32FloatGrad:
175
  case NVPTXISD::Tex1DArrayFloatS32:
176
  case NVPTXISD::Tex1DArrayFloatFloat:
177
  case NVPTXISD::Tex1DArrayFloatFloatLevel:
178
  case NVPTXISD::Tex1DArrayFloatFloatGrad:
179
  case NVPTXISD::Tex1DArrayS32S32:
180
  case NVPTXISD::Tex1DArrayS32Float:
181
  case NVPTXISD::Tex1DArrayS32FloatLevel:
182
  case NVPTXISD::Tex1DArrayS32FloatGrad:
183
  case NVPTXISD::Tex1DArrayU32S32:
184
  case NVPTXISD::Tex1DArrayU32Float:
185
  case NVPTXISD::Tex1DArrayU32FloatLevel:
186
  case NVPTXISD::Tex1DArrayU32FloatGrad:
187
  case NVPTXISD::Tex2DFloatS32:
188
  case NVPTXISD::Tex2DFloatFloat:
189
  case NVPTXISD::Tex2DFloatFloatLevel:
190
  case NVPTXISD::Tex2DFloatFloatGrad:
191
  case NVPTXISD::Tex2DS32S32:
192
  case NVPTXISD::Tex2DS32Float:
193
  case NVPTXISD::Tex2DS32FloatLevel:
194
  case NVPTXISD::Tex2DS32FloatGrad:
195
  case NVPTXISD::Tex2DU32S32:
196
  case NVPTXISD::Tex2DU32Float:
197
  case NVPTXISD::Tex2DU32FloatLevel:
198
  case NVPTXISD::Tex2DU32FloatGrad:
199
  case NVPTXISD::Tex2DArrayFloatS32:
200
  case NVPTXISD::Tex2DArrayFloatFloat:
201
  case NVPTXISD::Tex2DArrayFloatFloatLevel:
202
  case NVPTXISD::Tex2DArrayFloatFloatGrad:
203
  case NVPTXISD::Tex2DArrayS32S32:
204
  case NVPTXISD::Tex2DArrayS32Float:
205
  case NVPTXISD::Tex2DArrayS32FloatLevel:
206
  case NVPTXISD::Tex2DArrayS32FloatGrad:
207
  case NVPTXISD::Tex2DArrayU32S32:
208
  case NVPTXISD::Tex2DArrayU32Float:
209
  case NVPTXISD::Tex2DArrayU32FloatLevel:
210
  case NVPTXISD::Tex2DArrayU32FloatGrad:
211
  case NVPTXISD::Tex3DFloatS32:
212
  case NVPTXISD::Tex3DFloatFloat:
213
  case NVPTXISD::Tex3DFloatFloatLevel:
214
  case NVPTXISD::Tex3DFloatFloatGrad:
215
  case NVPTXISD::Tex3DS32S32:
216
  case NVPTXISD::Tex3DS32Float:
217
  case NVPTXISD::Tex3DS32FloatLevel:
218
  case NVPTXISD::Tex3DS32FloatGrad:
219
  case NVPTXISD::Tex3DU32S32:
220
  case NVPTXISD::Tex3DU32Float:
221
  case NVPTXISD::Tex3DU32FloatLevel:
222
  case NVPTXISD::Tex3DU32FloatGrad:
223
  case NVPTXISD::TexCubeFloatFloat:
224
  case NVPTXISD::TexCubeFloatFloatLevel:
225
  case NVPTXISD::TexCubeS32Float:
226
  case NVPTXISD::TexCubeS32FloatLevel:
227
  case NVPTXISD::TexCubeU32Float:
228
  case NVPTXISD::TexCubeU32FloatLevel:
229
  case NVPTXISD::TexCubeArrayFloatFloat:
230
  case NVPTXISD::TexCubeArrayFloatFloatLevel:
231
  case NVPTXISD::TexCubeArrayS32Float:
232
  case NVPTXISD::TexCubeArrayS32FloatLevel:
233
  case NVPTXISD::TexCubeArrayU32Float:
234
  case NVPTXISD::TexCubeArrayU32FloatLevel:
235
  case NVPTXISD::Tld4R2DFloatFloat:
236
  case NVPTXISD::Tld4G2DFloatFloat:
237
  case NVPTXISD::Tld4B2DFloatFloat:
238
  case NVPTXISD::Tld4A2DFloatFloat:
239
  case NVPTXISD::Tld4R2DS64Float:
240
  case NVPTXISD::Tld4G2DS64Float:
241
  case NVPTXISD::Tld4B2DS64Float:
242
  case NVPTXISD::Tld4A2DS64Float:
243
  case NVPTXISD::Tld4R2DU64Float:
244
  case NVPTXISD::Tld4G2DU64Float:
245
  case NVPTXISD::Tld4B2DU64Float:
246
  case NVPTXISD::Tld4A2DU64Float:
247
  case NVPTXISD::TexUnified1DFloatS32:
248
  case NVPTXISD::TexUnified1DFloatFloat:
249
  case NVPTXISD::TexUnified1DFloatFloatLevel:
250
  case NVPTXISD::TexUnified1DFloatFloatGrad:
251
  case NVPTXISD::TexUnified1DS32S32:
252
  case NVPTXISD::TexUnified1DS32Float:
253
  case NVPTXISD::TexUnified1DS32FloatLevel:
254
  case NVPTXISD::TexUnified1DS32FloatGrad:
255
  case NVPTXISD::TexUnified1DU32S32:
256
  case NVPTXISD::TexUnified1DU32Float:
257
  case NVPTXISD::TexUnified1DU32FloatLevel:
258
  case NVPTXISD::TexUnified1DU32FloatGrad:
259
  case NVPTXISD::TexUnified1DArrayFloatS32:
260
  case NVPTXISD::TexUnified1DArrayFloatFloat:
261
  case NVPTXISD::TexUnified1DArrayFloatFloatLevel:
262
  case NVPTXISD::TexUnified1DArrayFloatFloatGrad:
263
  case NVPTXISD::TexUnified1DArrayS32S32:
264
  case NVPTXISD::TexUnified1DArrayS32Float:
265
  case NVPTXISD::TexUnified1DArrayS32FloatLevel:
266
  case NVPTXISD::TexUnified1DArrayS32FloatGrad:
267
  case NVPTXISD::TexUnified1DArrayU32S32:
268
  case NVPTXISD::TexUnified1DArrayU32Float:
269
  case NVPTXISD::TexUnified1DArrayU32FloatLevel:
270
  case NVPTXISD::TexUnified1DArrayU32FloatGrad:
271
  case NVPTXISD::TexUnified2DFloatS32:
272
  case NVPTXISD::TexUnified2DFloatFloat:
273
  case NVPTXISD::TexUnified2DFloatFloatLevel:
274
  case NVPTXISD::TexUnified2DFloatFloatGrad:
275
  case NVPTXISD::TexUnified2DS32S32:
276
  case NVPTXISD::TexUnified2DS32Float:
277
  case NVPTXISD::TexUnified2DS32FloatLevel:
278
  case NVPTXISD::TexUnified2DS32FloatGrad:
279
  case NVPTXISD::TexUnified2DU32S32:
280
  case NVPTXISD::TexUnified2DU32Float:
281
  case NVPTXISD::TexUnified2DU32FloatLevel:
282
  case NVPTXISD::TexUnified2DU32FloatGrad:
283
  case NVPTXISD::TexUnified2DArrayFloatS32:
284
  case NVPTXISD::TexUnified2DArrayFloatFloat:
285
  case NVPTXISD::TexUnified2DArrayFloatFloatLevel:
286
  case NVPTXISD::TexUnified2DArrayFloatFloatGrad:
287
  case NVPTXISD::TexUnified2DArrayS32S32:
288
  case NVPTXISD::TexUnified2DArrayS32Float:
289
  case NVPTXISD::TexUnified2DArrayS32FloatLevel:
290
  case NVPTXISD::TexUnified2DArrayS32FloatGrad:
291
  case NVPTXISD::TexUnified2DArrayU32S32:
292
  case NVPTXISD::TexUnified2DArrayU32Float:
293
  case NVPTXISD::TexUnified2DArrayU32FloatLevel:
294
  case NVPTXISD::TexUnified2DArrayU32FloatGrad:
295
  case NVPTXISD::TexUnified3DFloatS32:
296
  case NVPTXISD::TexUnified3DFloatFloat:
297
  case NVPTXISD::TexUnified3DFloatFloatLevel:
298
  case NVPTXISD::TexUnified3DFloatFloatGrad:
299
  case NVPTXISD::TexUnified3DS32S32:
300
  case NVPTXISD::TexUnified3DS32Float:
301
  case NVPTXISD::TexUnified3DS32FloatLevel:
302
  case NVPTXISD::TexUnified3DS32FloatGrad:
303
  case NVPTXISD::TexUnified3DU32S32:
304
  case NVPTXISD::TexUnified3DU32Float:
305
  case NVPTXISD::TexUnified3DU32FloatLevel:
306
  case NVPTXISD::TexUnified3DU32FloatGrad:
307
  case NVPTXISD::TexUnifiedCubeFloatFloat:
308
  case NVPTXISD::TexUnifiedCubeFloatFloatLevel:
309
  case NVPTXISD::TexUnifiedCubeS32Float:
310
  case NVPTXISD::TexUnifiedCubeS32FloatLevel:
311
  case NVPTXISD::TexUnifiedCubeU32Float:
312
  case NVPTXISD::TexUnifiedCubeU32FloatLevel:
313
  case NVPTXISD::TexUnifiedCubeArrayFloatFloat:
314
  case NVPTXISD::TexUnifiedCubeArrayFloatFloatLevel:
315
  case NVPTXISD::TexUnifiedCubeArrayS32Float:
316
  case NVPTXISD::TexUnifiedCubeArrayS32FloatLevel:
317
  case NVPTXISD::TexUnifiedCubeArrayU32Float:
318
  case NVPTXISD::TexUnifiedCubeArrayU32FloatLevel:
319
  case NVPTXISD::TexUnifiedCubeFloatFloatGrad:
320
  case NVPTXISD::TexUnifiedCubeS32FloatGrad:
321
  case NVPTXISD::TexUnifiedCubeU32FloatGrad:
322
  case NVPTXISD::TexUnifiedCubeArrayFloatFloatGrad:
323
  case NVPTXISD::TexUnifiedCubeArrayS32FloatGrad:
324
  case NVPTXISD::TexUnifiedCubeArrayU32FloatGrad:
325
  case NVPTXISD::Tld4UnifiedR2DFloatFloat:
326
  case NVPTXISD::Tld4UnifiedG2DFloatFloat:
327
  case NVPTXISD::Tld4UnifiedB2DFloatFloat:
328
  case NVPTXISD::Tld4UnifiedA2DFloatFloat:
329
  case NVPTXISD::Tld4UnifiedR2DS64Float:
330
  case NVPTXISD::Tld4UnifiedG2DS64Float:
331
  case NVPTXISD::Tld4UnifiedB2DS64Float:
332
  case NVPTXISD::Tld4UnifiedA2DS64Float:
333
  case NVPTXISD::Tld4UnifiedR2DU64Float:
334
  case NVPTXISD::Tld4UnifiedG2DU64Float:
335
  case NVPTXISD::Tld4UnifiedB2DU64Float:
336
  case NVPTXISD::Tld4UnifiedA2DU64Float:
337
    if (tryTextureIntrinsic(N))
338
      return;
339
    break;
340
  case NVPTXISD::Suld1DI8Clamp:
341
  case NVPTXISD::Suld1DI16Clamp:
342
  case NVPTXISD::Suld1DI32Clamp:
343
  case NVPTXISD::Suld1DI64Clamp:
344
  case NVPTXISD::Suld1DV2I8Clamp:
345
  case NVPTXISD::Suld1DV2I16Clamp:
346
  case NVPTXISD::Suld1DV2I32Clamp:
347
  case NVPTXISD::Suld1DV2I64Clamp:
348
  case NVPTXISD::Suld1DV4I8Clamp:
349
  case NVPTXISD::Suld1DV4I16Clamp:
350
  case NVPTXISD::Suld1DV4I32Clamp:
351
  case NVPTXISD::Suld1DArrayI8Clamp:
352
  case NVPTXISD::Suld1DArrayI16Clamp:
353
  case NVPTXISD::Suld1DArrayI32Clamp:
354
  case NVPTXISD::Suld1DArrayI64Clamp:
355
  case NVPTXISD::Suld1DArrayV2I8Clamp:
356
  case NVPTXISD::Suld1DArrayV2I16Clamp:
357
  case NVPTXISD::Suld1DArrayV2I32Clamp:
358
  case NVPTXISD::Suld1DArrayV2I64Clamp:
359
  case NVPTXISD::Suld1DArrayV4I8Clamp:
360
  case NVPTXISD::Suld1DArrayV4I16Clamp:
361
  case NVPTXISD::Suld1DArrayV4I32Clamp:
362
  case NVPTXISD::Suld2DI8Clamp:
363
  case NVPTXISD::Suld2DI16Clamp:
364
  case NVPTXISD::Suld2DI32Clamp:
365
  case NVPTXISD::Suld2DI64Clamp:
366
  case NVPTXISD::Suld2DV2I8Clamp:
367
  case NVPTXISD::Suld2DV2I16Clamp:
368
  case NVPTXISD::Suld2DV2I32Clamp:
369
  case NVPTXISD::Suld2DV2I64Clamp:
370
  case NVPTXISD::Suld2DV4I8Clamp:
371
  case NVPTXISD::Suld2DV4I16Clamp:
372
  case NVPTXISD::Suld2DV4I32Clamp:
373
  case NVPTXISD::Suld2DArrayI8Clamp:
374
  case NVPTXISD::Suld2DArrayI16Clamp:
375
  case NVPTXISD::Suld2DArrayI32Clamp:
376
  case NVPTXISD::Suld2DArrayI64Clamp:
377
  case NVPTXISD::Suld2DArrayV2I8Clamp:
378
  case NVPTXISD::Suld2DArrayV2I16Clamp:
379
  case NVPTXISD::Suld2DArrayV2I32Clamp:
380
  case NVPTXISD::Suld2DArrayV2I64Clamp:
381
  case NVPTXISD::Suld2DArrayV4I8Clamp:
382
  case NVPTXISD::Suld2DArrayV4I16Clamp:
383
  case NVPTXISD::Suld2DArrayV4I32Clamp:
384
  case NVPTXISD::Suld3DI8Clamp:
385
  case NVPTXISD::Suld3DI16Clamp:
386
  case NVPTXISD::Suld3DI32Clamp:
387
  case NVPTXISD::Suld3DI64Clamp:
388
  case NVPTXISD::Suld3DV2I8Clamp:
389
  case NVPTXISD::Suld3DV2I16Clamp:
390
  case NVPTXISD::Suld3DV2I32Clamp:
391
  case NVPTXISD::Suld3DV2I64Clamp:
392
  case NVPTXISD::Suld3DV4I8Clamp:
393
  case NVPTXISD::Suld3DV4I16Clamp:
394
  case NVPTXISD::Suld3DV4I32Clamp:
395
  case NVPTXISD::Suld1DI8Trap:
396
  case NVPTXISD::Suld1DI16Trap:
397
  case NVPTXISD::Suld1DI32Trap:
398
  case NVPTXISD::Suld1DI64Trap:
399
  case NVPTXISD::Suld1DV2I8Trap:
400
  case NVPTXISD::Suld1DV2I16Trap:
401
  case NVPTXISD::Suld1DV2I32Trap:
402
  case NVPTXISD::Suld1DV2I64Trap:
403
  case NVPTXISD::Suld1DV4I8Trap:
404
  case NVPTXISD::Suld1DV4I16Trap:
405
  case NVPTXISD::Suld1DV4I32Trap:
406
  case NVPTXISD::Suld1DArrayI8Trap:
407
  case NVPTXISD::Suld1DArrayI16Trap:
408
  case NVPTXISD::Suld1DArrayI32Trap:
409
  case NVPTXISD::Suld1DArrayI64Trap:
410
  case NVPTXISD::Suld1DArrayV2I8Trap:
411
  case NVPTXISD::Suld1DArrayV2I16Trap:
412
  case NVPTXISD::Suld1DArrayV2I32Trap:
413
  case NVPTXISD::Suld1DArrayV2I64Trap:
414
  case NVPTXISD::Suld1DArrayV4I8Trap:
415
  case NVPTXISD::Suld1DArrayV4I16Trap:
416
  case NVPTXISD::Suld1DArrayV4I32Trap:
417
  case NVPTXISD::Suld2DI8Trap:
418
  case NVPTXISD::Suld2DI16Trap:
419
  case NVPTXISD::Suld2DI32Trap:
420
  case NVPTXISD::Suld2DI64Trap:
421
  case NVPTXISD::Suld2DV2I8Trap:
422
  case NVPTXISD::Suld2DV2I16Trap:
423
  case NVPTXISD::Suld2DV2I32Trap:
424
  case NVPTXISD::Suld2DV2I64Trap:
425
  case NVPTXISD::Suld2DV4I8Trap:
426
  case NVPTXISD::Suld2DV4I16Trap:
427
  case NVPTXISD::Suld2DV4I32Trap:
428
  case NVPTXISD::Suld2DArrayI8Trap:
429
  case NVPTXISD::Suld2DArrayI16Trap:
430
  case NVPTXISD::Suld2DArrayI32Trap:
431
  case NVPTXISD::Suld2DArrayI64Trap:
432
  case NVPTXISD::Suld2DArrayV2I8Trap:
433
  case NVPTXISD::Suld2DArrayV2I16Trap:
434
  case NVPTXISD::Suld2DArrayV2I32Trap:
435
  case NVPTXISD::Suld2DArrayV2I64Trap:
436
  case NVPTXISD::Suld2DArrayV4I8Trap:
437
  case NVPTXISD::Suld2DArrayV4I16Trap:
438
  case NVPTXISD::Suld2DArrayV4I32Trap:
439
  case NVPTXISD::Suld3DI8Trap:
440
  case NVPTXISD::Suld3DI16Trap:
441
  case NVPTXISD::Suld3DI32Trap:
442
  case NVPTXISD::Suld3DI64Trap:
443
  case NVPTXISD::Suld3DV2I8Trap:
444
  case NVPTXISD::Suld3DV2I16Trap:
445
  case NVPTXISD::Suld3DV2I32Trap:
446
  case NVPTXISD::Suld3DV2I64Trap:
447
  case NVPTXISD::Suld3DV4I8Trap:
448
  case NVPTXISD::Suld3DV4I16Trap:
449
  case NVPTXISD::Suld3DV4I32Trap:
450
  case NVPTXISD::Suld1DI8Zero:
451
  case NVPTXISD::Suld1DI16Zero:
452
  case NVPTXISD::Suld1DI32Zero:
453
  case NVPTXISD::Suld1DI64Zero:
454
  case NVPTXISD::Suld1DV2I8Zero:
455
  case NVPTXISD::Suld1DV2I16Zero:
456
  case NVPTXISD::Suld1DV2I32Zero:
457
  case NVPTXISD::Suld1DV2I64Zero:
458
  case NVPTXISD::Suld1DV4I8Zero:
459
  case NVPTXISD::Suld1DV4I16Zero:
460
  case NVPTXISD::Suld1DV4I32Zero:
461
  case NVPTXISD::Suld1DArrayI8Zero:
462
  case NVPTXISD::Suld1DArrayI16Zero:
463
  case NVPTXISD::Suld1DArrayI32Zero:
464
  case NVPTXISD::Suld1DArrayI64Zero:
465
  case NVPTXISD::Suld1DArrayV2I8Zero:
466
  case NVPTXISD::Suld1DArrayV2I16Zero:
467
  case NVPTXISD::Suld1DArrayV2I32Zero:
468
  case NVPTXISD::Suld1DArrayV2I64Zero:
469
  case NVPTXISD::Suld1DArrayV4I8Zero:
470
  case NVPTXISD::Suld1DArrayV4I16Zero:
471
  case NVPTXISD::Suld1DArrayV4I32Zero:
472
  case NVPTXISD::Suld2DI8Zero:
473
  case NVPTXISD::Suld2DI16Zero:
474
  case NVPTXISD::Suld2DI32Zero:
475
  case NVPTXISD::Suld2DI64Zero:
476
  case NVPTXISD::Suld2DV2I8Zero:
477
  case NVPTXISD::Suld2DV2I16Zero:
478
  case NVPTXISD::Suld2DV2I32Zero:
479
  case NVPTXISD::Suld2DV2I64Zero:
480
  case NVPTXISD::Suld2DV4I8Zero:
481
  case NVPTXISD::Suld2DV4I16Zero:
482
  case NVPTXISD::Suld2DV4I32Zero:
483
  case NVPTXISD::Suld2DArrayI8Zero:
484
  case NVPTXISD::Suld2DArrayI16Zero:
485
  case NVPTXISD::Suld2DArrayI32Zero:
486
  case NVPTXISD::Suld2DArrayI64Zero:
487
  case NVPTXISD::Suld2DArrayV2I8Zero:
488
  case NVPTXISD::Suld2DArrayV2I16Zero:
489
  case NVPTXISD::Suld2DArrayV2I32Zero:
490
  case NVPTXISD::Suld2DArrayV2I64Zero:
491
  case NVPTXISD::Suld2DArrayV4I8Zero:
492
  case NVPTXISD::Suld2DArrayV4I16Zero:
493
  case NVPTXISD::Suld2DArrayV4I32Zero:
494
  case NVPTXISD::Suld3DI8Zero:
495
  case NVPTXISD::Suld3DI16Zero:
496
  case NVPTXISD::Suld3DI32Zero:
497
  case NVPTXISD::Suld3DI64Zero:
498
  case NVPTXISD::Suld3DV2I8Zero:
499
  case NVPTXISD::Suld3DV2I16Zero:
500
  case NVPTXISD::Suld3DV2I32Zero:
501
  case NVPTXISD::Suld3DV2I64Zero:
502
  case NVPTXISD::Suld3DV4I8Zero:
503
  case NVPTXISD::Suld3DV4I16Zero:
504
  case NVPTXISD::Suld3DV4I32Zero:
505
    if (trySurfaceIntrinsic(N))
506
      return;
507
    break;
508
  case ISD::AND:
509
  case ISD::SRA:
510
  case ISD::SRL:
511
    // Try to select BFE
512
    if (tryBFE(N))
513
      return;
514
    break;
515
  case ISD::ADDRSPACECAST:
516
    SelectAddrSpaceCast(N);
517
    return;
518
  case ISD::ConstantFP:
519
    if (tryConstantFP(N))
520
      return;
521
    break;
522
  case ISD::CopyToReg: {
523
    if (N->getOperand(1).getValueType() == MVT::i128) {
524
      SelectV2I64toI128(N);
525
      return;
526
    }
527
    break;
528
  }
529
  case ISD::CopyFromReg: {
530
    if (N->getOperand(1).getValueType() == MVT::i128) {
531
      SelectI128toV2I64(N);
532
      return;
533
    }
534
    break;
535
  }
536
  default:
537
    break;
538
  }
539
  SelectCode(N);
540
}
541

542
bool NVPTXDAGToDAGISel::tryIntrinsicChain(SDNode *N) {
543
  unsigned IID = N->getConstantOperandVal(1);
544
  switch (IID) {
545
  default:
546
    return false;
547
  case Intrinsic::nvvm_ldg_global_f:
548
  case Intrinsic::nvvm_ldg_global_i:
549
  case Intrinsic::nvvm_ldg_global_p:
550
  case Intrinsic::nvvm_ldu_global_f:
551
  case Intrinsic::nvvm_ldu_global_i:
552
  case Intrinsic::nvvm_ldu_global_p:
553
    return tryLDGLDU(N);
554
  }
555
}
556

557
// There's no way to specify FP16 and BF16 immediates in .(b)f16 ops, so we
558
// have to load them into an .(b)f16 register first.
559
bool NVPTXDAGToDAGISel::tryConstantFP(SDNode *N) {
560
  if (N->getValueType(0) != MVT::f16 && N->getValueType(0) != MVT::bf16)
561
    return false;
562
  SDValue Val = CurDAG->getTargetConstantFP(
563
      cast<ConstantFPSDNode>(N)->getValueAPF(), SDLoc(N), N->getValueType(0));
564
  SDNode *LoadConstF16 = CurDAG->getMachineNode(
565
      (N->getValueType(0) == MVT::f16 ? NVPTX::LOAD_CONST_F16
566
                                      : NVPTX::LOAD_CONST_BF16),
567
      SDLoc(N), N->getValueType(0), Val);
568
  ReplaceNode(N, LoadConstF16);
569
  return true;
570
}
571

572
// Map ISD:CONDCODE value to appropriate CmpMode expected by
573
// NVPTXInstPrinter::printCmpMode()
574
static unsigned getPTXCmpMode(const CondCodeSDNode &CondCode, bool FTZ) {
575
  using NVPTX::PTXCmpMode::CmpMode;
576
  unsigned PTXCmpMode = [](ISD::CondCode CC) {
577
    switch (CC) {
578
    default:
579
      llvm_unreachable("Unexpected condition code.");
580
    case ISD::SETOEQ:
581
      return CmpMode::EQ;
582
    case ISD::SETOGT:
583
      return CmpMode::GT;
584
    case ISD::SETOGE:
585
      return CmpMode::GE;
586
    case ISD::SETOLT:
587
      return CmpMode::LT;
588
    case ISD::SETOLE:
589
      return CmpMode::LE;
590
    case ISD::SETONE:
591
      return CmpMode::NE;
592
    case ISD::SETO:
593
      return CmpMode::NUM;
594
    case ISD::SETUO:
595
      return CmpMode::NotANumber;
596
    case ISD::SETUEQ:
597
      return CmpMode::EQU;
598
    case ISD::SETUGT:
599
      return CmpMode::GTU;
600
    case ISD::SETUGE:
601
      return CmpMode::GEU;
602
    case ISD::SETULT:
603
      return CmpMode::LTU;
604
    case ISD::SETULE:
605
      return CmpMode::LEU;
606
    case ISD::SETUNE:
607
      return CmpMode::NEU;
608
    case ISD::SETEQ:
609
      return CmpMode::EQ;
610
    case ISD::SETGT:
611
      return CmpMode::GT;
612
    case ISD::SETGE:
613
      return CmpMode::GE;
614
    case ISD::SETLT:
615
      return CmpMode::LT;
616
    case ISD::SETLE:
617
      return CmpMode::LE;
618
    case ISD::SETNE:
619
      return CmpMode::NE;
620
    }
621
  }(CondCode.get());
622

623
  if (FTZ)
624
    PTXCmpMode |= NVPTX::PTXCmpMode::FTZ_FLAG;
625

626
  return PTXCmpMode;
627
}
628

629
bool NVPTXDAGToDAGISel::SelectSETP_F16X2(SDNode *N) {
630
  unsigned PTXCmpMode =
631
      getPTXCmpMode(*cast<CondCodeSDNode>(N->getOperand(2)), useF32FTZ());
632
  SDLoc DL(N);
633
  SDNode *SetP = CurDAG->getMachineNode(
634
      NVPTX::SETP_f16x2rr, DL, MVT::i1, MVT::i1, N->getOperand(0),
635
      N->getOperand(1), CurDAG->getTargetConstant(PTXCmpMode, DL, MVT::i32));
636
  ReplaceNode(N, SetP);
637
  return true;
638
}
639

640
bool NVPTXDAGToDAGISel::SelectSETP_BF16X2(SDNode *N) {
641
  unsigned PTXCmpMode =
642
      getPTXCmpMode(*cast<CondCodeSDNode>(N->getOperand(2)), useF32FTZ());
643
  SDLoc DL(N);
644
  SDNode *SetP = CurDAG->getMachineNode(
645
      NVPTX::SETP_bf16x2rr, DL, MVT::i1, MVT::i1, N->getOperand(0),
646
      N->getOperand(1), CurDAG->getTargetConstant(PTXCmpMode, DL, MVT::i32));
647
  ReplaceNode(N, SetP);
648
  return true;
649
}
650

651
// Find all instances of extract_vector_elt that use this v2f16 vector
652
// and coalesce them into a scattering move instruction.
653
bool NVPTXDAGToDAGISel::tryEXTRACT_VECTOR_ELEMENT(SDNode *N) {
654
  SDValue Vector = N->getOperand(0);
655

656
  // We only care about 16x2 as it's the only real vector type we
657
  // need to deal with.
658
  MVT VT = Vector.getSimpleValueType();
659
  if (!Isv2x16VT(VT))
660
    return false;
661
  // Find and record all uses of this vector that extract element 0 or 1.
662
  SmallVector<SDNode *, 4> E0, E1;
663
  for (auto *U : Vector.getNode()->uses()) {
664
    if (U->getOpcode() != ISD::EXTRACT_VECTOR_ELT)
665
      continue;
666
    if (U->getOperand(0) != Vector)
667
      continue;
668
    if (const ConstantSDNode *IdxConst =
669
            dyn_cast<ConstantSDNode>(U->getOperand(1))) {
670
      if (IdxConst->getZExtValue() == 0)
671
        E0.push_back(U);
672
      else if (IdxConst->getZExtValue() == 1)
673
        E1.push_back(U);
674
      else
675
        llvm_unreachable("Invalid vector index.");
676
    }
677
  }
678

679
  // There's no point scattering f16x2 if we only ever access one
680
  // element of it.
681
  if (E0.empty() || E1.empty())
682
    return false;
683

684
  // Merge (f16 extractelt(V, 0), f16 extractelt(V,1))
685
  // into f16,f16 SplitF16x2(V)
686
  MVT EltVT = VT.getVectorElementType();
687
  SDNode *ScatterOp =
688
      CurDAG->getMachineNode(NVPTX::I32toV2I16, SDLoc(N), EltVT, EltVT, Vector);
689
  for (auto *Node : E0)
690
    ReplaceUses(SDValue(Node, 0), SDValue(ScatterOp, 0));
691
  for (auto *Node : E1)
692
    ReplaceUses(SDValue(Node, 0), SDValue(ScatterOp, 1));
693

694
  return true;
695
}
696

697
static unsigned int getCodeAddrSpace(MemSDNode *N) {
698
  const Value *Src = N->getMemOperand()->getValue();
699

700
  if (!Src)
701
    return NVPTX::PTXLdStInstCode::GENERIC;
702

703
  if (auto *PT = dyn_cast<PointerType>(Src->getType())) {
704
    switch (PT->getAddressSpace()) {
705
    case llvm::ADDRESS_SPACE_LOCAL: return NVPTX::PTXLdStInstCode::LOCAL;
706
    case llvm::ADDRESS_SPACE_GLOBAL: return NVPTX::PTXLdStInstCode::GLOBAL;
707
    case llvm::ADDRESS_SPACE_SHARED: return NVPTX::PTXLdStInstCode::SHARED;
708
    case llvm::ADDRESS_SPACE_GENERIC: return NVPTX::PTXLdStInstCode::GENERIC;
709
    case llvm::ADDRESS_SPACE_PARAM: return NVPTX::PTXLdStInstCode::PARAM;
710
    case llvm::ADDRESS_SPACE_CONST: return NVPTX::PTXLdStInstCode::CONSTANT;
711
    default: break;
712
    }
713
  }
714
  return NVPTX::PTXLdStInstCode::GENERIC;
715
}
716

717
static bool canLowerToLDG(MemSDNode *N, const NVPTXSubtarget &Subtarget,
718
                          unsigned CodeAddrSpace, MachineFunction *F) {
719
  // We use ldg (i.e. ld.global.nc) for invariant loads from the global address
720
  // space.
721
  //
722
  // We have two ways of identifying invariant loads: Loads may be explicitly
723
  // marked as invariant, or we may infer them to be invariant.
724
  //
725
  // We currently infer invariance for loads from
726
  //  - constant global variables, and
727
  //  - kernel function pointer params that are noalias (i.e. __restrict) and
728
  //    never written to.
729
  //
730
  // TODO: Perform a more powerful invariance analysis (ideally IPO, and ideally
731
  // not during the SelectionDAG phase).
732
  //
733
  // TODO: Infer invariance only at -O2.  We still want to use ldg at -O0 for
734
  // explicitly invariant loads because these are how clang tells us to use ldg
735
  // when the user uses a builtin.
736
  if (!Subtarget.hasLDG() || CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL)
737
    return false;
738

739
  if (N->isInvariant())
740
    return true;
741

742
  bool IsKernelFn = isKernelFunction(F->getFunction());
743

744
  // We use getUnderlyingObjects() here instead of getUnderlyingObject() mainly
745
  // because the former looks through phi nodes while the latter does not. We
746
  // need to look through phi nodes to handle pointer induction variables.
747
  SmallVector<const Value *, 8> Objs;
748
  getUnderlyingObjects(N->getMemOperand()->getValue(), Objs);
749

750
  return all_of(Objs, [&](const Value *V) {
751
    if (auto *A = dyn_cast<const Argument>(V))
752
      return IsKernelFn && A->onlyReadsMemory() && A->hasNoAliasAttr();
753
    if (auto *GV = dyn_cast<const GlobalVariable>(V))
754
      return GV->isConstant();
755
    return false;
756
  });
757
}
758

759
bool NVPTXDAGToDAGISel::tryIntrinsicNoChain(SDNode *N) {
760
  unsigned IID = N->getConstantOperandVal(0);
761
  switch (IID) {
762
  default:
763
    return false;
764
  case Intrinsic::nvvm_texsurf_handle_internal:
765
    SelectTexSurfHandle(N);
766
    return true;
767
  }
768
}
769

770
void NVPTXDAGToDAGISel::SelectTexSurfHandle(SDNode *N) {
771
  // Op 0 is the intrinsic ID
772
  SDValue Wrapper = N->getOperand(1);
773
  SDValue GlobalVal = Wrapper.getOperand(0);
774
  ReplaceNode(N, CurDAG->getMachineNode(NVPTX::texsurf_handles, SDLoc(N),
775
                                        MVT::i64, GlobalVal));
776
}
777

778
void NVPTXDAGToDAGISel::SelectAddrSpaceCast(SDNode *N) {
779
  SDValue Src = N->getOperand(0);
780
  AddrSpaceCastSDNode *CastN = cast<AddrSpaceCastSDNode>(N);
781
  unsigned SrcAddrSpace = CastN->getSrcAddressSpace();
782
  unsigned DstAddrSpace = CastN->getDestAddressSpace();
783
  assert(SrcAddrSpace != DstAddrSpace &&
784
         "addrspacecast must be between different address spaces");
785

786
  if (DstAddrSpace == ADDRESS_SPACE_GENERIC) {
787
    // Specific to generic
788
    unsigned Opc;
789
    switch (SrcAddrSpace) {
790
    default: report_fatal_error("Bad address space in addrspacecast");
791
    case ADDRESS_SPACE_GLOBAL:
792
      Opc = TM.is64Bit() ? NVPTX::cvta_global_64 : NVPTX::cvta_global;
793
      break;
794
    case ADDRESS_SPACE_SHARED:
795
      Opc = TM.is64Bit() ? (TM.getPointerSizeInBits(SrcAddrSpace) == 32
796
                                ? NVPTX::cvta_shared_6432
797
                                : NVPTX::cvta_shared_64)
798
                         : NVPTX::cvta_shared;
799
      break;
800
    case ADDRESS_SPACE_CONST:
801
      Opc = TM.is64Bit() ? (TM.getPointerSizeInBits(SrcAddrSpace) == 32
802
                                ? NVPTX::cvta_const_6432
803
                                : NVPTX::cvta_const_64)
804
                         : NVPTX::cvta_const;
805
      break;
806
    case ADDRESS_SPACE_LOCAL:
807
      Opc = TM.is64Bit() ? (TM.getPointerSizeInBits(SrcAddrSpace) == 32
808
                                ? NVPTX::cvta_local_6432
809
                                : NVPTX::cvta_local_64)
810
                         : NVPTX::cvta_local;
811
      break;
812
    }
813
    ReplaceNode(N, CurDAG->getMachineNode(Opc, SDLoc(N), N->getValueType(0),
814
                                          Src));
815
    return;
816
  } else {
817
    // Generic to specific
818
    if (SrcAddrSpace != 0)
819
      report_fatal_error("Cannot cast between two non-generic address spaces");
820
    unsigned Opc;
821
    switch (DstAddrSpace) {
822
    default: report_fatal_error("Bad address space in addrspacecast");
823
    case ADDRESS_SPACE_GLOBAL:
824
      Opc = TM.is64Bit() ? NVPTX::cvta_to_global_64 : NVPTX::cvta_to_global;
825
      break;
826
    case ADDRESS_SPACE_SHARED:
827
      Opc = TM.is64Bit() ? (TM.getPointerSizeInBits(DstAddrSpace) == 32
828
                                ? NVPTX::cvta_to_shared_3264
829
                                : NVPTX::cvta_to_shared_64)
830
                         : NVPTX::cvta_to_shared;
831
      break;
832
    case ADDRESS_SPACE_CONST:
833
      Opc = TM.is64Bit() ? (TM.getPointerSizeInBits(DstAddrSpace) == 32
834
                                ? NVPTX::cvta_to_const_3264
835
                                : NVPTX::cvta_to_const_64)
836
                         : NVPTX::cvta_to_const;
837
      break;
838
    case ADDRESS_SPACE_LOCAL:
839
      Opc = TM.is64Bit() ? (TM.getPointerSizeInBits(DstAddrSpace) == 32
840
                                ? NVPTX::cvta_to_local_3264
841
                                : NVPTX::cvta_to_local_64)
842
                         : NVPTX::cvta_to_local;
843
      break;
844
    case ADDRESS_SPACE_PARAM:
845
      Opc = TM.is64Bit() ? NVPTX::nvvm_ptr_gen_to_param_64
846
                         : NVPTX::nvvm_ptr_gen_to_param;
847
      break;
848
    }
849
    ReplaceNode(N, CurDAG->getMachineNode(Opc, SDLoc(N), N->getValueType(0),
850
                                          Src));
851
    return;
852
  }
853
}
854

855
// Helper function template to reduce amount of boilerplate code for
856
// opcode selection.
857
static std::optional<unsigned>
858
pickOpcodeForVT(MVT::SimpleValueType VT, unsigned Opcode_i8,
859
                unsigned Opcode_i16, unsigned Opcode_i32,
860
                std::optional<unsigned> Opcode_i64, unsigned Opcode_f32,
861
                std::optional<unsigned> Opcode_f64) {
862
  switch (VT) {
863
  case MVT::i1:
864
  case MVT::i8:
865
    return Opcode_i8;
866
  case MVT::i16:
867
    return Opcode_i16;
868
  case MVT::i32:
869
    return Opcode_i32;
870
  case MVT::i64:
871
    return Opcode_i64;
872
  case MVT::f16:
873
  case MVT::bf16:
874
    return Opcode_i16;
875
  case MVT::v2f16:
876
  case MVT::v2bf16:
877
  case MVT::v2i16:
878
  case MVT::v4i8:
879
    return Opcode_i32;
880
  case MVT::f32:
881
    return Opcode_f32;
882
  case MVT::f64:
883
    return Opcode_f64;
884
  default:
885
    return std::nullopt;
886
  }
887
}
888

889
static int getLdStRegType(EVT VT) {
890
  if (VT.isFloatingPoint())
891
    switch (VT.getSimpleVT().SimpleTy) {
892
    case MVT::f16:
893
    case MVT::bf16:
894
    case MVT::v2f16:
895
    case MVT::v2bf16:
896
      return NVPTX::PTXLdStInstCode::Untyped;
897
    default:
898
      return NVPTX::PTXLdStInstCode::Float;
899
    }
900
  else
901
    return NVPTX::PTXLdStInstCode::Unsigned;
902
}
903

904
bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
905
  SDLoc dl(N);
906
  MemSDNode *LD = cast<MemSDNode>(N);
907
  assert(LD->readMem() && "Expected load");
908
  LoadSDNode *PlainLoad = dyn_cast<LoadSDNode>(N);
909
  EVT LoadedVT = LD->getMemoryVT();
910
  SDNode *NVPTXLD = nullptr;
911

912
  // do not support pre/post inc/dec
913
  if (PlainLoad && PlainLoad->isIndexed())
914
    return false;
915

916
  if (!LoadedVT.isSimple())
917
    return false;
918

919
  AtomicOrdering Ordering = LD->getSuccessOrdering();
920
  // In order to lower atomic loads with stronger guarantees we would need to
921
  // use load.acquire or insert fences. However these features were only added
922
  // with PTX ISA 6.0 / sm_70.
923
  // TODO: Check if we can actually use the new instructions and implement them.
924
  if (isStrongerThanMonotonic(Ordering))
925
    return false;
926

927
  // Address Space Setting
928
  unsigned int CodeAddrSpace = getCodeAddrSpace(LD);
929
  if (canLowerToLDG(LD, *Subtarget, CodeAddrSpace, MF)) {
930
    return tryLDGLDU(N);
931
  }
932

933
  unsigned int PointerSize =
934
      CurDAG->getDataLayout().getPointerSizeInBits(LD->getAddressSpace());
935

936
  // Volatile Setting
937
  // - .volatile is only available for .global and .shared
938
  // - .volatile has the same memory synchronization semantics as .relaxed.sys
939
  bool isVolatile = LD->isVolatile() || Ordering == AtomicOrdering::Monotonic;
940
  if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
941
      CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
942
      CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
943
    isVolatile = false;
944

945
  // Type Setting: fromType + fromTypeWidth
946
  //
947
  // Sign   : ISD::SEXTLOAD
948
  // Unsign : ISD::ZEXTLOAD, ISD::NON_EXTLOAD or ISD::EXTLOAD and the
949
  //          type is integer
950
  // Float  : ISD::NON_EXTLOAD or ISD::EXTLOAD and the type is float
951
  MVT SimpleVT = LoadedVT.getSimpleVT();
952
  MVT ScalarVT = SimpleVT.getScalarType();
953
  // Read at least 8 bits (predicates are stored as 8-bit values)
954
  unsigned fromTypeWidth = std::max(8U, (unsigned)ScalarVT.getSizeInBits());
955
  unsigned int fromType;
956

957
  // Vector Setting
958
  unsigned vecType = NVPTX::PTXLdStInstCode::Scalar;
959
  if (SimpleVT.isVector()) {
960
    assert((Isv2x16VT(LoadedVT) || LoadedVT == MVT::v4i8) &&
961
           "Unexpected vector type");
962
    // v2f16/v2bf16/v2i16 is loaded using ld.b32
963
    fromTypeWidth = 32;
964
  }
965

966
  if (PlainLoad && (PlainLoad->getExtensionType() == ISD::SEXTLOAD))
967
    fromType = NVPTX::PTXLdStInstCode::Signed;
968
  else
969
    fromType = getLdStRegType(ScalarVT);
970

971
  // Create the machine instruction DAG
972
  SDValue Chain = N->getOperand(0);
973
  SDValue N1 = N->getOperand(1);
974
  SDValue Addr;
975
  SDValue Offset, Base;
976
  std::optional<unsigned> Opcode;
977
  MVT::SimpleValueType TargetVT = LD->getSimpleValueType(0).SimpleTy;
978

979
  if (SelectDirectAddr(N1, Addr)) {
980
    Opcode = pickOpcodeForVT(TargetVT, NVPTX::LD_i8_avar, NVPTX::LD_i16_avar,
981
                             NVPTX::LD_i32_avar, NVPTX::LD_i64_avar,
982
                             NVPTX::LD_f32_avar, NVPTX::LD_f64_avar);
983
    if (!Opcode)
984
      return false;
985
    SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
986
                      getI32Imm(vecType, dl), getI32Imm(fromType, dl),
987
                      getI32Imm(fromTypeWidth, dl), Addr, Chain };
988
    NVPTXLD = CurDAG->getMachineNode(*Opcode, dl, TargetVT, MVT::Other, Ops);
989
  } else if (PointerSize == 64 ? SelectADDRsi64(N1.getNode(), N1, Base, Offset)
990
                               : SelectADDRsi(N1.getNode(), N1, Base, Offset)) {
991
    Opcode = pickOpcodeForVT(TargetVT, NVPTX::LD_i8_asi, NVPTX::LD_i16_asi,
992
                             NVPTX::LD_i32_asi, NVPTX::LD_i64_asi,
993
                             NVPTX::LD_f32_asi, NVPTX::LD_f64_asi);
994
    if (!Opcode)
995
      return false;
996
    SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
997
                      getI32Imm(vecType, dl), getI32Imm(fromType, dl),
998
                      getI32Imm(fromTypeWidth, dl), Base, Offset, Chain };
999
    NVPTXLD = CurDAG->getMachineNode(*Opcode, dl, TargetVT, MVT::Other, Ops);
1000
  } else if (PointerSize == 64 ? SelectADDRri64(N1.getNode(), N1, Base, Offset)
1001
                               : SelectADDRri(N1.getNode(), N1, Base, Offset)) {
1002
    if (PointerSize == 64)
1003
      Opcode =
1004
          pickOpcodeForVT(TargetVT, NVPTX::LD_i8_ari_64, NVPTX::LD_i16_ari_64,
1005
                          NVPTX::LD_i32_ari_64, NVPTX::LD_i64_ari_64,
1006
                          NVPTX::LD_f32_ari_64, NVPTX::LD_f64_ari_64);
1007
    else
1008
      Opcode = pickOpcodeForVT(TargetVT, NVPTX::LD_i8_ari, NVPTX::LD_i16_ari,
1009
                               NVPTX::LD_i32_ari, NVPTX::LD_i64_ari,
1010
                               NVPTX::LD_f32_ari, NVPTX::LD_f64_ari);
1011
    if (!Opcode)
1012
      return false;
1013
    SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
1014
                      getI32Imm(vecType, dl), getI32Imm(fromType, dl),
1015
                      getI32Imm(fromTypeWidth, dl), Base, Offset, Chain };
1016
    NVPTXLD = CurDAG->getMachineNode(*Opcode, dl, TargetVT, MVT::Other, Ops);
1017
  } else {
1018
    if (PointerSize == 64)
1019
      Opcode =
1020
          pickOpcodeForVT(TargetVT, NVPTX::LD_i8_areg_64, NVPTX::LD_i16_areg_64,
1021
                          NVPTX::LD_i32_areg_64, NVPTX::LD_i64_areg_64,
1022
                          NVPTX::LD_f32_areg_64, NVPTX::LD_f64_areg_64);
1023
    else
1024
      Opcode = pickOpcodeForVT(TargetVT, NVPTX::LD_i8_areg, NVPTX::LD_i16_areg,
1025
                               NVPTX::LD_i32_areg, NVPTX::LD_i64_areg,
1026
                               NVPTX::LD_f32_areg, NVPTX::LD_f64_areg);
1027
    if (!Opcode)
1028
      return false;
1029
    SDValue Ops[] = { getI32Imm(isVolatile, dl), getI32Imm(CodeAddrSpace, dl),
1030
                      getI32Imm(vecType, dl), getI32Imm(fromType, dl),
1031
                      getI32Imm(fromTypeWidth, dl), N1, Chain };
1032
    NVPTXLD = CurDAG->getMachineNode(*Opcode, dl, TargetVT, MVT::Other, Ops);
1033
  }
1034

1035
  if (!NVPTXLD)
1036
    return false;
1037

1038
  MachineMemOperand *MemRef = cast<MemSDNode>(N)->getMemOperand();
1039
  CurDAG->setNodeMemRefs(cast<MachineSDNode>(NVPTXLD), {MemRef});
1040

1041
  ReplaceNode(N, NVPTXLD);
1042
  return true;
1043
}
1044

1045
bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
1046

1047
  SDValue Chain = N->getOperand(0);
1048
  SDValue Op1 = N->getOperand(1);
1049
  SDValue Addr, Offset, Base;
1050
  std::optional<unsigned> Opcode;
1051
  SDLoc DL(N);
1052
  SDNode *LD;
1053
  MemSDNode *MemSD = cast<MemSDNode>(N);
1054
  EVT LoadedVT = MemSD->getMemoryVT();
1055

1056
  if (!LoadedVT.isSimple())
1057
    return false;
1058

1059
  // Address Space Setting
1060
  unsigned int CodeAddrSpace = getCodeAddrSpace(MemSD);
1061
  if (canLowerToLDG(MemSD, *Subtarget, CodeAddrSpace, MF)) {
1062
    return tryLDGLDU(N);
1063
  }
1064

1065
  unsigned int PointerSize =
1066
      CurDAG->getDataLayout().getPointerSizeInBits(MemSD->getAddressSpace());
1067

1068
  // Volatile Setting
1069
  // - .volatile is only availalble for .global and .shared
1070
  bool IsVolatile = MemSD->isVolatile();
1071
  if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
1072
      CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
1073
      CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
1074
    IsVolatile = false;
1075

1076
  // Vector Setting
1077
  MVT SimpleVT = LoadedVT.getSimpleVT();
1078

1079
  // Type Setting: fromType + fromTypeWidth
1080
  //
1081
  // Sign   : ISD::SEXTLOAD
1082
  // Unsign : ISD::ZEXTLOAD, ISD::NON_EXTLOAD or ISD::EXTLOAD and the
1083
  //          type is integer
1084
  // Float  : ISD::NON_EXTLOAD or ISD::EXTLOAD and the type is float
1085
  MVT ScalarVT = SimpleVT.getScalarType();
1086
  // Read at least 8 bits (predicates are stored as 8-bit values)
1087
  unsigned FromTypeWidth = std::max(8U, (unsigned)ScalarVT.getSizeInBits());
1088
  unsigned int FromType;
1089
  // The last operand holds the original LoadSDNode::getExtensionType() value
1090
  unsigned ExtensionType = cast<ConstantSDNode>(
1091
      N->getOperand(N->getNumOperands() - 1))->getZExtValue();
1092
  if (ExtensionType == ISD::SEXTLOAD)
1093
    FromType = NVPTX::PTXLdStInstCode::Signed;
1094
  else
1095
    FromType = getLdStRegType(ScalarVT);
1096

1097
  unsigned VecType;
1098

1099
  switch (N->getOpcode()) {
1100
  case NVPTXISD::LoadV2:
1101
    VecType = NVPTX::PTXLdStInstCode::V2;
1102
    break;
1103
  case NVPTXISD::LoadV4:
1104
    VecType = NVPTX::PTXLdStInstCode::V4;
1105
    break;
1106
  default:
1107
    return false;
1108
  }
1109

1110
  EVT EltVT = N->getValueType(0);
1111

1112
  // v8x16 is a special case. PTX doesn't have ld.v8.16
1113
  // instruction. Instead, we split the vector into v2x16 chunks and
1114
  // load them with ld.v4.b32.
1115
  if (Isv2x16VT(EltVT)) {
1116
    assert(N->getOpcode() == NVPTXISD::LoadV4 && "Unexpected load opcode.");
1117
    EltVT = MVT::i32;
1118
    FromType = NVPTX::PTXLdStInstCode::Untyped;
1119
    FromTypeWidth = 32;
1120
  }
1121

1122
  if (SelectDirectAddr(Op1, Addr)) {
1123
    switch (N->getOpcode()) {
1124
    default:
1125
      return false;
1126
    case NVPTXISD::LoadV2:
1127
      Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1128
                               NVPTX::LDV_i8_v2_avar, NVPTX::LDV_i16_v2_avar,
1129
                               NVPTX::LDV_i32_v2_avar, NVPTX::LDV_i64_v2_avar,
1130
                               NVPTX::LDV_f32_v2_avar, NVPTX::LDV_f64_v2_avar);
1131
      break;
1132
    case NVPTXISD::LoadV4:
1133
      Opcode =
1134
          pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v4_avar,
1135
                          NVPTX::LDV_i16_v4_avar, NVPTX::LDV_i32_v4_avar,
1136
                          std::nullopt, NVPTX::LDV_f32_v4_avar, std::nullopt);
1137
      break;
1138
    }
1139
    if (!Opcode)
1140
      return false;
1141
    SDValue Ops[] = { getI32Imm(IsVolatile, DL), getI32Imm(CodeAddrSpace, DL),
1142
                      getI32Imm(VecType, DL), getI32Imm(FromType, DL),
1143
                      getI32Imm(FromTypeWidth, DL), Addr, Chain };
1144
    LD = CurDAG->getMachineNode(*Opcode, DL, N->getVTList(), Ops);
1145
  } else if (PointerSize == 64
1146
                 ? SelectADDRsi64(Op1.getNode(), Op1, Base, Offset)
1147
                 : SelectADDRsi(Op1.getNode(), Op1, Base, Offset)) {
1148
    switch (N->getOpcode()) {
1149
    default:
1150
      return false;
1151
    case NVPTXISD::LoadV2:
1152
      Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1153
                               NVPTX::LDV_i8_v2_asi, NVPTX::LDV_i16_v2_asi,
1154
                               NVPTX::LDV_i32_v2_asi, NVPTX::LDV_i64_v2_asi,
1155
                               NVPTX::LDV_f32_v2_asi, NVPTX::LDV_f64_v2_asi);
1156
      break;
1157
    case NVPTXISD::LoadV4:
1158
      Opcode =
1159
          pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v4_asi,
1160
                          NVPTX::LDV_i16_v4_asi, NVPTX::LDV_i32_v4_asi,
1161
                          std::nullopt, NVPTX::LDV_f32_v4_asi, std::nullopt);
1162
      break;
1163
    }
1164
    if (!Opcode)
1165
      return false;
1166
    SDValue Ops[] = { getI32Imm(IsVolatile, DL), getI32Imm(CodeAddrSpace, DL),
1167
                      getI32Imm(VecType, DL), getI32Imm(FromType, DL),
1168
                      getI32Imm(FromTypeWidth, DL), Base, Offset, Chain };
1169
    LD = CurDAG->getMachineNode(*Opcode, DL, N->getVTList(), Ops);
1170
  } else if (PointerSize == 64
1171
                 ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset)
1172
                 : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) {
1173
    if (PointerSize == 64) {
1174
      switch (N->getOpcode()) {
1175
      default:
1176
        return false;
1177
      case NVPTXISD::LoadV2:
1178
        Opcode =
1179
            pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1180
                            NVPTX::LDV_i8_v2_ari_64, NVPTX::LDV_i16_v2_ari_64,
1181
                            NVPTX::LDV_i32_v2_ari_64, NVPTX::LDV_i64_v2_ari_64,
1182
                            NVPTX::LDV_f32_v2_ari_64, NVPTX::LDV_f64_v2_ari_64);
1183
        break;
1184
      case NVPTXISD::LoadV4:
1185
        Opcode = pickOpcodeForVT(
1186
            EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v4_ari_64,
1187
            NVPTX::LDV_i16_v4_ari_64, NVPTX::LDV_i32_v4_ari_64, std::nullopt,
1188
            NVPTX::LDV_f32_v4_ari_64, std::nullopt);
1189
        break;
1190
      }
1191
    } else {
1192
      switch (N->getOpcode()) {
1193
      default:
1194
        return false;
1195
      case NVPTXISD::LoadV2:
1196
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1197
                                 NVPTX::LDV_i8_v2_ari, NVPTX::LDV_i16_v2_ari,
1198
                                 NVPTX::LDV_i32_v2_ari, NVPTX::LDV_i64_v2_ari,
1199
                                 NVPTX::LDV_f32_v2_ari, NVPTX::LDV_f64_v2_ari);
1200
        break;
1201
      case NVPTXISD::LoadV4:
1202
        Opcode =
1203
            pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v4_ari,
1204
                            NVPTX::LDV_i16_v4_ari, NVPTX::LDV_i32_v4_ari,
1205
                            std::nullopt, NVPTX::LDV_f32_v4_ari, std::nullopt);
1206
        break;
1207
      }
1208
    }
1209
    if (!Opcode)
1210
      return false;
1211
    SDValue Ops[] = { getI32Imm(IsVolatile, DL), getI32Imm(CodeAddrSpace, DL),
1212
                      getI32Imm(VecType, DL), getI32Imm(FromType, DL),
1213
                      getI32Imm(FromTypeWidth, DL), Base, Offset, Chain };
1214

1215
    LD = CurDAG->getMachineNode(*Opcode, DL, N->getVTList(), Ops);
1216
  } else {
1217
    if (PointerSize == 64) {
1218
      switch (N->getOpcode()) {
1219
      default:
1220
        return false;
1221
      case NVPTXISD::LoadV2:
1222
        Opcode = pickOpcodeForVT(
1223
            EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v2_areg_64,
1224
            NVPTX::LDV_i16_v2_areg_64, NVPTX::LDV_i32_v2_areg_64,
1225
            NVPTX::LDV_i64_v2_areg_64, NVPTX::LDV_f32_v2_areg_64,
1226
            NVPTX::LDV_f64_v2_areg_64);
1227
        break;
1228
      case NVPTXISD::LoadV4:
1229
        Opcode = pickOpcodeForVT(
1230
            EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v4_areg_64,
1231
            NVPTX::LDV_i16_v4_areg_64, NVPTX::LDV_i32_v4_areg_64, std::nullopt,
1232
            NVPTX::LDV_f32_v4_areg_64, std::nullopt);
1233
        break;
1234
      }
1235
    } else {
1236
      switch (N->getOpcode()) {
1237
      default:
1238
        return false;
1239
      case NVPTXISD::LoadV2:
1240
        Opcode =
1241
            pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v2_areg,
1242
                            NVPTX::LDV_i16_v2_areg, NVPTX::LDV_i32_v2_areg,
1243
                            NVPTX::LDV_i64_v2_areg, NVPTX::LDV_f32_v2_areg,
1244
                            NVPTX::LDV_f64_v2_areg);
1245
        break;
1246
      case NVPTXISD::LoadV4:
1247
        Opcode =
1248
            pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::LDV_i8_v4_areg,
1249
                            NVPTX::LDV_i16_v4_areg, NVPTX::LDV_i32_v4_areg,
1250
                            std::nullopt, NVPTX::LDV_f32_v4_areg, std::nullopt);
1251
        break;
1252
      }
1253
    }
1254
    if (!Opcode)
1255
      return false;
1256
    SDValue Ops[] = { getI32Imm(IsVolatile, DL), getI32Imm(CodeAddrSpace, DL),
1257
                      getI32Imm(VecType, DL), getI32Imm(FromType, DL),
1258
                      getI32Imm(FromTypeWidth, DL), Op1, Chain };
1259
    LD = CurDAG->getMachineNode(*Opcode, DL, N->getVTList(), Ops);
1260
  }
1261

1262
  MachineMemOperand *MemRef = cast<MemSDNode>(N)->getMemOperand();
1263
  CurDAG->setNodeMemRefs(cast<MachineSDNode>(LD), {MemRef});
1264

1265
  ReplaceNode(N, LD);
1266
  return true;
1267
}
1268

1269
bool NVPTXDAGToDAGISel::tryLDGLDU(SDNode *N) {
1270

1271
  SDValue Chain = N->getOperand(0);
1272
  SDValue Op1;
1273
  MemSDNode *Mem;
1274
  bool IsLDG = true;
1275

1276
  // If this is an LDG intrinsic, the address is the third operand. If its an
1277
  // LDG/LDU SD node (from custom vector handling), then its the second operand
1278
  if (N->getOpcode() == ISD::INTRINSIC_W_CHAIN) {
1279
    Op1 = N->getOperand(2);
1280
    Mem = cast<MemIntrinsicSDNode>(N);
1281
    unsigned IID = N->getConstantOperandVal(1);
1282
    switch (IID) {
1283
    default:
1284
      return false;
1285
    case Intrinsic::nvvm_ldg_global_f:
1286
    case Intrinsic::nvvm_ldg_global_i:
1287
    case Intrinsic::nvvm_ldg_global_p:
1288
      IsLDG = true;
1289
      break;
1290
    case Intrinsic::nvvm_ldu_global_f:
1291
    case Intrinsic::nvvm_ldu_global_i:
1292
    case Intrinsic::nvvm_ldu_global_p:
1293
      IsLDG = false;
1294
      break;
1295
    }
1296
  } else {
1297
    Op1 = N->getOperand(1);
1298
    Mem = cast<MemSDNode>(N);
1299
  }
1300

1301
  std::optional<unsigned> Opcode;
1302
  SDLoc DL(N);
1303
  SDNode *LD;
1304
  SDValue Base, Offset, Addr;
1305
  EVT OrigType = N->getValueType(0);
1306

1307
  EVT EltVT = Mem->getMemoryVT();
1308
  unsigned NumElts = 1;
1309
  if (EltVT.isVector()) {
1310
    NumElts = EltVT.getVectorNumElements();
1311
    EltVT = EltVT.getVectorElementType();
1312
    // vectors of 16bits type are loaded/stored as multiples of v2x16 elements.
1313
    if ((EltVT == MVT::f16 && OrigType == MVT::v2f16) ||
1314
        (EltVT == MVT::bf16 && OrigType == MVT::v2bf16) ||
1315
        (EltVT == MVT::i16 && OrigType == MVT::v2i16)) {
1316
      assert(NumElts % 2 == 0 && "Vector must have even number of elements");
1317
      EltVT = OrigType;
1318
      NumElts /= 2;
1319
    } else if (OrigType == MVT::v4i8) {
1320
      EltVT = OrigType;
1321
      NumElts = 1;
1322
    }
1323
  }
1324

1325
  // Build the "promoted" result VTList for the load. If we are really loading
1326
  // i8s, then the return type will be promoted to i16 since we do not expose
1327
  // 8-bit registers in NVPTX.
1328
  EVT NodeVT = (EltVT == MVT::i8) ? MVT::i16 : EltVT;
1329
  SmallVector<EVT, 5> InstVTs;
1330
  for (unsigned i = 0; i != NumElts; ++i) {
1331
    InstVTs.push_back(NodeVT);
1332
  }
1333
  InstVTs.push_back(MVT::Other);
1334
  SDVTList InstVTList = CurDAG->getVTList(InstVTs);
1335

1336
  if (SelectDirectAddr(Op1, Addr)) {
1337
    switch (N->getOpcode()) {
1338
    default:
1339
      return false;
1340
    case ISD::LOAD:
1341
    case ISD::INTRINSIC_W_CHAIN:
1342
      if (IsLDG)
1343
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1344
                                 NVPTX::INT_PTX_LDG_GLOBAL_i8avar,
1345
                                 NVPTX::INT_PTX_LDG_GLOBAL_i16avar,
1346
                                 NVPTX::INT_PTX_LDG_GLOBAL_i32avar,
1347
                                 NVPTX::INT_PTX_LDG_GLOBAL_i64avar,
1348
                                 NVPTX::INT_PTX_LDG_GLOBAL_f32avar,
1349
                                 NVPTX::INT_PTX_LDG_GLOBAL_f64avar);
1350
      else
1351
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1352
                                 NVPTX::INT_PTX_LDU_GLOBAL_i8avar,
1353
                                 NVPTX::INT_PTX_LDU_GLOBAL_i16avar,
1354
                                 NVPTX::INT_PTX_LDU_GLOBAL_i32avar,
1355
                                 NVPTX::INT_PTX_LDU_GLOBAL_i64avar,
1356
                                 NVPTX::INT_PTX_LDU_GLOBAL_f32avar,
1357
                                 NVPTX::INT_PTX_LDU_GLOBAL_f64avar);
1358
      break;
1359
    case NVPTXISD::LoadV2:
1360
    case NVPTXISD::LDGV2:
1361
      Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1362
                               NVPTX::INT_PTX_LDG_G_v2i8_ELE_avar,
1363
                               NVPTX::INT_PTX_LDG_G_v2i16_ELE_avar,
1364
                               NVPTX::INT_PTX_LDG_G_v2i32_ELE_avar,
1365
                               NVPTX::INT_PTX_LDG_G_v2i64_ELE_avar,
1366
                               NVPTX::INT_PTX_LDG_G_v2f32_ELE_avar,
1367
                               NVPTX::INT_PTX_LDG_G_v2f64_ELE_avar);
1368
      break;
1369
    case NVPTXISD::LDUV2:
1370
      Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1371
                               NVPTX::INT_PTX_LDU_G_v2i8_ELE_avar,
1372
                               NVPTX::INT_PTX_LDU_G_v2i16_ELE_avar,
1373
                               NVPTX::INT_PTX_LDU_G_v2i32_ELE_avar,
1374
                               NVPTX::INT_PTX_LDU_G_v2i64_ELE_avar,
1375
                               NVPTX::INT_PTX_LDU_G_v2f32_ELE_avar,
1376
                               NVPTX::INT_PTX_LDU_G_v2f64_ELE_avar);
1377
      break;
1378
    case NVPTXISD::LoadV4:
1379
    case NVPTXISD::LDGV4:
1380
      Opcode = pickOpcodeForVT(
1381
          EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_G_v4i8_ELE_avar,
1382
          NVPTX::INT_PTX_LDG_G_v4i16_ELE_avar,
1383
          NVPTX::INT_PTX_LDG_G_v4i32_ELE_avar, std::nullopt,
1384
          NVPTX::INT_PTX_LDG_G_v4f32_ELE_avar, std::nullopt);
1385
      break;
1386
    case NVPTXISD::LDUV4:
1387
      Opcode = pickOpcodeForVT(
1388
          EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_G_v4i8_ELE_avar,
1389
          NVPTX::INT_PTX_LDU_G_v4i16_ELE_avar,
1390
          NVPTX::INT_PTX_LDU_G_v4i32_ELE_avar, std::nullopt,
1391
          NVPTX::INT_PTX_LDU_G_v4f32_ELE_avar, std::nullopt);
1392
      break;
1393
    }
1394
    if (!Opcode)
1395
      return false;
1396
    SDValue Ops[] = { Addr, Chain };
1397
    LD = CurDAG->getMachineNode(*Opcode, DL, InstVTList, Ops);
1398
  } else if (TM.is64Bit() ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset)
1399
                          : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) {
1400
    if (TM.is64Bit()) {
1401
      switch (N->getOpcode()) {
1402
      default:
1403
        return false;
1404
      case ISD::LOAD:
1405
      case ISD::INTRINSIC_W_CHAIN:
1406
        if (IsLDG)
1407
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1408
                                       NVPTX::INT_PTX_LDG_GLOBAL_i8ari64,
1409
                                       NVPTX::INT_PTX_LDG_GLOBAL_i16ari64,
1410
                                       NVPTX::INT_PTX_LDG_GLOBAL_i32ari64,
1411
                                       NVPTX::INT_PTX_LDG_GLOBAL_i64ari64,
1412
                                       NVPTX::INT_PTX_LDG_GLOBAL_f32ari64,
1413
                                       NVPTX::INT_PTX_LDG_GLOBAL_f64ari64);
1414
        else
1415
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1416
                                       NVPTX::INT_PTX_LDU_GLOBAL_i8ari64,
1417
                                       NVPTX::INT_PTX_LDU_GLOBAL_i16ari64,
1418
                                       NVPTX::INT_PTX_LDU_GLOBAL_i32ari64,
1419
                                       NVPTX::INT_PTX_LDU_GLOBAL_i64ari64,
1420
                                       NVPTX::INT_PTX_LDU_GLOBAL_f32ari64,
1421
                                       NVPTX::INT_PTX_LDU_GLOBAL_f64ari64);
1422
        break;
1423
      case NVPTXISD::LoadV2:
1424
      case NVPTXISD::LDGV2:
1425
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1426
                                     NVPTX::INT_PTX_LDG_G_v2i8_ELE_ari64,
1427
                                     NVPTX::INT_PTX_LDG_G_v2i16_ELE_ari64,
1428
                                     NVPTX::INT_PTX_LDG_G_v2i32_ELE_ari64,
1429
                                     NVPTX::INT_PTX_LDG_G_v2i64_ELE_ari64,
1430
                                     NVPTX::INT_PTX_LDG_G_v2f32_ELE_ari64,
1431
                                     NVPTX::INT_PTX_LDG_G_v2f64_ELE_ari64);
1432
        break;
1433
      case NVPTXISD::LDUV2:
1434
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1435
                                     NVPTX::INT_PTX_LDU_G_v2i8_ELE_ari64,
1436
                                     NVPTX::INT_PTX_LDU_G_v2i16_ELE_ari64,
1437
                                     NVPTX::INT_PTX_LDU_G_v2i32_ELE_ari64,
1438
                                     NVPTX::INT_PTX_LDU_G_v2i64_ELE_ari64,
1439
                                     NVPTX::INT_PTX_LDU_G_v2f32_ELE_ari64,
1440
                                     NVPTX::INT_PTX_LDU_G_v2f64_ELE_ari64);
1441
        break;
1442
      case NVPTXISD::LoadV4:
1443
      case NVPTXISD::LDGV4:
1444
        Opcode = pickOpcodeForVT(
1445
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_G_v4i8_ELE_ari64,
1446
            NVPTX::INT_PTX_LDG_G_v4i16_ELE_ari64,
1447
            NVPTX::INT_PTX_LDG_G_v4i32_ELE_ari64, std::nullopt,
1448
            NVPTX::INT_PTX_LDG_G_v4f32_ELE_ari64, std::nullopt);
1449
        break;
1450
      case NVPTXISD::LDUV4:
1451
        Opcode = pickOpcodeForVT(
1452
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_G_v4i8_ELE_ari64,
1453
            NVPTX::INT_PTX_LDU_G_v4i16_ELE_ari64,
1454
            NVPTX::INT_PTX_LDU_G_v4i32_ELE_ari64, std::nullopt,
1455
            NVPTX::INT_PTX_LDU_G_v4f32_ELE_ari64, std::nullopt);
1456
        break;
1457
      }
1458
    } else {
1459
      switch (N->getOpcode()) {
1460
      default:
1461
        return false;
1462
      case ISD::LOAD:
1463
      case ISD::INTRINSIC_W_CHAIN:
1464
        if (IsLDG)
1465
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1466
                                   NVPTX::INT_PTX_LDG_GLOBAL_i8ari,
1467
                                   NVPTX::INT_PTX_LDG_GLOBAL_i16ari,
1468
                                   NVPTX::INT_PTX_LDG_GLOBAL_i32ari,
1469
                                   NVPTX::INT_PTX_LDG_GLOBAL_i64ari,
1470
                                   NVPTX::INT_PTX_LDG_GLOBAL_f32ari,
1471
                                   NVPTX::INT_PTX_LDG_GLOBAL_f64ari);
1472
        else
1473
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1474
                                   NVPTX::INT_PTX_LDU_GLOBAL_i8ari,
1475
                                   NVPTX::INT_PTX_LDU_GLOBAL_i16ari,
1476
                                   NVPTX::INT_PTX_LDU_GLOBAL_i32ari,
1477
                                   NVPTX::INT_PTX_LDU_GLOBAL_i64ari,
1478
                                   NVPTX::INT_PTX_LDU_GLOBAL_f32ari,
1479
                                   NVPTX::INT_PTX_LDU_GLOBAL_f64ari);
1480
        break;
1481
      case NVPTXISD::LoadV2:
1482
      case NVPTXISD::LDGV2:
1483
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1484
                                 NVPTX::INT_PTX_LDG_G_v2i8_ELE_ari32,
1485
                                 NVPTX::INT_PTX_LDG_G_v2i16_ELE_ari32,
1486
                                 NVPTX::INT_PTX_LDG_G_v2i32_ELE_ari32,
1487
                                 NVPTX::INT_PTX_LDG_G_v2i64_ELE_ari32,
1488
                                 NVPTX::INT_PTX_LDG_G_v2f32_ELE_ari32,
1489
                                 NVPTX::INT_PTX_LDG_G_v2f64_ELE_ari32);
1490
        break;
1491
      case NVPTXISD::LDUV2:
1492
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1493
                                 NVPTX::INT_PTX_LDU_G_v2i8_ELE_ari32,
1494
                                 NVPTX::INT_PTX_LDU_G_v2i16_ELE_ari32,
1495
                                 NVPTX::INT_PTX_LDU_G_v2i32_ELE_ari32,
1496
                                 NVPTX::INT_PTX_LDU_G_v2i64_ELE_ari32,
1497
                                 NVPTX::INT_PTX_LDU_G_v2f32_ELE_ari32,
1498
                                 NVPTX::INT_PTX_LDU_G_v2f64_ELE_ari32);
1499
        break;
1500
      case NVPTXISD::LoadV4:
1501
      case NVPTXISD::LDGV4:
1502
        Opcode = pickOpcodeForVT(
1503
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_G_v4i8_ELE_ari32,
1504
            NVPTX::INT_PTX_LDG_G_v4i16_ELE_ari32,
1505
            NVPTX::INT_PTX_LDG_G_v4i32_ELE_ari32, std::nullopt,
1506
            NVPTX::INT_PTX_LDG_G_v4f32_ELE_ari32, std::nullopt);
1507
        break;
1508
      case NVPTXISD::LDUV4:
1509
        Opcode = pickOpcodeForVT(
1510
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_G_v4i8_ELE_ari32,
1511
            NVPTX::INT_PTX_LDU_G_v4i16_ELE_ari32,
1512
            NVPTX::INT_PTX_LDU_G_v4i32_ELE_ari32, std::nullopt,
1513
            NVPTX::INT_PTX_LDU_G_v4f32_ELE_ari32, std::nullopt);
1514
        break;
1515
      }
1516
    }
1517
    if (!Opcode)
1518
      return false;
1519
    SDValue Ops[] = {Base, Offset, Chain};
1520
    LD = CurDAG->getMachineNode(*Opcode, DL, InstVTList, Ops);
1521
  } else {
1522
    if (TM.is64Bit()) {
1523
      switch (N->getOpcode()) {
1524
      default:
1525
        return false;
1526
      case ISD::LOAD:
1527
      case ISD::INTRINSIC_W_CHAIN:
1528
        if (IsLDG)
1529
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1530
                                       NVPTX::INT_PTX_LDG_GLOBAL_i8areg64,
1531
                                       NVPTX::INT_PTX_LDG_GLOBAL_i16areg64,
1532
                                       NVPTX::INT_PTX_LDG_GLOBAL_i32areg64,
1533
                                       NVPTX::INT_PTX_LDG_GLOBAL_i64areg64,
1534
                                       NVPTX::INT_PTX_LDG_GLOBAL_f32areg64,
1535
                                       NVPTX::INT_PTX_LDG_GLOBAL_f64areg64);
1536
        else
1537
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1538
                                       NVPTX::INT_PTX_LDU_GLOBAL_i8areg64,
1539
                                       NVPTX::INT_PTX_LDU_GLOBAL_i16areg64,
1540
                                       NVPTX::INT_PTX_LDU_GLOBAL_i32areg64,
1541
                                       NVPTX::INT_PTX_LDU_GLOBAL_i64areg64,
1542
                                       NVPTX::INT_PTX_LDU_GLOBAL_f32areg64,
1543
                                       NVPTX::INT_PTX_LDU_GLOBAL_f64areg64);
1544
        break;
1545
      case NVPTXISD::LoadV2:
1546
      case NVPTXISD::LDGV2:
1547
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1548
                                     NVPTX::INT_PTX_LDG_G_v2i8_ELE_areg64,
1549
                                     NVPTX::INT_PTX_LDG_G_v2i16_ELE_areg64,
1550
                                     NVPTX::INT_PTX_LDG_G_v2i32_ELE_areg64,
1551
                                     NVPTX::INT_PTX_LDG_G_v2i64_ELE_areg64,
1552
                                     NVPTX::INT_PTX_LDG_G_v2f32_ELE_areg64,
1553
                                     NVPTX::INT_PTX_LDG_G_v2f64_ELE_areg64);
1554
        break;
1555
      case NVPTXISD::LDUV2:
1556
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1557
                                     NVPTX::INT_PTX_LDU_G_v2i8_ELE_areg64,
1558
                                     NVPTX::INT_PTX_LDU_G_v2i16_ELE_areg64,
1559
                                     NVPTX::INT_PTX_LDU_G_v2i32_ELE_areg64,
1560
                                     NVPTX::INT_PTX_LDU_G_v2i64_ELE_areg64,
1561
                                     NVPTX::INT_PTX_LDU_G_v2f32_ELE_areg64,
1562
                                     NVPTX::INT_PTX_LDU_G_v2f64_ELE_areg64);
1563
        break;
1564
      case NVPTXISD::LoadV4:
1565
      case NVPTXISD::LDGV4:
1566
        Opcode = pickOpcodeForVT(
1567
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_G_v4i8_ELE_areg64,
1568
            NVPTX::INT_PTX_LDG_G_v4i16_ELE_areg64,
1569
            NVPTX::INT_PTX_LDG_G_v4i32_ELE_areg64, std::nullopt,
1570
            NVPTX::INT_PTX_LDG_G_v4f32_ELE_areg64, std::nullopt);
1571
        break;
1572
      case NVPTXISD::LDUV4:
1573
        Opcode = pickOpcodeForVT(
1574
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_G_v4i8_ELE_areg64,
1575
            NVPTX::INT_PTX_LDU_G_v4i16_ELE_areg64,
1576
            NVPTX::INT_PTX_LDU_G_v4i32_ELE_areg64, std::nullopt,
1577
            NVPTX::INT_PTX_LDU_G_v4f32_ELE_areg64, std::nullopt);
1578
        break;
1579
      }
1580
    } else {
1581
      switch (N->getOpcode()) {
1582
      default:
1583
        return false;
1584
      case ISD::LOAD:
1585
      case ISD::INTRINSIC_W_CHAIN:
1586
        if (IsLDG)
1587
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1588
                                   NVPTX::INT_PTX_LDG_GLOBAL_i8areg,
1589
                                   NVPTX::INT_PTX_LDG_GLOBAL_i16areg,
1590
                                   NVPTX::INT_PTX_LDG_GLOBAL_i32areg,
1591
                                   NVPTX::INT_PTX_LDG_GLOBAL_i64areg,
1592
                                   NVPTX::INT_PTX_LDG_GLOBAL_f32areg,
1593
                                   NVPTX::INT_PTX_LDG_GLOBAL_f64areg);
1594
        else
1595
          Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1596
                                   NVPTX::INT_PTX_LDU_GLOBAL_i8areg,
1597
                                   NVPTX::INT_PTX_LDU_GLOBAL_i16areg,
1598
                                   NVPTX::INT_PTX_LDU_GLOBAL_i32areg,
1599
                                   NVPTX::INT_PTX_LDU_GLOBAL_i64areg,
1600
                                   NVPTX::INT_PTX_LDU_GLOBAL_f32areg,
1601
                                   NVPTX::INT_PTX_LDU_GLOBAL_f64areg);
1602
        break;
1603
      case NVPTXISD::LoadV2:
1604
      case NVPTXISD::LDGV2:
1605
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1606
                                 NVPTX::INT_PTX_LDG_G_v2i8_ELE_areg32,
1607
                                 NVPTX::INT_PTX_LDG_G_v2i16_ELE_areg32,
1608
                                 NVPTX::INT_PTX_LDG_G_v2i32_ELE_areg32,
1609
                                 NVPTX::INT_PTX_LDG_G_v2i64_ELE_areg32,
1610
                                 NVPTX::INT_PTX_LDG_G_v2f32_ELE_areg32,
1611
                                 NVPTX::INT_PTX_LDG_G_v2f64_ELE_areg32);
1612
        break;
1613
      case NVPTXISD::LDUV2:
1614
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1615
                                 NVPTX::INT_PTX_LDU_G_v2i8_ELE_areg32,
1616
                                 NVPTX::INT_PTX_LDU_G_v2i16_ELE_areg32,
1617
                                 NVPTX::INT_PTX_LDU_G_v2i32_ELE_areg32,
1618
                                 NVPTX::INT_PTX_LDU_G_v2i64_ELE_areg32,
1619
                                 NVPTX::INT_PTX_LDU_G_v2f32_ELE_areg32,
1620
                                 NVPTX::INT_PTX_LDU_G_v2f64_ELE_areg32);
1621
        break;
1622
      case NVPTXISD::LoadV4:
1623
      case NVPTXISD::LDGV4:
1624
        Opcode = pickOpcodeForVT(
1625
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDG_G_v4i8_ELE_areg32,
1626
            NVPTX::INT_PTX_LDG_G_v4i16_ELE_areg32,
1627
            NVPTX::INT_PTX_LDG_G_v4i32_ELE_areg32, std::nullopt,
1628
            NVPTX::INT_PTX_LDG_G_v4f32_ELE_areg32, std::nullopt);
1629
        break;
1630
      case NVPTXISD::LDUV4:
1631
        Opcode = pickOpcodeForVT(
1632
            EltVT.getSimpleVT().SimpleTy, NVPTX::INT_PTX_LDU_G_v4i8_ELE_areg32,
1633
            NVPTX::INT_PTX_LDU_G_v4i16_ELE_areg32,
1634
            NVPTX::INT_PTX_LDU_G_v4i32_ELE_areg32, std::nullopt,
1635
            NVPTX::INT_PTX_LDU_G_v4f32_ELE_areg32, std::nullopt);
1636
        break;
1637
      }
1638
    }
1639
    if (!Opcode)
1640
      return false;
1641
    SDValue Ops[] = { Op1, Chain };
1642
    LD = CurDAG->getMachineNode(*Opcode, DL, InstVTList, Ops);
1643
  }
1644

1645
  // For automatic generation of LDG (through SelectLoad[Vector], not the
1646
  // intrinsics), we may have an extending load like:
1647
  //
1648
  //   i32,ch = load<LD1[%data1(addrspace=1)], zext from i8> t0, t7, undef:i64
1649
  //
1650
  // In this case, the matching logic above will select a load for the original
1651
  // memory type (in this case, i8) and our types will not match (the node needs
1652
  // to return an i32 in this case). Our LDG/LDU nodes do not support the
1653
  // concept of sign-/zero-extension, so emulate it here by adding an explicit
1654
  // CVT instruction. Ptxas should clean up any redundancies here.
1655

1656
  LoadSDNode *LdNode = dyn_cast<LoadSDNode>(N);
1657

1658
  if (OrigType != EltVT &&
1659
      (LdNode || (OrigType.isFloatingPoint() && EltVT.isFloatingPoint()))) {
1660
    // We have an extending-load. The instruction we selected operates on the
1661
    // smaller type, but the SDNode we are replacing has the larger type. We
1662
    // need to emit a CVT to make the types match.
1663
    unsigned CvtOpc =
1664
        GetConvertOpcode(OrigType.getSimpleVT(), EltVT.getSimpleVT(), LdNode);
1665

1666
    // For each output value, apply the manual sign/zero-extension and make sure
1667
    // all users of the load go through that CVT.
1668
    for (unsigned i = 0; i != NumElts; ++i) {
1669
      SDValue Res(LD, i);
1670
      SDValue OrigVal(N, i);
1671

1672
      SDNode *CvtNode =
1673
        CurDAG->getMachineNode(CvtOpc, DL, OrigType, Res,
1674
                               CurDAG->getTargetConstant(NVPTX::PTXCvtMode::NONE,
1675
                                                         DL, MVT::i32));
1676
      ReplaceUses(OrigVal, SDValue(CvtNode, 0));
1677
    }
1678
  }
1679

1680
  ReplaceNode(N, LD);
1681
  return true;
1682
}
1683

1684
bool NVPTXDAGToDAGISel::tryStore(SDNode *N) {
1685
  SDLoc dl(N);
1686
  MemSDNode *ST = cast<MemSDNode>(N);
1687
  assert(ST->writeMem() && "Expected store");
1688
  StoreSDNode *PlainStore = dyn_cast<StoreSDNode>(N);
1689
  AtomicSDNode *AtomicStore = dyn_cast<AtomicSDNode>(N);
1690
  assert((PlainStore || AtomicStore) && "Expected store");
1691
  EVT StoreVT = ST->getMemoryVT();
1692
  SDNode *NVPTXST = nullptr;
1693

1694
  // do not support pre/post inc/dec
1695
  if (PlainStore && PlainStore->isIndexed())
1696
    return false;
1697

1698
  if (!StoreVT.isSimple())
1699
    return false;
1700

1701
  AtomicOrdering Ordering = ST->getSuccessOrdering();
1702
  // In order to lower atomic loads with stronger guarantees we would need to
1703
  // use store.release or insert fences. However these features were only added
1704
  // with PTX ISA 6.0 / sm_70.
1705
  // TODO: Check if we can actually use the new instructions and implement them.
1706
  if (isStrongerThanMonotonic(Ordering))
1707
    return false;
1708

1709
  // Address Space Setting
1710
  unsigned int CodeAddrSpace = getCodeAddrSpace(ST);
1711
  unsigned int PointerSize =
1712
      CurDAG->getDataLayout().getPointerSizeInBits(ST->getAddressSpace());
1713

1714
  // Volatile Setting
1715
  // - .volatile is only available for .global and .shared
1716
  // - .volatile has the same memory synchronization semantics as .relaxed.sys
1717
  bool isVolatile = ST->isVolatile() || Ordering == AtomicOrdering::Monotonic;
1718
  if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
1719
      CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
1720
      CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
1721
    isVolatile = false;
1722

1723
  // Vector Setting
1724
  MVT SimpleVT = StoreVT.getSimpleVT();
1725
  unsigned vecType = NVPTX::PTXLdStInstCode::Scalar;
1726

1727
  // Type Setting: toType + toTypeWidth
1728
  // - for integer type, always use 'u'
1729
  //
1730
  MVT ScalarVT = SimpleVT.getScalarType();
1731
  unsigned toTypeWidth = ScalarVT.getSizeInBits();
1732
  if (SimpleVT.isVector()) {
1733
    assert((Isv2x16VT(StoreVT) || StoreVT == MVT::v4i8) &&
1734
           "Unexpected vector type");
1735
    // v2x16 is stored using st.b32
1736
    toTypeWidth = 32;
1737
  }
1738

1739
  unsigned int toType = getLdStRegType(ScalarVT);
1740

1741
  // Create the machine instruction DAG
1742
  SDValue Chain = ST->getChain();
1743
  SDValue Value = PlainStore ? PlainStore->getValue() : AtomicStore->getVal();
1744
  SDValue BasePtr = ST->getBasePtr();
1745
  SDValue Addr;
1746
  SDValue Offset, Base;
1747
  std::optional<unsigned> Opcode;
1748
  MVT::SimpleValueType SourceVT =
1749
      Value.getNode()->getSimpleValueType(0).SimpleTy;
1750

1751
  if (SelectDirectAddr(BasePtr, Addr)) {
1752
    Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_avar, NVPTX::ST_i16_avar,
1753
                             NVPTX::ST_i32_avar, NVPTX::ST_i64_avar,
1754
                             NVPTX::ST_f32_avar, NVPTX::ST_f64_avar);
1755
    if (!Opcode)
1756
      return false;
1757
    SDValue Ops[] = {Value,
1758
                     getI32Imm(isVolatile, dl),
1759
                     getI32Imm(CodeAddrSpace, dl),
1760
                     getI32Imm(vecType, dl),
1761
                     getI32Imm(toType, dl),
1762
                     getI32Imm(toTypeWidth, dl),
1763
                     Addr,
1764
                     Chain};
1765
    NVPTXST = CurDAG->getMachineNode(*Opcode, dl, MVT::Other, Ops);
1766
  } else if (PointerSize == 64
1767
                 ? SelectADDRsi64(BasePtr.getNode(), BasePtr, Base, Offset)
1768
                 : SelectADDRsi(BasePtr.getNode(), BasePtr, Base, Offset)) {
1769
    Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_asi, NVPTX::ST_i16_asi,
1770
                             NVPTX::ST_i32_asi, NVPTX::ST_i64_asi,
1771
                             NVPTX::ST_f32_asi, NVPTX::ST_f64_asi);
1772
    if (!Opcode)
1773
      return false;
1774
    SDValue Ops[] = {Value,
1775
                     getI32Imm(isVolatile, dl),
1776
                     getI32Imm(CodeAddrSpace, dl),
1777
                     getI32Imm(vecType, dl),
1778
                     getI32Imm(toType, dl),
1779
                     getI32Imm(toTypeWidth, dl),
1780
                     Base,
1781
                     Offset,
1782
                     Chain};
1783
    NVPTXST = CurDAG->getMachineNode(*Opcode, dl, MVT::Other, Ops);
1784
  } else if (PointerSize == 64
1785
                 ? SelectADDRri64(BasePtr.getNode(), BasePtr, Base, Offset)
1786
                 : SelectADDRri(BasePtr.getNode(), BasePtr, Base, Offset)) {
1787
    if (PointerSize == 64)
1788
      Opcode =
1789
          pickOpcodeForVT(SourceVT, NVPTX::ST_i8_ari_64, NVPTX::ST_i16_ari_64,
1790
                          NVPTX::ST_i32_ari_64, NVPTX::ST_i64_ari_64,
1791
                          NVPTX::ST_f32_ari_64, NVPTX::ST_f64_ari_64);
1792
    else
1793
      Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_ari, NVPTX::ST_i16_ari,
1794
                               NVPTX::ST_i32_ari, NVPTX::ST_i64_ari,
1795
                               NVPTX::ST_f32_ari, NVPTX::ST_f64_ari);
1796
    if (!Opcode)
1797
      return false;
1798

1799
    SDValue Ops[] = {Value,
1800
                     getI32Imm(isVolatile, dl),
1801
                     getI32Imm(CodeAddrSpace, dl),
1802
                     getI32Imm(vecType, dl),
1803
                     getI32Imm(toType, dl),
1804
                     getI32Imm(toTypeWidth, dl),
1805
                     Base,
1806
                     Offset,
1807
                     Chain};
1808
    NVPTXST = CurDAG->getMachineNode(*Opcode, dl, MVT::Other, Ops);
1809
  } else {
1810
    if (PointerSize == 64)
1811
      Opcode =
1812
          pickOpcodeForVT(SourceVT, NVPTX::ST_i8_areg_64, NVPTX::ST_i16_areg_64,
1813
                          NVPTX::ST_i32_areg_64, NVPTX::ST_i64_areg_64,
1814
                          NVPTX::ST_f32_areg_64, NVPTX::ST_f64_areg_64);
1815
    else
1816
      Opcode = pickOpcodeForVT(SourceVT, NVPTX::ST_i8_areg, NVPTX::ST_i16_areg,
1817
                               NVPTX::ST_i32_areg, NVPTX::ST_i64_areg,
1818
                               NVPTX::ST_f32_areg, NVPTX::ST_f64_areg);
1819
    if (!Opcode)
1820
      return false;
1821
    SDValue Ops[] = {Value,
1822
                     getI32Imm(isVolatile, dl),
1823
                     getI32Imm(CodeAddrSpace, dl),
1824
                     getI32Imm(vecType, dl),
1825
                     getI32Imm(toType, dl),
1826
                     getI32Imm(toTypeWidth, dl),
1827
                     BasePtr,
1828
                     Chain};
1829
    NVPTXST = CurDAG->getMachineNode(*Opcode, dl, MVT::Other, Ops);
1830
  }
1831

1832
  if (!NVPTXST)
1833
    return false;
1834

1835
  MachineMemOperand *MemRef = cast<MemSDNode>(N)->getMemOperand();
1836
  CurDAG->setNodeMemRefs(cast<MachineSDNode>(NVPTXST), {MemRef});
1837
  ReplaceNode(N, NVPTXST);
1838
  return true;
1839
}
1840

1841
bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) {
1842
  SDValue Chain = N->getOperand(0);
1843
  SDValue Op1 = N->getOperand(1);
1844
  SDValue Addr, Offset, Base;
1845
  std::optional<unsigned> Opcode;
1846
  SDLoc DL(N);
1847
  SDNode *ST;
1848
  EVT EltVT = Op1.getValueType();
1849
  MemSDNode *MemSD = cast<MemSDNode>(N);
1850
  EVT StoreVT = MemSD->getMemoryVT();
1851

1852
  // Address Space Setting
1853
  unsigned CodeAddrSpace = getCodeAddrSpace(MemSD);
1854
  if (CodeAddrSpace == NVPTX::PTXLdStInstCode::CONSTANT) {
1855
    report_fatal_error("Cannot store to pointer that points to constant "
1856
                       "memory space");
1857
  }
1858
  unsigned int PointerSize =
1859
      CurDAG->getDataLayout().getPointerSizeInBits(MemSD->getAddressSpace());
1860

1861
  // Volatile Setting
1862
  // - .volatile is only availalble for .global and .shared
1863
  bool IsVolatile = MemSD->isVolatile();
1864
  if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
1865
      CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
1866
      CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
1867
    IsVolatile = false;
1868

1869
  // Type Setting: toType + toTypeWidth
1870
  // - for integer type, always use 'u'
1871
  assert(StoreVT.isSimple() && "Store value is not simple");
1872
  MVT ScalarVT = StoreVT.getSimpleVT().getScalarType();
1873
  unsigned ToTypeWidth = ScalarVT.getSizeInBits();
1874
  unsigned ToType = getLdStRegType(ScalarVT);
1875

1876
  SmallVector<SDValue, 12> StOps;
1877
  SDValue N2;
1878
  unsigned VecType;
1879

1880
  switch (N->getOpcode()) {
1881
  case NVPTXISD::StoreV2:
1882
    VecType = NVPTX::PTXLdStInstCode::V2;
1883
    StOps.push_back(N->getOperand(1));
1884
    StOps.push_back(N->getOperand(2));
1885
    N2 = N->getOperand(3);
1886
    break;
1887
  case NVPTXISD::StoreV4:
1888
    VecType = NVPTX::PTXLdStInstCode::V4;
1889
    StOps.push_back(N->getOperand(1));
1890
    StOps.push_back(N->getOperand(2));
1891
    StOps.push_back(N->getOperand(3));
1892
    StOps.push_back(N->getOperand(4));
1893
    N2 = N->getOperand(5);
1894
    break;
1895
  default:
1896
    return false;
1897
  }
1898

1899
  // v8x16 is a special case. PTX doesn't have st.v8.x16
1900
  // instruction. Instead, we split the vector into v2x16 chunks and
1901
  // store them with st.v4.b32.
1902
  if (Isv2x16VT(EltVT)) {
1903
    assert(N->getOpcode() == NVPTXISD::StoreV4 && "Unexpected load opcode.");
1904
    EltVT = MVT::i32;
1905
    ToType = NVPTX::PTXLdStInstCode::Untyped;
1906
    ToTypeWidth = 32;
1907
  }
1908

1909
  StOps.push_back(getI32Imm(IsVolatile, DL));
1910
  StOps.push_back(getI32Imm(CodeAddrSpace, DL));
1911
  StOps.push_back(getI32Imm(VecType, DL));
1912
  StOps.push_back(getI32Imm(ToType, DL));
1913
  StOps.push_back(getI32Imm(ToTypeWidth, DL));
1914

1915
  if (SelectDirectAddr(N2, Addr)) {
1916
    switch (N->getOpcode()) {
1917
    default:
1918
      return false;
1919
    case NVPTXISD::StoreV2:
1920
      Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1921
                               NVPTX::STV_i8_v2_avar, NVPTX::STV_i16_v2_avar,
1922
                               NVPTX::STV_i32_v2_avar, NVPTX::STV_i64_v2_avar,
1923
                               NVPTX::STV_f32_v2_avar, NVPTX::STV_f64_v2_avar);
1924
      break;
1925
    case NVPTXISD::StoreV4:
1926
      Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1927
                               NVPTX::STV_i8_v4_avar, NVPTX::STV_i16_v4_avar,
1928
                               NVPTX::STV_i32_v4_avar, std::nullopt,
1929
                               NVPTX::STV_f32_v4_avar, std::nullopt);
1930
      break;
1931
    }
1932
    StOps.push_back(Addr);
1933
  } else if (PointerSize == 64 ? SelectADDRsi64(N2.getNode(), N2, Base, Offset)
1934
                               : SelectADDRsi(N2.getNode(), N2, Base, Offset)) {
1935
    switch (N->getOpcode()) {
1936
    default:
1937
      return false;
1938
    case NVPTXISD::StoreV2:
1939
      Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1940
                               NVPTX::STV_i8_v2_asi, NVPTX::STV_i16_v2_asi,
1941
                               NVPTX::STV_i32_v2_asi, NVPTX::STV_i64_v2_asi,
1942
                               NVPTX::STV_f32_v2_asi, NVPTX::STV_f64_v2_asi);
1943
      break;
1944
    case NVPTXISD::StoreV4:
1945
      Opcode =
1946
          pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v4_asi,
1947
                          NVPTX::STV_i16_v4_asi, NVPTX::STV_i32_v4_asi,
1948
                          std::nullopt, NVPTX::STV_f32_v4_asi, std::nullopt);
1949
      break;
1950
    }
1951
    StOps.push_back(Base);
1952
    StOps.push_back(Offset);
1953
  } else if (PointerSize == 64 ? SelectADDRri64(N2.getNode(), N2, Base, Offset)
1954
                               : SelectADDRri(N2.getNode(), N2, Base, Offset)) {
1955
    if (PointerSize == 64) {
1956
      switch (N->getOpcode()) {
1957
      default:
1958
        return false;
1959
      case NVPTXISD::StoreV2:
1960
        Opcode =
1961
            pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1962
                            NVPTX::STV_i8_v2_ari_64, NVPTX::STV_i16_v2_ari_64,
1963
                            NVPTX::STV_i32_v2_ari_64, NVPTX::STV_i64_v2_ari_64,
1964
                            NVPTX::STV_f32_v2_ari_64, NVPTX::STV_f64_v2_ari_64);
1965
        break;
1966
      case NVPTXISD::StoreV4:
1967
        Opcode = pickOpcodeForVT(
1968
            EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v4_ari_64,
1969
            NVPTX::STV_i16_v4_ari_64, NVPTX::STV_i32_v4_ari_64, std::nullopt,
1970
            NVPTX::STV_f32_v4_ari_64, std::nullopt);
1971
        break;
1972
      }
1973
    } else {
1974
      switch (N->getOpcode()) {
1975
      default:
1976
        return false;
1977
      case NVPTXISD::StoreV2:
1978
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1979
                                 NVPTX::STV_i8_v2_ari, NVPTX::STV_i16_v2_ari,
1980
                                 NVPTX::STV_i32_v2_ari, NVPTX::STV_i64_v2_ari,
1981
                                 NVPTX::STV_f32_v2_ari, NVPTX::STV_f64_v2_ari);
1982
        break;
1983
      case NVPTXISD::StoreV4:
1984
        Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
1985
                                 NVPTX::STV_i8_v4_ari, NVPTX::STV_i16_v4_ari,
1986
                                 NVPTX::STV_i32_v4_ari, std::nullopt,
1987
                                 NVPTX::STV_f32_v4_ari, std::nullopt);
1988
        break;
1989
      }
1990
    }
1991
    StOps.push_back(Base);
1992
    StOps.push_back(Offset);
1993
  } else {
1994
    if (PointerSize == 64) {
1995
      switch (N->getOpcode()) {
1996
      default:
1997
        return false;
1998
      case NVPTXISD::StoreV2:
1999
        Opcode = pickOpcodeForVT(
2000
            EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v2_areg_64,
2001
            NVPTX::STV_i16_v2_areg_64, NVPTX::STV_i32_v2_areg_64,
2002
            NVPTX::STV_i64_v2_areg_64, NVPTX::STV_f32_v2_areg_64,
2003
            NVPTX::STV_f64_v2_areg_64);
2004
        break;
2005
      case NVPTXISD::StoreV4:
2006
        Opcode = pickOpcodeForVT(
2007
            EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v4_areg_64,
2008
            NVPTX::STV_i16_v4_areg_64, NVPTX::STV_i32_v4_areg_64, std::nullopt,
2009
            NVPTX::STV_f32_v4_areg_64, std::nullopt);
2010
        break;
2011
      }
2012
    } else {
2013
      switch (N->getOpcode()) {
2014
      default:
2015
        return false;
2016
      case NVPTXISD::StoreV2:
2017
        Opcode =
2018
            pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v2_areg,
2019
                            NVPTX::STV_i16_v2_areg, NVPTX::STV_i32_v2_areg,
2020
                            NVPTX::STV_i64_v2_areg, NVPTX::STV_f32_v2_areg,
2021
                            NVPTX::STV_f64_v2_areg);
2022
        break;
2023
      case NVPTXISD::StoreV4:
2024
        Opcode =
2025
            pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy, NVPTX::STV_i8_v4_areg,
2026
                            NVPTX::STV_i16_v4_areg, NVPTX::STV_i32_v4_areg,
2027
                            std::nullopt, NVPTX::STV_f32_v4_areg, std::nullopt);
2028
        break;
2029
      }
2030
    }
2031
    StOps.push_back(N2);
2032
  }
2033

2034
  if (!Opcode)
2035
    return false;
2036

2037
  StOps.push_back(Chain);
2038

2039
  ST = CurDAG->getMachineNode(*Opcode, DL, MVT::Other, StOps);
2040

2041
  MachineMemOperand *MemRef = cast<MemSDNode>(N)->getMemOperand();
2042
  CurDAG->setNodeMemRefs(cast<MachineSDNode>(ST), {MemRef});
2043

2044
  ReplaceNode(N, ST);
2045
  return true;
2046
}
2047

2048
bool NVPTXDAGToDAGISel::tryLoadParam(SDNode *Node) {
2049
  SDValue Chain = Node->getOperand(0);
2050
  SDValue Offset = Node->getOperand(2);
2051
  SDValue Glue = Node->getOperand(3);
2052
  SDLoc DL(Node);
2053
  MemSDNode *Mem = cast<MemSDNode>(Node);
2054

2055
  unsigned VecSize;
2056
  switch (Node->getOpcode()) {
2057
  default:
2058
    return false;
2059
  case NVPTXISD::LoadParam:
2060
    VecSize = 1;
2061
    break;
2062
  case NVPTXISD::LoadParamV2:
2063
    VecSize = 2;
2064
    break;
2065
  case NVPTXISD::LoadParamV4:
2066
    VecSize = 4;
2067
    break;
2068
  }
2069

2070
  EVT EltVT = Node->getValueType(0);
2071
  EVT MemVT = Mem->getMemoryVT();
2072

2073
  std::optional<unsigned> Opcode;
2074

2075
  switch (VecSize) {
2076
  default:
2077
    return false;
2078
  case 1:
2079
    Opcode = pickOpcodeForVT(MemVT.getSimpleVT().SimpleTy,
2080
                             NVPTX::LoadParamMemI8, NVPTX::LoadParamMemI16,
2081
                             NVPTX::LoadParamMemI32, NVPTX::LoadParamMemI64,
2082
                             NVPTX::LoadParamMemF32, NVPTX::LoadParamMemF64);
2083
    break;
2084
  case 2:
2085
    Opcode =
2086
        pickOpcodeForVT(MemVT.getSimpleVT().SimpleTy, NVPTX::LoadParamMemV2I8,
2087
                        NVPTX::LoadParamMemV2I16, NVPTX::LoadParamMemV2I32,
2088
                        NVPTX::LoadParamMemV2I64, NVPTX::LoadParamMemV2F32,
2089
                        NVPTX::LoadParamMemV2F64);
2090
    break;
2091
  case 4:
2092
    Opcode =
2093
        pickOpcodeForVT(MemVT.getSimpleVT().SimpleTy, NVPTX::LoadParamMemV4I8,
2094
                        NVPTX::LoadParamMemV4I16, NVPTX::LoadParamMemV4I32,
2095
                        std::nullopt, NVPTX::LoadParamMemV4F32, std::nullopt);
2096
    break;
2097
  }
2098
  if (!Opcode)
2099
    return false;
2100

2101
  SDVTList VTs;
2102
  if (VecSize == 1) {
2103
    VTs = CurDAG->getVTList(EltVT, MVT::Other, MVT::Glue);
2104
  } else if (VecSize == 2) {
2105
    VTs = CurDAG->getVTList(EltVT, EltVT, MVT::Other, MVT::Glue);
2106
  } else {
2107
    EVT EVTs[] = { EltVT, EltVT, EltVT, EltVT, MVT::Other, MVT::Glue };
2108
    VTs = CurDAG->getVTList(EVTs);
2109
  }
2110

2111
  unsigned OffsetVal = Offset->getAsZExtVal();
2112

2113
  SmallVector<SDValue, 2> Ops;
2114
  Ops.push_back(CurDAG->getTargetConstant(OffsetVal, DL, MVT::i32));
2115
  Ops.push_back(Chain);
2116
  Ops.push_back(Glue);
2117

2118
  ReplaceNode(Node, CurDAG->getMachineNode(*Opcode, DL, VTs, Ops));
2119
  return true;
2120
}
2121

2122
bool NVPTXDAGToDAGISel::tryStoreRetval(SDNode *N) {
2123
  SDLoc DL(N);
2124
  SDValue Chain = N->getOperand(0);
2125
  SDValue Offset = N->getOperand(1);
2126
  unsigned OffsetVal = Offset->getAsZExtVal();
2127
  MemSDNode *Mem = cast<MemSDNode>(N);
2128

2129
  // How many elements do we have?
2130
  unsigned NumElts = 1;
2131
  switch (N->getOpcode()) {
2132
  default:
2133
    return false;
2134
  case NVPTXISD::StoreRetval:
2135
    NumElts = 1;
2136
    break;
2137
  case NVPTXISD::StoreRetvalV2:
2138
    NumElts = 2;
2139
    break;
2140
  case NVPTXISD::StoreRetvalV4:
2141
    NumElts = 4;
2142
    break;
2143
  }
2144

2145
  // Build vector of operands
2146
  SmallVector<SDValue, 6> Ops;
2147
  for (unsigned i = 0; i < NumElts; ++i)
2148
    Ops.push_back(N->getOperand(i + 2));
2149
  Ops.push_back(CurDAG->getTargetConstant(OffsetVal, DL, MVT::i32));
2150
  Ops.push_back(Chain);
2151

2152
  // Determine target opcode
2153
  // If we have an i1, use an 8-bit store. The lowering code in
2154
  // NVPTXISelLowering will have already emitted an upcast.
2155
  std::optional<unsigned> Opcode = 0;
2156
  switch (NumElts) {
2157
  default:
2158
    return false;
2159
  case 1:
2160
    Opcode = pickOpcodeForVT(Mem->getMemoryVT().getSimpleVT().SimpleTy,
2161
                             NVPTX::StoreRetvalI8, NVPTX::StoreRetvalI16,
2162
                             NVPTX::StoreRetvalI32, NVPTX::StoreRetvalI64,
2163
                             NVPTX::StoreRetvalF32, NVPTX::StoreRetvalF64);
2164
    if (Opcode == NVPTX::StoreRetvalI8) {
2165
      // Fine tune the opcode depending on the size of the operand.
2166
      // This helps to avoid creating redundant COPY instructions in
2167
      // InstrEmitter::AddRegisterOperand().
2168
      switch (Ops[0].getSimpleValueType().SimpleTy) {
2169
      default:
2170
        break;
2171
      case MVT::i32:
2172
        Opcode = NVPTX::StoreRetvalI8TruncI32;
2173
        break;
2174
      case MVT::i64:
2175
        Opcode = NVPTX::StoreRetvalI8TruncI64;
2176
        break;
2177
      }
2178
    }
2179
    break;
2180
  case 2:
2181
    Opcode = pickOpcodeForVT(Mem->getMemoryVT().getSimpleVT().SimpleTy,
2182
                             NVPTX::StoreRetvalV2I8, NVPTX::StoreRetvalV2I16,
2183
                             NVPTX::StoreRetvalV2I32, NVPTX::StoreRetvalV2I64,
2184
                             NVPTX::StoreRetvalV2F32, NVPTX::StoreRetvalV2F64);
2185
    break;
2186
  case 4:
2187
    Opcode = pickOpcodeForVT(Mem->getMemoryVT().getSimpleVT().SimpleTy,
2188
                             NVPTX::StoreRetvalV4I8, NVPTX::StoreRetvalV4I16,
2189
                             NVPTX::StoreRetvalV4I32, std::nullopt,
2190
                             NVPTX::StoreRetvalV4F32, std::nullopt);
2191
    break;
2192
  }
2193
  if (!Opcode)
2194
    return false;
2195

2196
  SDNode *Ret = CurDAG->getMachineNode(*Opcode, DL, MVT::Other, Ops);
2197
  MachineMemOperand *MemRef = cast<MemSDNode>(N)->getMemOperand();
2198
  CurDAG->setNodeMemRefs(cast<MachineSDNode>(Ret), {MemRef});
2199

2200
  ReplaceNode(N, Ret);
2201
  return true;
2202
}
2203

2204
// Helpers for constructing opcode (ex: NVPTX::StoreParamV4F32_iiri)
2205
#define getOpcV2H(ty, opKind0, opKind1)                                        \
2206
  NVPTX::StoreParamV2##ty##_##opKind0##opKind1
2207

2208
#define getOpcV2H1(ty, opKind0, isImm1)                                        \
2209
  (isImm1) ? getOpcV2H(ty, opKind0, i) : getOpcV2H(ty, opKind0, r)
2210

2211
#define getOpcodeForVectorStParamV2(ty, isimm)                                 \
2212
  (isimm[0]) ? getOpcV2H1(ty, i, isimm[1]) : getOpcV2H1(ty, r, isimm[1])
2213

2214
#define getOpcV4H(ty, opKind0, opKind1, opKind2, opKind3)                      \
2215
  NVPTX::StoreParamV4##ty##_##opKind0##opKind1##opKind2##opKind3
2216

2217
#define getOpcV4H3(ty, opKind0, opKind1, opKind2, isImm3)                      \
2218
  (isImm3) ? getOpcV4H(ty, opKind0, opKind1, opKind2, i)                       \
2219
           : getOpcV4H(ty, opKind0, opKind1, opKind2, r)
2220

2221
#define getOpcV4H2(ty, opKind0, opKind1, isImm2, isImm3)                       \
2222
  (isImm2) ? getOpcV4H3(ty, opKind0, opKind1, i, isImm3)                       \
2223
           : getOpcV4H3(ty, opKind0, opKind1, r, isImm3)
2224

2225
#define getOpcV4H1(ty, opKind0, isImm1, isImm2, isImm3)                        \
2226
  (isImm1) ? getOpcV4H2(ty, opKind0, i, isImm2, isImm3)                        \
2227
           : getOpcV4H2(ty, opKind0, r, isImm2, isImm3)
2228

2229
#define getOpcodeForVectorStParamV4(ty, isimm)                                 \
2230
  (isimm[0]) ? getOpcV4H1(ty, i, isimm[1], isimm[2], isimm[3])                 \
2231
             : getOpcV4H1(ty, r, isimm[1], isimm[2], isimm[3])
2232

2233
#define getOpcodeForVectorStParam(n, ty, isimm)                                \
2234
  (n == 2) ? getOpcodeForVectorStParamV2(ty, isimm)                            \
2235
           : getOpcodeForVectorStParamV4(ty, isimm)
2236

2237
static unsigned pickOpcodeForVectorStParam(SmallVector<SDValue, 8> &Ops,
2238
                                           unsigned NumElts,
2239
                                           MVT::SimpleValueType MemTy,
2240
                                           SelectionDAG *CurDAG, SDLoc DL) {
2241
  // Determine which inputs are registers and immediates make new operators
2242
  // with constant values
2243
  SmallVector<bool, 4> IsImm(NumElts, false);
2244
  for (unsigned i = 0; i < NumElts; i++) {
2245
    IsImm[i] = (isa<ConstantSDNode>(Ops[i]) || isa<ConstantFPSDNode>(Ops[i]));
2246
    if (IsImm[i]) {
2247
      SDValue Imm = Ops[i];
2248
      if (MemTy == MVT::f32 || MemTy == MVT::f64) {
2249
        const ConstantFPSDNode *ConstImm = cast<ConstantFPSDNode>(Imm);
2250
        const ConstantFP *CF = ConstImm->getConstantFPValue();
2251
        Imm = CurDAG->getTargetConstantFP(*CF, DL, Imm->getValueType(0));
2252
      } else {
2253
        const ConstantSDNode *ConstImm = cast<ConstantSDNode>(Imm);
2254
        const ConstantInt *CI = ConstImm->getConstantIntValue();
2255
        Imm = CurDAG->getTargetConstant(*CI, DL, Imm->getValueType(0));
2256
      }
2257
      Ops[i] = Imm;
2258
    }
2259
  }
2260

2261
  // Get opcode for MemTy, size, and register/immediate operand ordering
2262
  switch (MemTy) {
2263
  case MVT::i8:
2264
    return getOpcodeForVectorStParam(NumElts, I8, IsImm);
2265
  case MVT::i16:
2266
    return getOpcodeForVectorStParam(NumElts, I16, IsImm);
2267
  case MVT::i32:
2268
    return getOpcodeForVectorStParam(NumElts, I32, IsImm);
2269
  case MVT::i64:
2270
    assert(NumElts == 2 && "MVT too large for NumElts > 2");
2271
    return getOpcodeForVectorStParamV2(I64, IsImm);
2272
  case MVT::f32:
2273
    return getOpcodeForVectorStParam(NumElts, F32, IsImm);
2274
  case MVT::f64:
2275
    assert(NumElts == 2 && "MVT too large for NumElts > 2");
2276
    return getOpcodeForVectorStParamV2(F64, IsImm);
2277

2278
  // These cases don't support immediates, just use the all register version
2279
  // and generate moves.
2280
  case MVT::i1:
2281
    return (NumElts == 2) ? NVPTX::StoreParamV2I8_rr
2282
                          : NVPTX::StoreParamV4I8_rrrr;
2283
  case MVT::f16:
2284
  case MVT::bf16:
2285
    return (NumElts == 2) ? NVPTX::StoreParamV2I16_rr
2286
                          : NVPTX::StoreParamV4I16_rrrr;
2287
  case MVT::v2f16:
2288
  case MVT::v2bf16:
2289
  case MVT::v2i16:
2290
  case MVT::v4i8:
2291
    return (NumElts == 2) ? NVPTX::StoreParamV2I32_rr
2292
                          : NVPTX::StoreParamV4I32_rrrr;
2293
  default:
2294
    llvm_unreachable("Cannot select st.param for unknown MemTy");
2295
  }
2296
}
2297

2298
bool NVPTXDAGToDAGISel::tryStoreParam(SDNode *N) {
2299
  SDLoc DL(N);
2300
  SDValue Chain = N->getOperand(0);
2301
  SDValue Param = N->getOperand(1);
2302
  unsigned ParamVal = Param->getAsZExtVal();
2303
  SDValue Offset = N->getOperand(2);
2304
  unsigned OffsetVal = Offset->getAsZExtVal();
2305
  MemSDNode *Mem = cast<MemSDNode>(N);
2306
  SDValue Glue = N->getOperand(N->getNumOperands() - 1);
2307

2308
  // How many elements do we have?
2309
  unsigned NumElts;
2310
  switch (N->getOpcode()) {
2311
  default:
2312
    llvm_unreachable("Unexpected opcode");
2313
  case NVPTXISD::StoreParamU32:
2314
  case NVPTXISD::StoreParamS32:
2315
  case NVPTXISD::StoreParam:
2316
    NumElts = 1;
2317
    break;
2318
  case NVPTXISD::StoreParamV2:
2319
    NumElts = 2;
2320
    break;
2321
  case NVPTXISD::StoreParamV4:
2322
    NumElts = 4;
2323
    break;
2324
  }
2325

2326
  // Build vector of operands
2327
  SmallVector<SDValue, 8> Ops;
2328
  for (unsigned i = 0; i < NumElts; ++i)
2329
    Ops.push_back(N->getOperand(i + 3));
2330
  Ops.push_back(CurDAG->getTargetConstant(ParamVal, DL, MVT::i32));
2331
  Ops.push_back(CurDAG->getTargetConstant(OffsetVal, DL, MVT::i32));
2332
  Ops.push_back(Chain);
2333
  Ops.push_back(Glue);
2334

2335
  // Determine target opcode
2336
  // If we have an i1, use an 8-bit store. The lowering code in
2337
  // NVPTXISelLowering will have already emitted an upcast.
2338
  std::optional<unsigned> Opcode;
2339
  switch (N->getOpcode()) {
2340
  default:
2341
    switch (NumElts) {
2342
    default:
2343
      llvm_unreachable("Unexpected NumElts");
2344
    case 1: {
2345
      MVT::SimpleValueType MemTy = Mem->getMemoryVT().getSimpleVT().SimpleTy;
2346
      SDValue Imm = Ops[0];
2347
      if (MemTy != MVT::f16 && MemTy != MVT::v2f16 &&
2348
          (isa<ConstantSDNode>(Imm) || isa<ConstantFPSDNode>(Imm))) {
2349
        // Convert immediate to target constant
2350
        if (MemTy == MVT::f32 || MemTy == MVT::f64) {
2351
          const ConstantFPSDNode *ConstImm = cast<ConstantFPSDNode>(Imm);
2352
          const ConstantFP *CF = ConstImm->getConstantFPValue();
2353
          Imm = CurDAG->getTargetConstantFP(*CF, DL, Imm->getValueType(0));
2354
        } else {
2355
          const ConstantSDNode *ConstImm = cast<ConstantSDNode>(Imm);
2356
          const ConstantInt *CI = ConstImm->getConstantIntValue();
2357
          Imm = CurDAG->getTargetConstant(*CI, DL, Imm->getValueType(0));
2358
        }
2359
        Ops[0] = Imm;
2360
        // Use immediate version of store param
2361
        Opcode = pickOpcodeForVT(MemTy, NVPTX::StoreParamI8_i,
2362
                                 NVPTX::StoreParamI16_i, NVPTX::StoreParamI32_i,
2363
                                 NVPTX::StoreParamI64_i, NVPTX::StoreParamF32_i,
2364
                                 NVPTX::StoreParamF64_i);
2365
      } else
2366
        Opcode =
2367
            pickOpcodeForVT(Mem->getMemoryVT().getSimpleVT().SimpleTy,
2368
                            NVPTX::StoreParamI8_r, NVPTX::StoreParamI16_r,
2369
                            NVPTX::StoreParamI32_r, NVPTX::StoreParamI64_r,
2370
                            NVPTX::StoreParamF32_r, NVPTX::StoreParamF64_r);
2371
      if (Opcode == NVPTX::StoreParamI8_r) {
2372
        // Fine tune the opcode depending on the size of the operand.
2373
        // This helps to avoid creating redundant COPY instructions in
2374
        // InstrEmitter::AddRegisterOperand().
2375
        switch (Ops[0].getSimpleValueType().SimpleTy) {
2376
        default:
2377
          break;
2378
        case MVT::i32:
2379
          Opcode = NVPTX::StoreParamI8TruncI32_r;
2380
          break;
2381
        case MVT::i64:
2382
          Opcode = NVPTX::StoreParamI8TruncI64_r;
2383
          break;
2384
        }
2385
      }
2386
      break;
2387
    }
2388
    case 2:
2389
    case 4: {
2390
      MVT::SimpleValueType MemTy = Mem->getMemoryVT().getSimpleVT().SimpleTy;
2391
      Opcode = pickOpcodeForVectorStParam(Ops, NumElts, MemTy, CurDAG, DL);
2392
      break;
2393
    }
2394
    }
2395
    break;
2396
  // Special case: if we have a sign-extend/zero-extend node, insert the
2397
  // conversion instruction first, and use that as the value operand to
2398
  // the selected StoreParam node.
2399
  case NVPTXISD::StoreParamU32: {
2400
    Opcode = NVPTX::StoreParamI32_r;
2401
    SDValue CvtNone = CurDAG->getTargetConstant(NVPTX::PTXCvtMode::NONE, DL,
2402
                                                MVT::i32);
2403
    SDNode *Cvt = CurDAG->getMachineNode(NVPTX::CVT_u32_u16, DL,
2404
                                         MVT::i32, Ops[0], CvtNone);
2405
    Ops[0] = SDValue(Cvt, 0);
2406
    break;
2407
  }
2408
  case NVPTXISD::StoreParamS32: {
2409
    Opcode = NVPTX::StoreParamI32_r;
2410
    SDValue CvtNone = CurDAG->getTargetConstant(NVPTX::PTXCvtMode::NONE, DL,
2411
                                                MVT::i32);
2412
    SDNode *Cvt = CurDAG->getMachineNode(NVPTX::CVT_s32_s16, DL,
2413
                                         MVT::i32, Ops[0], CvtNone);
2414
    Ops[0] = SDValue(Cvt, 0);
2415
    break;
2416
  }
2417
  }
2418

2419
  SDVTList RetVTs = CurDAG->getVTList(MVT::Other, MVT::Glue);
2420
  SDNode *Ret = CurDAG->getMachineNode(*Opcode, DL, RetVTs, Ops);
2421
  MachineMemOperand *MemRef = cast<MemSDNode>(N)->getMemOperand();
2422
  CurDAG->setNodeMemRefs(cast<MachineSDNode>(Ret), {MemRef});
2423

2424
  ReplaceNode(N, Ret);
2425
  return true;
2426
}
2427

2428
bool NVPTXDAGToDAGISel::tryTextureIntrinsic(SDNode *N) {
2429
  unsigned Opc = 0;
2430

2431
  switch (N->getOpcode()) {
2432
  default: return false;
2433
  case NVPTXISD::Tex1DFloatS32:
2434
    Opc = NVPTX::TEX_1D_F32_S32_RR;
2435
    break;
2436
  case NVPTXISD::Tex1DFloatFloat:
2437
    Opc = NVPTX::TEX_1D_F32_F32_RR;
2438
    break;
2439
  case NVPTXISD::Tex1DFloatFloatLevel:
2440
    Opc = NVPTX::TEX_1D_F32_F32_LEVEL_RR;
2441
    break;
2442
  case NVPTXISD::Tex1DFloatFloatGrad:
2443
    Opc = NVPTX::TEX_1D_F32_F32_GRAD_RR;
2444
    break;
2445
  case NVPTXISD::Tex1DS32S32:
2446
    Opc = NVPTX::TEX_1D_S32_S32_RR;
2447
    break;
2448
  case NVPTXISD::Tex1DS32Float:
2449
    Opc = NVPTX::TEX_1D_S32_F32_RR;
2450
    break;
2451
  case NVPTXISD::Tex1DS32FloatLevel:
2452
    Opc = NVPTX::TEX_1D_S32_F32_LEVEL_RR;
2453
    break;
2454
  case NVPTXISD::Tex1DS32FloatGrad:
2455
    Opc = NVPTX::TEX_1D_S32_F32_GRAD_RR;
2456
    break;
2457
  case NVPTXISD::Tex1DU32S32:
2458
    Opc = NVPTX::TEX_1D_U32_S32_RR;
2459
    break;
2460
  case NVPTXISD::Tex1DU32Float:
2461
    Opc = NVPTX::TEX_1D_U32_F32_RR;
2462
    break;
2463
  case NVPTXISD::Tex1DU32FloatLevel:
2464
    Opc = NVPTX::TEX_1D_U32_F32_LEVEL_RR;
2465
    break;
2466
  case NVPTXISD::Tex1DU32FloatGrad:
2467
    Opc = NVPTX::TEX_1D_U32_F32_GRAD_RR;
2468
    break;
2469
  case NVPTXISD::Tex1DArrayFloatS32:
2470
    Opc = NVPTX::TEX_1D_ARRAY_F32_S32_RR;
2471
    break;
2472
  case NVPTXISD::Tex1DArrayFloatFloat:
2473
    Opc = NVPTX::TEX_1D_ARRAY_F32_F32_RR;
2474
    break;
2475
  case NVPTXISD::Tex1DArrayFloatFloatLevel:
2476
    Opc = NVPTX::TEX_1D_ARRAY_F32_F32_LEVEL_RR;
2477
    break;
2478
  case NVPTXISD::Tex1DArrayFloatFloatGrad:
2479
    Opc = NVPTX::TEX_1D_ARRAY_F32_F32_GRAD_RR;
2480
    break;
2481
  case NVPTXISD::Tex1DArrayS32S32:
2482
    Opc = NVPTX::TEX_1D_ARRAY_S32_S32_RR;
2483
    break;
2484
  case NVPTXISD::Tex1DArrayS32Float:
2485
    Opc = NVPTX::TEX_1D_ARRAY_S32_F32_RR;
2486
    break;
2487
  case NVPTXISD::Tex1DArrayS32FloatLevel:
2488
    Opc = NVPTX::TEX_1D_ARRAY_S32_F32_LEVEL_RR;
2489
    break;
2490
  case NVPTXISD::Tex1DArrayS32FloatGrad:
2491
    Opc = NVPTX::TEX_1D_ARRAY_S32_F32_GRAD_RR;
2492
    break;
2493
  case NVPTXISD::Tex1DArrayU32S32:
2494
    Opc = NVPTX::TEX_1D_ARRAY_U32_S32_RR;
2495
    break;
2496
  case NVPTXISD::Tex1DArrayU32Float:
2497
    Opc = NVPTX::TEX_1D_ARRAY_U32_F32_RR;
2498
    break;
2499
  case NVPTXISD::Tex1DArrayU32FloatLevel:
2500
    Opc = NVPTX::TEX_1D_ARRAY_U32_F32_LEVEL_RR;
2501
    break;
2502
  case NVPTXISD::Tex1DArrayU32FloatGrad:
2503
    Opc = NVPTX::TEX_1D_ARRAY_U32_F32_GRAD_RR;
2504
    break;
2505
  case NVPTXISD::Tex2DFloatS32:
2506
    Opc = NVPTX::TEX_2D_F32_S32_RR;
2507
    break;
2508
  case NVPTXISD::Tex2DFloatFloat:
2509
    Opc = NVPTX::TEX_2D_F32_F32_RR;
2510
    break;
2511
  case NVPTXISD::Tex2DFloatFloatLevel:
2512
    Opc = NVPTX::TEX_2D_F32_F32_LEVEL_RR;
2513
    break;
2514
  case NVPTXISD::Tex2DFloatFloatGrad:
2515
    Opc = NVPTX::TEX_2D_F32_F32_GRAD_RR;
2516
    break;
2517
  case NVPTXISD::Tex2DS32S32:
2518
    Opc = NVPTX::TEX_2D_S32_S32_RR;
2519
    break;
2520
  case NVPTXISD::Tex2DS32Float:
2521
    Opc = NVPTX::TEX_2D_S32_F32_RR;
2522
    break;
2523
  case NVPTXISD::Tex2DS32FloatLevel:
2524
    Opc = NVPTX::TEX_2D_S32_F32_LEVEL_RR;
2525
    break;
2526
  case NVPTXISD::Tex2DS32FloatGrad:
2527
    Opc = NVPTX::TEX_2D_S32_F32_GRAD_RR;
2528
    break;
2529
  case NVPTXISD::Tex2DU32S32:
2530
    Opc = NVPTX::TEX_2D_U32_S32_RR;
2531
    break;
2532
  case NVPTXISD::Tex2DU32Float:
2533
    Opc = NVPTX::TEX_2D_U32_F32_RR;
2534
    break;
2535
  case NVPTXISD::Tex2DU32FloatLevel:
2536
    Opc = NVPTX::TEX_2D_U32_F32_LEVEL_RR;
2537
    break;
2538
  case NVPTXISD::Tex2DU32FloatGrad:
2539
    Opc = NVPTX::TEX_2D_U32_F32_GRAD_RR;
2540
    break;
2541
  case NVPTXISD::Tex2DArrayFloatS32:
2542
    Opc = NVPTX::TEX_2D_ARRAY_F32_S32_RR;
2543
    break;
2544
  case NVPTXISD::Tex2DArrayFloatFloat:
2545
    Opc = NVPTX::TEX_2D_ARRAY_F32_F32_RR;
2546
    break;
2547
  case NVPTXISD::Tex2DArrayFloatFloatLevel:
2548
    Opc = NVPTX::TEX_2D_ARRAY_F32_F32_LEVEL_RR;
2549
    break;
2550
  case NVPTXISD::Tex2DArrayFloatFloatGrad:
2551
    Opc = NVPTX::TEX_2D_ARRAY_F32_F32_GRAD_RR;
2552
    break;
2553
  case NVPTXISD::Tex2DArrayS32S32:
2554
    Opc = NVPTX::TEX_2D_ARRAY_S32_S32_RR;
2555
    break;
2556
  case NVPTXISD::Tex2DArrayS32Float:
2557
    Opc = NVPTX::TEX_2D_ARRAY_S32_F32_RR;
2558
    break;
2559
  case NVPTXISD::Tex2DArrayS32FloatLevel:
2560
    Opc = NVPTX::TEX_2D_ARRAY_S32_F32_LEVEL_RR;
2561
    break;
2562
  case NVPTXISD::Tex2DArrayS32FloatGrad:
2563
    Opc = NVPTX::TEX_2D_ARRAY_S32_F32_GRAD_RR;
2564
    break;
2565
  case NVPTXISD::Tex2DArrayU32S32:
2566
    Opc = NVPTX::TEX_2D_ARRAY_U32_S32_RR;
2567
    break;
2568
  case NVPTXISD::Tex2DArrayU32Float:
2569
    Opc = NVPTX::TEX_2D_ARRAY_U32_F32_RR;
2570
    break;
2571
  case NVPTXISD::Tex2DArrayU32FloatLevel:
2572
    Opc = NVPTX::TEX_2D_ARRAY_U32_F32_LEVEL_RR;
2573
    break;
2574
  case NVPTXISD::Tex2DArrayU32FloatGrad:
2575
    Opc = NVPTX::TEX_2D_ARRAY_U32_F32_GRAD_RR;
2576
    break;
2577
  case NVPTXISD::Tex3DFloatS32:
2578
    Opc = NVPTX::TEX_3D_F32_S32_RR;
2579
    break;
2580
  case NVPTXISD::Tex3DFloatFloat:
2581
    Opc = NVPTX::TEX_3D_F32_F32_RR;
2582
    break;
2583
  case NVPTXISD::Tex3DFloatFloatLevel:
2584
    Opc = NVPTX::TEX_3D_F32_F32_LEVEL_RR;
2585
    break;
2586
  case NVPTXISD::Tex3DFloatFloatGrad:
2587
    Opc = NVPTX::TEX_3D_F32_F32_GRAD_RR;
2588
    break;
2589
  case NVPTXISD::Tex3DS32S32:
2590
    Opc = NVPTX::TEX_3D_S32_S32_RR;
2591
    break;
2592
  case NVPTXISD::Tex3DS32Float:
2593
    Opc = NVPTX::TEX_3D_S32_F32_RR;
2594
    break;
2595
  case NVPTXISD::Tex3DS32FloatLevel:
2596
    Opc = NVPTX::TEX_3D_S32_F32_LEVEL_RR;
2597
    break;
2598
  case NVPTXISD::Tex3DS32FloatGrad:
2599
    Opc = NVPTX::TEX_3D_S32_F32_GRAD_RR;
2600
    break;
2601
  case NVPTXISD::Tex3DU32S32:
2602
    Opc = NVPTX::TEX_3D_U32_S32_RR;
2603
    break;
2604
  case NVPTXISD::Tex3DU32Float:
2605
    Opc = NVPTX::TEX_3D_U32_F32_RR;
2606
    break;
2607
  case NVPTXISD::Tex3DU32FloatLevel:
2608
    Opc = NVPTX::TEX_3D_U32_F32_LEVEL_RR;
2609
    break;
2610
  case NVPTXISD::Tex3DU32FloatGrad:
2611
    Opc = NVPTX::TEX_3D_U32_F32_GRAD_RR;
2612
    break;
2613
  case NVPTXISD::TexCubeFloatFloat:
2614
    Opc = NVPTX::TEX_CUBE_F32_F32_RR;
2615
    break;
2616
  case NVPTXISD::TexCubeFloatFloatLevel:
2617
    Opc = NVPTX::TEX_CUBE_F32_F32_LEVEL_RR;
2618
    break;
2619
  case NVPTXISD::TexCubeS32Float:
2620
    Opc = NVPTX::TEX_CUBE_S32_F32_RR;
2621
    break;
2622
  case NVPTXISD::TexCubeS32FloatLevel:
2623
    Opc = NVPTX::TEX_CUBE_S32_F32_LEVEL_RR;
2624
    break;
2625
  case NVPTXISD::TexCubeU32Float:
2626
    Opc = NVPTX::TEX_CUBE_U32_F32_RR;
2627
    break;
2628
  case NVPTXISD::TexCubeU32FloatLevel:
2629
    Opc = NVPTX::TEX_CUBE_U32_F32_LEVEL_RR;
2630
    break;
2631
  case NVPTXISD::TexCubeArrayFloatFloat:
2632
    Opc = NVPTX::TEX_CUBE_ARRAY_F32_F32_RR;
2633
    break;
2634
  case NVPTXISD::TexCubeArrayFloatFloatLevel:
2635
    Opc = NVPTX::TEX_CUBE_ARRAY_F32_F32_LEVEL_RR;
2636
    break;
2637
  case NVPTXISD::TexCubeArrayS32Float:
2638
    Opc = NVPTX::TEX_CUBE_ARRAY_S32_F32_RR;
2639
    break;
2640
  case NVPTXISD::TexCubeArrayS32FloatLevel:
2641
    Opc = NVPTX::TEX_CUBE_ARRAY_S32_F32_LEVEL_RR;
2642
    break;
2643
  case NVPTXISD::TexCubeArrayU32Float:
2644
    Opc = NVPTX::TEX_CUBE_ARRAY_U32_F32_RR;
2645
    break;
2646
  case NVPTXISD::TexCubeArrayU32FloatLevel:
2647
    Opc = NVPTX::TEX_CUBE_ARRAY_U32_F32_LEVEL_RR;
2648
    break;
2649
  case NVPTXISD::Tld4R2DFloatFloat:
2650
    Opc = NVPTX::TLD4_R_2D_F32_F32_RR;
2651
    break;
2652
  case NVPTXISD::Tld4G2DFloatFloat:
2653
    Opc = NVPTX::TLD4_G_2D_F32_F32_RR;
2654
    break;
2655
  case NVPTXISD::Tld4B2DFloatFloat:
2656
    Opc = NVPTX::TLD4_B_2D_F32_F32_RR;
2657
    break;
2658
  case NVPTXISD::Tld4A2DFloatFloat:
2659
    Opc = NVPTX::TLD4_A_2D_F32_F32_RR;
2660
    break;
2661
  case NVPTXISD::Tld4R2DS64Float:
2662
    Opc = NVPTX::TLD4_R_2D_S32_F32_RR;
2663
    break;
2664
  case NVPTXISD::Tld4G2DS64Float:
2665
    Opc = NVPTX::TLD4_G_2D_S32_F32_RR;
2666
    break;
2667
  case NVPTXISD::Tld4B2DS64Float:
2668
    Opc = NVPTX::TLD4_B_2D_S32_F32_RR;
2669
    break;
2670
  case NVPTXISD::Tld4A2DS64Float:
2671
    Opc = NVPTX::TLD4_A_2D_S32_F32_RR;
2672
    break;
2673
  case NVPTXISD::Tld4R2DU64Float:
2674
    Opc = NVPTX::TLD4_R_2D_U32_F32_RR;
2675
    break;
2676
  case NVPTXISD::Tld4G2DU64Float:
2677
    Opc = NVPTX::TLD4_G_2D_U32_F32_RR;
2678
    break;
2679
  case NVPTXISD::Tld4B2DU64Float:
2680
    Opc = NVPTX::TLD4_B_2D_U32_F32_RR;
2681
    break;
2682
  case NVPTXISD::Tld4A2DU64Float:
2683
    Opc = NVPTX::TLD4_A_2D_U32_F32_RR;
2684
    break;
2685
  case NVPTXISD::TexUnified1DFloatS32:
2686
    Opc = NVPTX::TEX_UNIFIED_1D_F32_S32_R;
2687
    break;
2688
  case NVPTXISD::TexUnified1DFloatFloat:
2689
    Opc = NVPTX::TEX_UNIFIED_1D_F32_F32_R;
2690
    break;
2691
  case NVPTXISD::TexUnified1DFloatFloatLevel:
2692
    Opc = NVPTX::TEX_UNIFIED_1D_F32_F32_LEVEL_R;
2693
    break;
2694
  case NVPTXISD::TexUnified1DFloatFloatGrad:
2695
    Opc = NVPTX::TEX_UNIFIED_1D_F32_F32_GRAD_R;
2696
    break;
2697
  case NVPTXISD::TexUnified1DS32S32:
2698
    Opc = NVPTX::TEX_UNIFIED_1D_S32_S32_R;
2699
    break;
2700
  case NVPTXISD::TexUnified1DS32Float:
2701
    Opc = NVPTX::TEX_UNIFIED_1D_S32_F32_R;
2702
    break;
2703
  case NVPTXISD::TexUnified1DS32FloatLevel:
2704
    Opc = NVPTX::TEX_UNIFIED_1D_S32_F32_LEVEL_R;
2705
    break;
2706
  case NVPTXISD::TexUnified1DS32FloatGrad:
2707
    Opc = NVPTX::TEX_UNIFIED_1D_S32_F32_GRAD_R;
2708
    break;
2709
  case NVPTXISD::TexUnified1DU32S32:
2710
    Opc = NVPTX::TEX_UNIFIED_1D_U32_S32_R;
2711
    break;
2712
  case NVPTXISD::TexUnified1DU32Float:
2713
    Opc = NVPTX::TEX_UNIFIED_1D_U32_F32_R;
2714
    break;
2715
  case NVPTXISD::TexUnified1DU32FloatLevel:
2716
    Opc = NVPTX::TEX_UNIFIED_1D_U32_F32_LEVEL_R;
2717
    break;
2718
  case NVPTXISD::TexUnified1DU32FloatGrad:
2719
    Opc = NVPTX::TEX_UNIFIED_1D_U32_F32_GRAD_R;
2720
    break;
2721
  case NVPTXISD::TexUnified1DArrayFloatS32:
2722
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_F32_S32_R;
2723
    break;
2724
  case NVPTXISD::TexUnified1DArrayFloatFloat:
2725
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_F32_F32_R;
2726
    break;
2727
  case NVPTXISD::TexUnified1DArrayFloatFloatLevel:
2728
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_F32_F32_LEVEL_R;
2729
    break;
2730
  case NVPTXISD::TexUnified1DArrayFloatFloatGrad:
2731
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_F32_F32_GRAD_R;
2732
    break;
2733
  case NVPTXISD::TexUnified1DArrayS32S32:
2734
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_S32_S32_R;
2735
    break;
2736
  case NVPTXISD::TexUnified1DArrayS32Float:
2737
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_S32_F32_R;
2738
    break;
2739
  case NVPTXISD::TexUnified1DArrayS32FloatLevel:
2740
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_S32_F32_LEVEL_R;
2741
    break;
2742
  case NVPTXISD::TexUnified1DArrayS32FloatGrad:
2743
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_S32_F32_GRAD_R;
2744
    break;
2745
  case NVPTXISD::TexUnified1DArrayU32S32:
2746
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_U32_S32_R;
2747
    break;
2748
  case NVPTXISD::TexUnified1DArrayU32Float:
2749
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_U32_F32_R;
2750
    break;
2751
  case NVPTXISD::TexUnified1DArrayU32FloatLevel:
2752
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_U32_F32_LEVEL_R;
2753
    break;
2754
  case NVPTXISD::TexUnified1DArrayU32FloatGrad:
2755
    Opc = NVPTX::TEX_UNIFIED_1D_ARRAY_U32_F32_GRAD_R;
2756
    break;
2757
  case NVPTXISD::TexUnified2DFloatS32:
2758
    Opc = NVPTX::TEX_UNIFIED_2D_F32_S32_R;
2759
    break;
2760
  case NVPTXISD::TexUnified2DFloatFloat:
2761
    Opc = NVPTX::TEX_UNIFIED_2D_F32_F32_R;
2762
    break;
2763
  case NVPTXISD::TexUnified2DFloatFloatLevel:
2764
    Opc = NVPTX::TEX_UNIFIED_2D_F32_F32_LEVEL_R;
2765
    break;
2766
  case NVPTXISD::TexUnified2DFloatFloatGrad:
2767
    Opc = NVPTX::TEX_UNIFIED_2D_F32_F32_GRAD_R;
2768
    break;
2769
  case NVPTXISD::TexUnified2DS32S32:
2770
    Opc = NVPTX::TEX_UNIFIED_2D_S32_S32_R;
2771
    break;
2772
  case NVPTXISD::TexUnified2DS32Float:
2773
    Opc = NVPTX::TEX_UNIFIED_2D_S32_F32_R;
2774
    break;
2775
  case NVPTXISD::TexUnified2DS32FloatLevel:
2776
    Opc = NVPTX::TEX_UNIFIED_2D_S32_F32_LEVEL_R;
2777
    break;
2778
  case NVPTXISD::TexUnified2DS32FloatGrad:
2779
    Opc = NVPTX::TEX_UNIFIED_2D_S32_F32_GRAD_R;
2780
    break;
2781
  case NVPTXISD::TexUnified2DU32S32:
2782
    Opc = NVPTX::TEX_UNIFIED_2D_U32_S32_R;
2783
    break;
2784
  case NVPTXISD::TexUnified2DU32Float:
2785
    Opc = NVPTX::TEX_UNIFIED_2D_U32_F32_R;
2786
    break;
2787
  case NVPTXISD::TexUnified2DU32FloatLevel:
2788
    Opc = NVPTX::TEX_UNIFIED_2D_U32_F32_LEVEL_R;
2789
    break;
2790
  case NVPTXISD::TexUnified2DU32FloatGrad:
2791
    Opc = NVPTX::TEX_UNIFIED_2D_U32_F32_GRAD_R;
2792
    break;
2793
  case NVPTXISD::TexUnified2DArrayFloatS32:
2794
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_F32_S32_R;
2795
    break;
2796
  case NVPTXISD::TexUnified2DArrayFloatFloat:
2797
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_F32_F32_R;
2798
    break;
2799
  case NVPTXISD::TexUnified2DArrayFloatFloatLevel:
2800
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_F32_F32_LEVEL_R;
2801
    break;
2802
  case NVPTXISD::TexUnified2DArrayFloatFloatGrad:
2803
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_F32_F32_GRAD_R;
2804
    break;
2805
  case NVPTXISD::TexUnified2DArrayS32S32:
2806
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_S32_S32_R;
2807
    break;
2808
  case NVPTXISD::TexUnified2DArrayS32Float:
2809
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_S32_F32_R;
2810
    break;
2811
  case NVPTXISD::TexUnified2DArrayS32FloatLevel:
2812
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_S32_F32_LEVEL_R;
2813
    break;
2814
  case NVPTXISD::TexUnified2DArrayS32FloatGrad:
2815
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_S32_F32_GRAD_R;
2816
    break;
2817
  case NVPTXISD::TexUnified2DArrayU32S32:
2818
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_U32_S32_R;
2819
    break;
2820
  case NVPTXISD::TexUnified2DArrayU32Float:
2821
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_U32_F32_R;
2822
    break;
2823
  case NVPTXISD::TexUnified2DArrayU32FloatLevel:
2824
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_U32_F32_LEVEL_R;
2825
    break;
2826
  case NVPTXISD::TexUnified2DArrayU32FloatGrad:
2827
    Opc = NVPTX::TEX_UNIFIED_2D_ARRAY_U32_F32_GRAD_R;
2828
    break;
2829
  case NVPTXISD::TexUnified3DFloatS32:
2830
    Opc = NVPTX::TEX_UNIFIED_3D_F32_S32_R;
2831
    break;
2832
  case NVPTXISD::TexUnified3DFloatFloat:
2833
    Opc = NVPTX::TEX_UNIFIED_3D_F32_F32_R;
2834
    break;
2835
  case NVPTXISD::TexUnified3DFloatFloatLevel:
2836
    Opc = NVPTX::TEX_UNIFIED_3D_F32_F32_LEVEL_R;
2837
    break;
2838
  case NVPTXISD::TexUnified3DFloatFloatGrad:
2839
    Opc = NVPTX::TEX_UNIFIED_3D_F32_F32_GRAD_R;
2840
    break;
2841
  case NVPTXISD::TexUnified3DS32S32:
2842
    Opc = NVPTX::TEX_UNIFIED_3D_S32_S32_R;
2843
    break;
2844
  case NVPTXISD::TexUnified3DS32Float:
2845
    Opc = NVPTX::TEX_UNIFIED_3D_S32_F32_R;
2846
    break;
2847
  case NVPTXISD::TexUnified3DS32FloatLevel:
2848
    Opc = NVPTX::TEX_UNIFIED_3D_S32_F32_LEVEL_R;
2849
    break;
2850
  case NVPTXISD::TexUnified3DS32FloatGrad:
2851
    Opc = NVPTX::TEX_UNIFIED_3D_S32_F32_GRAD_R;
2852
    break;
2853
  case NVPTXISD::TexUnified3DU32S32:
2854
    Opc = NVPTX::TEX_UNIFIED_3D_U32_S32_R;
2855
    break;
2856
  case NVPTXISD::TexUnified3DU32Float:
2857
    Opc = NVPTX::TEX_UNIFIED_3D_U32_F32_R;
2858
    break;
2859
  case NVPTXISD::TexUnified3DU32FloatLevel:
2860
    Opc = NVPTX::TEX_UNIFIED_3D_U32_F32_LEVEL_R;
2861
    break;
2862
  case NVPTXISD::TexUnified3DU32FloatGrad:
2863
    Opc = NVPTX::TEX_UNIFIED_3D_U32_F32_GRAD_R;
2864
    break;
2865
  case NVPTXISD::TexUnifiedCubeFloatFloat:
2866
    Opc = NVPTX::TEX_UNIFIED_CUBE_F32_F32_R;
2867
    break;
2868
  case NVPTXISD::TexUnifiedCubeFloatFloatLevel:
2869
    Opc = NVPTX::TEX_UNIFIED_CUBE_F32_F32_LEVEL_R;
2870
    break;
2871
  case NVPTXISD::TexUnifiedCubeS32Float:
2872
    Opc = NVPTX::TEX_UNIFIED_CUBE_S32_F32_R;
2873
    break;
2874
  case NVPTXISD::TexUnifiedCubeS32FloatLevel:
2875
    Opc = NVPTX::TEX_UNIFIED_CUBE_S32_F32_LEVEL_R;
2876
    break;
2877
  case NVPTXISD::TexUnifiedCubeU32Float:
2878
    Opc = NVPTX::TEX_UNIFIED_CUBE_U32_F32_R;
2879
    break;
2880
  case NVPTXISD::TexUnifiedCubeU32FloatLevel:
2881
    Opc = NVPTX::TEX_UNIFIED_CUBE_U32_F32_LEVEL_R;
2882
    break;
2883
  case NVPTXISD::TexUnifiedCubeArrayFloatFloat:
2884
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_F32_F32_R;
2885
    break;
2886
  case NVPTXISD::TexUnifiedCubeArrayFloatFloatLevel:
2887
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_F32_F32_LEVEL_R;
2888
    break;
2889
  case NVPTXISD::TexUnifiedCubeArrayS32Float:
2890
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_S32_F32_R;
2891
    break;
2892
  case NVPTXISD::TexUnifiedCubeArrayS32FloatLevel:
2893
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_S32_F32_LEVEL_R;
2894
    break;
2895
  case NVPTXISD::TexUnifiedCubeArrayU32Float:
2896
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_U32_F32_R;
2897
    break;
2898
  case NVPTXISD::TexUnifiedCubeArrayU32FloatLevel:
2899
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_U32_F32_LEVEL_R;
2900
    break;
2901
  case NVPTXISD::Tld4UnifiedR2DFloatFloat:
2902
    Opc = NVPTX::TLD4_UNIFIED_R_2D_F32_F32_R;
2903
    break;
2904
  case NVPTXISD::Tld4UnifiedG2DFloatFloat:
2905
    Opc = NVPTX::TLD4_UNIFIED_G_2D_F32_F32_R;
2906
    break;
2907
  case NVPTXISD::Tld4UnifiedB2DFloatFloat:
2908
    Opc = NVPTX::TLD4_UNIFIED_B_2D_F32_F32_R;
2909
    break;
2910
  case NVPTXISD::Tld4UnifiedA2DFloatFloat:
2911
    Opc = NVPTX::TLD4_UNIFIED_A_2D_F32_F32_R;
2912
    break;
2913
  case NVPTXISD::Tld4UnifiedR2DS64Float:
2914
    Opc = NVPTX::TLD4_UNIFIED_R_2D_S32_F32_R;
2915
    break;
2916
  case NVPTXISD::Tld4UnifiedG2DS64Float:
2917
    Opc = NVPTX::TLD4_UNIFIED_G_2D_S32_F32_R;
2918
    break;
2919
  case NVPTXISD::Tld4UnifiedB2DS64Float:
2920
    Opc = NVPTX::TLD4_UNIFIED_B_2D_S32_F32_R;
2921
    break;
2922
  case NVPTXISD::Tld4UnifiedA2DS64Float:
2923
    Opc = NVPTX::TLD4_UNIFIED_A_2D_S32_F32_R;
2924
    break;
2925
  case NVPTXISD::Tld4UnifiedR2DU64Float:
2926
    Opc = NVPTX::TLD4_UNIFIED_R_2D_U32_F32_R;
2927
    break;
2928
  case NVPTXISD::Tld4UnifiedG2DU64Float:
2929
    Opc = NVPTX::TLD4_UNIFIED_G_2D_U32_F32_R;
2930
    break;
2931
  case NVPTXISD::Tld4UnifiedB2DU64Float:
2932
    Opc = NVPTX::TLD4_UNIFIED_B_2D_U32_F32_R;
2933
    break;
2934
  case NVPTXISD::Tld4UnifiedA2DU64Float:
2935
    Opc = NVPTX::TLD4_UNIFIED_A_2D_U32_F32_R;
2936
    break;
2937
  case NVPTXISD::TexUnifiedCubeFloatFloatGrad:
2938
    Opc = NVPTX::TEX_UNIFIED_CUBE_F32_F32_GRAD_R;
2939
    break;
2940
  case NVPTXISD::TexUnifiedCubeS32FloatGrad:
2941
    Opc = NVPTX::TEX_UNIFIED_CUBE_S32_F32_GRAD_R;
2942
    break;
2943
  case NVPTXISD::TexUnifiedCubeU32FloatGrad:
2944
    Opc = NVPTX::TEX_UNIFIED_CUBE_U32_F32_GRAD_R;
2945
    break;
2946
  case NVPTXISD::TexUnifiedCubeArrayFloatFloatGrad:
2947
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_F32_F32_GRAD_R;
2948
    break;
2949
  case NVPTXISD::TexUnifiedCubeArrayS32FloatGrad:
2950
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_S32_F32_GRAD_R;
2951
    break;
2952
  case NVPTXISD::TexUnifiedCubeArrayU32FloatGrad:
2953
    Opc = NVPTX::TEX_UNIFIED_CUBE_ARRAY_U32_F32_GRAD_R;
2954
    break;
2955
  }
2956

2957
  // Copy over operands
2958
  SmallVector<SDValue, 8> Ops(drop_begin(N->ops()));
2959
  Ops.push_back(N->getOperand(0)); // Move chain to the back.
2960

2961
  ReplaceNode(N, CurDAG->getMachineNode(Opc, SDLoc(N), N->getVTList(), Ops));
2962
  return true;
2963
}
2964

2965
bool NVPTXDAGToDAGISel::trySurfaceIntrinsic(SDNode *N) {
2966
  unsigned Opc = 0;
2967
  switch (N->getOpcode()) {
2968
  default: return false;
2969
  case NVPTXISD::Suld1DI8Clamp:
2970
    Opc = NVPTX::SULD_1D_I8_CLAMP_R;
2971
    break;
2972
  case NVPTXISD::Suld1DI16Clamp:
2973
    Opc = NVPTX::SULD_1D_I16_CLAMP_R;
2974
    break;
2975
  case NVPTXISD::Suld1DI32Clamp:
2976
    Opc = NVPTX::SULD_1D_I32_CLAMP_R;
2977
    break;
2978
  case NVPTXISD::Suld1DI64Clamp:
2979
    Opc = NVPTX::SULD_1D_I64_CLAMP_R;
2980
    break;
2981
  case NVPTXISD::Suld1DV2I8Clamp:
2982
    Opc = NVPTX::SULD_1D_V2I8_CLAMP_R;
2983
    break;
2984
  case NVPTXISD::Suld1DV2I16Clamp:
2985
    Opc = NVPTX::SULD_1D_V2I16_CLAMP_R;
2986
    break;
2987
  case NVPTXISD::Suld1DV2I32Clamp:
2988
    Opc = NVPTX::SULD_1D_V2I32_CLAMP_R;
2989
    break;
2990
  case NVPTXISD::Suld1DV2I64Clamp:
2991
    Opc = NVPTX::SULD_1D_V2I64_CLAMP_R;
2992
    break;
2993
  case NVPTXISD::Suld1DV4I8Clamp:
2994
    Opc = NVPTX::SULD_1D_V4I8_CLAMP_R;
2995
    break;
2996
  case NVPTXISD::Suld1DV4I16Clamp:
2997
    Opc = NVPTX::SULD_1D_V4I16_CLAMP_R;
2998
    break;
2999
  case NVPTXISD::Suld1DV4I32Clamp:
3000
    Opc = NVPTX::SULD_1D_V4I32_CLAMP_R;
3001
    break;
3002
  case NVPTXISD::Suld1DArrayI8Clamp:
3003
    Opc = NVPTX::SULD_1D_ARRAY_I8_CLAMP_R;
3004
    break;
3005
  case NVPTXISD::Suld1DArrayI16Clamp:
3006
    Opc = NVPTX::SULD_1D_ARRAY_I16_CLAMP_R;
3007
    break;
3008
  case NVPTXISD::Suld1DArrayI32Clamp:
3009
    Opc = NVPTX::SULD_1D_ARRAY_I32_CLAMP_R;
3010
    break;
3011
  case NVPTXISD::Suld1DArrayI64Clamp:
3012
    Opc = NVPTX::SULD_1D_ARRAY_I64_CLAMP_R;
3013
    break;
3014
  case NVPTXISD::Suld1DArrayV2I8Clamp:
3015
    Opc = NVPTX::SULD_1D_ARRAY_V2I8_CLAMP_R;
3016
    break;
3017
  case NVPTXISD::Suld1DArrayV2I16Clamp:
3018
    Opc = NVPTX::SULD_1D_ARRAY_V2I16_CLAMP_R;
3019
    break;
3020
  case NVPTXISD::Suld1DArrayV2I32Clamp:
3021
    Opc = NVPTX::SULD_1D_ARRAY_V2I32_CLAMP_R;
3022
    break;
3023
  case NVPTXISD::Suld1DArrayV2I64Clamp:
3024
    Opc = NVPTX::SULD_1D_ARRAY_V2I64_CLAMP_R;
3025
    break;
3026
  case NVPTXISD::Suld1DArrayV4I8Clamp:
3027
    Opc = NVPTX::SULD_1D_ARRAY_V4I8_CLAMP_R;
3028
    break;
3029
  case NVPTXISD::Suld1DArrayV4I16Clamp:
3030
    Opc = NVPTX::SULD_1D_ARRAY_V4I16_CLAMP_R;
3031
    break;
3032
  case NVPTXISD::Suld1DArrayV4I32Clamp:
3033
    Opc = NVPTX::SULD_1D_ARRAY_V4I32_CLAMP_R;
3034
    break;
3035
  case NVPTXISD::Suld2DI8Clamp:
3036
    Opc = NVPTX::SULD_2D_I8_CLAMP_R;
3037
    break;
3038
  case NVPTXISD::Suld2DI16Clamp:
3039
    Opc = NVPTX::SULD_2D_I16_CLAMP_R;
3040
    break;
3041
  case NVPTXISD::Suld2DI32Clamp:
3042
    Opc = NVPTX::SULD_2D_I32_CLAMP_R;
3043
    break;
3044
  case NVPTXISD::Suld2DI64Clamp:
3045
    Opc = NVPTX::SULD_2D_I64_CLAMP_R;
3046
    break;
3047
  case NVPTXISD::Suld2DV2I8Clamp:
3048
    Opc = NVPTX::SULD_2D_V2I8_CLAMP_R;
3049
    break;
3050
  case NVPTXISD::Suld2DV2I16Clamp:
3051
    Opc = NVPTX::SULD_2D_V2I16_CLAMP_R;
3052
    break;
3053
  case NVPTXISD::Suld2DV2I32Clamp:
3054
    Opc = NVPTX::SULD_2D_V2I32_CLAMP_R;
3055
    break;
3056
  case NVPTXISD::Suld2DV2I64Clamp:
3057
    Opc = NVPTX::SULD_2D_V2I64_CLAMP_R;
3058
    break;
3059
  case NVPTXISD::Suld2DV4I8Clamp:
3060
    Opc = NVPTX::SULD_2D_V4I8_CLAMP_R;
3061
    break;
3062
  case NVPTXISD::Suld2DV4I16Clamp:
3063
    Opc = NVPTX::SULD_2D_V4I16_CLAMP_R;
3064
    break;
3065
  case NVPTXISD::Suld2DV4I32Clamp:
3066
    Opc = NVPTX::SULD_2D_V4I32_CLAMP_R;
3067
    break;
3068
  case NVPTXISD::Suld2DArrayI8Clamp:
3069
    Opc = NVPTX::SULD_2D_ARRAY_I8_CLAMP_R;
3070
    break;
3071
  case NVPTXISD::Suld2DArrayI16Clamp:
3072
    Opc = NVPTX::SULD_2D_ARRAY_I16_CLAMP_R;
3073
    break;
3074
  case NVPTXISD::Suld2DArrayI32Clamp:
3075
    Opc = NVPTX::SULD_2D_ARRAY_I32_CLAMP_R;
3076
    break;
3077
  case NVPTXISD::Suld2DArrayI64Clamp:
3078
    Opc = NVPTX::SULD_2D_ARRAY_I64_CLAMP_R;
3079
    break;
3080
  case NVPTXISD::Suld2DArrayV2I8Clamp:
3081
    Opc = NVPTX::SULD_2D_ARRAY_V2I8_CLAMP_R;
3082
    break;
3083
  case NVPTXISD::Suld2DArrayV2I16Clamp:
3084
    Opc = NVPTX::SULD_2D_ARRAY_V2I16_CLAMP_R;
3085
    break;
3086
  case NVPTXISD::Suld2DArrayV2I32Clamp:
3087
    Opc = NVPTX::SULD_2D_ARRAY_V2I32_CLAMP_R;
3088
    break;
3089
  case NVPTXISD::Suld2DArrayV2I64Clamp:
3090
    Opc = NVPTX::SULD_2D_ARRAY_V2I64_CLAMP_R;
3091
    break;
3092
  case NVPTXISD::Suld2DArrayV4I8Clamp:
3093
    Opc = NVPTX::SULD_2D_ARRAY_V4I8_CLAMP_R;
3094
    break;
3095
  case NVPTXISD::Suld2DArrayV4I16Clamp:
3096
    Opc = NVPTX::SULD_2D_ARRAY_V4I16_CLAMP_R;
3097
    break;
3098
  case NVPTXISD::Suld2DArrayV4I32Clamp:
3099
    Opc = NVPTX::SULD_2D_ARRAY_V4I32_CLAMP_R;
3100
    break;
3101
  case NVPTXISD::Suld3DI8Clamp:
3102
    Opc = NVPTX::SULD_3D_I8_CLAMP_R;
3103
    break;
3104
  case NVPTXISD::Suld3DI16Clamp:
3105
    Opc = NVPTX::SULD_3D_I16_CLAMP_R;
3106
    break;
3107
  case NVPTXISD::Suld3DI32Clamp:
3108
    Opc = NVPTX::SULD_3D_I32_CLAMP_R;
3109
    break;
3110
  case NVPTXISD::Suld3DI64Clamp:
3111
    Opc = NVPTX::SULD_3D_I64_CLAMP_R;
3112
    break;
3113
  case NVPTXISD::Suld3DV2I8Clamp:
3114
    Opc = NVPTX::SULD_3D_V2I8_CLAMP_R;
3115
    break;
3116
  case NVPTXISD::Suld3DV2I16Clamp:
3117
    Opc = NVPTX::SULD_3D_V2I16_CLAMP_R;
3118
    break;
3119
  case NVPTXISD::Suld3DV2I32Clamp:
3120
    Opc = NVPTX::SULD_3D_V2I32_CLAMP_R;
3121
    break;
3122
  case NVPTXISD::Suld3DV2I64Clamp:
3123
    Opc = NVPTX::SULD_3D_V2I64_CLAMP_R;
3124
    break;
3125
  case NVPTXISD::Suld3DV4I8Clamp:
3126
    Opc = NVPTX::SULD_3D_V4I8_CLAMP_R;
3127
    break;
3128
  case NVPTXISD::Suld3DV4I16Clamp:
3129
    Opc = NVPTX::SULD_3D_V4I16_CLAMP_R;
3130
    break;
3131
  case NVPTXISD::Suld3DV4I32Clamp:
3132
    Opc = NVPTX::SULD_3D_V4I32_CLAMP_R;
3133
    break;
3134
  case NVPTXISD::Suld1DI8Trap:
3135
    Opc = NVPTX::SULD_1D_I8_TRAP_R;
3136
    break;
3137
  case NVPTXISD::Suld1DI16Trap:
3138
    Opc = NVPTX::SULD_1D_I16_TRAP_R;
3139
    break;
3140
  case NVPTXISD::Suld1DI32Trap:
3141
    Opc = NVPTX::SULD_1D_I32_TRAP_R;
3142
    break;
3143
  case NVPTXISD::Suld1DI64Trap:
3144
    Opc = NVPTX::SULD_1D_I64_TRAP_R;
3145
    break;
3146
  case NVPTXISD::Suld1DV2I8Trap:
3147
    Opc = NVPTX::SULD_1D_V2I8_TRAP_R;
3148
    break;
3149
  case NVPTXISD::Suld1DV2I16Trap:
3150
    Opc = NVPTX::SULD_1D_V2I16_TRAP_R;
3151
    break;
3152
  case NVPTXISD::Suld1DV2I32Trap:
3153
    Opc = NVPTX::SULD_1D_V2I32_TRAP_R;
3154
    break;
3155
  case NVPTXISD::Suld1DV2I64Trap:
3156
    Opc = NVPTX::SULD_1D_V2I64_TRAP_R;
3157
    break;
3158
  case NVPTXISD::Suld1DV4I8Trap:
3159
    Opc = NVPTX::SULD_1D_V4I8_TRAP_R;
3160
    break;
3161
  case NVPTXISD::Suld1DV4I16Trap:
3162
    Opc = NVPTX::SULD_1D_V4I16_TRAP_R;
3163
    break;
3164
  case NVPTXISD::Suld1DV4I32Trap:
3165
    Opc = NVPTX::SULD_1D_V4I32_TRAP_R;
3166
    break;
3167
  case NVPTXISD::Suld1DArrayI8Trap:
3168
    Opc = NVPTX::SULD_1D_ARRAY_I8_TRAP_R;
3169
    break;
3170
  case NVPTXISD::Suld1DArrayI16Trap:
3171
    Opc = NVPTX::SULD_1D_ARRAY_I16_TRAP_R;
3172
    break;
3173
  case NVPTXISD::Suld1DArrayI32Trap:
3174
    Opc = NVPTX::SULD_1D_ARRAY_I32_TRAP_R;
3175
    break;
3176
  case NVPTXISD::Suld1DArrayI64Trap:
3177
    Opc = NVPTX::SULD_1D_ARRAY_I64_TRAP_R;
3178
    break;
3179
  case NVPTXISD::Suld1DArrayV2I8Trap:
3180
    Opc = NVPTX::SULD_1D_ARRAY_V2I8_TRAP_R;
3181
    break;
3182
  case NVPTXISD::Suld1DArrayV2I16Trap:
3183
    Opc = NVPTX::SULD_1D_ARRAY_V2I16_TRAP_R;
3184
    break;
3185
  case NVPTXISD::Suld1DArrayV2I32Trap:
3186
    Opc = NVPTX::SULD_1D_ARRAY_V2I32_TRAP_R;
3187
    break;
3188
  case NVPTXISD::Suld1DArrayV2I64Trap:
3189
    Opc = NVPTX::SULD_1D_ARRAY_V2I64_TRAP_R;
3190
    break;
3191
  case NVPTXISD::Suld1DArrayV4I8Trap:
3192
    Opc = NVPTX::SULD_1D_ARRAY_V4I8_TRAP_R;
3193
    break;
3194
  case NVPTXISD::Suld1DArrayV4I16Trap:
3195
    Opc = NVPTX::SULD_1D_ARRAY_V4I16_TRAP_R;
3196
    break;
3197
  case NVPTXISD::Suld1DArrayV4I32Trap:
3198
    Opc = NVPTX::SULD_1D_ARRAY_V4I32_TRAP_R;
3199
    break;
3200
  case NVPTXISD::Suld2DI8Trap:
3201
    Opc = NVPTX::SULD_2D_I8_TRAP_R;
3202
    break;
3203
  case NVPTXISD::Suld2DI16Trap:
3204
    Opc = NVPTX::SULD_2D_I16_TRAP_R;
3205
    break;
3206
  case NVPTXISD::Suld2DI32Trap:
3207
    Opc = NVPTX::SULD_2D_I32_TRAP_R;
3208
    break;
3209
  case NVPTXISD::Suld2DI64Trap:
3210
    Opc = NVPTX::SULD_2D_I64_TRAP_R;
3211
    break;
3212
  case NVPTXISD::Suld2DV2I8Trap:
3213
    Opc = NVPTX::SULD_2D_V2I8_TRAP_R;
3214
    break;
3215
  case NVPTXISD::Suld2DV2I16Trap:
3216
    Opc = NVPTX::SULD_2D_V2I16_TRAP_R;
3217
    break;
3218
  case NVPTXISD::Suld2DV2I32Trap:
3219
    Opc = NVPTX::SULD_2D_V2I32_TRAP_R;
3220
    break;
3221
  case NVPTXISD::Suld2DV2I64Trap:
3222
    Opc = NVPTX::SULD_2D_V2I64_TRAP_R;
3223
    break;
3224
  case NVPTXISD::Suld2DV4I8Trap:
3225
    Opc = NVPTX::SULD_2D_V4I8_TRAP_R;
3226
    break;
3227
  case NVPTXISD::Suld2DV4I16Trap:
3228
    Opc = NVPTX::SULD_2D_V4I16_TRAP_R;
3229
    break;
3230
  case NVPTXISD::Suld2DV4I32Trap:
3231
    Opc = NVPTX::SULD_2D_V4I32_TRAP_R;
3232
    break;
3233
  case NVPTXISD::Suld2DArrayI8Trap:
3234
    Opc = NVPTX::SULD_2D_ARRAY_I8_TRAP_R;
3235
    break;
3236
  case NVPTXISD::Suld2DArrayI16Trap:
3237
    Opc = NVPTX::SULD_2D_ARRAY_I16_TRAP_R;
3238
    break;
3239
  case NVPTXISD::Suld2DArrayI32Trap:
3240
    Opc = NVPTX::SULD_2D_ARRAY_I32_TRAP_R;
3241
    break;
3242
  case NVPTXISD::Suld2DArrayI64Trap:
3243
    Opc = NVPTX::SULD_2D_ARRAY_I64_TRAP_R;
3244
    break;
3245
  case NVPTXISD::Suld2DArrayV2I8Trap:
3246
    Opc = NVPTX::SULD_2D_ARRAY_V2I8_TRAP_R;
3247
    break;
3248
  case NVPTXISD::Suld2DArrayV2I16Trap:
3249
    Opc = NVPTX::SULD_2D_ARRAY_V2I16_TRAP_R;
3250
    break;
3251
  case NVPTXISD::Suld2DArrayV2I32Trap:
3252
    Opc = NVPTX::SULD_2D_ARRAY_V2I32_TRAP_R;
3253
    break;
3254
  case NVPTXISD::Suld2DArrayV2I64Trap:
3255
    Opc = NVPTX::SULD_2D_ARRAY_V2I64_TRAP_R;
3256
    break;
3257
  case NVPTXISD::Suld2DArrayV4I8Trap:
3258
    Opc = NVPTX::SULD_2D_ARRAY_V4I8_TRAP_R;
3259
    break;
3260
  case NVPTXISD::Suld2DArrayV4I16Trap:
3261
    Opc = NVPTX::SULD_2D_ARRAY_V4I16_TRAP_R;
3262
    break;
3263
  case NVPTXISD::Suld2DArrayV4I32Trap:
3264
    Opc = NVPTX::SULD_2D_ARRAY_V4I32_TRAP_R;
3265
    break;
3266
  case NVPTXISD::Suld3DI8Trap:
3267
    Opc = NVPTX::SULD_3D_I8_TRAP_R;
3268
    break;
3269
  case NVPTXISD::Suld3DI16Trap:
3270
    Opc = NVPTX::SULD_3D_I16_TRAP_R;
3271
    break;
3272
  case NVPTXISD::Suld3DI32Trap:
3273
    Opc = NVPTX::SULD_3D_I32_TRAP_R;
3274
    break;
3275
  case NVPTXISD::Suld3DI64Trap:
3276
    Opc = NVPTX::SULD_3D_I64_TRAP_R;
3277
    break;
3278
  case NVPTXISD::Suld3DV2I8Trap:
3279
    Opc = NVPTX::SULD_3D_V2I8_TRAP_R;
3280
    break;
3281
  case NVPTXISD::Suld3DV2I16Trap:
3282
    Opc = NVPTX::SULD_3D_V2I16_TRAP_R;
3283
    break;
3284
  case NVPTXISD::Suld3DV2I32Trap:
3285
    Opc = NVPTX::SULD_3D_V2I32_TRAP_R;
3286
    break;
3287
  case NVPTXISD::Suld3DV2I64Trap:
3288
    Opc = NVPTX::SULD_3D_V2I64_TRAP_R;
3289
    break;
3290
  case NVPTXISD::Suld3DV4I8Trap:
3291
    Opc = NVPTX::SULD_3D_V4I8_TRAP_R;
3292
    break;
3293
  case NVPTXISD::Suld3DV4I16Trap:
3294
    Opc = NVPTX::SULD_3D_V4I16_TRAP_R;
3295
    break;
3296
  case NVPTXISD::Suld3DV4I32Trap:
3297
    Opc = NVPTX::SULD_3D_V4I32_TRAP_R;
3298
    break;
3299
  case NVPTXISD::Suld1DI8Zero:
3300
    Opc = NVPTX::SULD_1D_I8_ZERO_R;
3301
    break;
3302
  case NVPTXISD::Suld1DI16Zero:
3303
    Opc = NVPTX::SULD_1D_I16_ZERO_R;
3304
    break;
3305
  case NVPTXISD::Suld1DI32Zero:
3306
    Opc = NVPTX::SULD_1D_I32_ZERO_R;
3307
    break;
3308
  case NVPTXISD::Suld1DI64Zero:
3309
    Opc = NVPTX::SULD_1D_I64_ZERO_R;
3310
    break;
3311
  case NVPTXISD::Suld1DV2I8Zero:
3312
    Opc = NVPTX::SULD_1D_V2I8_ZERO_R;
3313
    break;
3314
  case NVPTXISD::Suld1DV2I16Zero:
3315
    Opc = NVPTX::SULD_1D_V2I16_ZERO_R;
3316
    break;
3317
  case NVPTXISD::Suld1DV2I32Zero:
3318
    Opc = NVPTX::SULD_1D_V2I32_ZERO_R;
3319
    break;
3320
  case NVPTXISD::Suld1DV2I64Zero:
3321
    Opc = NVPTX::SULD_1D_V2I64_ZERO_R;
3322
    break;
3323
  case NVPTXISD::Suld1DV4I8Zero:
3324
    Opc = NVPTX::SULD_1D_V4I8_ZERO_R;
3325
    break;
3326
  case NVPTXISD::Suld1DV4I16Zero:
3327
    Opc = NVPTX::SULD_1D_V4I16_ZERO_R;
3328
    break;
3329
  case NVPTXISD::Suld1DV4I32Zero:
3330
    Opc = NVPTX::SULD_1D_V4I32_ZERO_R;
3331
    break;
3332
  case NVPTXISD::Suld1DArrayI8Zero:
3333
    Opc = NVPTX::SULD_1D_ARRAY_I8_ZERO_R;
3334
    break;
3335
  case NVPTXISD::Suld1DArrayI16Zero:
3336
    Opc = NVPTX::SULD_1D_ARRAY_I16_ZERO_R;
3337
    break;
3338
  case NVPTXISD::Suld1DArrayI32Zero:
3339
    Opc = NVPTX::SULD_1D_ARRAY_I32_ZERO_R;
3340
    break;
3341
  case NVPTXISD::Suld1DArrayI64Zero:
3342
    Opc = NVPTX::SULD_1D_ARRAY_I64_ZERO_R;
3343
    break;
3344
  case NVPTXISD::Suld1DArrayV2I8Zero:
3345
    Opc = NVPTX::SULD_1D_ARRAY_V2I8_ZERO_R;
3346
    break;
3347
  case NVPTXISD::Suld1DArrayV2I16Zero:
3348
    Opc = NVPTX::SULD_1D_ARRAY_V2I16_ZERO_R;
3349
    break;
3350
  case NVPTXISD::Suld1DArrayV2I32Zero:
3351
    Opc = NVPTX::SULD_1D_ARRAY_V2I32_ZERO_R;
3352
    break;
3353
  case NVPTXISD::Suld1DArrayV2I64Zero:
3354
    Opc = NVPTX::SULD_1D_ARRAY_V2I64_ZERO_R;
3355
    break;
3356
  case NVPTXISD::Suld1DArrayV4I8Zero:
3357
    Opc = NVPTX::SULD_1D_ARRAY_V4I8_ZERO_R;
3358
    break;
3359
  case NVPTXISD::Suld1DArrayV4I16Zero:
3360
    Opc = NVPTX::SULD_1D_ARRAY_V4I16_ZERO_R;
3361
    break;
3362
  case NVPTXISD::Suld1DArrayV4I32Zero:
3363
    Opc = NVPTX::SULD_1D_ARRAY_V4I32_ZERO_R;
3364
    break;
3365
  case NVPTXISD::Suld2DI8Zero:
3366
    Opc = NVPTX::SULD_2D_I8_ZERO_R;
3367
    break;
3368
  case NVPTXISD::Suld2DI16Zero:
3369
    Opc = NVPTX::SULD_2D_I16_ZERO_R;
3370
    break;
3371
  case NVPTXISD::Suld2DI32Zero:
3372
    Opc = NVPTX::SULD_2D_I32_ZERO_R;
3373
    break;
3374
  case NVPTXISD::Suld2DI64Zero:
3375
    Opc = NVPTX::SULD_2D_I64_ZERO_R;
3376
    break;
3377
  case NVPTXISD::Suld2DV2I8Zero:
3378
    Opc = NVPTX::SULD_2D_V2I8_ZERO_R;
3379
    break;
3380
  case NVPTXISD::Suld2DV2I16Zero:
3381
    Opc = NVPTX::SULD_2D_V2I16_ZERO_R;
3382
    break;
3383
  case NVPTXISD::Suld2DV2I32Zero:
3384
    Opc = NVPTX::SULD_2D_V2I32_ZERO_R;
3385
    break;
3386
  case NVPTXISD::Suld2DV2I64Zero:
3387
    Opc = NVPTX::SULD_2D_V2I64_ZERO_R;
3388
    break;
3389
  case NVPTXISD::Suld2DV4I8Zero:
3390
    Opc = NVPTX::SULD_2D_V4I8_ZERO_R;
3391
    break;
3392
  case NVPTXISD::Suld2DV4I16Zero:
3393
    Opc = NVPTX::SULD_2D_V4I16_ZERO_R;
3394
    break;
3395
  case NVPTXISD::Suld2DV4I32Zero:
3396
    Opc = NVPTX::SULD_2D_V4I32_ZERO_R;
3397
    break;
3398
  case NVPTXISD::Suld2DArrayI8Zero:
3399
    Opc = NVPTX::SULD_2D_ARRAY_I8_ZERO_R;
3400
    break;
3401
  case NVPTXISD::Suld2DArrayI16Zero:
3402
    Opc = NVPTX::SULD_2D_ARRAY_I16_ZERO_R;
3403
    break;
3404
  case NVPTXISD::Suld2DArrayI32Zero:
3405
    Opc = NVPTX::SULD_2D_ARRAY_I32_ZERO_R;
3406
    break;
3407
  case NVPTXISD::Suld2DArrayI64Zero:
3408
    Opc = NVPTX::SULD_2D_ARRAY_I64_ZERO_R;
3409
    break;
3410
  case NVPTXISD::Suld2DArrayV2I8Zero:
3411
    Opc = NVPTX::SULD_2D_ARRAY_V2I8_ZERO_R;
3412
    break;
3413
  case NVPTXISD::Suld2DArrayV2I16Zero:
3414
    Opc = NVPTX::SULD_2D_ARRAY_V2I16_ZERO_R;
3415
    break;
3416
  case NVPTXISD::Suld2DArrayV2I32Zero:
3417
    Opc = NVPTX::SULD_2D_ARRAY_V2I32_ZERO_R;
3418
    break;
3419
  case NVPTXISD::Suld2DArrayV2I64Zero:
3420
    Opc = NVPTX::SULD_2D_ARRAY_V2I64_ZERO_R;
3421
    break;
3422
  case NVPTXISD::Suld2DArrayV4I8Zero:
3423
    Opc = NVPTX::SULD_2D_ARRAY_V4I8_ZERO_R;
3424
    break;
3425
  case NVPTXISD::Suld2DArrayV4I16Zero:
3426
    Opc = NVPTX::SULD_2D_ARRAY_V4I16_ZERO_R;
3427
    break;
3428
  case NVPTXISD::Suld2DArrayV4I32Zero:
3429
    Opc = NVPTX::SULD_2D_ARRAY_V4I32_ZERO_R;
3430
    break;
3431
  case NVPTXISD::Suld3DI8Zero:
3432
    Opc = NVPTX::SULD_3D_I8_ZERO_R;
3433
    break;
3434
  case NVPTXISD::Suld3DI16Zero:
3435
    Opc = NVPTX::SULD_3D_I16_ZERO_R;
3436
    break;
3437
  case NVPTXISD::Suld3DI32Zero:
3438
    Opc = NVPTX::SULD_3D_I32_ZERO_R;
3439
    break;
3440
  case NVPTXISD::Suld3DI64Zero:
3441
    Opc = NVPTX::SULD_3D_I64_ZERO_R;
3442
    break;
3443
  case NVPTXISD::Suld3DV2I8Zero:
3444
    Opc = NVPTX::SULD_3D_V2I8_ZERO_R;
3445
    break;
3446
  case NVPTXISD::Suld3DV2I16Zero:
3447
    Opc = NVPTX::SULD_3D_V2I16_ZERO_R;
3448
    break;
3449
  case NVPTXISD::Suld3DV2I32Zero:
3450
    Opc = NVPTX::SULD_3D_V2I32_ZERO_R;
3451
    break;
3452
  case NVPTXISD::Suld3DV2I64Zero:
3453
    Opc = NVPTX::SULD_3D_V2I64_ZERO_R;
3454
    break;
3455
  case NVPTXISD::Suld3DV4I8Zero:
3456
    Opc = NVPTX::SULD_3D_V4I8_ZERO_R;
3457
    break;
3458
  case NVPTXISD::Suld3DV4I16Zero:
3459
    Opc = NVPTX::SULD_3D_V4I16_ZERO_R;
3460
    break;
3461
  case NVPTXISD::Suld3DV4I32Zero:
3462
    Opc = NVPTX::SULD_3D_V4I32_ZERO_R;
3463
    break;
3464
  }
3465

3466
  // Copy over operands
3467
  SmallVector<SDValue, 8> Ops(drop_begin(N->ops()));
3468
  Ops.push_back(N->getOperand(0)); // Move chain to the back.
3469

3470
  ReplaceNode(N, CurDAG->getMachineNode(Opc, SDLoc(N), N->getVTList(), Ops));
3471
  return true;
3472
}
3473

3474

3475
/// SelectBFE - Look for instruction sequences that can be made more efficient
3476
/// by using the 'bfe' (bit-field extract) PTX instruction
3477
bool NVPTXDAGToDAGISel::tryBFE(SDNode *N) {
3478
  SDLoc DL(N);
3479
  SDValue LHS = N->getOperand(0);
3480
  SDValue RHS = N->getOperand(1);
3481
  SDValue Len;
3482
  SDValue Start;
3483
  SDValue Val;
3484
  bool IsSigned = false;
3485

3486
  if (N->getOpcode() == ISD::AND) {
3487
    // Canonicalize the operands
3488
    // We want 'and %val, %mask'
3489
    if (isa<ConstantSDNode>(LHS) && !isa<ConstantSDNode>(RHS)) {
3490
      std::swap(LHS, RHS);
3491
    }
3492

3493
    ConstantSDNode *Mask = dyn_cast<ConstantSDNode>(RHS);
3494
    if (!Mask) {
3495
      // We need a constant mask on the RHS of the AND
3496
      return false;
3497
    }
3498

3499
    // Extract the mask bits
3500
    uint64_t MaskVal = Mask->getZExtValue();
3501
    if (!isMask_64(MaskVal)) {
3502
      // We *could* handle shifted masks here, but doing so would require an
3503
      // 'and' operation to fix up the low-order bits so we would trade
3504
      // shr+and for bfe+and, which has the same throughput
3505
      return false;
3506
    }
3507

3508
    // How many bits are in our mask?
3509
    int64_t NumBits = countr_one(MaskVal);
3510
    Len = CurDAG->getTargetConstant(NumBits, DL, MVT::i32);
3511

3512
    if (LHS.getOpcode() == ISD::SRL || LHS.getOpcode() == ISD::SRA) {
3513
      // We have a 'srl/and' pair, extract the effective start bit and length
3514
      Val = LHS.getNode()->getOperand(0);
3515
      Start = LHS.getNode()->getOperand(1);
3516
      ConstantSDNode *StartConst = dyn_cast<ConstantSDNode>(Start);
3517
      if (StartConst) {
3518
        uint64_t StartVal = StartConst->getZExtValue();
3519
        // How many "good" bits do we have left?  "good" is defined here as bits
3520
        // that exist in the original value, not shifted in.
3521
        int64_t GoodBits = Start.getValueSizeInBits() - StartVal;
3522
        if (NumBits > GoodBits) {
3523
          // Do not handle the case where bits have been shifted in. In theory
3524
          // we could handle this, but the cost is likely higher than just
3525
          // emitting the srl/and pair.
3526
          return false;
3527
        }
3528
        Start = CurDAG->getTargetConstant(StartVal, DL, MVT::i32);
3529
      } else {
3530
        // Do not handle the case where the shift amount (can be zero if no srl
3531
        // was found) is not constant. We could handle this case, but it would
3532
        // require run-time logic that would be more expensive than just
3533
        // emitting the srl/and pair.
3534
        return false;
3535
      }
3536
    } else {
3537
      // Do not handle the case where the LHS of the and is not a shift. While
3538
      // it would be trivial to handle this case, it would just transform
3539
      // 'and' -> 'bfe', but 'and' has higher-throughput.
3540
      return false;
3541
    }
3542
  } else if (N->getOpcode() == ISD::SRL || N->getOpcode() == ISD::SRA) {
3543
    if (LHS->getOpcode() == ISD::AND) {
3544
      ConstantSDNode *ShiftCnst = dyn_cast<ConstantSDNode>(RHS);
3545
      if (!ShiftCnst) {
3546
        // Shift amount must be constant
3547
        return false;
3548
      }
3549

3550
      uint64_t ShiftAmt = ShiftCnst->getZExtValue();
3551

3552
      SDValue AndLHS = LHS->getOperand(0);
3553
      SDValue AndRHS = LHS->getOperand(1);
3554

3555
      // Canonicalize the AND to have the mask on the RHS
3556
      if (isa<ConstantSDNode>(AndLHS)) {
3557
        std::swap(AndLHS, AndRHS);
3558
      }
3559

3560
      ConstantSDNode *MaskCnst = dyn_cast<ConstantSDNode>(AndRHS);
3561
      if (!MaskCnst) {
3562
        // Mask must be constant
3563
        return false;
3564
      }
3565

3566
      uint64_t MaskVal = MaskCnst->getZExtValue();
3567
      uint64_t NumZeros;
3568
      uint64_t NumBits;
3569
      if (isMask_64(MaskVal)) {
3570
        NumZeros = 0;
3571
        // The number of bits in the result bitfield will be the number of
3572
        // trailing ones (the AND) minus the number of bits we shift off
3573
        NumBits = llvm::countr_one(MaskVal) - ShiftAmt;
3574
      } else if (isShiftedMask_64(MaskVal)) {
3575
        NumZeros = llvm::countr_zero(MaskVal);
3576
        unsigned NumOnes = llvm::countr_one(MaskVal >> NumZeros);
3577
        // The number of bits in the result bitfield will be the number of
3578
        // trailing zeros plus the number of set bits in the mask minus the
3579
        // number of bits we shift off
3580
        NumBits = NumZeros + NumOnes - ShiftAmt;
3581
      } else {
3582
        // This is not a mask we can handle
3583
        return false;
3584
      }
3585

3586
      if (ShiftAmt < NumZeros) {
3587
        // Handling this case would require extra logic that would make this
3588
        // transformation non-profitable
3589
        return false;
3590
      }
3591

3592
      Val = AndLHS;
3593
      Start = CurDAG->getTargetConstant(ShiftAmt, DL, MVT::i32);
3594
      Len = CurDAG->getTargetConstant(NumBits, DL, MVT::i32);
3595
    } else if (LHS->getOpcode() == ISD::SHL) {
3596
      // Here, we have a pattern like:
3597
      //
3598
      // (sra (shl val, NN), MM)
3599
      // or
3600
      // (srl (shl val, NN), MM)
3601
      //
3602
      // If MM >= NN, we can efficiently optimize this with bfe
3603
      Val = LHS->getOperand(0);
3604

3605
      SDValue ShlRHS = LHS->getOperand(1);
3606
      ConstantSDNode *ShlCnst = dyn_cast<ConstantSDNode>(ShlRHS);
3607
      if (!ShlCnst) {
3608
        // Shift amount must be constant
3609
        return false;
3610
      }
3611
      uint64_t InnerShiftAmt = ShlCnst->getZExtValue();
3612

3613
      SDValue ShrRHS = RHS;
3614
      ConstantSDNode *ShrCnst = dyn_cast<ConstantSDNode>(ShrRHS);
3615
      if (!ShrCnst) {
3616
        // Shift amount must be constant
3617
        return false;
3618
      }
3619
      uint64_t OuterShiftAmt = ShrCnst->getZExtValue();
3620

3621
      // To avoid extra codegen and be profitable, we need Outer >= Inner
3622
      if (OuterShiftAmt < InnerShiftAmt) {
3623
        return false;
3624
      }
3625

3626
      // If the outer shift is more than the type size, we have no bitfield to
3627
      // extract (since we also check that the inner shift is <= the outer shift
3628
      // then this also implies that the inner shift is < the type size)
3629
      if (OuterShiftAmt >= Val.getValueSizeInBits()) {
3630
        return false;
3631
      }
3632

3633
      Start = CurDAG->getTargetConstant(OuterShiftAmt - InnerShiftAmt, DL,
3634
                                        MVT::i32);
3635
      Len = CurDAG->getTargetConstant(Val.getValueSizeInBits() - OuterShiftAmt,
3636
                                      DL, MVT::i32);
3637

3638
      if (N->getOpcode() == ISD::SRA) {
3639
        // If we have a arithmetic right shift, we need to use the signed bfe
3640
        // variant
3641
        IsSigned = true;
3642
      }
3643
    } else {
3644
      // No can do...
3645
      return false;
3646
    }
3647
  } else {
3648
    // No can do...
3649
    return false;
3650
  }
3651

3652

3653
  unsigned Opc;
3654
  // For the BFE operations we form here from "and" and "srl", always use the
3655
  // unsigned variants.
3656
  if (Val.getValueType() == MVT::i32) {
3657
    if (IsSigned) {
3658
      Opc = NVPTX::BFE_S32rii;
3659
    } else {
3660
      Opc = NVPTX::BFE_U32rii;
3661
    }
3662
  } else if (Val.getValueType() == MVT::i64) {
3663
    if (IsSigned) {
3664
      Opc = NVPTX::BFE_S64rii;
3665
    } else {
3666
      Opc = NVPTX::BFE_U64rii;
3667
    }
3668
  } else {
3669
    // We cannot handle this type
3670
    return false;
3671
  }
3672

3673
  SDValue Ops[] = {
3674
    Val, Start, Len
3675
  };
3676

3677
  ReplaceNode(N, CurDAG->getMachineNode(Opc, DL, N->getVTList(), Ops));
3678
  return true;
3679
}
3680

3681
// SelectDirectAddr - Match a direct address for DAG.
3682
// A direct address could be a globaladdress or externalsymbol.
3683
bool NVPTXDAGToDAGISel::SelectDirectAddr(SDValue N, SDValue &Address) {
3684
  // Return true if TGA or ES.
3685
  if (N.getOpcode() == ISD::TargetGlobalAddress ||
3686
      N.getOpcode() == ISD::TargetExternalSymbol) {
3687
    Address = N;
3688
    return true;
3689
  }
3690
  if (N.getOpcode() == NVPTXISD::Wrapper) {
3691
    Address = N.getOperand(0);
3692
    return true;
3693
  }
3694
  // addrspacecast(MoveParam(arg_symbol) to addrspace(PARAM)) -> arg_symbol
3695
  if (AddrSpaceCastSDNode *CastN = dyn_cast<AddrSpaceCastSDNode>(N)) {
3696
    if (CastN->getSrcAddressSpace() == ADDRESS_SPACE_GENERIC &&
3697
        CastN->getDestAddressSpace() == ADDRESS_SPACE_PARAM &&
3698
        CastN->getOperand(0).getOpcode() == NVPTXISD::MoveParam)
3699
      return SelectDirectAddr(CastN->getOperand(0).getOperand(0), Address);
3700
  }
3701
  return false;
3702
}
3703

3704
// symbol+offset
3705
bool NVPTXDAGToDAGISel::SelectADDRsi_imp(
3706
    SDNode *OpNode, SDValue Addr, SDValue &Base, SDValue &Offset, MVT mvt) {
3707
  if (Addr.getOpcode() == ISD::ADD) {
3708
    if (ConstantSDNode *CN = dyn_cast<ConstantSDNode>(Addr.getOperand(1))) {
3709
      SDValue base = Addr.getOperand(0);
3710
      if (SelectDirectAddr(base, Base)) {
3711
        Offset = CurDAG->getTargetConstant(CN->getZExtValue(), SDLoc(OpNode),
3712
                                           mvt);
3713
        return true;
3714
      }
3715
    }
3716
  }
3717
  return false;
3718
}
3719

3720
// symbol+offset
3721
bool NVPTXDAGToDAGISel::SelectADDRsi(SDNode *OpNode, SDValue Addr,
3722
                                     SDValue &Base, SDValue &Offset) {
3723
  return SelectADDRsi_imp(OpNode, Addr, Base, Offset, MVT::i32);
3724
}
3725

3726
// symbol+offset
3727
bool NVPTXDAGToDAGISel::SelectADDRsi64(SDNode *OpNode, SDValue Addr,
3728
                                       SDValue &Base, SDValue &Offset) {
3729
  return SelectADDRsi_imp(OpNode, Addr, Base, Offset, MVT::i64);
3730
}
3731

3732
// register+offset
3733
bool NVPTXDAGToDAGISel::SelectADDRri_imp(
3734
    SDNode *OpNode, SDValue Addr, SDValue &Base, SDValue &Offset, MVT mvt) {
3735
  if (FrameIndexSDNode *FIN = dyn_cast<FrameIndexSDNode>(Addr)) {
3736
    Base = CurDAG->getTargetFrameIndex(FIN->getIndex(), mvt);
3737
    Offset = CurDAG->getTargetConstant(0, SDLoc(OpNode), mvt);
3738
    return true;
3739
  }
3740
  if (Addr.getOpcode() == ISD::TargetExternalSymbol ||
3741
      Addr.getOpcode() == ISD::TargetGlobalAddress)
3742
    return false; // direct calls.
3743

3744
  if (Addr.getOpcode() == ISD::ADD) {
3745
    if (SelectDirectAddr(Addr.getOperand(0), Addr)) {
3746
      return false;
3747
    }
3748
    if (ConstantSDNode *CN = dyn_cast<ConstantSDNode>(Addr.getOperand(1))) {
3749
      if (FrameIndexSDNode *FIN =
3750
              dyn_cast<FrameIndexSDNode>(Addr.getOperand(0)))
3751
        // Constant offset from frame ref.
3752
        Base = CurDAG->getTargetFrameIndex(FIN->getIndex(), mvt);
3753
      else
3754
        Base = Addr.getOperand(0);
3755

3756
      // Offset must fit in a 32-bit signed int in PTX [register+offset] address
3757
      // mode
3758
      if (!CN->getAPIntValue().isSignedIntN(32))
3759
        return false;
3760

3761
      Offset = CurDAG->getTargetConstant(CN->getSExtValue(), SDLoc(OpNode),
3762
                                         MVT::i32);
3763
      return true;
3764
    }
3765
  }
3766
  return false;
3767
}
3768

3769
// register+offset
3770
bool NVPTXDAGToDAGISel::SelectADDRri(SDNode *OpNode, SDValue Addr,
3771
                                     SDValue &Base, SDValue &Offset) {
3772
  return SelectADDRri_imp(OpNode, Addr, Base, Offset, MVT::i32);
3773
}
3774

3775
// register+offset
3776
bool NVPTXDAGToDAGISel::SelectADDRri64(SDNode *OpNode, SDValue Addr,
3777
                                       SDValue &Base, SDValue &Offset) {
3778
  return SelectADDRri_imp(OpNode, Addr, Base, Offset, MVT::i64);
3779
}
3780

3781
bool NVPTXDAGToDAGISel::ChkMemSDNodeAddressSpace(SDNode *N,
3782
                                                 unsigned int spN) const {
3783
  const Value *Src = nullptr;
3784
  if (MemSDNode *mN = dyn_cast<MemSDNode>(N)) {
3785
    if (spN == 0 && mN->getMemOperand()->getPseudoValue())
3786
      return true;
3787
    Src = mN->getMemOperand()->getValue();
3788
  }
3789
  if (!Src)
3790
    return false;
3791
  if (auto *PT = dyn_cast<PointerType>(Src->getType()))
3792
    return (PT->getAddressSpace() == spN);
3793
  return false;
3794
}
3795

3796
/// SelectInlineAsmMemoryOperand - Implement addressing mode selection for
3797
/// inline asm expressions.
3798
bool NVPTXDAGToDAGISel::SelectInlineAsmMemoryOperand(
3799
    const SDValue &Op, InlineAsm::ConstraintCode ConstraintID,
3800
    std::vector<SDValue> &OutOps) {
3801
  SDValue Op0, Op1;
3802
  switch (ConstraintID) {
3803
  default:
3804
    return true;
3805
  case InlineAsm::ConstraintCode::m: // memory
3806
    if (SelectDirectAddr(Op, Op0)) {
3807
      OutOps.push_back(Op0);
3808
      OutOps.push_back(CurDAG->getTargetConstant(0, SDLoc(Op), MVT::i32));
3809
      return false;
3810
    }
3811
    if (SelectADDRri(Op.getNode(), Op, Op0, Op1)) {
3812
      OutOps.push_back(Op0);
3813
      OutOps.push_back(Op1);
3814
      return false;
3815
    }
3816
    break;
3817
  }
3818
  return true;
3819
}
3820

3821
void NVPTXDAGToDAGISel::SelectV2I64toI128(SDNode *N) {
3822
  // Lower a CopyToReg with two 64-bit inputs
3823
  // Dst:i128, lo:i64, hi:i64
3824
  //
3825
  // CopyToReg Dst, lo, hi;
3826
  //
3827
  // ==>
3828
  //
3829
  // tmp = V2I64toI128 {lo, hi};
3830
  // CopyToReg Dst, tmp;
3831
  SDValue Dst = N->getOperand(1);
3832
  SDValue Lo = N->getOperand(2);
3833
  SDValue Hi = N->getOperand(3);
3834

3835
  SDLoc DL(N);
3836
  SDNode *Mov =
3837
      CurDAG->getMachineNode(NVPTX::V2I64toI128, DL, MVT::i128, {Lo, Hi});
3838

3839
  SmallVector<SDValue, 4> NewOps(N->getNumOperands() - 1);
3840
  NewOps[0] = N->getOperand(0);
3841
  NewOps[1] = Dst;
3842
  NewOps[2] = SDValue(Mov, 0);
3843
  if (N->getNumOperands() == 5)
3844
    NewOps[3] = N->getOperand(4);
3845
  SDValue NewValue = CurDAG->getNode(ISD::CopyToReg, DL, SmallVector<EVT>(N->values()), NewOps);
3846

3847
  ReplaceNode(N, NewValue.getNode());
3848
}
3849

3850
void NVPTXDAGToDAGISel::SelectI128toV2I64(SDNode *N) {
3851
  // Lower CopyFromReg from a 128-bit regs to two 64-bit regs
3852
  // Dst:i128, Src:i128
3853
  //
3854
  // {lo, hi} = CopyFromReg Src
3855
  //
3856
  // ==>
3857
  //
3858
  // {lo, hi} = I128toV2I64 Src
3859
  //
3860
  SDValue Ch = N->getOperand(0);
3861
  SDValue Src = N->getOperand(1);
3862
  SDValue Glue = N->getOperand(2);
3863
  SDLoc DL(N);
3864

3865
  // Add Glue and Ch to the operands and results to avoid break the execution
3866
  // order
3867
  SDNode *Mov = CurDAG->getMachineNode(
3868
      NVPTX::I128toV2I64, DL,
3869
      {MVT::i64, MVT::i64, Ch.getValueType(), Glue.getValueType()},
3870
      {Src, Ch, Glue});
3871

3872
  ReplaceNode(N, Mov);
3873
}
3874

3875
/// GetConvertOpcode - Returns the CVT_ instruction opcode that implements a
3876
/// conversion from \p SrcTy to \p DestTy.
3877
unsigned NVPTXDAGToDAGISel::GetConvertOpcode(MVT DestTy, MVT SrcTy,
3878
                                             LoadSDNode *LdNode) {
3879
  bool IsSigned = LdNode && LdNode->getExtensionType() == ISD::SEXTLOAD;
3880
  switch (SrcTy.SimpleTy) {
3881
  default:
3882
    llvm_unreachable("Unhandled source type");
3883
  case MVT::i8:
3884
    switch (DestTy.SimpleTy) {
3885
    default:
3886
      llvm_unreachable("Unhandled dest type");
3887
    case MVT::i16:
3888
      return IsSigned ? NVPTX::CVT_s16_s8 : NVPTX::CVT_u16_u8;
3889
    case MVT::i32:
3890
      return IsSigned ? NVPTX::CVT_s32_s8 : NVPTX::CVT_u32_u8;
3891
    case MVT::i64:
3892
      return IsSigned ? NVPTX::CVT_s64_s8 : NVPTX::CVT_u64_u8;
3893
    }
3894
  case MVT::i16:
3895
    switch (DestTy.SimpleTy) {
3896
    default:
3897
      llvm_unreachable("Unhandled dest type");
3898
    case MVT::i8:
3899
      return IsSigned ? NVPTX::CVT_s8_s16 : NVPTX::CVT_u8_u16;
3900
    case MVT::i32:
3901
      return IsSigned ? NVPTX::CVT_s32_s16 : NVPTX::CVT_u32_u16;
3902
    case MVT::i64:
3903
      return IsSigned ? NVPTX::CVT_s64_s16 : NVPTX::CVT_u64_u16;
3904
    }
3905
  case MVT::i32:
3906
    switch (DestTy.SimpleTy) {
3907
    default:
3908
      llvm_unreachable("Unhandled dest type");
3909
    case MVT::i8:
3910
      return IsSigned ? NVPTX::CVT_s8_s32 : NVPTX::CVT_u8_u32;
3911
    case MVT::i16:
3912
      return IsSigned ? NVPTX::CVT_s16_s32 : NVPTX::CVT_u16_u32;
3913
    case MVT::i64:
3914
      return IsSigned ? NVPTX::CVT_s64_s32 : NVPTX::CVT_u64_u32;
3915
    }
3916
  case MVT::i64:
3917
    switch (DestTy.SimpleTy) {
3918
    default:
3919
      llvm_unreachable("Unhandled dest type");
3920
    case MVT::i8:
3921
      return IsSigned ? NVPTX::CVT_s8_s64 : NVPTX::CVT_u8_u64;
3922
    case MVT::i16:
3923
      return IsSigned ? NVPTX::CVT_s16_s64 : NVPTX::CVT_u16_u64;
3924
    case MVT::i32:
3925
      return IsSigned ? NVPTX::CVT_s32_s64 : NVPTX::CVT_u32_u64;
3926
    }
3927
  case MVT::f16:
3928
    switch (DestTy.SimpleTy) {
3929
    default:
3930
      llvm_unreachable("Unhandled dest type");
3931
    case MVT::f32:
3932
      return NVPTX::CVT_f32_f16;
3933
    case MVT::f64:
3934
      return NVPTX::CVT_f64_f16;
3935
    }
3936
  }
3937
}
3938

3939
Product

Resources

Company
