1//===-- SPIRVPrepareGlobals.cpp - Prepare IR SPIRV globals ------*- 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// The pass:
10// - transforms IR globals that cannot be trivially mapped to SPIRV into
11// something that is trival to lower;
12// - for AMDGCN flavoured SPIRV, it assigns unique IDs to the specialisation
13// constants associated with feature predicates, which were inserted by the
14// FE when expanding calls to __builtin_amdgcn_processor_is or
15// __builtin_amdgcn_is_invocable
16//
17//===----------------------------------------------------------------------===//
18
19#include "SPIRV.h"
20#include "SPIRVUtils.h"
21
22#include "llvm/ADT/STLExtras.h"
23#include "llvm/ADT/StringExtras.h"
24#include "llvm/ADT/StringMap.h"
25#include "llvm/IR/IntrinsicsSPIRV.h"
26#include "llvm/IR/Module.h"
27#include "llvm/Support/Debug.h"
28
29#include <string>
30
31#define DEBUG_TYPE "spirv-prepare-globals"
32
33using namespace llvm;
34
35namespace {
36
37struct SPIRVPrepareGlobalsImpl {
38 bool runOnModule(Module &M);
39};
40
41struct SPIRVPrepareGlobalsLegacy : public ModulePass {
42 static char ID;
43 SPIRVPrepareGlobalsLegacy() : ModulePass(ID) {}
44
45 StringRef getPassName() const override {
46 return "SPIRV prepare global variables";
47 }
48
49 bool runOnModule(Module &M) override {
50 return SPIRVPrepareGlobalsImpl().runOnModule(M);
51 }
52};
53
54// The backend does not support GlobalAlias. Replace aliases with their aliasees
55// when possible and remove them from the module.
56bool tryReplaceAliasWithAliasee(GlobalAlias &GA) {
57 // According to the lang ref, aliases cannot be replaced if either the alias
58 // or the aliasee are interposable. We only replace in the case that both
59 // are not interposable.
60 if (GA.isInterposable()) {
61 LLVM_DEBUG(dbgs() << "Skipping interposable alias: " << GA.getName()
62 << "\n");
63 return false;
64 }
65
66 auto *AO = dyn_cast<GlobalObject>(Val: GA.getAliasee());
67 if (!AO) {
68 LLVM_DEBUG(dbgs() << "Skipping alias whose aliasee is not a GlobalObject: "
69 << GA.getName() << "\n");
70 return false;
71 }
72
73 // noipa doesn't affect which definition is used at link time, so it doesn't
74 // prevent replacing the alias.
75 if (AO->isInterposable(/*CheckNoIPA=*/false)) {
76 LLVM_DEBUG(dbgs() << "Skipping interposable aliasee: " << AO->getName()
77 << "\n");
78 return false;
79 }
80
81 LLVM_DEBUG(dbgs() << "Replacing alias " << GA.getName()
82 << " with aliasee: " << AO->getName() << "\n");
83
84 GA.replaceAllUsesWith(V: AO);
85 if (GA.isDiscardableIfUnused()) {
86 GA.eraseFromParent();
87 }
88
89 return true;
90}
91
92bool tryAssignPredicateSpecConstIDs(Module &M, Function *F) {
93 StringMap<unsigned> IDs;
94 for (auto &&U : F->users()) {
95 auto *CI = dyn_cast<CallInst>(Val: U);
96 if (!CI)
97 continue;
98
99 auto *SpecID = dyn_cast<ConstantInt>(Val: CI->getArgOperand(i: 0));
100 if (!SpecID)
101 continue;
102
103 unsigned ID = SpecID->getZExtValue();
104 if (ID != UINT32_MAX)
105 continue;
106
107 // Replace placeholder Specialisation Constant IDs with unique IDs
108 // associated with the predicate being evaluated, which is encoded via
109 // spv_assign_name.
110 auto *MD =
111 cast<MDNode>(Val: cast<MetadataAsValue>(Val: CI->getOperand(i_nocapture: 2))->getMetadata());
112 auto *P = cast<MDString>(Val: MD->getOperand(I: 0));
113
114 ID = IDs.try_emplace(Key: P->getString(), Args: IDs.size()).first->second;
115 CI->setArgOperand(i: 0, v: ConstantInt::get(Ty: CI->getArgOperand(i: 0)->getType(), V: ID));
116 }
117
118 if (IDs.empty())
119 return false;
120
121 // Store the predicate -> ID mapping as a fixed format string
122 // (predicate ID\0...), for later use during SPIR-V consumption.
123 std::string Tmp;
124 for (auto &&[Predicate, SpecID] : IDs)
125 Tmp.append(svt: Predicate).append(s: " ").append(str: utostr(X: SpecID)).push_back(c: '\0');
126
127 Constant *PredSpecIDStr =
128 ConstantDataArray::getString(Context&: M.getContext(), Initializer: Tmp, AddNull: false);
129
130 new GlobalVariable(M, PredSpecIDStr->getType(), true,
131 GlobalVariable::LinkageTypes::ExternalLinkage,
132 PredSpecIDStr, "llvm.amdgcn.feature.predicate.ids");
133
134 return true;
135}
136
137bool SPIRVPrepareGlobalsImpl::runOnModule(Module &M) {
138 bool Changed = false;
139
140 for (GlobalAlias &GA : make_early_inc_range(Range: M.aliases())) {
141 Changed |= tryReplaceAliasWithAliasee(GA);
142 }
143
144 if (M.getTargetTriple().getVendor() != Triple::AMD)
145 return Changed;
146
147 // TODO: Currently, for AMDGCN flavoured SPIR-V, the symbol can only be
148 // inserted via feature predicate use, but in the future this will need
149 // revisiting if we start making more liberal use of the intrinsic.
150 if (Function *F = Intrinsic::getDeclarationIfExists(
151 M: &M, id: Intrinsic::spv_named_boolean_spec_constant))
152 Changed |= tryAssignPredicateSpecConstIDs(M, F);
153
154 return Changed;
155}
156char SPIRVPrepareGlobalsLegacy::ID = 0;
157
158} // namespace
159
160INITIALIZE_PASS(SPIRVPrepareGlobalsLegacy, "spirv-prepare-globals",
161 "SPIRV prepare global variables", false, false)
162
163PreservedAnalyses SPIRVPrepareGlobalsPass::run(Module &M,
164 ModuleAnalysisManager &AM) {
165 return SPIRVPrepareGlobalsImpl().runOnModule(M) ? PreservedAnalyses::none()
166 : PreservedAnalyses::all();
167}
168
169namespace llvm {
170ModulePass *createSPIRVPrepareGlobalsPass() {
171 return new SPIRVPrepareGlobalsLegacy();
172}
173} // namespace llvm
174