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