1//===-- NVPTXTargetTransformInfo.cpp - NVPTX specific TTI -----------------===//
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#include "NVPTXTargetTransformInfo.h"
10#include "NVVMProperties.h"
11#include "llvm/ADT/ArrayRef.h"
12#include "llvm/ADT/DenseMap.h"
13#include "llvm/ADT/STLExtras.h"
14#include "llvm/ADT/SmallVector.h"
15#include "llvm/Analysis/LoopInfo.h"
16#include "llvm/Analysis/TargetTransformInfo.h"
17#include "llvm/Analysis/ValueTracking.h"
18#include "llvm/CodeGen/BasicTTIImpl.h"
19#include "llvm/CodeGen/TargetLowering.h"
20#include "llvm/IR/Constants.h"
21#include "llvm/IR/IntrinsicInst.h"
22#include "llvm/IR/Intrinsics.h"
23#include "llvm/IR/IntrinsicsNVPTX.h"
24#include "llvm/IR/Value.h"
25#include "llvm/Support/Casting.h"
26#include "llvm/Support/ErrorHandling.h"
27#include "llvm/Support/NVPTXAddrSpace.h"
28#include "llvm/Transforms/InstCombine/InstCombiner.h"
29#include <optional>
30#include <utility>
31using namespace llvm;
32
33#define DEBUG_TYPE "NVPTXtti"
34
35// Whether the given intrinsic reads threadIdx.x/y/z.
36static bool readsThreadIndex(const IntrinsicInst *II) {
37 switch (II->getIntrinsicID()) {
38 default: return false;
39 case Intrinsic::nvvm_read_ptx_sreg_tid_x:
40 case Intrinsic::nvvm_read_ptx_sreg_tid_y:
41 case Intrinsic::nvvm_read_ptx_sreg_tid_z:
42 return true;
43 }
44}
45
46static bool readsLaneId(const IntrinsicInst *II) {
47 return II->getIntrinsicID() == Intrinsic::nvvm_read_ptx_sreg_laneid;
48}
49
50bool NVPTXTTIImpl::isSourceOfDivergence(const Value *V) const {
51 // Without inter-procedural analysis, we conservatively assume that arguments
52 // to __device__ functions are divergent.
53 if (const Argument *Arg = dyn_cast<Argument>(Val: V))
54 return !isKernelFunction(F: *Arg->getParent());
55
56 if (const Instruction *I = dyn_cast<Instruction>(Val: V)) {
57 // Without pointer analysis, we conservatively assume values loaded from
58 // generic or local address space are divergent.
59 if (const LoadInst *LI = dyn_cast<LoadInst>(Val: I)) {
60 unsigned AS = LI->getPointerAddressSpace();
61 return AS == ADDRESS_SPACE_GENERIC || AS == ADDRESS_SPACE_LOCAL;
62 }
63 // Atomic instructions may cause divergence. Atomic instructions are
64 // executed sequentially across all threads in a warp. Therefore, an earlier
65 // executed thread may see different memory inputs than a later executed
66 // thread. For example, suppose *a = 0 initially.
67 //
68 // atom.global.add.s32 d, [a], 1
69 //
70 // returns 0 for the first thread that enters the critical region, and 1 for
71 // the second thread.
72 if (I->isAtomic())
73 return true;
74 if (const IntrinsicInst *II = dyn_cast<IntrinsicInst>(Val: I)) {
75 // Instructions that read threadIdx are obviously divergent.
76 if (readsThreadIndex(II) || readsLaneId(II))
77 return true;
78 }
79 // Conservatively consider the return value of function calls as divergent.
80 // We could analyze callees with bodies more precisely using
81 // inter-procedural analysis.
82 if (isa<CallInst>(Val: I))
83 return true;
84 }
85
86 return false;
87}
88
89// Convert NVVM intrinsics to target-generic LLVM code where possible.
90static Instruction *convertNvvmIntrinsicToLlvm(InstCombiner &IC,
91 IntrinsicInst *II) {
92 // Each NVVM intrinsic we can simplify can be replaced with one of:
93 //
94 // * an LLVM intrinsic,
95 // * an LLVM cast operation,
96 // * an LLVM binary operation, or
97 // * ad-hoc LLVM IR for the particular operation.
98
99 // Some transformations are only valid when the module's
100 // flush-denormals-to-zero (ftz) setting is true/false, whereas other
101 // transformations are valid regardless of the module's ftz setting.
102 enum FtzRequirementTy {
103 FTZ_Any, // Any ftz setting is ok.
104 FTZ_MustBeOn, // Transformation is valid only if ftz is on.
105 FTZ_MustBeOff, // Transformation is valid only if ftz is off.
106 };
107 // Classes of NVVM intrinsics that can't be replaced one-to-one with a
108 // target-generic intrinsic, cast op, or binary op but that we can nonetheless
109 // simplify.
110 enum SpecialCase {
111 SPC_Reciprocal,
112 SCP_FunnelShiftClamp,
113 };
114
115 // SimplifyAction is a poor-man's variant (plus an additional flag) that
116 // represents how to replace an NVVM intrinsic with target-generic LLVM IR.
117 struct SimplifyAction {
118 // Invariant: At most one of these Optionals has a value.
119 std::optional<Intrinsic::ID> IID;
120 std::optional<Instruction::CastOps> CastOp;
121 std::optional<Instruction::BinaryOps> BinaryOp;
122 std::optional<SpecialCase> Special;
123
124 FtzRequirementTy FtzRequirement = FTZ_Any;
125 // Denormal handling is guarded by different attributes depending on the
126 // type (denormal-fp-math vs denormal-fp-math-f32), take note of halfs.
127 bool IsHalfTy = false;
128
129 SimplifyAction() = default;
130
131 SimplifyAction(Intrinsic::ID IID, FtzRequirementTy FtzReq,
132 bool IsHalfTy = false)
133 : IID(IID), FtzRequirement(FtzReq), IsHalfTy(IsHalfTy) {}
134
135 // Cast operations don't have anything to do with FTZ, so we skip that
136 // argument.
137 SimplifyAction(Instruction::CastOps CastOp) : CastOp(CastOp) {}
138
139 SimplifyAction(Instruction::BinaryOps BinaryOp, FtzRequirementTy FtzReq)
140 : BinaryOp(BinaryOp), FtzRequirement(FtzReq) {}
141
142 SimplifyAction(SpecialCase Special, FtzRequirementTy FtzReq)
143 : Special(Special), FtzRequirement(FtzReq) {}
144 };
145
146 // Try to generate a SimplifyAction describing how to replace our
147 // IntrinsicInstr with target-generic LLVM IR.
148 const SimplifyAction Action = [II]() -> SimplifyAction {
149 switch (II->getIntrinsicID()) {
150 // NVVM intrinsics that map directly to LLVM intrinsics.
151 case Intrinsic::nvvm_ceil_d:
152 return {Intrinsic::ceil, FTZ_Any};
153 case Intrinsic::nvvm_ceil_f:
154 return {Intrinsic::ceil, FTZ_MustBeOff};
155 case Intrinsic::nvvm_ceil_ftz_f:
156 return {Intrinsic::ceil, FTZ_MustBeOn};
157 case Intrinsic::nvvm_floor_d:
158 return {Intrinsic::floor, FTZ_Any};
159 case Intrinsic::nvvm_floor_f:
160 return {Intrinsic::floor, FTZ_MustBeOff};
161 case Intrinsic::nvvm_floor_ftz_f:
162 return {Intrinsic::floor, FTZ_MustBeOn};
163 case Intrinsic::nvvm_fma_rn_d:
164 return {Intrinsic::fma, FTZ_Any};
165 case Intrinsic::nvvm_fma_rn_f:
166 return {Intrinsic::fma, FTZ_MustBeOff};
167 case Intrinsic::nvvm_fma_rn_ftz_f:
168 return {Intrinsic::fma, FTZ_MustBeOn};
169 case Intrinsic::nvvm_fma_rn_f16:
170 return {Intrinsic::fma, FTZ_MustBeOff, true};
171 case Intrinsic::nvvm_fma_rn_ftz_f16:
172 return {Intrinsic::fma, FTZ_MustBeOn, true};
173 case Intrinsic::nvvm_fma_rn_f16x2:
174 return {Intrinsic::fma, FTZ_MustBeOff, true};
175 case Intrinsic::nvvm_fma_rn_ftz_f16x2:
176 return {Intrinsic::fma, FTZ_MustBeOn, true};
177 case Intrinsic::nvvm_fma_rn_bf16:
178 return {Intrinsic::fma, FTZ_MustBeOff, true};
179 case Intrinsic::nvvm_fma_rn_bf16x2:
180 return {Intrinsic::fma, FTZ_MustBeOff, true};
181 case Intrinsic::nvvm_fmax_d:
182 return {Intrinsic::maximumnum, FTZ_Any};
183 case Intrinsic::nvvm_fmax_f:
184 return {Intrinsic::maximumnum, FTZ_MustBeOff};
185 case Intrinsic::nvvm_fmax_ftz_f:
186 return {Intrinsic::maximumnum, FTZ_MustBeOn};
187 case Intrinsic::nvvm_fmax_nan_f:
188 return {Intrinsic::maximum, FTZ_MustBeOff};
189 case Intrinsic::nvvm_fmax_ftz_nan_f:
190 return {Intrinsic::maximum, FTZ_MustBeOn};
191 case Intrinsic::nvvm_fmax_f16:
192 return {Intrinsic::maximumnum, FTZ_MustBeOff, true};
193 case Intrinsic::nvvm_fmax_ftz_f16:
194 return {Intrinsic::maximumnum, FTZ_MustBeOn, true};
195 case Intrinsic::nvvm_fmax_f16x2:
196 return {Intrinsic::maximumnum, FTZ_MustBeOff, true};
197 case Intrinsic::nvvm_fmax_ftz_f16x2:
198 return {Intrinsic::maximumnum, FTZ_MustBeOn, true};
199 case Intrinsic::nvvm_fmax_nan_f16:
200 return {Intrinsic::maximum, FTZ_MustBeOff, true};
201 case Intrinsic::nvvm_fmax_ftz_nan_f16:
202 return {Intrinsic::maximum, FTZ_MustBeOn, true};
203 case Intrinsic::nvvm_fmax_nan_f16x2:
204 return {Intrinsic::maximum, FTZ_MustBeOff, true};
205 case Intrinsic::nvvm_fmax_ftz_nan_f16x2:
206 return {Intrinsic::maximum, FTZ_MustBeOn, true};
207 case Intrinsic::nvvm_fmin_d:
208 return {Intrinsic::minimumnum, FTZ_Any};
209 case Intrinsic::nvvm_fmin_f:
210 return {Intrinsic::minimumnum, FTZ_MustBeOff};
211 case Intrinsic::nvvm_fmin_ftz_f:
212 return {Intrinsic::minimumnum, FTZ_MustBeOn};
213 case Intrinsic::nvvm_fmin_nan_f:
214 return {Intrinsic::minimum, FTZ_MustBeOff};
215 case Intrinsic::nvvm_fmin_ftz_nan_f:
216 return {Intrinsic::minimum, FTZ_MustBeOn};
217 case Intrinsic::nvvm_fmin_f16:
218 return {Intrinsic::minimumnum, FTZ_MustBeOff, true};
219 case Intrinsic::nvvm_fmin_ftz_f16:
220 return {Intrinsic::minimumnum, FTZ_MustBeOn, true};
221 case Intrinsic::nvvm_fmin_f16x2:
222 return {Intrinsic::minimumnum, FTZ_MustBeOff, true};
223 case Intrinsic::nvvm_fmin_ftz_f16x2:
224 return {Intrinsic::minimumnum, FTZ_MustBeOn, true};
225 case Intrinsic::nvvm_fmin_nan_f16:
226 return {Intrinsic::minimum, FTZ_MustBeOff, true};
227 case Intrinsic::nvvm_fmin_ftz_nan_f16:
228 return {Intrinsic::minimum, FTZ_MustBeOn, true};
229 case Intrinsic::nvvm_fmin_nan_f16x2:
230 return {Intrinsic::minimum, FTZ_MustBeOff, true};
231 case Intrinsic::nvvm_fmin_ftz_nan_f16x2:
232 return {Intrinsic::minimum, FTZ_MustBeOn, true};
233 case Intrinsic::nvvm_sqrt_rn_d:
234 return {Intrinsic::sqrt, FTZ_Any};
235 case Intrinsic::nvvm_sqrt_f:
236 // nvvm_sqrt_f is a special case. For most intrinsics, foo_ftz_f is the
237 // ftz version, and foo_f is the non-ftz version. But nvvm_sqrt_f adopts
238 // the ftz-ness of the surrounding code. sqrt_rn_f and sqrt_rn_ftz_f are
239 // the versions with explicit ftz-ness.
240 return {Intrinsic::sqrt, FTZ_Any};
241 case Intrinsic::nvvm_trunc_d:
242 return {Intrinsic::trunc, FTZ_Any};
243 case Intrinsic::nvvm_trunc_f:
244 return {Intrinsic::trunc, FTZ_MustBeOff};
245 case Intrinsic::nvvm_trunc_ftz_f:
246 return {Intrinsic::trunc, FTZ_MustBeOn};
247
248 // NVVM intrinsics that map to LLVM cast operations.
249 // Note - we cannot map intrinsics like nvvm_d2ll_rz to LLVM's
250 // FPToSI, as NaN to int conversion with FPToSI is considered UB and is
251 // eliminated. NVVM conversion intrinsics are translated to PTX cvt
252 // instructions which define the outcome for NaN rather than leaving as UB.
253 // Therefore, translate NVVM intrinsics to sitofp/uitofp, but not to
254 // fptosi/fptoui.
255 case Intrinsic::nvvm_i2d_rn:
256 case Intrinsic::nvvm_i2f_rn:
257 case Intrinsic::nvvm_ll2d_rn:
258 case Intrinsic::nvvm_ll2f_rn:
259 return {Instruction::SIToFP};
260 case Intrinsic::nvvm_ui2d_rn:
261 case Intrinsic::nvvm_ui2f_rn:
262 case Intrinsic::nvvm_ull2d_rn:
263 case Intrinsic::nvvm_ull2f_rn:
264 return {Instruction::UIToFP};
265
266 // NVVM intrinsics that map to LLVM binary ops.
267 case Intrinsic::nvvm_div_rn_d:
268 return {Instruction::FDiv, FTZ_Any};
269
270 // The remainder of cases are NVVM intrinsics that map to LLVM idioms, but
271 // need special handling.
272 //
273 // We seem to be missing intrinsics for rcp.approx.{ftz.}f32, which is just
274 // as well.
275 case Intrinsic::nvvm_rcp_rn_d:
276 return {SPC_Reciprocal, FTZ_Any};
277
278 case Intrinsic::nvvm_fshl_clamp:
279 case Intrinsic::nvvm_fshr_clamp:
280 return {SCP_FunnelShiftClamp, FTZ_Any};
281
282 // We do not currently simplify intrinsics that give an approximate
283 // answer. These include:
284 //
285 // - nvvm_cos_approx_{f,ftz_f}
286 // - nvvm_ex2_approx(_ftz)
287 // - nvvm_lg2_approx_{d,f,ftz_f}
288 // - nvvm_sin_approx_{f,ftz_f}
289 // - nvvm_sqrt_approx_{f,ftz_f}
290 // - nvvm_rsqrt_approx_{d,f,ftz_f}
291 // - nvvm_div_approx_{ftz_d,ftz_f,f}
292 // - nvvm_rcp_approx_ftz_d
293 //
294 // Ideally we'd encode them as e.g. "fast call @llvm.cos", where "fast"
295 // means that fastmath is enabled in the intrinsic. Unfortunately only
296 // binary operators (currently) have a fastmath bit in SelectionDAG, so
297 // this information gets lost and we can't select on it.
298 //
299 // TODO: div and rcp are lowered to a binary op, so these we could in
300 // theory lower them to "fast fdiv".
301
302 default:
303 return {};
304 }
305 }();
306
307 // If Action.FtzRequirementTy is not satisfied by the module's ftz state, we
308 // can bail out now. (Notice that in the case that IID is not an NVVM
309 // intrinsic, we don't have to look up any module metadata, as
310 // FtzRequirementTy will be FTZ_Any.)
311 if (Action.FtzRequirement != FTZ_Any) {
312 // FIXME: Broken for f64
313 DenormalMode Mode = II->getFunction()->getDenormalMode(
314 FPType: Action.IsHalfTy ? APFloat::IEEEhalf() : APFloat::IEEEsingle());
315 bool FtzEnabled = Mode.Output == DenormalMode::PreserveSign;
316
317 if (FtzEnabled != (Action.FtzRequirement == FTZ_MustBeOn))
318 return nullptr;
319 }
320
321 // Simplify to target-generic intrinsic.
322 if (Action.IID) {
323 SmallVector<Value *, 4> Args(II->args());
324 // All the target-generic intrinsics currently of interest to us have one
325 // type argument, equal to that of the nvvm intrinsic's argument.
326 Type *Tys[] = {II->getArgOperand(i: 0)->getType()};
327 return CallInst::Create(
328 Func: Intrinsic::getOrInsertDeclaration(M: II->getModule(), id: *Action.IID, OverloadTys: Tys),
329 Args);
330 }
331
332 // Simplify to target-generic binary op.
333 if (Action.BinaryOp)
334 return BinaryOperator::Create(Op: *Action.BinaryOp, S1: II->getArgOperand(i: 0),
335 S2: II->getArgOperand(i: 1), Name: II->getName());
336
337 // Simplify to target-generic cast op.
338 if (Action.CastOp)
339 return CastInst::Create(*Action.CastOp, S: II->getArgOperand(i: 0), Ty: II->getType(),
340 Name: II->getName());
341
342 // All that's left are the special cases.
343 if (!Action.Special)
344 return nullptr;
345
346 switch (*Action.Special) {
347 case SPC_Reciprocal:
348 // Simplify reciprocal.
349 return BinaryOperator::Create(
350 Op: Instruction::FDiv, S1: ConstantFP::get(Ty: II->getArgOperand(i: 0)->getType(), V: 1),
351 S2: II->getArgOperand(i: 0), Name: II->getName());
352
353 case SCP_FunnelShiftClamp: {
354 // Canonicalize a clamping funnel shift to the generic llvm funnel shift
355 // when possible, as this is easier for llvm to optimize further.
356 if (const auto *ShiftConst = dyn_cast<ConstantInt>(Val: II->getArgOperand(i: 2))) {
357 const bool IsLeft = II->getIntrinsicID() == Intrinsic::nvvm_fshl_clamp;
358 if (ShiftConst->getZExtValue() >= II->getType()->getIntegerBitWidth())
359 return IC.replaceInstUsesWith(I&: *II, V: II->getArgOperand(i: IsLeft ? 1 : 0));
360
361 const unsigned FshIID = IsLeft ? Intrinsic::fshl : Intrinsic::fshr;
362 return CallInst::Create(Func: Intrinsic::getOrInsertDeclaration(
363 M: II->getModule(), id: FshIID, OverloadTys: II->getType()),
364 Args: SmallVector<Value *, 3>(II->args()));
365 }
366 return nullptr;
367 }
368 }
369 llvm_unreachable("All SpecialCase enumerators should be handled in switch.");
370}
371
372// Returns whether a pointer in AS is in the given specific address space, or
373// nullopt when this cannot be determined at compile time.
374static std::optional<bool> isInAddressSpace(unsigned AS, unsigned SpecificAS) {
375 if (AS == NVPTXAS::ADDRESS_SPACE_GENERIC ||
376 AS == NVPTXAS::ADDRESS_SPACE_ENTRY_PARAM)
377 return std::nullopt; // Got to check at run-time.
378 if (AS == NVPTXAS::ADDRESS_SPACE_SHARED_CLUSTER &&
379 SpecificAS == NVPTXAS::ADDRESS_SPACE_SHARED)
380 return std::nullopt;
381 return AS == SpecificAS ||
382 (AS == NVPTXAS::ADDRESS_SPACE_SHARED &&
383 SpecificAS == NVPTXAS::ADDRESS_SPACE_SHARED_CLUSTER);
384}
385
386// Returns true/false when we know the answer, nullopt otherwise.
387static std::optional<bool> evaluateIsSpace(Intrinsic::ID IID, unsigned AS) {
388 switch (IID) {
389 case Intrinsic::nvvm_isspacep_global:
390 return isInAddressSpace(AS, SpecificAS: NVPTXAS::ADDRESS_SPACE_GLOBAL);
391 case Intrinsic::nvvm_isspacep_local:
392 return isInAddressSpace(AS, SpecificAS: NVPTXAS::ADDRESS_SPACE_LOCAL);
393 case Intrinsic::nvvm_isspacep_shared:
394 return isInAddressSpace(AS, SpecificAS: NVPTXAS::ADDRESS_SPACE_SHARED);
395 case Intrinsic::nvvm_isspacep_shared_cluster:
396 return isInAddressSpace(AS, SpecificAS: NVPTXAS::ADDRESS_SPACE_SHARED_CLUSTER);
397 case Intrinsic::nvvm_isspacep_const:
398 return isInAddressSpace(AS, SpecificAS: NVPTXAS::ADDRESS_SPACE_CONST);
399 default:
400 llvm_unreachable("Unexpected intrinsic");
401 }
402}
403
404// Returns an instruction pointer (may be nullptr if we do not know the answer).
405// Returns nullopt if `II` is not one of the `isspacep` intrinsics.
406//
407// TODO: If InferAddressSpaces were run early enough in the pipeline this could
408// be removed in favor of the constant folding that occurs there through
409// rewriteIntrinsicWithAddressSpace
410static std::optional<Instruction *>
411handleSpaceCheckIntrinsics(InstCombiner &IC, IntrinsicInst &II) {
412
413 switch (auto IID = II.getIntrinsicID()) {
414 case Intrinsic::nvvm_isspacep_global:
415 case Intrinsic::nvvm_isspacep_local:
416 case Intrinsic::nvvm_isspacep_shared:
417 case Intrinsic::nvvm_isspacep_shared_cluster:
418 case Intrinsic::nvvm_isspacep_const: {
419 Value *Op0 = II.getArgOperand(i: 0);
420 unsigned AS = Op0->getType()->getPointerAddressSpace();
421 // Peek through ASC to generic AS.
422 // TODO: we could dig deeper through both ASCs and GEPs.
423 if (AS == NVPTXAS::ADDRESS_SPACE_GENERIC)
424 if (auto *ASCO = dyn_cast<AddrSpaceCastOperator>(Val: Op0))
425 AS = ASCO->getOperand(i_nocapture: 0)->getType()->getPointerAddressSpace();
426
427 if (std::optional<bool> Answer = evaluateIsSpace(IID, AS))
428 return IC.replaceInstUsesWith(I&: II,
429 V: ConstantInt::get(Ty: II.getType(), V: *Answer));
430 return nullptr; // Don't know the answer, got to check at run time.
431 }
432 default:
433 return std::nullopt;
434 }
435}
436
437static Instruction *foldAbsIntoRedux(InstCombiner &IC, IntrinsicInst &II) {
438 Intrinsic::ID AbsIID;
439 switch (II.getIntrinsicID()) {
440 case Intrinsic::nvvm_redux_sync_fmin:
441 case Intrinsic::nvvm_redux_sync_fmin_abs:
442 AbsIID = Intrinsic::nvvm_redux_sync_fmin_abs;
443 break;
444 case Intrinsic::nvvm_redux_sync_fmax:
445 case Intrinsic::nvvm_redux_sync_fmax_abs:
446 AbsIID = Intrinsic::nvvm_redux_sync_fmax_abs;
447 break;
448 case Intrinsic::nvvm_redux_sync_fmin_NaN:
449 case Intrinsic::nvvm_redux_sync_fmin_abs_NaN:
450 AbsIID = Intrinsic::nvvm_redux_sync_fmin_abs_NaN;
451 break;
452 case Intrinsic::nvvm_redux_sync_fmax_NaN:
453 case Intrinsic::nvvm_redux_sync_fmax_abs_NaN:
454 AbsIID = Intrinsic::nvvm_redux_sync_fmax_abs_NaN;
455 break;
456 default:
457 return nullptr;
458 }
459
460 auto *Abs = dyn_cast<IntrinsicInst>(Val: II.getArgOperand(i: 0));
461 if (!Abs || Abs->getIntrinsicID() != Intrinsic::fabs)
462 return nullptr;
463
464 II.setCalledFunction(
465 Intrinsic::getOrInsertDeclaration(M: II.getModule(), id: AbsIID));
466 // These attributes described the absolute value, not its source.
467 II.removeParamAttr(ArgNo: 0, Kind: Attribute::NoFPClass);
468 II.removeParamAttr(ArgNo: 0, Kind: Attribute::Returned);
469 return IC.replaceOperand(I&: II, OpNum: 0, V: Abs->getArgOperand(i: 0));
470}
471
472std::optional<Instruction *>
473NVPTXTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
474 if (std::optional<Instruction *> I = handleSpaceCheckIntrinsics(IC, II))
475 return *I;
476 if (Instruction *I = convertNvvmIntrinsicToLlvm(IC, II: &II))
477 return I;
478 if (Instruction *I = foldAbsIntoRedux(IC, II))
479 return I;
480 return std::nullopt;
481}
482
483InstructionCost
484NVPTXTTIImpl::getInstructionCost(const User *U,
485 ArrayRef<const Value *> Operands,
486 TTI::TargetCostKind CostKind) const {
487 if (const auto *CI = dyn_cast<CallInst>(Val: U))
488 if (const auto *IA = dyn_cast<InlineAsm>(Val: CI->getCalledOperand())) {
489 // Without this implementation getCallCost() would return the number
490 // of arguments+1 as the cost. Because the cost-model assumes it is a call
491 // since it is classified as a call in the IR. A better cost model would
492 // be to return the number of asm instructions embedded in the asm
493 // string.
494 StringRef AsmStr = IA->getAsmString();
495 const unsigned InstCount =
496 count_if(Range: split(Str: AsmStr, Separator: ';'), P: [](StringRef AsmInst) {
497 // Trim off scopes denoted by '{' and '}' as these can be ignored
498 AsmInst = AsmInst.trim().ltrim(Chars: "{} \t\n\v\f\r");
499 // This is pretty coarse but does a reasonably good job of
500 // identifying things that look like instructions, possibly with a
501 // predicate ("@").
502 return !AsmInst.empty() &&
503 (AsmInst[0] == '@' || isAlpha(C: AsmInst[0]) ||
504 AsmInst.contains(Other: ".pragma"));
505 });
506 return InstCount * TargetTransformInfo::TCC_Basic;
507 }
508
509 return BaseT::getInstructionCost(U, Operands, CostKind);
510}
511
512InstructionCost NVPTXTTIImpl::getArithmeticInstrCost(
513 unsigned Opcode, Type *Ty, TTI::TargetCostKind CostKind,
514 TTI::OperandValueInfo Op1Info, TTI::OperandValueInfo Op2Info,
515 ArrayRef<const Value *> Args, const Instruction *CtxI) const {
516 // Legalize the type.
517 std::pair<InstructionCost, MVT> LT = getTypeLegalizationCost(Ty);
518
519 int ISD = TLI->InstructionOpcodeToISD(Opcode);
520
521 switch (ISD) {
522 default:
523 return BaseT::getArithmeticInstrCost(Opcode, Ty, CostKind, Opd1Info: Op1Info,
524 Opd2Info: Op2Info);
525 case ISD::ADD:
526 case ISD::MUL:
527 case ISD::XOR:
528 case ISD::OR:
529 case ISD::AND:
530 // The machine code (SASS) simulates an i64 with two i32. Therefore, we
531 // estimate that arithmetic operations on i64 are twice as expensive as
532 // those on types that can fit into one machine register.
533 if (LT.second.SimpleTy == MVT::i64)
534 return 2 * LT.first;
535 // Delegate other cases to the basic TTI.
536 return BaseT::getArithmeticInstrCost(Opcode, Ty, CostKind, Opd1Info: Op1Info,
537 Opd2Info: Op2Info);
538 }
539}
540
541void NVPTXTTIImpl::getUnrollingPreferences(
542 Loop *L, ScalarEvolution &SE, TTI::UnrollingPreferences &UP,
543 OptimizationRemarkEmitter *ORE) const {
544 BaseT::getUnrollingPreferences(L, SE, UP, ORE);
545
546 // Enable partial unrolling and runtime unrolling, but reduce the
547 // threshold. This partially unrolls small loops which are often
548 // unrolled by the PTX to SASS compiler and unrolling earlier can be
549 // beneficial.
550 UP.Partial = UP.Runtime = true;
551 UP.PartialThreshold = UP.Threshold / 4;
552}
553
554void NVPTXTTIImpl::getPeelingPreferences(Loop *L, ScalarEvolution &SE,
555 TTI::PeelingPreferences &PP) const {
556 BaseT::getPeelingPreferences(L, SE, PP);
557}
558
559namespace {
560enum class InvalidAddressSpaceAction { LeaveUnchanged, Trap };
561
562// Describes how to specialize GenericID's PointerOperand from generic memory to
563// one of AddressSpaces, using SpecificID (which may be the same intrinsic).
564// Trap only when InvalidAction permits it and the pointer is provably outside
565// every listed address space; otherwise leave the intrinsic unchanged.
566struct IntrinsicAddressSpaceInfo {
567 Intrinsic::ID GenericID;
568 Intrinsic::ID SpecificID;
569 SmallVector<unsigned, 2> AddressSpaces;
570 InvalidAddressSpaceAction InvalidAction;
571 unsigned PointerOperand;
572
573 IntrinsicAddressSpaceInfo(
574 Intrinsic::ID GenericID, Intrinsic::ID SpecificID,
575 ArrayRef<unsigned> AddressSpaces = {NVPTXAS::ADDRESS_SPACE_SHARED},
576 InvalidAddressSpaceAction InvalidAction = InvalidAddressSpaceAction::Trap,
577 unsigned PointerOperand = 0)
578 : GenericID(GenericID), SpecificID(SpecificID),
579 AddressSpaces(AddressSpaces), InvalidAction(InvalidAction),
580 PointerOperand(PointerOperand) {}
581
582 IntrinsicAddressSpaceInfo(
583 Intrinsic::ID IID,
584 ArrayRef<unsigned> AddressSpaces = {NVPTXAS::ADDRESS_SPACE_SHARED},
585 InvalidAddressSpaceAction InvalidAction = InvalidAddressSpaceAction::Trap,
586 unsigned PointerOperand = 0)
587 : IntrinsicAddressSpaceInfo(IID, IID, AddressSpaces, InvalidAction,
588 PointerOperand) {}
589};
590} // namespace
591
592static const IntrinsicAddressSpaceInfo *
593getIntrinsicAddressSpaceInfo(Intrinsic::ID IID) {
594 static const auto Infos = [] {
595 const IntrinsicAddressSpaceInfo Entries[] = {
596 {Intrinsic::nvvm_mbarrier_init},
597 {Intrinsic::nvvm_mbarrier_inval, Intrinsic::nvvm_mbarrier_inval_shared},
598 {Intrinsic::nvvm_mbarrier_check_layout},
599 {Intrinsic::nvvm_mbarrier_arrive,
600 Intrinsic::nvvm_mbarrier_arrive_shared},
601 {Intrinsic::nvvm_mbarrier_arrive_noComplete,
602 Intrinsic::nvvm_mbarrier_arrive_noComplete_shared},
603 {Intrinsic::nvvm_mbarrier_arrive_drop,
604 Intrinsic::nvvm_mbarrier_arrive_drop_shared},
605 {Intrinsic::nvvm_mbarrier_arrive_drop_noComplete,
606 Intrinsic::nvvm_mbarrier_arrive_drop_noComplete_shared},
607 {Intrinsic::nvvm_mbarrier_test_wait,
608 Intrinsic::nvvm_mbarrier_test_wait_shared},
609 {Intrinsic::nvvm_cp_async_mbarrier_arrive,
610 Intrinsic::nvvm_cp_async_mbarrier_arrive_shared},
611 {Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc,
612 Intrinsic::nvvm_cp_async_mbarrier_arrive_noinc_shared},
613 // Other address spaces cannot be specialized but do not imply undefined
614 // behavior for prefetch.tensormap.
615 {Intrinsic::nvvm_prefetch_tensormap,
616 {NVPTXAS::ADDRESS_SPACE_CONST, NVPTXAS::ADDRESS_SPACE_ENTRY_PARAM},
617 InvalidAddressSpaceAction::LeaveUnchanged},
618 };
619 return SmallDenseMap<Intrinsic::ID, IntrinsicAddressSpaceInfo, 16>(
620 llvm::from_range, llvm::map_range(C: Entries, F: [](const auto &Info) {
621 return std::pair{Info.GenericID, Info};
622 }));
623 }();
624 auto It = Infos.find(Val: IID);
625 return It == Infos.end() ? nullptr : &It->second;
626}
627
628bool NVPTXTTIImpl::collectFlatAddressOperands(SmallVectorImpl<int> &OpIndexes,
629 Intrinsic::ID IID) const {
630 if (const auto *Info = getIntrinsicAddressSpaceInfo(IID)) {
631 OpIndexes.push_back(Elt: Info->PointerOperand);
632 return true;
633 }
634 switch (IID) {
635 case Intrinsic::nvvm_isspacep_const:
636 case Intrinsic::nvvm_isspacep_global:
637 case Intrinsic::nvvm_isspacep_local:
638 case Intrinsic::nvvm_isspacep_shared:
639 case Intrinsic::nvvm_isspacep_shared_cluster: {
640 OpIndexes.push_back(Elt: 0);
641 return true;
642 }
643 }
644 return false;
645}
646
647Value *NVPTXTTIImpl::rewriteIntrinsicWithAddressSpace(IntrinsicInst *II,
648 Value *OldV,
649 Value *NewV) const {
650 const Intrinsic::ID IID = II->getIntrinsicID();
651 if (const auto *Info = getIntrinsicAddressSpaceInfo(IID)) {
652 if (II->getArgOperand(i: Info->PointerOperand) != OldV)
653 return nullptr;
654
655 unsigned NewAS = NewV->getType()->getPointerAddressSpace();
656 if (!llvm::is_contained(Range: Info->AddressSpaces, Element: NewAS)) {
657 // A missing rewrite is not necessarily invalid. Trap only when the
658 // address is provably outside every allowed space and the entry permits
659 // it.
660 if (Info->InvalidAction == InvalidAddressSpaceAction::LeaveUnchanged ||
661 !llvm::all_of(Range: Info->AddressSpaces, P: [NewAS](unsigned AS) {
662 return isInAddressSpace(AS: NewAS, SpecificAS: AS) == false;
663 }))
664 return nullptr;
665 // InferAddressSpaces preserves the CFG. Emit a trap here and leave
666 // truncating the block at the noreturn call to SimplifyCFG.
667 Value *Trap = IRBuilder<>(II).CreateIntrinsic(ID: Intrinsic::trap, Args: {});
668 return II->getType()->isVoidTy() ? Trap : PoisonValue::get(T: II->getType());
669 }
670
671 SmallVector<Type *, 8> ArgTypes(II->getFunctionType()->params());
672 ArgTypes[Info->PointerOperand] = NewV->getType();
673 FunctionType *FT = FunctionType::get(Result: II->getType(), Params: ArgTypes, isVarArg: false);
674 SmallVector<Type *, 4> OverloadTypes;
675 if (!Intrinsic::isSignatureValid(ID: Info->SpecificID, FT, OverloadTys&: OverloadTypes))
676 return nullptr;
677 Function *Decl = Intrinsic::getOrInsertDeclaration(
678 M: II->getModule(), id: Info->SpecificID, OverloadTys: OverloadTypes);
679 II->setArgOperand(i: Info->PointerOperand, v: NewV);
680 II->setCalledFunction(Decl);
681 return II;
682 }
683 switch (IID) {
684 case Intrinsic::nvvm_isspacep_const:
685 case Intrinsic::nvvm_isspacep_global:
686 case Intrinsic::nvvm_isspacep_local:
687 case Intrinsic::nvvm_isspacep_shared:
688 case Intrinsic::nvvm_isspacep_shared_cluster: {
689 const unsigned NewAS = NewV->getType()->getPointerAddressSpace();
690 if (const auto R = evaluateIsSpace(IID, AS: NewAS))
691 return ConstantInt::get(Ty: II->getType(), V: *R);
692 return nullptr;
693 }
694 }
695 return nullptr;
696}
697
698bool NVPTXTTIImpl::isLegalMaskedStore(Type *DataTy, Align Alignment,
699 unsigned AddrSpace,
700 TTI::MaskKind MaskKind) const {
701 if (MaskKind != TTI::MaskKind::ConstantMask)
702 return false;
703
704 // We currently only support this feature for 256-bit vectors, so the
705 // alignment must be at least 32
706 if (Alignment < 32)
707 return false;
708
709 if (!ST->has256BitVectorLoadStore(AS: AddrSpace))
710 return false;
711
712 auto *VTy = dyn_cast<FixedVectorType>(Val: DataTy);
713 if (!VTy)
714 return false;
715
716 auto *ElemTy = VTy->getScalarType();
717 return (ElemTy->getScalarSizeInBits() == 32 && VTy->getNumElements() == 8) ||
718 (ElemTy->getScalarSizeInBits() == 64 && VTy->getNumElements() == 4);
719}
720
721bool NVPTXTTIImpl::isLegalMaskedLoad(Type *DataTy, Align Alignment,
722 unsigned /*AddrSpace*/,
723 TTI::MaskKind MaskKind) const {
724 if (MaskKind != TTI::MaskKind::ConstantMask)
725 return false;
726
727 if (Alignment < DL.getTypeStoreSize(Ty: DataTy))
728 return false;
729
730 // We do not support sub-byte element type masked loads.
731 auto *VTy = dyn_cast<FixedVectorType>(Val: DataTy);
732 if (!VTy)
733 return false;
734 return VTy->getElementType()->getScalarSizeInBits() >= 8;
735}
736
737unsigned NVPTXTTIImpl::getLoadStoreVecRegBitWidth(unsigned AddrSpace) const {
738 // 256 bit loads/stores are currently only supported for global address space
739 if (ST->has256BitVectorLoadStore(AS: AddrSpace))
740 return 256;
741 return 128;
742}
743
744unsigned NVPTXTTIImpl::getAssumedAddrSpace(const Value *V) const {
745 if (isa<AllocaInst>(Val: V))
746 return ADDRESS_SPACE_LOCAL;
747
748 if (const Argument *Arg = dyn_cast<Argument>(Val: V)) {
749 if (isKernelFunction(F: *Arg->getParent())) {
750 const NVPTXTargetMachine &TM =
751 static_cast<const NVPTXTargetMachine &>(getTLI()->getTargetMachine());
752 if (TM.getDrvInterface() == NVPTX::CUDA && !Arg->hasByValAttr())
753 return ADDRESS_SPACE_GLOBAL;
754 } else {
755 // We assume that all device parameters that are passed byval will be
756 // placed in the local AS. Very simple cases will be updated after ISel to
757 // use the device param space where possible.
758 if (Arg->hasByValAttr())
759 return ADDRESS_SPACE_LOCAL;
760 }
761 }
762
763 return -1;
764}
765
766void NVPTXTTIImpl::collectKernelLaunchBounds(
767 const Function &F,
768 SmallVectorImpl<std::pair<StringRef, int64_t>> &LB) const {
769 if (const auto Val = getMaxClusterRank(F))
770 LB.push_back(Elt: {"maxclusterrank", *Val});
771
772 const auto MaxNTID = getMaxNTID(F);
773 if (MaxNTID.size() > 0)
774 LB.push_back(Elt: {"maxntidx", MaxNTID[0]});
775 if (MaxNTID.size() > 1)
776 LB.push_back(Elt: {"maxntidy", MaxNTID[1]});
777 if (MaxNTID.size() > 2)
778 LB.push_back(Elt: {"maxntidz", MaxNTID[2]});
779}
780
781// Global addresses can only be materialized if they are in generic global or
782// constant space.
783bool NVPTXTTIImpl::shouldBuildLookupTablesForConstant(Constant *C) const {
784 if (const auto *GV = dyn_cast<GlobalValue>(Val: C)) {
785 const unsigned AS = GV->getAddressSpace();
786 return AS == NVPTXAS::ADDRESS_SPACE_GENERIC ||
787 AS == NVPTXAS::ADDRESS_SPACE_GLOBAL ||
788 AS == NVPTXAS::ADDRESS_SPACE_CONST;
789 }
790 return true;
791}
792
793ValueUniformity NVPTXTTIImpl::getValueUniformity(const Value *V) const {
794 if (isSourceOfDivergence(V))
795 return ValueUniformity::NeverUniform;
796
797 return ValueUniformity::Default;
798}
799