Google Git
Sign in
android / platform / external / spirv-llvm / ffc5a8a112667ac2e8e9194d603fc8b77e3fa7c8 / . / lib / SPIRV / SPIRVReader.cpp
blob: 65c93cee8000c87a46b21266de3280b146875bab [file] [log] [blame]
//===- 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);
}
#if 0
// this function is currently unneeded
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);
}
#endif
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, "", DICompileUnit::DebugEmissionKind::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 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));
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)) ?
GlobalValue::UnnamedAddr::Global :
GlobalValue::UnnamedAddr::None);
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();
(void) Set;
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;
}