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