1//===- SPIRVReader.cpp - Converts SPIR-V to LLVM ----------------*- C++ -*-===// 2// 3// The LLVM/SPIR-V Translator 4// 5// This file is distributed under the University of Illinois Open Source 6// License. See LICENSE.TXT for details. 7// 8// Copyright (c) 2014 Advanced Micro Devices, Inc. All rights reserved. 9// 10// Permission is hereby granted, free of charge, to any person obtaining a 11// copy of this software and associated documentation files (the "Software"), 12// to deal with the Software without restriction, including without limitation 13// the rights to use, copy, modify, merge, publish, distribute, sublicense, 14// and/or sell copies of the Software, and to permit persons to whom the 15// Software is furnished to do so, subject to the following conditions: 16// 17// Redistributions of source code must retain the above copyright notice, 18// this list of conditions and the following disclaimers. 19// Redistributions in binary form must reproduce the above copyright notice, 20// this list of conditions and the following disclaimers in the documentation 21// and/or other materials provided with the distribution. 22// Neither the names of Advanced Micro Devices, Inc., nor the names of its 23// contributors may be used to endorse or promote products derived from this 24// Software without specific prior written permission. 25// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 26// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 27// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE 28// CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 29// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, 30// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH 31// THE SOFTWARE. 32// 33//===----------------------------------------------------------------------===// 34/// \file 35/// 36/// This file implements conversion of SPIR-V binary to LLVM IR. 37/// 38//===----------------------------------------------------------------------===// 39#include "SPIRVUtil.h" 40#include "SPIRVType.h" 41#include "SPIRVValue.h" 42#include "SPIRVModule.h" 43#include "SPIRVFunction.h" 44#include "SPIRVBasicBlock.h" 45#include "SPIRVInstruction.h" 46#include "SPIRVExtInst.h" 47#include "SPIRVInternal.h" 48#include "SPIRVMDBuilder.h" 49#include "OCLUtil.h" 50 51#include "llvm/ADT/DenseMap.h" 52#include "llvm/ADT/StringSwitch.h" 53#include "llvm/IR/Constants.h" 54#include "llvm/IR/DerivedTypes.h" 55#include "llvm/IR/DIBuilder.h" 56#include "llvm/IR/Instructions.h" 57#include "llvm/IR/Metadata.h" 58#include "llvm/IR/Module.h" 59#include "llvm/IR/Operator.h" 60#include "llvm/IR/Type.h" 61#include "llvm/IR/LegacyPassManager.h" 62#include "llvm/Support/Casting.h" 63#include "llvm/Support/Debug.h" 64#include "llvm/Support/Dwarf.h" 65#include "llvm/Support/FileSystem.h" 66#include "llvm/Support/raw_ostream.h" 67#include "llvm/Support/CommandLine.h" 68 69#include <algorithm> 70#include <cstdlib> 71#include <functional> 72#include <fstream> 73#include <iostream> 74#include <iterator> 75#include <map> 76#include <set> 77#include <sstream> 78#include <string> 79 80#define DEBUG_TYPE "spirv" 81 82using namespace std; 83using namespace llvm; 84using namespace SPIRV; 85using namespace OCLUtil; 86 87namespace SPIRV{ 88 89cl::opt<bool> SPIRVEnableStepExpansion("spirv-expand-step", cl::init(true), 90 cl::desc("Enable expansion of OpenCL step and smoothstep function")); 91 92cl::opt<bool> SPIRVGenKernelArgNameMD("spirv-gen-kernel-arg-name-md", 93 cl::init(false), cl::desc("Enable generating OpenCL kernel argument name " 94 "metadata")); 95 96cl::opt<bool> SPIRVGenImgTypeAccQualPostfix("spirv-gen-image-type-acc-postfix", 97 cl::init(false), cl::desc("Enable generating access qualifier postfix" 98 " in OpenCL image type names")); 99 100// Prefix for placeholder global variable name. 101const char* kPlaceholderPrefix = "placeholder."; 102 103// Save the translated LLVM before validation for debugging purpose. 104static bool DbgSaveTmpLLVM = true; 105static const char *DbgTmpLLVMFileName = "_tmp_llvmbil.ll"; 106 107typedef std::pair < unsigned, AttributeSet > AttributeWithIndex; 108 109static bool 110isOpenCLKernel(SPIRVFunction *BF) { 111 return BF->getModule()->isEntryPoint(ExecutionModelKernel, BF->getId()); 112} 113 114static void 115dumpLLVM(Module *M, const std::string &FName) { 116 std::error_code EC; 117 raw_fd_ostream FS(FName, EC, sys::fs::F_None); 118 if (EC) { 119 FS << *M; 120 FS.close(); 121 } 122} 123 124static MDNode* 125getMDNodeStringIntVec(LLVMContext *Context, const std::string& Str, 126 const std::vector<SPIRVWord>& IntVals) { 127 std::vector<Metadata*> ValueVec; 128 ValueVec.push_back(MDString::get(*Context, Str)); 129 for (auto &I:IntVals) 130 ValueVec.push_back(ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), I))); 131 return MDNode::get(*Context, ValueVec); 132} 133 134static MDNode* 135getMDTwoInt(LLVMContext *Context, unsigned Int1, unsigned Int2) { 136 std::vector<Metadata*> ValueVec; 137 ValueVec.push_back(ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), Int1))); 138 ValueVec.push_back(ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), Int2))); 139 return MDNode::get(*Context, ValueVec); 140} 141 142#if 0 143// this function is currently unneeded 144static MDNode* 145getMDString(LLVMContext *Context, const std::string& Str) { 146 std::vector<Metadata*> ValueVec; 147 if (!Str.empty()) 148 ValueVec.push_back(MDString::get(*Context, Str)); 149 return MDNode::get(*Context, ValueVec); 150} 151#endif 152 153static void 154addOCLVersionMetadata(LLVMContext *Context, Module *M, 155 const std::string &MDName, unsigned Major, unsigned Minor) { 156 NamedMDNode *NamedMD = M->getOrInsertNamedMetadata(MDName); 157 NamedMD->addOperand(getMDTwoInt(Context, Major, Minor)); 158} 159 160static void 161addNamedMetadataStringSet(LLVMContext *Context, Module *M, 162 const std::string &MDName, const std::set<std::string> &StrSet) { 163 NamedMDNode *NamedMD = M->getOrInsertNamedMetadata(MDName); 164 std::vector<Metadata*> ValueVec; 165 for (auto &&Str : StrSet) { 166 ValueVec.push_back(MDString::get(*Context, Str)); 167 } 168 NamedMD->addOperand(MDNode::get(*Context, ValueVec)); 169} 170 171static void 172addOCLKernelArgumentMetadata(LLVMContext *Context, 173 std::vector<llvm::Metadata*> &KernelMD, const std::string &MDName, 174 SPIRVFunction *BF, std::function<Metadata *(SPIRVFunctionParameter *)>Func){ 175 std::vector<Metadata*> ValueVec; 176 ValueVec.push_back(MDString::get(*Context, MDName)); 177 BF->foreachArgument([&](SPIRVFunctionParameter *Arg) { 178 ValueVec.push_back(Func(Arg)); 179 }); 180 KernelMD.push_back(MDNode::get(*Context, ValueVec)); 181} 182 183class SPIRVToLLVMDbgTran { 184public: 185 SPIRVToLLVMDbgTran(SPIRVModule *TBM, Module *TM) 186 :BM(TBM), M(TM), SpDbg(BM), Builder(*M){ 187 Enable = BM->hasDebugInfo(); 188 } 189 190 void createCompileUnit() { 191 if (!Enable) 192 return; 193 auto File = SpDbg.getEntryPointFileStr(ExecutionModelKernel, 0); 194 std::string BaseName; 195 std::string Path; 196 splitFileName(File, BaseName, Path); 197 Builder.createCompileUnit(dwarf::DW_LANG_C99, 198 BaseName, Path, "spirv", false, "", 0, "", DICompileUnit::DebugEmissionKind::LineTablesOnly); 199 } 200 201 void addDbgInfoVersion() { 202 if (!Enable) 203 return; 204 M->addModuleFlag(Module::Warning, "Dwarf Version", 205 dwarf::DWARF_VERSION); 206 M->addModuleFlag(Module::Warning, "Debug Info Version", 207 DEBUG_METADATA_VERSION); 208 } 209 210 DIFile* getDIFile(const std::string &FileName){ 211 return getOrInsert(FileMap, FileName, [=](){ 212 std::string BaseName; 213 std::string Path; 214 splitFileName(FileName, BaseName, Path); 215 if (!BaseName.empty()) 216 return Builder.createFile(BaseName, Path); 217 else 218 return Builder.createFile("","");//DIFile(); 219 }); 220 } 221 222 DISubprogram* getDISubprogram(SPIRVFunction *SF, Function *F){ 223 return getOrInsert(FuncMap, F, [=](){ 224 auto DF = getDIFile(SpDbg.getFunctionFileStr(SF)); 225 auto FN = F->getName(); 226 auto LN = SpDbg.getFunctionLineNo(SF); 227 Metadata *Args[] = {Builder.createUnspecifiedType("")}; 228 return Builder.createFunction(static_cast<DIScope*>(DF), FN, FN, DF, LN, 229 Builder.createSubroutineType(Builder.getOrCreateTypeArray(Args)), 230 Function::isInternalLinkage(F->getLinkage()), 231 true, LN); 232 }); 233 } 234 235 void transDbgInfo(SPIRVValue *SV, Value *V) { 236 if (!Enable || !SV->hasLine()) 237 return; 238 if (auto I = dyn_cast<Instruction>(V)) { 239 assert(SV->isInst() && "Invalid instruction"); 240 auto SI = static_cast<SPIRVInstruction *>(SV); 241 assert(SI->getParent() && 242 SI->getParent()->getParent() && 243 "Invalid instruction"); 244 auto Line = SV->getLine(); 245 I->setDebugLoc(DebugLoc::get(Line->getLine(), Line->getColumn(), 246 getDISubprogram(SI->getParent()->getParent(), 247 I->getParent()->getParent()))); 248 } 249 } 250 251 void finalize() { 252 if (!Enable) 253 return; 254 Builder.finalize(); 255 } 256 257private: 258 SPIRVModule *BM; 259 Module *M; 260 SPIRVDbgInfo SpDbg; 261 DIBuilder Builder; 262 bool Enable; 263 std::unordered_map<std::string, DIFile*> FileMap; 264 std::unordered_map<Function *, DISubprogram*> FuncMap; 265 266 void splitFileName(const std::string &FileName, 267 std::string &BaseName, 268 std::string &Path) { 269 auto Loc = FileName.find_last_of("/\\"); 270 if (Loc != std::string::npos) { 271 BaseName = FileName.substr(Loc + 1); 272 Path = FileName.substr(0, Loc); 273 } else { 274 BaseName = FileName; 275 Path = "."; 276 } 277 } 278}; 279 280class SPIRVToLLVM { 281public: 282 SPIRVToLLVM(Module *LLVMModule, SPIRVModule *TheSPIRVModule) 283 :M(LLVMModule), BM(TheSPIRVModule), DbgTran(BM, M){ 284 assert(M); 285 Context = &M->getContext(); 286 } 287 288 std::string getOCLBuiltinName(SPIRVInstruction* BI); 289 std::string getOCLConvertBuiltinName(SPIRVInstruction *BI); 290 std::string getOCLGenericCastToPtrName(SPIRVInstruction *BI); 291 292 Type *transType(SPIRVType *BT, bool IsClassMember = false); 293 std::string transTypeToOCLTypeName(SPIRVType *BT, bool IsSigned = true); 294 std::vector<Type *> transTypeVector(const std::vector<SPIRVType *>&); 295 bool translate(); 296 bool transAddressingModel(); 297 298 Value *transValue(SPIRVValue *, Function *F, BasicBlock *, 299 bool CreatePlaceHolder = true); 300 Value *transValueWithoutDecoration(SPIRVValue *, Function *F, BasicBlock *, 301 bool CreatePlaceHolder = true); 302 bool transDecoration(SPIRVValue *, Value *); 303 bool transAlign(SPIRVValue *, Value *); 304 Instruction *transOCLBuiltinFromExtInst(SPIRVExtInst *BC, BasicBlock *BB); 305 std::vector<Value *> transValue(const std::vector<SPIRVValue *>&, Function *F, 306 BasicBlock *); 307 Function *transFunction(SPIRVFunction *F); 308 bool transFPContractMetadata(); 309 bool transKernelMetadata(); 310 bool transNonTemporalMetadata(Instruction *I); 311 bool transSourceLanguage(); 312 bool transSourceExtension(); 313 void transGeneratorMD(); 314 Value *transConvertInst(SPIRVValue* BV, Function* F, BasicBlock* BB); 315 Instruction *transBuiltinFromInst(const std::string& FuncName, 316 SPIRVInstruction* BI, BasicBlock* BB); 317 Instruction *transOCLBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB); 318 Instruction *transSPIRVBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB); 319 Instruction *transOCLBarrierFence(SPIRVInstruction* BI, BasicBlock *BB); 320 void transOCLVectorLoadStore(std::string& UnmangledName, 321 std::vector<SPIRVWord> &BArgs); 322 323 /// Post-process translated LLVM module for OpenCL. 324 bool postProcessOCL(); 325 326 /// \brief Post-process OpenCL builtin functions returning struct type. 327 /// 328 /// Some OpenCL builtin functions are translated to SPIR-V instructions with 329 /// struct type result, e.g. NDRange creation functions. Such functions 330 /// need to be post-processed to return the struct through sret argument. 331 bool postProcessOCLBuiltinReturnStruct(Function *F); 332 333 /// \brief Post-process OpenCL builtin functions having block argument. 334 /// 335 /// These functions are translated to functions with function pointer type 336 /// argument first, then post-processed to have block argument. 337 bool postProcessOCLBuiltinWithFuncPointer(Function *F, 338 Function::arg_iterator I); 339 340 /// \brief Post-process OpenCL builtin functions having array argument. 341 /// 342 /// These functions are translated to functions with array type argument 343 /// first, then post-processed to have pointer arguments. 344 bool postProcessOCLBuiltinWithArrayArguments(Function *F, 345 const std::string &DemangledName); 346 347 /// \brief Post-process OpImageSampleExplicitLod. 348 /// sampled_image = __spirv_SampledImage__(image, sampler); 349 /// return __spirv_ImageSampleExplicitLod__(sampled_image, image_operands, 350 /// ...); 351 /// => 352 /// read_image(image, sampler, ...) 353 /// \return transformed call instruction. 354 Instruction *postProcessOCLReadImage(SPIRVInstruction *BI, CallInst *CI, 355 const std::string &DemangledName); 356 357 /// \brief Post-process OpImageWrite. 358 /// return write_image(image, coord, color, image_operands, ...); 359 /// => 360 /// write_image(image, coord, ..., color) 361 /// \return transformed call instruction. 362 CallInst *postProcessOCLWriteImage(SPIRVInstruction *BI, CallInst *CI, 363 const std::string &DemangledName); 364 365 /// \brief Post-process OpBuildNDRange. 366 /// OpBuildNDRange GlobalWorkSize, LocalWorkSize, GlobalWorkOffset 367 /// => 368 /// call ndrange_XD(GlobalWorkOffset, GlobalWorkSize, LocalWorkSize) 369 /// \return transformed call instruction. 370 CallInst *postProcessOCLBuildNDRange(SPIRVInstruction *BI, CallInst *CI, 371 const std::string &DemangledName); 372 373 /// \brief Expand OCL builtin functions with scalar argument, e.g. 374 /// step, smoothstep. 375 /// gentype func (fp edge, gentype x) 376 /// => 377 /// gentype func (gentype edge, gentype x) 378 /// \return transformed call instruction. 379 CallInst *expandOCLBuiltinWithScalarArg(CallInst* CI, 380 const std::string &FuncName); 381 382 /// \brief Post-process OpGroupAll and OpGroupAny instructions translation. 383 /// i1 func (<n x i1> arg) 384 /// => 385 /// i32 func (<n x i32> arg) 386 /// \return transformed call instruction. 387 Instruction *postProcessGroupAllAny(CallInst *CI, 388 const std::string &DemangledName); 389 390 typedef DenseMap<SPIRVType *, Type *> SPIRVToLLVMTypeMap; 391 typedef DenseMap<SPIRVValue *, Value *> SPIRVToLLVMValueMap; 392 typedef DenseMap<SPIRVFunction *, Function *> SPIRVToLLVMFunctionMap; 393 typedef DenseMap<GlobalVariable *, SPIRVBuiltinVariableKind> BuiltinVarMap; 394 395 // A SPIRV value may be translated to a load instruction of a placeholder 396 // global variable. This map records load instruction of these placeholders 397 // which are supposed to be replaced by the real values later. 398 typedef std::map<SPIRVValue *, LoadInst*> SPIRVToLLVMPlaceholderMap; 399private: 400 Module *M; 401 BuiltinVarMap BuiltinGVMap; 402 LLVMContext *Context; 403 SPIRVModule *BM; 404 SPIRVToLLVMTypeMap TypeMap; 405 SPIRVToLLVMValueMap ValueMap; 406 SPIRVToLLVMFunctionMap FuncMap; 407 SPIRVToLLVMPlaceholderMap PlaceholderMap; 408 SPIRVToLLVMDbgTran DbgTran; 409 410 Type *mapType(SPIRVType *BT, Type *T) { 411 SPIRVDBG(dbgs() << *T << '\n';) 412 TypeMap[BT] = T; 413 return T; 414 } 415 416 // If a value is mapped twice, the existing mapped value is a placeholder, 417 // which must be a load instruction of a global variable whose name starts 418 // with kPlaceholderPrefix. 419 Value *mapValue(SPIRVValue *BV, Value *V) { 420 auto Loc = ValueMap.find(BV); 421 if (Loc != ValueMap.end()) { 422 if (Loc->second == V) 423 return V; 424 auto LD = dyn_cast<LoadInst>(Loc->second); 425 auto Placeholder = dyn_cast<GlobalVariable>(LD->getPointerOperand()); 426 assert (LD && Placeholder && 427 Placeholder->getName().startswith(kPlaceholderPrefix) && 428 "A value is translated twice"); 429 // Replaces placeholders for PHI nodes 430 LD->replaceAllUsesWith(V); 431 LD->dropAllReferences(); 432 LD->removeFromParent(); 433 Placeholder->dropAllReferences(); 434 Placeholder->removeFromParent(); 435 } 436 ValueMap[BV] = V; 437 return V; 438 } 439 440 bool isSPIRVBuiltinVariable(GlobalVariable *GV, 441 SPIRVBuiltinVariableKind *Kind = nullptr) { 442 auto Loc = BuiltinGVMap.find(GV); 443 if (Loc == BuiltinGVMap.end()) 444 return false; 445 if (Kind) 446 *Kind = Loc->second; 447 return true; 448 } 449 // OpenCL function always has NoUnwound attribute. 450 // Change this if it is no longer true. 451 bool isFuncNoUnwind() const { return true;} 452 bool isSPIRVCmpInstTransToLLVMInst(SPIRVInstruction *BI) const; 453 bool transOCLBuiltinsFromVariables(); 454 bool transOCLBuiltinFromVariable(GlobalVariable *GV, 455 SPIRVBuiltinVariableKind Kind); 456 MDString *transOCLKernelArgTypeName(SPIRVFunctionParameter *); 457 458 Value *mapFunction(SPIRVFunction *BF, Function *F) { 459 SPIRVDBG(spvdbgs() << "[mapFunction] " << *BF << " -> "; 460 dbgs() << *F << '\n';) 461 FuncMap[BF] = F; 462 return F; 463 } 464 465 Value *getTranslatedValue(SPIRVValue *BV); 466 Type *getTranslatedType(SPIRVType *BT); 467 468 SPIRVErrorLog &getErrorLog() { 469 return BM->getErrorLog(); 470 } 471 472 void setCallingConv(CallInst *Call) { 473 Function *F = Call->getCalledFunction(); 474 assert(F); 475 Call->setCallingConv(F->getCallingConv()); 476 } 477 478 void setAttrByCalledFunc(CallInst *Call); 479 Type *transFPType(SPIRVType* T); 480 BinaryOperator *transShiftLogicalBitwiseInst(SPIRVValue* BV, BasicBlock* BB, 481 Function* F); 482 void transFlags(llvm::Value* V); 483 Instruction *transCmpInst(SPIRVValue* BV, BasicBlock* BB, Function* F); 484 void transOCLBuiltinFromInstPreproc(SPIRVInstruction* BI, Type *&RetTy, 485 std::vector<SPIRVValue *> &Args); 486 Instruction* transOCLBuiltinPostproc(SPIRVInstruction* BI, 487 CallInst* CI, BasicBlock* BB, const std::string &DemangledName); 488 std::string transOCLImageTypeName(SPIRV::SPIRVTypeImage* ST); 489 std::string transOCLSampledImageTypeName(SPIRV::SPIRVTypeSampledImage* ST); 490 std::string transOCLPipeTypeName(SPIRV::SPIRVTypePipe* ST, 491 bool UseSPIRVFriendlyFormat = false, int PipeAccess = 0); 492 std::string transOCLPipeStorageTypeName(SPIRV::SPIRVTypePipeStorage* PST); 493 std::string transOCLImageTypeAccessQualifier(SPIRV::SPIRVTypeImage* ST); 494 std::string transOCLPipeTypeAccessQualifier(SPIRV::SPIRVTypePipe* ST); 495 496 Value *oclTransConstantSampler(SPIRV::SPIRVConstantSampler* BCS); 497 Value * oclTransConstantPipeStorage(SPIRV::SPIRVConstantPipeStorage* BCPS); 498 void setName(llvm::Value* V, SPIRVValue* BV); 499 void insertImageNameAccessQualifier(SPIRV::SPIRVTypeImage* ST, std::string &Name); 500 template<class Source, class Func> 501 bool foreachFuncCtlMask(Source, Func); 502 llvm::GlobalValue::LinkageTypes transLinkageType(const SPIRVValue* V); 503 Instruction *transOCLAllAny(SPIRVInstruction* BI, BasicBlock *BB); 504 Instruction *transOCLRelational(SPIRVInstruction* BI, BasicBlock *BB); 505 506 CallInst *transOCLBarrier(BasicBlock *BB, SPIRVWord ExecScope, 507 SPIRVWord MemSema, SPIRVWord MemScope); 508 509 CallInst *transOCLMemFence(BasicBlock *BB, 510 SPIRVWord MemSema, SPIRVWord MemScope); 511}; 512 513Type * 514SPIRVToLLVM::getTranslatedType(SPIRVType *BV){ 515 auto Loc = TypeMap.find(BV); 516 if (Loc != TypeMap.end()) 517 return Loc->second; 518 return nullptr; 519} 520 521Value * 522SPIRVToLLVM::getTranslatedValue(SPIRVValue *BV){ 523 auto Loc = ValueMap.find(BV); 524 if (Loc != ValueMap.end()) 525 return Loc->second; 526 return nullptr; 527} 528 529void 530SPIRVToLLVM::setAttrByCalledFunc(CallInst *Call) { 531 Function *F = Call->getCalledFunction(); 532 assert(F); 533 if (F->isIntrinsic()) { 534 return; 535 } 536 Call->setCallingConv(F->getCallingConv()); 537 Call->setAttributes(F->getAttributes()); 538} 539 540bool 541SPIRVToLLVM::transOCLBuiltinsFromVariables(){ 542 std::vector<GlobalVariable *> WorkList; 543 for (auto I = M->global_begin(), E = M->global_end(); I != E; ++I) { 544 SPIRVBuiltinVariableKind Kind; 545 auto I1 = static_cast<GlobalVariable*>(I); 546 if (!isSPIRVBuiltinVariable(I1, &Kind)) 547 continue; 548 if (!transOCLBuiltinFromVariable(I1, Kind)) 549 return false; 550 WorkList.push_back(I1); 551 } 552 for (auto &I:WorkList) { 553 I->dropAllReferences(); 554 I->removeFromParent(); 555 } 556 return true; 557} 558 559// For integer types shorter than 32 bit, unsigned/signedness can be inferred 560// from zext/sext attribute. 561MDString * 562SPIRVToLLVM::transOCLKernelArgTypeName(SPIRVFunctionParameter *Arg) { 563 auto Ty = Arg->isByVal() ? Arg->getType()->getPointerElementType() : 564 Arg->getType(); 565 return MDString::get(*Context, transTypeToOCLTypeName(Ty, !Arg->isZext())); 566} 567 568// Variable like GlobalInvolcationId[x] -> get_global_id(x). 569// Variable like WorkDim -> get_work_dim(). 570bool 571SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV, 572 SPIRVBuiltinVariableKind Kind) { 573 std::string FuncName = SPIRSPIRVBuiltinVariableMap::rmap(Kind); 574 std::string MangledName; 575 Type *ReturnTy = GV->getType()->getPointerElementType(); 576 bool IsVec = ReturnTy->isVectorTy(); 577 if (IsVec) 578 ReturnTy = cast<VectorType>(ReturnTy)->getElementType(); 579 std::vector<Type*> ArgTy; 580 if (IsVec) 581 ArgTy.push_back(Type::getInt32Ty(*Context)); 582 MangleOpenCLBuiltin(FuncName, ArgTy, MangledName); 583 Function *Func = M->getFunction(MangledName); 584 if (!Func) { 585 FunctionType *FT = FunctionType::get(ReturnTy, ArgTy, false); 586 Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); 587 Func->setCallingConv(CallingConv::SPIR_FUNC); 588 Func->addFnAttr(Attribute::NoUnwind); 589 Func->addFnAttr(Attribute::ReadNone); 590 } 591 std::vector<Instruction *> Deletes; 592 std::vector<Instruction *> Uses; 593 for (auto UI = GV->user_begin(), UE = GV->user_end(); UI != UE; ++UI) { 594 assert (isa<LoadInst>(*UI) && "Unsupported use"); 595 auto LD = dyn_cast<LoadInst>(*UI); 596 if (!IsVec) { 597 Uses.push_back(LD); 598 Deletes.push_back(LD); 599 continue; 600 } 601 for (auto LDUI = LD->user_begin(), LDUE = LD->user_end(); LDUI != LDUE; 602 ++LDUI) { 603 assert(isa<ExtractElementInst>(*LDUI) && "Unsupported use"); 604 auto EEI = dyn_cast<ExtractElementInst>(*LDUI); 605 Uses.push_back(EEI); 606 Deletes.push_back(EEI); 607 } 608 Deletes.push_back(LD); 609 } 610 for (auto &I:Uses) { 611 std::vector<Value *> Arg; 612 if (auto EEI = dyn_cast<ExtractElementInst>(I)) 613 Arg.push_back(EEI->getIndexOperand()); 614 auto Call = CallInst::Create(Func, Arg, "", I); 615 Call->takeName(I); 616 setAttrByCalledFunc(Call); 617 SPIRVDBG(dbgs() << "[transOCLBuiltinFromVariable] " << *I << " -> " << 618 *Call << '\n';) 619 I->replaceAllUsesWith(Call); 620 } 621 for (auto &I:Deletes) { 622 I->dropAllReferences(); 623 I->removeFromParent(); 624 } 625 return true; 626} 627 628Type * 629SPIRVToLLVM::transFPType(SPIRVType* T) { 630 switch(T->getFloatBitWidth()) { 631 case 16: return Type::getHalfTy(*Context); 632 case 32: return Type::getFloatTy(*Context); 633 case 64: return Type::getDoubleTy(*Context); 634 default: 635 llvm_unreachable("Invalid type"); 636 return nullptr; 637 } 638} 639 640std::string 641SPIRVToLLVM::transOCLImageTypeName(SPIRV::SPIRVTypeImage* ST) { 642 std::string Name = std::string(kSPR2TypeName::OCLPrefix) 643 + rmap<std::string>(ST->getDescriptor()); 644 if (SPIRVGenImgTypeAccQualPostfix) 645 SPIRVToLLVM::insertImageNameAccessQualifier(ST, Name); 646 return Name; 647} 648 649std::string 650SPIRVToLLVM::transOCLSampledImageTypeName(SPIRV::SPIRVTypeSampledImage* ST) { 651 return getSPIRVTypeName(kSPIRVTypeName::SampledImg, 652 getSPIRVImageTypePostfixes(getSPIRVImageSampledTypeName( 653 ST->getImageType()->getSampledType()), 654 ST->getImageType()->getDescriptor(), 655 ST->getImageType()->getAccessQualifier())); 656} 657 658std::string 659SPIRVToLLVM::transOCLPipeTypeName(SPIRV::SPIRVTypePipe* PT, 660 bool UseSPIRVFriendlyFormat, int PipeAccess){ 661 if (!UseSPIRVFriendlyFormat) 662 return kSPR2TypeName::Pipe; 663 else 664 return std::string(kSPIRVTypeName::PrefixAndDelim) 665 + kSPIRVTypeName::Pipe 666 + kSPIRVTypeName::Delimiter 667 + kSPIRVTypeName::PostfixDelim 668 + PipeAccess; 669} 670 671std::string 672SPIRVToLLVM::transOCLPipeStorageTypeName(SPIRV::SPIRVTypePipeStorage* PST) { 673 return std::string(kSPIRVTypeName::PrefixAndDelim) 674 + kSPIRVTypeName::PipeStorage; 675} 676 677Type * 678SPIRVToLLVM::transType(SPIRVType *T, bool IsClassMember) { 679 auto Loc = TypeMap.find(T); 680 if (Loc != TypeMap.end()) 681 return Loc->second; 682 683 SPIRVDBG(spvdbgs() << "[transType] " << *T << " -> ";) 684 T->validate(); 685 switch(T->getOpCode()) { 686 case OpTypeVoid: 687 return mapType(T, Type::getVoidTy(*Context)); 688 case OpTypeBool: 689 return mapType(T, Type::getInt1Ty(*Context)); 690 case OpTypeInt: 691 return mapType(T, Type::getIntNTy(*Context, T->getIntegerBitWidth())); 692 case OpTypeFloat: 693 return mapType(T, transFPType(T)); 694 case OpTypeArray: 695 return mapType(T, ArrayType::get(transType(T->getArrayElementType()), 696 T->getArrayLength())); 697 case OpTypePointer: 698 return mapType(T, PointerType::get(transType( 699 T->getPointerElementType(), IsClassMember), 700 SPIRSPIRVAddrSpaceMap::rmap(T->getPointerStorageClass()))); 701 case OpTypeVector: 702 return mapType(T, VectorType::get(transType(T->getVectorComponentType()), 703 T->getVectorComponentCount())); 704 case OpTypeOpaque: 705 return mapType(T, StructType::create(*Context, T->getName())); 706 case OpTypeFunction: { 707 auto FT = static_cast<SPIRVTypeFunction *>(T); 708 auto RT = transType(FT->getReturnType()); 709 std::vector<Type *> PT; 710 for (size_t I = 0, E = FT->getNumParameters(); I != E; ++I) 711 PT.push_back(transType(FT->getParameterType(I))); 712 return mapType(T, FunctionType::get(RT, PT, false)); 713 } 714 case OpTypeImage: { 715 auto ST = static_cast<SPIRVTypeImage *>(T); 716 if (ST->isOCLImage()) 717 return mapType(T, getOrCreateOpaquePtrType(M, 718 transOCLImageTypeName(ST))); 719 else 720 llvm_unreachable("Unsupported image type"); 721 return nullptr; 722 } 723 case OpTypeSampler: 724 return mapType(T, Type::getInt32Ty(*Context)); 725 case OpTypeSampledImage: { 726 auto ST = static_cast<SPIRVTypeSampledImage *>(T); 727 return mapType(T, getOrCreateOpaquePtrType(M, 728 transOCLSampledImageTypeName(ST))); 729 } 730 case OpTypeStruct: { 731 auto ST = static_cast<SPIRVTypeStruct *>(T); 732 auto Name = ST->getName(); 733 if (!Name.empty()) { 734 if (auto OldST = M->getTypeByName(Name)) 735 OldST->setName(""); 736 } 737 auto *StructTy = StructType::create(*Context, Name); 738 mapType(ST, StructTy); 739 SmallVector<Type *, 4> MT; 740 for (size_t I = 0, E = ST->getMemberCount(); I != E; ++I) 741 MT.push_back(transType(ST->getMemberType(I), true)); 742 StructTy->setBody(MT, ST->isPacked()); 743 return StructTy; 744 } 745 case OpTypePipe: { 746 auto PT = static_cast<SPIRVTypePipe *>(T); 747 return mapType(T, getOrCreateOpaquePtrType(M, 748 transOCLPipeTypeName(PT, IsClassMember, PT->getAccessQualifier()), 749 getOCLOpaqueTypeAddrSpace(T->getOpCode()))); 750 751 } 752 case OpTypePipeStorage: { 753 auto PST = static_cast<SPIRVTypePipeStorage *>(T); 754 return mapType(T, getOrCreateOpaquePtrType(M, 755 transOCLPipeStorageTypeName(PST), 756 getOCLOpaqueTypeAddrSpace(T->getOpCode()))); 757 } 758 default: { 759 auto OC = T->getOpCode(); 760 if (isOpaqueGenericTypeOpCode(OC)) 761 return mapType(T, getOrCreateOpaquePtrType(M, 762 OCLOpaqueTypeOpCodeMap::rmap(OC), 763 getOCLOpaqueTypeAddrSpace(OC))); 764 llvm_unreachable("Not implemented"); 765 } 766 } 767 return 0; 768} 769 770std::string 771SPIRVToLLVM::transTypeToOCLTypeName(SPIRVType *T, bool IsSigned) { 772 switch(T->getOpCode()) { 773 case OpTypeVoid: 774 return "void"; 775 case OpTypeBool: 776 return "bool"; 777 case OpTypeInt: { 778 std::string Prefix = IsSigned ? "" : "u"; 779 switch(T->getIntegerBitWidth()) { 780 case 8: 781 return Prefix + "char"; 782 case 16: 783 return Prefix + "short"; 784 case 32: 785 return Prefix + "int"; 786 case 64: 787 return Prefix + "long"; 788 default: 789 llvm_unreachable("invalid integer size"); 790 return Prefix + std::string("int") + T->getIntegerBitWidth() + "_t"; 791 } 792 } 793 break; 794 case OpTypeFloat: 795 switch(T->getFloatBitWidth()){ 796 case 16: 797 return "half"; 798 case 32: 799 return "float"; 800 case 64: 801 return "double"; 802 default: 803 llvm_unreachable("invalid floating pointer bitwidth"); 804 return std::string("float") + T->getFloatBitWidth() + "_t"; 805 } 806 break; 807 case OpTypeArray: 808 return "array"; 809 case OpTypePointer: 810 return transTypeToOCLTypeName(T->getPointerElementType()) + "*"; 811 case OpTypeVector: 812 return transTypeToOCLTypeName(T->getVectorComponentType()) + 813 T->getVectorComponentCount(); 814 case OpTypeOpaque: 815 return T->getName(); 816 case OpTypeFunction: 817 llvm_unreachable("Unsupported"); 818 return "function"; 819 case OpTypeStruct: { 820 auto Name = T->getName(); 821 if (Name.find("struct.") == 0) 822 Name[6] = ' '; 823 else if (Name.find("union.") == 0) 824 Name[5] = ' '; 825 return Name; 826 } 827 case OpTypePipe: 828 return "pipe"; 829 case OpTypeSampler: 830 return "sampler_t"; 831 case OpTypeImage: { 832 std::string Name; 833 Name = rmap<std::string>(static_cast<SPIRVTypeImage *>(T)->getDescriptor()); 834 if (SPIRVGenImgTypeAccQualPostfix) { 835 auto ST = static_cast<SPIRVTypeImage *>(T); 836 insertImageNameAccessQualifier(ST, Name); 837 } 838 return Name; 839 } 840 default: 841 if (isOpaqueGenericTypeOpCode(T->getOpCode())) { 842 return OCLOpaqueTypeOpCodeMap::rmap(T->getOpCode()); 843 } 844 llvm_unreachable("Not implemented"); 845 return "unknown"; 846 } 847} 848 849std::vector<Type *> 850SPIRVToLLVM::transTypeVector(const std::vector<SPIRVType *> &BT) { 851 std::vector<Type *> T; 852 for (auto I: BT) 853 T.push_back(transType(I)); 854 return T; 855} 856 857std::vector<Value *> 858SPIRVToLLVM::transValue(const std::vector<SPIRVValue *> &BV, Function *F, 859 BasicBlock *BB) { 860 std::vector<Value *> V; 861 for (auto I: BV) 862 V.push_back(transValue(I, F, BB)); 863 return V; 864} 865 866bool 867SPIRVToLLVM::isSPIRVCmpInstTransToLLVMInst(SPIRVInstruction* BI) const { 868 auto OC = BI->getOpCode(); 869 return isCmpOpCode(OC) && 870 !(OC >= OpLessOrGreater && OC <= OpUnordered); 871} 872 873void 874SPIRVToLLVM::transFlags(llvm::Value* V) { 875 if(!isa<Instruction>(V)) 876 return; 877 auto OC = cast<Instruction>(V)->getOpcode(); 878 if (OC == Instruction::AShr || OC == Instruction::LShr) { 879 cast<BinaryOperator>(V)->setIsExact(); 880 return; 881 } 882} 883 884void 885SPIRVToLLVM::setName(llvm::Value* V, SPIRVValue* BV) { 886 auto Name = BV->getName(); 887 if (!Name.empty() && (!V->hasName() || Name != V->getName())) 888 V->setName(Name); 889} 890 891void SPIRVToLLVM::insertImageNameAccessQualifier(SPIRV::SPIRVTypeImage* ST, std::string &Name) { 892 std::string QName = rmap<std::string>(ST->getAccessQualifier()); 893 // transform: read_only -> ro, write_only -> wo, read_write -> rw 894 QName = QName.substr(0,1) + QName.substr(QName.find("_") + 1, 1) + "_"; 895 assert(!Name.empty() && "image name should not be empty"); 896 Name.insert(Name.size() - 1, QName); 897} 898 899Value * 900SPIRVToLLVM::transValue(SPIRVValue *BV, Function *F, BasicBlock *BB, 901 bool CreatePlaceHolder){ 902 SPIRVToLLVMValueMap::iterator Loc = ValueMap.find(BV); 903 if (Loc != ValueMap.end() && (!PlaceholderMap.count(BV) || CreatePlaceHolder)) 904 return Loc->second; 905 906 SPIRVDBG(spvdbgs() << "[transValue] " << *BV << " -> ";) 907 BV->validate(); 908 909 auto V = transValueWithoutDecoration(BV, F, BB, CreatePlaceHolder); 910 if (!V) { 911 SPIRVDBG(dbgs() << " Warning ! nullptr\n";) 912 return nullptr; 913 } 914 setName(V, BV); 915 if (!transDecoration(BV, V)) { 916 assert (0 && "trans decoration fail"); 917 return nullptr; 918 } 919 transFlags(V); 920 921 SPIRVDBG(dbgs() << *V << '\n';) 922 923 return V; 924} 925 926Value * 927SPIRVToLLVM::transConvertInst(SPIRVValue* BV, Function* F, BasicBlock* BB) { 928 SPIRVUnary* BC = static_cast<SPIRVUnary*>(BV); 929 auto Src = transValue(BC->getOperand(0), F, BB, BB ? true : false); 930 auto Dst = transType(BC->getType()); 931 CastInst::CastOps CO = Instruction::BitCast; 932 bool IsExt = Dst->getScalarSizeInBits() 933 > Src->getType()->getScalarSizeInBits(); 934 switch (BC->getOpCode()) { 935 case OpPtrCastToGeneric: 936 case OpGenericCastToPtr: 937 CO = Instruction::AddrSpaceCast; 938 break; 939 case OpSConvert: 940 CO = IsExt ? Instruction::SExt : Instruction::Trunc; 941 break; 942 case OpUConvert: 943 CO = IsExt ? Instruction::ZExt : Instruction::Trunc; 944 break; 945 case OpFConvert: 946 CO = IsExt ? Instruction::FPExt : Instruction::FPTrunc; 947 break; 948 default: 949 CO = static_cast<CastInst::CastOps>(OpCodeMap::rmap(BC->getOpCode())); 950 } 951 assert(CastInst::isCast(CO) && "Invalid cast op code"); 952 SPIRVDBG(if (!CastInst::castIsValid(CO, Src, Dst)) { 953 spvdbgs() << "Invalid cast: " << *BV << " -> "; 954 dbgs() << "Op = " << CO << ", Src = " << *Src << " Dst = " << *Dst << '\n'; 955 }) 956 if (BB) 957 return CastInst::Create(CO, Src, Dst, BV->getName(), BB); 958 return ConstantExpr::getCast(CO, dyn_cast<Constant>(Src), Dst); 959} 960 961BinaryOperator *SPIRVToLLVM::transShiftLogicalBitwiseInst(SPIRVValue* BV, 962 BasicBlock* BB,Function* F) { 963 SPIRVBinary* BBN = static_cast<SPIRVBinary*>(BV); 964 assert(BB && "Invalid BB"); 965 Instruction::BinaryOps BO; 966 auto OP = BBN->getOpCode(); 967 if (isLogicalOpCode(OP)) 968 OP = IntBoolOpMap::rmap(OP); 969 BO = static_cast<Instruction::BinaryOps>(OpCodeMap::rmap(OP)); 970 auto Inst = BinaryOperator::Create(BO, 971 transValue(BBN->getOperand(0), F, BB), 972 transValue(BBN->getOperand(1), F, BB), BV->getName(), BB); 973 return Inst; 974} 975 976Instruction * 977SPIRVToLLVM::transCmpInst(SPIRVValue* BV, BasicBlock* BB, Function* F) { 978 SPIRVCompare* BC = static_cast<SPIRVCompare*>(BV); 979 assert(BB && "Invalid BB"); 980 SPIRVType* BT = BC->getOperand(0)->getType(); 981 Instruction* Inst = nullptr; 982 auto OP = BC->getOpCode(); 983 if (isLogicalOpCode(OP)) 984 OP = IntBoolOpMap::rmap(OP); 985 if (BT->isTypeVectorOrScalarInt() || BT->isTypeVectorOrScalarBool() || 986 BT->isTypePointer()) 987 Inst = new ICmpInst(*BB, CmpMap::rmap(OP), 988 transValue(BC->getOperand(0), F, BB), 989 transValue(BC->getOperand(1), F, BB)); 990 else if (BT->isTypeVectorOrScalarFloat()) 991 Inst = new FCmpInst(*BB, CmpMap::rmap(OP), 992 transValue(BC->getOperand(0), F, BB), 993 transValue(BC->getOperand(1), F, BB)); 994 assert(Inst && "not implemented"); 995 return Inst; 996} 997 998bool 999SPIRVToLLVM::postProcessOCL() { 1000 std::string DemangledName; 1001 SPIRVWord SrcLangVer = 0; 1002 BM->getSourceLanguage(&SrcLangVer); 1003 bool isCPP = SrcLangVer == kOCLVer::CL21; 1004 for (auto I = M->begin(), E = M->end(); I != E;) { 1005 auto F = I++; 1006 if (F->hasName() && F->isDeclaration()) { 1007 DEBUG(dbgs() << "[postProcessOCL sret] " << *F << '\n'); 1008 if (F->getReturnType()->isStructTy() && 1009 oclIsBuiltin(F->getName(), &DemangledName, isCPP)) { 1010 if (!postProcessOCLBuiltinReturnStruct(static_cast<Function*>(F))) 1011 return false; 1012 } 1013 } 1014 } 1015 for (auto I = M->begin(), E = M->end(); I != E;) { 1016 auto F = static_cast<Function*>(I++); 1017 if (F->hasName() && F->isDeclaration()) { 1018 DEBUG(dbgs() << "[postProcessOCL func ptr] " << *F << '\n'); 1019 auto AI = F->arg_begin(); 1020 if (hasFunctionPointerArg(F, AI) && isDecoratedSPIRVFunc(F)) 1021 if (!postProcessOCLBuiltinWithFuncPointer(F, AI)) 1022 return false; 1023 } 1024 } 1025 for (auto I = M->begin(), E = M->end(); I != E;) { 1026 auto F = static_cast<Function*>(I++); 1027 if (F->hasName() && F->isDeclaration()) { 1028 DEBUG(dbgs() << "[postProcessOCL array arg] " << *F << '\n'); 1029 if (hasArrayArg(F) && oclIsBuiltin(F->getName(), &DemangledName, isCPP)) 1030 if (!postProcessOCLBuiltinWithArrayArguments(F, DemangledName)) 1031 return false; 1032 } 1033 } 1034 return true; 1035} 1036 1037bool 1038SPIRVToLLVM::postProcessOCLBuiltinReturnStruct(Function *F) { 1039 std::string Name = F->getName(); 1040 F->setName(Name + ".old"); 1041 for (auto I = F->user_begin(), E = F->user_end(); I != E;) { 1042 if (auto CI = dyn_cast<CallInst>(*I++)) { 1043 auto ST = dyn_cast<StoreInst>(*(CI->user_begin())); 1044 assert(ST); 1045 std::vector<Type *> ArgTys; 1046 getFunctionTypeParameterTypes(F->getFunctionType(), ArgTys); 1047 ArgTys.insert(ArgTys.begin(), PointerType::get(F->getReturnType(), 1048 SPIRAS_Private)); 1049 auto newF = getOrCreateFunction(M, Type::getVoidTy(*Context), 1050 ArgTys, Name); 1051 newF->setCallingConv(F->getCallingConv()); 1052 auto Args = getArguments(CI); 1053 Args.insert(Args.begin(), ST->getPointerOperand()); 1054 auto NewCI = CallInst::Create(newF, Args, CI->getName(), CI); 1055 NewCI->setCallingConv(CI->getCallingConv()); 1056 ST->dropAllReferences(); 1057 ST->removeFromParent(); 1058 CI->dropAllReferences(); 1059 CI->removeFromParent(); 1060 } 1061 } 1062 F->dropAllReferences(); 1063 F->removeFromParent(); 1064 return true; 1065} 1066 1067bool 1068SPIRVToLLVM::postProcessOCLBuiltinWithFuncPointer(Function* F, 1069 Function::arg_iterator I) { 1070 auto Name = undecorateSPIRVFunction(F->getName()); 1071 std::set<Value *> InvokeFuncPtrs; 1072 mutateFunctionOCL (F, [=, &InvokeFuncPtrs]( 1073 CallInst *CI, std::vector<Value *> &Args) { 1074 auto ALoc = std::find_if(Args.begin(), Args.end(), [](Value * elem) { 1075 return isFunctionPointerType(elem->getType()); 1076 }); 1077 assert(ALoc != Args.end() && "Buit-in must accept a pointer to function"); 1078 assert(isa<Function>(*ALoc) && "Invalid function pointer usage"); 1079 Value *Ctx = ALoc[1]; 1080 Value *CtxLen = ALoc[2]; 1081 Value *CtxAlign = ALoc[3]; 1082 if (Name == kOCLBuiltinName::EnqueueKernel) 1083 assert(Args.end() - ALoc > 3); 1084 else 1085 assert(Args.end() - ALoc > 0); 1086 // Erase arguments what are hanled by "spir_block_bind" according to SPIR 2.0 1087 Args.erase(ALoc + 1, ALoc + 4); 1088 1089 InvokeFuncPtrs.insert(*ALoc); 1090 // There will be as many calls to spir_block_bind as how much device execution 1091 // bult-ins using this block. This doesn't contradict SPIR 2.0 specification. 1092 *ALoc = addBlockBind(M, cast<Function>(removeCast(*ALoc)), 1093 Ctx, CtxLen, CtxAlign, CI); 1094 return Name; 1095 }); 1096 for (auto &I:InvokeFuncPtrs) 1097 eraseIfNoUse(I); 1098 return true; 1099} 1100 1101bool 1102SPIRVToLLVM::postProcessOCLBuiltinWithArrayArguments(Function* F, 1103 const std::string &DemangledName) { 1104 DEBUG(dbgs() << "[postProcessOCLBuiltinWithArrayArguments] " << *F << '\n'); 1105 auto Attrs = F->getAttributes(); 1106 auto Name = F->getName(); 1107 mutateFunction(F, [=](CallInst *CI, std::vector<Value *> &Args) { 1108 auto FBegin = CI->getParent()->getParent()->begin()->getFirstInsertionPt(); 1109 for (auto &I:Args) { 1110 auto T = I->getType(); 1111 if (!T->isArrayTy()) 1112 continue; 1113 auto Alloca = new AllocaInst(T, "", static_cast<Instruction*>(FBegin)); 1114 new StoreInst(I, Alloca, false, CI); 1115 auto Zero = ConstantInt::getNullValue(Type::getInt32Ty(T->getContext())); 1116 Value *Index[] = {Zero, Zero}; 1117 I = GetElementPtrInst::CreateInBounds(Alloca, Index, "", CI); 1118 } 1119 return Name; 1120 }, nullptr, &Attrs); 1121 return true; 1122} 1123 1124// ToDo: Handle unsigned integer return type. May need spec change. 1125Instruction * 1126SPIRVToLLVM::postProcessOCLReadImage(SPIRVInstruction *BI, CallInst* CI, 1127 const std::string &FuncName) { 1128 AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); 1129 StringRef ImageTypeName; 1130 bool isDepthImage = false; 1131 if (isOCLImageType( 1132 (cast<CallInst>(CI->getOperand(0)))->getArgOperand(0)->getType(), 1133 &ImageTypeName)) 1134 isDepthImage = ImageTypeName.endswith("depth_t"); 1135 return mutateCallInstOCL( 1136 M, CI, 1137 [=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { 1138 CallInst *CallSampledImg = cast<CallInst>(Args[0]); 1139 auto Img = CallSampledImg->getArgOperand(0); 1140 assert(isOCLImageType(Img->getType())); 1141 auto Sampler = CallSampledImg->getArgOperand(1); 1142 Args[0] = Img; 1143 Args.insert(Args.begin() + 1, Sampler); 1144 if(Args.size() > 4 ) { 1145 ConstantInt* ImOp = dyn_cast<ConstantInt>(Args[3]); 1146 ConstantFP* LodVal = dyn_cast<ConstantFP>(Args[4]); 1147 // Drop "Image Operands" argument. 1148 Args.erase(Args.begin() + 3, Args.begin() + 4); 1149 // If the image operand is LOD and its value is zero, drop it too. 1150 if (ImOp && LodVal && LodVal->isNullValue() && 1151 ImOp->getZExtValue() == ImageOperandsMask::ImageOperandsLodMask ) 1152 Args.erase(Args.begin() + 3, Args.end()); 1153 } 1154 if (CallSampledImg->hasOneUse()) { 1155 CallSampledImg->replaceAllUsesWith( 1156 UndefValue::get(CallSampledImg->getType())); 1157 CallSampledImg->dropAllReferences(); 1158 CallSampledImg->eraseFromParent(); 1159 } 1160 Type *T = CI->getType(); 1161 if (auto VT = dyn_cast<VectorType>(T)) 1162 T = VT->getElementType(); 1163 RetTy = isDepthImage ? T : CI->getType(); 1164 return std::string(kOCLBuiltinName::SampledReadImage) + 1165 (T->isFloatingPointTy() ? 'f' : 'i'); 1166 }, 1167 [=](CallInst *NewCI) -> Instruction * { 1168 if (isDepthImage) 1169 return InsertElementInst::Create( 1170 UndefValue::get(VectorType::get(NewCI->getType(), 4)), NewCI, 1171 getSizet(M, 0), "", NewCI->getParent()); 1172 return NewCI; 1173 }, 1174 &Attrs); 1175} 1176 1177CallInst* 1178SPIRVToLLVM::postProcessOCLWriteImage(SPIRVInstruction *BI, CallInst *CI, 1179 const std::string &DemangledName) { 1180 AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); 1181 return mutateCallInstOCL(M, CI, [=](CallInst *, std::vector<Value *> &Args) { 1182 llvm::Type *T = Args[2]->getType(); 1183 if (Args.size() > 4) { 1184 ConstantInt* ImOp = dyn_cast<ConstantInt>(Args[3]); 1185 ConstantFP* LodVal = dyn_cast<ConstantFP>(Args[4]); 1186 // Drop "Image Operands" argument. 1187 Args.erase(Args.begin() + 3, Args.begin() + 4); 1188 // If the image operand is LOD and its value is zero, drop it too. 1189 if (ImOp && LodVal && LodVal->isNullValue() && 1190 ImOp->getZExtValue() == ImageOperandsMask::ImageOperandsLodMask ) 1191 Args.erase(Args.begin() + 3, Args.end()); 1192 else 1193 std::swap(Args[2], Args[3]); 1194 } 1195 return std::string(kOCLBuiltinName::WriteImage) + 1196 (T->isFPOrFPVectorTy() ? 'f' : 'i'); 1197 }, &Attrs); 1198} 1199 1200CallInst * 1201SPIRVToLLVM::postProcessOCLBuildNDRange(SPIRVInstruction *BI, CallInst *CI, 1202 const std::string &FuncName) { 1203 assert(CI->getNumArgOperands() == 3); 1204 auto GWS = CI->getArgOperand(0); 1205 auto LWS = CI->getArgOperand(1); 1206 auto GWO = CI->getArgOperand(2); 1207 CI->setArgOperand(0, GWO); 1208 CI->setArgOperand(1, GWS); 1209 CI->setArgOperand(2, LWS); 1210 return CI; 1211} 1212 1213Instruction * 1214SPIRVToLLVM::postProcessGroupAllAny(CallInst *CI, 1215 const std::string &DemangledName) { 1216 AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); 1217 return mutateCallInstSPIRV( 1218 M, CI, 1219 [=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { 1220 Type *Int32Ty = Type::getInt32Ty(*Context); 1221 RetTy = Int32Ty; 1222 Args[1] = CastInst::CreateZExtOrBitCast(Args[1], Int32Ty, "", CI); 1223 return DemangledName; 1224 }, 1225 [=](CallInst *NewCI) -> Instruction * { 1226 Type *RetTy = Type::getInt1Ty(*Context); 1227 return CastInst::CreateTruncOrBitCast(NewCI, RetTy, "", 1228 NewCI->getNextNode()); 1229 }, 1230 &Attrs); 1231} 1232 1233CallInst * 1234SPIRVToLLVM::expandOCLBuiltinWithScalarArg(CallInst* CI, 1235 const std::string &FuncName) { 1236 AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); 1237 if (!CI->getOperand(0)->getType()->isVectorTy() && 1238 CI->getOperand(1)->getType()->isVectorTy()) { 1239 return mutateCallInstOCL(M, CI, [=](CallInst *, std::vector<Value *> &Args){ 1240 unsigned vecSize = CI->getOperand(1)->getType()->getVectorNumElements(); 1241 Value *NewVec = nullptr; 1242 if (auto CA = dyn_cast<Constant>(Args[0])) 1243 NewVec = ConstantVector::getSplat(vecSize, CA); 1244 else { 1245 NewVec = ConstantVector::getSplat(vecSize, 1246 Constant::getNullValue(Args[0]->getType())); 1247 NewVec = InsertElementInst::Create(NewVec, Args[0], getInt32(M, 0), "", 1248 CI); 1249 NewVec = new ShuffleVectorInst(NewVec, NewVec, 1250 ConstantVector::getSplat(vecSize, getInt32(M, 0)), "", CI); 1251 } 1252 NewVec->takeName(Args[0]); 1253 Args[0] = NewVec; 1254 return FuncName; 1255 }, &Attrs); 1256 } 1257 return CI; 1258} 1259 1260std::string 1261SPIRVToLLVM::transOCLPipeTypeAccessQualifier(SPIRV::SPIRVTypePipe* ST) { 1262 return SPIRSPIRVAccessQualifierMap::rmap(ST->getAccessQualifier()); 1263} 1264 1265void 1266SPIRVToLLVM::transGeneratorMD() { 1267 SPIRVMDBuilder B(*M); 1268 B.addNamedMD(kSPIRVMD::Generator) 1269 .addOp() 1270 .addU16(BM->getGeneratorId()) 1271 .addU16(BM->getGeneratorVer()) 1272 .done(); 1273} 1274 1275Value * 1276SPIRVToLLVM::oclTransConstantSampler(SPIRV::SPIRVConstantSampler* BCS) { 1277 auto Lit = (BCS->getAddrMode() << 1) | 1278 BCS->getNormalized() | 1279 ((BCS->getFilterMode() + 1) << 4); 1280 auto Ty = IntegerType::getInt32Ty(*Context); 1281 return ConstantInt::get(Ty, Lit); 1282} 1283 1284Value * 1285SPIRVToLLVM::oclTransConstantPipeStorage( 1286 SPIRV::SPIRVConstantPipeStorage* BCPS) { 1287 1288 string CPSName = string(kSPIRVTypeName::PrefixAndDelim) 1289 + kSPIRVTypeName::ConstantPipeStorage; 1290 1291 auto Int32Ty = IntegerType::getInt32Ty(*Context); 1292 auto CPSTy = M->getTypeByName(CPSName); 1293 if (!CPSTy) { 1294 Type* CPSElemsTy[] = { Int32Ty, Int32Ty, Int32Ty }; 1295 CPSTy = StructType::create(*Context, CPSElemsTy, CPSName); 1296 } 1297 1298 assert(CPSTy != nullptr && "Could not create spirv.ConstantPipeStorage"); 1299 1300 Constant* CPSElems[] = { 1301 ConstantInt::get(Int32Ty, BCPS->getPacketSize()), 1302 ConstantInt::get(Int32Ty, BCPS->getPacketAlign()), 1303 ConstantInt::get(Int32Ty, BCPS->getCapacity()) 1304 }; 1305 1306 return new GlobalVariable(*M, CPSTy, false, GlobalValue::LinkOnceODRLinkage, 1307 ConstantStruct::get(CPSTy, CPSElems), BCPS->getName(), 1308 nullptr, GlobalValue::NotThreadLocal, SPIRAS_Global); 1309} 1310 1311/// For instructions, this function assumes they are created in order 1312/// and appended to the given basic block. An instruction may use a 1313/// instruction from another BB which has not been translated. Such 1314/// instructions should be translated to place holders at the point 1315/// of first use, then replaced by real instructions when they are 1316/// created. 1317/// 1318/// When CreatePlaceHolder is true, create a load instruction of a 1319/// global variable as placeholder for SPIRV instruction. Otherwise, 1320/// create instruction and replace placeholder if there is one. 1321Value * 1322SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F, 1323 BasicBlock *BB, bool CreatePlaceHolder){ 1324 1325 auto OC = BV->getOpCode(); 1326 IntBoolOpMap::rfind(OC, &OC); 1327 1328 // Translation of non-instruction values 1329 switch(OC) { 1330 case OpConstant: { 1331 SPIRVConstant *BConst = static_cast<SPIRVConstant *>(BV); 1332 SPIRVType *BT = BV->getType(); 1333 Type *LT = transType(BT); 1334 switch(BT->getOpCode()) { 1335 case OpTypeBool: 1336 case OpTypeInt: 1337 return mapValue(BV, ConstantInt::get(LT, BConst->getZExtIntValue(), 1338 static_cast<SPIRVTypeInt*>(BT)->isSigned())); 1339 case OpTypeFloat: { 1340 const llvm::fltSemantics *FS = nullptr; 1341 switch (BT->getFloatBitWidth()) { 1342 case 16: 1343 FS = &APFloat::IEEEhalf; 1344 break; 1345 case 32: 1346 FS = &APFloat::IEEEsingle; 1347 break; 1348 case 64: 1349 FS = &APFloat::IEEEdouble; 1350 break; 1351 default: 1352 llvm_unreachable("invalid float type"); 1353 } 1354 return mapValue(BV, ConstantFP::get(*Context, APFloat(*FS, 1355 APInt(BT->getFloatBitWidth(), BConst->getZExtIntValue())))); 1356 } 1357 default: 1358 llvm_unreachable("Not implemented"); 1359 return nullptr; 1360 } 1361 } 1362 1363 case OpConstantTrue: 1364 return mapValue(BV, ConstantInt::getTrue(*Context)); 1365 1366 case OpConstantFalse: 1367 return mapValue(BV, ConstantInt::getFalse(*Context)); 1368 1369 case OpConstantNull: { 1370 auto LT = transType(BV->getType()); 1371 return mapValue(BV, Constant::getNullValue(LT)); 1372 } 1373 1374 case OpConstantComposite: { 1375 auto BCC = static_cast<SPIRVConstantComposite*>(BV); 1376 std::vector<Constant *> CV; 1377 for (auto &I:BCC->getElements()) 1378 CV.push_back(dyn_cast<Constant>(transValue(I, F, BB))); 1379 switch(BV->getType()->getOpCode()) { 1380 case OpTypeVector: 1381 return mapValue(BV, ConstantVector::get(CV)); 1382 case OpTypeArray: 1383 return mapValue(BV, ConstantArray::get( 1384 dyn_cast<ArrayType>(transType(BCC->getType())), CV)); 1385 case OpTypeStruct: { 1386 auto BCCTy = dyn_cast<StructType>(transType(BCC->getType())); 1387 auto Members = BCCTy->getNumElements(); 1388 auto Constants = CV.size(); 1389 //if we try to initialize constant TypeStruct, add bitcasts 1390 //if src and dst types are both pointers but to different types 1391 if (Members == Constants) { 1392 for (unsigned i = 0; i < Members; ++i) { 1393 if (CV[i]->getType() == BCCTy->getElementType(i)) 1394 continue; 1395 if (!CV[i]->getType()->isPointerTy() || 1396 !BCCTy->getElementType(i)->isPointerTy()) 1397 continue; 1398 1399 CV[i] = ConstantExpr::getBitCast(CV[i], BCCTy->getElementType(i)); 1400 } 1401 } 1402 1403 return mapValue(BV, ConstantStruct::get( 1404 dyn_cast<StructType>(transType(BCC->getType())), CV)); 1405 } 1406 default: 1407 llvm_unreachable("not implemented"); 1408 return nullptr; 1409 } 1410 } 1411 1412 case OpConstantSampler: { 1413 auto BCS = static_cast<SPIRVConstantSampler*>(BV); 1414 return mapValue(BV, oclTransConstantSampler(BCS)); 1415 } 1416 1417 case OpConstantPipeStorage: { 1418 auto BCPS = static_cast<SPIRVConstantPipeStorage*>(BV); 1419 return mapValue(BV, oclTransConstantPipeStorage(BCPS)); 1420 } 1421 1422 case OpSpecConstantOp: { 1423 auto BI = createInstFromSpecConstantOp( 1424 static_cast<SPIRVSpecConstantOp*>(BV)); 1425 return mapValue(BV, transValue(BI, nullptr, nullptr, false)); 1426 } 1427 1428 case OpUndef: 1429 return mapValue(BV, UndefValue::get(transType(BV->getType()))); 1430 1431 case OpVariable: { 1432 auto BVar = static_cast<SPIRVVariable *>(BV); 1433 auto Ty = transType(BVar->getType()->getPointerElementType()); 1434 bool IsConst = BVar->isConstant(); 1435 llvm::GlobalValue::LinkageTypes LinkageTy = transLinkageType(BVar); 1436 Constant *Initializer = nullptr; 1437 SPIRVValue *Init = BVar->getInitializer(); 1438 if (Init) 1439 Initializer = dyn_cast<Constant>(transValue(Init, F, BB, false)); 1440 else if (LinkageTy == GlobalValue::CommonLinkage) 1441 // In LLVM variables with common linkage type must be initilized by 0 1442 Initializer = Constant::getNullValue(Ty); 1443 1444 SPIRVStorageClassKind BS = BVar->getStorageClass(); 1445 if (BS == StorageClassFunction && !Init) { 1446 assert (BB && "Invalid BB"); 1447 return mapValue(BV, new AllocaInst(Ty, BV->getName(), BB)); 1448 } 1449 auto AddrSpace = SPIRSPIRVAddrSpaceMap::rmap(BS); 1450 auto LVar = new GlobalVariable(*M, Ty, IsConst, LinkageTy, Initializer, 1451 BV->getName(), 0, GlobalVariable::NotThreadLocal, AddrSpace); 1452 LVar->setUnnamedAddr((IsConst && Ty->isArrayTy() && 1453 Ty->getArrayElementType()->isIntegerTy(8)) ? 1454 GlobalValue::UnnamedAddr::Global : 1455 GlobalValue::UnnamedAddr::None); 1456 SPIRVBuiltinVariableKind BVKind; 1457 if (BVar->isBuiltin(&BVKind)) 1458 BuiltinGVMap[LVar] = BVKind; 1459 return mapValue(BV, LVar); 1460 } 1461 1462 case OpFunctionParameter: { 1463 auto BA = static_cast<SPIRVFunctionParameter*>(BV); 1464 assert (F && "Invalid function"); 1465 unsigned ArgNo = 0; 1466 for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(); I != E; 1467 ++I, ++ArgNo) { 1468 if (ArgNo == BA->getArgNo()) 1469 return mapValue(BV, static_cast<Argument*>(I)); 1470 } 1471 llvm_unreachable("Invalid argument"); 1472 return nullptr; 1473 } 1474 1475 case OpFunction: 1476 return mapValue(BV, transFunction(static_cast<SPIRVFunction *>(BV))); 1477 1478 case OpLabel: 1479 return mapValue(BV, BasicBlock::Create(*Context, BV->getName(), F)); 1480 1481 case OpBitcast: // Can be translated without BB pointer 1482 if(!CreatePlaceHolder) // May be a placeholder 1483 return mapValue(BV, transConvertInst(BV, F, BB)); 1484 1485 default: 1486 // do nothing 1487 break; 1488 } 1489 1490 // All other values require valid BB pointer. 1491 assert(BB && "Invalid BB"); 1492 1493 // Creation of place holder 1494 if (CreatePlaceHolder) { 1495 auto GV = new GlobalVariable(*M, 1496 transType(BV->getType()), 1497 false, 1498 GlobalValue::PrivateLinkage, 1499 nullptr, 1500 std::string(kPlaceholderPrefix) + BV->getName(), 1501 0, GlobalVariable::NotThreadLocal, 0); 1502 auto LD = new LoadInst(GV, BV->getName(), BB); 1503 PlaceholderMap[BV] = LD; 1504 return mapValue(BV, LD); 1505 } 1506 1507 // Translation of instructions 1508 switch (BV->getOpCode()) { 1509 case OpBranch: { 1510 auto BR = static_cast<SPIRVBranch *>(BV); 1511 return mapValue(BV, BranchInst::Create( 1512 dyn_cast<BasicBlock>(transValue(BR->getTargetLabel(), F, BB)), BB)); 1513 } 1514 1515 case OpBranchConditional: { 1516 auto BR = static_cast<SPIRVBranchConditional *>(BV); 1517 return mapValue( 1518 BV, BranchInst::Create( 1519 dyn_cast<BasicBlock>(transValue(BR->getTrueLabel(), F, BB)), 1520 dyn_cast<BasicBlock>(transValue(BR->getFalseLabel(), F, BB)), 1521 transValue(BR->getCondition(), F, BB), BB)); 1522 } 1523 1524 case OpPhi: { 1525 auto Phi = static_cast<SPIRVPhi *>(BV); 1526 auto LPhi = dyn_cast<PHINode>(mapValue( 1527 BV, PHINode::Create(transType(Phi->getType()), 1528 Phi->getPairs().size() / 2, Phi->getName(), BB))); 1529 Phi->foreachPair([&](SPIRVValue *IncomingV, SPIRVBasicBlock *IncomingBB, 1530 size_t Index) { 1531 auto Translated = transValue(IncomingV, F, BB); 1532 LPhi->addIncoming(Translated, 1533 dyn_cast<BasicBlock>(transValue(IncomingBB, F, BB))); 1534 }); 1535 return LPhi; 1536 } 1537 1538 case OpReturn: 1539 return mapValue(BV, ReturnInst::Create(*Context, BB)); 1540 1541 case OpReturnValue: { 1542 auto RV = static_cast<SPIRVReturnValue *>(BV); 1543 return mapValue( 1544 BV, ReturnInst::Create(*Context, 1545 transValue(RV->getReturnValue(), F, BB), BB)); 1546 } 1547 1548 case OpStore: { 1549 SPIRVStore *BS = static_cast<SPIRVStore*>(BV); 1550 StoreInst *SI = new StoreInst(transValue(BS->getSrc(), F, BB), 1551 transValue(BS->getDst(), F, BB), 1552 BS->SPIRVMemoryAccess::isVolatile(), 1553 BS->SPIRVMemoryAccess::getAlignment(), BB); 1554 if (BS->SPIRVMemoryAccess::isNonTemporal()) 1555 transNonTemporalMetadata(SI); 1556 return mapValue(BV, SI); 1557 } 1558 1559 case OpLoad: { 1560 SPIRVLoad *BL = static_cast<SPIRVLoad*>(BV); 1561 LoadInst *LI = new LoadInst(transValue(BL->getSrc(), F, BB), BV->getName(), 1562 BL->SPIRVMemoryAccess::isVolatile(), 1563 BL->SPIRVMemoryAccess::getAlignment(), BB); 1564 if (BL->SPIRVMemoryAccess::isNonTemporal()) 1565 transNonTemporalMetadata(LI); 1566 return mapValue(BV, LI); 1567 } 1568 1569 case OpCopyMemorySized: { 1570 SPIRVCopyMemorySized *BC = static_cast<SPIRVCopyMemorySized *>(BV); 1571 std::string FuncName = "llvm.memcpy"; 1572 SPIRVType* BS = BC->getSource()->getType(); 1573 SPIRVType* BT = BC->getTarget()->getType(); 1574 Type *Int1Ty = Type::getInt1Ty(*Context); 1575 Type* Int32Ty = Type::getInt32Ty(*Context); 1576 Type* VoidTy = Type::getVoidTy(*Context); 1577 Type* SrcTy = transType(BS); 1578 Type* TrgTy = transType(BT); 1579 Type* SizeTy = transType(BC->getSize()->getType()); 1580 Type* ArgTy[] = { TrgTy, SrcTy, SizeTy, Int32Ty, Int1Ty }; 1581 1582 ostringstream TempName; 1583 TempName << ".p" << SPIRSPIRVAddrSpaceMap::rmap(BT->getPointerStorageClass()) << "i8"; 1584 TempName << ".p" << SPIRSPIRVAddrSpaceMap::rmap(BS->getPointerStorageClass()) << "i8"; 1585 FuncName += TempName.str(); 1586 if (BC->getSize()->getType()->getBitWidth() == 32) 1587 FuncName += ".i32"; 1588 else 1589 FuncName += ".i64"; 1590 1591 FunctionType *FT = FunctionType::get(VoidTy, ArgTy, false); 1592 Function *Func = dyn_cast<Function>(M->getOrInsertFunction(FuncName, FT)); 1593 assert(Func && Func->getFunctionType() == FT && "Function type mismatch"); 1594 Func->setLinkage(GlobalValue::ExternalLinkage); 1595 1596 if (isFuncNoUnwind()) 1597 Func->addFnAttr(Attribute::NoUnwind); 1598 1599 Value *Arg[] = { transValue(BC->getTarget(), Func, BB), 1600 transValue(BC->getSource(), Func, BB), 1601 dyn_cast<llvm::ConstantInt>(transValue(BC->getSize(), 1602 Func, BB)), 1603 ConstantInt::get(Int32Ty, 1604 BC->SPIRVMemoryAccess::getAlignment()), 1605 ConstantInt::get(Int1Ty, 1606 BC->SPIRVMemoryAccess::isVolatile())}; 1607 return mapValue( BV, CallInst::Create(Func, Arg, "", BB)); 1608 } 1609 1610 case OpSelect: { 1611 SPIRVSelect *BS = static_cast<SPIRVSelect*>(BV); 1612 return mapValue(BV, 1613 SelectInst::Create(transValue(BS->getCondition(), F, BB), 1614 transValue(BS->getTrueValue(), F, BB), 1615 transValue(BS->getFalseValue(), F, BB), 1616 BV->getName(), BB)); 1617 } 1618 1619 case OpSwitch: { 1620 auto BS = static_cast<SPIRVSwitch *>(BV); 1621 auto Select = transValue(BS->getSelect(), F, BB); 1622 auto LS = SwitchInst::Create( 1623 Select, dyn_cast<BasicBlock>(transValue(BS->getDefault(), F, BB)), 1624 BS->getNumPairs(), BB); 1625 BS->foreachPair( 1626 [&](SPIRVWord Literal, SPIRVBasicBlock *Label, size_t Index) { 1627 LS->addCase(ConstantInt::get(dyn_cast<IntegerType>(Select->getType()), 1628 Literal), 1629 dyn_cast<BasicBlock>(transValue(Label, F, BB))); 1630 }); 1631 return mapValue(BV, LS); 1632 } 1633 1634 case OpAccessChain: 1635 case OpInBoundsAccessChain: 1636 case OpPtrAccessChain: 1637 case OpInBoundsPtrAccessChain: { 1638 auto AC = static_cast<SPIRVAccessChainBase *>(BV); 1639 auto Base = transValue(AC->getBase(), F, BB); 1640 auto Index = transValue(AC->getIndices(), F, BB); 1641 if (!AC->hasPtrIndex()) 1642 Index.insert(Index.begin(), getInt32(M, 0)); 1643 auto IsInbound = AC->isInBounds(); 1644 Value *V = nullptr; 1645 if (BB) { 1646 auto GEP = GetElementPtrInst::Create(nullptr, Base, Index, 1647 BV->getName(), BB); 1648 GEP->setIsInBounds(IsInbound); 1649 V = GEP; 1650 } else { 1651 V = ConstantExpr::getGetElementPtr(Base->getType(), 1652 dyn_cast<Constant>(Base), 1653 Index, 1654 IsInbound); 1655 } 1656 return mapValue(BV, V); 1657 } 1658 1659 case OpCompositeExtract: { 1660 SPIRVCompositeExtract *CE = static_cast<SPIRVCompositeExtract *>(BV); 1661 if (CE->getComposite()->getType()->isTypeVector()) { 1662 assert(CE->getIndices().size() == 1 && "Invalid index"); 1663 return mapValue( 1664 BV, ExtractElementInst::Create( 1665 transValue(CE->getComposite(), F, BB), 1666 ConstantInt::get(*Context, APInt(32, CE->getIndices()[0])), 1667 BV->getName(), BB)); 1668 } 1669 return mapValue( 1670 BV, ExtractValueInst::Create( 1671 transValue(CE->getComposite(), F, BB), 1672 CE->getIndices(), BV->getName(), BB)); 1673 } 1674 1675 case OpVectorExtractDynamic: { 1676 auto CE = static_cast<SPIRVVectorExtractDynamic *>(BV); 1677 return mapValue( 1678 BV, ExtractElementInst::Create(transValue(CE->getVector(), F, BB), 1679 transValue(CE->getIndex(), F, BB), 1680 BV->getName(), BB)); 1681 } 1682 1683 case OpCompositeInsert: { 1684 auto CI = static_cast<SPIRVCompositeInsert *>(BV); 1685 if (CI->getComposite()->getType()->isTypeVector()) { 1686 assert(CI->getIndices().size() == 1 && "Invalid index"); 1687 return mapValue( 1688 BV, InsertElementInst::Create( 1689 transValue(CI->getComposite(), F, BB), 1690 transValue(CI->getObject(), F, BB), 1691 ConstantInt::get(*Context, APInt(32, CI->getIndices()[0])), 1692 BV->getName(), BB)); 1693 } 1694 return mapValue( 1695 BV, InsertValueInst::Create( 1696 transValue(CI->getComposite(), F, BB), 1697 transValue(CI->getObject(), F, BB), 1698 CI->getIndices(), BV->getName(), BB)); 1699 } 1700 1701 case OpVectorInsertDynamic: { 1702 auto CI = static_cast<SPIRVVectorInsertDynamic *>(BV); 1703 return mapValue( 1704 BV, InsertElementInst::Create(transValue(CI->getVector(), F, BB), 1705 transValue(CI->getComponent(), F, BB), 1706 transValue(CI->getIndex(), F, BB), 1707 BV->getName(), BB)); 1708 } 1709 1710 case OpVectorShuffle: { 1711 auto VS = static_cast<SPIRVVectorShuffle *>(BV); 1712 std::vector<Constant *> Components; 1713 IntegerType *Int32Ty = IntegerType::get(*Context, 32); 1714 for (auto I : VS->getComponents()) { 1715 if (I == static_cast<SPIRVWord>(-1)) 1716 Components.push_back(UndefValue::get(Int32Ty)); 1717 else 1718 Components.push_back(ConstantInt::get(Int32Ty, I)); 1719 } 1720 return mapValue(BV, 1721 new ShuffleVectorInst(transValue(VS->getVector1(), F, BB), 1722 transValue(VS->getVector2(), F, BB), 1723 ConstantVector::get(Components), 1724 BV->getName(), BB)); 1725 } 1726 1727 case OpFunctionCall: { 1728 SPIRVFunctionCall *BC = static_cast<SPIRVFunctionCall *>(BV); 1729 auto Call = CallInst::Create(transFunction(BC->getFunction()), 1730 transValue(BC->getArgumentValues(), F, BB), 1731 BC->getName(), BB); 1732 setCallingConv(Call); 1733 setAttrByCalledFunc(Call); 1734 return mapValue(BV, Call); 1735 } 1736 1737 case OpExtInst: 1738 return mapValue( 1739 BV, transOCLBuiltinFromExtInst(static_cast<SPIRVExtInst *>(BV), BB)); 1740 1741 case OpControlBarrier: 1742 case OpMemoryBarrier: 1743 return mapValue( 1744 BV, transOCLBarrierFence(static_cast<SPIRVInstruction *>(BV), BB)); 1745 1746 case OpSNegate: { 1747 SPIRVUnary *BC = static_cast<SPIRVUnary *>(BV); 1748 return mapValue( 1749 BV, BinaryOperator::CreateNSWNeg(transValue(BC->getOperand(0), F, BB), 1750 BV->getName(), BB)); 1751 } 1752 1753 case OpFNegate: { 1754 SPIRVUnary *BC = static_cast<SPIRVUnary *>(BV); 1755 return mapValue( 1756 BV, BinaryOperator::CreateFNeg(transValue(BC->getOperand(0), F, BB), 1757 BV->getName(), BB)); 1758 } 1759 1760 case OpNot: { 1761 SPIRVUnary *BC = static_cast<SPIRVUnary *>(BV); 1762 return mapValue( 1763 BV, BinaryOperator::CreateNot(transValue(BC->getOperand(0), F, BB), 1764 BV->getName(), BB)); 1765 } 1766 1767 case OpAll : 1768 case OpAny : 1769 return mapValue(BV, 1770 transOCLAllAny(static_cast<SPIRVInstruction *>(BV), BB)); 1771 1772 case OpIsFinite : 1773 case OpIsInf : 1774 case OpIsNan : 1775 case OpIsNormal : 1776 case OpSignBitSet : 1777 return mapValue(BV, 1778 transOCLRelational(static_cast<SPIRVInstruction *>(BV), BB)); 1779 1780 default: { 1781 auto OC = BV->getOpCode(); 1782 if (isSPIRVCmpInstTransToLLVMInst(static_cast<SPIRVInstruction*>(BV))) { 1783 return mapValue(BV, transCmpInst(BV, BB, F)); 1784 } else if (OCLSPIRVBuiltinMap::rfind(OC, nullptr) && 1785 !isAtomicOpCode(OC) && 1786 !isGroupOpCode(OC) && 1787 !isPipeOpCode(OC)) { 1788 return mapValue(BV, transOCLBuiltinFromInst( 1789 static_cast<SPIRVInstruction *>(BV), BB)); 1790 } else if (isBinaryShiftLogicalBitwiseOpCode(OC) || 1791 isLogicalOpCode(OC)) { 1792 return mapValue(BV, transShiftLogicalBitwiseInst(BV, BB, F)); 1793 } else if (isCvtOpCode(OC)) { 1794 auto BI = static_cast<SPIRVInstruction *>(BV); 1795 Value *Inst = nullptr; 1796 if (BI->hasFPRoundingMode() || BI->isSaturatedConversion()) 1797 Inst = transOCLBuiltinFromInst(BI, BB); 1798 else 1799 Inst = transConvertInst(BV, F, BB); 1800 return mapValue(BV, Inst); 1801 } 1802 return mapValue(BV, transSPIRVBuiltinFromInst( 1803 static_cast<SPIRVInstruction *>(BV), BB)); 1804 } 1805 1806 SPIRVDBG(spvdbgs() << "Cannot translate " << *BV << '\n';) 1807 llvm_unreachable("Translation of SPIRV instruction not implemented"); 1808 return NULL; 1809 } 1810} 1811 1812template<class SourceTy, class FuncTy> 1813bool 1814SPIRVToLLVM::foreachFuncCtlMask(SourceTy Source, FuncTy Func) { 1815 SPIRVWord FCM = Source->getFuncCtlMask(); 1816 SPIRSPIRVFuncCtlMaskMap::foreach([&](Attribute::AttrKind Attr, 1817 SPIRVFunctionControlMaskKind Mask){ 1818 if (FCM & Mask) 1819 Func(Attr); 1820 }); 1821 return true; 1822} 1823 1824Function * 1825SPIRVToLLVM::transFunction(SPIRVFunction *BF) { 1826 auto Loc = FuncMap.find(BF); 1827 if (Loc != FuncMap.end()) 1828 return Loc->second; 1829 1830 auto IsKernel = BM->isEntryPoint(ExecutionModelKernel, BF->getId()); 1831 auto Linkage = IsKernel ? GlobalValue::ExternalLinkage : transLinkageType(BF); 1832 FunctionType *FT = dyn_cast<FunctionType>(transType(BF->getFunctionType())); 1833 Function *F = dyn_cast<Function>(mapValue(BF, Function::Create(FT, Linkage, 1834 BF->getName(), M))); 1835 assert(F); 1836 mapFunction(BF, F); 1837 if (!F->isIntrinsic()) { 1838 F->setCallingConv(IsKernel ? CallingConv::SPIR_KERNEL : 1839 CallingConv::SPIR_FUNC); 1840 if (isFuncNoUnwind()) 1841 F->addFnAttr(Attribute::NoUnwind); 1842 foreachFuncCtlMask(BF, [&](Attribute::AttrKind Attr){ 1843 F->addFnAttr(Attr); 1844 }); 1845 } 1846 1847 for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(); I != E; 1848 ++I) { 1849 auto BA = BF->getArgument(I->getArgNo()); 1850 mapValue(BA, static_cast<Argument*>(I)); 1851 setName(static_cast<Argument*>(I), BA); 1852 BA->foreachAttr([&](SPIRVFuncParamAttrKind Kind){ 1853 if (Kind == FunctionParameterAttributeNoWrite) 1854 return; 1855 F->addAttribute(I->getArgNo() + 1, SPIRSPIRVFuncParamAttrMap::rmap(Kind)); 1856 }); 1857 1858 SPIRVWord MaxOffset = 0; 1859 if (BA->hasDecorate(DecorationMaxByteOffset, 0, &MaxOffset)) { 1860 AttrBuilder Builder; 1861 Builder.addDereferenceableAttr(MaxOffset); 1862 I->addAttr(AttributeSet::get(*Context, I->getArgNo() + 1, Builder)); 1863 } 1864 } 1865 BF->foreachReturnValueAttr([&](SPIRVFuncParamAttrKind Kind){ 1866 if (Kind == FunctionParameterAttributeNoWrite) 1867 return; 1868 F->addAttribute(AttributeSet::ReturnIndex, 1869 SPIRSPIRVFuncParamAttrMap::rmap(Kind)); 1870 }); 1871 1872 // Creating all basic blocks before creating instructions. 1873 for (size_t I = 0, E = BF->getNumBasicBlock(); I != E; ++I) { 1874 transValue(BF->getBasicBlock(I), F, nullptr); 1875 } 1876 1877 for (size_t I = 0, E = BF->getNumBasicBlock(); I != E; ++I) { 1878 SPIRVBasicBlock *BBB = BF->getBasicBlock(I); 1879 BasicBlock *BB = dyn_cast<BasicBlock>(transValue(BBB, F, nullptr)); 1880 for (size_t BI = 0, BE = BBB->getNumInst(); BI != BE; ++BI) { 1881 SPIRVInstruction *BInst = BBB->getInst(BI); 1882 transValue(BInst, F, BB, false); 1883 } 1884 } 1885 return F; 1886} 1887 1888/// LLVM convert builtin functions is translated to two instructions: 1889/// y = i32 islessgreater(float x, float z) -> 1890/// y = i32 ZExt(bool LessGreater(float x, float z)) 1891/// When translating back, for simplicity, a trunc instruction is inserted 1892/// w = bool LessGreater(float x, float z) -> 1893/// w = bool Trunc(i32 islessgreater(float x, float z)) 1894/// Optimizer should be able to remove the redundant trunc/zext 1895void 1896SPIRVToLLVM::transOCLBuiltinFromInstPreproc(SPIRVInstruction* BI, Type *&RetTy, 1897 std::vector<SPIRVValue *> &Args) { 1898 if (!BI->hasType()) 1899 return; 1900 auto BT = BI->getType(); 1901 auto OC = BI->getOpCode(); 1902 if (isCmpOpCode(BI->getOpCode())) { 1903 if (BT->isTypeBool()) 1904 RetTy = IntegerType::getInt32Ty(*Context); 1905 else if (BT->isTypeVectorBool()) 1906 RetTy = VectorType::get(IntegerType::get(*Context, 1907 Args[0]->getType()->getVectorComponentType()->isTypeFloat(64)?64:32), 1908 BT->getVectorComponentCount()); 1909 else 1910 llvm_unreachable("invalid compare instruction"); 1911 } else if (OC == OpGenericCastToPtrExplicit) 1912 Args.pop_back(); 1913 else if (OC == OpImageRead && Args.size() > 2) { 1914 // Drop "Image operands" argument 1915 Args.erase(Args.begin() + 2); 1916 } 1917} 1918 1919Instruction* 1920SPIRVToLLVM::transOCLBuiltinPostproc(SPIRVInstruction* BI, 1921 CallInst* CI, BasicBlock* BB, const std::string &DemangledName) { 1922 auto OC = BI->getOpCode(); 1923 if (isCmpOpCode(OC) && 1924 BI->getType()->isTypeVectorOrScalarBool()) { 1925 return CastInst::Create(Instruction::Trunc, CI, transType(BI->getType()), 1926 "cvt", BB); 1927 } 1928 if (OC == OpImageSampleExplicitLod) 1929 return postProcessOCLReadImage(BI, CI, DemangledName); 1930 if (OC == OpImageWrite) { 1931 return postProcessOCLWriteImage(BI, CI, DemangledName); 1932 } 1933 if (OC == OpGenericPtrMemSemantics) 1934 return BinaryOperator::CreateShl(CI, getInt32(M, 8), "", BB); 1935 if (OC == OpImageQueryFormat) 1936 return BinaryOperator::CreateSub( 1937 CI, getInt32(M, OCLImageChannelDataTypeOffset), "", BB); 1938 if (OC == OpImageQueryOrder) 1939 return BinaryOperator::CreateSub( 1940 CI, getInt32(M, OCLImageChannelOrderOffset), "", BB); 1941 if (OC == OpBuildNDRange) 1942 return postProcessOCLBuildNDRange(BI, CI, DemangledName); 1943 if (OC == OpGroupAll || OC == OpGroupAny) 1944 return postProcessGroupAllAny(CI, DemangledName); 1945 if (SPIRVEnableStepExpansion && 1946 (DemangledName == "smoothstep" || 1947 DemangledName == "step")) 1948 return expandOCLBuiltinWithScalarArg(CI, DemangledName); 1949 return CI; 1950} 1951 1952Instruction * 1953SPIRVToLLVM::transBuiltinFromInst(const std::string& FuncName, 1954 SPIRVInstruction* BI, BasicBlock* BB) { 1955 std::string MangledName; 1956 auto Ops = BI->getOperands(); 1957 Type* RetTy = BI->hasType() ? transType(BI->getType()) : 1958 Type::getVoidTy(*Context); 1959 transOCLBuiltinFromInstPreproc(BI, RetTy, Ops); 1960 std::vector<Type*> ArgTys = transTypeVector( 1961 SPIRVInstruction::getOperandTypes(Ops)); 1962 bool HasFuncPtrArg = false; 1963 for (auto& I:ArgTys) { 1964 if (isa<FunctionType>(I)) { 1965 I = PointerType::get(I, SPIRAS_Private); 1966 HasFuncPtrArg = true; 1967 } 1968 } 1969 if (!HasFuncPtrArg) 1970 MangleOpenCLBuiltin(FuncName, ArgTys, MangledName); 1971 else 1972 MangledName = decorateSPIRVFunction(FuncName); 1973 Function* Func = M->getFunction(MangledName); 1974 FunctionType* FT = FunctionType::get(RetTy, ArgTys, false); 1975 // ToDo: Some intermediate functions have duplicate names with 1976 // different function types. This is OK if the function name 1977 // is used internally and finally translated to unique function 1978 // names. However it is better to have a way to differentiate 1979 // between intermidiate functions and final functions and make 1980 // sure final functions have unique names. 1981 SPIRVDBG( 1982 if (!HasFuncPtrArg && Func && Func->getFunctionType() != FT) { 1983 dbgs() << "Warning: Function name conflict:\n" 1984 << *Func << '\n' 1985 << " => " << *FT << '\n'; 1986 } 1987 ) 1988 if (!Func || Func->getFunctionType() != FT) { 1989 DEBUG(for (auto& I:ArgTys) { 1990 dbgs() << *I << '\n'; 1991 }); 1992 Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); 1993 Func->setCallingConv(CallingConv::SPIR_FUNC); 1994 if (isFuncNoUnwind()) 1995 Func->addFnAttr(Attribute::NoUnwind); 1996 } 1997 auto Call = CallInst::Create(Func, 1998 transValue(Ops, BB->getParent(), BB), "", BB); 1999 setName(Call, BI); 2000 setAttrByCalledFunc(Call); 2001 SPIRVDBG(spvdbgs() << "[transInstToBuiltinCall] " << *BI << " -> "; dbgs() << 2002 *Call << '\n';) 2003 Instruction *Inst = Call; 2004 Inst = transOCLBuiltinPostproc(BI, Call, BB, FuncName); 2005 return Inst; 2006} 2007 2008std::string 2009SPIRVToLLVM::getOCLBuiltinName(SPIRVInstruction* BI) { 2010 auto OC = BI->getOpCode(); 2011 if (OC == OpGenericCastToPtrExplicit) 2012 return getOCLGenericCastToPtrName(BI); 2013 if (isCvtOpCode(OC)) 2014 return getOCLConvertBuiltinName(BI); 2015 if (OC == OpBuildNDRange) { 2016 auto NDRangeInst = static_cast<SPIRVBuildNDRange *>(BI); 2017 auto EleTy = ((NDRangeInst->getOperands())[0])->getType(); 2018 int Dim = EleTy->isTypeArray() ? EleTy->getArrayLength() : 1; 2019 // cygwin does not have std::to_string 2020 ostringstream OS; 2021 OS << Dim; 2022 assert((EleTy->isTypeInt() && Dim == 1) || 2023 (EleTy->isTypeArray() && Dim >= 2 && Dim <= 3)); 2024 return std::string(kOCLBuiltinName::NDRangePrefix) + OS.str() + "D"; 2025 } 2026 auto Name = OCLSPIRVBuiltinMap::rmap(OC); 2027 2028 SPIRVType *T = nullptr; 2029 switch(OC) { 2030 case OpImageRead: 2031 T = BI->getType(); 2032 break; 2033 case OpImageWrite: 2034 T = BI->getOperands()[2]->getType(); 2035 break; 2036 default: 2037 // do nothing 2038 break; 2039 } 2040 if (T && T->isTypeVector()) 2041 T = T->getVectorComponentType(); 2042 if (T) 2043 Name += T->isTypeFloat()?'f':'i'; 2044 2045 return Name; 2046} 2047 2048Instruction * 2049SPIRVToLLVM::transOCLBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB) { 2050 assert(BB && "Invalid BB"); 2051 auto FuncName = getOCLBuiltinName(BI); 2052 return transBuiltinFromInst(FuncName, BI, BB); 2053} 2054 2055Instruction * 2056SPIRVToLLVM::transSPIRVBuiltinFromInst(SPIRVInstruction *BI, BasicBlock *BB) { 2057 assert(BB && "Invalid BB"); 2058 string Suffix = ""; 2059 if (BI->getOpCode() == OpCreatePipeFromPipeStorage) { 2060 auto CPFPS = static_cast<SPIRVCreatePipeFromPipeStorage*>(BI); 2061 assert(CPFPS->getType()->isTypePipe() && 2062 "Invalid type of CreatePipeFromStorage"); 2063 auto PipeType = static_cast<SPIRVTypePipe*>(CPFPS->getType()); 2064 switch (PipeType->getAccessQualifier()) { 2065 case AccessQualifierReadOnly: Suffix = "_read"; break; 2066 case AccessQualifierWriteOnly: Suffix = "_write"; break; 2067 case AccessQualifierReadWrite: Suffix = "_read_write"; break; 2068 } 2069 } 2070 2071 return transBuiltinFromInst(getSPIRVFuncName(BI->getOpCode(), Suffix), BI, BB); 2072} 2073 2074bool 2075SPIRVToLLVM::translate() { 2076 if (!transAddressingModel()) 2077 return false; 2078 2079 DbgTran.createCompileUnit(); 2080 DbgTran.addDbgInfoVersion(); 2081 2082 for (unsigned I = 0, E = BM->getNumVariables(); I != E; ++I) { 2083 auto BV = BM->getVariable(I); 2084 if (BV->getStorageClass() != StorageClassFunction) 2085 transValue(BV, nullptr, nullptr); 2086 } 2087 2088 for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) { 2089 transFunction(BM->getFunction(I)); 2090 } 2091 if (!transKernelMetadata()) 2092 return false; 2093 if (!transFPContractMetadata()) 2094 return false; 2095 if (!transSourceLanguage()) 2096 return false; 2097 if (!transSourceExtension()) 2098 return false; 2099 transGeneratorMD(); 2100 if (!transOCLBuiltinsFromVariables()) 2101 return false; 2102 if (!postProcessOCL()) 2103 return false; 2104 eraseUselessFunctions(M); 2105 DbgTran.finalize(); 2106 return true; 2107} 2108 2109bool 2110SPIRVToLLVM::transAddressingModel() { 2111 switch (BM->getAddressingModel()) { 2112 case AddressingModelPhysical64: 2113 M->setTargetTriple(SPIR_TARGETTRIPLE64); 2114 M->setDataLayout(SPIR_DATALAYOUT64); 2115 break; 2116 case AddressingModelPhysical32: 2117 M->setTargetTriple(SPIR_TARGETTRIPLE32); 2118 M->setDataLayout(SPIR_DATALAYOUT32); 2119 break; 2120 case AddressingModelLogical: 2121 // Do not set target triple and data layout 2122 break; 2123 default: 2124 SPIRVCKRT(0, InvalidAddressingModel, "Actual addressing mode is " + 2125 (unsigned)BM->getAddressingModel()); 2126 } 2127 return true; 2128} 2129 2130bool 2131SPIRVToLLVM::transDecoration(SPIRVValue *BV, Value *V) { 2132 if (!transAlign(BV, V)) 2133 return false; 2134 DbgTran.transDbgInfo(BV, V); 2135 return true; 2136} 2137 2138bool 2139SPIRVToLLVM::transFPContractMetadata() { 2140 bool ContractOff = false; 2141 for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) { 2142 SPIRVFunction *BF = BM->getFunction(I); 2143 if (!isOpenCLKernel(BF)) 2144 continue; 2145 if (BF->getExecutionMode(ExecutionModeContractionOff)) { 2146 ContractOff = true; 2147 break; 2148 } 2149 } 2150 if (!ContractOff) 2151 M->getOrInsertNamedMetadata(kSPIR2MD::FPContract); 2152 return true; 2153} 2154 2155std::string SPIRVToLLVM::transOCLImageTypeAccessQualifier( 2156 SPIRV::SPIRVTypeImage* ST) { 2157 return SPIRSPIRVAccessQualifierMap::rmap(ST->getAccessQualifier()); 2158} 2159 2160bool 2161SPIRVToLLVM::transNonTemporalMetadata(Instruction *I) { 2162 Constant* One = ConstantInt::get(Type::getInt32Ty(*Context), 1); 2163 MDNode *Node = MDNode::get(*Context, ConstantAsMetadata::get(One)); 2164 I->setMetadata(M->getMDKindID("nontemporal"), Node); 2165 return true; 2166} 2167 2168bool 2169SPIRVToLLVM::transKernelMetadata() { 2170 NamedMDNode *KernelMDs = M->getOrInsertNamedMetadata(SPIR_MD_KERNELS); 2171 for (unsigned I = 0, E = BM->getNumFunctions(); I != E; ++I) { 2172 SPIRVFunction *BF = BM->getFunction(I); 2173 Function *F = static_cast<Function *>(getTranslatedValue(BF)); 2174 assert(F && "Invalid translated function"); 2175 if (F->getCallingConv() != CallingConv::SPIR_KERNEL) 2176 continue; 2177 std::vector<llvm::Metadata*> KernelMD; 2178 KernelMD.push_back(ValueAsMetadata::get(F)); 2179 2180 // Generate metadata for kernel_arg_address_spaces 2181 addOCLKernelArgumentMetadata(Context, KernelMD, 2182 SPIR_MD_KERNEL_ARG_ADDR_SPACE, BF, 2183 [=](SPIRVFunctionParameter *Arg){ 2184 SPIRVType *ArgTy = Arg->getType(); 2185 SPIRAddressSpace AS = SPIRAS_Private; 2186 if (ArgTy->isTypePointer()) 2187 AS = SPIRSPIRVAddrSpaceMap::rmap(ArgTy->getPointerStorageClass()); 2188 else if (ArgTy->isTypeOCLImage() || ArgTy->isTypePipe()) 2189 AS = SPIRAS_Global; 2190 return ConstantAsMetadata::get( 2191 ConstantInt::get(Type::getInt32Ty(*Context), AS)); 2192 }); 2193 // Generate metadata for kernel_arg_access_qual 2194 addOCLKernelArgumentMetadata(Context, KernelMD, 2195 SPIR_MD_KERNEL_ARG_ACCESS_QUAL, BF, 2196 [=](SPIRVFunctionParameter *Arg){ 2197 std::string Qual; 2198 auto T = Arg->getType(); 2199 if (T->isTypeOCLImage()) { 2200 auto ST = static_cast<SPIRVTypeImage *>(T); 2201 Qual = transOCLImageTypeAccessQualifier(ST); 2202 } else if (T->isTypePipe()){ 2203 auto PT = static_cast<SPIRVTypePipe *>(T); 2204 Qual = transOCLPipeTypeAccessQualifier(PT); 2205 } else 2206 Qual = "none"; 2207 return MDString::get(*Context, Qual); 2208 }); 2209 // Generate metadata for kernel_arg_type 2210 addOCLKernelArgumentMetadata(Context, KernelMD, 2211 SPIR_MD_KERNEL_ARG_TYPE, BF, 2212 [=](SPIRVFunctionParameter *Arg){ 2213 return transOCLKernelArgTypeName(Arg); 2214 }); 2215 // Generate metadata for kernel_arg_type_qual 2216 addOCLKernelArgumentMetadata(Context, KernelMD, 2217 SPIR_MD_KERNEL_ARG_TYPE_QUAL, BF, 2218 [=](SPIRVFunctionParameter *Arg){ 2219 std::string Qual; 2220 if (Arg->hasDecorate(DecorationVolatile)) 2221 Qual = kOCLTypeQualifierName::Volatile; 2222 Arg->foreachAttr([&](SPIRVFuncParamAttrKind Kind){ 2223 Qual += Qual.empty() ? "" : " "; 2224 switch(Kind){ 2225 case FunctionParameterAttributeNoAlias: 2226 Qual += kOCLTypeQualifierName::Restrict; 2227 break; 2228 case FunctionParameterAttributeNoWrite: 2229 Qual += kOCLTypeQualifierName::Const; 2230 break; 2231 default: 2232 // do nothing. 2233 break; 2234 } 2235 }); 2236 if (Arg->getType()->isTypePipe()) { 2237 Qual += Qual.empty() ? "" : " "; 2238 Qual += kOCLTypeQualifierName::Pipe; 2239 } 2240 return MDString::get(*Context, Qual); 2241 }); 2242 // Generate metadata for kernel_arg_base_type 2243 addOCLKernelArgumentMetadata(Context, KernelMD, 2244 SPIR_MD_KERNEL_ARG_BASE_TYPE, BF, 2245 [=](SPIRVFunctionParameter *Arg){ 2246 return transOCLKernelArgTypeName(Arg); 2247 }); 2248 // Generate metadata for kernel_arg_name 2249 if (SPIRVGenKernelArgNameMD) { 2250 bool ArgHasName = true; 2251 BF->foreachArgument([&](SPIRVFunctionParameter *Arg){ 2252 ArgHasName &= !Arg->getName().empty(); 2253 }); 2254 if (ArgHasName) 2255 addOCLKernelArgumentMetadata(Context, KernelMD, 2256 SPIR_MD_KERNEL_ARG_NAME, BF, 2257 [=](SPIRVFunctionParameter *Arg){ 2258 return MDString::get(*Context, Arg->getName()); 2259 }); 2260 } 2261 // Generate metadata for reqd_work_group_size 2262 if (auto EM = BF->getExecutionMode(ExecutionModeLocalSize)) { 2263 KernelMD.push_back(getMDNodeStringIntVec(Context, 2264 kSPIR2MD::WGSize, EM->getLiterals())); 2265 } 2266 // Generate metadata for work_group_size_hint 2267 if (auto EM = BF->getExecutionMode(ExecutionModeLocalSizeHint)) { 2268 KernelMD.push_back(getMDNodeStringIntVec(Context, 2269 kSPIR2MD::WGSizeHint, EM->getLiterals())); 2270 } 2271 // Generate metadata for vec_type_hint 2272 if (auto EM = BF->getExecutionMode(ExecutionModeVecTypeHint)) { 2273 std::vector<Metadata*> MetadataVec; 2274 MetadataVec.push_back(MDString::get(*Context, kSPIR2MD::VecTyHint)); 2275 Type *VecHintTy = decodeVecTypeHint(*Context, EM->getLiterals()[0]); 2276 assert(VecHintTy); 2277 MetadataVec.push_back(ValueAsMetadata::get(UndefValue::get(VecHintTy))); 2278 MetadataVec.push_back( 2279 ConstantAsMetadata::get(ConstantInt::get(Type::getInt32Ty(*Context), 2280 1))); 2281 KernelMD.push_back(MDNode::get(*Context, MetadataVec)); 2282 } 2283 2284 llvm::MDNode *Node = MDNode::get(*Context, KernelMD); 2285 KernelMDs->addOperand(Node); 2286 } 2287 return true; 2288} 2289 2290bool 2291SPIRVToLLVM::transAlign(SPIRVValue *BV, Value *V) { 2292 if (auto AL = dyn_cast<AllocaInst>(V)) { 2293 SPIRVWord Align = 0; 2294 if (BV->hasAlignment(&Align)) 2295 AL->setAlignment(Align); 2296 return true; 2297 } 2298 if (auto GV = dyn_cast<GlobalVariable>(V)) { 2299 SPIRVWord Align = 0; 2300 if (BV->hasAlignment(&Align)) 2301 GV->setAlignment(Align); 2302 return true; 2303 } 2304 return true; 2305} 2306 2307void 2308SPIRVToLLVM::transOCLVectorLoadStore(std::string& UnmangledName, 2309 std::vector<SPIRVWord> &BArgs) { 2310 if (UnmangledName.find("vload") == 0 && 2311 UnmangledName.find("n") != std::string::npos) { 2312 if (BArgs.back() != 1) { 2313 std::stringstream SS; 2314 SS << BArgs.back(); 2315 UnmangledName.replace(UnmangledName.find("n"), 1, SS.str()); 2316 } else { 2317 UnmangledName.erase(UnmangledName.find("n"), 1); 2318 } 2319 BArgs.pop_back(); 2320 } else if (UnmangledName.find("vstore") == 0) { 2321 if (UnmangledName.find("n") != std::string::npos) { 2322 auto T = BM->getValueType(BArgs[0]); 2323 if (T->isTypeVector()) { 2324 auto W = T->getVectorComponentCount(); 2325 std::stringstream SS; 2326 SS << W; 2327 UnmangledName.replace(UnmangledName.find("n"), 1, SS.str()); 2328 } else { 2329 UnmangledName.erase(UnmangledName.find("n"), 1); 2330 } 2331 } 2332 if (UnmangledName.find("_r") != std::string::npos) { 2333 UnmangledName.replace(UnmangledName.find("_r"), 2, std::string("_") + 2334 SPIRSPIRVFPRoundingModeMap::rmap(static_cast<SPIRVFPRoundingModeKind>( 2335 BArgs.back()))); 2336 BArgs.pop_back(); 2337 } 2338 } 2339} 2340 2341// printf is not mangled. The function type should have just one argument. 2342// read_image*: the second argument should be mangled as sampler. 2343Instruction * 2344SPIRVToLLVM::transOCLBuiltinFromExtInst(SPIRVExtInst *BC, BasicBlock *BB) { 2345 assert(BB && "Invalid BB"); 2346 std::string MangledName; 2347 SPIRVWord EntryPoint = BC->getExtOp(); 2348 SPIRVExtInstSetKind Set = BM->getBuiltinSet(BC->getExtSetId()); 2349 bool IsVarArg = false; 2350 bool IsPrintf = false; 2351 std::string UnmangledName; 2352 auto BArgs = BC->getArguments(); 2353 2354 (void) Set; 2355 assert (Set == SPIRVEIS_OpenCL && "Not OpenCL extended instruction"); 2356 if (EntryPoint == OpenCLLIB::Printf) 2357 IsPrintf = true; 2358 else { 2359 UnmangledName = OCLExtOpMap::map(static_cast<OCLExtOpKind>( 2360 EntryPoint)); 2361 } 2362 2363 SPIRVDBG(spvdbgs() << "[transOCLBuiltinFromExtInst] OrigUnmangledName: " << 2364 UnmangledName << '\n'); 2365 transOCLVectorLoadStore(UnmangledName, BArgs); 2366 2367 std::vector<Type *> ArgTypes = transTypeVector(BC->getValueTypes(BArgs)); 2368 2369 if (IsPrintf) { 2370 MangledName = "printf"; 2371 IsVarArg = true; 2372 ArgTypes.resize(1); 2373 } else if (UnmangledName.find("read_image") == 0) { 2374 auto ModifiedArgTypes = ArgTypes; 2375 ModifiedArgTypes[1] = getOrCreateOpaquePtrType(M, "opencl.sampler_t"); 2376 MangleOpenCLBuiltin(UnmangledName, ModifiedArgTypes, MangledName); 2377 } else { 2378 MangleOpenCLBuiltin(UnmangledName, ArgTypes, MangledName); 2379 } 2380 SPIRVDBG(spvdbgs() << "[transOCLBuiltinFromExtInst] ModifiedUnmangledName: " << 2381 UnmangledName << " MangledName: " << MangledName << '\n'); 2382 2383 FunctionType *FT = FunctionType::get( 2384 transType(BC->getType()), 2385 ArgTypes, 2386 IsVarArg); 2387 Function *F = M->getFunction(MangledName); 2388 if (!F) { 2389 F = Function::Create(FT, 2390 GlobalValue::ExternalLinkage, 2391 MangledName, 2392 M); 2393 F->setCallingConv(CallingConv::SPIR_FUNC); 2394 if (isFuncNoUnwind()) 2395 F->addFnAttr(Attribute::NoUnwind); 2396 } 2397 auto Args = transValue(BC->getValues(BArgs), F, BB); 2398 SPIRVDBG(dbgs() << "[transOCLBuiltinFromExtInst] Function: " << *F << 2399 ", Args: "; 2400 for (auto &I:Args) dbgs() << *I << ", "; dbgs() << '\n'); 2401 CallInst *Call = CallInst::Create(F, 2402 Args, 2403 BC->getName(), 2404 BB); 2405 setCallingConv(Call); 2406 addFnAttr(Context, Call, Attribute::NoUnwind); 2407 return transOCLBuiltinPostproc(BC, Call, BB, UnmangledName); 2408} 2409 2410CallInst * 2411SPIRVToLLVM::transOCLBarrier(BasicBlock *BB, SPIRVWord ExecScope, 2412 SPIRVWord MemSema, SPIRVWord MemScope) { 2413 SPIRVWord Ver = 0; 2414 BM->getSourceLanguage(&Ver); 2415 2416 Type* Int32Ty = Type::getInt32Ty(*Context); 2417 Type* VoidTy = Type::getVoidTy(*Context); 2418 2419 std::string FuncName; 2420 SmallVector<Type *, 2> ArgTy; 2421 SmallVector<Value *, 2> Arg; 2422 2423 Constant *MemFenceFlags = 2424 ConstantInt::get(Int32Ty, rmapBitMask<OCLMemFenceMap>(MemSema)); 2425 2426 FuncName = (ExecScope == ScopeWorkgroup) ? kOCLBuiltinName::WorkGroupBarrier 2427 : kOCLBuiltinName::SubGroupBarrier; 2428 2429 if (ExecScope == ScopeWorkgroup && Ver > 0 && Ver <= kOCLVer::CL12) { 2430 FuncName = kOCLBuiltinName::Barrier; 2431 ArgTy.push_back(Int32Ty); 2432 Arg.push_back(MemFenceFlags); 2433 } else { 2434 Constant *Scope = ConstantInt::get(Int32Ty, OCLMemScopeMap::rmap( 2435 static_cast<spv::Scope>(MemScope))); 2436 2437 ArgTy.append(2, Int32Ty); 2438 Arg.push_back(MemFenceFlags); 2439 Arg.push_back(Scope); 2440 } 2441 2442 std::string MangledName; 2443 2444 MangleOpenCLBuiltin(FuncName, ArgTy, MangledName); 2445 Function *Func = M->getFunction(MangledName); 2446 if (!Func) { 2447 FunctionType *FT = FunctionType::get(VoidTy, ArgTy, false); 2448 Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); 2449 Func->setCallingConv(CallingConv::SPIR_FUNC); 2450 if (isFuncNoUnwind()) 2451 Func->addFnAttr(Attribute::NoUnwind); 2452 } 2453 2454 return CallInst::Create(Func, Arg, "", BB); 2455} 2456 2457CallInst * 2458SPIRVToLLVM::transOCLMemFence(BasicBlock *BB, 2459 SPIRVWord MemSema, SPIRVWord MemScope) { 2460 SPIRVWord Ver = 0; 2461 BM->getSourceLanguage(&Ver); 2462 2463 Type* Int32Ty = Type::getInt32Ty(*Context); 2464 Type* VoidTy = Type::getVoidTy(*Context); 2465 2466 std::string FuncName; 2467 SmallVector<Type *, 3> ArgTy; 2468 SmallVector<Value *, 3> Arg; 2469 2470 Constant *MemFenceFlags = 2471 ConstantInt::get(Int32Ty, rmapBitMask<OCLMemFenceMap>(MemSema)); 2472 2473 if (Ver > 0 && Ver <= kOCLVer::CL12) { 2474 FuncName = kOCLBuiltinName::MemFence; 2475 ArgTy.push_back(Int32Ty); 2476 Arg.push_back(MemFenceFlags); 2477 } else { 2478 Constant *Order = 2479 ConstantInt::get(Int32Ty, mapSPIRVMemOrderToOCL(MemSema)); 2480 2481 Constant *Scope = ConstantInt::get(Int32Ty, OCLMemScopeMap::rmap( 2482 static_cast<spv::Scope>(MemScope))); 2483 2484 FuncName = kOCLBuiltinName::AtomicWorkItemFence; 2485 ArgTy.append(3, Int32Ty); 2486 Arg.push_back(MemFenceFlags); 2487 Arg.push_back(Order); 2488 Arg.push_back(Scope); 2489 } 2490 2491 std::string MangledName; 2492 2493 MangleOpenCLBuiltin(FuncName, ArgTy, MangledName); 2494 Function *Func = M->getFunction(MangledName); 2495 if (!Func) { 2496 FunctionType *FT = FunctionType::get(VoidTy, ArgTy, false); 2497 Func = Function::Create(FT, GlobalValue::ExternalLinkage, MangledName, M); 2498 Func->setCallingConv(CallingConv::SPIR_FUNC); 2499 if (isFuncNoUnwind()) 2500 Func->addFnAttr(Attribute::NoUnwind); 2501 } 2502 2503 return CallInst::Create(Func, Arg, "", BB); 2504} 2505 2506Instruction * 2507SPIRVToLLVM::transOCLBarrierFence(SPIRVInstruction *MB, BasicBlock *BB) { 2508 assert(BB && "Invalid BB"); 2509 std::string FuncName; 2510 auto getIntVal = [](SPIRVValue *value){ 2511 return static_cast<SPIRVConstant*>(value)->getZExtIntValue(); 2512 }; 2513 2514 CallInst* Call = nullptr; 2515 2516 if (MB->getOpCode() == OpMemoryBarrier) { 2517 auto MemB = static_cast<SPIRVMemoryBarrier*>(MB); 2518 2519 SPIRVWord MemScope = getIntVal(MemB->getOpValue(0)); 2520 SPIRVWord MemSema = getIntVal(MemB->getOpValue(1)); 2521 2522 Call = transOCLMemFence(BB, MemSema, MemScope); 2523 } else if (MB->getOpCode() == OpControlBarrier) { 2524 auto CtlB = static_cast<SPIRVControlBarrier*>(MB); 2525 2526 SPIRVWord ExecScope = getIntVal(CtlB->getExecScope()); 2527 SPIRVWord MemSema = getIntVal(CtlB->getMemSemantic()); 2528 SPIRVWord MemScope = getIntVal(CtlB->getMemScope()); 2529 2530 Call = transOCLBarrier(BB, ExecScope, MemSema, MemScope); 2531 } else { 2532 llvm_unreachable("Invalid instruction"); 2533 } 2534 2535 setName(Call, MB); 2536 setAttrByCalledFunc(Call); 2537 SPIRVDBG(spvdbgs() << "[transBarrier] " << *MB << " -> "; 2538 dbgs() << *Call << '\n';) 2539 2540 return Call; 2541} 2542 2543// SPIR-V only contains language version. Use OpenCL language version as 2544// SPIR version. 2545bool 2546SPIRVToLLVM::transSourceLanguage() { 2547 SPIRVWord Ver = 0; 2548 SourceLanguage Lang = BM->getSourceLanguage(&Ver); 2549 assert((Lang == SourceLanguageOpenCL_C || 2550 Lang == SourceLanguageOpenCL_CPP) && "Unsupported source language"); 2551 unsigned short Major = 0; 2552 unsigned char Minor = 0; 2553 unsigned char Rev = 0; 2554 std::tie(Major, Minor, Rev) = decodeOCLVer(Ver); 2555 SPIRVMDBuilder Builder(*M); 2556 Builder.addNamedMD(kSPIRVMD::Source) 2557 .addOp() 2558 .add(Lang) 2559 .add(Ver) 2560 .done(); 2561 // ToDo: Phasing out usage of old SPIR metadata 2562 if (Ver <= kOCLVer::CL12) 2563 addOCLVersionMetadata(Context, M, kSPIR2MD::SPIRVer, 1, 2); 2564 else 2565 addOCLVersionMetadata(Context, M, kSPIR2MD::SPIRVer, 2, 0); 2566 2567 addOCLVersionMetadata(Context, M, kSPIR2MD::OCLVer, Major, Minor); 2568 return true; 2569} 2570 2571bool 2572SPIRVToLLVM::transSourceExtension() { 2573 auto ExtSet = rmap<OclExt::Kind>(BM->getExtension()); 2574 auto CapSet = rmap<OclExt::Kind>(BM->getCapability()); 2575 ExtSet.insert(CapSet.begin(), CapSet.end()); 2576 auto OCLExtensions = map<std::string>(ExtSet); 2577 std::set<std::string> OCLOptionalCoreFeatures; 2578 static const char *OCLOptCoreFeatureNames[] = { 2579 "cl_images", "cl_doubles", 2580 }; 2581 for (auto &I : OCLOptCoreFeatureNames) { 2582 auto Loc = OCLExtensions.find(I); 2583 if (Loc != OCLExtensions.end()) { 2584 OCLExtensions.erase(Loc); 2585 OCLOptionalCoreFeatures.insert(I); 2586 } 2587 } 2588 addNamedMetadataStringSet(Context, M, kSPIR2MD::Extensions, OCLExtensions); 2589 addNamedMetadataStringSet(Context, M, kSPIR2MD::OptFeatures, 2590 OCLOptionalCoreFeatures); 2591 return true; 2592} 2593 2594// If the argument is unsigned return uconvert*, otherwise return convert*. 2595std::string 2596SPIRVToLLVM::getOCLConvertBuiltinName(SPIRVInstruction* BI) { 2597 auto OC = BI->getOpCode(); 2598 assert(isCvtOpCode(OC) && "Not convert instruction"); 2599 auto U = static_cast<SPIRVUnary *>(BI); 2600 std::string Name; 2601 if (isCvtFromUnsignedOpCode(OC)) 2602 Name = "u"; 2603 Name += "convert_"; 2604 Name += mapSPIRVTypeToOCLType(U->getType(), 2605 !isCvtToUnsignedOpCode(OC)); 2606 SPIRVFPRoundingModeKind Rounding; 2607 if (U->isSaturatedConversion()) 2608 Name += "_sat"; 2609 if (U->hasFPRoundingMode(&Rounding)) { 2610 Name += "_"; 2611 Name += SPIRSPIRVFPRoundingModeMap::rmap(Rounding); 2612 } 2613 return Name; 2614} 2615 2616//Check Address Space of the Pointer Type 2617std::string 2618SPIRVToLLVM::getOCLGenericCastToPtrName(SPIRVInstruction* BI) { 2619 auto GenericCastToPtrInst = BI->getType()->getPointerStorageClass(); 2620 switch (GenericCastToPtrInst) { 2621 case StorageClassCrossWorkgroup: 2622 return std::string(kOCLBuiltinName::ToGlobal); 2623 case StorageClassWorkgroup: 2624 return std::string(kOCLBuiltinName::ToLocal); 2625 case StorageClassFunction: 2626 return std::string(kOCLBuiltinName::ToPrivate); 2627 default: 2628 llvm_unreachable("Invalid address space"); 2629 return ""; 2630 } 2631} 2632 2633llvm::GlobalValue::LinkageTypes 2634SPIRVToLLVM::transLinkageType(const SPIRVValue* V) { 2635 if (V->getLinkageType() == LinkageTypeInternal) { 2636 return GlobalValue::InternalLinkage; 2637 } 2638 else if (V->getLinkageType() == LinkageTypeImport) { 2639 // Function declaration 2640 if (V->getOpCode() == OpFunction) { 2641 if (static_cast<const SPIRVFunction*>(V)->getNumBasicBlock() == 0) 2642 return GlobalValue::ExternalLinkage; 2643 } 2644 // Variable declaration 2645 if (V->getOpCode() == OpVariable) { 2646 if (static_cast<const SPIRVVariable*>(V)->getInitializer() == 0) 2647 return GlobalValue::ExternalLinkage; 2648 } 2649 // Definition 2650 return GlobalValue::AvailableExternallyLinkage; 2651 } 2652 else {// LinkageTypeExport 2653 if (V->getOpCode() == OpVariable) { 2654 if (static_cast<const SPIRVVariable*>(V)->getInitializer() == 0 ) 2655 // Tentative definition 2656 return GlobalValue::CommonLinkage; 2657 } 2658 return GlobalValue::ExternalLinkage; 2659 } 2660} 2661 2662Instruction *SPIRVToLLVM::transOCLAllAny(SPIRVInstruction *I, BasicBlock *BB) { 2663 CallInst *CI = cast<CallInst>(transSPIRVBuiltinFromInst(I, BB)); 2664 AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); 2665 return cast<Instruction>(mapValue( 2666 I, mutateCallInstOCL( 2667 M, CI, 2668 [=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { 2669 Type *Int32Ty = Type::getInt32Ty(*Context); 2670 auto OldArg = CI->getOperand(0); 2671 auto NewArgTy = VectorType::get( 2672 Int32Ty, OldArg->getType()->getVectorNumElements()); 2673 auto NewArg = 2674 CastInst::CreateSExtOrBitCast(OldArg, NewArgTy, "", CI); 2675 Args[0] = NewArg; 2676 RetTy = Int32Ty; 2677 return CI->getCalledFunction()->getName(); 2678 }, 2679 [=](CallInst *NewCI) -> Instruction * { 2680 return CastInst::CreateTruncOrBitCast( 2681 NewCI, Type::getInt1Ty(*Context), "", NewCI->getNextNode()); 2682 }, 2683 &Attrs))); 2684} 2685 2686Instruction *SPIRVToLLVM::transOCLRelational(SPIRVInstruction *I, BasicBlock *BB) { 2687 CallInst *CI = cast<CallInst>(transSPIRVBuiltinFromInst(I, BB)); 2688 AttributeSet Attrs = CI->getCalledFunction()->getAttributes(); 2689 return cast<Instruction>(mapValue( 2690 I, mutateCallInstOCL( 2691 M, CI, 2692 [=](CallInst *, std::vector<Value *> &Args, llvm::Type *&RetTy) { 2693 Type *IntTy = Type::getInt32Ty(*Context); 2694 RetTy = IntTy; 2695 if (CI->getType()->isVectorTy()) { 2696 if (cast<VectorType>(CI->getOperand(0)->getType()) 2697 ->getElementType() 2698 ->isDoubleTy()) 2699 IntTy = Type::getInt64Ty(*Context); 2700 if (cast<VectorType>(CI->getOperand(0)->getType()) 2701 ->getElementType() 2702 ->isHalfTy()) 2703 IntTy = Type::getInt16Ty(*Context); 2704 RetTy = VectorType::get(IntTy, 2705 CI->getType()->getVectorNumElements()); 2706 } 2707 return CI->getCalledFunction()->getName(); 2708 }, 2709 [=](CallInst *NewCI) -> Instruction * { 2710 Type *RetTy = Type::getInt1Ty(*Context); 2711 if (NewCI->getType()->isVectorTy()) 2712 RetTy = 2713 VectorType::get(Type::getInt1Ty(*Context), 2714 NewCI->getType()->getVectorNumElements()); 2715 return CastInst::CreateTruncOrBitCast(NewCI, RetTy, "", 2716 NewCI->getNextNode()); 2717 }, 2718 &Attrs))); 2719} 2720} 2721 2722bool 2723llvm::ReadSPIRV(LLVMContext &C, std::istream &IS, Module *&M, 2724 std::string &ErrMsg) { 2725 M = new Module("", C); 2726 std::unique_ptr<SPIRVModule> BM(SPIRVModule::createSPIRVModule()); 2727 2728 IS >> *BM; 2729 2730 SPIRVToLLVM BTL(M, BM.get()); 2731 bool Succeed = true; 2732 if (!BTL.translate()) { 2733 BM->getError(ErrMsg); 2734 Succeed = false; 2735 } 2736 legacy::PassManager PassMgr; 2737 PassMgr.add(createSPIRVToOCL20()); 2738 PassMgr.add(createOCL20To12()); 2739 PassMgr.run(*M); 2740 2741 if (DbgSaveTmpLLVM) 2742 dumpLLVM(M, DbgTmpLLVMFileName); 2743 if (!Succeed) { 2744 delete M; 2745 M = nullptr; 2746 } 2747 return Succeed; 2748} 2749