[Pkg-opencl-devel] [beignet] 21/66: Imported Upstream version 0.1+git20130502+63e60ed
Andreas Beckmann
anbe at moszumanska.debian.org
Fri Oct 31 07:27:03 UTC 2014
This is an automated email from the git hooks/post-receive script.
anbe pushed a commit to branch master
in repository beignet.
commit 4c2e469a3119d35eaec4158929090deaf6cfce13
Author: Simon Richter <sjr at debian.org>
Date: Mon May 6 06:30:32 2013 +0200
Imported Upstream version 0.1+git20130502+63e60ed
---
CMake/FindDRM.cmake | 1 +
CMake/FindDRMIntel.cmake | 1 +
CMake/FindEGL.cmake | 1 +
CMake/FindGBE.cmake | 1 +
CMake/FindGBM.cmake | 1 +
CMake/FindOCLIcd.cmake | 1 +
CMake/FindXext.cmake | 1 +
CMake/FindXfixes.cmake | 1 +
backend/src/backend/context.cpp | 12 ++++-
backend/src/backend/gen_reg_allocation.cpp | 1 +
backend/src/backend/program.cpp | 20 ++++++-
backend/src/backend/program.h | 10 ++++
backend/src/backend/program.hpp | 8 +++
backend/src/ir/constant.hpp | 39 +++++++++++++-
backend/src/ir/profile.cpp | 2 +
backend/src/ir/profile.hpp | 3 +-
backend/src/ir/unit.hpp | 4 ++
backend/src/llvm/llvm_gen_backend.cpp | 86 ++++++++++++++++++++++++++++++
backend/src/llvm/llvm_passes.cpp | 20 -------
backend/src/ocl_stdlib.h | 2 +-
kernels/compiler_ceil.cl | 4 ++
kernels/compiler_global_constant.cl | 8 +++
kernels/compiler_global_constant_2.cl | 9 ++++
src/cl_command_queue_gen7.c | 8 +++
utests/CMakeLists.txt | 3 ++
utests/compiler_ceil.cpp | 43 +++++++++++++++
utests/compiler_global_constant.cpp | 29 ++++++++++
utests/compiler_global_constant_2.cpp | 30 +++++++++++
28 files changed, 323 insertions(+), 26 deletions(-)
diff --git a/CMake/FindDRM.cmake b/CMake/FindDRM.cmake
index f65c457..3de35bf 100644
--- a/CMake/FindDRM.cmake
+++ b/CMake/FindDRM.cmake
@@ -27,6 +27,7 @@ FIND_LIBRARY(DRM_LIBRARY
DOC "The DRM library")
IF(DRM_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${DRM_INCLUDE_PATH})
SET(DRM_FOUND 1 CACHE STRING "Set to 1 if DRM is found, 0 otherwise")
ELSE(DRM_INCLUDE_PATH)
SET(DRM_FOUND 0 CACHE STRING "Set to 1 if DRM is found, 0 otherwise")
diff --git a/CMake/FindDRMIntel.cmake b/CMake/FindDRMIntel.cmake
index 2ab9c1f..2d45c64 100644
--- a/CMake/FindDRMIntel.cmake
+++ b/CMake/FindDRMIntel.cmake
@@ -28,6 +28,7 @@ FIND_LIBRARY(DRM_INTEL_LIBRARY
DOC "The DRM_INTEL library")
IF(DRM_INTEL_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${DRM_INTEL_INCLUDE_PATH})
SET(DRM_INTEL_FOUND 1 CACHE STRING "Set to 1 if DRM_INTEL is found, 0 otherwise")
ELSE(DRM_INTEL_INCLUDE_PATH)
SET(DRM_INTEL_FOUND 0 CACHE STRING "Set to 1 if DRM_INTEL is found, 0 otherwise")
diff --git a/CMake/FindEGL.cmake b/CMake/FindEGL.cmake
index 703ff64..69d4852 100644
--- a/CMake/FindEGL.cmake
+++ b/CMake/FindEGL.cmake
@@ -27,6 +27,7 @@ FIND_LIBRARY(EGL_LIBRARY
DOC "The EGL library")
IF(EGL_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${EGL_INCLUDE_PATH})
SET(EGL_FOUND 1 CACHE STRING "Set to 1 if EGL is found, 0 otherwise")
ELSE(EGL_INCLUDE_PATH)
SET(EGL_FOUND 0 CACHE STRING "Set to 1 if EGL is found, 0 otherwise")
diff --git a/CMake/FindGBE.cmake b/CMake/FindGBE.cmake
index 4670483..db938c7 100644
--- a/CMake/FindGBE.cmake
+++ b/CMake/FindGBE.cmake
@@ -27,6 +27,7 @@ FIND_LIBRARY(GBE_LIBRARY
DOC "The GBE library")
IF(GBE_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${GBE_INCLUDE_PATH})
SET(GBE_FOUND 1 CACHE STRING "Set to 1 if GBE is found, 0 otherwise")
ELSE(GBE_INCLUDE_PATH)
SET(GBE_FOUND 0 CACHE STRING "Set to 1 if GBE is found, 0 otherwise")
diff --git a/CMake/FindGBM.cmake b/CMake/FindGBM.cmake
index d3b6086..f20f4b2 100644
--- a/CMake/FindGBM.cmake
+++ b/CMake/FindGBM.cmake
@@ -27,6 +27,7 @@ FIND_LIBRARY(GBM_LIBRARY
DOC "The GBM library")
IF(GBM_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${GBM_INCLUDE_PATH})
SET(GBM_FOUND 1 CACHE STRING "Set to 1 if GBM is found, 0 otherwise")
ELSE(GBM_INCLUDE_PATH)
SET(GBM_FOUND 0 CACHE STRING "Set to 1 if GBM is found, 0 otherwise")
diff --git a/CMake/FindOCLIcd.cmake b/CMake/FindOCLIcd.cmake
index 076f00e..b0a8ad7 100644
--- a/CMake/FindOCLIcd.cmake
+++ b/CMake/FindOCLIcd.cmake
@@ -15,6 +15,7 @@ FIND_PATH(OCLIcd_INCLUDE_PATH ocl_icd.h
DOC "The directory where ocl_icd.h resides")
IF(OCLIcd_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${OCLIcd_INCLUDE_PATH})
SET(OCLIcd_FOUND 1 CACHE STRING "Set to 1 if OCLIcd is found, 0 otherwise")
ELSE(OCLIcd_INCLUDE_PATH)
SET(OCLIcd_FOUND 0 CACHE STRING "Set to 1 if OCLIcd is found, 0 otherwise")
diff --git a/CMake/FindXext.cmake b/CMake/FindXext.cmake
index 3025da8..5bbd719 100644
--- a/CMake/FindXext.cmake
+++ b/CMake/FindXext.cmake
@@ -25,6 +25,7 @@ FIND_LIBRARY(XEXT_LIBRARY
DOC "The XEXT library")
IF(XEXT_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${XEXT_INCLUDE_PATH})
SET(XEXT_FOUND 1 CACHE STRING "Set to 1 if XEXT is found, 0 otherwise")
ELSE(XEXT_INCLUDE_PATH)
SET(XEXT_FOUND 0 CACHE STRING "Set to 1 if XEXT is found, 0 otherwise")
diff --git a/CMake/FindXfixes.cmake b/CMake/FindXfixes.cmake
index 182e479..47259e1 100644
--- a/CMake/FindXfixes.cmake
+++ b/CMake/FindXfixes.cmake
@@ -25,6 +25,7 @@ FIND_LIBRARY(XFIXES_LIBRARY
DOC "The XFIXES library")
IF(XFIXES_INCLUDE_PATH)
+ INCLUDE_DIRECTORIES(${XFIXES_INCLUDE_PATH})
SET(XFIXES_FOUND 1 CACHE STRING "Set to 1 if XFIXES is found, 0 otherwise")
ELSE(XFIXES_INCLUDE_PATH)
SET(XFIXES_FOUND 0 CACHE STRING "Set to 1 if XFIXES is found, 0 otherwise")
diff --git a/backend/src/backend/context.cpp b/backend/src/backend/context.cpp
index 91d8d8c..c636b48 100644
--- a/backend/src/backend/context.cpp
+++ b/backend/src/backend/context.cpp
@@ -419,6 +419,15 @@ namespace gbe
}
});
#undef INSERT_REG
+ this->newCurbeEntry(GBE_CURBE_GLOBAL_CONSTANT_OFFSET, 0, sizeof(int));
+ specialRegs.insert(ir::ocl::constoffst);
+
+ // Insert serialized global constant arrays if used
+ const ir::ConstantSet& constantSet = unit.getConstantSet();
+ if (constantSet.getConstantNum()) {
+ size_t size = constantSet.getDataSize();
+ this->newCurbeEntry(GBE_CURBE_GLOBAL_CONSTANT_DATA, 0, size);
+ }
// Insert the number of threads
this->newCurbeEntry(GBE_CURBE_THREAD_NUM, 0, sizeof(uint32_t));
@@ -591,7 +600,8 @@ namespace gbe
reg == ir::ocl::gsize2 ||
reg == ir::ocl::goffset0 ||
reg == ir::ocl::goffset1 ||
- reg == ir::ocl::goffset2)
+ reg == ir::ocl::goffset2 ||
+ reg == ir::ocl::constoffst)
return true;
return false;
}
diff --git a/backend/src/backend/gen_reg_allocation.cpp b/backend/src/backend/gen_reg_allocation.cpp
index 10e4ab6..8c9f358 100644
--- a/backend/src/backend/gen_reg_allocation.cpp
+++ b/backend/src/backend/gen_reg_allocation.cpp
@@ -523,6 +523,7 @@ namespace gbe
allocatePayloadReg(GBE_CURBE_GROUP_NUM_Z, ocl::numgroup2);
allocatePayloadReg(GBE_CURBE_STACK_POINTER, ocl::stackptr);
allocatePayloadReg(GBE_CURBE_THREAD_NUM, ocl::threadn);
+ allocatePayloadReg(GBE_CURBE_GLOBAL_CONSTANT_OFFSET, ocl::constoffst);
// Group and barrier IDs are always allocated by the hardware in r0
RA.insert(std::make_pair(ocl::groupid0, 1*sizeof(float))); // r0.1
diff --git a/backend/src/backend/program.cpp b/backend/src/backend/program.cpp
index 0882e5a..38cc236 100644
--- a/backend/src/backend/program.cpp
+++ b/backend/src/backend/program.cpp
@@ -63,9 +63,10 @@ namespace gbe {
return it->offset; // we found it!
}
- Program::Program(void) {}
+ Program::Program(void) : constantSet(NULL) {}
Program::~Program(void) {
for (auto &kernel : kernels) GBE_DELETE(kernel.second);
+ if (constantSet) delete constantSet;
}
BVAR(OCL_OUTPUT_GEN_IR, false);
@@ -81,6 +82,7 @@ namespace gbe {
}
bool Program::buildFromUnit(const ir::Unit &unit, std::string &error) {
+ constantSet = new ir::ConstantSet(unit.getConstantSet());
const auto &set = unit.getFunctionSet();
const uint32_t kernelNum = set.size();
if (OCL_OUTPUT_GEN_IR) std::cout << unit;
@@ -146,6 +148,18 @@ namespace gbe {
return p;
}
+ static size_t programGetGlobalConstantSize(gbe_program gbeProgram) {
+ if (gbeProgram == NULL) return 0;
+ const gbe::Program *program = (const gbe::Program*) gbeProgram;
+ return program->getGlobalConstantSize();
+ }
+
+ static void programGetGlobalConstantData(gbe_program gbeProgram, char *mem) {
+ if (gbeProgram == NULL) return;
+ const gbe::Program *program = (const gbe::Program*) gbeProgram;
+ program->getGlobalConstantData(mem);
+ }
+
static uint32_t programGetKernelNum(gbe_program gbeProgram) {
if (gbeProgram == NULL) return 0;
const gbe::Program *program = (const gbe::Program*) gbeProgram;
@@ -244,6 +258,8 @@ namespace gbe {
GBE_EXPORT_SYMBOL gbe_program_new_from_source_cb *gbe_program_new_from_source = NULL;
GBE_EXPORT_SYMBOL gbe_program_new_from_binary_cb *gbe_program_new_from_binary = NULL;
GBE_EXPORT_SYMBOL gbe_program_new_from_llvm_cb *gbe_program_new_from_llvm = NULL;
+GBE_EXPORT_SYMBOL gbe_program_get_global_constant_size_cb *gbe_program_get_global_constant_size = NULL;
+GBE_EXPORT_SYMBOL gbe_program_get_global_constant_data_cb *gbe_program_get_global_constant_data = NULL;
GBE_EXPORT_SYMBOL gbe_program_delete_cb *gbe_program_delete = NULL;
GBE_EXPORT_SYMBOL gbe_program_get_kernel_num_cb *gbe_program_get_kernel_num = NULL;
GBE_EXPORT_SYMBOL gbe_program_get_kernel_by_name_cb *gbe_program_get_kernel_by_name = NULL;
@@ -269,6 +285,8 @@ namespace gbe
{
CallBackInitializer(void) {
gbe_program_new_from_source = gbe::programNewFromSource;
+ gbe_program_get_global_constant_size = gbe::programGetGlobalConstantSize;
+ gbe_program_get_global_constant_data = gbe::programGetGlobalConstantData;
gbe_program_delete = gbe::programDelete;
gbe_program_get_kernel_num = gbe::programGetKernelNum;
gbe_program_get_kernel_by_name = gbe::programGetKernelByName;
diff --git a/backend/src/backend/program.h b/backend/src/backend/program.h
index 4273a77..575196a 100644
--- a/backend/src/backend/program.h
+++ b/backend/src/backend/program.h
@@ -69,6 +69,8 @@ enum gbe_curbe_type {
GBE_CURBE_GROUP_NUM_X,
GBE_CURBE_GROUP_NUM_Y,
GBE_CURBE_GROUP_NUM_Z,
+ GBE_CURBE_GLOBAL_CONSTANT_OFFSET,
+ GBE_CURBE_GLOBAL_CONSTANT_DATA,
GBE_CURBE_IMAGE_WIDTH,
GBE_CURBE_IMAGE_HEIGHT,
GBE_CURBE_IMAGE_DEPTH,
@@ -104,6 +106,14 @@ typedef gbe_program (gbe_program_new_from_llvm_cb)(const char *fileName,
size_t *err_size);
extern gbe_program_new_from_llvm_cb *gbe_program_new_from_llvm;
+/*! Get the size of global constants */
+typedef size_t (gbe_program_get_global_constant_size_cb)(gbe_program gbeProgram);
+extern gbe_program_get_global_constant_size_cb *gbe_program_get_global_constant_size;
+
+/*! Get the content of global constants */
+typedef void (gbe_program_get_global_constant_data_cb)(gbe_program gbeProgram, char *mem);
+extern gbe_program_get_global_constant_data_cb *gbe_program_get_global_constant_data;
+
/*! Destroy and deallocate the given program */
typedef void (gbe_program_delete_cb)(gbe_program);
extern gbe_program_delete_cb *gbe_program_delete;
diff --git a/backend/src/backend/program.hpp b/backend/src/backend/program.hpp
index e7584d9..e754899 100644
--- a/backend/src/backend/program.hpp
+++ b/backend/src/backend/program.hpp
@@ -27,6 +27,8 @@
#include "backend/program.h"
#include "backend/context.hpp"
+#include "ir/constant.hpp"
+#include "ir/unit.hpp"
#include "sys/hash_map.hpp"
#include "sys/vector.hpp"
#include <string>
@@ -157,11 +159,17 @@ namespace gbe {
bool buildFromLLVMFile(const char *fileName, std::string &error);
/*! Buils a program from a OCL string */
bool buildFromSource(const char *source, std::string &error);
+ /*! Get size of the global constant arrays */
+ size_t getGlobalConstantSize(void) const { return constantSet->getDataSize(); }
+ /*! Get the content of global constant arrays */
+ void getGlobalConstantData(char *mem) const { constantSet->getData(mem); }
protected:
/*! Compile a kernel */
virtual Kernel *compileKernel(const ir::Unit &unit, const std::string &name) = 0;
/*! Kernels sorted by their name */
hash_map<std::string, Kernel*> kernels;
+ /*! Global (constants) outside any kernel */
+ ir::ConstantSet *constantSet;
/*! Use custom allocators */
GBE_CLASS(Program);
};
diff --git a/backend/src/ir/constant.hpp b/backend/src/ir/constant.hpp
index 3a23dc2..0717391 100644
--- a/backend/src/ir/constant.hpp
+++ b/backend/src/ir/constant.hpp
@@ -36,25 +36,31 @@ namespace ir {
public:
/*! Build a constant description */
INLINE Constant(const std::string &name, uint32_t size, uint32_t alignment, uint32_t offset) :
- name(name), size(size), alignment(alignment), offset(offset) {}
+ name(name), size(size), alignment(alignment), offset(offset), reg(0) {}
/*! Copy constructor */
INLINE Constant(const Constant &other) :
- name(other.name), size(other.size), alignment(other.alignment), offset(other.offset) {}
+ name(other.name), size(other.size), alignment(other.alignment), offset(other.offset), reg(other.reg) {}
/*! Copy operator */
INLINE Constant& operator= (const Constant &other) {
this->name = other.name;
this->size = other.size;
this->alignment = other.alignment;
this->offset = other.offset;
+ this->reg = other.reg;
return *this;
}
/*! Nothing happens here */
INLINE ~Constant(void) {}
+ const std::string& getName(void) const { return name; }
+ uint32_t getOffset(void) const { return offset; }
+ uint16_t getReg(void) const { return reg; }
+ void setReg(uint16_t reg) { this->reg = reg; }
private:
std::string name; //!< Optional name of the constant
uint32_t size; //!< Size of the constant
uint32_t alignment; //!< Alignment required for each constant
uint32_t offset; //!< Offset of the constant in the data segment
+ uint16_t reg; //!< Virtual register number
GBE_CLASS(Constant);
};
@@ -66,6 +72,35 @@ namespace ir {
public:
/*! Append a new constant in the constant set */
void append(const char*, const std::string&, uint32_t size, uint32_t alignment);
+ /*! Number of constants */
+ size_t getConstantNum(void) const { return constants.size(); }
+ /*! Get a special constant */
+ Constant& getConstant(size_t i) { return constants[i]; }
+ /*! Get a special constant */
+ Constant& getConstant(const std::string & name) {
+ for (auto & c : constants) {
+ if (c.getName() == name)
+ return c;
+ }
+ GBE_ASSERT(false);
+ return *(Constant *)nullptr;
+ }
+ /*! Number of bytes of serialized constant data */
+ size_t getDataSize(void) const { return data.size(); }
+ /*! Store serialized constant data into an array */
+ void getData(char *mem) const {
+ for (size_t i = 0; i < data.size(); i ++)
+ mem[i] = data[i];
+ }
+ ConstantSet() {}
+ ConstantSet(const ConstantSet& other) : data(other.data), constants(other.constants) {}
+ ConstantSet & operator = (const ConstantSet& other) {
+ if (&other != this) {
+ data = other.data;
+ constants = other.constants;
+ }
+ return *this;
+ }
private:
vector<char> data; //!< The constant data serialized in one array
vector<Constant> constants;//!< Each constant description
diff --git a/backend/src/ir/profile.cpp b/backend/src/ir/profile.cpp
index 61531be..c1dc650 100644
--- a/backend/src/ir/profile.cpp
+++ b/backend/src/ir/profile.cpp
@@ -40,6 +40,7 @@ namespace ir {
"stack_pointer",
"block_ip",
"barrier_id", "thread_number",
+ "const_curbe_offset",
};
#if GBE_DEBUG
@@ -74,6 +75,7 @@ namespace ir {
DECL_NEW_REG(FAMILY_WORD, blockip);
DECL_NEW_REG(FAMILY_DWORD, barrierid);
DECL_NEW_REG(FAMILY_DWORD, threadn);
+ DECL_NEW_REG(FAMILY_DWORD, constoffst);
}
#undef DECL_NEW_REG
diff --git a/backend/src/ir/profile.hpp b/backend/src/ir/profile.hpp
index 12050ff..32dd149 100644
--- a/backend/src/ir/profile.hpp
+++ b/backend/src/ir/profile.hpp
@@ -63,7 +63,8 @@ namespace ir {
static const Register blockip = Register(19); // blockip
static const Register barrierid = Register(20);// barrierid
static const Register threadn = Register(21); // number of threads
- static const uint32_t regNum = 22; // number of special registers
+ static const Register constoffst = Register(22); // offset of global constant array's curbe
+ static const uint32_t regNum = 23; // number of special registers
extern const char *specialRegMean[]; // special register name.
} /* namespace ocl */
diff --git a/backend/src/ir/unit.hpp b/backend/src/ir/unit.hpp
index 37a5dbf..ae78638 100644
--- a/backend/src/ir/unit.hpp
+++ b/backend/src/ir/unit.hpp
@@ -67,6 +67,10 @@ namespace ir {
else
return FAMILY_QWORD;
}
+ /*! Return the constant set */
+ ConstantSet& getConstantSet(void) { return constantSet; }
+ /*! Return the constant set */
+ const ConstantSet& getConstantSet(void) const { return constantSet; }
private:
friend class ContextInterface; //!< Can free modify the unit
hash_map<std::string, Function*> functions; //!< All the defined functions
diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 42265ee..637e7be 100644
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -455,6 +455,8 @@ namespace gbe
virtual bool doInitialization(Module &M);
+ void collectGlobalConstant(void) const;
+
bool runOnFunction(Function &F) {
// Do not codegen any 'available_externally' functions at all, they have
// definitions outside the translation unit.
@@ -550,11 +552,60 @@ namespace gbe
char GenWriter::ID = 0;
+ void GenWriter::collectGlobalConstant(void) const {
+ const Module::GlobalListType &globalList = TheModule->getGlobalList();
+ for(auto i = globalList.begin(); i != globalList.end(); i ++) {
+ const GlobalVariable &v = *i;
+ const char *name = v.getName().data();
+ unsigned addrSpace = v.getType()->getAddressSpace();
+ if(addrSpace == ir::AddressSpace::MEM_CONSTANT) {
+ GBE_ASSERT(v.hasInitializer());
+ const Constant *c = v.getInitializer();
+ GBE_ASSERT(c->getType()->getTypeID() == Type::ArrayTyID);
+ const ConstantDataArray *cda = dyn_cast<ConstantDataArray>(c);
+ GBE_ASSERT(cda);
+ unsigned len = cda->getNumElements();
+ uint64_t elementSize = cda->getElementByteSize();
+ Type::TypeID typeID = cda->getElementType()->getTypeID();
+ if(typeID == Type::TypeID::IntegerTyID)
+ elementSize = sizeof(unsigned);
+ void *mem = malloc(elementSize * len);
+ for(unsigned j = 0; j < len; j ++) {
+ switch(typeID) {
+ case Type::TypeID::FloatTyID:
+ {
+ float f = cda->getElementAsFloat(j);
+ memcpy((float *)mem + j, &f, elementSize);
+ }
+ break;
+ case Type::TypeID::DoubleTyID:
+ {
+ double d = cda->getElementAsDouble(j);
+ memcpy((double *)mem + j, &d, elementSize);
+ }
+ break;
+ case Type::TypeID::IntegerTyID:
+ {
+ unsigned u = (unsigned) cda->getElementAsInteger(j);
+ memcpy((unsigned *)mem + j, &u, elementSize);
+ }
+ break;
+ default:
+ NOT_IMPLEMENTED;
+ }
+ }
+ unit.newConstant((char *)mem, name, elementSize * len, sizeof(unsigned));
+ free(mem);
+ }
+ }
+ }
+
bool GenWriter::doInitialization(Module &M) {
FunctionPass::doInitialization(M);
// Initialize
TheModule = &M;
+ collectGlobalConstant();
return false;
}
@@ -704,6 +755,17 @@ namespace gbe
}
ir::Register GenWriter::getRegister(Value *value, uint32_t elemID) {
+ if (dyn_cast<ConstantExpr>(value)) {
+ ConstantExpr *ce = dyn_cast<ConstantExpr>(value);
+ if(ce->isCast()) {
+ GBE_ASSERT(ce->getOpcode() == Instruction::PtrToInt);
+ const Value *pointer = ce->getOperand(0);
+ GBE_ASSERT(pointer->hasName());
+ auto name = pointer->getName().str();
+ uint16_t reg = unit.getConstantSet().getConstant(name).getReg();
+ return ir::Register(reg);
+ }
+ }
Constant *CPV = dyn_cast<Constant>(value);
if (CPV) {
GBE_ASSERT(isa<GlobalValue>(CPV) == false);
@@ -1076,6 +1138,30 @@ namespace gbe
this->labelMap.clear();
this->emitFunctionPrototype(F);
+ // Allocate a virtual register for each global constant array
+ const Module::GlobalListType &globalList = TheModule->getGlobalList();
+ size_t j = 0;
+ for(auto i = globalList.begin(); i != globalList.end(); i ++) {
+ const GlobalVariable &v = *i;
+ unsigned addrSpace = v.getType()->getAddressSpace();
+ if(addrSpace != ir::AddressSpace::MEM_CONSTANT)
+ continue;
+ GBE_ASSERT(v.hasInitializer());
+ const Constant *c = v.getInitializer();
+ GBE_ASSERT(c->getType()->getTypeID() == Type::ArrayTyID);
+ const ConstantDataArray *cda = dyn_cast<ConstantDataArray>(c);
+ GBE_ASSERT(cda);
+ ir::Register reg = ctx.reg(ir::RegisterFamily::FAMILY_DWORD);
+ ir::Constant &con = unit.getConstantSet().getConstant(j ++);
+ con.setReg(reg.value());
+ if(con.getOffset() != 0) {
+ ctx.LOADI(ir::TYPE_S32, reg, ctx.newIntegerImmediate(con.getOffset(), ir::TYPE_S32));
+ ctx.ADD(ir::TYPE_S32, reg, ir::ocl::constoffst, reg);
+ } else {
+ ctx.MOV(ir::TYPE_S32, reg, ir::ocl::constoffst);
+ }
+ }
+
// Visit all the instructions and emit the IR registers or the value to
// value mapping when a new register is not needed
pass = PASS_EMIT_REGISTERS;
diff --git a/backend/src/llvm/llvm_passes.cpp b/backend/src/llvm/llvm_passes.cpp
index b38ef58..1a7a658 100644
--- a/backend/src/llvm/llvm_passes.cpp
+++ b/backend/src/llvm/llvm_passes.cpp
@@ -228,26 +228,6 @@ namespace gbe
#endif
CompositeType* CompTy = cast<CompositeType>(parentPointer->getType());
- if(isa<GlobalVariable>(parentPointer)) //HACK: !!!!
- {
-#if 1//FORMER_VERSION
- Function *constWrapper =
- Function::Create(FunctionType::get(parentPointer->getType(),true),
- GlobalValue::ExternalLinkage,
- Twine("__gen_ocl_const_wrapper"));
-
- llvm::ArrayRef<Value*> params(parentPointer);
- // params.push_back(parentPointer);
-
- //create and insert wrapper call
- CallInst * wrapperCall =
- CallInst::Create(constWrapper,params,"",GEPInst);
- parentPointer = wrapperCall;
-#else
- // NOT_IMPLEMENTED;
-#endif
- }
-
Value* currentAddrInst =
new PtrToIntInst(parentPointer, IntegerType::get(GEPInst->getContext(), ptrSize), "", GEPInst);
diff --git a/backend/src/ocl_stdlib.h b/backend/src/ocl_stdlib.h
index 79616b8..a21f332 100644
--- a/backend/src/ocl_stdlib.h
+++ b/backend/src/ocl_stdlib.h
@@ -728,7 +728,7 @@ INLINE OVERLOADABLE float16 mix(float16 x, float16 y, float a) { return mix(x,y,
#define trunc __gen_ocl_internal_trunc
#define round __gen_ocl_internal_round
#define floor __gen_ocl_internal_floor
-#define ceil __gen_ocl_internal_ceil,
+#define ceil __gen_ocl_internal_ceil
#define log __gen_ocl_internal_log
#define log2 __gen_ocl_internal_log2
#define log10 __gen_ocl_internal_log10
diff --git a/kernels/compiler_ceil.cl b/kernels/compiler_ceil.cl
new file mode 100644
index 0000000..cf27483
--- /dev/null
+++ b/kernels/compiler_ceil.cl
@@ -0,0 +1,4 @@
+kernel void compiler_ceil(global float *src, global float *dst) {
+ int i = get_global_id(0);
+ dst[i] = ceil(src[i]);
+}
diff --git a/kernels/compiler_global_constant.cl b/kernels/compiler_global_constant.cl
new file mode 100644
index 0000000..af3c7b1
--- /dev/null
+++ b/kernels/compiler_global_constant.cl
@@ -0,0 +1,8 @@
+constant int m[3] = {71,72,73};
+
+__kernel void
+compiler_global_constant(__global int *dst, int e, int r)
+{
+ int id = (int)get_global_id(0);
+ dst[id] = m[id%3] + e + r;
+}
diff --git a/kernels/compiler_global_constant_2.cl b/kernels/compiler_global_constant_2.cl
new file mode 100644
index 0000000..353ebd7
--- /dev/null
+++ b/kernels/compiler_global_constant_2.cl
@@ -0,0 +1,9 @@
+constant int m[3] = {0x15b,0x25b,0x35b};
+constant int t[5] = {0x45b,0x55b,0x65b,0x75b,0x85b};
+
+__kernel void
+compiler_global_constant_2(__global int *dst, int e, int r)
+{
+ int id = (int)get_global_id(0);
+ dst[id] = m[id%3] + t[id%5] + e + r;
+}
diff --git a/src/cl_command_queue_gen7.c b/src/cl_command_queue_gen7.c
index 9402549..108684f 100644
--- a/src/cl_command_queue_gen7.c
+++ b/src/cl_command_queue_gen7.c
@@ -120,6 +120,7 @@ cl_curbe_fill(cl_kernel ker,
UPLOAD(GBE_CURBE_GROUP_NUM_Y, global_wk_sz[1]/local_wk_sz[1]);
UPLOAD(GBE_CURBE_GROUP_NUM_Z, global_wk_sz[2]/local_wk_sz[2]);
UPLOAD(GBE_CURBE_THREAD_NUM, thread_n);
+ UPLOAD(GBE_CURBE_GLOBAL_CONSTANT_OFFSET, gbe_kernel_get_curbe_offset(ker->opaque, GBE_CURBE_GLOBAL_CONSTANT_DATA, 0) + 32);
#undef UPLOAD
/* Write identity for the stack pointer. This is required by the stack pointer
@@ -132,6 +133,13 @@ cl_curbe_fill(cl_kernel ker,
for (i = 0; i < (int32_t) simd_sz; ++i) stackptr[i] = i;
}
+ /* Write global constant arrays */
+ if ((offset = gbe_kernel_get_curbe_offset(ker->opaque, GBE_CURBE_GLOBAL_CONSTANT_DATA, 0)) >= 0) {
+ /* Write the global constant arrays */
+ gbe_program prog = ker->program->opaque;
+ gbe_program_get_global_constant_data(prog, ker->curbe + offset);
+ }
+
/* Handle the various offsets to SLM */
const int32_t arg_n = gbe_kernel_get_arg_num(ker->opaque);
int32_t arg, slm_offset = 0;
diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt
index b2e3c97..1275faa 100644
--- a/utests/CMakeLists.txt
+++ b/utests/CMakeLists.txt
@@ -19,6 +19,7 @@ set (utests_sources
compiler_array2.cpp
compiler_array3.cpp
compiler_byte_scatter.cpp
+ compiler_ceil.cpp
compiler_convert_uchar_sat.cpp
compiler_copy_buffer.cpp
compiler_copy_image.cpp
@@ -32,6 +33,8 @@ set (utests_sources
compiler_function_constant0.cpp
compiler_function_constant1.cpp
compiler_function_constant.cpp
+ compiler_global_constant.cpp
+ compiler_global_constant_2.cpp
compiler_if_else.cpp
compiler_integer_division.cpp
compiler_integer_remainder.cpp
diff --git a/utests/compiler_ceil.cpp b/utests/compiler_ceil.cpp
new file mode 100644
index 0000000..29c7551
--- /dev/null
+++ b/utests/compiler_ceil.cpp
@@ -0,0 +1,43 @@
+#include <cmath>
+#include "utest_helper.hpp"
+
+static void cpu(int global_id, float *src, float *dst) {
+ dst[global_id] = ceilf(src[global_id]);
+}
+
+void compiler_ceil(void)
+{
+ const size_t n = 16;
+ float cpu_dst[16], cpu_src[16];
+
+ // Setup kernel and buffers
+ OCL_CREATE_KERNEL("compiler_ceil");
+ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(float), NULL);
+ OCL_CREATE_BUFFER(buf[1], 0, n * sizeof(float), NULL);
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(cl_mem), &buf[1]);
+ globals[0] = 16;
+ locals[0] = 16;
+
+ // Run random tests
+ for (uint32_t pass = 0; pass < 8; ++pass) {
+ OCL_MAP_BUFFER(0);
+ for (int32_t i = 0; i < (int32_t) n; ++i)
+ cpu_src[i] = ((float*)buf_data[0])[i] = .1f * (rand() & 15) - .75f;
+ OCL_UNMAP_BUFFER(0);
+
+ // Run the kernel on GPU
+ OCL_NDRANGE(1);
+
+ // Run on CPU
+ for (int32_t i = 0; i < (int32_t) n; ++i) cpu(i, cpu_src, cpu_dst);
+
+ // Compare
+ OCL_MAP_BUFFER(1);
+ for (int32_t i = 0; i < (int32_t) n; ++i)
+ OCL_ASSERT(((float *)buf_data[1])[i] == cpu_dst[i]);
+ OCL_UNMAP_BUFFER(1);
+ }
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_ceil);
diff --git a/utests/compiler_global_constant.cpp b/utests/compiler_global_constant.cpp
new file mode 100644
index 0000000..1547020
--- /dev/null
+++ b/utests/compiler_global_constant.cpp
@@ -0,0 +1,29 @@
+#include "utest_helper.hpp"
+
+void compiler_global_constant(void)
+{
+ const size_t n = 2048;
+ const uint32_t e = 34, r = 77;
+
+ // Setup kernel and buffers
+ OCL_CREATE_KERNEL("compiler_global_constant");
+ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(uint32_t), &e);
+ OCL_SET_ARG(2, sizeof(uint32_t), &r);
+
+ // Run the kernel
+ globals[0] = n;
+ locals[0] = 16;
+ OCL_NDRANGE(1);
+
+ unsigned int m[3] = {71,72,73};
+
+ // Check results
+ OCL_MAP_BUFFER(0);
+ for (uint32_t i = 0; i < n; ++i)
+ OCL_ASSERT(((uint32_t *)buf_data[0])[i] == m[i%3] + e + r);
+ OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_global_constant);
diff --git a/utests/compiler_global_constant_2.cpp b/utests/compiler_global_constant_2.cpp
new file mode 100644
index 0000000..56fccb5
--- /dev/null
+++ b/utests/compiler_global_constant_2.cpp
@@ -0,0 +1,30 @@
+#include "utest_helper.hpp"
+
+void compiler_global_constant_2(void)
+{
+ const size_t n = 2048;
+ const uint32_t e = 34, r = 77;
+
+ // Setup kernel and buffers
+ OCL_CREATE_KERNEL("compiler_global_constant_2");
+ OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL);
+ OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]);
+ OCL_SET_ARG(1, sizeof(uint32_t), &e);
+ OCL_SET_ARG(2, sizeof(uint32_t), &r);
+
+ // Run the kernel
+ globals[0] = n;
+ locals[0] = 16;
+ OCL_NDRANGE(1);
+
+ unsigned int m[3] = {0x15b,0x25b,0x35b};
+ unsigned int t[5] = {0x45b,0x55b,0x65b,0x75b,0x85b};
+
+ // Check results
+ OCL_MAP_BUFFER(0);
+ for (uint32_t i = 0; i < n; ++i)
+ OCL_ASSERT(((uint32_t *)buf_data[0])[i] == m[i%3] + t[i%5] + e + r);
+ OCL_UNMAP_BUFFER(0);
+}
+
+MAKE_UTEST_FROM_FUNCTION(compiler_global_constant_2);
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/beignet.git
More information about the Pkg-opencl-devel
mailing list