//===- OCL20ToSPIRV.cpp - Transform OCL20 to SPIR-V builtins -----*- C++ -*-===//
//
// The LLVM/SPIRV Translator
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
// Copyright (c) 2014 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a
// copy of this software and associated documentation files (the "Software"),
// to deal with the Software without restriction, including without limitation
// the rights to use, copy, modify, merge, publish, distribute, sublicense,
// and/or sell copies of the Software, and to permit persons to whom the
// Software is furnished to do so, subject to the following conditions:
//
// Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimers in the documentation
// and/or other materials provided with the distribution.
// Neither the names of Advanced Micro Devices, Inc., nor the names of its
// contributors may be used to endorse or promote products derived from this
// Software without specific prior written permission.
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH
// THE SOFTWARE.
//
//===----------------------------------------------------------------------===//
//
// This file implements translation of OCL20 builtin functions.
//
//===----------------------------------------------------------------------===//
#define DEBUG_TYPE "cl20tospv"
#include "SPIRVInternal.h"
#include "OCLUtil.h"
#include "OCLTypeToSPIRV.h"
#include "llvm/ADT/StringSwitch.h"
#include "llvm/IR/InstVisitor.h"
#include "llvm/IR/Instructions.h"
#include "llvm/IR/Instruction.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/Verifier.h"
#include "llvm/Pass.h"
#include "llvm/PassSupport.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/raw_ostream.h"
#include <set>
using namespace llvm;
using namespace SPIRV;
using namespace OCLUtil;
namespace SPIRV {
static size_t
getOCLCpp11AtomicMaxNumOps(StringRef Name) {
return StringSwitch<size_t>(Name)
.Cases("load", "flag_test_and_set", "flag_clear", 3)
.Cases("store", "exchange", 4)
.StartsWith("compare_exchange", 6)
.StartsWith("fetch", 4)
.Default(0);
}
class OCL20ToSPIRV: public ModulePass,
public InstVisitor<OCL20ToSPIRV> {
public:
OCL20ToSPIRV():ModulePass(ID), M(nullptr), Ctx(nullptr), CLVer(0) {
initializeOCL20ToSPIRVPass(*PassRegistry::getPassRegistry());
}
virtual bool runOnModule(Module &M);
void getAnalysisUsage(AnalysisUsage &AU) const {
AU.addRequired<OCLTypeToSPIRV>();
}
virtual void visitCallInst(CallInst &CI);
/// Transform barrier/work_group_barrier/sub_group_barrier
/// to __spirv_ControlBarrier.
/// barrier(flag) =>
/// __spirv_ControlBarrier(workgroup, workgroup, map(flag))
/// work_group_barrier(scope, flag) =>
/// __spirv_ControlBarrier(workgroup, map(scope), map(flag))
/// sub_group_barrier(scope, flag) =>
/// __spirv_ControlBarrier(subgroup, map(scope), map(flag))
void visitCallBarrier(CallInst *CI);
/// Erase useless convert functions.
/// \return true if the call instruction is erased.
bool eraseUselessConvert(CallInst *Call, const std::string &MangledName,
const std::string &DeMangledName);
/// Transform convert_ to
/// __spirv_{CastOpName}_R{TargeTyName}{_sat}{_rt[p|n|z|e]}
void visitCallConvert(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform async_work_group{_strided}_copy.
/// async_work_group_copy(dst, src, n, event)
/// => async_work_group_strided_copy(dst, src, n, 1, event)
/// async_work_group_strided_copy(dst, src, n, stride, event)
/// => __spirv_AsyncGroupCopy(ScopeWorkGroup, dst, src, n, stride, event)
void visitCallAsyncWorkGroupCopy(CallInst *CI,
const std::string &DemangledName);
/// Transform OCL builtin function to SPIR-V builtin function.
void transBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info);
/// Transform OCL work item builtin functions to SPIR-V builtin variables.
void transWorkItemBuiltinsToVariables();
/// Transform atomic_work_item_fence/mem_fence to __spirv_MemoryBarrier.
/// func(flag, order, scope) =>
/// __spirv_MemoryBarrier(map(scope), map(flag)|map(order))
void transMemoryBarrier(CallInst *CI, AtomicWorkItemFenceLiterals);
/// Transform all to __spirv_Op(All|Any). Note that the types mismatch so
// some extra code is emitted to convert between the two.
void visitCallAllAny(spv::Op OC, CallInst *CI);
/// Transform atomic_* to __spirv_Atomic*.
/// atomic_x(ptr_arg, args, order, scope) =>
/// __spirv_AtomicY(ptr_arg, map(order), map(scope), args)
void transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info);
/// Transform atomic_work_item_fence to __spirv_MemoryBarrier.
/// atomic_work_item_fence(flag, order, scope) =>
/// __spirv_MemoryBarrier(map(scope), map(flag)|map(order))
void visitCallAtomicWorkItemFence(CallInst *CI);
/// Transform atomic_compare_exchange call.
/// In atomic_compare_exchange, the expected value parameter is a pointer.
/// However in SPIR-V it is a value. The transformation adds a load
/// instruction, result of which is passed to atomic_compare_exchange as
/// argument.
/// The transformation adds a store instruction after the call, to update the
/// value in expected with the value pointed to by object. Though, it is not
/// necessary in case they are equal, this approach makes result code simpler.
/// Also ICmp instruction is added, because the call must return result of
/// comparison.
/// \returns the call instruction of atomic_compare_exchange_strong.
CallInst *visitCallAtomicCmpXchg(CallInst *CI,
const std::string &DemangledName);
/// Transform atomic_init.
/// atomic_init(p, x) => store p, x
void visitCallAtomicInit(CallInst *CI);
/// Transform legacy OCL 1.x atomic builtins to SPIR-V builtins for extensions
/// cl_khr_int64_base_atomics
/// cl_khr_int64_extended_atomics
/// Do nothing if the called function is not a legacy atomic builtin.
void visitCallAtomicLegacy(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform OCL 2.0 C++11 atomic builtins to SPIR-V builtins.
/// Do nothing if the called function is not a C++11 atomic builtin.
void visitCallAtomicCpp11(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform OCL builtin function to SPIR-V builtin function.
/// Assuming there is a simple name mapping without argument changes.
/// Should be called at last.
void visitCallBuiltinSimple(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform get_image_{width|height|depth|dim}.
/// get_image_xxx(...) =>
/// dimension = __spirv_ImageQuerySizeLod_R{ReturnType}(...);
/// return dimension.{x|y|z};
void visitCallGetImageSize(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform {work|sub}_group_x =>
/// __spirv_{OpName}
///
/// Special handling of work_group_broadcast.
/// work_group_broadcast(a, x, y, z)
/// =>
/// __spirv_GroupBroadcast(a, vec3(x, y, z))
void visitCallGroupBuiltin(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform mem_fence to __spirv_MemoryBarrier.
/// mem_fence(flag) => __spirv_MemoryBarrier(Workgroup, map(flag))
void visitCallMemFence(CallInst *CI);
void visitCallNDRange(CallInst *CI, const std::string &DemangledName);
/// Transform OCL pipe builtin function to SPIR-V pipe builtin function.
void visitCallPipeBuiltin(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform read_image with sampler arguments.
/// read_image(image, sampler, ...) =>
/// sampled_image = __spirv_SampledImage(image, sampler);
/// return __spirv_ImageSampleExplicitLod_R{ReturnType}(sampled_image, ...);
void visitCallReadImageWithSampler(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform read_image with msaa image arguments.
/// Sample argument must be acoded as Image Operand.
void visitCallReadImageMSAA(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform {read|write}_image without sampler arguments.
void visitCallReadWriteImage(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform to_{global|local|private}.
///
/// T* a = ...;
/// addr T* b = to_addr(a);
/// =>
/// i8* x = cast<i8*>(a);
/// addr i8* y = __spirv_GenericCastToPtr_ToAddr(x);
/// addr T* b = cast<addr T*>(y);
void visitCallToAddr(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform return type of relatinal built-in functions like isnan, isfinite
/// to boolean values.
void visitCallRelational(CallInst *CI, const std::string &DemangledName);
/// Transform vector load/store functions to SPIR-V extended builtin
/// functions
/// {vload|vstore{a}}{_half}{n}{_rte|_rtz|_rtp|_rtn} =>
/// __spirv_ocl_{ExtendedInstructionOpCodeName}__R{ReturnType}
void visitCallVecLoadStore(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transforms get_mem_fence built-in to SPIR-V function and aligns result values with SPIR 1.2.
/// get_mem_fence(ptr) => __spirv_GenericPtrMemSemantics
/// GenericPtrMemSemantics valid values are 0x100, 0x200 and 0x300, where is
/// SPIR 1.2 defines them as 0x1, 0x2 and 0x3, so this function adjusts
/// GenericPtrMemSemantics results to SPIR 1.2 values.
void visitCallGetFence(CallInst *CI, StringRef MangledName, const std::string& DemangledName);
/// Transforms OpDot instructions with a scalar type to a fmul instruction
void visitCallDot(CallInst *CI);
/// Fixes for built-in functions with vector+scalar arguments that are
/// translated to the SPIR-V instructions where all arguments must have the
/// same type.
void visitCallScalToVec(CallInst *CI, StringRef MangledName,
const std::string &DemangledName);
/// Transform get_image_channel_{order|data_type} built-in functions to
/// __spirv_ocl_{ImageQueryOrder|ImageQueryFormat}
void visitCallGetImageChannel(CallInst *CI, StringRef MangledName,
const std::string &DemangledName,
unsigned int Offset);
void visitDbgInfoIntrinsic(DbgInfoIntrinsic &I){
I.dropAllReferences();
I.eraseFromParent();
}
static char ID;
private:
Module *M;
LLVMContext *Ctx;
unsigned CLVer; /// OpenCL version as major*10+minor
std::set<Value *> ValuesToDelete;
ConstantInt *addInt32(int I) {
return getInt32(M, I);
}
ConstantInt *addSizet(uint64_t I) {
return getSizet(M, I);
}
/// Get vector width from OpenCL vload* function name.
SPIRVWord getVecLoadWidth(const std::string& DemangledName) {
SPIRVWord Width = 0;
if (DemangledName == "vloada_half")
Width = 1;
else {
unsigned Loc = 5;
if (DemangledName.find("vload_half") == 0)
Loc = 10;
else if (DemangledName.find("vloada_half") == 0)
Loc = 11;
std::stringstream SS(DemangledName.substr(Loc));
SS >> Width;
}
return Width;
}
/// Transform OpenCL vload/vstore function name.
void transVecLoadStoreName(std::string& DemangledName,
const std::string &Stem, bool AlwaysN) {
auto HalfStem = Stem + "_half";
auto HalfStemR = HalfStem + "_r";
if (!AlwaysN && DemangledName == HalfStem)
return;
if (!AlwaysN && DemangledName.find(HalfStemR) == 0) {
DemangledName = HalfStemR;
return;
}
if (DemangledName.find(HalfStem) == 0) {
auto OldName = DemangledName;
DemangledName = HalfStem + "n";
if (OldName.find("_r") != std::string::npos)
DemangledName += "_r";
return;
}
if (DemangledName.find(Stem) == 0) {
DemangledName = Stem + "n";
return;
}
}
};
char OCL20ToSPIRV::ID = 0;
bool
OCL20ToSPIRV::runOnModule(Module& Module) {
M = &Module;
Ctx = &M->getContext();
auto Src = getSPIRVSource(&Module);
if (std::get<0>(Src) != spv::SourceLanguageOpenCL_C)
return false;
CLVer = std::get<1>(Src);
if (CLVer > kOCLVer::CL20)
return false;
DEBUG(dbgs() << "Enter OCL20ToSPIRV:\n");
transWorkItemBuiltinsToVariables();
visit(*M);
for (auto &I:ValuesToDelete)
if (auto Inst = dyn_cast<Instruction>(I))
Inst->eraseFromParent();
for (auto &I:ValuesToDelete)
if (auto GV = dyn_cast<GlobalValue>(I))
GV->eraseFromParent();
DEBUG(dbgs() << "After OCL20ToSPIRV:\n" << *M);
std::string Err;
raw_string_ostream ErrorOS(Err);
if (verifyModule(*M, &ErrorOS)){
DEBUG(errs() << "Fails to verify module: " << ErrorOS.str());
}
return true;
}
// The order of handling OCL builtin functions is important.
// Workgroup functions need to be handled before pipe functions since
// there are functions fall into both categories.
void
OCL20ToSPIRV::visitCallInst(CallInst& CI) {
DEBUG(dbgs() << "[visistCallInst] " << CI << '\n');
auto F = CI.getCalledFunction();
if (!F)
return;
auto MangledName = F->getName();
std::string DemangledName;
if (!oclIsBuiltin(MangledName, &DemangledName))
return;
DEBUG(dbgs() << "DemangledName: " << DemangledName << '\n');
if (DemangledName.find(kOCLBuiltinName::NDRangePrefix) == 0) {
visitCallNDRange(&CI, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::All) {
visitCallAllAny(OpAll, &CI);
return;
}
if (DemangledName == kOCLBuiltinName::Any) {
visitCallAllAny(OpAny, &CI);
return;
}
if (DemangledName.find(kOCLBuiltinName::AsyncWorkGroupCopy) == 0 ||
DemangledName.find(kOCLBuiltinName::AsyncWorkGroupStridedCopy) == 0) {
visitCallAsyncWorkGroupCopy(&CI, DemangledName);
return;
}
if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 ||
DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) {
auto PCI = &CI;
if (DemangledName == kOCLBuiltinName::AtomicInit) {
visitCallAtomicInit(PCI);
return;
}
if (DemangledName == kOCLBuiltinName::AtomicWorkItemFence) {
visitCallAtomicWorkItemFence(PCI);
return;
}
if (DemangledName == kOCLBuiltinName::AtomicCmpXchgWeak ||
DemangledName == kOCLBuiltinName::AtomicCmpXchgStrong ||
DemangledName == kOCLBuiltinName::AtomicCmpXchgWeakExplicit ||
DemangledName == kOCLBuiltinName::AtomicCmpXchgStrongExplicit) {
assert(CLVer == kOCLVer::CL20 && "Wrong version of OpenCL");
PCI = visitCallAtomicCmpXchg(PCI, DemangledName);
}
visitCallAtomicLegacy(PCI, MangledName, DemangledName);
visitCallAtomicCpp11(PCI, MangledName, DemangledName);
return;
}
if (DemangledName.find(kOCLBuiltinName::ConvertPrefix) == 0) {
visitCallConvert(&CI, MangledName, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::GetImageWidth ||
DemangledName == kOCLBuiltinName::GetImageHeight ||
DemangledName == kOCLBuiltinName::GetImageDepth ||
DemangledName == kOCLBuiltinName::GetImageDim ||
DemangledName == kOCLBuiltinName::GetImageArraySize) {
visitCallGetImageSize(&CI, MangledName, DemangledName);
return;
}
if ((DemangledName.find(kOCLBuiltinName::WorkGroupPrefix) == 0 &&
DemangledName != kOCLBuiltinName::WorkGroupBarrier) ||
DemangledName == kOCLBuiltinName::WaitGroupEvent ||
(DemangledName.find(kOCLBuiltinName::SubGroupPrefix) == 0 &&
DemangledName != kOCLBuiltinName::SubGroupBarrier)) {
visitCallGroupBuiltin(&CI, MangledName, DemangledName);
return;
}
if (DemangledName.find(kOCLBuiltinName::Pipe) != std::string::npos) {
visitCallPipeBuiltin(&CI, MangledName, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::MemFence) {
visitCallMemFence(&CI);
return;
}
if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0) {
if (MangledName.find(kMangledName::Sampler) != StringRef::npos) {
visitCallReadImageWithSampler(&CI, MangledName, DemangledName);
return;
}
if (MangledName.find("msaa") != StringRef::npos) {
visitCallReadImageMSAA(&CI, MangledName, DemangledName);
return;
}
}
if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0 ||
DemangledName.find(kOCLBuiltinName::WriteImage) == 0) {
visitCallReadWriteImage(&CI, MangledName, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::ToGlobal ||
DemangledName == kOCLBuiltinName::ToLocal ||
DemangledName == kOCLBuiltinName::ToPrivate) {
visitCallToAddr(&CI, MangledName, DemangledName);
return;
}
if (DemangledName.find(kOCLBuiltinName::VLoadPrefix) == 0 ||
DemangledName.find(kOCLBuiltinName::VStorePrefix) == 0) {
visitCallVecLoadStore(&CI, MangledName, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::IsFinite ||
DemangledName == kOCLBuiltinName::IsInf ||
DemangledName == kOCLBuiltinName::IsNan ||
DemangledName == kOCLBuiltinName::IsNormal ||
DemangledName == kOCLBuiltinName::Signbit) {
visitCallRelational(&CI, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::WorkGroupBarrier ||
DemangledName == kOCLBuiltinName::Barrier) {
visitCallBarrier(&CI);
return;
}
if (DemangledName == kOCLBuiltinName::GetFence) {
visitCallGetFence(&CI, MangledName, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::Dot &&
!(CI.getOperand(0)->getType()->isVectorTy())) {
visitCallDot(&CI);
return;
}
if (DemangledName == kOCLBuiltinName::FMin ||
DemangledName == kOCLBuiltinName::FMax ||
DemangledName == kOCLBuiltinName::Min ||
DemangledName == kOCLBuiltinName::Max ||
DemangledName == kOCLBuiltinName::Step ||
DemangledName == kOCLBuiltinName::SmoothStep ||
DemangledName == kOCLBuiltinName::Clamp ||
DemangledName == kOCLBuiltinName::Mix) {
visitCallScalToVec(&CI, MangledName, DemangledName);
return;
}
if (DemangledName == kOCLBuiltinName::GetImageChannelDataType) {
visitCallGetImageChannel(&CI, MangledName, DemangledName,
OCLImageChannelDataTypeOffset);
return;
}
if (DemangledName == kOCLBuiltinName::GetImageChannelOrder) {
visitCallGetImageChannel(&CI, MangledName, DemangledName,
OCLImageChannelOrderOffset);
return;
}
visitCallBuiltinSimple(&CI, MangledName, DemangledName);
}
void
OCL20ToSPIRV::visitCallNDRange(CallInst *CI,
const std::string &DemangledName) {
assert(DemangledName.find(kOCLBuiltinName::NDRangePrefix) == 0);
std::string lenStr = DemangledName.substr(8, 1);
auto Len = atoi(lenStr.c_str());
assert (Len >= 1 && Len <= 3);
// SPIR-V ndrange structure requires 3 members in the following order:
// global work offset
// global work size
// local work size
// The arguments need to add missing members.
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
for (size_t I = 1, E = Args.size(); I != E; ++I)
Args[I] = getScalarOrArray(Args[I], Len, CI);
switch (Args.size()) {
case 2: {
// Has global work size.
auto T = Args[1]->getType();
auto C = getScalarOrArrayConstantInt(CI, T, Len, 0);
Args.push_back(C);
Args.push_back(C);
}
break;
case 3: {
// Has global and local work size.
auto T = Args[1]->getType();
Args.push_back(getScalarOrArrayConstantInt(CI, T, Len, 0));
}
break;
case 4: {
// Move offset arg to the end
auto OffsetPos = Args.begin() + 1;
Value* OffsetVal = *OffsetPos;
Args.erase(OffsetPos);
Args.push_back(OffsetVal);
}
break;
default:
assert(0 && "Invalid number of arguments");
}
// Translate ndrange_ND into differently named SPIR-V decorated functions because
// they have array arugments of different dimension which mangled the same way.
return getSPIRVFuncName(OpBuildNDRange, "_" + lenStr + "D");
}, &Attrs);
}
void
OCL20ToSPIRV::visitCallAsyncWorkGroupCopy(CallInst* CI,
const std::string &DemangledName) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
if (DemangledName == OCLUtil::kOCLBuiltinName::AsyncWorkGroupCopy) {
Args.insert(Args.begin()+3, addSizet(1));
}
Args.insert(Args.begin(), addInt32(ScopeWorkgroup));
return getSPIRVFuncName(OpGroupAsyncCopy);
}, &Attrs);
}
CallInst *
OCL20ToSPIRV::visitCallAtomicCmpXchg(CallInst* CI,
const std::string& DemangledName) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
Value *Expected = nullptr;
CallInst *NewCI = nullptr;
mutateCallInstOCL(M, CI, [&](CallInst * CI, std::vector<Value *> &Args,
Type *&RetTy){
Expected = Args[1]; // temporary save second argument.
Args[1] = new LoadInst(Args[1], "exp", false, CI);
RetTy = Args[2]->getType();
assert(Args[0]->getType()->getPointerElementType()->isIntegerTy() &&
Args[1]->getType()->isIntegerTy() && Args[2]->getType()->isIntegerTy() &&
"In SPIR-V 1.0 arguments of OpAtomicCompareExchange must be "
"an integer type scalars");
return kOCLBuiltinName::AtomicCmpXchgStrong;
},
[&](CallInst *NCI)->Instruction * {
NewCI = NCI;
Instruction* Store = new StoreInst(NCI, Expected, NCI->getNextNode());
return new ICmpInst(Store->getNextNode(), CmpInst::ICMP_EQ, NCI,
NCI->getArgOperand(1));
},
&Attrs);
return NewCI;
}
void
OCL20ToSPIRV::visitCallAtomicInit(CallInst* CI) {
auto ST = new StoreInst(CI->getArgOperand(1), CI->getArgOperand(0), CI);
ST->takeName(CI);
CI->dropAllReferences();
CI->eraseFromParent();
}
void
OCL20ToSPIRV::visitCallAllAny(spv::Op OC, CallInst* CI) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
auto Args = getArguments(CI);
assert(Args.size() == 1);
auto *ArgTy = Args[0]->getType();
auto Zero = Constant::getNullValue(Args[0]->getType());
auto *Cmp = CmpInst::Create(CmpInst::ICmp, CmpInst::ICMP_SLT, Args[0], Zero,
"cast", CI);
if (!isa<VectorType>(ArgTy)) {
auto *Cast = CastInst::CreateZExtOrBitCast(Cmp, Type::getInt32Ty(*Ctx),
"", Cmp->getNextNode());
CI->replaceAllUsesWith(Cast);
CI->eraseFromParent();
} else {
mutateCallInstSPIRV(
M, CI,
[&](CallInst *, std::vector<Value *> &Args, Type *&Ret) {
Args[0] = Cmp;
Ret = Type::getInt1Ty(*Ctx);
return getSPIRVFuncName(OC);
},
[&](CallInst *CI) -> Instruction * {
return CastInst::CreateZExtOrBitCast(CI, Type::getInt32Ty(*Ctx), "",
CI->getNextNode());
},
&Attrs);
}
}
void
OCL20ToSPIRV::visitCallAtomicWorkItemFence(CallInst* CI) {
transMemoryBarrier(CI, getAtomicWorkItemFenceLiterals(CI));
}
void
OCL20ToSPIRV::visitCallMemFence(CallInst* CI) {
transMemoryBarrier(CI, std::make_tuple(
cast<ConstantInt>(CI->getArgOperand(0))->getZExtValue(),
OCLMO_relaxed,
OCLMS_work_group));
}
void OCL20ToSPIRV::transMemoryBarrier(CallInst* CI,
AtomicWorkItemFenceLiterals Lit) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
Args.resize(2);
Args[0] = addInt32(map<Scope>(std::get<2>(Lit)));
Args[1] = addInt32(mapOCLMemSemanticToSPIRV(std::get<0>(Lit),
std::get<1>(Lit)));
return getSPIRVFuncName(OpMemoryBarrier);
}, &Attrs);
}
void
OCL20ToSPIRV::visitCallAtomicLegacy(CallInst* CI,
StringRef MangledName, const std::string& DemangledName) {
StringRef Stem = DemangledName;
if (Stem.startswith("atom_"))
Stem = Stem.drop_front(strlen("atom_"));
else if (Stem.startswith("atomic_"))
Stem = Stem.drop_front(strlen("atomic_"));
else
return;
std::string Sign;
std::string Postfix;
std::string Prefix;
if (Stem == "add" ||
Stem == "sub" ||
Stem == "and" ||
Stem == "or" ||
Stem == "xor" ||
Stem == "min" ||
Stem == "max") {
if ((Stem == "min" || Stem == "max") &&
isMangledTypeUnsigned(MangledName.back()))
Sign = 'u';
Prefix = "fetch_";
Postfix = "_explicit";
} else if (Stem == "xchg") {
Stem = "exchange";
Postfix = "_explicit";
}
else if (Stem == "cmpxchg") {
Stem = "compare_exchange_strong";
Postfix = "_explicit";
}
else if (Stem == "inc" ||
Stem == "dec") {
// do nothing
} else
return;
OCLBuiltinTransInfo Info;
Info.UniqName = "atomic_" + Prefix + Sign + Stem.str() + Postfix;
std::vector<int> PostOps;
PostOps.push_back(OCLLegacyAtomicMemOrder);
if (Stem.startswith("compare_exchange"))
PostOps.push_back(OCLLegacyAtomicMemOrder);
PostOps.push_back(OCLLegacyAtomicMemScope);
Info.PostProc = [=](std::vector<Value *> &Ops){
for (auto &I:PostOps){
Ops.push_back(addInt32(I));
}
};
transAtomicBuiltin(CI, Info);
}
void
OCL20ToSPIRV::visitCallAtomicCpp11(CallInst* CI,
StringRef MangledName, const std::string& DemangledName) {
StringRef Stem = DemangledName;
if (Stem.startswith("atomic_"))
Stem = Stem.drop_front(strlen("atomic_"));
else
return;
std::string NewStem = Stem;
std::vector<int> PostOps;
if (Stem.startswith("store") ||
Stem.startswith("load") ||
Stem.startswith("exchange") ||
Stem.startswith("compare_exchange") ||
Stem.startswith("fetch") ||
Stem.startswith("flag")) {
if ((Stem.startswith("fetch_min") ||
Stem.startswith("fetch_max")) &&
containsUnsignedAtomicType(MangledName))
NewStem.insert(NewStem.begin() + strlen("fetch_"), 'u');
if (!Stem.endswith("_explicit")) {
NewStem = NewStem + "_explicit";
PostOps.push_back(OCLMO_seq_cst);
if (Stem.startswith("compare_exchange"))
PostOps.push_back(OCLMO_seq_cst);
PostOps.push_back(OCLMS_device);
} else {
auto MaxOps = getOCLCpp11AtomicMaxNumOps(
Stem.drop_back(strlen("_explicit")));
if (CI->getNumArgOperands() < MaxOps)
PostOps.push_back(OCLMS_device);
}
} else if (Stem == "work_item_fence") {
// do nothing
} else
return;
OCLBuiltinTransInfo Info;
Info.UniqName = std::string("atomic_") + NewStem;
Info.PostProc = [=](std::vector<Value *> &Ops){
for (auto &I:PostOps){
Ops.push_back(addInt32(I));
}
};
transAtomicBuiltin(CI, Info);
}
void
OCL20ToSPIRV::transAtomicBuiltin(CallInst* CI,
OCLBuiltinTransInfo& Info) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(M, CI, [=](CallInst * CI, std::vector<Value *> &Args){
Info.PostProc(Args);
// Order of args in OCL20:
// object, 0-2 other args, 1-2 order, scope
const size_t NumOrder = getAtomicBuiltinNumMemoryOrderArgs(Info.UniqName);
const size_t ArgsCount = Args.size();
const size_t ScopeIdx = ArgsCount - 1;
const size_t OrderIdx = ScopeIdx - NumOrder;
Args[ScopeIdx] = mapUInt(M, cast<ConstantInt>(Args[ScopeIdx]),
[](unsigned I){
return map<Scope>(static_cast<OCLScopeKind>(I));
});
for (size_t I = 0; I < NumOrder; ++I)
Args[OrderIdx + I] = mapUInt(M, cast<ConstantInt>(Args[OrderIdx + I]),
[](unsigned Ord) {
return mapOCLMemSemanticToSPIRV(0, static_cast<OCLMemOrderKind>(Ord));
});
// Order of args in SPIR-V:
// object, scope, 1-2 order, 0-2 other args
std::swap(Args[1], Args[ScopeIdx]);
if(OrderIdx > 2) {
// For atomic_compare_exchange the swap above puts Comparator/Expected
// argument just where it should be, so don't move the last argument then.
int offset = Info.UniqName.find("atomic_compare_exchange") == 0 ? 1 : 0;
std::rotate(Args.begin() + 2, Args.begin() + OrderIdx,
Args.end() - offset);
}
return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
}, &Attrs);
}
void
OCL20ToSPIRV::visitCallBarrier(CallInst* CI) {
auto Lit = getBarrierLiterals(CI);
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
Args.resize(3);
Args[0] = addInt32(map<Scope>(std::get<2>(Lit)));
Args[1] = addInt32(map<Scope>(std::get<1>(Lit)));
Args[2] = addInt32(mapOCLMemFenceFlagToSPIRV(std::get<0>(Lit)));
return getSPIRVFuncName(OpControlBarrier);
}, &Attrs);
}
void OCL20ToSPIRV::visitCallConvert(CallInst* CI,
StringRef MangledName, const std::string& DemangledName) {
if (eraseUselessConvert(CI, MangledName, DemangledName))
return;
Op OC = OpNop;
auto TargetTy = CI->getType();
auto SrcTy = CI->getArgOperand(0)->getType();
if (isa<VectorType>(TargetTy))
TargetTy = TargetTy->getVectorElementType();
if (isa<VectorType>(SrcTy))
SrcTy = SrcTy->getVectorElementType();
auto IsTargetInt = isa<IntegerType>(TargetTy);
std::string TargetTyName = DemangledName.substr(
strlen(kOCLBuiltinName::ConvertPrefix));
auto FirstUnderscoreLoc = TargetTyName.find('_');
if (FirstUnderscoreLoc != std::string::npos)
TargetTyName = TargetTyName.substr(0, FirstUnderscoreLoc);
TargetTyName = std::string("_R") + TargetTyName;
std::string Sat = DemangledName.find("_sat") != std::string::npos ?
"_sat" : "";
auto TargetSigned = DemangledName[8] != 'u';
if (isa<IntegerType>(SrcTy)) {
bool Signed = isLastFuncParamSigned(MangledName);
if (IsTargetInt) {
if (!Sat.empty() && TargetSigned != Signed) {
OC = Signed ? OpSatConvertSToU : OpSatConvertUToS;
Sat = "";
} else
OC = Signed ? OpSConvert : OpUConvert;
} else
OC = Signed ? OpConvertSToF : OpConvertUToF;
} else {
if (IsTargetInt) {
OC = TargetSigned ? OpConvertFToS : OpConvertFToU;
} else
OC = OpFConvert;
}
auto Loc = DemangledName.find("_rt");
std::string Rounding;
if (Loc != std::string::npos &&
!(isa<IntegerType>(SrcTy) && IsTargetInt)) {
Rounding = DemangledName.substr(Loc, 4);
}
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args){
return getSPIRVFuncName(OC, TargetTyName + Sat + Rounding);
}, &Attrs);
}
void OCL20ToSPIRV::visitCallGroupBuiltin(CallInst* CI,
StringRef MangledName, const std::string& OrigDemangledName) {
auto F = CI->getCalledFunction();
std::vector<int> PreOps;
std::string DemangledName = OrigDemangledName;
if (DemangledName == kOCLBuiltinName::WorkGroupBarrier)
return;
if (DemangledName == kOCLBuiltinName::WaitGroupEvent) {
PreOps.push_back(ScopeWorkgroup);
} else if (DemangledName.find(kOCLBuiltinName::WorkGroupPrefix) == 0) {
DemangledName.erase(0, strlen(kOCLBuiltinName::WorkPrefix));
PreOps.push_back(ScopeWorkgroup);
} else if (DemangledName.find(kOCLBuiltinName::SubGroupPrefix) == 0) {
DemangledName.erase(0, strlen(kOCLBuiltinName::SubPrefix));
PreOps.push_back(ScopeSubgroup);
} else
return;
if (DemangledName != kOCLBuiltinName::WaitGroupEvent) {
StringRef GroupOp = DemangledName;
GroupOp = GroupOp.drop_front(strlen(kSPIRVName::GroupPrefix));
SPIRSPIRVGroupOperationMap::foreach_conditional([&](const std::string &S,
SPIRVGroupOperationKind G){
if (!GroupOp.startswith(S))
return true; // continue
PreOps.push_back(G);
StringRef Op = GroupOp.drop_front(S.size() + 1);
assert(!Op.empty() && "Invalid OpenCL group builtin function");
char OpTyC = 0;
auto NeedSign = Op == "max" || Op == "min";
auto OpTy = F->getReturnType();
if (OpTy->isFloatingPointTy())
OpTyC = 'f';
else if (OpTy->isIntegerTy()) {
if (!NeedSign)
OpTyC = 'i';
else {
if (isLastFuncParamSigned(F->getName()))
OpTyC = 's';
else
OpTyC = 'u';
}
} else
llvm_unreachable("Invalid OpenCL group builtin argument type");
DemangledName = std::string(kSPIRVName::GroupPrefix) + OpTyC + Op.str();
return false; // break out of loop
});
}
bool IsGroupAllAny = (DemangledName.find("_all") != std::string::npos ||
DemangledName.find("_any") != std::string::npos);
auto Consts = getInt32(M, PreOps);
OCLBuiltinTransInfo Info;
if (IsGroupAllAny)
Info.RetTy = Type::getInt1Ty(*Ctx);
Info.UniqName = DemangledName;
Info.PostProc = [=](std::vector<Value *> &Ops) {
if (IsGroupAllAny) {
IRBuilder<> IRB(CI);
Ops[0] =
IRB.CreateICmpNE(Ops[0], ConstantInt::get(Type::getInt32Ty(*Ctx), 0));
}
size_t E = Ops.size();
if (DemangledName == "group_broadcast" && E > 2) {
assert(E == 3 || E == 4);
makeVector(CI, Ops, std::make_pair(Ops.begin() + 1, Ops.end()));
}
Ops.insert(Ops.begin(), Consts.begin(), Consts.end());
};
transBuiltin(CI, Info);
}
void
OCL20ToSPIRV::transBuiltin(CallInst* CI,
OCLBuiltinTransInfo& Info) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
Op OC = OpNop;
unsigned ExtOp = ~0U;
if (StringRef(Info.UniqName).startswith(kSPIRVName::Prefix))
return;
if (OCLSPIRVBuiltinMap::find(Info.UniqName, &OC))
Info.UniqName = getSPIRVFuncName(OC);
else if ((ExtOp = getExtOp(Info.MangledName, Info.UniqName)) != ~0U)
Info.UniqName = getSPIRVExtFuncName(SPIRVEIS_OpenCL, ExtOp);
else
return;
if (!Info.RetTy)
mutateCallInstSPIRV(M, CI,
[=](CallInst *, std::vector<Value *> &Args) {
Info.PostProc(Args);
return Info.UniqName + Info.Postfix;
},
&Attrs);
else
mutateCallInstSPIRV(
M, CI,
[=](CallInst *, std::vector<Value *> &Args, Type *&RetTy) {
Info.PostProc(Args);
RetTy = Info.RetTy;
return Info.UniqName + Info.Postfix;
},
[=](CallInst *NewCI) -> Instruction * {
if (NewCI->getType()->isIntegerTy() && CI->getType()->isIntegerTy())
return CastInst::CreateIntegerCast(NewCI, CI->getType(),
Info.isRetSigned, "", CI);
else
return CastInst::CreatePointerBitCastOrAddrSpaceCast(
NewCI, CI->getType(), "", CI);
},
&Attrs);
}
void
OCL20ToSPIRV::visitCallPipeBuiltin(CallInst* CI,
StringRef MangledName, const std::string& DemangledName) {
std::string NewName = DemangledName;
// Transform OpenCL read_pipe/write_pipe builtin function names
// with reserve_id argument to reserved_read_pipe/reserved_write_pipe.
if ((DemangledName.find(kOCLBuiltinName::ReadPipe) == 0 ||
DemangledName.find(kOCLBuiltinName::WritePipe) == 0)
&& CI->getNumArgOperands() > 4)
NewName = std::string(kSPIRVName::ReservedPrefix) + DemangledName;
OCLBuiltinTransInfo Info;
Info.UniqName = NewName;
transBuiltin(CI, Info);
}
void OCL20ToSPIRV::visitCallReadImageMSAA(CallInst *CI, StringRef MangledName,
const std::string &DemangledName) {
assert(MangledName.find("msaa") != StringRef::npos);
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(
M, CI,
[=](CallInst *, std::vector<Value *> &Args) {
Args.insert(Args.begin() + 2, getInt32(M, ImageOperandsSampleMask));
return getSPIRVFuncName(OpImageRead,
std::string(kSPIRVPostfix::ExtDivider) +
getPostfixForReturnType(CI));
},
&Attrs);
}
void OCL20ToSPIRV::visitCallReadImageWithSampler(
CallInst *CI, StringRef MangledName, const std::string &DemangledName) {
assert (MangledName.find(kMangledName::Sampler) != StringRef::npos);
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
bool isRetScalar = !CI->getType()->isVectorTy();
mutateCallInstSPIRV(
M, CI,
[=](CallInst *, std::vector<Value *> &Args, Type *&Ret) {
auto ImageTy = getAnalysis<OCLTypeToSPIRV>().getAdaptedType(Args[0]);
if (isOCLImageType(ImageTy))
ImageTy = getSPIRVImageTypeFromOCL(M, ImageTy);
auto SampledImgTy = getSPIRVTypeByChangeBaseTypeName(
M, ImageTy, kSPIRVTypeName::Image, kSPIRVTypeName::SampledImg);
Value *SampledImgArgs[] = {Args[0], Args[1]};
auto SampledImg = addCallInstSPIRV(
M, getSPIRVFuncName(OpSampledImage), SampledImgTy, SampledImgArgs,
nullptr, CI, kSPIRVName::TempSampledImage);
Args[0] = SampledImg;
Args.erase(Args.begin() + 1, Args.begin() + 2);
switch (Args.size()) {
case 2: // no lod
Args.push_back(getInt32(M, ImageOperandsMask::ImageOperandsLodMask));
Args.push_back(getFloat32(M, 0.f));
break;
case 3: // explicit lod
Args.insert(Args.begin() + 2,
getInt32(M, ImageOperandsMask::ImageOperandsLodMask));
break;
case 4: // gradient
Args.insert(Args.begin() + 2,
getInt32(M, ImageOperandsMask::ImageOperandsGradMask));
break;
default:
assert(0 && "read_image* with unhandled number of args!");
}
// SPIR-V intruction always returns 4-element vector
if (isRetScalar)
Ret = VectorType::get(Ret, 4);
return getSPIRVFuncName(OpImageSampleExplicitLod,
std::string(kSPIRVPostfix::ExtDivider) +
getPostfixForReturnType(Ret));
},
[&](CallInst *CI) -> Instruction * {
if (isRetScalar)
return ExtractElementInst::Create(CI, getSizet(M, 0), "",
CI->getNextNode());
return CI;
},
&Attrs);
}
void
OCL20ToSPIRV::visitCallGetImageSize(CallInst* CI,
StringRef MangledName, const std::string& DemangledName) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
StringRef TyName;
SmallVector<StringRef, 4> SubStrs;
auto IsImg = isOCLImageType(CI->getArgOperand(0)->getType(), &TyName);
(void)IsImg; // prevent warning about unused variable in NDEBUG build
assert(IsImg);
std::string ImageTyName = TyName.str();
if (hasAccessQualifiedName(TyName))
ImageTyName.erase(ImageTyName.size() - 5, 3);
auto Desc = map<SPIRVTypeImageDescriptor>(ImageTyName);
unsigned Dim = getImageDimension(Desc.Dim) + Desc.Arrayed;
assert(Dim > 0 && "Invalid image dimension.");
mutateCallInstSPIRV(M, CI,
[&](CallInst *, std::vector<Value *> &Args, Type *&Ret){
assert(Args.size() == 1);
Ret = CI->getType()->isIntegerTy(64) ? Type::getInt64Ty(*Ctx)
: Type::getInt32Ty(*Ctx);
if (Dim > 1)
Ret = VectorType::get(Ret, Dim);
if (Desc.Dim == DimBuffer)
return getSPIRVFuncName(OpImageQuerySize, CI->getType());
else {
Args.push_back(getInt32(M, 0));
return getSPIRVFuncName(OpImageQuerySizeLod, CI->getType());
}
},
[&](CallInst *NCI)->Instruction * {
if (Dim == 1)
return NCI;
if (DemangledName == kOCLBuiltinName::GetImageDim) {
if (Desc.Dim == Dim3D) {
auto ZeroVec = ConstantVector::getSplat(3,
Constant::getNullValue(NCI->getType()->getVectorElementType()));
Constant *Index[] = {getInt32(M, 0), getInt32(M, 1),
getInt32(M, 2), getInt32(M, 3)};
return new ShuffleVectorInst(NCI, ZeroVec,
ConstantVector::get(Index), "", CI);
} else if (Desc.Dim == Dim2D && Desc.Arrayed) {
Constant *Index[] = {getInt32(M, 0), getInt32(M, 1)};
Constant *mask = ConstantVector::get(Index);
return new ShuffleVectorInst(NCI, UndefValue::get(NCI->getType()),
mask, NCI->getName(), CI);
}
return NCI;
}
unsigned I = StringSwitch<unsigned>(DemangledName)
.Case(kOCLBuiltinName::GetImageWidth, 0)
.Case(kOCLBuiltinName::GetImageHeight, 1)
.Case(kOCLBuiltinName::GetImageDepth, 2)
.Case(kOCLBuiltinName::GetImageArraySize, Dim - 1);
return ExtractElementInst::Create(NCI, getUInt32(M, I), "",
NCI->getNextNode());
},
&Attrs);
}
/// Remove trivial conversion functions
bool
OCL20ToSPIRV::eraseUselessConvert(CallInst *CI,
const std::string &MangledName,
const std::string &DemangledName) {
auto TargetTy = CI->getType();
auto SrcTy = CI->getArgOperand(0)->getType();
if (isa<VectorType>(TargetTy))
TargetTy = TargetTy->getVectorElementType();
if (isa<VectorType>(SrcTy))
SrcTy = SrcTy->getVectorElementType();
if (TargetTy == SrcTy) {
if (isa<IntegerType>(TargetTy) &&
DemangledName.find("_sat") != std::string::npos &&
isLastFuncParamSigned(MangledName) != (DemangledName[8] != 'u'))
return false;
CI->getArgOperand(0)->takeName(CI);
SPIRVDBG(dbgs() << "[regularizeOCLConvert] " << *CI << " <- " <<
*CI->getArgOperand(0) << '\n');
CI->replaceAllUsesWith(CI->getArgOperand(0));
ValuesToDelete.insert(CI);
ValuesToDelete.insert(CI->getCalledFunction());
return true;
}
return false;
}
void
OCL20ToSPIRV::visitCallBuiltinSimple(CallInst* CI,
StringRef MangledName, const std::string& DemangledName) {
OCLBuiltinTransInfo Info;
Info.MangledName = MangledName.str();
Info.UniqName = DemangledName;
transBuiltin(CI, Info);
}
/// Translates OCL work-item builtin functions to SPIRV builtin variables.
/// Function like get_global_id(i) -> x = load GlobalInvocationId; extract x, i
/// Function like get_work_dim() -> load WorkDim
void OCL20ToSPIRV::transWorkItemBuiltinsToVariables() {
DEBUG(dbgs() << "Enter transWorkItemBuiltinsToVariables\n");
std::vector<Function *> WorkList;
for (auto I = M->begin(), E = M->end(); I != E; ++I) {
std::string DemangledName;
if (!oclIsBuiltin(I->getName(), &DemangledName))
continue;
DEBUG(dbgs() << "Function demangled name: " << DemangledName << '\n');
std::string BuiltinVarName;
SPIRVBuiltinVariableKind BVKind;
if (!SPIRSPIRVBuiltinVariableMap::find(DemangledName, &BVKind))
continue;
BuiltinVarName = std::string(kSPIRVName::Prefix) +
SPIRVBuiltInNameMap::map(BVKind);
DEBUG(dbgs() << "builtin variable name: " << BuiltinVarName << '\n');
bool IsVec = I->getFunctionType()->getNumParams() > 0;
Type *GVType = IsVec ? VectorType::get(I->getReturnType(),3) :
I->getReturnType();
auto BV = new GlobalVariable(*M, GVType,
true,
GlobalValue::ExternalLinkage,
nullptr, BuiltinVarName,
0,
GlobalVariable::NotThreadLocal,
SPIRAS_Constant);
std::vector<Instruction *> InstList;
for (auto UI = I->user_begin(), UE = I->user_end(); UI != UE; ++UI) {
auto CI = dyn_cast<CallInst>(*UI);
assert(CI && "invalid instruction");
Value * NewValue = new LoadInst(BV, "", CI);
DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n');
if (IsVec) {
NewValue = ExtractElementInst::Create(NewValue,
CI->getArgOperand(0),
"", CI);
DEBUG(dbgs() << *NewValue << '\n');
}
NewValue->takeName(CI);
CI->replaceAllUsesWith(NewValue);
InstList.push_back(CI);
}
for (auto &Inst:InstList) {
Inst->dropAllReferences();
Inst->removeFromParent();
}
WorkList.push_back(static_cast<Function*>(I));
}
for (auto &I:WorkList) {
I->dropAllReferences();
I->removeFromParent();
}
}
void
OCL20ToSPIRV::visitCallReadWriteImage(CallInst* CI,
StringRef MangledName, const std::string& DemangledName) {
OCLBuiltinTransInfo Info;
if (DemangledName.find(kOCLBuiltinName::ReadImage) == 0)
Info.UniqName = kOCLBuiltinName::ReadImage;
if (DemangledName.find(kOCLBuiltinName::WriteImage) == 0)
{
Info.UniqName = kOCLBuiltinName::WriteImage;
Info.PostProc = [&](std::vector<Value*> &Args) {
if (Args.size() == 4) // write with lod
{
auto Lod = Args[2];
Args.erase(Args.begin() + 2);
Args.push_back(getInt32(M, ImageOperandsMask::ImageOperandsLodMask));
Args.push_back(Lod);
}
};
}
transBuiltin(CI, Info);
}
void
OCL20ToSPIRV::visitCallToAddr(CallInst* CI, StringRef MangledName,
const std::string &DemangledName) {
auto AddrSpace = static_cast<SPIRAddressSpace>(
CI->getType()->getPointerAddressSpace());
OCLBuiltinTransInfo Info;
Info.UniqName = DemangledName;
Info.Postfix = std::string(kSPIRVPostfix::Divider) + "To" +
SPIRAddrSpaceCapitalizedNameMap::map(AddrSpace);
auto StorageClass = addInt32(SPIRSPIRVAddrSpaceMap::map(AddrSpace));
Info.RetTy = getInt8PtrTy(cast<PointerType>(CI->getType()));
Info.PostProc = [=](std::vector<Value *> &Ops){
auto P = Ops.back();
Ops.pop_back();
Ops.push_back(castToInt8Ptr(P, CI));
Ops.push_back(StorageClass);
};
transBuiltin(CI, Info);
}
void OCL20ToSPIRV::visitCallRelational(CallInst *CI,
const std::string &DemangledName) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
Op OC = OpNop;
OCLSPIRVBuiltinMap::find(DemangledName, &OC);
std::string SPIRVName = getSPIRVFuncName(OC);
mutateCallInstSPIRV(
M, CI,
[=](CallInst *, std::vector<Value *> &Args, Type *&Ret) {
Ret = Type::getInt1Ty(*Ctx);
if (CI->getOperand(0)->getType()->isVectorTy())
Ret = VectorType::get(
Type::getInt1Ty(*Ctx),
CI->getOperand(0)->getType()->getVectorNumElements());
return SPIRVName;
},
[=](CallInst *NewCI) -> Instruction * {
Value *False = nullptr, *True = nullptr;
if (NewCI->getType()->isVectorTy()) {
Type *IntTy = Type::getInt32Ty(*Ctx);
if (cast<VectorType>(NewCI->getOperand(0)->getType())
->getElementType()
->isDoubleTy())
IntTy = Type::getInt64Ty(*Ctx);
if (cast<VectorType>(NewCI->getOperand(0)->getType())
->getElementType()
->isHalfTy())
IntTy = Type::getInt16Ty(*Ctx);
Type *VTy = VectorType::get(IntTy,
NewCI->getType()->getVectorNumElements());
False = Constant::getNullValue(VTy);
True = Constant::getAllOnesValue(VTy);
} else {
False = getInt32(M, 0);
True = getInt32(M, 1);
}
return SelectInst::Create(NewCI, True, False, "", NewCI->getNextNode());
},
&Attrs);
}
void
OCL20ToSPIRV::visitCallVecLoadStore(CallInst* CI,
StringRef MangledName, const std::string& OrigDemangledName) {
std::vector<int> PreOps;
std::string DemangledName = OrigDemangledName;
if (DemangledName.find(kOCLBuiltinName::VLoadPrefix) == 0 &&
DemangledName != kOCLBuiltinName::VLoadHalf) {
SPIRVWord Width = getVecLoadWidth(DemangledName);
SPIRVDBG(spvdbgs() << "[visitCallVecLoadStore] DemangledName: " <<
DemangledName << " Width: " << Width << '\n');
PreOps.push_back(Width);
} else if (DemangledName.find(kOCLBuiltinName::RoundingPrefix)
!= std::string::npos) {
auto R = SPIRSPIRVFPRoundingModeMap::map(DemangledName.substr(
DemangledName.find(kOCLBuiltinName::RoundingPrefix) + 1, 3));
PreOps.push_back(R);
}
if (DemangledName.find(kOCLBuiltinName::VLoadAPrefix) == 0)
transVecLoadStoreName(DemangledName, kOCLBuiltinName::VLoadAPrefix, true);
else
transVecLoadStoreName(DemangledName, kOCLBuiltinName::VLoadPrefix, false);
if (DemangledName.find(kOCLBuiltinName::VStoreAPrefix) == 0)
transVecLoadStoreName(DemangledName, kOCLBuiltinName::VStoreAPrefix, true);
else
transVecLoadStoreName(DemangledName, kOCLBuiltinName::VStorePrefix, false);
auto Consts = getInt32(M, PreOps);
OCLBuiltinTransInfo Info;
Info.MangledName = MangledName;
Info.UniqName = DemangledName;
if (DemangledName.find(kOCLBuiltinName::VLoadPrefix) == 0)
Info.Postfix = std::string(kSPIRVPostfix::ExtDivider) +
getPostfixForReturnType(CI);
Info.PostProc = [=](std::vector<Value *> &Ops){
Ops.insert(Ops.end(), Consts.begin(), Consts.end());
};
transBuiltin(CI, Info);
}
void OCL20ToSPIRV::visitCallGetFence(CallInst *CI, StringRef MangledName,
const std::string &DemangledName) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
Op OC = OpNop;
OCLSPIRVBuiltinMap::find(DemangledName, &OC);
std::string SPIRVName = getSPIRVFuncName(OC);
mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args,
Type *&Ret) { return SPIRVName; },
[=](CallInst *NewCI) -> Instruction * {
return BinaryOperator::CreateLShr(NewCI, getInt32(M, 8), "", CI);
},
&Attrs);
}
void OCL20ToSPIRV::visitCallDot(CallInst *CI) {
IRBuilder<> Builder(CI);
Value *FMulVal = Builder.CreateFMul(CI->getOperand(0), CI->getOperand(1));
CI->replaceAllUsesWith(FMulVal);
CI->dropAllReferences();
CI->removeFromParent();
}
void OCL20ToSPIRV::visitCallScalToVec(CallInst *CI, StringRef MangledName,
const std::string &DemangledName) {
// Check if all arguments have the same type - it's simple case.
auto Uniform = true;
auto IsArg0Vector = isa<VectorType>(CI->getOperand(0)->getType());
for (unsigned I = 1, E = CI->getNumArgOperands(); Uniform && (I != E); ++I) {
Uniform = isa<VectorType>(CI->getOperand(I)->getType()) == IsArg0Vector;
}
if (Uniform) {
visitCallBuiltinSimple(CI, MangledName, DemangledName);
return;
}
std::vector<unsigned int> VecPos;
std::vector<unsigned int> ScalarPos;
if (DemangledName == kOCLBuiltinName::FMin ||
DemangledName == kOCLBuiltinName::FMax ||
DemangledName == kOCLBuiltinName::Min ||
DemangledName == kOCLBuiltinName::Max) {
VecPos.push_back(0);
ScalarPos.push_back(1);
} else if (DemangledName == kOCLBuiltinName::Clamp) {
VecPos.push_back(0);
ScalarPos.push_back(1);
ScalarPos.push_back(2);
} else if (DemangledName == kOCLBuiltinName::Mix) {
VecPos.push_back(0);
VecPos.push_back(1);
ScalarPos.push_back(2);
} else if (DemangledName == kOCLBuiltinName::Step) {
VecPos.push_back(1);
ScalarPos.push_back(0);
} else if (DemangledName == kOCLBuiltinName::SmoothStep) {
VecPos.push_back(2);
ScalarPos.push_back(0);
ScalarPos.push_back(1);
}
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
mutateCallInstSPIRV(
M, CI,
[=](CallInst *, std::vector<Value *> &Args) {
Args.resize(VecPos.size() + ScalarPos.size());
for (auto I : VecPos) {
Args[I] = CI->getOperand(I);
}
auto VecArgWidth =
CI->getOperand(VecPos[0])->getType()->getVectorNumElements();
for (auto I : ScalarPos) {
Instruction *Inst = InsertElementInst::Create(
UndefValue::get(CI->getOperand(VecPos[0])->getType()),
CI->getOperand(I), getInt32(M, 0), "", CI);
Value *NewVec = new ShuffleVectorInst(
Inst, UndefValue::get(CI->getOperand(VecPos[0])->getType()),
ConstantVector::getSplat(VecArgWidth, getInt32(M, 0)), "", CI);
Args[I] = NewVec;
}
return getSPIRVExtFuncName(SPIRVEIS_OpenCL,
getExtOp(MangledName, DemangledName));
},
&Attrs);
}
void OCL20ToSPIRV::visitCallGetImageChannel(CallInst *CI, StringRef MangledName,
const std::string &DemangledName,
unsigned int Offset) {
AttributeSet Attrs = CI->getCalledFunction()->getAttributes();
Op OC = OpNop;
OCLSPIRVBuiltinMap::find(DemangledName, &OC);
std::string SPIRVName = getSPIRVFuncName(OC);
mutateCallInstSPIRV(M, CI, [=](CallInst *, std::vector<Value *> &Args,
Type *&Ret) { return SPIRVName; },
[=](CallInst *NewCI) -> Instruction * {
return BinaryOperator::CreateAdd(
NewCI, getInt32(M, Offset), "", CI);
},
&Attrs);
}
}
INITIALIZE_PASS_BEGIN(OCL20ToSPIRV, "cl20tospv", "Transform OCL 2.0 to SPIR-V",
false, false)
INITIALIZE_PASS_DEPENDENCY(OCLTypeToSPIRV)
INITIALIZE_PASS_END(OCL20ToSPIRV, "cl20tospv", "Transform OCL 2.0 to SPIR-V",
false, false)
ModulePass *llvm::createOCL20ToSPIRV() {
return new OCL20ToSPIRV();
}