//===- SPIRVReader.cpp - Converts SPIR-V to LLVM ----------------*- C++ -*-===// | |
// | |
// The LLVM/SPIR-V Translator | |
// | |
// This file is distributed under the University of Illinois Open Source | |
// License. See LICENSE.TXT for details. | |
// | |
// Copyright (c) 2014 Advanced Micro Devices, Inc. All rights reserved. | |
// | |
// Permission is hereby granted, free of charge, to any person obtaining a | |
// copy of this software and associated documentation files (the "Software"), | |
// to deal with the Software without restriction, including without limitation | |
// the rights to use, copy, modify, merge, publish, distribute, sublicense, | |
// and/or sell copies of the Software, and to permit persons to whom the | |
// Software is furnished to do so, subject to the following conditions: | |
// | |
// Redistributions of source code must retain the above copyright notice, | |
// this list of conditions and the following disclaimers. | |
// Redistributions in binary form must reproduce the above copyright notice, | |
// this list of conditions and the following disclaimers in the documentation | |
// and/or other materials provided with the distribution. | |
// Neither the names of Advanced Micro Devices, Inc., nor the names of its | |
// contributors may be used to endorse or promote products derived from this | |
// Software without specific prior written permission. | |
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR | |
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, | |
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE | |
// CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER | |
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, | |
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH | |
// THE SOFTWARE. | |
// | |
//===----------------------------------------------------------------------===// | |
/// \file | |
/// | |
/// This file implements conversion of SPIR-V binary to LLVM IR. | |
/// | |
//===----------------------------------------------------------------------===// | |
#include "SPIRVUtil.h" | |
#include "SPIRVType.h" | |
#include "SPIRVValue.h" | |
#include "SPIRVModule.h" | |
#include "SPIRVFunction.h" | |
#include "SPIRVBasicBlock.h" | |
#include "SPIRVInstruction.h" | |
#include "SPIRVExtInst.h" | |
#include "SPIRVInternal.h" | |
#include "SPIRVMDBuilder.h" | |
#include "OCLUtil.h" | |
#include "llvm/ADT/DenseMap.h" | |
#include "llvm/ADT/StringSwitch.h" | |
#include "llvm/IR/Constants.h" | |
#include "llvm/IR/DerivedTypes.h" | |
#include "llvm/IR/DIBuilder.h" | |
#include "llvm/IR/Instructions.h" | |
#include "llvm/IR/Metadata.h" | |
#include "llvm/IR/Module.h" | |
#include "llvm/IR/Operator.h" | |
#include "llvm/IR/Type.h" | |
#include "llvm/IR/LegacyPassManager.h" | |
#include "llvm/Support/Casting.h" | |
#include "llvm/Support/Debug.h" | |
#include "llvm/Support/Dwarf.h" | |
#include "llvm/Support/FileSystem.h" | |
#include "llvm/Support/raw_ostream.h" | |
#include "llvm/Support/CommandLine.h" | |
#include <algorithm> | |
#include <cstdlib> | |
#include <functional> | |
#include <fstream> | |
#include <iostream> | |
#include <iterator> | |
#include <map> | |
#include <set> | |
#include <sstream> | |
#include <string> | |
#define DEBUG_TYPE "spirv" | |
using namespace std; | |
using namespace llvm; | |
using namespace SPIRV; | |
using namespace OCLUtil; | |
namespace SPIRV{ | |
cl::opt<bool> SPIRVEnableStepExpansion("spirv-expand-step", cl::init(true), | |
cl::desc("Enable expansion of OpenCL step and smoothstep function")); | |
cl::opt<bool> SPIRVGenKernelArgNameMD("spirv-gen-kernel-arg-name-md", | |
cl::init(false), cl::desc("Enable generating OpenCL kernel argument name " | |
"metadata")); | |
cl::opt<bool> SPIRVGenImgTypeAccQualPostfix("spirv-gen-image-type-acc-postfix", | |
cl::init(false), cl::desc("Enable generating access qualifier postfix" | |
" in OpenCL image type names")); | |
// Prefix for placeholder global variable name. | |
const char* kPlaceholderPrefix = "placeholder."; | |
// Save the translated LLVM before validation for debugging purpose. | |
static bool DbgSaveTmpLLVM = true; | |
static const char *DbgTmpLLVMFileName = "_tmp_llvmbil.ll"; | |
typedef std::pair < unsigned, AttributeSet > AttributeWithIndex; | |
static bool | |
isOpenCLKernel(SPIRVFunction *BF) { | |
return BF->getModule()->isEntryPoint(ExecutionModelKernel, BF->getId()); | |
} | |
static void | |
dumpLLVM(Module *M, const std::string &FName) { | |
std::error_code EC; | |
raw_fd_ostream FS(FName, EC, sys::fs::F_None); | |
if (EC) { | |
FS << *M; | |
FS.close(); | |
} | |
} | |
static MDNode* | |
getMDNodeStringIntVec(LLVMContext *Context, const std::string& Str, | |
const std::vector<SPIRVWord>& IntVals) { | |
std::vector<Metadata*> ValueVec; | |
ValueVec.push_back(MDString::get(*Context, Str)); | |
for (auto &I:IntVals) | |
ValueVec.push_back(ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), I))); | |
return MDNode::get(*Context, ValueVec); | |
} | |
static MDNode* | |
getMDTwoInt(LLVMContext *Context, unsigned Int1, unsigned Int2) { | |
std::vector<Metadata*> ValueVec; | |
ValueVec.push_back(ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), Int1))); | |
ValueVec.push_back(ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), Int2))); | |
return MDNode::get(*Context, ValueVec); | |
} | |
static MDNode* | |
getMDString(LLVMContext *Context, const std::string& Str) { | |
std::vector<Metadata*> ValueVec; | |
if (!Str.empty()) | |
ValueVec.push_back(MDString::get(*Context, Str)); | |
return MDNode::get(*Context, ValueVec); | |
} | |
static void | |
addOCLVersionMetadata(LLVMContext *Context, Module *M, | |
const std::string &MDName, unsigned Major, unsigned Minor) { | |
NamedMDNode *NamedMD = M->getOrInsertNamedMetadata(MDName); | |
NamedMD->addOperand(getMDTwoInt(Context, Major, Minor)); | |
} | |
static void | |
addNamedMetadataStringSet(LLVMContext *Context, Module *M, | |
const std::string &MDName, const std::set<std::string> &StrSet) { | |
NamedMDNode *NamedMD = M->getOrInsertNamedMetadata(MDName); | |
std::vector<Metadata*> ValueVec; | |
for (auto &&Str : StrSet) { | |
ValueVec.push_back(MDString::get(*Context, Str)); | |
} | |
NamedMD->addOperand(MDNode::get(*Context, ValueVec)); | |
} | |
static void | |
addOCLKernelArgumentMetadata(LLVMContext *Context, | |
std::vector<llvm::Metadata*> &KernelMD, const std::string &MDName, | |
SPIRVFunction *BF, std::function<Metadata *(SPIRVFunctionParameter *)>Func){ | |
std::vector<Metadata*> ValueVec; | |
ValueVec.push_back(MDString::get(*Context, MDName)); | |
BF->foreachArgument([&](SPIRVFunctionParameter *Arg) { | |
ValueVec.push_back(Func(Arg)); | |
}); | |
KernelMD.push_back(MDNode::get(*Context, ValueVec)); | |
} | |
class SPIRVToLLVMDbgTran { | |
public: | |
SPIRVToLLVMDbgTran(SPIRVModule *TBM, Module *TM) | |
:BM(TBM), M(TM), SpDbg(BM), Builder(*M){ | |
Enable = BM->hasDebugInfo(); | |
} | |
void createCompileUnit() { | |
if (!Enable) | |
return; | |
auto File = SpDbg.getEntryPointFileStr(ExecutionModelKernel, 0); | |
std::string BaseName; | |
std::string Path; | |
splitFileName(File, BaseName, Path); | |
Builder.createCompileUnit(dwarf::DW_LANG_C99, | |
BaseName, Path, "spirv", false, "", 0, "", DIBuilder::LineTablesOnly); | |
} | |
void addDbgInfoVersion() { | |
if (!Enable) | |
return; | |
M->addModuleFlag(Module::Warning, "Dwarf Version", | |
dwarf::DWARF_VERSION); | |
M->addModuleFlag(Module::Warning, "Debug Info Version", | |
DEBUG_METADATA_VERSION); | |
} | |
DIFile* getDIFile(const std::string &FileName){ | |
return getOrInsert(FileMap, FileName, [=](){ | |
std::string BaseName; | |
std::string Path; | |
splitFileName(FileName, BaseName, Path); | |
if (!BaseName.empty()) | |
return Builder.createFile(BaseName, Path); | |
else | |
return Builder.createFile("","");//DIFile(); | |
}); | |
} | |
DISubprogram* getDISubprogram(SPIRVFunction *SF, Function *F){ | |
return getOrInsert(FuncMap, F, [=](){ | |
auto DF = getDIFile(SpDbg.getFunctionFileStr(SF)); | |
auto FN = F->getName(); | |
auto LN = SpDbg.getFunctionLineNo(SF); | |
Metadata *Args[] = {Builder.createUnspecifiedType("")}; | |
return Builder.createFunction(static_cast<DIScope*>(DF), FN, FN, DF, LN, | |
Builder.createSubroutineType(Builder.getOrCreateTypeArray(Args)), | |
Function::isInternalLinkage(F->getLinkage()), | |
true, LN); | |
}); | |
} | |
void transDbgInfo(SPIRVValue *SV, Value *V) { | |
if (!Enable || !SV->hasLine()) | |
return; | |
if (auto I = dyn_cast<Instruction>(V)) { | |
assert(SV->isInst() && "Invalid instruction"); | |
auto SI = static_cast<SPIRVInstruction *>(SV); | |
assert(SI->getParent() && | |
SI->getParent()->getParent() && | |
"Invalid instruction"); | |
auto Line = SV->getLine(); | |
I->setDebugLoc(DebugLoc::get(Line->getLine(), Line->getColumn(), | |
getDISubprogram(SI->getParent()->getParent(), | |
I->getParent()->getParent()))); | |
} | |
} | |
void finalize() { | |
if (!Enable) | |
return; | |
Builder.finalize(); | |
} | |
private: | |
SPIRVModule *BM; | |
Module *M; | |
SPIRVDbgInfo SpDbg; | |
DIBuilder Builder; | |
bool Enable; | |
std::unordered_map<std::string, DIFile*> FileMap; | |
std::unordered_map<Function *, DISubprogram*> FuncMap; | |
void splitFileName(const std::string &FileName, | |
std::string &BaseName, | |
std::string &Path) { | |
auto Loc = FileName.find_last_of("/\\"); | |
if (Loc != std::string::npos) { | |
BaseName = FileName.substr(Loc + 1); | |
Path = FileName.substr(0, Loc); | |
} else { | |
BaseName = FileName; | |
Path = "."; | |
} | |
} | |
}; | |
class SPIRVToLLVM { | |
public: | |
SPIRVToLLVM(Module *LLVMModule, SPIRVModule *TheSPIRVModule) | |
:M(LLVMModule), BM(TheSPIRVModule), DbgTran(BM, M){ | |
assert(M); | |
Context = &M->getContext(); | |
} | |
std::string getOCLBuiltinName(SPIRVInstruction* BI); | |
std::string getOCLConvertBuiltinName(SPIRVInstruction *BI); | |
std::string getOCLGenericCastToPtrName(SPIRVInstruction *BI); | |
Type *transType(SPIRVType *BT, bool IsClassMember = false); | |
std::string transTypeToOCLTypeName(SPIRVType *BT, bool IsSigned = true); | |
std::vector<Type *> transTypeVector(const std::vector<SPIRVType *>&); | |
bool translate(); | |
bool transAddressingModel(); | |
Value *transValue(SPIRVValue *, Function *F, BasicBlock *, | |
bool CreatePlaceHolder = true); | |
Value *transValueWithoutDecoration(SPIRVValue *, Function *F, BasicBlock *, | |
bool CreatePlaceHolder = true); | |
bool transDecoration(SPIRVValue *, Value *); | |
bool transAlign(SPIRVValue *, Value *); | |
Instruction *transOCLBuiltinFromExtInst(SPIRVExtInst *BC, BasicBlock *BB); | |
std::vector<Value *> transValue(const std::vector<SPIRVValue *>&, Function *F, | |
BasicBlock *); | |
Function *transFunction(SPIRVFunction *F); | |
bool transFPContractMetadata(); | |
bool transKernelMetadata(); | |
bool transNonTemporalMetadata(Instruction *I); | |
bool transSourceLanguage(); | |
bool transSourceExtension(); | |
void transGeneratorMD(); | |
Value *transConvertInst(SPIRVValue* BV, Function* F, BasicBlock* BB); | |
Instruction *transBuiltinFromInst(const std::string& FuncName, | |
SPIRVInstruction* BI, BasicBlock* BB); | |
Instruction *transOCLBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB); | |
Instruction *transSPIRVBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB); | |
Instruction *transOCLBarrierFence(SPIRVInstruction* BI, BasicBlock *BB); | |
void transOCLVectorLoadStore(std::string& UnmangledName, | |
std::vector<SPIRVWord> &BArgs); | |
/// Post-process translated LLVM module for OpenCL. | |
bool postProcessOCL(); | |
/// \brief Post-process OpenCL builtin functions returning struct type. | |
/// | |
/// Some OpenCL builtin functions are translated to SPIR-V instructions with | |
/// struct type result, e.g. NDRange creation functions. Such functions | |
/// need to be post-processed to return the struct through sret argument. | |
bool postProcessOCLBuiltinReturnStruct(Function *F); | |
/// \brief Post-process OpenCL builtin functions having block argument. | |
/// | |
/// These functions are translated to functions with function pointer type | |
/// argument first, then post-processed to have block argument. | |
bool postProcessOCLBuiltinWithFuncPointer(Function *F, | |
Function::arg_iterator I); | |
/// \brief Post-process OpenCL builtin functions having array argument. | |
/// | |
/// These functions are translated to functions with array type argument | |
/// first, then post-processed to have pointer arguments. | |
bool postProcessOCLBuiltinWithArrayArguments(Function *F, | |
const std::string &DemangledName); | |
/// \brief Post-process OpImageSampleExplicitLod. | |
/// sampled_image = __spirv_SampledImage__(image, sampler); | |
/// return __spirv_ImageSampleExplicitLod__(sampled_image, image_operands, | |
/// ...); | |
/// => | |
/// read_image(image, sampler, ...) | |
/// \return transformed call instruction. | |
Instruction *postProcessOCLReadImage(SPIRVInstruction *BI, CallInst *CI, | |
const std::string &DemangledName); | |
/// \brief Post-process OpImageWrite. | |
/// return write_image(image, coord, color, image_operands, ...); | |
/// => | |
/// write_image(image, coord, ..., color) | |
/// \return transformed call instruction. | |
CallInst *postProcessOCLWriteImage(SPIRVInstruction *BI, CallInst *CI, | |
const std::string &DemangledName); | |
/// \brief Post-process OpBuildNDRange. | |
/// OpBuildNDRange GlobalWorkSize, LocalWorkSize, GlobalWorkOffset | |
/// => | |
/// call ndrange_XD(GlobalWorkOffset, GlobalWorkSize, LocalWorkSize) | |
/// \return transformed call instruction. | |
CallInst *postProcessOCLBuildNDRange(SPIRVInstruction *BI, CallInst *CI, | |
const std::string &DemangledName); | |
/// \brief Expand OCL builtin functions with scalar argument, e.g. | |
/// step, smoothstep. | |
/// gentype func (fp edge, gentype x) | |
/// => | |
/// gentype func (gentype edge, gentype x) | |
/// \return transformed call instruction. | |
CallInst *expandOCLBuiltinWithScalarArg(CallInst* CI, | |
const std::string &FuncName); | |
/// \brief Post-process OpGroupAll and OpGroupAny instructions translation. | |
/// i1 func (<n x i1> arg) | |
/// => | |
/// i32 func (<n x i32> arg) | |
/// \return transformed call instruction. | |
Instruction *postProcessGroupAllAny(CallInst *CI, | |
const std::string &DemangledName); | |
typedef DenseMap<SPIRVType *, Type *> SPIRVToLLVMTypeMap; | |
typedef DenseMap<SPIRVValue *, Value *> SPIRVToLLVMValueMap; | |
typedef DenseMap<SPIRVFunction *, Function *> SPIRVToLLVMFunctionMap; | |
typedef DenseMap<GlobalVariable *, SPIRVBuiltinVariableKind> BuiltinVarMap; | |
// A SPIRV value may be translated to a load instruction of a placeholder | |
// global variable. This map records load instruction of these placeholders | |
// which are supposed to be replaced by the real values later. | |
typedef std::map<SPIRVValue *, LoadInst*> SPIRVToLLVMPlaceholderMap; | |
private: | |
Module *M; | |
BuiltinVarMap BuiltinGVMap; | |
LLVMContext *Context; | |
SPIRVModule *BM; | |
SPIRVToLLVMTypeMap TypeMap; | |
SPIRVToLLVMValueMap ValueMap; | |
SPIRVToLLVMFunctionMap FuncMap; | |
SPIRVToLLVMPlaceholderMap PlaceholderMap; | |
SPIRVToLLVMDbgTran DbgTran; | |
Type *mapType(SPIRVType *BT, Type *T) { | |
SPIRVDBG(dbgs() << *T << '\n';) | |
TypeMap[BT] = T; | |
return T; | |
} | |
// If a value is mapped twice, the existing mapped value is a placeholder, | |
// which must be a load instruction of a global variable whose name starts | |
// with kPlaceholderPrefix. | |
Value *mapValue(SPIRVValue *BV, Value *V) { | |
auto Loc = ValueMap.find(BV); | |
if (Loc != ValueMap.end()) { | |
if (Loc->second == V) | |
return V; | |
auto LD = dyn_cast<LoadInst>(Loc->second); | |
auto Placeholder = dyn_cast<GlobalVariable>(LD->getPointerOperand()); | |
assert (LD && Placeholder && | |
Placeholder->getName().startswith(kPlaceholderPrefix) && | |
"A value is translated twice"); | |
// Replaces placeholders for PHI nodes | |
LD->replaceAllUsesWith(V); | |
LD->dropAllReferences(); | |
LD->removeFromParent(); | |
Placeholder->dropAllReferences(); | |
Placeholder->removeFromParent(); | |
} | |
ValueMap[BV] = V; | |
return V; | |
} | |
bool isSPIRVBuiltinVariable(GlobalVariable *GV, | |
SPIRVBuiltinVariableKind *Kind = nullptr) { | |
auto Loc = BuiltinGVMap.find(GV); | |
if (Loc == BuiltinGVMap.end()) | |
return false; | |
if (Kind) | |
*Kind = Loc->second; | |
return true; | |
} | |
// OpenCL function always has NoUnwound attribute. | |
// Change this if it is no longer true. | |
bool isFuncNoUnwind() const { return true;} | |
bool isSPIRVCmpInstTransToLLVMInst(SPIRVInstruction *BI) const; | |
bool transOCLBuiltinsFromVariables(); | |
bool transOCLBuiltinFromVariable(GlobalVariable *GV, | |
SPIRVBuiltinVariableKind Kind); | |
MDString *transOCLKernelArgTypeName(SPIRVFunctionParameter *); | |
Value *mapFunction(SPIRVFunction *BF, Function *F) { | |
SPIRVDBG(spvdbgs() << "[mapFunction] " << *BF << " -> "; | |
dbgs() << *F << '\n';) | |
FuncMap[BF] = F; | |
return F; | |
} | |
Value *getTranslatedValue(SPIRVValue *BV); | |
Type *getTranslatedType(SPIRVType *BT); | |
SPIRVErrorLog &getErrorLog() { | |
return BM->getErrorLog(); | |
} | |
void setCallingConv(CallInst *Call) { | |
Function *F = Call->getCalledFunction(); | |
assert(F); | |
Call->setCallingConv(F->getCallingConv()); | |
} | |
void setAttrByCalledFunc(CallInst *Call); | |
Type *transFPType(SPIRVType* T); | |
BinaryOperator *transShiftLogicalBitwiseInst(SPIRVValue* BV, BasicBlock* BB, | |
Function* F); | |
void transFlags(llvm::Value* V); | |
Instruction *transCmpInst(SPIRVValue* BV, BasicBlock* BB, Function* F); | |
void transOCLBuiltinFromInstPreproc(SPIRVInstruction* BI, Type *&RetTy, | |
std::vector<SPIRVValue *> &Args); | |
Instruction* transOCLBuiltinPostproc(SPIRVInstruction* BI, | |
CallInst* CI, BasicBlock* BB, const std::string &DemangledName); | |
std::string transOCLImageTypeName(SPIRV::SPIRVTypeImage* ST); | |
std::string transOCLSampledImageTypeName(SPIRV::SPIRVTypeSampledImage* ST); | |
std::string transOCLPipeTypeName(SPIRV::SPIRVTypePipe* ST, | |
bool UseSPIRVFriendlyFormat = false, int PipeAccess = 0); | |
std::string transOCLPipeStorageTypeName(SPIRV::SPIRVTypePipeStorage* PST); | |
std::string transOCLImageTypeAccessQualifier(SPIRV::SPIRVTypeImage* ST); | |
std::string transOCLPipeTypeAccessQualifier(SPIRV::SPIRVTypePipe* ST); | |
Value *oclTransConstantSampler(SPIRV::SPIRVConstantSampler* BCS); | |
Value * oclTransConstantPipeStorage(SPIRV::SPIRVConstantPipeStorage* BCPS); | |
void setName(llvm::Value* V, SPIRVValue* BV); | |
void insertImageNameAccessQualifier(SPIRV::SPIRVTypeImage* ST, std::string &Name); | |
template<class Source, class Func> | |
bool foreachFuncCtlMask(Source, Func); | |
llvm::GlobalValue::LinkageTypes transLinkageType(const SPIRVValue* V); | |
Instruction *transOCLAllAny(SPIRVInstruction* BI, BasicBlock *BB); | |
Instruction *transOCLRelational(SPIRVInstruction* BI, BasicBlock *BB); | |
CallInst *transOCLBarrier(BasicBlock *BB, SPIRVWord ExecScope, | |
SPIRVWord MemSema, SPIRVWord MemScope); | |
CallInst *transOCLMemFence(BasicBlock *BB, | |
SPIRVWord MemSema, SPIRVWord MemScope); | |
}; | |
Type * | |
SPIRVToLLVM::getTranslatedType(SPIRVType *BV){ | |
auto Loc = TypeMap.find(BV); | |
if (Loc != TypeMap.end()) | |
return Loc->second; | |
return nullptr; | |
} | |
Value * | |
SPIRVToLLVM::getTranslatedValue(SPIRVValue *BV){ | |
auto Loc = ValueMap.find(BV); | |
if (Loc != ValueMap.end()) | |
return Loc->second; | |
return nullptr; | |
} | |
void | |
SPIRVToLLVM::setAttrByCalledFunc(CallInst *Call) { | |
Function *F = Call->getCalledFunction(); | |
assert(F); | |
if (F->isIntrinsic()) { | |
return; | |
} | |
Call->setCallingConv(F->getCallingConv()); | |
Call->setAttributes(F->getAttributes()); | |
} | |
bool | |
SPIRVToLLVM::transOCLBuiltinsFromVariables(){ | |
std::vector<GlobalVariable *> WorkList; | |
for (auto I = M->global_begin(), E = M->global_end(); I != E; ++I) { | |
SPIRVBuiltinVariableKind Kind; | |
auto I1 = static_cast<GlobalVariable*>(I); | |
if (!isSPIRVBuiltinVariable(I1, &Kind)) | |
continue; | |
if (!transOCLBuiltinFromVariable(I1, Kind)) | |
return false; | |
WorkList.push_back(I1); | |
} | |
for (auto &I:WorkList) { | |
I->dropAllReferences(); | |
I->removeFromParent(); | |
} | |
return true; | |
} | |
// For integer types shorter than 32 bit, unsigned/signedness can be inferred | |
// from zext/sext attribute. | |
MDString * | |
SPIRVToLLVM::transOCLKernelArgTypeName(SPIRVFunctionParameter *Arg) { | |
auto Ty = Arg->isByVal() ? Arg->getType()->getPointerElementType() : | |
Arg->getType(); | |
return MDString::get(*Context, transTypeToOCLTypeName(Ty, !Arg->isZext())); | |
} | |
// Variable like GlobalInvolcationId[x] -> get_global_id(x). | |
// Variable like WorkDim -> get_work_dim(). | |
bool | |
SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV, | |
SPIRVBuiltinVariableKind Kind) { | |
std::string FuncName = SPIRSPIRVBuiltinVariableMap::rmap(Kind); | |
std::string MangledName; | |
Type *ReturnTy = GV->getType()->getPointerElementType(); | |
bool IsVec = ReturnTy->isVectorTy(); | |
if (IsVec) | |
ReturnTy = cast<VectorType>(ReturnTy)->getElementType(); | |
std::vector<Type*> ArgTy; | |
if (IsVec) | |
ArgTy.push_back(Type::getInt32Ty(*Context)); | |
MangleOpenCLBuiltin(FuncName, ArgTy, MangledName); | |
Function *Func = M->getFunction(MangledName); | |
if (!Func) { | |
FunctionType *FT = FunctionType::get(ReturnTy, ArgTy, false); | |
Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); | |
Func->setCallingConv(CallingConv::SPIR_FUNC); | |
Func->addFnAttr(Attribute::NoUnwind); | |
Func->addFnAttr(Attribute::ReadNone); | |
} | |
std::vector<Instruction *> Deletes; | |
std::vector<Instruction *> Uses; | |
for (auto UI = GV->user_begin(), UE = GV->user_end(); UI != UE; ++UI) { | |
assert (isa<LoadInst>(*UI) && "Unsupported use"); | |
auto LD = dyn_cast<LoadInst>(*UI); | |
if (!IsVec) { | |
Uses.push_back(LD); | |
Deletes.push_back(LD); | |
continue; | |
} | |
for (auto LDUI = LD->user_begin(), LDUE = LD->user_end(); LDUI != LDUE; | |
++LDUI) { | |
assert(isa<ExtractElementInst>(*LDUI) && "Unsupported use"); | |
auto EEI = dyn_cast<ExtractElementInst>(*LDUI); | |
Uses.push_back(EEI); | |
Deletes.push_back(EEI); | |
} | |
Deletes.push_back(LD); | |
} | |
for (auto &I:Uses) { | |
std::vector<Value *> Arg; | |
if (auto EEI = dyn_cast<ExtractElementInst>(I)) | |
Arg.push_back(EEI->getIndexOperand()); | |
auto Call = CallInst::Create(Func, Arg, "", I); | |
Call->takeName(I); | |
setAttrByCalledFunc(Call); | |
SPIRVDBG(dbgs() << "[transOCLBuiltinFromVariable] " << *I << " -> " << | |
*Call << '\n';) | |
I->replaceAllUsesWith(Call); | |
} | |
for (auto &I:Deletes) { | |
I->dropAllReferences(); | |
I->removeFromParent(); | |
} | |
return true; | |
} | |
Type * | |
SPIRVToLLVM::transFPType(SPIRVType* T) { | |
switch(T->getFloatBitWidth()) { | |
case 16: return Type::getHalfTy(*Context); | |
case 32: return Type::getFloatTy(*Context); | |
case 64: return Type::getDoubleTy(*Context); | |
default: | |
llvm_unreachable("Invalid type"); | |
return nullptr; | |
} | |
} | |
std::string | |
SPIRVToLLVM::transOCLImageTypeName(SPIRV::SPIRVTypeImage* ST) { | |
std::string Name = std::string(kSPR2TypeName::OCLPrefix) | |
+ rmap<std::string>(ST->getDescriptor()); | |
if (SPIRVGenImgTypeAccQualPostfix) | |
SPIRVToLLVM::insertImageNameAccessQualifier(ST, Name); | |
return std::move(Name); | |
} | |
std::string | |
SPIRVToLLVM::transOCLSampledImageTypeName(SPIRV::SPIRVTypeSampledImage* ST) { | |
return getSPIRVTypeName(kSPIRVTypeName::SampledImg, | |
getSPIRVImageTypePostfixes(getSPIRVImageSampledTypeName( | |
ST->getImageType()->getSampledType()), | |
ST->getImageType()->getDescriptor(), | |
ST->getImageType()->getAccessQualifier())); | |
} | |
std::string | |
SPIRVToLLVM::transOCLPipeTypeName(SPIRV::SPIRVTypePipe* PT, | |
bool UseSPIRVFriendlyFormat, int PipeAccess){ | |
if (!UseSPIRVFriendlyFormat) | |
return kSPR2TypeName::Pipe; | |
else | |
return std::string(kSPIRVTypeName::PrefixAndDelim) | |
+ kSPIRVTypeName::Pipe | |
+ kSPIRVTypeName::Delimiter | |
+ kSPIRVTypeName::PostfixDelim | |
+ PipeAccess; | |
} | |
std::string | |
SPIRVToLLVM::transOCLPipeStorageTypeName(SPIRV::SPIRVTypePipeStorage* PST) { | |
return std::string(kSPIRVTypeName::PrefixAndDelim) | |
+ kSPIRVTypeName::PipeStorage; | |
} | |
Type * | |
SPIRVToLLVM::transType(SPIRVType *T, bool IsClassMember) { | |
auto Loc = TypeMap.find(T); | |
if (Loc != TypeMap.end()) | |
return Loc->second; | |
SPIRVDBG(spvdbgs() << "[transType] " << *T << " -> ";) | |
T->validate(); | |
switch(T->getOpCode()) { | |
case OpTypeVoid: | |
return mapType(T, Type::getVoidTy(*Context)); | |
case OpTypeBool: | |
return mapType(T, Type::getInt1Ty(*Context)); | |
case OpTypeInt: | |
return mapType(T, Type::getIntNTy(*Context, T->getIntegerBitWidth())); | |
case OpTypeFloat: | |
return mapType(T, transFPType(T)); | |
case OpTypeArray: | |
return mapType(T, ArrayType::get(transType(T->getArrayElementType()), | |
T->getArrayLength())); | |
case OpTypePointer: | |
return mapType(T, PointerType::get(transType( | |
T->getPointerElementType(), IsClassMember), | |
SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass()))); | |
case OpTypeVector: | |
return mapType(T, VectorType::get(transType(T->getVectorComponentType()), | |
T->getVectorComponentCount())); | |
case OpTypeOpaque: | |
return mapType(T, StructType::create(*Context, T->getName())); | |
case OpTypeFunction: { | |
auto FT = static_cast<SPIRVTypeFunction *>(T); | |
auto RT = transType(FT->getReturnType()); | |
std::vector<Type *> PT; | |
for (size_t I = 0, E = FT->getNumParameters(); I != E; ++I) | |
PT.push_back(transType(FT->getParameterType(I))); | |
return mapType(T, FunctionType::get(RT, PT, false)); | |
} | |
case OpTypeImage: { | |
auto ST = static_cast<SPIRVTypeImage *>(T); | |
if (ST->isOCLImage()) | |
return mapType(T, getOrCreateOpaquePtrType(M, | |
transOCLImageTypeName(ST))); | |
else | |
llvm_unreachable("Unsupported image type"); | |
return nullptr; | |
} | |
case OpTypeSampler: | |
return mapType(T, Type::getInt32Ty(*Context)); | |
case OpTypeSampledImage: { | |
auto ST = static_cast<SPIRVTypeSampledImage *>(T); | |
return mapType(T, getOrCreateOpaquePtrType(M, | |
transOCLSampledImageTypeName(ST))); | |
} | |
case OpTypeStruct: { | |
auto ST = static_cast<SPIRVTypeStruct *>(T); | |
auto Name = ST->getName(); | |
if (!Name.empty()) { | |
if (auto OldST = M->getTypeByName(Name)) | |
OldST->setName(""); | |
} | |
auto *StructTy = StructType::create(*Context, Name); | |
mapType(ST, StructTy); | |
SmallVector<Type *, 4> MT; | |
for (size_t I = 0, E = ST->getMemberCount(); I != E; ++I) | |
MT.push_back(transType(ST->getMemberType(I), true)); | |
StructTy->setBody(MT, ST->isPacked()); | |
return StructTy; | |
} | |
case OpTypePipe: { | |
auto PT = static_cast<SPIRVTypePipe *>(T); | |
return mapType(T, getOrCreateOpaquePtrType(M, | |
transOCLPipeTypeName(PT, IsClassMember, PT->getAccessQualifier()), | |
getOCLOpaqueTypeAddrSpace(T->getOpCode()))); | |
} | |
case OpTypePipeStorage: { | |
auto PST = static_cast<SPIRVTypePipeStorage *>(T); | |
return mapType(T, getOrCreateOpaquePtrType(M, | |
transOCLPipeStorageTypeName(PST), | |
getOCLOpaqueTypeAddrSpace(T->getOpCode()))); | |
} | |
default: { | |
auto OC = T->getOpCode(); | |
if (isOpaqueGenericTypeOpCode(OC)) | |
return mapType(T, getOrCreateOpaquePtrType(M, | |
OCLOpaqueTypeOpCodeMap::rmap(OC), | |
getOCLOpaqueTypeAddrSpace(OC))); | |
llvm_unreachable("Not implemented"); | |
} | |
} | |
return 0; | |
} | |
std::string | |
SPIRVToLLVM::transTypeToOCLTypeName(SPIRVType *T, bool IsSigned) { | |
switch(T->getOpCode()) { | |
case OpTypeVoid: | |
return "void"; | |
case OpTypeBool: | |
return "bool"; | |
case OpTypeInt: { | |
std::string Prefix = IsSigned ? "" : "u"; | |
switch(T->getIntegerBitWidth()) { | |
case 8: | |
return Prefix + "char"; | |
case 16: | |
return Prefix + "short"; | |
case 32: | |
return Prefix + "int"; | |
case 64: | |
return Prefix + "long"; | |
default: | |
llvm_unreachable("invalid integer size"); | |
return Prefix + std::string("int") + T->getIntegerBitWidth() + "_t"; | |
} | |
} | |
break; | |
case OpTypeFloat: | |
switch(T->getFloatBitWidth()){ | |
case 16: | |
return "half"; | |
case 32: | |
return "float"; | |
case 64: | |
return "double"; | |
default: | |
llvm_unreachable("invalid floating pointer bitwidth"); | |
return std::string("float") + T->getFloatBitWidth() + "_t"; | |
} | |
break; | |
case OpTypeArray: | |
return "array"; | |
case OpTypePointer: | |
return transTypeToOCLTypeName(T->getPointerElementType()) + "*"; | |
case OpTypeVector: | |
return transTypeToOCLTypeName(T->getVectorComponentType()) + | |
T->getVectorComponentCount(); | |
case OpTypeOpaque: | |
return T->getName(); | |
case OpTypeFunction: | |
llvm_unreachable("Unsupported"); | |
return "function"; | |
case OpTypeStruct: { | |
auto Name = T->getName(); | |
if (Name.find("struct.") == 0) | |
Name[6] = ' '; | |
else if (Name.find("union.") == 0) | |
Name[5] = ' '; | |
return Name; | |
} | |
case OpTypePipe: | |
return "pipe"; | |
case OpTypeSampler: | |
return "sampler_t"; | |
case OpTypeImage: { | |
std::string Name; | |
Name = rmap<std::string>(static_cast<SPIRVTypeImage *>(T)->getDescriptor()); | |
if (SPIRVGenImgTypeAccQualPostfix) { | |
auto ST = static_cast<SPIRVTypeImage *>(T); | |
insertImageNameAccessQualifier(ST, Name); | |
} | |
return Name; | |
} | |
default: | |
if (isOpaqueGenericTypeOpCode(T->getOpCode())) { | |
return OCLOpaqueTypeOpCodeMap::rmap(T->getOpCode()); | |
} | |
llvm_unreachable("Not implemented"); | |
return "unknown"; | |
} | |
} | |
std::vector<Type *> | |
SPIRVToLLVM::transTypeVector(const std::vector<SPIRVType *> &BT) { | |
std::vector<Type *> T; | |
for (auto I: BT) | |
T.push_back(transType(I)); | |
return T; | |
} | |
std::vector<Value *> | |
SPIRVToLLVM::transValue(const std::vector<SPIRVValue *> &BV, Function *F, | |
BasicBlock *BB) { | |
std::vector<Value *> V; | |
for (auto I: BV) | |
V.push_back(transValue(I, F, BB)); | |
return V; | |
} | |
bool | |
SPIRVToLLVM::isSPIRVCmpInstTransToLLVMInst(SPIRVInstruction* BI) const { | |
auto OC = BI->getOpCode(); | |
return isCmpOpCode(OC) && | |
!(OC >= OpLessOrGreater && OC <= OpUnordered); | |
} | |
void | |
SPIRVToLLVM::transFlags(llvm::Value* V) { | |
if(!isa<Instruction>(V)) | |
return; | |
auto OC = cast<Instruction>(V)->getOpcode(); | |
if (OC == Instruction::AShr || OC == Instruction::LShr) { | |
cast<BinaryOperator>(V)->setIsExact(); | |
return; | |
} | |
} | |
void | |
SPIRVToLLVM::setName(llvm::Value* V, SPIRVValue* BV) { | |
auto Name = BV->getName(); | |
if (!Name.empty() && (!V->hasName() || Name != V->getName())) | |
V->setName(Name); | |
} | |
void SPIRVToLLVM::insertImageNameAccessQualifier(SPIRV::SPIRVTypeImage* ST, std::string &Name) { | |
std::string QName = rmap<std::string>(ST->getAccessQualifier()); | |
// transform: read_only -> ro, write_only -> wo, read_write -> rw | |
QName = QName.substr(0,1) + QName.substr(QName.find("_") + 1, 1) + "_"; | |
assert(!Name.empty() && "image name should not be empty"); | |
Name.insert(Name.size() - 1, QName); | |
} | |
Value * | |
SPIRVToLLVM::transValue(SPIRVValue *BV, Function *F, BasicBlock *BB, | |
bool CreatePlaceHolder){ | |
SPIRVToLLVMValueMap::iterator Loc = ValueMap.find(BV); | |
if (Loc != ValueMap.end() && (!PlaceholderMap.count(BV) || CreatePlaceHolder)) | |
return Loc->second; | |
SPIRVDBG(spvdbgs() << "[transValue] " << *BV << " -> ";) | |
BV->validate(); | |
auto V = transValueWithoutDecoration(BV, F, BB, CreatePlaceHolder); | |
if (!V) { | |
SPIRVDBG(dbgs() << " Warning ! nullptr\n";) | |
return nullptr; | |
} | |
setName(V, BV); | |
if (!transDecoration(BV, V)) { | |
assert (0 && "trans decoration fail"); | |
return nullptr; | |
} | |
transFlags(V); | |
SPIRVDBG(dbgs() << *V << '\n';) | |
return V; | |
} | |
Value * | |
SPIRVToLLVM::transConvertInst(SPIRVValue* BV, Function* F, BasicBlock* BB) { | |
SPIRVUnary* BC = static_cast<SPIRVUnary*>(BV); | |
auto Src = transValue(BC->getOperand(0), F, BB, BB ? true : false); | |
auto Dst = transType(BC->getType()); | |
CastInst::CastOps CO = Instruction::BitCast; | |
bool IsExt = Dst->getScalarSizeInBits() | |
> Src->getType()->getScalarSizeInBits(); | |
switch (BC->getOpCode()) { | |
case OpPtrCastToGeneric: | |
case OpGenericCastToPtr: | |
CO = Instruction::AddrSpaceCast; | |
break; | |
case OpSConvert: | |
CO = IsExt ? Instruction::SExt : Instruction::Trunc; | |
break; | |
case OpUConvert: | |
CO = IsExt ? Instruction::ZExt : Instruction::Trunc; | |
break; | |
case OpFConvert: | |
CO = IsExt ? Instruction::FPExt : Instruction::FPTrunc; | |
break; | |
default: | |
CO = static_cast<CastInst::CastOps>(OpCodeMap::rmap(BC->getOpCode())); | |
} | |
assert(CastInst::isCast(CO) && "Invalid cast op code"); | |
SPIRVDBG(if (!CastInst::castIsValid(CO, Src, Dst)) { | |
spvdbgs() << "Invalid cast: " << *BV << " -> "; | |
dbgs() << "Op = " << CO << ", Src = " << *Src << " Dst = " << *Dst << '\n'; | |
}) | |
if (BB) | |
return CastInst::Create(CO, Src, Dst, BV->getName(), BB); | |
return ConstantExpr::getCast(CO, dyn_cast<Constant>(Src), Dst); | |
} | |
BinaryOperator *SPIRVToLLVM::transShiftLogicalBitwiseInst(SPIRVValue* BV, | |
BasicBlock* BB,Function* F) { | |
SPIRVBinary* BBN = static_cast<SPIRVBinary*>(BV); | |
assert(BB && "Invalid BB"); | |
Instruction::BinaryOps BO; | |
auto OP = BBN->getOpCode(); | |
if (isLogicalOpCode(OP)) | |
OP = IntBoolOpMap::rmap(OP); | |
BO = static_cast<Instruction::BinaryOps>(OpCodeMap::rmap(OP)); | |
auto Inst = BinaryOperator::Create(BO, | |
transValue(BBN->getOperand(0), F, BB), | |
transValue(BBN->getOperand(1), F, BB), BV->getName(), BB); | |
return Inst; | |
} | |
Instruction * | |
SPIRVToLLVM::transCmpInst(SPIRVValue* BV, BasicBlock* BB, Function* F) { | |
SPIRVCompare* BC = static_cast<SPIRVCompare*>(BV); | |
assert(BB && "Invalid BB"); | |
SPIRVType* BT = BC->getOperand(0)->getType(); | |
Instruction* Inst = nullptr; | |
auto OP = BC->getOpCode(); | |
if (isLogicalOpCode(OP)) | |
OP = IntBoolOpMap::rmap(OP); | |
if (BT->isTypeVectorOrScalarInt() || BT->isTypeVectorOrScalarBool() || | |
BT->isTypePointer()) | |
Inst = new ICmpInst(*BB, CmpMap::rmap(OP), | |
transValue(BC->getOperand(0), F, BB), | |
transValue(BC->getOperand(1), F, BB)); | |
else if (BT->isTypeVectorOrScalarFloat()) | |
Inst = new FCmpInst(*BB, CmpMap::rmap(OP), | |
transValue(BC->getOperand(0), F, BB), | |
transValue(BC->getOperand(1), F, BB)); | |
assert(Inst && "not implemented"); | |
return Inst; | |
} | |
bool | |
SPIRVToLLVM::postProcessOCL() { | |
std::string DemangledName; | |
SPIRVWord SrcLangVer = 0; | |
BM->getSourceLanguage(&SrcLangVer); | |
bool isCPP = SrcLangVer == kOCLVer::CL21; | |
for (auto I = M->begin(), E = M->end(); I != E;) { | |
auto F = I++; | |
if (F->hasName() && F->isDeclaration()) { | |
DEBUG(dbgs() << "[postProcessOCL sret] " << *F << '\n'); | |
if (F->getReturnType()->isStructTy() && | |
oclIsBuiltin(F->getName(), &DemangledName, isCPP)) { | |
if (!postProcessOCLBuiltinReturnStruct(static_cast<Function*>(F))) | |
return false; | |
} | |
} | |
} | |
for (auto I = M->begin(), E = M->end(); I != E;) { | |
auto F = static_cast<Function*>(I++); | |
if (F->hasName() && F->isDeclaration()) { | |
DEBUG(dbgs() << "[postProcessOCL func ptr] " << *F << '\n'); | |
auto AI = F->arg_begin(); | |
if (hasFunctionPointerArg(F, AI) && isDecoratedSPIRVFunc(F)) | |
if (!postProcessOCLBuiltinWithFuncPointer(F, AI)) | |
return false; | |
} | |
} | |
for (auto I = M->begin(), E = M->end(); I != E;) { | |
auto F = static_cast<Function*>(I++); | |
if (F->hasName() && F->isDeclaration()) { | |
DEBUG(dbgs() << "[postProcessOCL array arg] " << *F << '\n'); | |
if (hasArrayArg(F) && oclIsBuiltin(F->getName(), &DemangledName, isCPP)) | |
if (!postProcessOCLBuiltinWithArrayArguments(F, DemangledName)) | |
return false; | |
} | |
} | |
return true; | |
} | |
bool | |
SPIRVToLLVM::postProcessOCLBuiltinReturnStruct(Function *F) { | |
std::string Name = F->getName(); | |
F->setName(Name + ".old"); | |
for (auto I = F->user_begin(), E = F->user_end(); I != E;) { | |
if (auto CI = dyn_cast<CallInst>(*I++)) { | |
auto ST = dyn_cast<StoreInst>(*(CI->user_begin())); | |
assert(ST); | |
std::vector<Type *> ArgTys; | |
getFunctionTypeParameterTypes(F->getFunctionType(), ArgTys); | |
ArgTys.insert(ArgTys.begin(), PointerType::get(F->getReturnType(), | |
SPIRAS_Private)); | |
auto newF = getOrCreateFunction(M, Type::getVoidTy(*Context), | |
ArgTys, Name); | |
newF->setCallingConv(F->getCallingConv()); | |
auto Args = getArguments(CI); | |
Args.insert(Args.begin(), ST->getPointerOperand()); | |
auto NewCI = CallInst::Create(newF, Args, CI->getName(), CI); | |
NewCI->setCallingConv(CI->getCallingConv()); | |
ST->dropAllReferences(); | |
ST->removeFromParent(); | |
CI->dropAllReferences(); | |
CI->removeFromParent(); | |
} | |
} | |
F->dropAllReferences(); | |
F->removeFromParent(); | |
return true; | |
} | |
bool | |
SPIRVToLLVM::postProcessOCLBuiltinWithFuncPointer(Function* F, | |
Function::arg_iterator I) { | |
auto Name = undecorateSPIRVFunction(F->getName()); | |
std::set<Value *> InvokeFuncPtrs; | |
mutateFunctionOCL (F, [=, &InvokeFuncPtrs]( | |
CallInst *CI, std::vector<Value *> &Args) { | |
auto ALoc = std::find_if(Args.begin(), Args.end(), [](Value * elem) { | |
return isFunctionPointerType(elem->getType()); | |
}); | |
assert(ALoc != Args.end() && "Buit-in must accept a pointer to function"); | |
assert(isa<Function>(*ALoc) && "Invalid function pointer usage"); | |
Value *Ctx = ALoc[1]; | |
Value *CtxLen = ALoc[2]; | |
Value *CtxAlign = ALoc[3]; | |
if (Name == kOCLBuiltinName::EnqueueKernel) | |
assert(Args.end() - ALoc > 3); | |
else | |
assert(Args.end() - ALoc > 0); | |
// Erase arguments what are hanled by "spir_block_bind" according to SPIR 2.0 | |
Args.erase(ALoc + 1, ALoc + 4); | |
InvokeFuncPtrs.insert(*ALoc); | |
// There will be as many calls to spir_block_bind as how much device execution | |
// bult-ins using this block. This doesn't contradict SPIR 2.0 specification. | |
*ALoc = addBlockBind(M, cast<Function>(removeCast(*ALoc)), | |
Ctx, CtxLen, CtxAlign, CI); | |
return Name; | |
}); | |
for (auto &I:InvokeFuncPtrs) | |
eraseIfNoUse(I); | |
return true; | |
} | |
bool | |
SPIRVToLLVM::postProcessOCLBuiltinWithArrayArguments(Function* F, | |
const std::string &DemangledName) { | |
DEBUG(dbgs() << "[postProcessOCLBuiltinWithArrayArguments] " << *F << '\n'); | |
auto Attrs = F->getAttributes(); | |
auto Name = F->getName(); | |
mutateFunction(F, [=](CallInst *CI, std::vector<Value *> &Args) { | |
auto FBegin = CI->getParent()->getParent()->begin()->getFirstInsertionPt(); | |
for (auto &I:Args) { | |
auto T = I->getType(); | |
if (!T->isArrayTy()) | |
continue; | |
auto Alloca = new AllocaInst(T, "", static_cast<Instruction*>(FBegin)); | |
auto Store = new StoreInst(I, Alloca, false, CI); | |
auto Zero = ConstantInt::getNullValue(Type::getInt32Ty(T->getContext())); | |
Value *Index[] = {Zero, Zero}; | |
I = GetElementPtrInst::CreateInBounds(Alloca, Index, "", CI); | |
} | |
return Name; | |
}, nullptr, &Attrs); | |
return true; | |
} | |
// ToDo: Handle unsigned integer return type. May need spec change. | |
Instruction * | |
SPIRVToLLVM::postProcessOCLReadImage(SPIRVInstruction *BI, CallInst* CI, | |
const std::string &FuncName) { | |
AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); | |
StringRef ImageTypeName; | |
bool isDepthImage = false; | |
if (isOCLImageType( | |
(cast<CallInst>(CI->getOperand(0)))->getArgOperand(0)->getType(), | |
&ImageTypeName)) | |
isDepthImage = ImageTypeName.endswith("depth_t"); | |
return mutateCallInstOCL( | |
M, CI, | |
[=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { | |
CallInst *CallSampledImg = cast<CallInst>(Args[0]); | |
auto Img = CallSampledImg->getArgOperand(0); | |
assert(isOCLImageType(Img->getType())); | |
auto Sampler = CallSampledImg->getArgOperand(1); | |
Args[0] = Img; | |
Args.insert(Args.begin() + 1, Sampler); | |
if(Args.size() > 4 ) { | |
ConstantInt* ImOp = dyn_cast<ConstantInt>(Args[3]); | |
ConstantFP* LodVal = dyn_cast<ConstantFP>(Args[4]); | |
// Drop "Image Operands" argument. | |
Args.erase(Args.begin() + 3, Args.begin() + 4); | |
// If the image operand is LOD and its value is zero, drop it too. | |
if (ImOp && LodVal && LodVal->isNullValue() && | |
ImOp->getZExtValue() == ImageOperandsMask::ImageOperandsLodMask ) | |
Args.erase(Args.begin() + 3, Args.end()); | |
} | |
if (CallSampledImg->hasOneUse()) { | |
CallSampledImg->replaceAllUsesWith( | |
UndefValue::get(CallSampledImg->getType())); | |
CallSampledImg->dropAllReferences(); | |
CallSampledImg->eraseFromParent(); | |
} | |
Type *T = CI->getType(); | |
if (auto VT = dyn_cast<VectorType>(T)) | |
T = VT->getElementType(); | |
RetTy = isDepthImage ? T : CI->getType(); | |
return std::string(kOCLBuiltinName::SampledReadImage) + | |
(T->isFloatingPointTy() ? 'f' : 'i'); | |
}, | |
[=](CallInst *NewCI) -> Instruction * { | |
if (isDepthImage) | |
return InsertElementInst::Create( | |
UndefValue::get(VectorType::get(NewCI->getType(), 4)), NewCI, | |
getSizet(M, 0), "", NewCI->getParent()); | |
return NewCI; | |
}, | |
&Attrs); | |
} | |
CallInst* | |
SPIRVToLLVM::postProcessOCLWriteImage(SPIRVInstruction *BI, CallInst *CI, | |
const std::string &DemangledName) { | |
AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); | |
return mutateCallInstOCL(M, CI, [=](CallInst *, std::vector<Value *> &Args) { | |
llvm::Type *T = Args[2]->getType(); | |
if (Args.size() > 4) { | |
ConstantInt* ImOp = dyn_cast<ConstantInt>(Args[3]); | |
ConstantFP* LodVal = dyn_cast<ConstantFP>(Args[4]); | |
// Drop "Image Operands" argument. | |
Args.erase(Args.begin() + 3, Args.begin() + 4); | |
// If the image operand is LOD and its value is zero, drop it too. | |
if (ImOp && LodVal && LodVal->isNullValue() && | |
ImOp->getZExtValue() == ImageOperandsMask::ImageOperandsLodMask ) | |
Args.erase(Args.begin() + 3, Args.end()); | |
else | |
std::swap(Args[2], Args[3]); | |
} | |
return std::string(kOCLBuiltinName::WriteImage) + | |
(T->isFPOrFPVectorTy() ? 'f' : 'i'); | |
}, &Attrs); | |
} | |
CallInst * | |
SPIRVToLLVM::postProcessOCLBuildNDRange(SPIRVInstruction *BI, CallInst *CI, | |
const std::string &FuncName) { | |
assert(CI->getNumArgOperands() == 3); | |
auto GWS = CI->getArgOperand(0); | |
auto LWS = CI->getArgOperand(1); | |
auto GWO = CI->getArgOperand(2); | |
CI->setArgOperand(0, GWO); | |
CI->setArgOperand(1, GWS); | |
CI->setArgOperand(2, LWS); | |
return CI; | |
} | |
Instruction * | |
SPIRVToLLVM::postProcessGroupAllAny(CallInst *CI, | |
const std::string &DemangledName) { | |
AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); | |
return mutateCallInstSPIRV( | |
M, CI, | |
[=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { | |
Type *Int32Ty = Type::getInt32Ty(*Context); | |
RetTy = Int32Ty; | |
Args[1] = CastInst::CreateZExtOrBitCast(Args[1], Int32Ty, "", CI); | |
return DemangledName; | |
}, | |
[=](CallInst *NewCI) -> Instruction * { | |
Type *RetTy = Type::getInt1Ty(*Context); | |
return CastInst::CreateTruncOrBitCast(NewCI, RetTy, "", | |
NewCI->getNextNode()); | |
}, | |
&Attrs); | |
} | |
CallInst * | |
SPIRVToLLVM::expandOCLBuiltinWithScalarArg(CallInst* CI, | |
const std::string &FuncName) { | |
AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); | |
if (!CI->getOperand(0)->getType()->isVectorTy() && | |
CI->getOperand(1)->getType()->isVectorTy()) { | |
return mutateCallInstOCL(M, CI, [=](CallInst *, std::vector<Value *> &Args){ | |
unsigned vecSize = CI->getOperand(1)->getType()->getVectorNumElements(); | |
Value *NewVec = nullptr; | |
if (auto CA = dyn_cast<Constant>(Args[0])) | |
NewVec = ConstantVector::getSplat(vecSize, CA); | |
else { | |
NewVec = ConstantVector::getSplat(vecSize, | |
Constant::getNullValue(Args[0]->getType())); | |
NewVec = InsertElementInst::Create(NewVec, Args[0], getInt32(M, 0), "", | |
CI); | |
NewVec = new ShuffleVectorInst(NewVec, NewVec, | |
ConstantVector::getSplat(vecSize, getInt32(M, 0)), "", CI); | |
} | |
NewVec->takeName(Args[0]); | |
Args[0] = NewVec; | |
return FuncName; | |
}, &Attrs); | |
} | |
return CI; | |
} | |
std::string | |
SPIRVToLLVM::transOCLPipeTypeAccessQualifier(SPIRV::SPIRVTypePipe* ST) { | |
return SPIRSPIRVAccessQualifierMap::rmap(ST->getAccessQualifier()); | |
} | |
void | |
SPIRVToLLVM::transGeneratorMD() { | |
SPIRVMDBuilder B(*M); | |
B.addNamedMD(kSPIRVMD::Generator) | |
.addOp() | |
.addU16(BM->getGeneratorId()) | |
.addU16(BM->getGeneratorVer()) | |
.done(); | |
} | |
Value * | |
SPIRVToLLVM::oclTransConstantSampler(SPIRV::SPIRVConstantSampler* BCS) { | |
auto Lit = (BCS->getAddrMode() << 1) | | |
BCS->getNormalized() | | |
((BCS->getFilterMode() + 1) << 4); | |
auto Ty = IntegerType::getInt32Ty(*Context); | |
return ConstantInt::get(Ty, Lit); | |
} | |
Value * | |
SPIRVToLLVM::oclTransConstantPipeStorage( | |
SPIRV::SPIRVConstantPipeStorage* BCPS) { | |
string CPSName = string(kSPIRVTypeName::PrefixAndDelim) | |
+ kSPIRVTypeName::ConstantPipeStorage; | |
auto Int32Ty = IntegerType::getInt32Ty(*Context); | |
auto CPSTy = M->getTypeByName(CPSName); | |
if (!CPSTy) { | |
Type* CPSElemsTy[] = { Int32Ty, Int32Ty, Int32Ty }; | |
CPSTy = StructType::create(*Context, CPSElemsTy, CPSName); | |
} | |
assert(CPSTy != nullptr && "Could not create spirv.ConstantPipeStorage"); | |
Constant* CPSElems[] = { | |
ConstantInt::get(Int32Ty, BCPS->getPacketSize()), | |
ConstantInt::get(Int32Ty, BCPS->getPacketAlign()), | |
ConstantInt::get(Int32Ty, BCPS->getCapacity()) | |
}; | |
return new GlobalVariable(*M, CPSTy, false, GlobalValue::LinkOnceODRLinkage, | |
ConstantStruct::get(CPSTy, CPSElems), BCPS->getName(), | |
nullptr, GlobalValue::NotThreadLocal, SPIRAS_Global); | |
} | |
/// For instructions, this function assumes they are created in order | |
/// and appended to the given basic block. An instruction may use a | |
/// instruction from another BB which has not been translated. Such | |
/// instructions should be translated to place holders at the point | |
/// of first use, then replaced by real instructions when they are | |
/// created. | |
/// | |
/// When CreatePlaceHolder is true, create a load instruction of a | |
/// global variable as placeholder for SPIRV instruction. Otherwise, | |
/// create instruction and replace placeholder if there is one. | |
Value * | |
SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, | |
BasicBlock *BB, bool CreatePlaceHolder){ | |
auto OC = BV->getOpCode(); | |
IntBoolOpMap::rfind(OC, &OC); | |
// Translation of non-instruction values | |
switch(OC) { | |
case OpConstant: { | |
SPIRVConstant *BConst = static_cast<SPIRVConstant *>(BV); | |
SPIRVType *BT = BV->getType(); | |
Type *LT = transType(BT); | |
switch(BT->getOpCode()) { | |
case OpTypeBool: | |
case OpTypeInt: | |
return mapValue(BV, ConstantInt::get(LT, BConst->getZExtIntValue(), | |
static_cast<SPIRVTypeInt*>(BT)->isSigned())); | |
case OpTypeFloat: { | |
const llvm::fltSemantics *FS = nullptr; | |
switch (BT->getFloatBitWidth()) { | |
case 16: | |
FS = &APFloat::IEEEhalf; | |
break; | |
case 32: | |
FS = &APFloat::IEEEsingle; | |
break; | |
case 64: | |
FS = &APFloat::IEEEdouble; | |
break; | |
default: | |
llvm_unreachable("invalid float type"); | |
} | |
return mapValue(BV, ConstantFP::get(*Context, APFloat(*FS, | |
APInt(BT->getFloatBitWidth(), BConst->getZExtIntValue())))); | |
} | |
default: | |
llvm_unreachable("Not implemented"); | |
return nullptr; | |
} | |
} | |
case OpConstantTrue: | |
return mapValue(BV, ConstantInt::getTrue(*Context)); | |
case OpConstantFalse: | |
return mapValue(BV, ConstantInt::getFalse(*Context)); | |
case OpConstantNull: { | |
auto LT = transType(BV->getType()); | |
return mapValue(BV, Constant::getNullValue(LT)); | |
} | |
case OpConstantComposite: { | |
auto BCC = static_cast<SPIRVConstantComposite*>(BV); | |
std::vector<Constant *> CV; | |
for (auto &I:BCC->getElements()) | |
CV.push_back(dyn_cast<Constant>(transValue(I, F, BB))); | |
switch(BV->getType()->getOpCode()) { | |
case OpTypeVector: | |
return mapValue(BV, ConstantVector::get(CV)); | |
case OpTypeArray: | |
return mapValue(BV, ConstantArray::get( | |
dyn_cast<ArrayType>(transType(BCC->getType())), CV)); | |
case OpTypeStruct: { | |
auto BCCTy = dyn_cast<StructType>(transType(BCC->getType())); | |
auto Members = BCCTy->getNumElements(); | |
auto Constants = CV.size(); | |
//if we try to initialize constant TypeStruct, add bitcasts | |
//if src and dst types are both pointers but to different types | |
if (Members == Constants) { | |
for (unsigned i = 0; i < Members; ++i) { | |
if (CV[i]->getType() == BCCTy->getElementType(i)) | |
continue; | |
if (!CV[i]->getType()->isPointerTy() || | |
!BCCTy->getElementType(i)->isPointerTy()) | |
continue; | |
CV[i] = ConstantExpr::getBitCast(CV[i], BCCTy->getElementType(i)); | |
} | |
} | |
return mapValue(BV, ConstantStruct::get( | |
dyn_cast<StructType>(transType(BCC->getType())), CV)); | |
} | |
default: | |
llvm_unreachable("not implemented"); | |
return nullptr; | |
} | |
} | |
case OpConstantSampler: { | |
auto BCS = static_cast<SPIRVConstantSampler*>(BV); | |
return mapValue(BV, oclTransConstantSampler(BCS)); | |
} | |
case OpConstantPipeStorage: { | |
auto BCPS = static_cast<SPIRVConstantPipeStorage*>(BV); | |
return mapValue(BV, oclTransConstantPipeStorage(BCPS)); | |
} | |
case OpSpecConstantOp: { | |
auto BI = createInstFromSpecConstantOp( | |
static_cast<SPIRVSpecConstantOp*>(BV)); | |
return mapValue(BV, transValue(BI, nullptr, nullptr, false)); | |
} | |
case OpUndef: | |
return mapValue(BV, UndefValue::get(transType(BV->getType()))); | |
case OpVariable: { | |
auto BVar = static_cast<SPIRVVariable *>(BV); | |
auto Ty = transType(BVar->getType()->getPointerElementType()); | |
bool IsConst = BVar->isConstant(); | |
llvm::GlobalValue::LinkageTypes LinkageTy = transLinkageType(BVar); | |
Constant *Initializer = nullptr; | |
SPIRVValue *Init = BVar->getInitializer(); | |
if (Init) | |
Initializer = dyn_cast<Constant>(transValue(Init, F, BB, false)); | |
else if (LinkageTy == GlobalValue::CommonLinkage) | |
// In LLVM variables with common linkage type must be initilized by 0 | |
Initializer = Constant::getNullValue(Ty); | |
SPIRVStorageClassKind BS = BVar->getStorageClass(); | |
if (BS == StorageClassFunction && !Init) { | |
assert (BB && "Invalid BB"); | |
return mapValue(BV, new AllocaInst(Ty, BV->getName(), BB)); | |
} | |
auto AddrSpace = SPIRSPIRVAddrSpaceMap::rmap(BS); | |
auto LVar = new GlobalVariable(*M, Ty, IsConst, LinkageTy, Initializer, | |
BV->getName(), 0, GlobalVariable::NotThreadLocal, AddrSpace); | |
LVar->setUnnamedAddr(IsConst && Ty->isArrayTy() && | |
Ty->getArrayElementType()->isIntegerTy(8)); | |
SPIRVBuiltinVariableKind BVKind; | |
if (BVar->isBuiltin(&BVKind)) | |
BuiltinGVMap[LVar] = BVKind; | |
return mapValue(BV, LVar); | |
} | |
case OpFunctionParameter: { | |
auto BA = static_cast<SPIRVFunctionParameter*>(BV); | |
assert (F && "Invalid function"); | |
unsigned ArgNo = 0; | |
for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(); I != E; | |
++I, ++ArgNo) { | |
if (ArgNo == BA->getArgNo()) | |
return mapValue(BV, static_cast<Argument*>(I)); | |
} | |
llvm_unreachable("Invalid argument"); | |
return nullptr; | |
} | |
case OpFunction: | |
return mapValue(BV, transFunction(static_cast<SPIRVFunction *>(BV))); | |
case OpLabel: | |
return mapValue(BV, BasicBlock::Create(*Context, BV->getName(), F)); | |
case OpBitcast: // Can be translated without BB pointer | |
if(!CreatePlaceHolder) // May be a placeholder | |
return mapValue(BV, transConvertInst(BV, F, BB)); | |
default: | |
// do nothing | |
break; | |
} | |
// All other values require valid BB pointer. | |
assert(BB && "Invalid BB"); | |
// Creation of place holder | |
if (CreatePlaceHolder) { | |
auto GV = new GlobalVariable(*M, | |
transType(BV->getType()), | |
false, | |
GlobalValue::PrivateLinkage, | |
nullptr, | |
std::string(kPlaceholderPrefix) + BV->getName(), | |
0, GlobalVariable::NotThreadLocal, 0); | |
auto LD = new LoadInst(GV, BV->getName(), BB); | |
PlaceholderMap[BV] = LD; | |
return mapValue(BV, LD); | |
} | |
// Translation of instructions | |
switch (BV->getOpCode()) { | |
case OpBranch: { | |
auto BR = static_cast<SPIRVBranch *>(BV); | |
return mapValue(BV, BranchInst::Create( | |
dyn_cast<BasicBlock>(transValue(BR->getTargetLabel(), F, BB)), BB)); | |
} | |
case OpBranchConditional: { | |
auto BR = static_cast<SPIRVBranchConditional *>(BV); | |
return mapValue( | |
BV, BranchInst::Create( | |
dyn_cast<BasicBlock>(transValue(BR->getTrueLabel(), F, BB)), | |
dyn_cast<BasicBlock>(transValue(BR->getFalseLabel(), F, BB)), | |
transValue(BR->getCondition(), F, BB), BB)); | |
} | |
case OpPhi: { | |
auto Phi = static_cast<SPIRVPhi *>(BV); | |
auto LPhi = dyn_cast<PHINode>(mapValue( | |
BV, PHINode::Create(transType(Phi->getType()), | |
Phi->getPairs().size() / 2, Phi->getName(), BB))); | |
Phi->foreachPair([&](SPIRVValue *IncomingV, SPIRVBasicBlock *IncomingBB, | |
size_t Index) { | |
auto Translated = transValue(IncomingV, F, BB); | |
LPhi->addIncoming(Translated, | |
dyn_cast<BasicBlock>(transValue(IncomingBB, F, BB))); | |
}); | |
return LPhi; | |
} | |
case OpReturn: | |
return mapValue(BV, ReturnInst::Create(*Context, BB)); | |
case OpReturnValue: { | |
auto RV = static_cast<SPIRVReturnValue *>(BV); | |
return mapValue( | |
BV, ReturnInst::Create(*Context, | |
transValue(RV->getReturnValue(), F, BB), BB)); | |
} | |
case OpStore: { | |
SPIRVStore *BS = static_cast<SPIRVStore*>(BV); | |
StoreInst *SI = new StoreInst(transValue(BS->getSrc(), F, BB), | |
transValue(BS->getDst(), F, BB), | |
BS->SPIRVMemoryAccess::isVolatile(), | |
BS->SPIRVMemoryAccess::getAlignment(), BB); | |
if (BS->SPIRVMemoryAccess::isNonTemporal()) | |
transNonTemporalMetadata(SI); | |
return mapValue(BV, SI); | |
} | |
case OpLoad: { | |
SPIRVLoad *BL = static_cast<SPIRVLoad*>(BV); | |
LoadInst *LI = new LoadInst(transValue(BL->getSrc(), F, BB), BV->getName(), | |
BL->SPIRVMemoryAccess::isVolatile(), | |
BL->SPIRVMemoryAccess::getAlignment(), BB); | |
if (BL->SPIRVMemoryAccess::isNonTemporal()) | |
transNonTemporalMetadata(LI); | |
return mapValue(BV, LI); | |
} | |
case OpCopyMemorySized: { | |
SPIRVCopyMemorySized *BC = static_cast<SPIRVCopyMemorySized *>(BV); | |
std::string FuncName = "llvm.memcpy"; | |
SPIRVType* BS = BC->getSource()->getType(); | |
SPIRVType* BT = BC->getTarget()->getType(); | |
Type *Int1Ty = Type::getInt1Ty(*Context); | |
Type* Int32Ty = Type::getInt32Ty(*Context); | |
Type* VoidTy = Type::getVoidTy(*Context); | |
Type* SrcTy = transType(BS); | |
Type* TrgTy = transType(BT); | |
Type* SizeTy = transType(BC->getSize()->getType()); | |
Type* ArgTy[] = { TrgTy, SrcTy, SizeTy, Int32Ty, Int1Ty }; | |
ostringstream TempName; | |
TempName << ".p" << SPIRSPIRVAddrSpaceMap::rmap(BT->getPointerStorageClass()) << "i8"; | |
TempName << ".p" << SPIRSPIRVAddrSpaceMap::rmap(BS->getPointerStorageClass()) << "i8"; | |
FuncName += TempName.str(); | |
if (BC->getSize()->getType()->getBitWidth() == 32) | |
FuncName += ".i32"; | |
else | |
FuncName += ".i64"; | |
FunctionType *FT = FunctionType::get(VoidTy, ArgTy, false); | |
Function *Func = dyn_cast<Function>(M->getOrInsertFunction(FuncName, FT)); | |
assert(Func && Func->getFunctionType() == FT && "Function type mismatch"); | |
Func->setLinkage(GlobalValue::ExternalLinkage); | |
if (isFuncNoUnwind()) | |
Func->addFnAttr(Attribute::NoUnwind); | |
Value *Arg[] = { transValue(BC->getTarget(), Func, BB), | |
transValue(BC->getSource(), Func, BB), | |
dyn_cast<llvm::ConstantInt>(transValue(BC->getSize(), | |
Func, BB)), | |
ConstantInt::get(Int32Ty, | |
BC->SPIRVMemoryAccess::getAlignment()), | |
ConstantInt::get(Int1Ty, | |
BC->SPIRVMemoryAccess::isVolatile())}; | |
return mapValue( BV, CallInst::Create(Func, Arg, "", BB)); | |
} | |
case OpSelect: { | |
SPIRVSelect *BS = static_cast<SPIRVSelect*>(BV); | |
return mapValue(BV, | |
SelectInst::Create(transValue(BS->getCondition(), F, BB), | |
transValue(BS->getTrueValue(), F, BB), | |
transValue(BS->getFalseValue(), F, BB), | |
BV->getName(), BB)); | |
} | |
case OpSwitch: { | |
auto BS = static_cast<SPIRVSwitch *>(BV); | |
auto Select = transValue(BS->getSelect(), F, BB); | |
auto LS = SwitchInst::Create( | |
Select, dyn_cast<BasicBlock>(transValue(BS->getDefault(), F, BB)), | |
BS->getNumPairs(), BB); | |
BS->foreachPair( | |
[&](SPIRVWord Literal, SPIRVBasicBlock *Label, size_t Index) { | |
LS->addCase(ConstantInt::get(dyn_cast<IntegerType>(Select->getType()), | |
Literal), | |
dyn_cast<BasicBlock>(transValue(Label, F, BB))); | |
}); | |
return mapValue(BV, LS); | |
} | |
case OpAccessChain: | |
case OpInBoundsAccessChain: | |
case OpPtrAccessChain: | |
case OpInBoundsPtrAccessChain: { | |
auto AC = static_cast<SPIRVAccessChainBase *>(BV); | |
auto Base = transValue(AC->getBase(), F, BB); | |
auto Index = transValue(AC->getIndices(), F, BB); | |
if (!AC->hasPtrIndex()) | |
Index.insert(Index.begin(), getInt32(M, 0)); | |
auto IsInbound = AC->isInBounds(); | |
Value *V = nullptr; | |
if (BB) { | |
auto GEP = GetElementPtrInst::Create(nullptr, Base, Index, | |
BV->getName(), BB); | |
GEP->setIsInBounds(IsInbound); | |
V = GEP; | |
} else { | |
V = ConstantExpr::getGetElementPtr(Base->getType(), | |
dyn_cast<Constant>(Base), | |
Index, | |
IsInbound); | |
} | |
return mapValue(BV, V); | |
} | |
case OpCompositeExtract: { | |
SPIRVCompositeExtract *CE = static_cast<SPIRVCompositeExtract *>(BV); | |
if (CE->getComposite()->getType()->isTypeVector()) { | |
assert(CE->getIndices().size() == 1 && "Invalid index"); | |
return mapValue( | |
BV, ExtractElementInst::Create( | |
transValue(CE->getComposite(), F, BB), | |
ConstantInt::get(*Context, APInt(32, CE->getIndices()[0])), | |
BV->getName(), BB)); | |
} | |
return mapValue( | |
BV, ExtractValueInst::Create( | |
transValue(CE->getComposite(), F, BB), | |
CE->getIndices(), BV->getName(), BB)); | |
} | |
case OpVectorExtractDynamic: { | |
auto CE = static_cast<SPIRVVectorExtractDynamic *>(BV); | |
return mapValue( | |
BV, ExtractElementInst::Create(transValue(CE->getVector(), F, BB), | |
transValue(CE->getIndex(), F, BB), | |
BV->getName(), BB)); | |
} | |
case OpCompositeInsert: { | |
auto CI = static_cast<SPIRVCompositeInsert *>(BV); | |
if (CI->getComposite()->getType()->isTypeVector()) { | |
assert(CI->getIndices().size() == 1 && "Invalid index"); | |
return mapValue( | |
BV, InsertElementInst::Create( | |
transValue(CI->getComposite(), F, BB), | |
transValue(CI->getObject(), F, BB), | |
ConstantInt::get(*Context, APInt(32, CI->getIndices()[0])), | |
BV->getName(), BB)); | |
} | |
return mapValue( | |
BV, InsertValueInst::Create( | |
transValue(CI->getComposite(), F, BB), | |
transValue(CI->getObject(), F, BB), | |
CI->getIndices(), BV->getName(), BB)); | |
} | |
case OpVectorInsertDynamic: { | |
auto CI = static_cast<SPIRVVectorInsertDynamic *>(BV); | |
return mapValue( | |
BV, InsertElementInst::Create(transValue(CI->getVector(), F, BB), | |
transValue(CI->getComponent(), F, BB), | |
transValue(CI->getIndex(), F, BB), | |
BV->getName(), BB)); | |
} | |
case OpVectorShuffle: { | |
auto VS = static_cast<SPIRVVectorShuffle *>(BV); | |
std::vector<Constant *> Components; | |
IntegerType *Int32Ty = IntegerType::get(*Context, 32); | |
for (auto I : VS->getComponents()) { | |
if (I == static_cast<SPIRVWord>(-1)) | |
Components.push_back(UndefValue::get(Int32Ty)); | |
else | |
Components.push_back(ConstantInt::get(Int32Ty, I)); | |
} | |
return mapValue(BV, | |
new ShuffleVectorInst(transValue(VS->getVector1(), F, BB), | |
transValue(VS->getVector2(), F, BB), | |
ConstantVector::get(Components), | |
BV->getName(), BB)); | |
} | |
case OpFunctionCall: { | |
SPIRVFunctionCall *BC = static_cast<SPIRVFunctionCall *>(BV); | |
auto Call = CallInst::Create(transFunction(BC->getFunction()), | |
transValue(BC->getArgumentValues(), F, BB), | |
BC->getName(), BB); | |
setCallingConv(Call); | |
setAttrByCalledFunc(Call); | |
return mapValue(BV, Call); | |
} | |
case OpExtInst: | |
return mapValue( | |
BV, transOCLBuiltinFromExtInst(static_cast<SPIRVExtInst *>(BV), BB)); | |
case OpControlBarrier: | |
case OpMemoryBarrier: | |
return mapValue( | |
BV, transOCLBarrierFence(static_cast<SPIRVInstruction *>(BV), BB)); | |
case OpSNegate: { | |
SPIRVUnary *BC = static_cast<SPIRVUnary *>(BV); | |
return mapValue( | |
BV, BinaryOperator::CreateNSWNeg(transValue(BC->getOperand(0), F, BB), | |
BV->getName(), BB)); | |
} | |
case OpFNegate: { | |
SPIRVUnary *BC = static_cast<SPIRVUnary *>(BV); | |
return mapValue( | |
BV, BinaryOperator::CreateFNeg(transValue(BC->getOperand(0), F, BB), | |
BV->getName(), BB)); | |
} | |
case OpNot: { | |
SPIRVUnary *BC = static_cast<SPIRVUnary *>(BV); | |
return mapValue( | |
BV, BinaryOperator::CreateNot(transValue(BC->getOperand(0), F, BB), | |
BV->getName(), BB)); | |
} | |
case OpAll : | |
case OpAny : | |
return mapValue(BV, | |
transOCLAllAny(static_cast<SPIRVInstruction *>(BV), BB)); | |
case OpIsFinite : | |
case OpIsInf : | |
case OpIsNan : | |
case OpIsNormal : | |
case OpSignBitSet : | |
return mapValue(BV, | |
transOCLRelational(static_cast<SPIRVInstruction *>(BV), BB)); | |
default: { | |
auto OC = BV->getOpCode(); | |
if (isSPIRVCmpInstTransToLLVMInst(static_cast<SPIRVInstruction*>(BV))) { | |
return mapValue(BV, transCmpInst(BV, BB, F)); | |
} else if (OCLSPIRVBuiltinMap::rfind(OC, nullptr) && | |
!isAtomicOpCode(OC) && | |
!isGroupOpCode(OC) && | |
!isPipeOpCode(OC)) { | |
return mapValue(BV, transOCLBuiltinFromInst( | |
static_cast<SPIRVInstruction *>(BV), BB)); | |
} else if (isBinaryShiftLogicalBitwiseOpCode(OC) || | |
isLogicalOpCode(OC)) { | |
return mapValue(BV, transShiftLogicalBitwiseInst(BV, BB, F)); | |
} else if (isCvtOpCode(OC)) { | |
auto BI = static_cast<SPIRVInstruction *>(BV); | |
Value *Inst = nullptr; | |
if (BI->hasFPRoundingMode() || BI->isSaturatedConversion()) | |
Inst = transOCLBuiltinFromInst(BI, BB); | |
else | |
Inst = transConvertInst(BV, F, BB); | |
return mapValue(BV, Inst); | |
} | |
return mapValue(BV, transSPIRVBuiltinFromInst( | |
static_cast<SPIRVInstruction *>(BV), BB)); | |
} | |
SPIRVDBG(spvdbgs() << "Cannot translate " << *BV << '\n';) | |
llvm_unreachable("Translation of SPIRV instruction not implemented"); | |
return NULL; | |
} | |
} | |
template<class SourceTy, class FuncTy> | |
bool | |
SPIRVToLLVM::foreachFuncCtlMask(SourceTy Source, FuncTy Func) { | |
SPIRVWord FCM = Source->getFuncCtlMask(); | |
SPIRSPIRVFuncCtlMaskMap::foreach([&](Attribute::AttrKind Attr, | |
SPIRVFunctionControlMaskKind Mask){ | |
if (FCM & Mask) | |
Func(Attr); | |
}); | |
return true; | |
} | |
Function * | |
SPIRVToLLVM::transFunction(SPIRVFunction *BF) { | |
auto Loc = FuncMap.find(BF); | |
if (Loc != FuncMap.end()) | |
return Loc->second; | |
auto IsKernel = BM->isEntryPoint(ExecutionModelKernel, BF->getId()); | |
auto Linkage = IsKernel ? GlobalValue::ExternalLinkage : transLinkageType(BF); | |
FunctionType *FT = dyn_cast<FunctionType>(transType(BF->getFunctionType())); | |
Function *F = dyn_cast<Function>(mapValue(BF, Function::Create(FT, Linkage, | |
BF->getName(), M))); | |
assert(F); | |
mapFunction(BF, F); | |
if (!F->isIntrinsic()) { | |
F->setCallingConv(IsKernel ? CallingConv::SPIR_KERNEL : | |
CallingConv::SPIR_FUNC); | |
if (isFuncNoUnwind()) | |
F->addFnAttr(Attribute::NoUnwind); | |
foreachFuncCtlMask(BF, [&](Attribute::AttrKind Attr){ | |
F->addFnAttr(Attr); | |
}); | |
} | |
for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(); I != E; | |
++I) { | |
auto BA = BF->getArgument(I->getArgNo()); | |
mapValue(BA, static_cast<Argument*>(I)); | |
setName(static_cast<Argument*>(I), BA); | |
BA->foreachAttr([&](SPIRVFuncParamAttrKind Kind){ | |
if (Kind == FunctionParameterAttributeNoWrite) | |
return; | |
F->addAttribute(I->getArgNo() + 1, SPIRSPIRVFuncParamAttrMap::rmap(Kind)); | |
}); | |
SPIRVWord MaxOffset = 0; | |
if (BA->hasDecorate(DecorationMaxByteOffset, 0, &MaxOffset)) { | |
AttrBuilder Builder; | |
Builder.addDereferenceableAttr(MaxOffset); | |
I->addAttr(AttributeSet::get(*Context, I->getArgNo() + 1, Builder)); | |
} | |
} | |
BF->foreachReturnValueAttr([&](SPIRVFuncParamAttrKind Kind){ | |
if (Kind == FunctionParameterAttributeNoWrite) | |
return; | |
F->addAttribute(AttributeSet::ReturnIndex, | |
SPIRSPIRVFuncParamAttrMap::rmap(Kind)); | |
}); | |
// Creating all basic blocks before creating instructions. | |
for (size_t I = 0, E = BF->getNumBasicBlock(); I != E; ++I) { | |
transValue(BF->getBasicBlock(I), F, nullptr); | |
} | |
for (size_t I = 0, E = BF->getNumBasicBlock(); I != E; ++I) { | |
SPIRVBasicBlock *BBB = BF->getBasicBlock(I); | |
BasicBlock *BB = dyn_cast<BasicBlock>(transValue(BBB, F, nullptr)); | |
for (size_t BI = 0, BE = BBB->getNumInst(); BI != BE; ++BI) { | |
SPIRVInstruction *BInst = BBB->getInst(BI); | |
transValue(BInst, F, BB, false); | |
} | |
} | |
return F; | |
} | |
/// LLVM convert builtin functions is translated to two instructions: | |
/// y = i32 islessgreater(float x, float z) -> | |
/// y = i32 ZExt(bool LessGreater(float x, float z)) | |
/// When translating back, for simplicity, a trunc instruction is inserted | |
/// w = bool LessGreater(float x, float z) -> | |
/// w = bool Trunc(i32 islessgreater(float x, float z)) | |
/// Optimizer should be able to remove the redundant trunc/zext | |
void | |
SPIRVToLLVM::transOCLBuiltinFromInstPreproc(SPIRVInstruction* BI, Type *&RetTy, | |
std::vector<SPIRVValue *> &Args) { | |
if (!BI->hasType()) | |
return; | |
auto BT = BI->getType(); | |
auto OC = BI->getOpCode(); | |
if (isCmpOpCode(BI->getOpCode())) { | |
if (BT->isTypeBool()) | |
RetTy = IntegerType::getInt32Ty(*Context); | |
else if (BT->isTypeVectorBool()) | |
RetTy = VectorType::get(IntegerType::get(*Context, | |
Args[0]->getType()->getVectorComponentType()->isTypeFloat(64)?64:32), | |
BT->getVectorComponentCount()); | |
else | |
llvm_unreachable("invalid compare instruction"); | |
} else if (OC == OpGenericCastToPtrExplicit) | |
Args.pop_back(); | |
else if (OC == OpImageRead && Args.size() > 2) { | |
// Drop "Image operands" argument | |
Args.erase(Args.begin() + 2); | |
} | |
} | |
Instruction* | |
SPIRVToLLVM::transOCLBuiltinPostproc(SPIRVInstruction* BI, | |
CallInst* CI, BasicBlock* BB, const std::string &DemangledName) { | |
auto OC = BI->getOpCode(); | |
if (isCmpOpCode(OC) && | |
BI->getType()->isTypeVectorOrScalarBool()) { | |
return CastInst::Create(Instruction::Trunc, CI, transType(BI->getType()), | |
"cvt", BB); | |
} | |
if (OC == OpImageSampleExplicitLod) | |
return postProcessOCLReadImage(BI, CI, DemangledName); | |
if (OC == OpImageWrite) { | |
return postProcessOCLWriteImage(BI, CI, DemangledName); | |
} | |
if (OC == OpGenericPtrMemSemantics) | |
return BinaryOperator::CreateShl(CI, getInt32(M, 8), "", BB); | |
if (OC == OpImageQueryFormat) | |
return BinaryOperator::CreateSub( | |
CI, getInt32(M, OCLImageChannelDataTypeOffset), "", BB); | |
if (OC == OpImageQueryOrder) | |
return BinaryOperator::CreateSub( | |
CI, getInt32(M, OCLImageChannelOrderOffset), "", BB); | |
if (OC == OpBuildNDRange) | |
return postProcessOCLBuildNDRange(BI, CI, DemangledName); | |
if (OC == OpGroupAll || OC == OpGroupAny) | |
return postProcessGroupAllAny(CI, DemangledName); | |
if (SPIRVEnableStepExpansion && | |
(DemangledName == "smoothstep" || | |
DemangledName == "step")) | |
return expandOCLBuiltinWithScalarArg(CI, DemangledName); | |
return CI; | |
} | |
Instruction * | |
SPIRVToLLVM::transBuiltinFromInst(const std::string& FuncName, | |
SPIRVInstruction* BI, BasicBlock* BB) { | |
std::string MangledName; | |
auto Ops = BI->getOperands(); | |
Type* RetTy = BI->hasType() ? transType(BI->getType()) : | |
Type::getVoidTy(*Context); | |
transOCLBuiltinFromInstPreproc(BI, RetTy, Ops); | |
std::vector<Type*> ArgTys = transTypeVector( | |
SPIRVInstruction::getOperandTypes(Ops)); | |
bool HasFuncPtrArg = false; | |
for (auto& I:ArgTys) { | |
if (isa<FunctionType>(I)) { | |
I = PointerType::get(I, SPIRAS_Private); | |
HasFuncPtrArg = true; | |
} | |
} | |
if (!HasFuncPtrArg) | |
MangleOpenCLBuiltin(FuncName, ArgTys, MangledName); | |
else | |
MangledName = decorateSPIRVFunction(FuncName); | |
Function* Func = M->getFunction(MangledName); | |
FunctionType* FT = FunctionType::get(RetTy, ArgTys, false); | |
// ToDo: Some intermediate functions have duplicate names with | |
// different function types. This is OK if the function name | |
// is used internally and finally translated to unique function | |
// names. However it is better to have a way to differentiate | |
// between intermidiate functions and final functions and make | |
// sure final functions have unique names. | |
SPIRVDBG( | |
if (!HasFuncPtrArg && Func && Func->getFunctionType() != FT) { | |
dbgs() << "Warning: Function name conflict:\n" | |
<< *Func << '\n' | |
<< " => " << *FT << '\n'; | |
} | |
) | |
if (!Func || Func->getFunctionType() != FT) { | |
DEBUG(for (auto& I:ArgTys) { | |
dbgs() << *I << '\n'; | |
}); | |
Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); | |
Func->setCallingConv(CallingConv::SPIR_FUNC); | |
if (isFuncNoUnwind()) | |
Func->addFnAttr(Attribute::NoUnwind); | |
} | |
auto Call = CallInst::Create(Func, | |
transValue(Ops, BB->getParent(), BB), "", BB); | |
setName(Call, BI); | |
setAttrByCalledFunc(Call); | |
SPIRVDBG(spvdbgs() << "[transInstToBuiltinCall] " << *BI << " -> "; dbgs() << | |
*Call << '\n';) | |
Instruction *Inst = Call; | |
Inst = transOCLBuiltinPostproc(BI, Call, BB, FuncName); | |
return Inst; | |
} | |
std::string | |
SPIRVToLLVM::getOCLBuiltinName(SPIRVInstruction* BI) { | |
auto OC = BI->getOpCode(); | |
if (OC == OpGenericCastToPtrExplicit) | |
return getOCLGenericCastToPtrName(BI); | |
if (isCvtOpCode(OC)) | |
return getOCLConvertBuiltinName(BI); | |
if (OC == OpBuildNDRange) { | |
auto NDRangeInst = static_cast<SPIRVBuildNDRange *>(BI); | |
auto EleTy = ((NDRangeInst->getOperands())[0])->getType(); | |
int Dim = EleTy->isTypeArray() ? EleTy->getArrayLength() : 1; | |
// cygwin does not have std::to_string | |
ostringstream OS; | |
OS << Dim; | |
assert((EleTy->isTypeInt() && Dim == 1) || | |
(EleTy->isTypeArray() && Dim >= 2 && Dim <= 3)); | |
return std::string(kOCLBuiltinName::NDRangePrefix) + OS.str() + "D"; | |
} | |
auto Name = OCLSPIRVBuiltinMap::rmap(OC); | |
SPIRVType *T = nullptr; | |
switch(OC) { | |
case OpImageRead: | |
T = BI->getType(); | |
break; | |
case OpImageWrite: | |
T = BI->getOperands()[2]->getType(); | |
break; | |
default: | |
// do nothing | |
break; | |
} | |
if (T && T->isTypeVector()) | |
T = T->getVectorComponentType(); | |
if (T) | |
Name += T->isTypeFloat()?'f':'i'; | |
return Name; | |
} | |
Instruction * | |
SPIRVToLLVM::transOCLBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB) { | |
assert(BB && "Invalid BB"); | |
auto FuncName = getOCLBuiltinName(BI); | |
return transBuiltinFromInst(FuncName, BI, BB); | |
} | |
Instruction * | |
SPIRVToLLVM::transSPIRVBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB) { | |
assert(BB && "Invalid BB"); | |
string Suffix = ""; | |
if (BI->getOpCode() == OpCreatePipeFromPipeStorage) { | |
auto CPFPS = static_cast<SPIRVCreatePipeFromPipeStorage*>(BI); | |
assert(CPFPS->getType()->isTypePipe() && | |
"Invalid type of CreatePipeFromStorage"); | |
auto PipeType = static_cast<SPIRVTypePipe*>(CPFPS->getType()); | |
switch (PipeType->getAccessQualifier()) { | |
case AccessQualifierReadOnly: Suffix = "_read"; break; | |
case AccessQualifierWriteOnly: Suffix = "_write"; break; | |
case AccessQualifierReadWrite: Suffix = "_read_write"; break; | |
} | |
} | |
return transBuiltinFromInst(getSPIRVFuncName(BI->getOpCode(), Suffix), BI, BB); | |
} | |
bool | |
SPIRVToLLVM::translate() { | |
if (!transAddressingModel()) | |
return false; | |
DbgTran.createCompileUnit(); | |
DbgTran.addDbgInfoVersion(); | |
for (unsigned I = 0, E = BM->getNumVariables(); I != E; ++I) { | |
auto BV = BM->getVariable(I); | |
if (BV->getStorageClass() != StorageClassFunction) | |
transValue(BV, nullptr, nullptr); | |
} | |
for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) { | |
transFunction(BM->getFunction(I)); | |
} | |
if (!transKernelMetadata()) | |
return false; | |
if (!transFPContractMetadata()) | |
return false; | |
if (!transSourceLanguage()) | |
return false; | |
if (!transSourceExtension()) | |
return false; | |
transGeneratorMD(); | |
if (!transOCLBuiltinsFromVariables()) | |
return false; | |
if (!postProcessOCL()) | |
return false; | |
eraseUselessFunctions(M); | |
DbgTran.finalize(); | |
return true; | |
} | |
bool | |
SPIRVToLLVM::transAddressingModel() { | |
switch (BM->getAddressingModel()) { | |
case AddressingModelPhysical64: | |
M->setTargetTriple(SPIR_TARGETTRIPLE64); | |
M->setDataLayout(SPIR_DATALAYOUT64); | |
break; | |
case AddressingModelPhysical32: | |
M->setTargetTriple(SPIR_TARGETTRIPLE32); | |
M->setDataLayout(SPIR_DATALAYOUT32); | |
break; | |
case AddressingModelLogical: | |
// Do not set target triple and data layout | |
break; | |
default: | |
SPIRVCKRT(0, InvalidAddressingModel, "Actual addressing mode is " + | |
(unsigned)BM->getAddressingModel()); | |
} | |
return true; | |
} | |
bool | |
SPIRVToLLVM::transDecoration(SPIRVValue *BV, Value *V) { | |
if (!transAlign(BV, V)) | |
return false; | |
DbgTran.transDbgInfo(BV, V); | |
return true; | |
} | |
bool | |
SPIRVToLLVM::transFPContractMetadata() { | |
bool ContractOff = false; | |
for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) { | |
SPIRVFunction *BF = BM->getFunction(I); | |
if (!isOpenCLKernel(BF)) | |
continue; | |
if (BF->getExecutionMode(ExecutionModeContractionOff)) { | |
ContractOff = true; | |
break; | |
} | |
} | |
if (!ContractOff) | |
M->getOrInsertNamedMetadata(kSPIR2MD::FPContract); | |
return true; | |
} | |
std::string SPIRVToLLVM::transOCLImageTypeAccessQualifier( | |
SPIRV::SPIRVTypeImage* ST) { | |
return SPIRSPIRVAccessQualifierMap::rmap(ST->getAccessQualifier()); | |
} | |
bool | |
SPIRVToLLVM::transNonTemporalMetadata(Instruction *I) { | |
Constant* One = ConstantInt::get(Type::getInt32Ty(*Context), 1); | |
MDNode *Node = MDNode::get(*Context, ConstantAsMetadata::get(One)); | |
I->setMetadata(M->getMDKindID("nontemporal"), Node); | |
return true; | |
} | |
bool | |
SPIRVToLLVM::transKernelMetadata() { | |
NamedMDNode *KernelMDs = M->getOrInsertNamedMetadata(SPIR_MD_KERNELS); | |
for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) { | |
SPIRVFunction *BF = BM->getFunction(I); | |
Function *F = static_cast<Function *>(getTranslatedValue(BF)); | |
assert(F && "Invalid translated function"); | |
if (F->getCallingConv() != CallingConv::SPIR_KERNEL) | |
continue; | |
std::vector<llvm::Metadata*> KernelMD; | |
KernelMD.push_back(ValueAsMetadata::get(F)); | |
// Generate metadata for kernel_arg_address_spaces | |
addOCLKernelArgumentMetadata(Context, KernelMD, | |
SPIR_MD_KERNEL_ARG_ADDR_SPACE, BF, | |
[=](SPIRVFunctionParameter *Arg){ | |
SPIRVType *ArgTy = Arg->getType(); | |
SPIRAddressSpace AS = SPIRAS_Private; | |
if (ArgTy->isTypePointer()) | |
AS = SPIRSPIRVAddrSpaceMap::rmap(ArgTy->getPointerStorageClass()); | |
else if (ArgTy->isTypeOCLImage() || ArgTy->isTypePipe()) | |
AS = SPIRAS_Global; | |
return ConstantAsMetadata::get( | |
ConstantInt::get(Type::getInt32Ty(*Context), AS)); | |
}); | |
// Generate metadata for kernel_arg_access_qual | |
addOCLKernelArgumentMetadata(Context, KernelMD, | |
SPIR_MD_KERNEL_ARG_ACCESS_QUAL, BF, | |
[=](SPIRVFunctionParameter *Arg){ | |
std::string Qual; | |
auto T = Arg->getType(); | |
if (T->isTypeOCLImage()) { | |
auto ST = static_cast<SPIRVTypeImage *>(T); | |
Qual = transOCLImageTypeAccessQualifier(ST); | |
} else if (T->isTypePipe()){ | |
auto PT = static_cast<SPIRVTypePipe *>(T); | |
Qual = transOCLPipeTypeAccessQualifier(PT); | |
} else | |
Qual = "none"; | |
return MDString::get(*Context, Qual); | |
}); | |
// Generate metadata for kernel_arg_type | |
addOCLKernelArgumentMetadata(Context, KernelMD, | |
SPIR_MD_KERNEL_ARG_TYPE, BF, | |
[=](SPIRVFunctionParameter *Arg){ | |
return transOCLKernelArgTypeName(Arg); | |
}); | |
// Generate metadata for kernel_arg_type_qual | |
addOCLKernelArgumentMetadata(Context, KernelMD, | |
SPIR_MD_KERNEL_ARG_TYPE_QUAL, BF, | |
[=](SPIRVFunctionParameter *Arg){ | |
std::string Qual; | |
if (Arg->hasDecorate(DecorationVolatile)) | |
Qual = kOCLTypeQualifierName::Volatile; | |
Arg->foreachAttr([&](SPIRVFuncParamAttrKind Kind){ | |
Qual += Qual.empty() ? "" : " "; | |
switch(Kind){ | |
case FunctionParameterAttributeNoAlias: | |
Qual += kOCLTypeQualifierName::Restrict; | |
break; | |
case FunctionParameterAttributeNoWrite: | |
Qual += kOCLTypeQualifierName::Const; | |
break; | |
default: | |
// do nothing. | |
break; | |
} | |
}); | |
if (Arg->getType()->isTypePipe()) { | |
Qual += Qual.empty() ? "" : " "; | |
Qual += kOCLTypeQualifierName::Pipe; | |
} | |
return MDString::get(*Context, Qual); | |
}); | |
// Generate metadata for kernel_arg_base_type | |
addOCLKernelArgumentMetadata(Context, KernelMD, | |
SPIR_MD_KERNEL_ARG_BASE_TYPE, BF, | |
[=](SPIRVFunctionParameter *Arg){ | |
return transOCLKernelArgTypeName(Arg); | |
}); | |
// Generate metadata for kernel_arg_name | |
if (SPIRVGenKernelArgNameMD) { | |
bool ArgHasName = true; | |
BF->foreachArgument([&](SPIRVFunctionParameter *Arg){ | |
ArgHasName &= !Arg->getName().empty(); | |
}); | |
if (ArgHasName) | |
addOCLKernelArgumentMetadata(Context, KernelMD, | |
SPIR_MD_KERNEL_ARG_NAME, BF, | |
[=](SPIRVFunctionParameter *Arg){ | |
return MDString::get(*Context, Arg->getName()); | |
}); | |
} | |
// Generate metadata for reqd_work_group_size | |
if (auto EM = BF->getExecutionMode(ExecutionModeLocalSize)) { | |
KernelMD.push_back(getMDNodeStringIntVec(Context, | |
kSPIR2MD::WGSize, EM->getLiterals())); | |
} | |
// Generate metadata for work_group_size_hint | |
if (auto EM = BF->getExecutionMode(ExecutionModeLocalSizeHint)) { | |
KernelMD.push_back(getMDNodeStringIntVec(Context, | |
kSPIR2MD::WGSizeHint, EM->getLiterals())); | |
} | |
// Generate metadata for vec_type_hint | |
if (auto EM = BF->getExecutionMode(ExecutionModeVecTypeHint)) { | |
std::vector<Metadata*> MetadataVec; | |
MetadataVec.push_back(MDString::get(*Context, kSPIR2MD::VecTyHint)); | |
Type *VecHintTy = decodeVecTypeHint(*Context, EM->getLiterals()[0]); | |
assert(VecHintTy); | |
MetadataVec.push_back(ValueAsMetadata::get(UndefValue::get(VecHintTy))); | |
MetadataVec.push_back( | |
ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), | |
1))); | |
KernelMD.push_back(MDNode::get(*Context, MetadataVec)); | |
} | |
llvm::MDNode *Node = MDNode::get(*Context, KernelMD); | |
KernelMDs->addOperand(Node); | |
} | |
return true; | |
} | |
bool | |
SPIRVToLLVM::transAlign(SPIRVValue *BV, Value *V) { | |
if (auto AL = dyn_cast<AllocaInst>(V)) { | |
SPIRVWord Align = 0; | |
if (BV->hasAlignment(&Align)) | |
AL->setAlignment(Align); | |
return true; | |
} | |
if (auto GV = dyn_cast<GlobalVariable>(V)) { | |
SPIRVWord Align = 0; | |
if (BV->hasAlignment(&Align)) | |
GV->setAlignment(Align); | |
return true; | |
} | |
return true; | |
} | |
void | |
SPIRVToLLVM::transOCLVectorLoadStore(std::string& UnmangledName, | |
std::vector<SPIRVWord> &BArgs) { | |
if (UnmangledName.find("vload") == 0 && | |
UnmangledName.find("n") != std::string::npos) { | |
if (BArgs.back() != 1) { | |
std::stringstream SS; | |
SS << BArgs.back(); | |
UnmangledName.replace(UnmangledName.find("n"), 1, SS.str()); | |
} else { | |
UnmangledName.erase(UnmangledName.find("n"), 1); | |
} | |
BArgs.pop_back(); | |
} else if (UnmangledName.find("vstore") == 0) { | |
if (UnmangledName.find("n") != std::string::npos) { | |
auto T = BM->getValueType(BArgs[0]); | |
if (T->isTypeVector()) { | |
auto W = T->getVectorComponentCount(); | |
std::stringstream SS; | |
SS << W; | |
UnmangledName.replace(UnmangledName.find("n"), 1, SS.str()); | |
} else { | |
UnmangledName.erase(UnmangledName.find("n"), 1); | |
} | |
} | |
if (UnmangledName.find("_r") != std::string::npos) { | |
UnmangledName.replace(UnmangledName.find("_r"), 2, std::string("_") + | |
SPIRSPIRVFPRoundingModeMap::rmap(static_cast<SPIRVFPRoundingModeKind>( | |
BArgs.back()))); | |
BArgs.pop_back(); | |
} | |
} | |
} | |
// printf is not mangled. The function type should have just one argument. | |
// read_image*: the second argument should be mangled as sampler. | |
Instruction * | |
SPIRVToLLVM::transOCLBuiltinFromExtInst(SPIRVExtInst *BC, BasicBlock *BB) { | |
assert(BB && "Invalid BB"); | |
std::string MangledName; | |
SPIRVWord EntryPoint = BC->getExtOp(); | |
SPIRVExtInstSetKind Set = BM->getBuiltinSet(BC->getExtSetId()); | |
bool IsVarArg = false; | |
bool IsPrintf = false; | |
std::string UnmangledName; | |
auto BArgs = BC->getArguments(); | |
assert (Set == SPIRVEIS_OpenCL && "Not OpenCL extended instruction"); | |
if (EntryPoint == OpenCLLIB::Printf) | |
IsPrintf = true; | |
else { | |
UnmangledName = OCLExtOpMap::map(static_cast<OCLExtOpKind>( | |
EntryPoint)); | |
} | |
SPIRVDBG(spvdbgs() << "[transOCLBuiltinFromExtInst] OrigUnmangledName: " << | |
UnmangledName << '\n'); | |
transOCLVectorLoadStore(UnmangledName, BArgs); | |
std::vector<Type *> ArgTypes = transTypeVector(BC->getValueTypes(BArgs)); | |
if (IsPrintf) { | |
MangledName = "printf"; | |
IsVarArg = true; | |
ArgTypes.resize(1); | |
} else if (UnmangledName.find("read_image") == 0) { | |
auto ModifiedArgTypes = ArgTypes; | |
ModifiedArgTypes[1] = getOrCreateOpaquePtrType(M, "opencl.sampler_t"); | |
MangleOpenCLBuiltin(UnmangledName, ModifiedArgTypes, MangledName); | |
} else { | |
MangleOpenCLBuiltin(UnmangledName, ArgTypes, MangledName); | |
} | |
SPIRVDBG(spvdbgs() << "[transOCLBuiltinFromExtInst] ModifiedUnmangledName: " << | |
UnmangledName << " MangledName: " << MangledName << '\n'); | |
FunctionType *FT = FunctionType::get( | |
transType(BC->getType()), | |
ArgTypes, | |
IsVarArg); | |
Function *F = M->getFunction(MangledName); | |
if (!F) { | |
F = Function::Create(FT, | |
GlobalValue::ExternalLinkage, | |
MangledName, | |
M); | |
F->setCallingConv(CallingConv::SPIR_FUNC); | |
if (isFuncNoUnwind()) | |
F->addFnAttr(Attribute::NoUnwind); | |
} | |
auto Args = transValue(BC->getValues(BArgs), F, BB); | |
SPIRVDBG(dbgs() << "[transOCLBuiltinFromExtInst] Function: " << *F << | |
", Args: "; | |
for (auto &I:Args) dbgs() << *I << ", "; dbgs() << '\n'); | |
CallInst *Call = CallInst::Create(F, | |
Args, | |
BC->getName(), | |
BB); | |
setCallingConv(Call); | |
addFnAttr(Context, Call, Attribute::NoUnwind); | |
return transOCLBuiltinPostproc(BC, Call, BB, UnmangledName); | |
} | |
CallInst * | |
SPIRVToLLVM::transOCLBarrier(BasicBlock *BB, SPIRVWord ExecScope, | |
SPIRVWord MemSema, SPIRVWord MemScope) { | |
SPIRVWord Ver = 0; | |
BM->getSourceLanguage(&Ver); | |
Type* Int32Ty = Type::getInt32Ty(*Context); | |
Type* VoidTy = Type::getVoidTy(*Context); | |
std::string FuncName; | |
SmallVector<Type *, 2> ArgTy; | |
SmallVector<Value *, 2> Arg; | |
Constant *MemFenceFlags = | |
ConstantInt::get(Int32Ty, rmapBitMask<OCLMemFenceMap>(MemSema)); | |
FuncName = (ExecScope == ScopeWorkgroup) ? kOCLBuiltinName::WorkGroupBarrier | |
: kOCLBuiltinName::SubGroupBarrier; | |
if (ExecScope == ScopeWorkgroup && Ver > 0 && Ver <= kOCLVer::CL12) { | |
FuncName = kOCLBuiltinName::Barrier; | |
ArgTy.push_back(Int32Ty); | |
Arg.push_back(MemFenceFlags); | |
} else { | |
Constant *Scope = ConstantInt::get(Int32Ty, OCLMemScopeMap::rmap( | |
static_cast<spv::Scope>(MemScope))); | |
ArgTy.append(2, Int32Ty); | |
Arg.push_back(MemFenceFlags); | |
Arg.push_back(Scope); | |
} | |
std::string MangledName; | |
MangleOpenCLBuiltin(FuncName, ArgTy, MangledName); | |
Function *Func = M->getFunction(MangledName); | |
if (!Func) { | |
FunctionType *FT = FunctionType::get(VoidTy, ArgTy, false); | |
Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); | |
Func->setCallingConv(CallingConv::SPIR_FUNC); | |
if (isFuncNoUnwind()) | |
Func->addFnAttr(Attribute::NoUnwind); | |
} | |
return CallInst::Create(Func, Arg, "", BB); | |
} | |
CallInst * | |
SPIRVToLLVM::transOCLMemFence(BasicBlock *BB, | |
SPIRVWord MemSema, SPIRVWord MemScope) { | |
SPIRVWord Ver = 0; | |
BM->getSourceLanguage(&Ver); | |
Type* Int32Ty = Type::getInt32Ty(*Context); | |
Type* VoidTy = Type::getVoidTy(*Context); | |
std::string FuncName; | |
SmallVector<Type *, 3> ArgTy; | |
SmallVector<Value *, 3> Arg; | |
Constant *MemFenceFlags = | |
ConstantInt::get(Int32Ty, rmapBitMask<OCLMemFenceMap>(MemSema)); | |
if (Ver > 0 && Ver <= kOCLVer::CL12) { | |
FuncName = kOCLBuiltinName::MemFence; | |
ArgTy.push_back(Int32Ty); | |
Arg.push_back(MemFenceFlags); | |
} else { | |
Constant *Order = | |
ConstantInt::get(Int32Ty, mapSPIRVMemOrderToOCL(MemSema)); | |
Constant *Scope = ConstantInt::get(Int32Ty, OCLMemScopeMap::rmap( | |
static_cast<spv::Scope>(MemScope))); | |
FuncName = kOCLBuiltinName::AtomicWorkItemFence; | |
ArgTy.append(3, Int32Ty); | |
Arg.push_back(MemFenceFlags); | |
Arg.push_back(Order); | |
Arg.push_back(Scope); | |
} | |
std::string MangledName; | |
MangleOpenCLBuiltin(FuncName, ArgTy, MangledName); | |
Function *Func = M->getFunction(MangledName); | |
if (!Func) { | |
FunctionType *FT = FunctionType::get(VoidTy, ArgTy, false); | |
Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); | |
Func->setCallingConv(CallingConv::SPIR_FUNC); | |
if (isFuncNoUnwind()) | |
Func->addFnAttr(Attribute::NoUnwind); | |
} | |
return CallInst::Create(Func, Arg, "", BB); | |
} | |
Instruction * | |
SPIRVToLLVM::transOCLBarrierFence(SPIRVInstruction *MB, BasicBlock *BB) { | |
assert(BB && "Invalid BB"); | |
std::string FuncName; | |
auto getIntVal = [](SPIRVValue *value){ | |
return static_cast<SPIRVConstant*>(value)->getZExtIntValue(); | |
}; | |
CallInst* Call = nullptr; | |
if (MB->getOpCode() == OpMemoryBarrier) { | |
auto MemB = static_cast<SPIRVMemoryBarrier*>(MB); | |
SPIRVWord MemScope = getIntVal(MemB->getOpValue(0)); | |
SPIRVWord MemSema = getIntVal(MemB->getOpValue(1)); | |
Call = transOCLMemFence(BB, MemSema, MemScope); | |
} else if (MB->getOpCode() == OpControlBarrier) { | |
auto CtlB = static_cast<SPIRVControlBarrier*>(MB); | |
SPIRVWord ExecScope = getIntVal(CtlB->getExecScope()); | |
SPIRVWord MemSema = getIntVal(CtlB->getMemSemantic()); | |
SPIRVWord MemScope = getIntVal(CtlB->getMemScope()); | |
Call = transOCLBarrier(BB, ExecScope, MemSema, MemScope); | |
} else { | |
llvm_unreachable("Invalid instruction"); | |
} | |
setName(Call, MB); | |
setAttrByCalledFunc(Call); | |
SPIRVDBG(spvdbgs() << "[transBarrier] " << *MB << " -> "; | |
dbgs() << *Call << '\n';) | |
return Call; | |
} | |
// SPIR-V only contains language version. Use OpenCL language version as | |
// SPIR version. | |
bool | |
SPIRVToLLVM::transSourceLanguage() { | |
SPIRVWord Ver = 0; | |
SourceLanguage Lang = BM->getSourceLanguage(&Ver); | |
assert((Lang == SourceLanguageOpenCL_C || | |
Lang == SourceLanguageOpenCL_CPP) && "Unsupported source language"); | |
unsigned short Major = 0; | |
unsigned char Minor = 0; | |
unsigned char Rev = 0; | |
std::tie(Major, Minor, Rev) = decodeOCLVer(Ver); | |
SPIRVMDBuilder Builder(*M); | |
Builder.addNamedMD(kSPIRVMD::Source) | |
.addOp() | |
.add(Lang) | |
.add(Ver) | |
.done(); | |
// ToDo: Phasing out usage of old SPIR metadata | |
if (Ver <= kOCLVer::CL12) | |
addOCLVersionMetadata(Context, M, kSPIR2MD::SPIRVer, 1, 2); | |
else | |
addOCLVersionMetadata(Context, M, kSPIR2MD::SPIRVer, 2, 0); | |
addOCLVersionMetadata(Context, M, kSPIR2MD::OCLVer, Major, Minor); | |
return true; | |
} | |
bool | |
SPIRVToLLVM::transSourceExtension() { | |
auto ExtSet = rmap<OclExt::Kind>(BM->getExtension()); | |
auto CapSet = rmap<OclExt::Kind>(BM->getCapability()); | |
ExtSet.insert(CapSet.begin(), CapSet.end()); | |
auto OCLExtensions = map<std::string>(ExtSet); | |
std::set<std::string> OCLOptionalCoreFeatures; | |
static const char *OCLOptCoreFeatureNames[] = { | |
"cl_images", "cl_doubles", | |
}; | |
for (auto &I : OCLOptCoreFeatureNames) { | |
auto Loc = OCLExtensions.find(I); | |
if (Loc != OCLExtensions.end()) { | |
OCLExtensions.erase(Loc); | |
OCLOptionalCoreFeatures.insert(I); | |
} | |
} | |
addNamedMetadataStringSet(Context, M, kSPIR2MD::Extensions, OCLExtensions); | |
addNamedMetadataStringSet(Context, M, kSPIR2MD::OptFeatures, | |
OCLOptionalCoreFeatures); | |
return true; | |
} | |
// If the argument is unsigned return uconvert*, otherwise return convert*. | |
std::string | |
SPIRVToLLVM::getOCLConvertBuiltinName(SPIRVInstruction* BI) { | |
auto OC = BI->getOpCode(); | |
assert(isCvtOpCode(OC) && "Not convert instruction"); | |
auto U = static_cast<SPIRVUnary *>(BI); | |
std::string Name; | |
if (isCvtFromUnsignedOpCode(OC)) | |
Name = "u"; | |
Name += "convert_"; | |
Name += mapSPIRVTypeToOCLType(U->getType(), | |
!isCvtToUnsignedOpCode(OC)); | |
SPIRVFPRoundingModeKind Rounding; | |
if (U->isSaturatedConversion()) | |
Name += "_sat"; | |
if (U->hasFPRoundingMode(&Rounding)) { | |
Name += "_"; | |
Name += SPIRSPIRVFPRoundingModeMap::rmap(Rounding); | |
} | |
return Name; | |
} | |
//Check Address Space of the Pointer Type | |
std::string | |
SPIRVToLLVM::getOCLGenericCastToPtrName(SPIRVInstruction* BI) { | |
auto GenericCastToPtrInst = BI->getType()->getPointerStorageClass(); | |
switch (GenericCastToPtrInst) { | |
case StorageClassCrossWorkgroup: | |
return std::string(kOCLBuiltinName::ToGlobal); | |
case StorageClassWorkgroup: | |
return std::string(kOCLBuiltinName::ToLocal); | |
case StorageClassFunction: | |
return std::string(kOCLBuiltinName::ToPrivate); | |
default: | |
llvm_unreachable("Invalid address space"); | |
return ""; | |
} | |
} | |
llvm::GlobalValue::LinkageTypes | |
SPIRVToLLVM::transLinkageType(const SPIRVValue* V) { | |
if (V->getLinkageType() == LinkageTypeInternal) { | |
return GlobalValue::InternalLinkage; | |
} | |
else if (V->getLinkageType() == LinkageTypeImport) { | |
// Function declaration | |
if (V->getOpCode() == OpFunction) { | |
if (static_cast<const SPIRVFunction*>(V)->getNumBasicBlock() == 0) | |
return GlobalValue::ExternalLinkage; | |
} | |
// Variable declaration | |
if (V->getOpCode() == OpVariable) { | |
if (static_cast<const SPIRVVariable*>(V)->getInitializer() == 0) | |
return GlobalValue::ExternalLinkage; | |
} | |
// Definition | |
return GlobalValue::AvailableExternallyLinkage; | |
} | |
else {// LinkageTypeExport | |
if (V->getOpCode() == OpVariable) { | |
if (static_cast<const SPIRVVariable*>(V)->getInitializer() == 0 ) | |
// Tentative definition | |
return GlobalValue::CommonLinkage; | |
} | |
return GlobalValue::ExternalLinkage; | |
} | |
} | |
Instruction *SPIRVToLLVM::transOCLAllAny(SPIRVInstruction *I, BasicBlock *BB) { | |
CallInst *CI = cast<CallInst>(transSPIRVBuiltinFromInst(I, BB)); | |
AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); | |
return cast<Instruction>(mapValue( | |
I, mutateCallInstOCL( | |
M, CI, | |
[=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { | |
Type *Int32Ty = Type::getInt32Ty(*Context); | |
auto OldArg = CI->getOperand(0); | |
auto NewArgTy = VectorType::get( | |
Int32Ty, OldArg->getType()->getVectorNumElements()); | |
auto NewArg = | |
CastInst::CreateSExtOrBitCast(OldArg, NewArgTy, "", CI); | |
Args[0] = NewArg; | |
RetTy = Int32Ty; | |
return CI->getCalledFunction()->getName(); | |
}, | |
[=](CallInst *NewCI) -> Instruction * { | |
return CastInst::CreateTruncOrBitCast( | |
NewCI, Type::getInt1Ty(*Context), "", NewCI->getNextNode()); | |
}, | |
&Attrs))); | |
} | |
Instruction *SPIRVToLLVM::transOCLRelational(SPIRVInstruction *I, BasicBlock *BB) { | |
CallInst *CI = cast<CallInst>(transSPIRVBuiltinFromInst(I, BB)); | |
AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); | |
return cast<Instruction>(mapValue( | |
I, mutateCallInstOCL( | |
M, CI, | |
[=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { | |
Type *IntTy = Type::getInt32Ty(*Context); | |
RetTy = IntTy; | |
if (CI->getType()->isVectorTy()) { | |
if (cast<VectorType>(CI->getOperand(0)->getType()) | |
->getElementType() | |
->isDoubleTy()) | |
IntTy = Type::getInt64Ty(*Context); | |
if (cast<VectorType>(CI->getOperand(0)->getType()) | |
->getElementType() | |
->isHalfTy()) | |
IntTy = Type::getInt16Ty(*Context); | |
RetTy = VectorType::get(IntTy, | |
CI->getType()->getVectorNumElements()); | |
} | |
return CI->getCalledFunction()->getName(); | |
}, | |
[=](CallInst *NewCI) -> Instruction * { | |
Type *RetTy = Type::getInt1Ty(*Context); | |
if (NewCI->getType()->isVectorTy()) | |
RetTy = | |
VectorType::get(Type::getInt1Ty(*Context), | |
NewCI->getType()->getVectorNumElements()); | |
return CastInst::CreateTruncOrBitCast(NewCI, RetTy, "", | |
NewCI->getNextNode()); | |
}, | |
&Attrs))); | |
} | |
} | |
bool | |
llvm::ReadSPIRV(LLVMContext &C, std::istream &IS, Module *&M, | |
std::string &ErrMsg) { | |
M = new Module("", C); | |
std::unique_ptr<SPIRVModule> BM(SPIRVModule::createSPIRVModule()); | |
IS >> *BM; | |
SPIRVToLLVM BTL(M, BM.get()); | |
bool Succeed = true; | |
if (!BTL.translate()) { | |
BM->getError(ErrMsg); | |
Succeed = false; | |
} | |
legacy::PassManager PassMgr; | |
PassMgr.add(createSPIRVToOCL20()); | |
PassMgr.add(createOCL20To12()); | |
PassMgr.run(*M); | |
if (DbgSaveTmpLLVM) | |
dumpLLVM(M, DbgTmpLLVMFileName); | |
if (!Succeed) { | |
delete M; | |
M = nullptr; | |
} | |
return Succeed; | |
} |