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