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