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 "SPIRVPrepareGlobals.h"
20#include "SPIRV.h"
21#include "SPIRVUtils.h"
22
23#include "llvm/ADT/STLExtras.h"
24#include "llvm/ADT/StringExtras.h"
25#include "llvm/ADT/StringMap.h"
26#include "llvm/IR/IntrinsicsSPIRV.h"
27#include "llvm/IR/Module.h"
28#include "llvm/Support/Debug.h"
29
30#include <climits>
31#include <string>
32
33#define DEBUG_TYPE "spirv-prepare-globals"
34
35using namespace llvm;
36
37namespace {
38
39struct SPIRVPrepareGlobalsImpl {
40 bool runOnModule(Module &M);
41};
42
43struct SPIRVPrepareGlobalsLegacy : public ModulePass {
44 static char ID;
45 SPIRVPrepareGlobalsLegacy() : ModulePass(ID) {}
46
47 StringRef getPassName() const override {
48 return "SPIRV prepare global variables";
49 }
50
51 bool runOnModule(Module &M) override {
52 return SPIRVPrepareGlobalsImpl().runOnModule(M);
53 }
54};
55
56// The backend does not support GlobalAlias. Replace aliases with their aliasees
57// when possible and remove them from the module.
58bool tryReplaceAliasWithAliasee(GlobalAlias &GA) {
59 // According to the lang ref, aliases cannot be replaced if either the alias
60 // or the aliasee are interposable. We only replace in the case that both
61 // are not interposable.
62 if (GA.isInterposable()) {
63 LLVM_DEBUG(dbgs() << "Skipping interposable alias: " << GA.getName()
64 << "\n");
65 return false;
66 }
67
68 auto *AO = dyn_cast<GlobalObject>(Val: GA.getAliasee());
69 if (!AO) {
70 LLVM_DEBUG(dbgs() << "Skipping alias whose aliasee is not a GlobalObject: "
71 << GA.getName() << "\n");
72 return false;
73 }
74
75 if (AO->isInterposable()) {
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 SPIRVPrepareGlobals::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