106 lines · cpp
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-exception6//7//===----------------------------------------------------------------------===//8//9// The pass transforms IR globals that cannot be trivially mapped to SPIRV10// into something that is trival to lower.11//12//===----------------------------------------------------------------------===//13 14#include "SPIRV.h"15#include "SPIRVUtils.h"16 17#include "llvm/ADT/STLExtras.h"18#include "llvm/IR/Module.h"19 20using namespace llvm;21 22namespace {23 24struct SPIRVPrepareGlobals : public ModulePass {25 static char ID;26 SPIRVPrepareGlobals() : ModulePass(ID) {}27 28 StringRef getPassName() const override {29 return "SPIRV prepare global variables";30 }31 32 bool runOnModule(Module &M) override;33};34 35bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) {36 assert(Bitcode.getName() == "llvm.embedded.module");37 38 ArrayType *AT = cast<ArrayType>(Bitcode.getValueType());39 if (AT->getNumElements() != 0)40 return false;41 42 ArrayType *AT1 = ArrayType::get(AT->getElementType(), 1);43 Constant *OneEltInit = Constant::getNullValue(AT1);44 Bitcode.replaceInitializer(OneEltInit);45 return true;46}47 48// In HIP, dynamic LDS variables are represented using 0-element global arrays49// in the __shared__ language address-space.50//51// extern __shared__ int LDS[];52//53// These are not representable in SPIRV directly.54// To represent them, for AMD, we use an array with UINT32_MAX-elements.55// These are reverse translated to 0-element arrays.56bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) {57 constexpr unsigned WorkgroupAS =58 storageClassToAddressSpace(SPIRV::StorageClass::Workgroup);59 const bool IsWorkgroupExternal =60 GV.hasExternalLinkage() && GV.getAddressSpace() == WorkgroupAS;61 if (!IsWorkgroupExternal)62 return false;63 64 const ArrayType *AT = dyn_cast<ArrayType>(GV.getValueType());65 if (!AT || AT->getNumElements() != 0)66 return false;67 68 constexpr auto UInt32Max = std::numeric_limits<uint32_t>::max();69 ArrayType *NewAT = ArrayType::get(AT->getElementType(), UInt32Max);70 GlobalVariable *NewGV = new GlobalVariable(71 *GV.getParent(), NewAT, GV.isConstant(), GV.getLinkage(), nullptr, "",72 &GV, GV.getThreadLocalMode(), WorkgroupAS, GV.isExternallyInitialized());73 NewGV->takeName(&GV);74 GV.replaceAllUsesWith(NewGV);75 GV.eraseFromParent();76 77 return true;78}79 80bool SPIRVPrepareGlobals::runOnModule(Module &M) {81 const bool IsAMD = M.getTargetTriple().getVendor() == Triple::AMD;82 if (!IsAMD)83 return false;84 85 bool Changed = false;86 if (GlobalVariable *Bitcode = M.getNamedGlobal("llvm.embedded.module"))87 Changed |= tryExtendLLVMBitcodeMarker(*Bitcode);88 89 for (GlobalVariable &GV : make_early_inc_range(M.globals()))90 Changed |= tryExtendDynamicLDSGlobal(GV);91 92 return Changed;93}94char SPIRVPrepareGlobals::ID = 0;95 96} // namespace97 98INITIALIZE_PASS(SPIRVPrepareGlobals, "prepare-globals",99 "SPIRV prepare global variables", false, false)100 101namespace llvm {102ModulePass *createSPIRVPrepareGlobalsPass() {103 return new SPIRVPrepareGlobals();104}105} // namespace llvm106