1//=====-- NVPTXSubtarget.h - Define Subtarget for the NVPTX ---*- C++ -*--====//
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 declares the NVPTX specific subclass of TargetSubtarget.
10//
11//===----------------------------------------------------------------------===//
12
13#ifndef LLVM_LIB_TARGET_NVPTX_NVPTXSUBTARGET_H
14#define LLVM_LIB_TARGET_NVPTX_NVPTXSUBTARGET_H
15
16#include "NVPTX.h"
17#include "NVPTXFrameLowering.h"
18#include "NVPTXISelLowering.h"
19#include "NVPTXInstrInfo.h"
20#include "NVPTXRegisterInfo.h"
21#include "llvm/CodeGen/TargetSubtargetInfo.h"
22#include "llvm/IR/DataLayout.h"
23#include "llvm/IR/NVVMIntrinsicUtils.h"
24#include "llvm/Support/NVPTXAddrSpace.h"
25#include "llvm/TargetParser/NVPTXTargetParser.h"
26
27#define GET_SUBTARGETINFO_HEADER
28#include "NVPTXGenSubtargetInfo.inc"
29
30namespace llvm {
31
32class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
33 virtual void anchor();
34
35 // PTX version x.y is represented as 10*x+y, e.g. 3.1 == 31
36 unsigned PTXVersion;
37
38 NVPTX::GPUKind Arch = NVPTX::GK_NONE;
39
40 // Set by every architecture feature. Their bits are the point, so this is
41 // only here because a subtarget feature must name a field.
42 bool HasArchitecture = false;
43
44 NVPTXInstrInfo InstrInfo;
45 NVPTXTargetLowering TLInfo;
46 std::unique_ptr<const SelectionDAGTargetInfo> TSInfo;
47
48 // NVPTX does not have any call stack frame, but need a NVPTX specific
49 // FrameLowering class because TargetFrameLowering is abstract.
50 NVPTXFrameLowering FrameLowering;
51
52public:
53 /// This constructor initializes the data members to match that
54 /// of the specified module.
55 ///
56 NVPTXSubtarget(const Triple &TT, StringRef CPU, StringRef FS,
57 const NVPTXTargetMachine &TM);
58
59 ~NVPTXSubtarget() override;
60
61 const TargetFrameLowering *getFrameLowering() const override {
62 return &FrameLowering;
63 }
64 const NVPTXInstrInfo *getInstrInfo() const override { return &InstrInfo; }
65 const NVPTXRegisterInfo *getRegisterInfo() const override {
66 return &InstrInfo.getRegisterInfo();
67 }
68 const NVPTXTargetLowering *getTargetLowering() const override {
69 return &TLInfo;
70 }
71
72 const SelectionDAGTargetInfo *getSelectionDAGInfo() const override {
73 return TSInfo.get();
74 }
75
76 // True when any of `Features` is enabled.
77 bool hasAnyFeature(ArrayRef<unsigned> Features) const {
78 return llvm::any_of(Range&: Features, P: [this](unsigned F) { return hasFeature(Feature: F); });
79 }
80
81 bool has256BitVectorLoadStore(unsigned AS) const {
82 return hasFeature(Feature: NVPTX::SM100) && hasFeature(Feature: NVPTX::PTX88) &&
83 AS == NVPTXAS::ADDRESS_SPACE_GLOBAL;
84 }
85 bool hasUsedBytesMaskPragma() const {
86 return hasFeature(Feature: NVPTX::SM50) && hasFeature(Feature: NVPTX::PTX83);
87 }
88 bool hasAtomAddF64() const { return hasFeature(Feature: NVPTX::SM60); }
89 bool hasAtomScope() const { return hasFeature(Feature: NVPTX::SM60); }
90 bool hasAtomBitwise64() const { return hasFeature(Feature: NVPTX::SM32); }
91 bool hasAtomMinMax64() const { return hasFeature(Feature: NVPTX::SM32); }
92 bool hasAtomCas16() const {
93 return hasFeature(Feature: NVPTX::SM70) && hasFeature(Feature: NVPTX::PTX63);
94 }
95 bool hasAtomSwap128() const {
96 return hasFeature(Feature: NVPTX::SM90) && hasFeature(Feature: NVPTX::PTX83);
97 }
98 bool hasClusters() const { return hasFeature(Feature: NVPTX::SM90); }
99 bool hasLDG() const { return hasFeature(Feature: NVPTX::SM32); }
100 bool hasHWROT32() const { return hasFeature(Feature: NVPTX::SM32); }
101 bool hasBrx() const {
102 return hasFeature(Feature: NVPTX::SM30) && hasFeature(Feature: NVPTX::PTX60);
103 }
104 bool hasFP16Math() const { return hasFeature(Feature: NVPTX::SM53); }
105 bool hasBF16Math() const { return hasFeature(Feature: NVPTX::SM80); }
106 bool allowFP16Math() const;
107 bool hasMaskOperator() const { return hasFeature(Feature: NVPTX::PTX71); }
108 bool hasNoReturn() const {
109 return hasFeature(Feature: NVPTX::SM30) && hasFeature(Feature: NVPTX::PTX64);
110 }
111 // Does SM & PTX support memory orderings (weak and atomic: relaxed, acquire,
112 // release, acq_rel, sc) ?
113 bool hasMemoryOrdering() const { return hasFeature(Feature: NVPTX::SM70); }
114 // Does SM & PTX support .acquire and .release qualifiers for fence?
115 bool hasSplitAcquireAndReleaseFences() const {
116 return hasFeature(Feature: NVPTX::SM90) && hasFeature(Feature: NVPTX::PTX86);
117 }
118 // Does SM & PTX support atomic relaxed MMIO operations ?
119 bool hasRelaxedMMIO() const {
120 return hasFeature(Feature: NVPTX::SM70) && hasFeature(Feature: NVPTX::PTX82);
121 }
122 bool hasDotInstructions() const { return hasFeature(Feature: NVPTX::SM61); }
123 // Cache hint SM/PTX version requirements
124 bool hasL1EvictionHint() const {
125 return hasFeature(Feature: NVPTX::SM70) && hasFeature(Feature: NVPTX::PTX74);
126 }
127 bool hasL2EvictionHint() const {
128 return hasFeature(Feature: NVPTX::SM100) && hasFeature(Feature: NVPTX::PTX88);
129 }
130 bool hasL2Prefetch64B() const {
131 return hasFeature(Feature: NVPTX::SM75) && hasFeature(Feature: NVPTX::PTX74);
132 }
133 bool hasL2Prefetch128B() const {
134 return hasFeature(Feature: NVPTX::SM75) && hasFeature(Feature: NVPTX::PTX74);
135 }
136 bool hasL2Prefetch256B() const {
137 return hasFeature(Feature: NVPTX::SM80) && hasFeature(Feature: NVPTX::PTX74);
138 }
139 bool hasL2CacheHint() const {
140 return hasFeature(Feature: NVPTX::SM80) && hasFeature(Feature: NVPTX::PTX74);
141 }
142
143 // Checks following instructions support:
144 // - tcgen05.ld/st
145 // - tcgen05.alloc/dealloc/relinquish
146 // - tcgen05.cp
147 // - tcgen05.fence/wait
148 // - tcgen05.commit
149 // - tcgen05.mma
150 bool hasTcgen05InstSupport() const {
151 return hasAnyFeature(Features: {NVPTX::SM100f, NVPTX::SM110f});
152 }
153
154 // f32x2 instructions in Blackwell family
155 bool hasF32x2Instructions() const;
156
157 bool hasTensormapReplaceSupport() const {
158 return hasAnyFeature(Features: {NVPTX::SM100f, NVPTX::SM110f, NVPTX::SM120f}) ||
159 (hasFeature(Feature: NVPTX::PTX83) && hasAnyFeature(Features: {NVPTX::SM90a}));
160 }
161
162 bool hasTensormapReplaceElemtypeSupport(unsigned ElemType) const {
163 if (ElemType >= static_cast<unsigned>(nvvm::TensormapElemType::B4x16))
164 return (hasFeature(Feature: NVPTX::PTX88) &&
165 hasAnyFeature(Features: {NVPTX::SM100f, NVPTX::SM110f, NVPTX::SM120f})) ||
166 (hasFeature(Feature: NVPTX::PTX87) &&
167 hasAnyFeature(Features: {NVPTX::SM100a, NVPTX::SM110a, NVPTX::SM120a}));
168
169 return hasTensormapReplaceSupport();
170 }
171
172 bool hasTensormapReplaceSwizzleModeSupport(unsigned SwizzleMode) const {
173 if (SwizzleMode ==
174 static_cast<unsigned>(nvvm::TensormapSwizzleMode::SWIZZLE_96B))
175 return hasAnyFeature(Features: {NVPTX::SM103a});
176
177 return hasTensormapReplaceSupport();
178 }
179
180 // Prior to CUDA 12.3 ptxas did not recognize that the trap instruction
181 // terminates a basic block. Instead, it would assume that control flow
182 // continued to the next instruction. The next instruction could be in the
183 // block that's lexically below it. This would lead to a phantom CFG edges
184 // being created within ptxas. This issue was fixed in CUDA 12.3. Thus, when
185 // PTX ISA versions 8.3+ we can confidently say that the bug will not be
186 // present.
187 bool hasPTXASUnreachableBug() const { return !hasFeature(Feature: NVPTX::PTX83); }
188 bool hasCvtaParam() const {
189 return hasFeature(Feature: NVPTX::SM70) && hasFeature(Feature: NVPTX::PTX77);
190 }
191 // The compute capability as a number, for __CUDA_ARCH__. This is the one
192 // place an architecture needs to be a number, and it is not an identity:
193 // sm_100, sm_100f and sm_100a all report 100.
194 unsigned getSmVersion() const { return NVPTX::getSmVersion(Kind: Arch) / 10; }
195
196 // Whether -mcpu named a target at all, as opposed to falling back to the
197 // default architecture.
198 bool hasTargetName() const { return !getCPU().empty(); }
199
200 // The architecture's name, which is what `.target` is emitted from.
201 StringRef getTargetName() const;
202
203 bool hasNativeBF16Support(unsigned Opcode) const;
204
205 // Get maximum value of required alignments among the supported data types.
206 // From the PTX ISA doc, section 8.2.3:
207 // The memory consistency model relates operations executed on memory
208 // locations with scalar data-types, which have a maximum size and alignment
209 // of 64 bits. Memory operations with a vector data-type are modelled as a
210 // set of equivalent memory operations with a scalar data-type, executed in
211 // an unspecified order on the elements in the vector.
212 unsigned getMaxRequiredAlignment() const { return 8; }
213 // Get the smallest cmpxchg word size that the hardware supports.
214 unsigned getMinCmpXchgSizeInBits() const { return 32; }
215
216 unsigned getPTXVersion() const { return PTXVersion; }
217
218 NVPTXSubtarget &initializeSubtargetDependencies(StringRef CPU, StringRef FS);
219 void ParseSubtargetFeatures(StringRef CPU, StringRef TuneCPU, StringRef FS);
220};
221
222} // namespace llvm
223
224#endif
225