{
struct big p;
p.a = x + y;
- p.b = x - y;
+ p.b = x - y + 10;
return p;
}
define ptx_kernel void @add(%struct.big* noalias nocapture sret %agg.result, i32 %x, i32 %y) nounwind noinline {
entry:
%add = add i32 %y, %x
- %sub = sub i32 %x, %y
+ %sub = add i32 %x, 10
+ %add1 = sub i32 %sub, %y
%agg.result.0 = getelementptr inbounds %struct.big* %agg.result, i32 0, i32 0
store i32 %add, i32* %agg.result.0, align 4
%agg.result.1 = getelementptr inbounds %struct.big* %agg.result, i32 0, i32 1
- store i32 %sub, i32* %agg.result.1, align 4
+ store i32 %add1, i32* %agg.result.1, align 4
ret void
}
--- /dev/null
+/*
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+/* Provide Declarations */
+#include <stdarg.h>
+#include <setjmp.h>
+#include <limits.h>
+/* get a declaration for alloca */
+#if defined(__CYGWIN__) || defined(__MINGW32__)
+#define alloca(x) __builtin_alloca((x))
+#define _alloca(x) __builtin_alloca((x))
+#elif defined(__APPLE__)
+extern void *__builtin_alloca(unsigned long);
+#define alloca(x) __builtin_alloca(x)
+#define longjmp _longjmp
+#define setjmp _setjmp
+#elif defined(__sun__)
+#if defined(__sparcv9)
+extern void *__builtin_alloca(unsigned long);
+#else
+extern void *__builtin_alloca(unsigned int);
+#endif
+#define alloca(x) __builtin_alloca(x)
+#elif defined(__FreeBSD__) || defined(__NetBSD__) || defined(__OpenBSD__) || defined(__DragonFly__) || defined(__arm__)
+#define alloca(x) __builtin_alloca(x)
+#elif defined(_MSC_VER)
+#define inline _inline
+#define alloca(x) _alloca(x)
+#else
+#include <alloca.h>
+#endif
+
+#ifndef __GNUC__ /* Can only support "linkonce" vars with GCC */
+#define __attribute__(X)
+#endif
+
+#if defined(__GNUC__) && defined(__APPLE_CC__)
+#define __EXTERNAL_WEAK__ __attribute__((weak_import))
+#elif defined(__GNUC__)
+#define __EXTERNAL_WEAK__ __attribute__((weak))
+#else
+#define __EXTERNAL_WEAK__
+#endif
+
+#if defined(__GNUC__) && defined(__APPLE_CC__)
+#define __ATTRIBUTE_WEAK__
+#elif defined(__GNUC__)
+#define __ATTRIBUTE_WEAK__ __attribute__((weak))
+#else
+#define __ATTRIBUTE_WEAK__
+#endif
+
+#if defined(__GNUC__)
+#define __HIDDEN__ __attribute__((visibility("hidden")))
+#endif
+
+#ifdef __GNUC__
+#define LLVM_NAN(NanStr) __builtin_nan(NanStr) /* Double */
+#define LLVM_NANF(NanStr) __builtin_nanf(NanStr) /* Float */
+#define LLVM_NANS(NanStr) __builtin_nans(NanStr) /* Double */
+#define LLVM_NANSF(NanStr) __builtin_nansf(NanStr) /* Float */
+#define LLVM_INF __builtin_inf() /* Double */
+#define LLVM_INFF __builtin_inff() /* Float */
+#define LLVM_PREFETCH(addr,rw,locality) __builtin_prefetch(addr,rw,locality)
+#define __ATTRIBUTE_CTOR__ __attribute__((constructor))
+#define __ATTRIBUTE_DTOR__ __attribute__((destructor))
+#define LLVM_ASM __asm__
+#else
+#define LLVM_NAN(NanStr) ((double)0.0) /* Double */
+#define LLVM_NANF(NanStr) 0.0F /* Float */
+#define LLVM_NANS(NanStr) ((double)0.0) /* Double */
+#define LLVM_NANSF(NanStr) 0.0F /* Float */
+#define LLVM_INF ((double)0.0) /* Double */
+#define LLVM_INFF 0.0F /* Float */
+#define LLVM_PREFETCH(addr,rw,locality) /* PREFETCH */
+#define __ATTRIBUTE_CTOR__
+#define __ATTRIBUTE_DTOR__
+#define LLVM_ASM(X)
+#endif
+
+#if __GNUC__ < 4 /* Old GCC's, or compilers not GCC */
+#define __builtin_stack_save() 0 /* not implemented */
+#define __builtin_stack_restore(X) /* noop */
+#endif
+
+#if __GNUC__ && __LP64__ /* 128-bit integer types */
+typedef int __attribute__((mode(TI))) llvmInt128;
+typedef unsigned __attribute__((mode(TI))) llvmUInt128;
+#endif
+
+#define CODE_FOR_MAIN() /* Any target-specific code for main()*/
+
+#ifndef __cplusplus
+typedef unsigned char bool;
+#endif
+
+
+/* Support for floating point constants */
+typedef unsigned long long ConstantDoubleTy;
+typedef unsigned int ConstantFloatTy;
+typedef struct { unsigned long long f1; unsigned short f2; unsigned short pad[3]; } ConstantFP80Ty;
+typedef struct { unsigned long long f1; unsigned long long f2; } ConstantFP128Ty;
+
+
+/* Global Declarations */
+/* Helper union for bitcasts */
+typedef union {
+ unsigned int Int32;
+ unsigned long long Int64;
+ float Float;
+ double Double;
+} llvmBitCastUnion;
+
+/* Function Declarations */
+double fmod(double, double);
+float fmodf(float, float);
+long double fmodl(long double, long double);
+void test_global_id(unsigned int *llvm_cbe_dst);
+unsigned int __gen_get_global_id0(void);
+void abort(void);
+
+
+/* Function Bodies */
+static inline int llvm_fcmp_ord(double X, double Y) { return X == X && Y == Y; }
+static inline int llvm_fcmp_uno(double X, double Y) { return X != X || Y != Y; }
+static inline int llvm_fcmp_ueq(double X, double Y) { return X == Y || llvm_fcmp_uno(X, Y); }
+static inline int llvm_fcmp_une(double X, double Y) { return X != Y; }
+static inline int llvm_fcmp_ult(double X, double Y) { return X < Y || llvm_fcmp_uno(X, Y); }
+static inline int llvm_fcmp_ugt(double X, double Y) { return X > Y || llvm_fcmp_uno(X, Y); }
+static inline int llvm_fcmp_ule(double X, double Y) { return X <= Y || llvm_fcmp_uno(X, Y); }
+static inline int llvm_fcmp_uge(double X, double Y) { return X >= Y || llvm_fcmp_uno(X, Y); }
+static inline int llvm_fcmp_oeq(double X, double Y) { return X == Y ; }
+static inline int llvm_fcmp_one(double X, double Y) { return X != Y && llvm_fcmp_ord(X, Y); }
+static inline int llvm_fcmp_olt(double X, double Y) { return X < Y ; }
+static inline int llvm_fcmp_ogt(double X, double Y) { return X > Y ; }
+static inline int llvm_fcmp_ole(double X, double Y) { return X <= Y ; }
+static inline int llvm_fcmp_oge(double X, double Y) { return X >= Y ; }
+
+void test_global_id(unsigned int *llvm_cbe_dst) {
+ unsigned int llvm_cbe_call_2e_i;
+
+ llvm_cbe_call_2e_i = /*tail*/ __gen_get_global_id0();
+ *((&llvm_cbe_dst[((signed int )llvm_cbe_call_2e_i)])) = 1u;
+ return;
+}
+
--- /dev/null
+__attribute__((pure)) unsigned int __gen_get_global_id0(void);
+__attribute__((pure)) unsigned int __gen_get_global_id1(void);
+__attribute__((pure)) unsigned int __gen_get_global_id2(void);
+
+inline unsigned get_global_id(unsigned int dim) {
+ if (dim == 0) return __gen_get_global_id0();
+ else if (dim == 1) return __gen_get_global_id1();
+ else if (dim == 2) return __gen_get_global_id2();
+ else return 0;
+}
+
+__kernel void test_global_id(__global int *dst)
+{
+ short hop = get_global_id(0);
+ dst[get_global_id(0)] = hop;
+}
+
--- /dev/null
+; ModuleID = 'get_global_id.o'
+target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
+target triple = "ptx32--"
+
+define ptx_kernel void @test_global_id(i32* nocapture %dst) nounwind noinline {
+get_global_id.exit5:
+ %call.i = tail call ptx_device i32 @__gen_get_global_id0() nounwind readonly
+ %sext = shl i32 %call.i, 16
+ %conv1 = ashr exact i32 %sext, 16
+ %arrayidx = getelementptr inbounds i32* %dst, i32 %call.i
+ store i32 %conv1, i32* %arrayidx, align 4, !tbaa !1
+ ret void
+}
+
+declare ptx_device i32 @__gen_get_global_id0() nounwind readonly
+
+!opencl.kernels = !{!0}
+
+!0 = metadata !{void (i32*)* @test_global_id}
+!1 = metadata !{metadata !"int", metadata !2}
+!2 = metadata !{metadata !"omnipotent char", metadata !3}
+!3 = metadata !{metadata !"Simple C/C++ TBAA", null}
sys/platform.hpp
ir/context.cpp
ir/context.hpp
+ ir/profile.cpp
+ ir/profile.hpp
ir/type.cpp
ir/type.hpp
ir/unit.cpp
Register Context::reg(RegisterData::Family family) {
GBE_ASSERTM(fn != NULL, "No function currently defined");
- return fn->file.append(family);
+ return fn->newRegister(family);
}
LabelIndex Context::label(void) {
#include "ir/instruction.hpp"
#include "ir/function.hpp"
#include "ir/register.hpp"
+#include "ir/value.hpp"
#include "ir/unit.hpp"
#include "sys/vector.hpp"
#include <tuple>
void endFunction(void);
/*! Create a new register with the given family for the current function */
Register reg(RegisterData::Family family);
+ /*! Create a new register holding the given value. A LOADI is pushed */
+ template <typename T> INLINE Register immReg(T value) {
+ GBE_ASSERTM(fn != NULL, "No function currently defined");
+ const Immediate imm(value);
+ const ImmediateIndex index = fn->newImmediate(imm);
+ const RegisterData::Family family = getFamily(imm.type);
+ const Register reg = this->reg(family);
+ this->LOADI(imm.type, reg, index);
+ return reg;
+ }
/*! Create a new label for the current function */
LabelIndex label(void);
/*! Append a new input register for the function */
/*! Get the current processed function */
Function &getFunction(void);
/*! Append a new tuple */
- template <typename... Args> INLINE Tuple tuple(Args...args);
+ template <typename... Args> INLINE Tuple tuple(Args...args) {
+ GBE_ASSERTM(fn != NULL, "No function currently defined");
+ return fn->file.appendTuple(args...);
+ }
/*! We just use variadic templates to forward instruction functions */
#define DECL_INSN(NAME, FAMILY) \
template <typename... Args> INLINE void NAME(Args...args);
GBE_CLASS(Context);
};
- template <typename... Args>
- INLINE Tuple Context::tuple(Args...args) {
- GBE_ASSERTM(fn != NULL, "No function currently defined");
- return fn->file.appendTuple(args...);
- }
-
// Use argument checker to assert argument value correctness
#define DECL_INSN(NAME, FAMILY) \
template <typename... Args> \
namespace gbe {
namespace ir {
- Function::Function(const std::string &name) :
- name(name), structReturned(false) {}
+ Function::Function(const std::string &name, Profile profile) :
+ name(name), structReturned(false), profile(profile)
+ {
+ initProfile(*this);
+ }
Function::~Function(void) {
for (auto it = blocks.begin(); it != blocks.end(); ++it)
return index;
}
- std::ostream &Function::outImmediate(std::ostream &out, ImmediateIndex index) const {
+ void Function::outImmediate(std::ostream &out, ImmediateIndex index) const {
GBE_ASSERT(index < immediates.size());
const Immediate imm = immediates[index];
switch (imm.type) {
- case TYPE_BOOL: return out << !!imm.data.u8;
- case TYPE_S8: return out << imm.data.s8;
- case TYPE_U8: return out << imm.data.u8;
- case TYPE_S16: return out << imm.data.s16;
- case TYPE_U16: return out << imm.data.u16;
- case TYPE_S32: return out << imm.data.s32;
- case TYPE_U32: return out << imm.data.u32;
- case TYPE_S64: return out << imm.data.s64;
- case TYPE_U64: return out << imm.data.u64;
- case TYPE_HALF: return out << "half(" << imm.data.u16 << ")";
- case TYPE_FLOAT: return out << imm.data.f32;
- case TYPE_DOUBLE: return out << imm.data.f64;
+ case TYPE_BOOL: out << !!imm.data.u8; break;
+ case TYPE_S8: out << imm.data.s8; break;
+ case TYPE_U8: out << imm.data.u8; break;
+ case TYPE_S16: out << imm.data.s16; break;
+ case TYPE_U16: out << imm.data.u16; break;
+ case TYPE_S32: out << imm.data.s32; break;
+ case TYPE_U32: out << imm.data.u32; break;
+ case TYPE_S64: out << imm.data.s64; break;
+ case TYPE_U64: out << imm.data.u64; break;
+ case TYPE_HALF: out << "half(" << imm.data.u16 << ")"; break;
+ case TYPE_FLOAT: out << imm.data.f32; break;
+ case TYPE_DOUBLE: out << imm.data.f64; break;
};
- return out;
}
std::ostream &operator<< (std::ostream &out, const Function &fn)
#include "ir/value.hpp"
#include "ir/register.hpp"
#include "ir/instruction.hpp"
+#include "ir/profile.hpp"
#include "sys/vector.hpp"
#include "sys/list.hpp"
#include "sys/alloc.hpp"
{
public:
/*! Create an empty function */
- Function(const std::string &name);
+ Function(const std::string &name, Profile profile = PROFILE_OCL);
/*! Release everything *including* the basic block pointers */
~Function(void);
+ /*! Get the function profile */
+ INLINE Profile getProfile(void) const { return profile; }
+ /*! Get a new valid register */
+ INLINE Register newRegister(RegisterData::Family family) {
+ return this->file.append(family);
+ }
/*! Get the function name */
const std::string &getName(void) const { return name; }
/*! Extract the register from the register file */
GBE_ASSERT(ID < immediateNum());
return immediates[ID];
}
+ /*! Create a new immediate and returns its index */
+ INLINE ImmediateIndex newImmediate(const Immediate &imm) {
+ const ImmediateIndex index(this->immediateNum());
+ this->immediates.push_back(imm);
+ return index;
+ }
/*! Allocate a new instruction (with the growing pool) */
INLINE Instruction *newInstruction(void) {
return new (insnPool.allocate()) Instruction();
/*! Number of blocks in the function */
INLINE uint32_t blockNum(void) const { return blocks.size(); }
/*! Output an immediate value in a stream */
- std::ostream &outImmediate(std::ostream &out, ImmediateIndex index) const;
+ void outImmediate(std::ostream &out, ImmediateIndex index) const;
private:
friend class Context; //!< Can freely modify a function
std::string name; //!< Function name
RegisterFile file; //!< RegisterDatas used by the instructions
GrowingPool<Instruction> insnPool; //!< For fast instruction allocation
bool structReturned; //!< First argument is pointer to struct
+ Profile profile; //!< Current function profile
GBE_CLASS(Function);
};
INLINE void LoadImmInstruction::out(std::ostream &out, const Function &fn) const {
this->outOpcode(out);
out << "." << type;
- out << " %" << this->getSrcIndex(fn,0);
- out << " " << fn.outImmediate(out, immediateIndex);
+ out << " %" << this->getDstIndex(fn,0) << " ";
+ fn.outImmediate(out, immediateIndex);
}
} /* namespace internal */
/*! A label is identified with an unsigned short */
TYPE_SAFE(LabelIndex, uint16_t)
- /*! A value is stored in a per-function vector. This is the index to it */
- TYPE_SAFE(ImmediateIndex, uint16_t)
-
/*! Function class contains the register file and the register tuple. Any
* information related to the registers may therefore require a function
*/
--- /dev/null
+/*
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+/**
+ * \file profile.hpp
+ * \author Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+#include "ir/profile.hpp"
+#include "ir/function.hpp"
+#include "sys/platform.hpp"
+
+namespace gbe {
+namespace ir {
+
+ namespace ocl
+ {
+ static void init(Function &fn) {
+ IF_DEBUG(Register r);
+ IF_DEBUG(r = ) fn.newRegister(RegisterData::WORD);
+ GBE_ASSERT(r == lid0);
+ IF_DEBUG(r = ) fn.newRegister(RegisterData::WORD);
+ GBE_ASSERT(r == lid1);
+ IF_DEBUG(r = ) fn.newRegister(RegisterData::WORD);
+ GBE_ASSERT(r == lid2);
+ IF_DEBUG(r = ) fn.newRegister(RegisterData::DWORD);
+ GBE_ASSERT(r == gid0);
+ IF_DEBUG(r = ) fn.newRegister(RegisterData::DWORD);
+ GBE_ASSERT(r == gid1);
+ IF_DEBUG(r = ) fn.newRegister(RegisterData::DWORD);
+ GBE_ASSERT(r == gid2);
+ }
+ } /* namespace ocl */
+
+ void initProfile(Function &fn) {
+ const Profile profile = fn.getProfile();
+ switch (profile) {
+ case PROFILE_C: GBE_ASSERTM(false, "Unsupported profile"); break;
+ case PROFILE_OCL: ocl::init(fn);
+ };
+ }
+
+} /* namespace ir */
+} /* namespace gbe */
+
+
--- /dev/null
+/*
+ * Copyright © 2012 Intel Corporation
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ *
+ * Author: Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+
+/**
+ * \file profile.hpp
+ * \author Benjamin Segovia <benjamin.segovia@intel.com>
+ */
+#ifndef __GBE_IR_PROFILE_HPP__
+#define __GBE_IR_PROFILE_HPP__
+
+#include "ir/register.hpp"
+
+namespace gbe {
+namespace ir {
+
+ /*! Profile is defined *per-function* and mostly predefined registers */
+ enum Profile : uint32_t {
+ PROFILE_C = 0, // Not used now
+ PROFILE_OCL = 1
+ };
+
+ // Will be pre-initialized
+ class Function;
+
+ /*! Registers used for ocl */
+ namespace ocl
+ {
+ static const Register lid0 = Register(0); // get_local_id(0)
+ static const Register lid1 = Register(1); // get_local_id(1)
+ static const Register lid2 = Register(2); // get_local_id(2)
+ static const Register gid0 = Register(3); // get_global_id(0)
+ static const Register gid1 = Register(4); // get_global_id(1)
+ static const Register gid2 = Register(5); // get_global_id(2)
+ static const uint32_t regNum = 6; // number of special registers
+ } /* namespace ocl */
+
+ /*! Initialize the profile of the given function */
+ void initProfile(Function &fn);
+
+} /* namespace ir */
+} /* namespace gbe */
+
+#endif /* __GBE_IR_PROFILE_HPP__ */
+
class Immediate
{
public:
-#define DECL_CONSTRUCTOR(TYPE, FIELD) \
- Immediate(TYPE FIELD) { this->data.u64 = 0llu; this->data.FIELD = FIELD; }
- DECL_CONSTRUCTOR(int8_t, s8)
- DECL_CONSTRUCTOR(uint8_t, u8)
+#define DECL_CONSTRUCTOR(TYPE, FIELD, IR_TYPE) \
+ Immediate(TYPE FIELD) { \
+ this->type = IR_TYPE; \
+ this->data.u64 = 0llu; \
+ this->data.FIELD = FIELD; \
+ }
+ DECL_CONSTRUCTOR(bool, b, TYPE_BOOL)
+ DECL_CONSTRUCTOR(int8_t, s8, TYPE_S8)
+ DECL_CONSTRUCTOR(uint8_t, u8, TYPE_U8)
+ DECL_CONSTRUCTOR(int16_t, s16, TYPE_S16)
+ DECL_CONSTRUCTOR(uint16_t, u16, TYPE_S16)
+ DECL_CONSTRUCTOR(int32_t, s32, TYPE_S32)
+ DECL_CONSTRUCTOR(uint32_t, u32, TYPE_S32)
+ DECL_CONSTRUCTOR(int64_t, s64, TYPE_S64)
+ DECL_CONSTRUCTOR(uint64_t, u64, TYPE_S64)
+ DECL_CONSTRUCTOR(float, f32, TYPE_FLOAT)
+ DECL_CONSTRUCTOR(double, f64, TYPE_DOUBLE)
#undef DECL_CONSTRUCTOR
union {
+ bool b;
int8_t s8;
uint8_t u8;
int16_t s16;
Type type; //!< Type of the value
};
+ /*! A value is stored in a per-function vector. This is the index to it */
+ TYPE_SAFE(ImmediateIndex, uint16_t)
+
} /* namespace ir */
} /* namespace gbe */
bool writeInstructionCast(const Instruction &I);
private :
- std::string InterpretASMConstraint(InlineAsm::ConstraintInfo& c);
void lowerIntrinsics(Function &F);
/// Prints the definition of the intrinsic function F. Supports the
INLINE void newRegister(const Value *value);
/*! Return a valid register from an operand (can use LOADI to make one) */
INLINE ir::Register getRegister(Value *value);
+ /*! Return a valid register for a constant value */
+ INLINE ir::Register getConstantRegister(Constant *CPV);
/*! Insert a new label index when this is a scalar value */
INLINE void newLabelIndex(const Value *value);
/*! int / float / double / bool are scalars */
}
}
+ ir::Register GenWriter::getConstantRegister(Constant *CPV) {
+ if (dyn_cast<ConstantExpr>(CPV))
+ GBE_ASSERTM(false, "Unsupported constant expression");
+ else if (isa<UndefValue>(CPV) && CPV->getType()->isSingleValueType())
+ GBE_ASSERTM(false, "Unsupported constant expression");
+ if (ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
+ Type* Ty = CI->getType();
+ if (Ty == Type::getInt1Ty(CPV->getContext())) {
+ const bool b = CI->getZExtValue();
+ return ctx.immReg(b);
+ } else if (Ty == Type::getInt8Ty(CPV->getContext())) {
+ const uint8_t u8 = CI->getZExtValue();
+ return ctx.immReg(u8);
+ } else if (Ty == Type::getInt16Ty(CPV->getContext())) {
+ const uint16_t u16 = CI->getZExtValue();
+ return ctx.immReg(u16);
+ } else if (Ty == Type::getInt32Ty(CPV->getContext())) {
+ const uint32_t u32 = CI->getZExtValue();
+ return ctx.immReg(u32);
+ } else if (Ty == Type::getInt64Ty(CPV->getContext())) {
+ const uint64_t u64 = CI->getZExtValue();
+ return ctx.immReg(u64);
+ } else {
+ GBE_ASSERTM(false, "Unsupported integer size");
+ return ctx.immReg(uint64_t(0));
+ }
+ }
+
+ switch (CPV->getType()->getTypeID()) {
+ case Type::FloatTyID:
+ case Type::DoubleTyID:
+ {
+ ConstantFP *FPC = cast<ConstantFP>(CPV);
+ if (FPC->getType() == Type::getFloatTy(CPV->getContext())) {
+ const float f32 = FPC->getValueAPF().convertToFloat();
+ return ctx.immReg(f32);
+ } else {
+ const double f64 = FPC->getValueAPF().convertToDouble();
+ return ctx.immReg(f64);
+ }
+ }
+ break;
+ default:
+ GBE_ASSERTM(false, "Unsupported constant type");
+ }
+ return ctx.immReg(uint64_t(0));
+ }
+
ir::Register GenWriter::getRegister(Value *value) {
Constant *CPV = dyn_cast<Constant>(value);
- if (CPV && !isa<GlobalValue>(CPV)) {
- GBE_ASSERT(0);
- // printConstant(CPV, Static);
- } else {
+ if (CPV && !isa<GlobalValue>(CPV))
+ return getConstantRegister(CPV);
+ else {
GBE_ASSERT(this->registerMap.find(value) != this->registerMap.end());
return this->registerMap[value];
}
}
}
- void GenWriter::visitCallInst(CallInst &I) {
+ void GenWriter::visitCallInst(CallInst &I)
+ {
+#if 0
if (isa<InlineAsm>(I.getCalledValue()))
return visitInlineAsm(I);
PrintedArg = true;
}
Out << ')';
+#endif
}
/// visitBuiltinCall - Handle the call to the specified builtin. Returns true
/// if the entire call is handled, return false if it wasn't handled, and
/// optionally set 'WroteCallee' if the callee has already been printed out.
bool GenWriter::visitBuiltinCall(CallInst &I, Intrinsic::ID ID, bool &WroteCallee) {
- switch (ID) {
- default: {
- // If this is an intrinsic that directly corresponds to a GCC
- // builtin, we emit it here.
- const char *BuiltinName = "";
- Function *F = I.getCalledFunction();
-#define GET_GCC_BUILTIN_NAME
-#include "llvm/Intrinsics.gen"
-#undef GET_GCC_BUILTIN_NAME
- assert(BuiltinName[0] && "Unknown LLVM intrinsic!");
-
- Out << BuiltinName;
- WroteCallee = true;
- return false;
- }
- case Intrinsic::vastart:
- Out << "0; ";
-
- Out << "va_start(*(va_list*)";
- writeOperand(I.getArgOperand(0));
- Out << ", ";
- // Output the last argument to the enclosing function.
- if (I.getParent()->getParent()->arg_empty())
- Out << "vararg_dummy_arg";
- else
- writeOperand(--I.getParent()->getParent()->arg_end());
- Out << ')';
- return true;
- case Intrinsic::vaend:
- if (!isa<ConstantPointerNull>(I.getArgOperand(0))) {
- Out << "0; va_end(*(va_list*)";
- writeOperand(I.getArgOperand(0));
- Out << ')';
- } else {
- Out << "va_end(*(va_list*)0)";
- }
- return true;
- case Intrinsic::vacopy:
- Out << "0; ";
- Out << "va_copy(*(va_list*)";
- writeOperand(I.getArgOperand(0));
- Out << ", *(va_list*)";
- writeOperand(I.getArgOperand(1));
- Out << ')';
- return true;
- case Intrinsic::returnaddress:
- Out << "__builtin_return_address(";
- writeOperand(I.getArgOperand(0));
- Out << ')';
- return true;
- case Intrinsic::frameaddress:
- Out << "__builtin_frame_address(";
- writeOperand(I.getArgOperand(0));
- Out << ')';
- return true;
- case Intrinsic::powi:
- Out << "__builtin_powi(";
- writeOperand(I.getArgOperand(0));
- Out << ", ";
- writeOperand(I.getArgOperand(1));
- Out << ')';
- return true;
- case Intrinsic::setjmp:
- Out << "setjmp(*(jmp_buf*)";
- writeOperand(I.getArgOperand(0));
- Out << ')';
- return true;
- case Intrinsic::longjmp:
- Out << "longjmp(*(jmp_buf*)";
- writeOperand(I.getArgOperand(0));
- Out << ", ";
- writeOperand(I.getArgOperand(1));
- Out << ')';
- return true;
- case Intrinsic::prefetch:
- Out << "LLVM_PREFETCH((const void *)";
- writeOperand(I.getArgOperand(0));
- Out << ", ";
- writeOperand(I.getArgOperand(1));
- Out << ", ";
- writeOperand(I.getArgOperand(2));
- Out << ")";
- return true;
- case Intrinsic::stacksave:
- // Emit this as: Val = 0; *((void**)&Val) = __builtin_stack_save()
- // to work around GCC bugs (see PR1809).
- Out << "0; *((void**)&" << GetValueName(&I)
- << ") = __builtin_stack_save()";
- return true;
- case Intrinsic::x86_sse_cmp_ss:
- case Intrinsic::x86_sse_cmp_ps:
- case Intrinsic::x86_sse2_cmp_sd:
- case Intrinsic::x86_sse2_cmp_pd:
- Out << '(';
- printType(Out, I.getType());
- Out << ')';
- // Multiple GCC builtins multiplex onto this intrinsic.
- switch (cast<ConstantInt>(I.getArgOperand(2))->getZExtValue()) {
- default: llvm_unreachable("Invalid llvm.x86.sse.cmp!");
- case 0: Out << "__builtin_ia32_cmpeq"; break;
- case 1: Out << "__builtin_ia32_cmplt"; break;
- case 2: Out << "__builtin_ia32_cmple"; break;
- case 3: Out << "__builtin_ia32_cmpunord"; break;
- case 4: Out << "__builtin_ia32_cmpneq"; break;
- case 5: Out << "__builtin_ia32_cmpnlt"; break;
- case 6: Out << "__builtin_ia32_cmpnle"; break;
- case 7: Out << "__builtin_ia32_cmpord"; break;
- }
- if (ID == Intrinsic::x86_sse_cmp_ps || ID == Intrinsic::x86_sse2_cmp_pd)
- Out << 'p';
- else
- Out << 's';
- if (ID == Intrinsic::x86_sse_cmp_ss || ID == Intrinsic::x86_sse_cmp_ps)
- Out << 's';
- else
- Out << 'd';
-
- Out << "(";
- writeOperand(I.getArgOperand(0));
- Out << ", ";
- writeOperand(I.getArgOperand(1));
- Out << ")";
- return true;
- case Intrinsic::ppc_altivec_lvsl:
- Out << '(';
- printType(Out, I.getType());
- Out << ')';
- Out << "__builtin_altivec_lvsl(0, (void*)";
- writeOperand(I.getArgOperand(0));
- Out << ")";
- return true;
- case Intrinsic::uadd_with_overflow:
- case Intrinsic::sadd_with_overflow:
- Out << GetValueName(I.getCalledFunction()) << "(";
- writeOperand(I.getArgOperand(0));
- Out << ", ";
- writeOperand(I.getArgOperand(1));
- Out << ")";
- return true;
- }
- }
-
- //This converts the llvm constraint string to something gcc is expecting.
- //TODO: work out platform independent constraints and factor those out
- // of the per target tables
- // handle multiple constraint codes
- std::string GenWriter::InterpretASMConstraint(InlineAsm::ConstraintInfo& c) {
- assert(c.Codes.size() == 1 && "Too many asm constraint codes to handle");
-
- // Grab the translation table from MCAsmInfo if it exists.
- const MCAsmInfo *TargetAsm;
- std::string Triple = TheModule->getTargetTriple();
- if (Triple.empty())
- Triple = llvm::sys::getHostTriple();
-
- std::string E;
- if (const Target *Match = TargetRegistry::lookupTarget(Triple, E))
- TargetAsm = Match->createMCAsmInfo(Triple);
- else
- return c.Codes[0];
-
- const char *const *table = TargetAsm->getAsmCBE();
-
- // Search the translation table if it exists.
- for (int i = 0; table && table[i]; i += 2)
- if (c.Codes[0] == table[i]) {
- delete TargetAsm;
- return table[i+1];
- }
-
- // Default is identity.
- delete TargetAsm;
- return c.Codes[0];
+ GBE_ASSERTM(false, "builtin call is not supported");
+ return false;
}
void GenWriter::visitAllocaInst(AllocaInst &I) {
/*! Debug syntactic sugar */
#if GBE_DEBUG
-#define IF_DEBUG(EXPR)
-#else
#define IF_DEBUG(EXPR) EXPR
+#else
+#define IF_DEBUG(EXPR)
#endif /* GBE_DEBUG */
/*! Debug printing macros */
fclose(dummyKernel);
//UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add.ll"));
- UTEST_EXPECT_SUCCESS(utestLLVM2Gen("load_store.ll"));
+ //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("load_store.ll"));
//UTEST_EXPECT_SUCCESS(utestLLVM2Gen("add2.ll"));
+ UTEST_EXPECT_SUCCESS(utestLLVM2Gen("get_global_id.ll"));
//UTEST_EXPECT_SUCCESS(utestLLVM2Gen("loop.ll"));
//UTEST_EXPECT_SUCCESS(utestLLVM2Gen("void.ll"));
}
projects / contrib / beignet.git / commitdiff