261 lines · c
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-exception6//7//===----------------------------------------------------------------------===//8//9// This file declares the NVPTX specific subclass of TargetSubtarget.10//11//===----------------------------------------------------------------------===//12 13#ifndef LLVM_LIB_TARGET_NVPTX_NVPTXSUBTARGET_H14#define LLVM_LIB_TARGET_NVPTX_NVPTXSUBTARGET_H15 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/Support/NVPTXAddrSpace.h"24#include <string>25 26#define GET_SUBTARGETINFO_HEADER27#include "NVPTXGenSubtargetInfo.inc"28 29namespace llvm {30 31class NVPTXSubtarget : public NVPTXGenSubtargetInfo {32 virtual void anchor();33 std::string TargetName;34 35 // PTX version x.y is represented as 10*x+y, e.g. 3.1 == 3136 unsigned PTXVersion;37 38 // Full SM version x.y is represented as 100*x+10*y+feature, e.g. 3.1 == 31039 // sm_90a == 90140 unsigned int FullSmVersion;41 42 // SM version x.y is represented as 10*x+y, e.g. 3.1 == 31. Derived from43 // FullSmVersion.44 unsigned int SmVersion;45 46 NVPTXInstrInfo InstrInfo;47 NVPTXTargetLowering TLInfo;48 std::unique_ptr<const SelectionDAGTargetInfo> TSInfo;49 50 // NVPTX does not have any call stack frame, but need a NVPTX specific51 // FrameLowering class because TargetFrameLowering is abstract.52 NVPTXFrameLowering FrameLowering;53 54public:55 /// This constructor initializes the data members to match that56 /// of the specified module.57 ///58 NVPTXSubtarget(const Triple &TT, const std::string &CPU,59 const std::string &FS, const NVPTXTargetMachine &TM);60 61 ~NVPTXSubtarget() override;62 63 const TargetFrameLowering *getFrameLowering() const override {64 return &FrameLowering;65 }66 const NVPTXInstrInfo *getInstrInfo() const override { return &InstrInfo; }67 const NVPTXRegisterInfo *getRegisterInfo() const override {68 return &InstrInfo.getRegisterInfo();69 }70 const NVPTXTargetLowering *getTargetLowering() const override {71 return &TLInfo;72 }73 74 const SelectionDAGTargetInfo *getSelectionDAGInfo() const override;75 76 // Checks PTX version and family-specific and architecture-specific SM77 // versions. For example, sm_100{f/a} and any future variants in the same78 // family will match for any PTX version greater than or equal to79 // `PTXVersion`.80 bool hasPTXWithFamilySMs(unsigned PTXVersion,81 ArrayRef<unsigned> SMVersions) const;82 // Checks PTX version and architecture-specific SM versions.83 // For example, sm_100{a} will match for any PTX version greater than or equal84 // to `PTXVersion`.85 bool hasPTXWithAccelSMs(unsigned PTXVersion,86 ArrayRef<unsigned> SMVersions) const;87 88 bool has256BitVectorLoadStore(unsigned AS) const {89 return SmVersion >= 100 && PTXVersion >= 88 &&90 AS == NVPTXAS::ADDRESS_SPACE_GLOBAL;91 }92 bool hasAtomAddF64() const { return SmVersion >= 60; }93 bool hasAtomScope() const { return SmVersion >= 60; }94 bool hasAtomBitwise64() const { return SmVersion >= 32; }95 bool hasAtomMinMax64() const { return SmVersion >= 32; }96 bool hasAtomCas16() const { return SmVersion >= 70 && PTXVersion >= 63; }97 bool hasAtomSwap128() const { return SmVersion >= 90 && PTXVersion >= 83; }98 bool hasClusters() const { return SmVersion >= 90 && PTXVersion >= 78; }99 bool hasLDG() const { return SmVersion >= 32; }100 bool hasHWROT32() const { return SmVersion >= 32; }101 bool hasFP16Math() const { return SmVersion >= 53; }102 bool hasBF16Math() const { return SmVersion >= 80; }103 bool allowFP16Math() const;104 bool hasMaskOperator() const { return PTXVersion >= 71; }105 bool hasNoReturn() const { return SmVersion >= 30 && PTXVersion >= 64; }106 // Does SM & PTX support memory orderings (weak and atomic: relaxed, acquire,107 // release, acq_rel, sc) ?108 bool hasMemoryOrdering() const { return SmVersion >= 70 && PTXVersion >= 60; }109 // Does SM & PTX support .acquire and .release qualifiers for fence?110 bool hasSplitAcquireAndReleaseFences() const {111 return SmVersion >= 90 && PTXVersion >= 86;112 }113 // Does SM & PTX support atomic relaxed MMIO operations ?114 bool hasRelaxedMMIO() const { return SmVersion >= 70 && PTXVersion >= 82; }115 bool hasDotInstructions() const {116 return SmVersion >= 61 && PTXVersion >= 50;117 }118 // Tcgen05 instructions in Blackwell family119 bool hasTcgen05Instructions() const {120 bool HasTcgen05 = false;121 unsigned MinPTXVersion = 86;122 switch (FullSmVersion) {123 default:124 break;125 case 1003: // sm_100a126 case 1013: // sm_101a127 HasTcgen05 = true;128 break;129 case 1103: // sm_110a130 HasTcgen05 = true;131 MinPTXVersion = 90;132 break;133 case 1033: // sm_103a134 HasTcgen05 = true;135 MinPTXVersion = 88;136 break;137 }138 139 return HasTcgen05 && PTXVersion >= MinPTXVersion;140 }141 142 // Checks following instructions support:143 // - tcgen05.ld/st144 // - tcgen05.alloc/dealloc/relinquish145 // - tcgen05.cp146 // - tcgen05.fence/wait147 // - tcgen05.commit148 bool hasTcgen05InstSupport() const {149 // sm_101 renamed to sm_110 in PTX 9.0150 return hasPTXWithFamilySMs(90, {100, 110}) ||151 hasPTXWithFamilySMs(88, {100, 101}) ||152 hasPTXWithAccelSMs(86, {100, 101});153 }154 155 // Checks tcgen05.shift instruction support.156 bool hasTcgen05ShiftSupport() const {157 // sm_101 renamed to sm_110 in PTX 9.0158 return hasPTXWithAccelSMs(90, {100, 110, 103}) ||159 hasPTXWithAccelSMs(88, {100, 101, 103}) ||160 hasPTXWithAccelSMs(86, {100, 101});161 }162 163 bool hasTcgen05MMAScaleInputDImm() const {164 return FullSmVersion == 1003 && PTXVersion >= 86;165 }166 // f32x2 instructions in Blackwell family167 bool hasF32x2Instructions() const;168 169 // Checks support for following in TMA:170 // - cta_group::1/2 support171 // - im2col_w/w_128 mode support172 // - tile_gather4 mode support173 // - tile_scatter4 mode support174 bool hasTMABlackwellSupport() const {175 return hasPTXWithFamilySMs(90, {100, 110}) ||176 hasPTXWithFamilySMs(88, {100, 101}) ||177 hasPTXWithAccelSMs(86, {100, 101});178 }179 180 // Checks support for conversions involving e4m3x2 and e5m2x2.181 bool hasFP8ConversionSupport() const {182 if (PTXVersion >= 81)183 return SmVersion >= 89;184 185 if (PTXVersion >= 78)186 return SmVersion >= 90;187 188 return false;189 }190 191 // Checks support for conversions involving the following types:192 // - e2m3x2/e3m2x2193 // - e2m1x2194 // - ue8m0x2195 bool hasNarrowFPConversionSupport() const {196 return hasPTXWithFamilySMs(90, {100, 110, 120}) ||197 hasPTXWithFamilySMs(88, {100, 101, 120}) ||198 hasPTXWithAccelSMs(86, {100, 101, 120});199 }200 201 // Prior to CUDA 12.3 ptxas did not recognize that the trap instruction202 // terminates a basic block. Instead, it would assume that control flow203 // continued to the next instruction. The next instruction could be in the204 // block that's lexically below it. This would lead to a phantom CFG edges205 // being created within ptxas. This issue was fixed in CUDA 12.3. Thus, when206 // PTX ISA versions 8.3+ we can confidently say that the bug will not be207 // present.208 bool hasPTXASUnreachableBug() const { return PTXVersion < 83; }209 bool hasCvtaParam() const { return SmVersion >= 70 && PTXVersion >= 77; }210 unsigned int getFullSmVersion() const { return FullSmVersion; }211 unsigned int getSmVersion() const { return getFullSmVersion() / 10; }212 unsigned int getSmFamilyVersion() const { return getFullSmVersion() / 100; }213 // GPUs with "a" suffix have architecture-accelerated features that are214 // supported on the specified architecture only, hence such targets do not215 // follow the onion layer model. hasArchAccelFeatures() allows distinguishing216 // such GPU variants from the base GPU architecture.217 // - false represents non-accelerated architecture.218 // - true represents architecture-accelerated variant.219 bool hasArchAccelFeatures() const {220 return (getFullSmVersion() & 1) && PTXVersion >= 80;221 }222 // GPUs with 'f' suffix have architecture-accelerated features which are223 // portable across all future architectures under same SM major. For example,224 // sm_100f features will work for sm_10X*f*/sm_10X*a* future architectures.225 // - false represents non-family-specific architecture.226 // - true represents family-specific variant.227 bool hasFamilySpecificFeatures() const {228 return getFullSmVersion() % 10 == 2 ? PTXVersion >= 88229 : hasArchAccelFeatures();230 }231 // If the user did not provide a target we default to the `sm_30` target.232 std::string getTargetName() const {233 return TargetName.empty() ? "sm_30" : TargetName;234 }235 bool hasTargetName() const { return !TargetName.empty(); }236 237 bool hasNativeBF16Support(int Opcode) const;238 239 // Get maximum value of required alignments among the supported data types.240 // From the PTX ISA doc, section 8.2.3:241 // The memory consistency model relates operations executed on memory242 // locations with scalar data-types, which have a maximum size and alignment243 // of 64 bits. Memory operations with a vector data-type are modelled as a244 // set of equivalent memory operations with a scalar data-type, executed in245 // an unspecified order on the elements in the vector.246 unsigned getMaxRequiredAlignment() const { return 8; }247 // Get the smallest cmpxchg word size that the hardware supports.248 unsigned getMinCmpXchgSizeInBits() const { return 32; }249 250 unsigned getPTXVersion() const { return PTXVersion; }251 252 NVPTXSubtarget &initializeSubtargetDependencies(StringRef CPU, StringRef FS);253 void ParseSubtargetFeatures(StringRef CPU, StringRef TuneCPU, StringRef FS);254 255 void failIfClustersUnsupported(std::string const &FailureMessage) const;256};257 258} // End llvm namespace259 260#endif261