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