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