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