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 "MCTargetDesc/NVPTXBaseInfo.h"
14#include "NVPTX.h"
15#include "NVPTXISelLowering.h"
16#include "NVPTXSelectionDAGInfo.h"
17#include "NVPTXTargetMachine.h"
18#include "NVPTXUtilities.h"
19#include "llvm/ADT/APInt.h"
20#include "llvm/ADT/MapVector.h"
21#include "llvm/ADT/StringSwitch.h"
22#include "llvm/ADT/Twine.h"
23#include "llvm/Analysis/ValueTracking.h"
24#include "llvm/CodeGen/ISDOpcodes.h"
25#include "llvm/CodeGen/SelectionDAG.h"
26#include "llvm/CodeGen/SelectionDAGISel.h"
27#include "llvm/CodeGen/SelectionDAGNodes.h"
28#include "llvm/IR/Constants.h"
29#include "llvm/IR/DiagnosticInfo.h"
30#include "llvm/IR/InlineAsm.h"
31#include "llvm/IR/Instructions.h"
32#include "llvm/IR/Intrinsics.h"
33#include "llvm/IR/IntrinsicsNVPTX.h"
34#include "llvm/IR/LLVMContext.h"
35#include "llvm/IR/Metadata.h"
36#include "llvm/Support/AtomicOrdering.h"
37#include "llvm/Support/CommandLine.h"
38#include "llvm/Support/ErrorHandling.h"
39#include "llvm/Support/FormatVariadic.h"
40#include "llvm/Support/KnownFPClass.h"
41#include "llvm/Support/MathExtras.h"
42#include "llvm/TargetParser/AtomicScope.h"
43#include <optional>
44
45using namespace llvm;
46
47#define DEBUG_TYPE "nvptx-isel"
48#define PASS_NAME "NVPTX DAG->DAG Pattern Instruction Selection"
49
50static cl::opt<bool>
51 EnableRsqrtOpt("nvptx-rsqrt-approx-opt", cl::init(Val: true), cl::Hidden,
52 cl::desc("Enable reciprocal sqrt optimization"));
53
54// FIXME: This is a WAR to recover lost performance from #155024.
55// We still need to investigate the regression and find a more permanent
56// solution.
57static cl::opt<bool> EnableMADWide("nvptx-mad-wide-opt", cl::init(Val: false),
58 cl::Hidden,
59 cl::desc("Enable MAD wide optimization"));
60
61namespace {
62
63struct NVPTXScopes {
64 NVPTXScopes() = default;
65 NVPTXScopes(LLVMContext &C, const Triple &T);
66 NVPTX::Scope operator[](SyncScope::ID ID) const;
67 bool empty() const;
68
69private:
70 SmallMapVector<SyncScope::ID, NVPTX::Scope, 8> Scopes{};
71 LLVMContext *Context = nullptr;
72};
73
74enum class NVPTXMemCacheHintInstruction { Ld, St, Atom };
75
76struct NVPTXMemCacheHintAccess {
77 NVPTXMemCacheHintInstruction Instruction;
78 NVPTX::AddressSpace AddrSpace;
79 unsigned NumElts;
80 unsigned EltWidth;
81 bool IsVolatile;
82};
83
84struct NVPTXMemCacheHintOperands {
85 SDValue EvictionAndPrefetchHint;
86 SDValue CachePolicyReg;
87};
88
89class NVPTXDAGToDAGISel : public SelectionDAGISel {
90 const NVPTXTargetMachine &TM;
91
92 NVPTX::DivPrecisionLevel getDivF32Level(const SDNode *N) const;
93 bool usePrecSqrtF32(const SDNode *N) const;
94 bool useF32FTZ() const;
95 bool allowFMA() const;
96 bool doRsqrtOpt() const;
97 bool doMADWideOpt() const;
98
99 NVPTXScopes Scopes{};
100
101public:
102 NVPTXDAGToDAGISel() = delete;
103
104 explicit NVPTXDAGToDAGISel(NVPTXTargetMachine &tm, CodeGenOptLevel OptLevel);
105
106 bool runOnMachineFunction(MachineFunction &MF) override;
107 const NVPTXSubtarget *Subtarget = nullptr;
108
109 bool SelectInlineAsmMemoryOperand(const SDValue &Op,
110 InlineAsm::ConstraintCode ConstraintID,
111 std::vector<SDValue> &OutOps) override;
112
113private:
114// Include the pieces autogenerated from the target description.
115#include "NVPTXGenDAGISel.inc"
116
117 void Select(SDNode *N) override;
118 bool tryIntrinsicChain(SDNode *N);
119 bool tryIntrinsicVoid(SDNode *N);
120 void SelectTexSurfHandle(SDNode *N);
121 bool tryLoad(SDNode *N);
122 bool tryLoadVector(SDNode *N);
123 bool tryLDU(SDNode *N);
124 bool tryLDG(MemSDNode *N);
125 bool tryStore(SDNode *N);
126 bool tryStoreVector(SDNode *N);
127 bool tryFence(SDNode *N);
128 bool tryBFE(SDNode *N);
129 bool tryBF16ArithToFMA(SDNode *N);
130 bool tryConstantFP(SDNode *N);
131 bool SelectSETP_F16X2(SDNode *N);
132 bool SelectSETP_BF16X2(SDNode *N);
133 bool tryUNPACK_VECTOR(SDNode *N);
134 bool tryEXTRACT_VECTOR_ELEMENT(SDNode *N);
135 void SelectV2I64toI128(SDNode *N);
136 void SelectI128toV2I64(SDNode *N);
137 void SelectCpAsyncBulkTensorReduceCommon(SDNode *N, unsigned RedOp,
138 bool IsIm2Col = false);
139 void SelectTcgen05Ld(SDNode *N, bool hasOffset = false);
140 void SelectTcgen05St(SDNode *N, bool hasOffset = false);
141 void selectAtomicSwap128(SDNode *N);
142
143 inline SDValue getI32Imm(unsigned Imm, const SDLoc &DL) {
144 return CurDAG->getTargetConstant(Val: Imm, DL, VT: MVT::i32);
145 }
146 NVPTX::Ordering getMemOrder(const MemSDNode *N) const;
147 NVPTX::Scope getAtomicScope(const MemSDNode *N) const;
148
149 bool SelectADDR(SDValue Addr, SDValue &Base, SDValue &Offset);
150 bool SelectFAbs(SDValue N, SDValue &Src);
151 SDValue getPTXCmpMode(const CondCodeSDNode &CondCode);
152 SDValue selectPossiblyImm(SDValue V);
153
154 // Returns the encoded eviction/prefetch hint and cache policy register for a
155 // memory operation. Hints unsupported by the subtarget or address space are
156 // dropped. If L2::cache_hint is active, returns the hint with
157 // L2CacheHintBit set and a register containing the 64-bit cache policy
158 // value. Otherwise returns NOREG for the policy operand.
159 NVPTXMemCacheHintOperands
160 getMemCacheHintOperands(const MemSDNode *N, NVPTXMemCacheHintAccess Access,
161 const SDLoc &DL, bool EmitDiagnostics = true);
162
163 // Returns the Memory Order and Scope that the PTX memory instruction should
164 // use, and inserts appropriate fence instruction before the memory
165 // instruction, if needed to implement the instructions memory order. Required
166 // fences after the instruction need to be handled elsewhere.
167 std::pair<NVPTX::Ordering, NVPTX::Scope>
168 insertMemoryInstructionFence(SDLoc DL, SDValue &Chain, MemSDNode *N);
169 NVPTX::Scope getOperationScope(MemSDNode *N, NVPTX::Ordering O) const;
170
171public:
172 static NVPTX::AddressSpace getAddrSpace(const MemSDNode *N);
173};
174
175class NVPTXDAGToDAGISelLegacy : public SelectionDAGISelLegacy {
176public:
177 static char ID;
178 explicit NVPTXDAGToDAGISelLegacy(NVPTXTargetMachine &tm,
179 CodeGenOptLevel OptLevel);
180};
181
182} // end anonymous namespace
183
184/// createNVPTXISelDag - This pass converts a legalized DAG into a
185/// NVPTX-specific DAG, ready for instruction scheduling.
186FunctionPass *llvm::createNVPTXISelDag(NVPTXTargetMachine &TM,
187 llvm::CodeGenOptLevel OptLevel) {
188 return new NVPTXDAGToDAGISelLegacy(TM, OptLevel);
189}
190
191NVPTXDAGToDAGISelLegacy::NVPTXDAGToDAGISelLegacy(NVPTXTargetMachine &tm,
192 CodeGenOptLevel OptLevel)
193 : SelectionDAGISelLegacy(
194 ID, std::make_unique<NVPTXDAGToDAGISel>(args&: tm, args&: OptLevel)) {}
195
196char NVPTXDAGToDAGISelLegacy::ID = 0;
197
198INITIALIZE_PASS(NVPTXDAGToDAGISelLegacy, DEBUG_TYPE, PASS_NAME, false, false)
199
200NVPTXISelDAGToDAGPass::NVPTXISelDAGToDAGPass(NVPTXTargetMachine &TM,
201 CodeGenOptLevel OptLevel)
202 : SelectionDAGISelPass(std::make_unique<NVPTXDAGToDAGISel>(args&: TM, args&: OptLevel)) {}
203
204NVPTXDAGToDAGISel::NVPTXDAGToDAGISel(NVPTXTargetMachine &tm,
205 CodeGenOptLevel OptLevel)
206 : SelectionDAGISel(tm, OptLevel), TM(tm) {}
207
208bool NVPTXDAGToDAGISel::runOnMachineFunction(MachineFunction &MF) {
209 Subtarget = &MF.getSubtarget<NVPTXSubtarget>();
210 Scopes = NVPTXScopes(MF.getFunction().getContext(),
211 MF.getTarget().getTargetTriple());
212 return SelectionDAGISel::runOnMachineFunction(mf&: MF);
213}
214
215NVPTX::DivPrecisionLevel
216NVPTXDAGToDAGISel::getDivF32Level(const SDNode *N) const {
217 return Subtarget->getTargetLowering()->getDivF32Level(MF: *MF, N: *N);
218}
219
220bool NVPTXDAGToDAGISel::usePrecSqrtF32(const SDNode *N) const {
221 return Subtarget->getTargetLowering()->usePrecSqrtF32(N);
222}
223
224bool NVPTXDAGToDAGISel::useF32FTZ() const {
225 return Subtarget->getTargetLowering()->useF32FTZ(MF: *MF);
226}
227
228bool NVPTXDAGToDAGISel::allowFMA() const {
229 const NVPTXTargetLowering *TL = Subtarget->getTargetLowering();
230 return TL->allowFMA(MF&: *MF, OptLevel);
231}
232
233bool NVPTXDAGToDAGISel::doRsqrtOpt() const { return EnableRsqrtOpt; }
234
235bool NVPTXDAGToDAGISel::doMADWideOpt() const { return EnableMADWide; }
236
237/// Select - Select instructions not customized! Used for
238/// expanded, promoted and normal instructions.
239void NVPTXDAGToDAGISel::Select(SDNode *N) {
240
241 if (N->isMachineOpcode()) {
242 N->setNodeId(-1);
243 return; // Already selected.
244 }
245
246 switch (N->getOpcode()) {
247 case ISD::LOAD:
248 case ISD::ATOMIC_LOAD:
249 case NVPTXISD::MLoad:
250 if (tryLoad(N))
251 return;
252 break;
253 case ISD::STORE:
254 case ISD::ATOMIC_STORE:
255 if (tryStore(N))
256 return;
257 break;
258 case ISD::ATOMIC_FENCE:
259 if (tryFence(N))
260 return;
261 break;
262 case NVPTXISD::UNPACK_VECTOR:
263 tryUNPACK_VECTOR(N);
264 return;
265 case ISD::EXTRACT_VECTOR_ELT:
266 if (tryEXTRACT_VECTOR_ELEMENT(N))
267 return;
268 break;
269 case NVPTXISD::SETP_F16X2:
270 SelectSETP_F16X2(N);
271 return;
272 case NVPTXISD::SETP_BF16X2:
273 SelectSETP_BF16X2(N);
274 return;
275 case NVPTXISD::LoadV2:
276 case NVPTXISD::LoadV4:
277 case NVPTXISD::LoadV8:
278 if (tryLoadVector(N))
279 return;
280 break;
281 case NVPTXISD::LDUV2:
282 case NVPTXISD::LDUV4:
283 if (tryLDU(N))
284 return;
285 break;
286 case NVPTXISD::StoreV2:
287 case NVPTXISD::StoreV4:
288 case NVPTXISD::StoreV8:
289 if (tryStoreVector(N))
290 return;
291 break;
292 case ISD::INTRINSIC_W_CHAIN:
293 if (tryIntrinsicChain(N))
294 return;
295 break;
296 case ISD::INTRINSIC_VOID:
297 if (tryIntrinsicVoid(N))
298 return;
299 break;
300 case ISD::AND:
301 case ISD::SRA:
302 case ISD::SRL:
303 // Try to select BFE
304 if (tryBFE(N))
305 return;
306 break;
307 case ISD::CopyToReg: {
308 if (N->getOperand(Num: 1).getValueType() == MVT::i128) {
309 SelectV2I64toI128(N);
310 return;
311 }
312 break;
313 }
314 case ISD::CopyFromReg: {
315 if (N->getOperand(Num: 1).getValueType() == MVT::i128) {
316 SelectI128toV2I64(N);
317 return;
318 }
319 break;
320 }
321 case NVPTXISD::ATOMIC_CMP_SWAP_B128:
322 case NVPTXISD::ATOMIC_SWAP_B128:
323 selectAtomicSwap128(N);
324 return;
325 case ISD::FADD:
326 case ISD::FMUL:
327 case ISD::FSUB:
328 if (tryBF16ArithToFMA(N))
329 return;
330 break;
331 default:
332 break;
333 }
334 SelectCode(N);
335}
336
337#define TCGEN05_LD_OPCODE(SHAPE, NUM) \
338 (enablePack ? NVPTX::TCGEN05_LD_##SHAPE##_##NUM##_PACK \
339 : NVPTX::TCGEN05_LD_##SHAPE##_##NUM)
340
341static unsigned getTcgen05LdOpcode(unsigned IID, bool enablePack) {
342 switch (IID) {
343 case Intrinsic::nvvm_tcgen05_ld_16x64b_x1:
344 return TCGEN05_LD_OPCODE(16x64b, x1);
345 case Intrinsic::nvvm_tcgen05_ld_16x64b_x2:
346 return TCGEN05_LD_OPCODE(16x64b, x2);
347 case Intrinsic::nvvm_tcgen05_ld_16x64b_x4:
348 return TCGEN05_LD_OPCODE(16x64b, x4);
349 case Intrinsic::nvvm_tcgen05_ld_16x64b_x8:
350 return TCGEN05_LD_OPCODE(16x64b, x8);
351 case Intrinsic::nvvm_tcgen05_ld_16x64b_x16:
352 return TCGEN05_LD_OPCODE(16x64b, x16);
353 case Intrinsic::nvvm_tcgen05_ld_16x64b_x32:
354 return TCGEN05_LD_OPCODE(16x64b, x32);
355 case Intrinsic::nvvm_tcgen05_ld_16x64b_x64:
356 return TCGEN05_LD_OPCODE(16x64b, x64);
357 case Intrinsic::nvvm_tcgen05_ld_16x64b_x128:
358 return TCGEN05_LD_OPCODE(16x64b, x128);
359 case Intrinsic::nvvm_tcgen05_ld_16x128b_x1:
360 return TCGEN05_LD_OPCODE(16x128b, x1);
361 case Intrinsic::nvvm_tcgen05_ld_16x128b_x2:
362 return TCGEN05_LD_OPCODE(16x128b, x2);
363 case Intrinsic::nvvm_tcgen05_ld_16x128b_x4:
364 return TCGEN05_LD_OPCODE(16x128b, x4);
365 case Intrinsic::nvvm_tcgen05_ld_16x128b_x8:
366 return TCGEN05_LD_OPCODE(16x128b, x8);
367 case Intrinsic::nvvm_tcgen05_ld_16x128b_x16:
368 return TCGEN05_LD_OPCODE(16x128b, x16);
369 case Intrinsic::nvvm_tcgen05_ld_16x128b_x32:
370 return TCGEN05_LD_OPCODE(16x128b, x32);
371 case Intrinsic::nvvm_tcgen05_ld_16x128b_x64:
372 return TCGEN05_LD_OPCODE(16x128b, x64);
373 case Intrinsic::nvvm_tcgen05_ld_16x256b_x1:
374 return TCGEN05_LD_OPCODE(16x256b, x1);
375 case Intrinsic::nvvm_tcgen05_ld_16x256b_x2:
376 return TCGEN05_LD_OPCODE(16x256b, x2);
377 case Intrinsic::nvvm_tcgen05_ld_16x256b_x4:
378 return TCGEN05_LD_OPCODE(16x256b, x4);
379 case Intrinsic::nvvm_tcgen05_ld_16x256b_x8:
380 return TCGEN05_LD_OPCODE(16x256b, x8);
381 case Intrinsic::nvvm_tcgen05_ld_16x256b_x16:
382 return TCGEN05_LD_OPCODE(16x256b, x16);
383 case Intrinsic::nvvm_tcgen05_ld_16x256b_x32:
384 return TCGEN05_LD_OPCODE(16x256b, x32);
385 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x1:
386 return TCGEN05_LD_OPCODE(16x32bx2, x1);
387 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x2:
388 return TCGEN05_LD_OPCODE(16x32bx2, x2);
389 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x4:
390 return TCGEN05_LD_OPCODE(16x32bx2, x4);
391 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x8:
392 return TCGEN05_LD_OPCODE(16x32bx2, x8);
393 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x16:
394 return TCGEN05_LD_OPCODE(16x32bx2, x16);
395 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x32:
396 return TCGEN05_LD_OPCODE(16x32bx2, x32);
397 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x64:
398 return TCGEN05_LD_OPCODE(16x32bx2, x64);
399 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x128:
400 return TCGEN05_LD_OPCODE(16x32bx2, x128);
401 case Intrinsic::nvvm_tcgen05_ld_32x32b_x1:
402 return TCGEN05_LD_OPCODE(32x32b, x1);
403 case Intrinsic::nvvm_tcgen05_ld_32x32b_x2:
404 return TCGEN05_LD_OPCODE(32x32b, x2);
405 case Intrinsic::nvvm_tcgen05_ld_32x32b_x4:
406 return TCGEN05_LD_OPCODE(32x32b, x4);
407 case Intrinsic::nvvm_tcgen05_ld_32x32b_x8:
408 return TCGEN05_LD_OPCODE(32x32b, x8);
409 case Intrinsic::nvvm_tcgen05_ld_32x32b_x16:
410 return TCGEN05_LD_OPCODE(32x32b, x16);
411 case Intrinsic::nvvm_tcgen05_ld_32x32b_x32:
412 return TCGEN05_LD_OPCODE(32x32b, x32);
413 case Intrinsic::nvvm_tcgen05_ld_32x32b_x64:
414 return TCGEN05_LD_OPCODE(32x32b, x64);
415 case Intrinsic::nvvm_tcgen05_ld_32x32b_x128:
416 return TCGEN05_LD_OPCODE(32x32b, x128);
417 }
418 llvm_unreachable("unhandled tcgen05.ld lowering");
419}
420
421void NVPTXDAGToDAGISel::SelectTcgen05Ld(SDNode *N, bool hasOffset) {
422 if (!Subtarget->hasTcgen05InstSupport())
423 report_fatal_error(
424 reason: "tcgen05.ld is not supported on this architecture variant");
425
426 SDLoc DL(N);
427 unsigned IID = cast<ConstantSDNode>(Val: N->getOperand(Num: 1))->getZExtValue();
428
429 if (hasOffset) {
430 bool enablePack = cast<ConstantSDNode>(Val: N->getOperand(Num: 4))->getZExtValue();
431 auto OffsetNode = CurDAG->getTargetConstant(
432 Val: cast<ConstantSDNode>(Val: N->getOperand(Num: 3))->getZExtValue(), DL, VT: MVT::i32);
433 ReplaceNode(F: N, T: CurDAG->getMachineNode(
434 Opcode: getTcgen05LdOpcode(IID, enablePack), dl: DL, VTs: N->getVTList(),
435 Ops: {N->getOperand(Num: 2), OffsetNode, N->getOperand(Num: 0)}));
436 } else {
437 bool enablePack = cast<ConstantSDNode>(Val: N->getOperand(Num: 3))->getZExtValue();
438 ReplaceNode(F: N, T: CurDAG->getMachineNode(
439 Opcode: getTcgen05LdOpcode(IID, enablePack), dl: DL, VTs: N->getVTList(),
440 Ops: {N->getOperand(Num: 2), N->getOperand(Num: 0)}));
441 }
442}
443
444bool NVPTXDAGToDAGISel::tryIntrinsicChain(SDNode *N) {
445 unsigned IID = N->getConstantOperandVal(Num: 1);
446 switch (IID) {
447 default:
448 return false;
449 case Intrinsic::nvvm_ldu_global_f:
450 case Intrinsic::nvvm_ldu_global_i:
451 case Intrinsic::nvvm_ldu_global_p:
452 return tryLDU(N);
453
454 case Intrinsic::nvvm_tcgen05_ld_16x64b_x1:
455 case Intrinsic::nvvm_tcgen05_ld_16x64b_x2:
456 case Intrinsic::nvvm_tcgen05_ld_16x64b_x4:
457 case Intrinsic::nvvm_tcgen05_ld_16x64b_x8:
458 case Intrinsic::nvvm_tcgen05_ld_16x64b_x16:
459 case Intrinsic::nvvm_tcgen05_ld_16x64b_x32:
460 case Intrinsic::nvvm_tcgen05_ld_16x64b_x64:
461 case Intrinsic::nvvm_tcgen05_ld_16x64b_x128:
462 case Intrinsic::nvvm_tcgen05_ld_16x128b_x1:
463 case Intrinsic::nvvm_tcgen05_ld_16x128b_x2:
464 case Intrinsic::nvvm_tcgen05_ld_16x128b_x4:
465 case Intrinsic::nvvm_tcgen05_ld_16x128b_x16:
466 case Intrinsic::nvvm_tcgen05_ld_16x128b_x32:
467 case Intrinsic::nvvm_tcgen05_ld_16x128b_x64:
468 case Intrinsic::nvvm_tcgen05_ld_16x256b_x1:
469 case Intrinsic::nvvm_tcgen05_ld_16x128b_x8:
470 case Intrinsic::nvvm_tcgen05_ld_16x256b_x2:
471 case Intrinsic::nvvm_tcgen05_ld_16x256b_x4:
472 case Intrinsic::nvvm_tcgen05_ld_16x256b_x8:
473 case Intrinsic::nvvm_tcgen05_ld_16x256b_x16:
474 case Intrinsic::nvvm_tcgen05_ld_16x256b_x32:
475 case Intrinsic::nvvm_tcgen05_ld_32x32b_x1:
476 case Intrinsic::nvvm_tcgen05_ld_32x32b_x2:
477 case Intrinsic::nvvm_tcgen05_ld_32x32b_x4:
478 case Intrinsic::nvvm_tcgen05_ld_32x32b_x8:
479 case Intrinsic::nvvm_tcgen05_ld_32x32b_x16:
480 case Intrinsic::nvvm_tcgen05_ld_32x32b_x32:
481 case Intrinsic::nvvm_tcgen05_ld_32x32b_x64:
482 case Intrinsic::nvvm_tcgen05_ld_32x32b_x128: {
483 SelectTcgen05Ld(N);
484 return true;
485 }
486
487 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x1:
488 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x2:
489 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x4:
490 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x8:
491 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x16:
492 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x32:
493 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x64:
494 case Intrinsic::nvvm_tcgen05_ld_16x32bx2_x128: {
495 SelectTcgen05Ld(N, /* hasOffset */ true);
496 return true;
497 }
498 }
499}
500
501// Map ISD:CONDCODE value to appropriate CmpMode expected by
502// NVPTXInstPrinter::printCmpMode()
503SDValue NVPTXDAGToDAGISel::getPTXCmpMode(const CondCodeSDNode &CondCode) {
504 using NVPTX::PTXCmpMode::CmpMode;
505 const unsigned PTXCmpMode = [](ISD::CondCode CC) {
506 switch (CC) {
507 default:
508 llvm_unreachable("Unexpected condition code.");
509 case ISD::SETOEQ:
510 case ISD::SETEQ:
511 return CmpMode::EQ;
512 case ISD::SETOGT:
513 case ISD::SETGT:
514 return CmpMode::GT;
515 case ISD::SETOGE:
516 case ISD::SETGE:
517 return CmpMode::GE;
518 case ISD::SETOLT:
519 case ISD::SETLT:
520 return CmpMode::LT;
521 case ISD::SETOLE:
522 case ISD::SETLE:
523 return CmpMode::LE;
524 case ISD::SETONE:
525 case ISD::SETNE:
526 return CmpMode::NE;
527 case ISD::SETO:
528 return CmpMode::NUM;
529 case ISD::SETUO:
530 return CmpMode::NotANumber;
531 case ISD::SETUEQ:
532 return CmpMode::EQU;
533 case ISD::SETUGT:
534 return CmpMode::GTU;
535 case ISD::SETUGE:
536 return CmpMode::GEU;
537 case ISD::SETULT:
538 return CmpMode::LTU;
539 case ISD::SETULE:
540 return CmpMode::LEU;
541 case ISD::SETUNE:
542 return CmpMode::NEU;
543 }
544 }(CondCode.get());
545 return CurDAG->getTargetConstant(Val: PTXCmpMode, DL: SDLoc(), VT: MVT::i32);
546}
547
548bool NVPTXDAGToDAGISel::SelectSETP_F16X2(SDNode *N) {
549 SDValue PTXCmpMode = getPTXCmpMode(CondCode: *cast<CondCodeSDNode>(Val: N->getOperand(Num: 2)));
550 SDLoc DL(N);
551 SDNode *SetP = CurDAG->getMachineNode(
552 Opcode: NVPTX::SETP_f16x2rr, dl: DL, VT1: MVT::i1, VT2: MVT::i1,
553 Ops: {N->getOperand(Num: 0), N->getOperand(Num: 1), PTXCmpMode,
554 CurDAG->getTargetConstant(Val: useF32FTZ() ? 1 : 0, DL, VT: MVT::i1)});
555 ReplaceNode(F: N, T: SetP);
556 return true;
557}
558
559bool NVPTXDAGToDAGISel::SelectSETP_BF16X2(SDNode *N) {
560 SDValue PTXCmpMode = getPTXCmpMode(CondCode: *cast<CondCodeSDNode>(Val: N->getOperand(Num: 2)));
561 SDLoc DL(N);
562 SDNode *SetP =
563 CurDAG->getMachineNode(Opcode: NVPTX::SETP_bf16x2rr, dl: DL, VT1: MVT::i1, VT2: MVT::i1,
564 Ops: {N->getOperand(Num: 0), N->getOperand(Num: 1), PTXCmpMode});
565 ReplaceNode(F: N, T: SetP);
566 return true;
567}
568
569bool NVPTXDAGToDAGISel::tryUNPACK_VECTOR(SDNode *N) {
570 SDValue Vector = N->getOperand(Num: 0);
571 MVT EltVT = N->getSimpleValueType(ResNo: 0);
572
573 MachineSDNode *N2 =
574 CurDAG->getMachineNode(Opcode: NVPTX::I64toV2I32, dl: SDLoc(N), VT1: EltVT, VT2: EltVT, Ops: Vector);
575
576 ReplaceNode(F: N, T: N2);
577 return true;
578}
579
580// Find all instances of extract_vector_elt that use this v2f16 vector
581// and coalesce them into a scattering move instruction.
582bool NVPTXDAGToDAGISel::tryEXTRACT_VECTOR_ELEMENT(SDNode *N) {
583 SDValue Vector = N->getOperand(Num: 0);
584
585 MVT VT = Vector.getSimpleValueType();
586 if (!(NVPTX::isPackedVectorTy(VT) && VT.getVectorNumElements() == 2))
587 return false;
588
589 unsigned Opcode;
590 if (VT.is32BitVector())
591 Opcode = NVPTX::I32toV2I16;
592 else if (VT.is64BitVector())
593 Opcode = NVPTX::I64toV2I32;
594 else
595 llvm_unreachable("Unhandled packed type");
596
597 // Find and record all uses of this vector that extract element 0 or 1.
598 SmallVector<SDNode *, 4> E0, E1;
599 for (auto *U : Vector.getNode()->users()) {
600 if (U->getOpcode() != ISD::EXTRACT_VECTOR_ELT)
601 continue;
602 if (U->getOperand(Num: 0) != Vector)
603 continue;
604 if (const ConstantSDNode *IdxConst =
605 dyn_cast<ConstantSDNode>(Val: U->getOperand(Num: 1))) {
606 if (IdxConst->getZExtValue() == 0)
607 E0.push_back(Elt: U);
608 else if (IdxConst->getZExtValue() == 1)
609 E1.push_back(Elt: U);
610 else
611 llvm_unreachable("Invalid vector index.");
612 }
613 }
614
615 // There's no point scattering f16x2 if we only ever access one
616 // element of it.
617 if (E0.empty() || E1.empty())
618 return false;
619
620 // Merge (EltTy extractelt(V, 0), EltTy extractelt(V,1))
621 // into EltTy,EltTy Split[EltTy]x2(V)
622 MVT EltVT = VT.getVectorElementType();
623 SDNode *ScatterOp =
624 CurDAG->getMachineNode(Opcode, dl: SDLoc(N), VT1: EltVT, VT2: EltVT, Ops: Vector);
625 for (auto *Node : E0)
626 ReplaceUses(F: SDValue(Node, 0), T: SDValue(ScatterOp, 0));
627 for (auto *Node : E1)
628 ReplaceUses(F: SDValue(Node, 0), T: SDValue(ScatterOp, 1));
629
630 return true;
631}
632
633NVPTX::AddressSpace NVPTXDAGToDAGISel::getAddrSpace(const MemSDNode *N) {
634 auto AS =
635 static_cast<NVPTX::AddressSpace>(N->getMemOperand()->getAddrSpace());
636 switch (AS) {
637 case NVPTX::AddressSpace::Generic:
638 case NVPTX::AddressSpace::Global:
639 case NVPTX::AddressSpace::Shared:
640 case NVPTX::AddressSpace::Const:
641 case NVPTX::AddressSpace::Local:
642 case NVPTX::AddressSpace::SharedCluster:
643 case NVPTX::AddressSpace::EntryParam:
644 case NVPTX::AddressSpace::DeviceParam:
645 return AS;
646 }
647 llvm_unreachable("Unexpected address space");
648}
649
650NVPTX::Ordering NVPTXDAGToDAGISel::getMemOrder(const MemSDNode *N) const {
651 // No "sem" orderings for SM/PTX versions which do not support memory ordering
652 if (!Subtarget->hasMemoryOrdering())
653 return NVPTX::Ordering::NotAtomic;
654 auto Ordering = N->getMergedOrdering();
655 switch (Ordering) {
656 case AtomicOrdering::NotAtomic:
657 return NVPTX::Ordering::NotAtomic;
658 case AtomicOrdering::Unordered:
659 case AtomicOrdering::Monotonic:
660 return NVPTX::Ordering::Relaxed;
661 case AtomicOrdering::Acquire:
662 return NVPTX::Ordering::Acquire;
663 case AtomicOrdering::Release:
664 return NVPTX::Ordering::Release;
665 case AtomicOrdering::AcquireRelease:
666 return NVPTX::Ordering::AcquireRelease;
667 case AtomicOrdering::SequentiallyConsistent:
668 return NVPTX::Ordering::SequentiallyConsistent;
669 }
670 llvm_unreachable("Invalid atomic ordering");
671}
672
673// Clusters contain exactly 1 block on targets without cluster support.
674static NVPTX::Scope resolveScope(NVPTX::Scope S, const NVPTXSubtarget *T) {
675 if (S == NVPTX::Scope::Cluster && !T->hasClusters())
676 return NVPTX::Scope::Block;
677 return S;
678}
679
680NVPTX::Scope NVPTXDAGToDAGISel::getAtomicScope(const MemSDNode *N) const {
681 NVPTX::Scope Scope = resolveScope(S: Scopes[N->getSyncScopeID()], T: Subtarget);
682 if (!Subtarget->hasAtomScope()) {
683 if (Scope == NVPTX::Scope::System)
684 CurDAG->getContext()->diagnose(DI: DiagnosticInfoUnsupported(
685 CurDAG->getMachineFunction().getFunction(),
686 "NVPTX system scope atomics require sm_60 or later",
687 N->getDebugLoc()));
688 return NVPTX::Scope::DefaultDevice;
689 }
690 return Scope;
691}
692
693namespace {
694
695struct OperationOrderings {
696 NVPTX::Ordering InstructionOrdering, FenceOrdering;
697 OperationOrderings(NVPTX::Ordering IO = NVPTX::Ordering::NotAtomic,
698 NVPTX::Ordering FO = NVPTX::Ordering::NotAtomic)
699 : InstructionOrdering(IO), FenceOrdering(FO) {}
700};
701
702static OperationOrderings
703getOperationOrderings(MemSDNode *N, const NVPTXSubtarget *Subtarget) {
704 AtomicOrdering Ordering = N->getSuccessOrdering();
705 auto CodeAddrSpace = NVPTXDAGToDAGISel::getAddrSpace(N);
706
707 bool HasMemoryOrdering = Subtarget->hasMemoryOrdering();
708 bool HasRelaxedMMIO = Subtarget->hasRelaxedMMIO();
709 bool IsSupportedLocalVolatile = CodeAddrSpace == NVPTX::AddressSpace::Local &&
710 Subtarget->hasFeature(Feature: NVPTX::PTX91) &&
711 N->isVolatile() &&
712 (Ordering == AtomicOrdering::NotAtomic ||
713 Ordering == AtomicOrdering::Unordered ||
714 Ordering == AtomicOrdering::Monotonic);
715
716 // clang-format off
717
718 // Lowering for Load/Store Operations (note: AcquireRelease Loads or Stores error).
719 // Note: uses of Relaxed in the Atomic column of this table refer
720 // to LLVM AtomicOrdering::Monotonic.
721 //
722 // | Atomic | Volatile | Statespace | PTX sm_60- | PTX sm_70+ |
723 // |---------|----------|--------------------|------------|------------------------------|
724 // | No | No | All | plain | .weak |
725 // | No | Yes | Generic,Shared, | .volatile | .volatile |
726 // | | | Global [0] | | |
727 // | No | Yes | Local (PTX 9.0-) | plain [1] | .weak [1] |
728 // | No | Yes | Local (PTX 9.1+) | .volatile | .volatile |
729 // | No | Yes | Const,Param | plain [1] | .weak [1] |
730 // | Unorder | Yes/No | All | == Relaxed | == Relaxed |
731 // | Relaxed | No | Generic,Shared, | .volatile | <atomic sem> |
732 // | | | Global [0] | | |
733 // | Other | No | Generic,Shared, | Error [2] | <atomic sem> |
734 // | | | Global [0] | | |
735 // | Yes | No | Local,Const,Param | plain [1] | .weak [1] |
736 // | Relaxed | Yes | Generic,Shared [0] | .volatile | .volatile |
737 // | Relaxed | Yes | Global [0] | .volatile | .mmio.relaxed.sys (PTX 8.2+) |
738 // | | | | | or .volatile (PTX 8.1-) |
739 // | Relaxed | Yes | Local (PTX 9.0-) | plain [1] | .weak [1] |
740 // | Relaxed | Yes | Local (PTX 9.1+) | .volatile | .volatile |
741 // | Relaxed | Yes | Const,Param | plain [1] | .weak [1] |
742 // | Other | Yes | Generic, Shared, | Error [2] | <atomic sem> [3] |
743 // | | | / Global [0] | | |
744
745 // Lowering of CUDA C++ SequentiallyConsistent Operations and Fences to PTX
746 // by following the ABI proven sound in:
747 // Lustig et al, A Formal Analysis of the NVIDIA PTX Memory Consistency Model, ASPLOS’19.
748 // https://dl.acm.org/doi/pdf/10.1145/3297858.3304043
749 //
750 // | CUDA C++ Atomic Operation or Atomic Fence | PTX Atomic Operation or Fence |
751 // |------------------------------------------------------|-------------------------------|
752 // | cuda::atomic_thread_fence | fence.sc.<scope>; |
753 // | (memory_order_seq_cst, cuda::thread_scope_<scope>) | |
754 // |------------------------------------------------------|-------------------------------|
755 // | cuda::atomic_load | fence.sc.<scope>; |
756 // | (memory_order_seq_cst, cuda::thread_scope_<scope>) | ld.acquire.<scope>; |
757 // |------------------------------------------------------|-------------------------------|
758 // | cuda::atomic_store | fence.sc.<scope>; |
759 // | (memory_order_seq_cst, cuda::thread_scope_<scope>) | st.release.<scope>; |
760 // |------------------------------------------------------|-------------------------------|
761 // | cuda::atomic_fetch_<op> | fence.sc.<scope>; |
762 // | (memory_order_seq_cst, cuda::thread_scope_<scope>) | atom.acq_rel.<scope>; |
763
764 // clang-format on
765
766 // [0]: volatile and atomics are only supported on global or shared
767 // memory locations, accessed via generic/shared/global pointers.
768 // PTX 9.1 adds volatile support on local ld/st.
769 // MMIO is only supported on global memory locations,
770 // accessed via generic/global pointers.
771 // TODO: Implement MMIO access via generic pointer to global.
772 // Currently implemented for global pointers only.
773
774 // [1]: Lowering volatile/atomic operations to non-volatile/non-atomic
775 // PTX instructions fails to preserve their C++ side-effects.
776 //
777 // Example (https://github.com/llvm/llvm-project/issues/62057):
778 //
779 // void example() {
780 // std::atomic<bool> True = true;
781 // while (True.load(std::memory_order_relaxed));
782 // }
783 //
784 // A C++ program that calls "example" is well-defined: the infinite loop
785 // performs an atomic operation. By lowering volatile/atomics to
786 // "weak" memory operations, we are transforming the above into:
787 //
788 // void undefined_behavior() {
789 // bool True = true;
790 // while (True);
791 // }
792 //
793 // which exhibits undefined behavior in both C++ and PTX.
794 //
795 // Calling "example" in CUDA C++ compiled for sm_60- exhibits undefined
796 // behavior due to lack of Independent Forward Progress. Lowering these
797 // to weak memory operations in sm_60- is therefore fine.
798 //
799 // TODO: Where direct volatile or atomic operations are unsupported,
800 // preserve the side-effect using the weak memory instruction and
801 // another instruction, such as a dead dummy volatile load.
802
803 if ((CodeAddrSpace == NVPTX::AddressSpace::Local &&
804 !IsSupportedLocalVolatile) ||
805 CodeAddrSpace == NVPTX::AddressSpace::Const ||
806 CodeAddrSpace == NVPTX::AddressSpace::EntryParam ||
807 CodeAddrSpace == NVPTX::AddressSpace::DeviceParam) {
808 return NVPTX::Ordering::NotAtomic;
809 }
810
811 // [2]: Atomics with Ordering different than Unordered or Relaxed are not
812 // supported on sm_60 and older; this includes volatile atomics.
813 if (!(Ordering == AtomicOrdering::NotAtomic ||
814 Ordering == AtomicOrdering::Unordered ||
815 Ordering == AtomicOrdering::Monotonic) &&
816 !HasMemoryOrdering) {
817 report_fatal_error(
818 reason: formatv(Fmt: "PTX does not support \"atomic\" for orderings different than"
819 "\"NotAtomic\" or \"Monotonic\" for sm_60 or older, but order "
820 "is: \"{}\".",
821 Vals: toIRString(ao: Ordering)));
822 }
823
824 // [3]: TODO: these should eventually use .mmio<.atomic sem>; for now we drop
825 // the volatile semantics and preserve the atomic ones.
826
827 // PTX atomics are not available outside generic, global, or shared memory.
828 // PTX volatile operations additionally support local memory in PTX 9.1+.
829 bool AddrSupportsVolatileOrAtomic =
830 (IsSupportedLocalVolatile ||
831 CodeAddrSpace == NVPTX::AddressSpace::Generic ||
832 CodeAddrSpace == NVPTX::AddressSpace::Global ||
833 CodeAddrSpace == NVPTX::AddressSpace::Shared ||
834 CodeAddrSpace == NVPTX::AddressSpace::SharedCluster);
835 if (!AddrSupportsVolatileOrAtomic)
836 return NVPTX::Ordering::NotAtomic;
837
838 bool UseRelaxedMMIO =
839 HasRelaxedMMIO && CodeAddrSpace == NVPTX::AddressSpace::Global;
840
841 switch (Ordering) {
842 case AtomicOrdering::NotAtomic:
843 return N->isVolatile() ? NVPTX::Ordering::Volatile
844 : NVPTX::Ordering::NotAtomic;
845 case AtomicOrdering::Unordered:
846 // We lower unordered in the exact same way as 'monotonic' to respect
847 // LLVM IR atomicity requirements.
848 case AtomicOrdering::Monotonic:
849 if (N->isVolatile())
850 return UseRelaxedMMIO ? NVPTX::Ordering::RelaxedMMIO
851 : NVPTX::Ordering::Volatile;
852 else
853 return HasMemoryOrdering ? NVPTX::Ordering::Relaxed
854 : NVPTX::Ordering::Volatile;
855 // case AtomicOrdering::Consume: // If LLVM ever provides this, lower it to
856 // Acquire.
857 case AtomicOrdering::Acquire:
858 if (!N->readMem())
859 report_fatal_error(
860 reason: formatv(Fmt: "PTX only supports Acquire Ordering on reads: {}",
861 Vals: N->getOperationName()));
862 return NVPTX::Ordering::Acquire;
863 case AtomicOrdering::Release:
864 if (!N->writeMem())
865 report_fatal_error(
866 reason: formatv(Fmt: "PTX only supports Release Ordering on writes: {}",
867 Vals: N->getOperationName()));
868 return NVPTX::Ordering::Release;
869 case AtomicOrdering::AcquireRelease: {
870 report_fatal_error(
871 reason: formatv(Fmt: "NVPTX does not support AcquireRelease Ordering on "
872 "read-modify-write "
873 "yet and PTX does not support it on loads or stores: {}",
874 Vals: N->getOperationName()));
875 }
876 case AtomicOrdering::SequentiallyConsistent: {
877 // LLVM-IR SequentiallyConsistent atomics map to a two-instruction PTX
878 // sequence including a "fence.sc.sco" and the memory instruction with an
879 // Ordering that differs from "sc": acq, rel, or acq_rel, depending on
880 // whether the memory operation is a read, write, or read-modify-write.
881 //
882 // This sets the ordering of the fence to SequentiallyConsistent, and
883 // sets the corresponding ordering for the instruction.
884 NVPTX::Ordering InstrOrder;
885 if (N->readMem())
886 InstrOrder = NVPTX::Ordering::Acquire;
887 else if (N->writeMem())
888 InstrOrder = NVPTX::Ordering::Release;
889 else
890 report_fatal_error(
891 reason: formatv(Fmt: "NVPTX does not support SequentiallyConsistent Ordering on "
892 "read-modify-writes yet: {}",
893 Vals: N->getOperationName()));
894 return OperationOrderings(InstrOrder,
895 NVPTX::Ordering::SequentiallyConsistent);
896 }
897 }
898 report_fatal_error(
899 reason: formatv(Fmt: "NVPTX backend does not support AtomicOrdering \"{}\" yet.",
900 Vals: toIRString(ao: Ordering)));
901}
902
903} // namespace
904
905NVPTX::Scope NVPTXDAGToDAGISel::getOperationScope(MemSDNode *N,
906 NVPTX::Ordering O) const {
907 switch (O) {
908 case NVPTX::Ordering::NotAtomic:
909 case NVPTX::Ordering::Volatile: // Non-atomic volatile operations
910 // NVPTX uses Thread scope as the scope of non-atomic operations.
911 return NVPTX::Scope::Thread;
912 case NVPTX::Ordering::RelaxedMMIO:
913 // RelaxedMMIO operations are always system scope.
914 // If a RelaxedMMIO order was generated from an atomic volatile operation
915 // with a smaller thread scope, we bump it here to system scope.
916 return NVPTX::Scope::System;
917 case NVPTX::Ordering::Relaxed:
918 case NVPTX::Ordering::Acquire:
919 case NVPTX::Ordering::Release:
920 case NVPTX::Ordering::AcquireRelease:
921 case NVPTX::Ordering::SequentiallyConsistent:
922 auto S = Scopes[N->getSyncScopeID()];
923
924 S = resolveScope(S, T: Subtarget);
925
926 // If operation is volatile, then its scope is system.
927 return N->isVolatile() ? NVPTX::Scope::System : S;
928 }
929 llvm_unreachable("unhandled ordering");
930}
931
932static bool canLowerToLDG(const MemSDNode &N, const NVPTXSubtarget &Subtarget,
933 NVPTX::AddressSpace CodeAddrSpace) {
934 // We use ldg (i.e. ld.global.nc) for invariant loads from the global address
935 // space.
936 return Subtarget.hasLDG() && CodeAddrSpace == NVPTX::AddressSpace::Global &&
937 N.isInvariant();
938}
939
940static unsigned int getFenceOp(NVPTX::Ordering O, NVPTX::Scope S,
941 NVPTXSubtarget const *T) {
942 S = resolveScope(S, T);
943
944 // Fall back to .acq_rel if .acquire, .release is not supported.
945 if (!T->hasSplitAcquireAndReleaseFences() &&
946 (O == NVPTX::Ordering::Acquire || O == NVPTX::Ordering::Release))
947 O = NVPTX::Ordering::AcquireRelease;
948
949 switch (O) {
950 case NVPTX::Ordering::Acquire:
951 switch (S) {
952 case NVPTX::Scope::System:
953 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_acquire_sys
954 : NVPTX::INT_MEMBAR_SYS;
955 case NVPTX::Scope::Block:
956 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_acquire_cta
957 : NVPTX::INT_MEMBAR_CTA;
958 case NVPTX::Scope::Cluster:
959 return NVPTX::atomic_thread_fence_acquire_cluster;
960 case NVPTX::Scope::Device:
961 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_acquire_gpu
962 : NVPTX::INT_MEMBAR_GL;
963 case NVPTX::Scope::Thread:
964 case NVPTX::Scope::DefaultDevice:
965 report_fatal_error(
966 reason: formatv(Fmt: "Unsupported scope \"{}\" for acquire/release/acq_rel fence.",
967 Vals: ScopeToString(S)));
968 }
969 break;
970 case NVPTX::Ordering::Release:
971 switch (S) {
972 case NVPTX::Scope::System:
973 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_release_sys
974 : NVPTX::INT_MEMBAR_SYS;
975 case NVPTX::Scope::Block:
976 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_release_cta
977 : NVPTX::INT_MEMBAR_CTA;
978 case NVPTX::Scope::Cluster:
979 return NVPTX::atomic_thread_fence_release_cluster;
980 case NVPTX::Scope::Device:
981 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_release_gpu
982 : NVPTX::INT_MEMBAR_GL;
983 case NVPTX::Scope::Thread:
984 case NVPTX::Scope::DefaultDevice:
985 report_fatal_error(
986 reason: formatv(Fmt: "Unsupported scope \"{}\" for acquire/release/acq_rel fence.",
987 Vals: ScopeToString(S)));
988 }
989 break;
990 case NVPTX::Ordering::AcquireRelease: {
991 switch (S) {
992 case NVPTX::Scope::System:
993 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_acq_rel_sys
994 : NVPTX::INT_MEMBAR_SYS;
995 case NVPTX::Scope::Block:
996 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_acq_rel_cta
997 : NVPTX::INT_MEMBAR_CTA;
998 case NVPTX::Scope::Cluster:
999 return NVPTX::atomic_thread_fence_acq_rel_cluster;
1000 case NVPTX::Scope::Device:
1001 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_acq_rel_gpu
1002 : NVPTX::INT_MEMBAR_GL;
1003 case NVPTX::Scope::Thread:
1004 case NVPTX::Scope::DefaultDevice:
1005 report_fatal_error(
1006 reason: formatv(Fmt: "Unsupported scope \"{}\" for acquire/release/acq_rel fence.",
1007 Vals: ScopeToString(S)));
1008 }
1009 break;
1010 }
1011 case NVPTX::Ordering::SequentiallyConsistent: {
1012 switch (S) {
1013 case NVPTX::Scope::System:
1014 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_seq_cst_sys
1015 : NVPTX::INT_MEMBAR_SYS;
1016 case NVPTX::Scope::Block:
1017 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_seq_cst_cta
1018 : NVPTX::INT_MEMBAR_CTA;
1019 case NVPTX::Scope::Cluster:
1020 return NVPTX::atomic_thread_fence_seq_cst_cluster;
1021 case NVPTX::Scope::Device:
1022 return T->hasMemoryOrdering() ? NVPTX::atomic_thread_fence_seq_cst_gpu
1023 : NVPTX::INT_MEMBAR_GL;
1024 case NVPTX::Scope::Thread:
1025 case NVPTX::Scope::DefaultDevice:
1026 report_fatal_error(reason: formatv(Fmt: "Unsupported scope \"{}\" for seq_cst fence.",
1027 Vals: ScopeToString(S)));
1028 }
1029 break;
1030 }
1031 case NVPTX::Ordering::NotAtomic:
1032 case NVPTX::Ordering::Relaxed:
1033 case NVPTX::Ordering::Volatile:
1034 case NVPTX::Ordering::RelaxedMMIO:
1035 report_fatal_error(
1036 reason: formatv(Fmt: "Unsupported \"{}\" ordering and \"{}\" scope for fence.",
1037 Vals: OrderingToString(Order: O), Vals: ScopeToString(S)));
1038 }
1039 llvm_unreachable("unhandled ordering");
1040}
1041
1042// Returns Memory Order and Scope of a memory instruction, and
1043// inserts any fence before the instruction that's required to
1044// implement its memory ordering.
1045std::pair<NVPTX::Ordering, NVPTX::Scope>
1046NVPTXDAGToDAGISel::insertMemoryInstructionFence(SDLoc DL, SDValue &Chain,
1047 MemSDNode *N) {
1048 auto [InstructionOrdering, FenceOrdering] =
1049 getOperationOrderings(N, Subtarget);
1050 auto Scope = getOperationScope(N, O: InstructionOrdering);
1051
1052 // Singlethread scope has no inter-thread synchronization requirements, so
1053 // the atomic operation is lowered as plain and the fence is skipped.
1054 // NotAtomic and Volatile operations naturally have Thread scope and must
1055 // preserve their ordering.
1056 if (Scope == NVPTX::Scope::Thread &&
1057 InstructionOrdering != NVPTX::Ordering::NotAtomic &&
1058 InstructionOrdering != NVPTX::Ordering::Volatile)
1059 return {NVPTX::Ordering::NotAtomic, Scope};
1060
1061 // If a fence is required before the operation, insert it:
1062 switch (NVPTX::Ordering(FenceOrdering)) {
1063 case NVPTX::Ordering::NotAtomic:
1064 break;
1065 case NVPTX::Ordering::SequentiallyConsistent: {
1066 auto Op = getFenceOp(O: FenceOrdering, S: Scope, T: Subtarget);
1067 Chain = SDValue(CurDAG->getMachineNode(Opcode: Op, dl: DL, VT: MVT::Other, Op1: Chain), 0);
1068 break;
1069 }
1070 default:
1071 report_fatal_error(
1072 reason: formatv(Fmt: "Unexpected fence ordering: \"{}\".",
1073 Vals: OrderingToString(Order: NVPTX::Ordering(FenceOrdering))));
1074 }
1075 return {InstructionOrdering, Scope};
1076}
1077
1078// Helper function template to reduce amount of boilerplate code for
1079// opcode selection.
1080static std::optional<unsigned>
1081pickOpcodeForVT(MVT::SimpleValueType VT, std::optional<unsigned> Opcode_i16,
1082 std::optional<unsigned> Opcode_i32,
1083 std::optional<unsigned> Opcode_i64) {
1084 switch (VT) {
1085 case MVT::f16:
1086 case MVT::i16:
1087 case MVT::bf16:
1088 return Opcode_i16;
1089 case MVT::v2f16:
1090 case MVT::v2bf16:
1091 case MVT::v2i16:
1092 case MVT::v4i8:
1093 case MVT::i32:
1094 case MVT::f32:
1095 return Opcode_i32;
1096 case MVT::v2f32:
1097 case MVT::v2i32:
1098 case MVT::i64:
1099 case MVT::f64:
1100 return Opcode_i64;
1101 default:
1102 return std::nullopt;
1103 }
1104}
1105
1106static inline bool isAddLike(const SDValue V) {
1107 return V.getOpcode() == ISD::ADD ||
1108 (V->getOpcode() == ISD::OR && V->getFlags().hasDisjoint());
1109}
1110
1111static SDValue stripAssertAlign(SDValue N) {
1112 if (N.getOpcode() == ISD::AssertAlign)
1113 N = N.getOperand(i: 0);
1114 return N;
1115}
1116
1117// selectBaseADDR - Match a dag node which will serve as the base address for an
1118// ADDR operand pair.
1119static SDValue selectBaseADDR(SDValue N, SelectionDAG *DAG) {
1120 N = stripAssertAlign(N);
1121 if (const auto *GA = dyn_cast<GlobalAddressSDNode>(Val&: N))
1122 return DAG->getTargetGlobalAddress(GV: GA->getGlobal(), DL: SDLoc(N),
1123 VT: GA->getValueType(ResNo: 0), offset: GA->getOffset(),
1124 TargetFlags: GA->getTargetFlags());
1125 if (const auto *ES = dyn_cast<ExternalSymbolSDNode>(Val&: N))
1126 return DAG->getTargetExternalSymbol(Sym: ES->getSymbol(), VT: ES->getValueType(ResNo: 0),
1127 TargetFlags: ES->getTargetFlags());
1128 if (const auto *FIN = dyn_cast<FrameIndexSDNode>(Val&: N))
1129 return DAG->getTargetFrameIndex(FI: FIN->getIndex(), VT: FIN->getValueType(ResNo: 0));
1130 if (N.getOpcode() == NVPTXISD::Symbol)
1131 return N.getOperand(i: 0);
1132
1133 return N;
1134}
1135
1136static SDValue accumulateOffset(SDValue &Addr, SDLoc DL, SelectionDAG *DAG) {
1137 Addr = stripAssertAlign(N: Addr);
1138 APInt AccumulatedOffset(64u, 0);
1139 while (isAddLike(V: Addr)) {
1140 const auto *CN = dyn_cast<ConstantSDNode>(Val: Addr.getOperand(i: 1));
1141 if (!CN)
1142 break;
1143
1144 const APInt CI = CN->getAPIntValue().sext(width: 64);
1145 if (!(CI + AccumulatedOffset).isSignedIntN(N: 32))
1146 break;
1147
1148 AccumulatedOffset += CI;
1149 Addr = stripAssertAlign(N: Addr->getOperand(Num: 0));
1150 }
1151 return DAG->getSignedTargetConstant(Val: AccumulatedOffset.getSExtValue(), DL,
1152 VT: MVT::i32);
1153}
1154
1155static std::pair<SDValue, SDValue> selectADDR(SDValue Addr, SelectionDAG *DAG) {
1156 SDValue Offset = accumulateOffset(Addr, DL: SDLoc(Addr), DAG);
1157 SDValue Base = selectBaseADDR(N: Addr, DAG);
1158 return {Base, Offset};
1159}
1160
1161// Select a pair of operands which represent a valid PTX address, this could be
1162// one of the following things:
1163// - [var] - Offset is simply set to 0
1164// - [reg] - Offset is simply set to 0
1165// - [reg+immOff]
1166// - [var+immOff]
1167// Note that immOff must fit into a 32-bit signed integer.
1168bool NVPTXDAGToDAGISel::SelectADDR(SDValue Addr, SDValue &Base,
1169 SDValue &Offset) {
1170 std::tie(args&: Base, args&: Offset) = selectADDR(Addr, DAG: CurDAG);
1171 return true;
1172}
1173
1174static void emitInvalidMemCacheHint(LLVMContext &Ctx, const Twine &Msg) {
1175 Ctx.diagnose(DI: DiagnosticInfoGeneric(
1176 Twine("invalid NVPTX !mem.cache_hint metadata: ") + Msg, DS_Warning));
1177}
1178
1179static std::optional<NVPTX::L1Eviction> parseL1Eviction(StringRef Str) {
1180 return StringSwitch<std::optional<NVPTX::L1Eviction>>(Str)
1181 .Case(S: "normal", Value: NVPTX::L1Eviction::Normal)
1182 .Case(S: "unchanged", Value: NVPTX::L1Eviction::Unchanged)
1183 .Case(S: "first", Value: NVPTX::L1Eviction::First)
1184 .Case(S: "last", Value: NVPTX::L1Eviction::Last)
1185 .Case(S: "no_allocate", Value: NVPTX::L1Eviction::NoAllocate)
1186 .Default(Value: std::nullopt);
1187}
1188
1189static std::optional<NVPTX::L2Eviction> parseL2Eviction(StringRef Str) {
1190 return StringSwitch<std::optional<NVPTX::L2Eviction>>(Str)
1191 .Case(S: "normal", Value: NVPTX::L2Eviction::Normal)
1192 .Case(S: "first", Value: NVPTX::L2Eviction::First)
1193 .Case(S: "last", Value: NVPTX::L2Eviction::Last)
1194 .Default(Value: std::nullopt);
1195}
1196
1197static std::optional<NVPTX::L2Prefetch> parseL2Prefetch(StringRef Str) {
1198 return StringSwitch<std::optional<NVPTX::L2Prefetch>>(Str)
1199 .Case(S: "64B", Value: NVPTX::L2Prefetch::Bytes64)
1200 .Case(S: "128B", Value: NVPTX::L2Prefetch::Bytes128)
1201 .Case(S: "256B", Value: NVPTX::L2Prefetch::Bytes256)
1202 .Default(Value: std::nullopt);
1203}
1204
1205template <typename T>
1206static std::optional<T> parseMemCacheHintStringValue(
1207 LLVMContext &Ctx, StringRef Key, const Metadata *Value,
1208 std::optional<T> (*Parse)(StringRef), bool EmitDiagnostics) {
1209 const auto *Val = dyn_cast<MDString>(Val: Value);
1210 if (!Val) {
1211 if (EmitDiagnostics)
1212 emitInvalidMemCacheHint(Ctx,
1213 Msg: Twine("'") + Key + "' expects a string value");
1214 return std::nullopt;
1215 }
1216
1217 StringRef ValStr = Val->getString();
1218 auto Parsed = Parse(ValStr);
1219 if (!Parsed && EmitDiagnostics)
1220 emitInvalidMemCacheHint(Ctx, Msg: Twine("unknown value '") + ValStr + "' for '" +
1221 Key + "'");
1222 return Parsed;
1223}
1224
1225static bool isGlobalOrGeneric(NVPTX::AddressSpace AddrSpace) {
1226 return AddrSpace == NVPTX::AddressSpace::Global ||
1227 AddrSpace == NVPTX::AddressSpace::Generic;
1228}
1229
1230static bool isLdOrSt(NVPTXMemCacheHintAccess Access) {
1231 return Access.Instruction == NVPTXMemCacheHintInstruction::Ld ||
1232 Access.Instruction == NVPTXMemCacheHintInstruction::St;
1233}
1234
1235static bool isL1EvictionSupported(const NVPTXSubtarget &Subtarget,
1236 NVPTX::L1Eviction Eviction,
1237 NVPTXMemCacheHintAccess Access) {
1238 if (Eviction == NVPTX::L1Eviction::Normal)
1239 return true;
1240
1241 return isLdOrSt(Access) && !Access.IsVolatile &&
1242 Subtarget.hasL1EvictionHint();
1243}
1244
1245static bool isL2PrefetchSupported(const NVPTXSubtarget &Subtarget,
1246 NVPTX::L2Prefetch Prefetch,
1247 NVPTXMemCacheHintAccess Access) {
1248 switch (Prefetch) {
1249 case NVPTX::L2Prefetch::None:
1250 return true;
1251 case NVPTX::L2Prefetch::Bytes64:
1252 return Access.Instruction == NVPTXMemCacheHintInstruction::Ld &&
1253 isGlobalOrGeneric(AddrSpace: Access.AddrSpace) && Subtarget.hasL2Prefetch64B();
1254 case NVPTX::L2Prefetch::Bytes128:
1255 return Access.Instruction == NVPTXMemCacheHintInstruction::Ld &&
1256 isGlobalOrGeneric(AddrSpace: Access.AddrSpace) && Subtarget.hasL2Prefetch128B();
1257 case NVPTX::L2Prefetch::Bytes256:
1258 return Access.Instruction == NVPTXMemCacheHintInstruction::Ld &&
1259 isGlobalOrGeneric(AddrSpace: Access.AddrSpace) && Subtarget.hasL2Prefetch256B();
1260 }
1261 llvm_unreachable("Unexpected L2 prefetch hint");
1262}
1263
1264static bool isL2EvictionSupported(const NVPTXSubtarget &Subtarget,
1265 NVPTX::L2Eviction Eviction,
1266 NVPTXMemCacheHintAccess Access) {
1267 if (Eviction == NVPTX::L2Eviction::Normal)
1268 return true;
1269
1270 return isLdOrSt(Access) && !Access.IsVolatile &&
1271 Subtarget.hasL2EvictionHint() && isGlobalOrGeneric(AddrSpace: Access.AddrSpace) &&
1272 ((Access.NumElts == 8 && Access.EltWidth == 32) ||
1273 (Access.NumElts == 4 && Access.EltWidth == 64));
1274}
1275
1276static bool isCachePolicySupported(const NVPTXSubtarget &Subtarget,
1277 NVPTXMemCacheHintAccess Access) {
1278 return !Access.IsVolatile && isGlobalOrGeneric(AddrSpace: Access.AddrSpace) &&
1279 Subtarget.hasL2CacheHint();
1280}
1281
1282NVPTXMemCacheHintOperands NVPTXDAGToDAGISel::getMemCacheHintOperands(
1283 const MemSDNode *N, NVPTXMemCacheHintAccess Access, const SDLoc &DL,
1284 bool EmitDiagnostics) {
1285 LLVMContext &Ctx = *CurDAG->getContext();
1286 const MDNode *Node = N->getMemCacheHint();
1287 SDValue PolicyReg = CurDAG->getRegister(Reg: NVPTX::NoRegister, VT: MVT::i64);
1288 if (!Node)
1289 return {.EvictionAndPrefetchHint: getI32Imm(Imm: 0, DL), .CachePolicyReg: PolicyReg};
1290 if (Node->getNumOperands() == 0) {
1291 if (EmitDiagnostics)
1292 emitInvalidMemCacheHint(Ctx, Msg: "empty hint node");
1293 return {.EvictionAndPrefetchHint: getI32Imm(Imm: 0, DL), .CachePolicyReg: PolicyReg};
1294 }
1295
1296 NVPTX::L1Eviction L1 = NVPTX::L1Eviction::Normal;
1297 NVPTX::L2Eviction L2 = NVPTX::L2Eviction::Normal;
1298 NVPTX::L2Prefetch Prefetch = NVPTX::L2Prefetch::None;
1299 std::optional<uint64_t> CachePolicy;
1300
1301 for (unsigned I = 0; I + 1 < Node->getNumOperands(); I += 2) {
1302 const auto *Key = cast<MDString>(Val: Node->getOperand(I));
1303 StringRef KeyStr = Key->getString();
1304 const Metadata *Value = Node->getOperand(I: I + 1).get();
1305
1306 if (KeyStr == "nvvm.l1_eviction") {
1307 auto ParsedL1 = parseMemCacheHintStringValue(
1308 Ctx, Key: KeyStr, Value, Parse: parseL1Eviction, EmitDiagnostics);
1309 if (ParsedL1 && isL1EvictionSupported(Subtarget: *Subtarget, Eviction: *ParsedL1, Access))
1310 L1 = *ParsedL1;
1311 continue;
1312 }
1313
1314 if (KeyStr == "nvvm.l2_eviction") {
1315 auto ParsedL2 = parseMemCacheHintStringValue(
1316 Ctx, Key: KeyStr, Value, Parse: parseL2Eviction, EmitDiagnostics);
1317 if (ParsedL2 && isL2EvictionSupported(Subtarget: *Subtarget, Eviction: *ParsedL2, Access))
1318 L2 = *ParsedL2;
1319 continue;
1320 }
1321
1322 if (KeyStr == "nvvm.l2_prefetch_size") {
1323 auto ParsedPrefetch = parseMemCacheHintStringValue(
1324 Ctx, Key: KeyStr, Value, Parse: parseL2Prefetch, EmitDiagnostics);
1325 if (ParsedPrefetch &&
1326 isL2PrefetchSupported(Subtarget: *Subtarget, Prefetch: *ParsedPrefetch, Access))
1327 Prefetch = *ParsedPrefetch;
1328 continue;
1329 }
1330
1331 if (KeyStr == "nvvm.l2_cache_hint") {
1332 const auto *ValCI = mdconst::dyn_extract<ConstantInt>(MD&: Value);
1333 if (!ValCI) {
1334 if (EmitDiagnostics)
1335 emitInvalidMemCacheHint(
1336 Ctx, Msg: "'nvvm.l2_cache_hint' expects an integer value");
1337 } else if (isCachePolicySupported(Subtarget: *Subtarget, Access)) {
1338 CachePolicy = ValCI->getZExtValue();
1339 }
1340 continue;
1341 }
1342
1343 if (EmitDiagnostics)
1344 emitInvalidMemCacheHint(Ctx, Msg: Twine("unknown key '") + KeyStr + "'");
1345 }
1346
1347 unsigned EvictionAndPrefetchHint =
1348 NVPTX::encodeEvictionAndPrefetchHint(L1, L2, P: Prefetch);
1349 if (CachePolicy) {
1350 SDValue PolicyConst = CurDAG->getTargetConstant(Val: *CachePolicy, DL, VT: MVT::i64);
1351 PolicyReg = SDValue(
1352 CurDAG->getMachineNode(Opcode: NVPTX::MOV_B64_i, dl: DL, VT: MVT::i64, Op1: PolicyConst), 0);
1353 Bitfield::set<NVPTX::L2CacheHintBit>(Packed&: EvictionAndPrefetchHint, Value: true);
1354 }
1355
1356 return {.EvictionAndPrefetchHint: getI32Imm(Imm: EvictionAndPrefetchHint, DL), .CachePolicyReg: PolicyReg};
1357}
1358
1359bool NVPTXDAGToDAGISel::tryLoad(SDNode *N) {
1360 MemSDNode *LD = cast<MemSDNode>(Val: N);
1361 assert(LD->readMem() && "Expected load");
1362
1363 // do not support pre/post inc/dec
1364 const LoadSDNode *PlainLoad = dyn_cast<LoadSDNode>(Val: LD);
1365 if (PlainLoad && PlainLoad->isIndexed())
1366 return false;
1367
1368 // Address Space Setting
1369 const auto CodeAddrSpace = getAddrSpace(N: LD);
1370 if (canLowerToLDG(N: *LD, Subtarget: *Subtarget, CodeAddrSpace))
1371 return tryLDG(N: LD);
1372
1373 SDLoc DL(LD);
1374 SDValue Chain = N->getOperand(Num: 0);
1375 const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, N: LD);
1376
1377 const unsigned FromTypeWidth = LD->getMemoryVT().getSizeInBits();
1378
1379 // Vector Setting
1380 const unsigned FromType =
1381 (PlainLoad && (PlainLoad->getExtensionType() == ISD::SEXTLOAD))
1382 ? NVPTX::PTXLdStInstCode::Signed
1383 : NVPTX::PTXLdStInstCode::Untyped;
1384
1385 uint32_t UsedBytesMask;
1386 switch (N->getOpcode()) {
1387 case ISD::LOAD:
1388 case ISD::ATOMIC_LOAD:
1389 UsedBytesMask = UINT32_MAX;
1390 break;
1391 case NVPTXISD::MLoad:
1392 UsedBytesMask = N->getConstantOperandVal(Num: 3);
1393 break;
1394 default:
1395 llvm_unreachable("Unexpected opcode");
1396 }
1397
1398 assert(isPowerOf2_32(FromTypeWidth) && FromTypeWidth >= 8 &&
1399 FromTypeWidth <= 128 && "Invalid width for load");
1400
1401 const auto [Base, Offset] = selectADDR(Addr: N->getOperand(Num: 1), DAG: CurDAG);
1402 const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
1403 N: LD,
1404 Access: {.Instruction: NVPTXMemCacheHintInstruction::Ld, .AddrSpace: CodeAddrSpace,
1405 /*NumElts=*/1, /*EltWidth=*/FromTypeWidth, .IsVolatile: LD->isVolatile()},
1406 DL);
1407
1408 // Create the machine instruction DAG
1409 SDValue Ops[] = {getI32Imm(Imm: Ordering, DL),
1410 getI32Imm(Imm: Scope, DL),
1411 getI32Imm(Imm: CodeAddrSpace, DL),
1412 getI32Imm(Imm: FromType, DL),
1413 getI32Imm(Imm: FromTypeWidth, DL),
1414 getI32Imm(Imm: UsedBytesMask, DL),
1415 Base,
1416 Offset,
1417 EvictionAndPrefetchHint,
1418 PolicyReg,
1419 Chain};
1420
1421 const MVT::SimpleValueType TargetVT = LD->getSimpleValueType(ResNo: 0).SimpleTy;
1422 const std::optional<unsigned> Opcode =
1423 pickOpcodeForVT(VT: TargetVT, Opcode_i16: NVPTX::LD_i16, Opcode_i32: NVPTX::LD_i32, Opcode_i64: NVPTX::LD_i64);
1424 if (!Opcode)
1425 return false;
1426
1427 SDNode *NVPTXLD = CurDAG->getMachineNode(Opcode: *Opcode, dl: DL, VTs: LD->getVTList(), Ops);
1428 if (!NVPTXLD)
1429 return false;
1430
1431 MachineMemOperand *MemRef = LD->getMemOperand();
1432 CurDAG->setNodeMemRefs(N: cast<MachineSDNode>(Val: NVPTXLD), NewMemRefs: {MemRef});
1433
1434 ReplaceNode(F: LD, T: NVPTXLD);
1435 return true;
1436}
1437
1438static unsigned getStoreVectorNumElts(SDNode *N) {
1439 switch (N->getOpcode()) {
1440 case NVPTXISD::StoreV2:
1441 return 2;
1442 case NVPTXISD::StoreV4:
1443 return 4;
1444 case NVPTXISD::StoreV8:
1445 return 8;
1446 default:
1447 llvm_unreachable("Unexpected opcode");
1448 }
1449}
1450
1451bool NVPTXDAGToDAGISel::tryLoadVector(SDNode *N) {
1452 MemSDNode *LD = cast<MemSDNode>(Val: N);
1453
1454 // Address Space Setting
1455 const auto CodeAddrSpace = getAddrSpace(N: LD);
1456 if (canLowerToLDG(N: *LD, Subtarget: *Subtarget, CodeAddrSpace))
1457 return tryLDG(N: LD);
1458
1459 const MVT EltVT = LD->getSimpleValueType(ResNo: 0);
1460 SDLoc DL(LD);
1461 SDValue Chain = LD->getChain();
1462 const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, N: LD);
1463
1464 // Type Setting: fromType + fromTypeWidth
1465 //
1466 // Sign : ISD::SEXTLOAD
1467 // Unsign : ISD::ZEXTLOAD, ISD::NON_EXTLOAD or ISD::EXTLOAD and the
1468 // type is integer
1469 // Float : ISD::NON_EXTLOAD or ISD::EXTLOAD and the type is float
1470 // Read at least 8 bits (predicates are stored as 8-bit values)
1471 // Get the original LoadSDNode::getExtensionType() value
1472 const unsigned ExtensionType = N->getConstantOperandVal(Num: 4);
1473 const unsigned FromType = (ExtensionType == ISD::SEXTLOAD)
1474 ? NVPTX::PTXLdStInstCode::Signed
1475 : NVPTX::PTXLdStInstCode::Untyped;
1476
1477 const unsigned FromTypeWidth = getFromTypeWidthForLoad(Mem: LD);
1478 const uint32_t UsedBytesMask = N->getConstantOperandVal(Num: 3);
1479
1480 assert(!(EltVT.isVector() && ExtensionType != ISD::NON_EXTLOAD));
1481
1482 const auto [EvictionAndPrefetchHint, PolicyReg] =
1483 getMemCacheHintOperands(N: LD,
1484 Access: {.Instruction: NVPTXMemCacheHintInstruction::Ld, .AddrSpace: CodeAddrSpace,
1485 /*NumElts=*/LD->getNumValues() - 1,
1486 /*EltWidth=*/FromTypeWidth, .IsVolatile: LD->isVolatile()},
1487 DL);
1488 const auto [Base, Offset] = selectADDR(Addr: N->getOperand(Num: 1), DAG: CurDAG);
1489 SDValue Ops[] = {getI32Imm(Imm: Ordering, DL),
1490 getI32Imm(Imm: Scope, DL),
1491 getI32Imm(Imm: CodeAddrSpace, DL),
1492 getI32Imm(Imm: FromType, DL),
1493 getI32Imm(Imm: FromTypeWidth, DL),
1494 getI32Imm(Imm: UsedBytesMask, DL),
1495 Base,
1496 Offset,
1497 EvictionAndPrefetchHint,
1498 PolicyReg,
1499 Chain};
1500
1501 std::optional<unsigned> Opcode;
1502 switch (N->getOpcode()) {
1503 default:
1504 llvm_unreachable("Unexpected opcode");
1505 case NVPTXISD::LoadV2:
1506 Opcode = pickOpcodeForVT(VT: EltVT.SimpleTy, Opcode_i16: NVPTX::LDV_i16_v2,
1507 Opcode_i32: NVPTX::LDV_i32_v2, Opcode_i64: NVPTX::LDV_i64_v2);
1508 break;
1509 case NVPTXISD::LoadV4:
1510 Opcode = pickOpcodeForVT(VT: EltVT.SimpleTy, Opcode_i16: NVPTX::LDV_i16_v4,
1511 Opcode_i32: NVPTX::LDV_i32_v4, Opcode_i64: NVPTX::LDV_i64_v4);
1512 break;
1513 case NVPTXISD::LoadV8:
1514 Opcode = pickOpcodeForVT(VT: EltVT.SimpleTy, Opcode_i16: {/* no v8i16 */},
1515 Opcode_i32: NVPTX::LDV_i32_v8, Opcode_i64: {/* no v8i64 */});
1516 break;
1517 }
1518 if (!Opcode)
1519 return false;
1520
1521 SDNode *NVPTXLD = CurDAG->getMachineNode(Opcode: *Opcode, dl: DL, VTs: LD->getVTList(), Ops);
1522
1523 MachineMemOperand *MemRef = LD->getMemOperand();
1524 CurDAG->setNodeMemRefs(N: cast<MachineSDNode>(Val: NVPTXLD), NewMemRefs: {MemRef});
1525
1526 ReplaceNode(F: LD, T: NVPTXLD);
1527 return true;
1528}
1529
1530bool NVPTXDAGToDAGISel::tryLDG(MemSDNode *LD) {
1531 SDLoc DL(LD);
1532
1533 unsigned ExtensionType;
1534 uint32_t UsedBytesMask;
1535 if (const auto *Load = dyn_cast<LoadSDNode>(Val: LD)) {
1536 ExtensionType = Load->getExtensionType();
1537 UsedBytesMask = UINT32_MAX;
1538 } else {
1539 ExtensionType = LD->getConstantOperandVal(Num: 4);
1540 UsedBytesMask = LD->getConstantOperandVal(Num: 3);
1541 }
1542 const unsigned FromType = (ExtensionType == ISD::SEXTLOAD)
1543 ? NVPTX::PTXLdStInstCode::Signed
1544 : NVPTX::PTXLdStInstCode::Untyped;
1545
1546 const unsigned FromTypeWidth = getFromTypeWidthForLoad(Mem: LD);
1547
1548 assert(!(LD->getSimpleValueType(0).isVector() &&
1549 ExtensionType != ISD::NON_EXTLOAD));
1550
1551 const auto [Base, Offset] = selectADDR(Addr: LD->getOperand(Num: 1), DAG: CurDAG);
1552 const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
1553 N: LD,
1554 Access: {.Instruction: NVPTXMemCacheHintInstruction::Ld, .AddrSpace: NVPTX::AddressSpace::Global,
1555 .NumElts: LD->getNumValues() - 1, .EltWidth: FromTypeWidth, .IsVolatile: LD->isVolatile()},
1556 DL);
1557 SDValue Ops[] = {getI32Imm(Imm: FromType, DL),
1558 getI32Imm(Imm: FromTypeWidth, DL),
1559 getI32Imm(Imm: UsedBytesMask, DL),
1560 Base,
1561 Offset,
1562 EvictionAndPrefetchHint,
1563 PolicyReg,
1564 LD->getChain()};
1565
1566 const MVT::SimpleValueType TargetVT = LD->getSimpleValueType(ResNo: 0).SimpleTy;
1567 std::optional<unsigned> Opcode;
1568 switch (LD->getOpcode()) {
1569 default:
1570 llvm_unreachable("Unexpected opcode");
1571 case ISD::LOAD:
1572 Opcode = pickOpcodeForVT(VT: TargetVT, Opcode_i16: NVPTX::LD_GLOBAL_NC_i16,
1573 Opcode_i32: NVPTX::LD_GLOBAL_NC_i32, Opcode_i64: NVPTX::LD_GLOBAL_NC_i64);
1574 break;
1575 case NVPTXISD::MLoad:
1576 Opcode = pickOpcodeForVT(VT: TargetVT, Opcode_i16: std::nullopt, Opcode_i32: NVPTX::LD_GLOBAL_NC_i32,
1577 Opcode_i64: NVPTX::LD_GLOBAL_NC_i64);
1578 break;
1579 case NVPTXISD::LoadV2:
1580 Opcode =
1581 pickOpcodeForVT(VT: TargetVT, Opcode_i16: NVPTX::LD_GLOBAL_NC_v2i16,
1582 Opcode_i32: NVPTX::LD_GLOBAL_NC_v2i32, Opcode_i64: NVPTX::LD_GLOBAL_NC_v2i64);
1583 break;
1584 case NVPTXISD::LoadV4:
1585 Opcode =
1586 pickOpcodeForVT(VT: TargetVT, Opcode_i16: NVPTX::LD_GLOBAL_NC_v4i16,
1587 Opcode_i32: NVPTX::LD_GLOBAL_NC_v4i32, Opcode_i64: NVPTX::LD_GLOBAL_NC_v4i64);
1588 break;
1589 case NVPTXISD::LoadV8:
1590 Opcode = pickOpcodeForVT(VT: TargetVT, Opcode_i16: {/* no v8i16 */},
1591 Opcode_i32: NVPTX::LD_GLOBAL_NC_v8i32, Opcode_i64: {/* no v8i64 */});
1592 break;
1593 }
1594 if (!Opcode)
1595 return false;
1596
1597 SDNode *NVPTXLDG = CurDAG->getMachineNode(Opcode: *Opcode, dl: DL, VTs: LD->getVTList(), Ops);
1598
1599 ReplaceNode(F: LD, T: NVPTXLDG);
1600 return true;
1601}
1602
1603bool NVPTXDAGToDAGISel::tryLDU(SDNode *N) {
1604 auto *LD = cast<MemSDNode>(Val: N);
1605
1606 SDLoc DL(N);
1607 const unsigned FromTypeWidth = getFromTypeWidthForLoad(Mem: LD);
1608 const MVT::SimpleValueType TargetVT = LD->getSimpleValueType(ResNo: 0).SimpleTy;
1609
1610 // If this is an LDU intrinsic, the address is the third operand. If its an
1611 // LDU SD node (from custom vector handling), then its the second operand
1612 SDValue Addr =
1613 LD->getOperand(Num: LD->getOpcode() == ISD::INTRINSIC_W_CHAIN ? 2 : 1);
1614
1615 const auto [Base, Offset] = selectADDR(Addr, DAG: CurDAG);
1616 SDValue Ops[] = {getI32Imm(Imm: FromTypeWidth, DL), Base, Offset, LD->getChain()};
1617
1618 std::optional<unsigned> Opcode;
1619 switch (N->getOpcode()) {
1620 default:
1621 llvm_unreachable("Unexpected opcode");
1622 case ISD::INTRINSIC_W_CHAIN:
1623 Opcode = pickOpcodeForVT(VT: TargetVT, Opcode_i16: NVPTX::LDU_GLOBAL_i16,
1624 Opcode_i32: NVPTX::LDU_GLOBAL_i32, Opcode_i64: NVPTX::LDU_GLOBAL_i64);
1625 break;
1626 case NVPTXISD::LDUV2:
1627 Opcode = pickOpcodeForVT(VT: TargetVT, Opcode_i16: NVPTX::LDU_GLOBAL_v2i16,
1628 Opcode_i32: NVPTX::LDU_GLOBAL_v2i32, Opcode_i64: NVPTX::LDU_GLOBAL_v2i64);
1629 break;
1630 case NVPTXISD::LDUV4:
1631 Opcode = pickOpcodeForVT(VT: TargetVT, Opcode_i16: NVPTX::LDU_GLOBAL_v4i16,
1632 Opcode_i32: NVPTX::LDU_GLOBAL_v4i32, Opcode_i64: {/* no v4i64 */});
1633 break;
1634 }
1635 if (!Opcode)
1636 return false;
1637
1638 SDNode *NVPTXLDU = CurDAG->getMachineNode(Opcode: *Opcode, dl: DL, VTs: LD->getVTList(), Ops);
1639
1640 ReplaceNode(F: LD, T: NVPTXLDU);
1641 return true;
1642}
1643
1644bool NVPTXDAGToDAGISel::tryStore(SDNode *N) {
1645 MemSDNode *ST = cast<MemSDNode>(Val: N);
1646 assert(ST->writeMem() && "Expected store");
1647 StoreSDNode *PlainStore = dyn_cast<StoreSDNode>(Val: ST);
1648 AtomicSDNode *AtomicStore = dyn_cast<AtomicSDNode>(Val: ST);
1649 assert((PlainStore || AtomicStore) && "Expected store");
1650
1651 // do not support pre/post inc/dec
1652 if (PlainStore && PlainStore->isIndexed())
1653 return false;
1654
1655 // Address Space Setting
1656 const auto CodeAddrSpace = getAddrSpace(N: ST);
1657
1658 SDLoc DL(ST);
1659 SDValue Chain = ST->getChain();
1660 const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, N: ST);
1661
1662 // Vector Setting
1663 const unsigned ToTypeWidth = ST->getMemoryVT().getSizeInBits();
1664
1665 // Create the machine instruction DAG
1666 SDValue Value = PlainStore ? PlainStore->getValue() : AtomicStore->getVal();
1667
1668 assert(isPowerOf2_32(ToTypeWidth) && ToTypeWidth >= 8 && ToTypeWidth <= 128 &&
1669 "Invalid width for store");
1670
1671 const auto [Base, Offset] = selectADDR(Addr: ST->getBasePtr(), DAG: CurDAG);
1672
1673 // Extract eviction/prefetch hint and cache policy register.
1674 const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
1675 N: ST,
1676 Access: {.Instruction: NVPTXMemCacheHintInstruction::St, .AddrSpace: CodeAddrSpace,
1677 /*NumElts=*/1, /*EltWidth=*/ToTypeWidth, .IsVolatile: ST->isVolatile()},
1678 DL);
1679
1680 SDValue Ops[] = {selectPossiblyImm(V: Value),
1681 getI32Imm(Imm: Ordering, DL),
1682 getI32Imm(Imm: Scope, DL),
1683 getI32Imm(Imm: CodeAddrSpace, DL),
1684 getI32Imm(Imm: ToTypeWidth, DL),
1685 Base,
1686 Offset,
1687 EvictionAndPrefetchHint,
1688 PolicyReg,
1689 Chain};
1690
1691 const std::optional<unsigned> Opcode =
1692 pickOpcodeForVT(VT: Value.getSimpleValueType().SimpleTy, Opcode_i16: NVPTX::ST_i16,
1693 Opcode_i32: NVPTX::ST_i32, Opcode_i64: NVPTX::ST_i64);
1694 if (!Opcode)
1695 return false;
1696
1697 SDNode *NVPTXST = CurDAG->getMachineNode(Opcode: *Opcode, dl: DL, VT: MVT::Other, Ops);
1698
1699 if (!NVPTXST)
1700 return false;
1701
1702 MachineMemOperand *MemRef = ST->getMemOperand();
1703 CurDAG->setNodeMemRefs(N: cast<MachineSDNode>(Val: NVPTXST), NewMemRefs: {MemRef});
1704 ReplaceNode(F: ST, T: NVPTXST);
1705 return true;
1706}
1707
1708bool NVPTXDAGToDAGISel::tryStoreVector(SDNode *N) {
1709 MemSDNode *ST = cast<MemSDNode>(Val: N);
1710 const unsigned TotalWidth = ST->getMemoryVT().getSizeInBits();
1711
1712 // Address Space Setting
1713 const auto CodeAddrSpace = getAddrSpace(N: ST);
1714 if (CodeAddrSpace == NVPTX::AddressSpace::Const) {
1715 report_fatal_error(reason: "Cannot store to pointer that points to constant "
1716 "memory space");
1717 }
1718
1719 SDLoc DL(ST);
1720 SDValue Chain = ST->getChain();
1721 const auto [Ordering, Scope] = insertMemoryInstructionFence(DL, Chain, N: ST);
1722
1723 const unsigned NumElts = getStoreVectorNumElts(N: ST);
1724
1725 SmallVector<SDValue, 16> Ops;
1726 for (auto &V : ST->ops().slice(N: 1, M: NumElts))
1727 Ops.push_back(Elt: selectPossiblyImm(V));
1728 SDValue Addr = N->getOperand(Num: NumElts + 1);
1729 const unsigned ToTypeWidth = TotalWidth / NumElts;
1730
1731 assert(isPowerOf2_32(ToTypeWidth) && ToTypeWidth >= 8 && ToTypeWidth <= 128 &&
1732 TotalWidth <= 256 && "Invalid width for store");
1733
1734 // Extract eviction/prefetch hint and cache policy register.
1735 const auto [EvictionAndPrefetchHint, PolicyReg] = getMemCacheHintOperands(
1736 N: ST,
1737 Access: {.Instruction: NVPTXMemCacheHintInstruction::St, .AddrSpace: CodeAddrSpace,
1738 /*NumElts=*/NumElts, /*EltWidth=*/ToTypeWidth, .IsVolatile: ST->isVolatile()},
1739 DL);
1740
1741 const auto [Base, Offset] = selectADDR(Addr, DAG: CurDAG);
1742 Ops.append(IL: {getI32Imm(Imm: Ordering, DL), getI32Imm(Imm: Scope, DL),
1743 getI32Imm(Imm: CodeAddrSpace, DL), getI32Imm(Imm: ToTypeWidth, DL), Base,
1744 Offset, EvictionAndPrefetchHint, PolicyReg, Chain});
1745
1746 const MVT::SimpleValueType EltVT =
1747 ST->getOperand(Num: 1).getSimpleValueType().SimpleTy;
1748 std::optional<unsigned> Opcode;
1749 switch (ST->getOpcode()) {
1750 default:
1751 return false;
1752 case NVPTXISD::StoreV2:
1753 Opcode = pickOpcodeForVT(VT: EltVT, Opcode_i16: NVPTX::STV_i16_v2, Opcode_i32: NVPTX::STV_i32_v2,
1754 Opcode_i64: NVPTX::STV_i64_v2);
1755 break;
1756 case NVPTXISD::StoreV4:
1757 Opcode = pickOpcodeForVT(VT: EltVT, Opcode_i16: NVPTX::STV_i16_v4, Opcode_i32: NVPTX::STV_i32_v4,
1758 Opcode_i64: NVPTX::STV_i64_v4);
1759 break;
1760 case NVPTXISD::StoreV8:
1761 Opcode = pickOpcodeForVT(VT: EltVT, Opcode_i16: {/* no v8i16 */}, Opcode_i32: NVPTX::STV_i32_v8,
1762 Opcode_i64: {/* no v8i64 */});
1763 break;
1764 }
1765
1766 if (!Opcode)
1767 return false;
1768
1769 SDNode *NVPTXST = CurDAG->getMachineNode(Opcode: *Opcode, dl: DL, VT: MVT::Other, Ops);
1770
1771 MachineMemOperand *MemRef = ST->getMemOperand();
1772 CurDAG->setNodeMemRefs(N: cast<MachineSDNode>(Val: NVPTXST), NewMemRefs: {MemRef});
1773
1774 ReplaceNode(F: ST, T: NVPTXST);
1775 return true;
1776}
1777
1778/// SelectBFE - Look for instruction sequences that can be made more efficient
1779/// by using the 'bfe' (bit-field extract) PTX instruction
1780bool NVPTXDAGToDAGISel::tryBFE(SDNode *N) {
1781 SDLoc DL(N);
1782 SDValue LHS = N->getOperand(Num: 0);
1783 SDValue RHS = N->getOperand(Num: 1);
1784 SDValue Len;
1785 SDValue Start;
1786 SDValue Val;
1787 bool IsSigned = false;
1788
1789 if (N->getOpcode() == ISD::AND) {
1790 // Canonicalize the operands
1791 // We want 'and %val, %mask'
1792 if (isa<ConstantSDNode>(Val: LHS) && !isa<ConstantSDNode>(Val: RHS)) {
1793 std::swap(a&: LHS, b&: RHS);
1794 }
1795
1796 ConstantSDNode *Mask = dyn_cast<ConstantSDNode>(Val&: RHS);
1797 if (!Mask) {
1798 // We need a constant mask on the RHS of the AND
1799 return false;
1800 }
1801
1802 // Extract the mask bits
1803 uint64_t MaskVal = Mask->getZExtValue();
1804 if (!isMask_64(Value: MaskVal)) {
1805 // We *could* handle shifted masks here, but doing so would require an
1806 // 'and' operation to fix up the low-order bits so we would trade
1807 // shr+and for bfe+and, which has the same throughput
1808 return false;
1809 }
1810
1811 // How many bits are in our mask?
1812 int64_t NumBits = countr_one(Value: MaskVal);
1813 Len = CurDAG->getTargetConstant(Val: NumBits, DL, VT: MVT::i32);
1814
1815 if (LHS.getOpcode() == ISD::SRL || LHS.getOpcode() == ISD::SRA) {
1816 // We have a 'srl/and' pair, extract the effective start bit and length
1817 Val = LHS.getNode()->getOperand(Num: 0);
1818 Start = LHS.getNode()->getOperand(Num: 1);
1819 ConstantSDNode *StartConst = dyn_cast<ConstantSDNode>(Val&: Start);
1820 if (StartConst) {
1821 uint64_t StartVal = StartConst->getZExtValue();
1822 // How many "good" bits do we have left? "good" is defined here as bits
1823 // that exist in the original value, not shifted in.
1824 int64_t GoodBits = Start.getValueSizeInBits() - StartVal;
1825 if (NumBits > GoodBits) {
1826 // Do not handle the case where bits have been shifted in. In theory
1827 // we could handle this, but the cost is likely higher than just
1828 // emitting the srl/and pair.
1829 return false;
1830 }
1831 Start = CurDAG->getTargetConstant(Val: StartVal, DL, VT: MVT::i32);
1832 } else {
1833 // Do not handle the case where the shift amount (can be zero if no srl
1834 // was found) is not constant. We could handle this case, but it would
1835 // require run-time logic that would be more expensive than just
1836 // emitting the srl/and pair.
1837 return false;
1838 }
1839 } else {
1840 // Do not handle the case where the LHS of the and is not a shift. While
1841 // it would be trivial to handle this case, it would just transform
1842 // 'and' -> 'bfe', but 'and' has higher-throughput.
1843 return false;
1844 }
1845 } else if (N->getOpcode() == ISD::SRL || N->getOpcode() == ISD::SRA) {
1846 if (LHS->getOpcode() == ISD::AND) {
1847 ConstantSDNode *ShiftCnst = dyn_cast<ConstantSDNode>(Val&: RHS);
1848 if (!ShiftCnst) {
1849 // Shift amount must be constant
1850 return false;
1851 }
1852
1853 uint64_t ShiftAmt = ShiftCnst->getZExtValue();
1854
1855 SDValue AndLHS = LHS->getOperand(Num: 0);
1856 SDValue AndRHS = LHS->getOperand(Num: 1);
1857
1858 // Canonicalize the AND to have the mask on the RHS
1859 if (isa<ConstantSDNode>(Val: AndLHS)) {
1860 std::swap(a&: AndLHS, b&: AndRHS);
1861 }
1862
1863 ConstantSDNode *MaskCnst = dyn_cast<ConstantSDNode>(Val&: AndRHS);
1864 if (!MaskCnst) {
1865 // Mask must be constant
1866 return false;
1867 }
1868
1869 uint64_t MaskVal = MaskCnst->getZExtValue();
1870 uint64_t NumZeros;
1871 uint64_t NumBits;
1872 if (isMask_64(Value: MaskVal)) {
1873 NumZeros = 0;
1874 // The number of bits in the result bitfield will be the number of
1875 // trailing ones (the AND) minus the number of bits we shift off
1876 NumBits = llvm::countr_one(Value: MaskVal) - ShiftAmt;
1877 } else if (isShiftedMask_64(Value: MaskVal)) {
1878 NumZeros = llvm::countr_zero(Val: MaskVal);
1879 unsigned NumOnes = llvm::countr_one(Value: MaskVal >> NumZeros);
1880 // The number of bits in the result bitfield will be the number of
1881 // trailing zeros plus the number of set bits in the mask minus the
1882 // number of bits we shift off
1883 NumBits = NumZeros + NumOnes - ShiftAmt;
1884 } else {
1885 // This is not a mask we can handle
1886 return false;
1887 }
1888
1889 if (ShiftAmt < NumZeros) {
1890 // Handling this case would require extra logic that would make this
1891 // transformation non-profitable
1892 return false;
1893 }
1894
1895 Val = AndLHS;
1896 Start = CurDAG->getTargetConstant(Val: ShiftAmt, DL, VT: MVT::i32);
1897 Len = CurDAG->getTargetConstant(Val: NumBits, DL, VT: MVT::i32);
1898
1899 // If pre-shift AND includes the sign bit in the bitfield, we must use
1900 // signed BFE to replicate that bit during bitfield extraction. If the
1901 // sign bit is not part of the mask, unsigned BFE will zero out upper bits
1902 // of the result
1903 if (N->getOpcode() == ISD::SRA)
1904 IsSigned = (ShiftAmt + NumBits) == Val.getValueSizeInBits();
1905 } else if (LHS->getOpcode() == ISD::SHL) {
1906 // Here, we have a pattern like:
1907 //
1908 // (sra (shl val, NN), MM)
1909 // or
1910 // (srl (shl val, NN), MM)
1911 //
1912 // If MM >= NN, we can efficiently optimize this with bfe
1913 Val = LHS->getOperand(Num: 0);
1914
1915 SDValue ShlRHS = LHS->getOperand(Num: 1);
1916 ConstantSDNode *ShlCnst = dyn_cast<ConstantSDNode>(Val&: ShlRHS);
1917 if (!ShlCnst) {
1918 // Shift amount must be constant
1919 return false;
1920 }
1921 uint64_t InnerShiftAmt = ShlCnst->getZExtValue();
1922
1923 SDValue ShrRHS = RHS;
1924 ConstantSDNode *ShrCnst = dyn_cast<ConstantSDNode>(Val&: ShrRHS);
1925 if (!ShrCnst) {
1926 // Shift amount must be constant
1927 return false;
1928 }
1929 uint64_t OuterShiftAmt = ShrCnst->getZExtValue();
1930
1931 // To avoid extra codegen and be profitable, we need Outer >= Inner
1932 if (OuterShiftAmt < InnerShiftAmt) {
1933 return false;
1934 }
1935
1936 // If the outer shift is more than the type size, we have no bitfield to
1937 // extract (since we also check that the inner shift is <= the outer shift
1938 // then this also implies that the inner shift is < the type size)
1939 if (OuterShiftAmt >= Val.getValueSizeInBits()) {
1940 return false;
1941 }
1942
1943 Start = CurDAG->getTargetConstant(Val: OuterShiftAmt - InnerShiftAmt, DL,
1944 VT: MVT::i32);
1945 Len = CurDAG->getTargetConstant(Val: Val.getValueSizeInBits() - OuterShiftAmt,
1946 DL, VT: MVT::i32);
1947
1948 if (N->getOpcode() == ISD::SRA) {
1949 // If we have a arithmetic right shift, we need to use the signed bfe
1950 // variant
1951 IsSigned = true;
1952 }
1953 } else {
1954 // No can do...
1955 return false;
1956 }
1957 } else {
1958 // No can do...
1959 return false;
1960 }
1961
1962
1963 unsigned Opc;
1964 // For the BFE operations we form here from "and" and "srl", always use the
1965 // unsigned variants.
1966 if (Val.getValueType() == MVT::i32) {
1967 if (IsSigned) {
1968 Opc = NVPTX::BFE_S32rii;
1969 } else {
1970 Opc = NVPTX::BFE_U32rii;
1971 }
1972 } else if (Val.getValueType() == MVT::i64) {
1973 if (IsSigned) {
1974 Opc = NVPTX::BFE_S64rii;
1975 } else {
1976 Opc = NVPTX::BFE_U64rii;
1977 }
1978 } else {
1979 // We cannot handle this type
1980 return false;
1981 }
1982
1983 SDValue Ops[] = {
1984 Val, Start, Len
1985 };
1986
1987 ReplaceNode(F: N, T: CurDAG->getMachineNode(Opcode: Opc, dl: DL, VTs: N->getVTList(), Ops));
1988 return true;
1989}
1990
1991// Select bf16/bf16v2 FADD, FSUB, FMUL as fma on targets with only fma
1992bool NVPTXDAGToDAGISel::tryBF16ArithToFMA(SDNode *N) {
1993 EVT VT = SDValue(N, 0).getValueType();
1994 if (VT.getScalarType() != MVT::bf16)
1995 return false;
1996
1997 const NVPTXSubtarget *STI = TM.getSubtargetImpl();
1998 if (STI->hasNativeBF16Support(Opcode: N->getOpcode()))
1999 return false;
2000
2001 const bool IsVec = VT.isVector();
2002 assert(!IsVec || VT.getVectorNumElements() == 2);
2003 SDLoc DL(N);
2004 SDValue N0 = N->getOperand(Num: 0);
2005 SDValue N1 = N->getOperand(Num: 1);
2006 SmallVector<SDValue, 3> Operands;
2007 auto GetConstant = [&](float Value) -> SDValue {
2008 // BF16 immediates must be legalized to integer register values
2009 APFloat APF(Value);
2010 bool LosesInfo;
2011 APF.convert(ToSemantics: APFloat::BFloat(), RM: APFloat::rmNearestTiesToEven, losesInfo: &LosesInfo);
2012 assert(!LosesInfo);
2013 if (IsVec) {
2014 auto API = APF.bitcastToAPInt();
2015 API = API.concat(NewLSB: API);
2016 auto Const = CurDAG->getTargetConstant(Val: API, DL, VT: MVT::i32);
2017 return SDValue(CurDAG->getMachineNode(Opcode: NVPTX::MOV_B32_i, dl: DL, VT, Op1: Const),
2018 0);
2019 }
2020 auto Const = CurDAG->getTargetConstantFP(Val: APF, DL, VT);
2021 return SDValue(CurDAG->getMachineNode(Opcode: NVPTX::MOV_BF16_i, dl: DL, VT, Op1: Const), 0);
2022 };
2023
2024 switch (N->getOpcode()) {
2025 case ISD::FADD:
2026 // add(a, b) -> fma(a, 1.0, b)
2027 Operands = {N0, GetConstant(1.0), N1};
2028 break;
2029 case ISD::FSUB:
2030 // sub(a, b) -> fma(b, -1.0, a)
2031 Operands = {N1, GetConstant(-1.0), N0};
2032 break;
2033 case ISD::FMUL:
2034 // mul(a, b) -> fma(a, b, -0.0)
2035 // NOTE: The identity is -0, not 0, because -0 + 0 == 0 for floats
2036 Operands = {N0, N1, GetConstant(-0.0)};
2037 break;
2038 default:
2039 llvm_unreachable("Unexpected opcode");
2040 };
2041
2042 int Opcode = IsVec ? NVPTX::FMA_BF16x2rrr : NVPTX::FMA_BF16rrr;
2043 MachineSDNode *FMA = CurDAG->getMachineNode(Opcode, dl: DL, VT, Ops: Operands);
2044 ReplaceNode(F: N, T: FMA);
2045 return true;
2046}
2047
2048// The min/max .abs modifier also accepts operands already known to have no
2049// negative values (not even -0). NaN signs are immaterial to these
2050// instructions.
2051bool NVPTXDAGToDAGISel::SelectFAbs(SDValue N, SDValue &Src) {
2052 if (N.getOpcode() == ISD::FABS)
2053 Src = N.getOperand(i: 0);
2054 else if (CurDAG->computeKnownFPClass(Op: N, InterestedClasses: fcNegative).signBitIsZeroOrNaN())
2055 Src = N;
2056 else
2057 return false;
2058 Src = selectPossiblyImm(V: Src);
2059 return true;
2060}
2061
2062SDValue NVPTXDAGToDAGISel::selectPossiblyImm(SDValue V) {
2063 if (V.getOpcode() == ISD::BITCAST)
2064 V = V.getOperand(i: 0);
2065
2066 if (auto *CN = dyn_cast<ConstantSDNode>(Val&: V))
2067 return CurDAG->getTargetConstant(Val: CN->getAPIntValue(), DL: SDLoc(V),
2068 VT: V.getValueType());
2069 if (auto *CN = dyn_cast<ConstantFPSDNode>(Val&: V))
2070 return CurDAG->getTargetConstantFP(Val: CN->getValueAPF(), DL: SDLoc(V),
2071 VT: V.getValueType());
2072 return V;
2073}
2074
2075/// SelectInlineAsmMemoryOperand - Implement addressing mode selection for
2076/// inline asm expressions.
2077bool NVPTXDAGToDAGISel::SelectInlineAsmMemoryOperand(
2078 const SDValue &Op, InlineAsm::ConstraintCode ConstraintID,
2079 std::vector<SDValue> &OutOps) {
2080 switch (ConstraintID) {
2081 default:
2082 return true;
2083 case InlineAsm::ConstraintCode::m: { // memory
2084 const auto [Base, Offset] = selectADDR(Addr: Op, DAG: CurDAG);
2085 OutOps.push_back(x: Base);
2086 OutOps.push_back(x: Offset);
2087 return false;
2088 }
2089 }
2090 return true;
2091}
2092
2093void NVPTXDAGToDAGISel::SelectV2I64toI128(SDNode *N) {
2094 // Lower a CopyToReg with two 64-bit inputs
2095 // Dst:i128, lo:i64, hi:i64
2096 //
2097 // CopyToReg Dst, lo, hi;
2098 //
2099 // ==>
2100 //
2101 // tmp = V2I64toI128 {lo, hi};
2102 // CopyToReg Dst, tmp;
2103 SDValue Dst = N->getOperand(Num: 1);
2104 SDValue Lo = N->getOperand(Num: 2);
2105 SDValue Hi = N->getOperand(Num: 3);
2106
2107 SDLoc DL(N);
2108 SDNode *Mov =
2109 CurDAG->getMachineNode(Opcode: NVPTX::V2I64toI128, dl: DL, VT: MVT::i128, Ops: {Lo, Hi});
2110
2111 SmallVector<SDValue, 4> NewOps(N->getNumOperands() - 1);
2112 NewOps[0] = N->getOperand(Num: 0);
2113 NewOps[1] = Dst;
2114 NewOps[2] = SDValue(Mov, 0);
2115 if (N->getNumOperands() == 5)
2116 NewOps[3] = N->getOperand(Num: 4);
2117 SDValue NewValue = CurDAG->getNode(Opcode: ISD::CopyToReg, DL, ResultTys: SmallVector<EVT>(N->values()), Ops: NewOps);
2118
2119 ReplaceNode(F: N, T: NewValue.getNode());
2120}
2121
2122void NVPTXDAGToDAGISel::SelectI128toV2I64(SDNode *N) {
2123 // Lower CopyFromReg from a 128-bit regs to two 64-bit regs
2124 // Dst:i128, Src:i128
2125 //
2126 // {lo, hi} = CopyFromReg Src
2127 //
2128 // ==>
2129 //
2130 // {lo, hi} = I128toV2I64 Src
2131 //
2132 SDValue Ch = N->getOperand(Num: 0);
2133 SDValue Src = N->getOperand(Num: 1);
2134 SDValue Glue = N->getOperand(Num: 2);
2135 SDLoc DL(N);
2136
2137 // Add Glue and Ch to the operands and results to avoid break the execution
2138 // order
2139 SDNode *Mov = CurDAG->getMachineNode(
2140 Opcode: NVPTX::I128toV2I64, dl: DL,
2141 ResultTys: {MVT::i64, MVT::i64, Ch.getValueType(), Glue.getValueType()},
2142 Ops: {Src, Ch, Glue});
2143
2144 ReplaceNode(F: N, T: Mov);
2145}
2146
2147bool NVPTXDAGToDAGISel::tryFence(SDNode *N) {
2148 SDLoc DL(N);
2149 assert(N->getOpcode() == ISD::ATOMIC_FENCE);
2150 auto Scope = Scopes[N->getConstantOperandVal(Num: 2)];
2151
2152 // Singlethread fences have no inter-thread synchronization requirements.
2153 // Note: std::atomic_signal_fence lowers to singlethread LLVM IR fences;
2154 // this intentionally drops these before emitting PTX.
2155 if (Scope == NVPTX::Scope::Thread) {
2156 CurDAG->ReplaceAllUsesOfValueWith(From: SDValue(N, 0), To: N->getOperand(Num: 0));
2157 CurDAG->RemoveDeadNode(N);
2158 return true;
2159 }
2160
2161 unsigned int FenceOp = getFenceOp(
2162 O: NVPTX::Ordering(N->getConstantOperandVal(Num: 1)), S: Scope, T: Subtarget);
2163 SDValue Chain = N->getOperand(Num: 0);
2164 SDNode *FenceNode = CurDAG->getMachineNode(Opcode: FenceOp, dl: DL, VT: MVT::Other, Op1: Chain);
2165 ReplaceNode(F: N, T: FenceNode);
2166 return true;
2167}
2168
2169NVPTXScopes::NVPTXScopes(LLVMContext &C, const Triple &T) : Context(&C) {
2170 auto ScopeID = [&](AtomicScope Scope) {
2171 return C.getOrInsertSyncScopeID(SSN: *getAtomicScopeIRString(T, S: Scope));
2172 };
2173 Scopes[ScopeID(AtomicScope::Single)] = NVPTX::Scope::Thread;
2174 Scopes[ScopeID(AtomicScope::System)] = NVPTX::Scope::System;
2175 Scopes[ScopeID(AtomicScope::Workgroup)] = NVPTX::Scope::Block;
2176 Scopes[ScopeID(AtomicScope::Cluster)] = NVPTX::Scope::Cluster;
2177 Scopes[ScopeID(AtomicScope::Device)] = NVPTX::Scope::Device;
2178}
2179
2180NVPTX::Scope NVPTXScopes::operator[](SyncScope::ID ID) const {
2181 if (Scopes.empty())
2182 llvm_unreachable("NVPTX Scopes must be initialized before calling "
2183 "NVPTXScopes::operator[]");
2184
2185 auto S = Scopes.find(Key: ID);
2186 if (S == Scopes.end()) {
2187 auto scopeName = Context->getSyncScopeName(Id: ID);
2188 assert(scopeName.has_value() && "Scope name must exist.");
2189
2190 // Build list of supported syncscopes programmatically
2191 SmallVector<StringRef> supportedScopes;
2192 for (const auto &Entry : Scopes) {
2193 if (auto name = Context->getSyncScopeName(Id: Entry.first))
2194 supportedScopes.push_back(Elt: name->empty() ? "<empty string>" : *name);
2195 }
2196
2197 reportFatalUsageError(
2198 reason: formatv(Fmt: "NVPTX backend does not support syncscope \"{0}\" (ID={1}).\n"
2199 "Supported syncscopes are: {2}.",
2200 Vals&: scopeName.value(), Vals: int(ID),
2201 Vals: make_range(x: supportedScopes.begin(), y: supportedScopes.end())));
2202 }
2203 return S->second;
2204}
2205
2206bool NVPTXScopes::empty() const { return Scopes.size() == 0; }
2207
2208#define TCGEN05_ST_OPCODE(SHAPE, NUM) \
2209 (enableUnpack ? NVPTX::TCGEN05_ST_##SHAPE##_##NUM##_UNPACK \
2210 : NVPTX::TCGEN05_ST_##SHAPE##_##NUM)
2211
2212static unsigned getTcgen05StOpcode(unsigned IID, bool enableUnpack) {
2213 switch (IID) {
2214 case Intrinsic::nvvm_tcgen05_st_16x64b_x1:
2215 return TCGEN05_ST_OPCODE(16x64b, x1);
2216 case Intrinsic::nvvm_tcgen05_st_16x64b_x2:
2217 return TCGEN05_ST_OPCODE(16x64b, x2);
2218 case Intrinsic::nvvm_tcgen05_st_16x64b_x4:
2219 return TCGEN05_ST_OPCODE(16x64b, x4);
2220 case Intrinsic::nvvm_tcgen05_st_16x64b_x8:
2221 return TCGEN05_ST_OPCODE(16x64b, x8);
2222 case Intrinsic::nvvm_tcgen05_st_16x64b_x16:
2223 return TCGEN05_ST_OPCODE(16x64b, x16);
2224 case Intrinsic::nvvm_tcgen05_st_16x64b_x32:
2225 return TCGEN05_ST_OPCODE(16x64b, x32);
2226 case Intrinsic::nvvm_tcgen05_st_16x64b_x64:
2227 return TCGEN05_ST_OPCODE(16x64b, x64);
2228 case Intrinsic::nvvm_tcgen05_st_16x64b_x128:
2229 return TCGEN05_ST_OPCODE(16x64b, x128);
2230 case Intrinsic::nvvm_tcgen05_st_16x128b_x1:
2231 return TCGEN05_ST_OPCODE(16x128b, x1);
2232 case Intrinsic::nvvm_tcgen05_st_16x128b_x2:
2233 return TCGEN05_ST_OPCODE(16x128b, x2);
2234 case Intrinsic::nvvm_tcgen05_st_16x128b_x4:
2235 return TCGEN05_ST_OPCODE(16x128b, x4);
2236 case Intrinsic::nvvm_tcgen05_st_16x128b_x8:
2237 return TCGEN05_ST_OPCODE(16x128b, x8);
2238 case Intrinsic::nvvm_tcgen05_st_16x128b_x16:
2239 return TCGEN05_ST_OPCODE(16x128b, x16);
2240 case Intrinsic::nvvm_tcgen05_st_16x128b_x32:
2241 return TCGEN05_ST_OPCODE(16x128b, x32);
2242 case Intrinsic::nvvm_tcgen05_st_16x128b_x64:
2243 return TCGEN05_ST_OPCODE(16x128b, x64);
2244 case Intrinsic::nvvm_tcgen05_st_16x256b_x1:
2245 return TCGEN05_ST_OPCODE(16x256b, x1);
2246 case Intrinsic::nvvm_tcgen05_st_16x256b_x2:
2247 return TCGEN05_ST_OPCODE(16x256b, x2);
2248 case Intrinsic::nvvm_tcgen05_st_16x256b_x4:
2249 return TCGEN05_ST_OPCODE(16x256b, x4);
2250 case Intrinsic::nvvm_tcgen05_st_16x256b_x8:
2251 return TCGEN05_ST_OPCODE(16x256b, x8);
2252 case Intrinsic::nvvm_tcgen05_st_16x256b_x16:
2253 return TCGEN05_ST_OPCODE(16x256b, x16);
2254 case Intrinsic::nvvm_tcgen05_st_16x256b_x32:
2255 return TCGEN05_ST_OPCODE(16x256b, x32);
2256 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1:
2257 return TCGEN05_ST_OPCODE(16x32bx2, x1);
2258 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2:
2259 return TCGEN05_ST_OPCODE(16x32bx2, x2);
2260 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4:
2261 return TCGEN05_ST_OPCODE(16x32bx2, x4);
2262 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8:
2263 return TCGEN05_ST_OPCODE(16x32bx2, x8);
2264 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x16:
2265 return TCGEN05_ST_OPCODE(16x32bx2, x16);
2266 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32:
2267 return TCGEN05_ST_OPCODE(16x32bx2, x32);
2268 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64:
2269 return TCGEN05_ST_OPCODE(16x32bx2, x64);
2270 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128:
2271 return TCGEN05_ST_OPCODE(16x32bx2, x128);
2272 case Intrinsic::nvvm_tcgen05_st_32x32b_x1:
2273 return TCGEN05_ST_OPCODE(32x32b, x1);
2274 case Intrinsic::nvvm_tcgen05_st_32x32b_x2:
2275 return TCGEN05_ST_OPCODE(32x32b, x2);
2276 case Intrinsic::nvvm_tcgen05_st_32x32b_x4:
2277 return TCGEN05_ST_OPCODE(32x32b, x4);
2278 case Intrinsic::nvvm_tcgen05_st_32x32b_x8:
2279 return TCGEN05_ST_OPCODE(32x32b, x8);
2280 case Intrinsic::nvvm_tcgen05_st_32x32b_x16:
2281 return TCGEN05_ST_OPCODE(32x32b, x16);
2282 case Intrinsic::nvvm_tcgen05_st_32x32b_x32:
2283 return TCGEN05_ST_OPCODE(32x32b, x32);
2284 case Intrinsic::nvvm_tcgen05_st_32x32b_x64:
2285 return TCGEN05_ST_OPCODE(32x32b, x64);
2286 case Intrinsic::nvvm_tcgen05_st_32x32b_x128:
2287 return TCGEN05_ST_OPCODE(32x32b, x128);
2288 }
2289 llvm_unreachable("unhandled tcgen05.st lowering");
2290}
2291
2292void NVPTXDAGToDAGISel::SelectTcgen05St(SDNode *N, bool hasOffset) {
2293 if (!Subtarget->hasTcgen05InstSupport())
2294 report_fatal_error(
2295 reason: "tcgen05.st is not supported on this architecture variant");
2296
2297 SDLoc DL(N);
2298 unsigned IID = cast<ConstantSDNode>(Val: N->getOperand(Num: 1))->getZExtValue();
2299
2300 SmallVector<SDValue, 128> Operands = {
2301 N->getOperand(Num: 2) // taddr
2302 };
2303
2304 if (hasOffset)
2305 Operands.push_back(Elt: CurDAG->getTargetConstant(
2306 Val: cast<ConstantSDNode>(Val: N->getOperand(Num: 3))->getZExtValue(), DL,
2307 VT: MVT::i32)); // Offset
2308
2309 for (unsigned I = hasOffset ? 4 : 3; I < (N->getNumOperands() - 1); I++)
2310 Operands.push_back(Elt: N->getOperand(Num: I));
2311
2312 bool enableUnpack =
2313 cast<ConstantSDNode>(Val: N->getOperand(Num: N->getNumOperands() - 1))
2314 ->getZExtValue();
2315
2316 Operands.push_back(Elt: N->getOperand(Num: 0)); // Chain
2317 ReplaceNode(F: N, T: CurDAG->getMachineNode(Opcode: getTcgen05StOpcode(IID, enableUnpack),
2318 dl: DL, VTs: N->getVTList(), Ops: Operands));
2319}
2320
2321bool NVPTXDAGToDAGISel::tryIntrinsicVoid(SDNode *N) {
2322 unsigned IID = N->getConstantOperandVal(Num: 1);
2323 switch (IID) {
2324 default:
2325 return false;
2326 case Intrinsic::nvvm_tcgen05_st_16x64b_x1:
2327 case Intrinsic::nvvm_tcgen05_st_16x64b_x2:
2328 case Intrinsic::nvvm_tcgen05_st_16x64b_x4:
2329 case Intrinsic::nvvm_tcgen05_st_16x64b_x8:
2330 case Intrinsic::nvvm_tcgen05_st_16x64b_x16:
2331 case Intrinsic::nvvm_tcgen05_st_16x64b_x32:
2332 case Intrinsic::nvvm_tcgen05_st_16x64b_x64:
2333 case Intrinsic::nvvm_tcgen05_st_16x64b_x128:
2334 case Intrinsic::nvvm_tcgen05_st_32x32b_x1:
2335 case Intrinsic::nvvm_tcgen05_st_32x32b_x2:
2336 case Intrinsic::nvvm_tcgen05_st_32x32b_x4:
2337 case Intrinsic::nvvm_tcgen05_st_32x32b_x8:
2338 case Intrinsic::nvvm_tcgen05_st_32x32b_x16:
2339 case Intrinsic::nvvm_tcgen05_st_32x32b_x32:
2340 case Intrinsic::nvvm_tcgen05_st_32x32b_x64:
2341 case Intrinsic::nvvm_tcgen05_st_32x32b_x128:
2342 case Intrinsic::nvvm_tcgen05_st_16x128b_x1:
2343 case Intrinsic::nvvm_tcgen05_st_16x128b_x2:
2344 case Intrinsic::nvvm_tcgen05_st_16x128b_x4:
2345 case Intrinsic::nvvm_tcgen05_st_16x128b_x8:
2346 case Intrinsic::nvvm_tcgen05_st_16x128b_x16:
2347 case Intrinsic::nvvm_tcgen05_st_16x128b_x32:
2348 case Intrinsic::nvvm_tcgen05_st_16x128b_x64:
2349 case Intrinsic::nvvm_tcgen05_st_16x256b_x1:
2350 case Intrinsic::nvvm_tcgen05_st_16x256b_x2:
2351 case Intrinsic::nvvm_tcgen05_st_16x256b_x4:
2352 case Intrinsic::nvvm_tcgen05_st_16x256b_x8:
2353 case Intrinsic::nvvm_tcgen05_st_16x256b_x16:
2354 case Intrinsic::nvvm_tcgen05_st_16x256b_x32: {
2355 SelectTcgen05St(N);
2356 return true;
2357 }
2358
2359 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x1:
2360 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x2:
2361 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x4:
2362 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x8:
2363 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x16:
2364 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x32:
2365 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x64:
2366 case Intrinsic::nvvm_tcgen05_st_16x32bx2_x128: {
2367 SelectTcgen05St(N, /* hasOffset */ true);
2368 return true;
2369 }
2370 }
2371}
2372
2373void NVPTXDAGToDAGISel::selectAtomicSwap128(SDNode *N) {
2374 MemSDNode *AN = cast<MemSDNode>(Val: N);
2375 SDLoc dl(N);
2376
2377 const SDValue Chain = N->getOperand(Num: 0);
2378 const auto [Base, Offset] = selectADDR(Addr: N->getOperand(Num: 1), DAG: CurDAG);
2379 SmallVector<SDValue, 10> Ops{Base, Offset};
2380 Ops.append(in_start: N->op_begin() + 2, in_end: N->op_end());
2381 Ops.append(IL: {getI32Imm(Imm: getMemOrder(N: AN), DL: dl), getI32Imm(Imm: getAtomicScope(N: AN), DL: dl),
2382 getI32Imm(Imm: getAddrSpace(N: AN), DL: dl)});
2383
2384 if (N->getOpcode() == NVPTXISD::ATOMIC_SWAP_B128) {
2385 unsigned EltWidth = AN->getMemoryVT().getFixedSizeInBits();
2386 NVPTXMemCacheHintAccess Access{.Instruction: NVPTXMemCacheHintInstruction::Atom,
2387 .AddrSpace: getAddrSpace(N: AN),
2388 /*NumElts=*/1, .EltWidth: EltWidth, .IsVolatile: AN->isVolatile()};
2389 const auto [EvictionAndPrefetchHint, CachePolicyReg] =
2390 getMemCacheHintOperands(N: AN, Access, DL: dl);
2391 Ops.push_back(Elt: EvictionAndPrefetchHint);
2392 Ops.push_back(Elt: CachePolicyReg);
2393 }
2394
2395 Ops.push_back(Elt: Chain);
2396
2397 assert(N->getOpcode() == NVPTXISD::ATOMIC_CMP_SWAP_B128 ||
2398 N->getOpcode() == NVPTXISD::ATOMIC_SWAP_B128);
2399 unsigned Opcode = N->getOpcode() == NVPTXISD::ATOMIC_SWAP_B128
2400 ? NVPTX::ATOM_EXCH_B128
2401 : NVPTX::ATOM_CAS_B128;
2402
2403 auto *ATOM = CurDAG->getMachineNode(Opcode, dl, VTs: N->getVTList(), Ops);
2404 CurDAG->setNodeMemRefs(N: ATOM, NewMemRefs: AN->getMemOperand());
2405
2406 ReplaceNode(F: N, T: ATOM);
2407}
2408