From: Benjamin Segovia Date: Wed, 29 Feb 2012 20:37:17 +0000 (+0000) Subject: Added first support for immediates Started to add support for builtin functions X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=ca43f8273cf5aa82b8cf77e26ada33846d772325;p=contrib%2Fbeignet.git Added first support for immediates Started to add support for builtin functions --- diff --git a/backend/kernels/add2.cl b/backend/kernels/add2.cl index c43e2c3..8070576 100644 --- a/backend/kernels/add2.cl +++ b/backend/kernels/add2.cl @@ -6,7 +6,7 @@ __kernel struct big add(unsigned int x, unsigned int y) { struct big p; p.a = x + y; - p.b = x - y; + p.b = x - y + 10; return p; } diff --git a/backend/kernels/add2.ll b/backend/kernels/add2.ll index 37cf7a3..4ca1125 100644 --- a/backend/kernels/add2.ll +++ b/backend/kernels/add2.ll @@ -7,11 +7,12 @@ target triple = "ptx32--" 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 } diff --git a/backend/kernels/add2.o b/backend/kernels/add2.o index 1feb035..8b5ebb4 100644 Binary files a/backend/kernels/add2.o and b/backend/kernels/add2.o differ diff --git a/backend/kernels/get_global_id.cbe.c b/backend/kernels/get_global_id.cbe.c new file mode 100644 index 0000000..4dbae41 --- /dev/null +++ b/backend/kernels/get_global_id.cbe.c @@ -0,0 +1,162 @@ +/* + * 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 . + * + * Author: Benjamin Segovia + */ + +/* Provide Declarations */ +#include +#include +#include +/* 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 +#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; +} + diff --git a/backend/kernels/get_global_id.cl b/backend/kernels/get_global_id.cl new file mode 100644 index 0000000..86500ad --- /dev/null +++ b/backend/kernels/get_global_id.cl @@ -0,0 +1,17 @@ +__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; +} + diff --git a/backend/kernels/get_global_id.ll b/backend/kernels/get_global_id.ll new file mode 100644 index 0000000..965739a --- /dev/null +++ b/backend/kernels/get_global_id.ll @@ -0,0 +1,22 @@ +; 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} diff --git a/backend/kernels/get_global_id.o b/backend/kernels/get_global_id.o new file mode 100644 index 0000000..e21b2e1 Binary files /dev/null and b/backend/kernels/get_global_id.o differ diff --git a/backend/src/CMakeLists.txt b/backend/src/CMakeLists.txt index 86a37fc..0a56316 100644 --- a/backend/src/CMakeLists.txt +++ b/backend/src/CMakeLists.txt @@ -25,6 +25,8 @@ else (GBE_USE_BLOB) sys/platform.hpp ir/context.cpp ir/context.hpp + ir/profile.cpp + ir/profile.hpp ir/type.cpp ir/type.hpp ir/unit.cpp diff --git a/backend/src/ir/context.cpp b/backend/src/ir/context.cpp index d4d1ce8..6be17d6 100644 --- a/backend/src/ir/context.cpp +++ b/backend/src/ir/context.cpp @@ -67,7 +67,7 @@ namespace ir { 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) { diff --git a/backend/src/ir/context.hpp b/backend/src/ir/context.hpp index bbdb8c6..f65751f 100644 --- a/backend/src/ir/context.hpp +++ b/backend/src/ir/context.hpp @@ -27,6 +27,7 @@ #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 @@ -53,6 +54,16 @@ namespace ir { 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 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 */ @@ -62,7 +73,10 @@ namespace ir { /*! Get the current processed function */ Function &getFunction(void); /*! Append a new tuple */ - template INLINE Tuple tuple(Args...args); + template 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 INLINE void NAME(Args...args); @@ -129,12 +143,6 @@ namespace ir { GBE_CLASS(Context); }; - template - 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 \ diff --git a/backend/src/ir/function.cpp b/backend/src/ir/function.cpp index eeb47d6..f8a7282 100644 --- a/backend/src/ir/function.cpp +++ b/backend/src/ir/function.cpp @@ -27,8 +27,11 @@ 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) @@ -43,24 +46,23 @@ namespace ir { 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) diff --git a/backend/src/ir/function.hpp b/backend/src/ir/function.hpp index 6af43a4..8c2cc82 100644 --- a/backend/src/ir/function.hpp +++ b/backend/src/ir/function.hpp @@ -27,6 +27,7 @@ #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" @@ -72,9 +73,15 @@ namespace ir { { 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 */ @@ -94,6 +101,12 @@ namespace ir { 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(); @@ -139,7 +152,7 @@ namespace ir { /*! 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 @@ -151,6 +164,7 @@ namespace ir { RegisterFile file; //!< RegisterDatas used by the instructions GrowingPool insnPool; //!< For fast instruction allocation bool structReturned; //!< First argument is pointer to struct + Profile profile; //!< Current function profile GBE_CLASS(Function); }; diff --git a/backend/src/ir/instruction.cpp b/backend/src/ir/instruction.cpp index fec684b..4c950f0 100644 --- a/backend/src/ir/instruction.cpp +++ b/backend/src/ir/instruction.cpp @@ -651,8 +651,8 @@ namespace ir { 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 */ diff --git a/backend/src/ir/instruction.hpp b/backend/src/ir/instruction.hpp index 57853ff..fb63a62 100644 --- a/backend/src/ir/instruction.hpp +++ b/backend/src/ir/instruction.hpp @@ -55,9 +55,6 @@ namespace ir { /*! 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 */ diff --git a/backend/src/ir/profile.cpp b/backend/src/ir/profile.cpp new file mode 100644 index 0000000..378e63f --- /dev/null +++ b/backend/src/ir/profile.cpp @@ -0,0 +1,61 @@ +/* + * 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 . + * + * Author: Benjamin Segovia + */ + +/** + * \file profile.hpp + * \author Benjamin Segovia + */ +#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 */ + + diff --git a/backend/src/ir/profile.hpp b/backend/src/ir/profile.hpp new file mode 100644 index 0000000..beb9dd3 --- /dev/null +++ b/backend/src/ir/profile.hpp @@ -0,0 +1,60 @@ +/* + * 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 . + * + * Author: Benjamin Segovia + */ + +/** + * \file profile.hpp + * \author Benjamin Segovia + */ +#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__ */ + diff --git a/backend/src/ir/value.hpp b/backend/src/ir/value.hpp index b4cdf3d..a9ac133 100644 --- a/backend/src/ir/value.hpp +++ b/backend/src/ir/value.hpp @@ -35,12 +35,26 @@ namespace ir { 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; @@ -55,6 +69,9 @@ namespace ir { 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 */ diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index ad902fc..858f925 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -210,7 +210,6 @@ 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 @@ -236,6 +235,8 @@ namespace gbe 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 */ @@ -1708,12 +1709,59 @@ static std::string CBEMangle(const std::string &S) { } } + ir::Register GenWriter::getConstantRegister(Constant *CPV) { + if (dyn_cast(CPV)) + GBE_ASSERTM(false, "Unsupported constant expression"); + else if (isa(CPV) && CPV->getType()->isSingleValueType()) + GBE_ASSERTM(false, "Unsupported constant expression"); + if (ConstantInt *CI = dyn_cast(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(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(value); - if (CPV && !isa(CPV)) { - GBE_ASSERT(0); - // printConstant(CPV, Static); - } else { + if (CPV && !isa(CPV)) + return getConstantRegister(CPV); + else { GBE_ASSERT(this->registerMap.find(value) != this->registerMap.end()); return this->registerMap[value]; } @@ -2432,7 +2480,9 @@ static std::string CBEMangle(const std::string &S) { } } - void GenWriter::visitCallInst(CallInst &I) { + void GenWriter::visitCallInst(CallInst &I) + { +#if 0 if (isa(I.getCalledValue())) return visitInlineAsm(I); @@ -2536,185 +2586,15 @@ static std::string CBEMangle(const std::string &S) { 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(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(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) { diff --git a/backend/src/sys/platform.hpp b/backend/src/sys/platform.hpp index b887aca..50aae1b 100644 --- a/backend/src/sys/platform.hpp +++ b/backend/src/sys/platform.hpp @@ -188,9 +188,9 @@ /*! 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 */ diff --git a/backend/src/utest/utest_llvm.cpp b/backend/src/utest/utest_llvm.cpp index f0d58e8..ff36d14 100644 --- a/backend/src/utest/utest_llvm.cpp +++ b/backend/src/utest/utest_llvm.cpp @@ -77,8 +77,9 @@ runTests: 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")); }