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 bool hasCLMAD() const {
124 return hasFeature(Feature: NVPTX::SM80) && hasFeature(Feature: NVPTX::PTX93);
125 }
126 // Cache hint SM/PTX version requirements
127 bool hasL1EvictionHint() const {
128 return hasFeature(Feature: NVPTX::SM70) && hasFeature(Feature: NVPTX::PTX74);
129 }
130 bool hasL2EvictionHint() const {
131 return hasFeature(Feature: NVPTX::SM100) && hasFeature(Feature: NVPTX::PTX88);
132 }
133 bool hasL2Prefetch64B() const {
134 return hasFeature(Feature: NVPTX::SM75) && hasFeature(Feature: NVPTX::PTX74);
135 }
136 bool hasL2Prefetch128B() const {
137 return hasFeature(Feature: NVPTX::SM75) && hasFeature(Feature: NVPTX::PTX74);
138 }
139 bool hasL2Prefetch256B() const {
140 return hasFeature(Feature: NVPTX::SM80) && hasFeature(Feature: NVPTX::PTX74);
141 }
142 bool hasL2CacheHint() const {
143 return hasFeature(Feature: NVPTX::SM80) && hasFeature(Feature: NVPTX::PTX74);
144 }
145
146 // Checks following instructions support:
147 // - tcgen05.ld/st
148 // - tcgen05.alloc/dealloc/relinquish
149 // - tcgen05.cp
150 // - tcgen05.fence/wait
151 // - tcgen05.commit
152 // - tcgen05.mma
153 bool hasTcgen05InstSupport() const {
154 return hasAnyFeature(Features: {NVPTX::SM100f, NVPTX::SM110f});
155 }
156
157 // f32x2 instructions in Blackwell family
158 bool hasF32x2Instructions() const;
159
160 bool hasTensormapReplaceSupport() const {
161 return hasAnyFeature(Features: {NVPTX::SM100f, NVPTX::SM110f, NVPTX::SM120f}) ||
162 (hasFeature(Feature: NVPTX::PTX83) && hasAnyFeature(Features: {NVPTX::SM90a}));
163 }
164
165 bool hasTensormapReplaceElemtypeSupport(unsigned ElemType) const {
166 if (ElemType >= static_cast<unsigned>(nvvm::TensormapElemType::B4x16))
167 return (hasFeature(Feature: NVPTX::PTX88) &&
168 hasAnyFeature(Features: {NVPTX::SM100f, NVPTX::SM110f, NVPTX::SM120f})) ||
169 (hasFeature(Feature: NVPTX::PTX87) &&
170 hasAnyFeature(Features: {NVPTX::SM100a, NVPTX::SM110a, NVPTX::SM120a}));
171
172 return hasTensormapReplaceSupport();
173 }
174
175 bool hasTensormapReplaceSwizzleModeSupport(unsigned SwizzleMode) const {
176 if (SwizzleMode ==
177 static_cast<unsigned>(nvvm::TensormapSwizzleMode::SWIZZLE_96B))
178 return hasAnyFeature(Features: {NVPTX::SM103a});
179
180 return hasTensormapReplaceSupport();
181 }
182
183 // Prior to CUDA 12.3 ptxas did not recognize that the trap instruction
184 // terminates a basic block. Instead, it would assume that control flow
185 // continued to the next instruction. The next instruction could be in the
186 // block that's lexically below it. This would lead to a phantom CFG edges
187 // being created within ptxas. This issue was fixed in CUDA 12.3. Thus, when
188 // PTX ISA versions 8.3+ we can confidently say that the bug will not be
189 // present.
190 bool hasPTXASUnreachableBug() const { return !hasFeature(Feature: NVPTX::PTX83); }
191 bool hasCvtaParam() const {
192 return hasFeature(Feature: NVPTX::SM70) && hasFeature(Feature: NVPTX::PTX77);
193 }
194 // The compute capability as a number, for __CUDA_ARCH__. This is the one
195 // place an architecture needs to be a number, and it is not an identity:
196 // sm_100, sm_100f and sm_100a all report 100.
197 unsigned getSmVersion() const { return NVPTX::getSmVersion(Kind: Arch) / 10; }
198
199 // Whether -mcpu named a target at all, as opposed to falling back to the
200 // default architecture.
201 bool hasTargetName() const { return !getCPU().empty(); }
202
203 // The architecture's name, which is what `.target` is emitted from.
204 StringRef getTargetName() const;
205
206 bool hasNativeBF16Support(unsigned Opcode) const;
207
208 // Get maximum value of required alignments among the supported data types.
209 // From the PTX ISA doc, section 8.2.3:
210 // The memory consistency model relates operations executed on memory
211 // locations with scalar data-types, which have a maximum size and alignment
212 // of 64 bits. Memory operations with a vector data-type are modelled as a
213 // set of equivalent memory operations with a scalar data-type, executed in
214 // an unspecified order on the elements in the vector.
215 unsigned getMaxRequiredAlignment() const { return 8; }
216 // Get the smallest cmpxchg word size that the hardware supports.
217 unsigned getMinCmpXchgSizeInBits() const { return 32; }
218
219 unsigned getPTXVersion() const { return PTXVersion; }
220
221 NVPTXSubtarget &initializeSubtargetDependencies(StringRef CPU, StringRef FS);
222 void ParseSubtargetFeatures(StringRef CPU, StringRef TuneCPU, StringRef FS);
223};
224
225} // namespace llvm
226
227#endif
228