Pushed, thanks.
On Thu, Oct 10, 2013 at 03:13:50PM +0800, Ruiling Song wrote: > As Clang treat local variable in similar way like global constant, > (they are treated as Global variable in each own address space) > we refine the previous constant implementation in order to > share same code between local variable and global constant. > > We will allocate an address register for each GlobalVariable > (constant or local) through calling newRegister(). > In later step, through getRegister() we will get a proper > register derived from the allocated address register. > > Signed-off-by: Ruiling Song <ruiling.s...@intel.com> > --- > backend/src/backend/context.cpp | 1 + > backend/src/backend/program.cpp | 10 ++- > backend/src/backend/program.h | 3 + > backend/src/backend/program.hpp | 3 + > backend/src/ir/function.cpp | 2 +- > backend/src/ir/function.hpp | 5 ++ > backend/src/llvm/llvm_gen_backend.cpp | 157 > ++++++++++++++++++++++----------- > kernels/compiler_local_slm.cl | 28 ++++-- > src/cl_command_queue_gen7.c | 3 +- > utests/CMakeLists.txt | 1 + > utests/compiler_local_slm.cpp | 30 ++++++- > 11 files changed, 179 insertions(+), 64 deletions(-) > > diff --git a/backend/src/backend/context.cpp b/backend/src/backend/context.cpp > index cbd38f1..bc15761 100644 > --- a/backend/src/backend/context.cpp > +++ b/backend/src/backend/context.cpp > @@ -632,6 +632,7 @@ namespace gbe > void Context::handleSLM(void) { > const bool useSLM = fn.getUseSLM(); > kernel->useSLM = useSLM; > + kernel->slmSize = fn.getSLMSize(); > } > > bool Context::isScalarReg(const ir::Register ®) const { > diff --git a/backend/src/backend/program.cpp b/backend/src/backend/program.cpp > index 6550eac..24029c7 100644 > --- a/backend/src/backend/program.cpp > +++ b/backend/src/backend/program.cpp > @@ -75,7 +75,7 @@ > namespace gbe { > > Kernel::Kernel(const std::string &name) : > - name(name), args(NULL), argNum(0), curbeSize(0), stackSize(0), > useSLM(false), ctx(NULL), samplerSet(NULL), imageSet(NULL) > + name(name), args(NULL), argNum(0), curbeSize(0), stackSize(0), > useSLM(false), slmSize(0), ctx(NULL), samplerSet(NULL), imageSet(NULL) > {} > Kernel::~Kernel(void) { > if(ctx) GBE_DELETE(ctx); > @@ -709,6 +709,12 @@ namespace gbe { > return kernel->getUseSLM() ? 1 : 0; > } > > + static int32_t kernelGetSLMSize(gbe_kernel genKernel) { > + if (genKernel == NULL) return 0; > + const gbe::Kernel *kernel = (const gbe::Kernel*) genKernel; > + return kernel->getSLMSize(); > + } > + > static int32_t kernelSetConstBufSize(gbe_kernel genKernel, uint32_t argID, > size_t sz) { > if (genKernel == NULL) return -1; > gbe::Kernel *kernel = (gbe::Kernel*) genKernel; > @@ -776,6 +782,7 @@ GBE_EXPORT_SYMBOL gbe_kernel_get_scratch_size_cb > *gbe_kernel_get_scratch_size = > GBE_EXPORT_SYMBOL gbe_kernel_set_const_buffer_size_cb > *gbe_kernel_set_const_buffer_size = NULL; > GBE_EXPORT_SYMBOL gbe_kernel_get_required_work_group_size_cb > *gbe_kernel_get_required_work_group_size = NULL; > GBE_EXPORT_SYMBOL gbe_kernel_use_slm_cb *gbe_kernel_use_slm = NULL; > +GBE_EXPORT_SYMBOL gbe_kernel_get_slm_size_cb *gbe_kernel_get_slm_size = NULL; > GBE_EXPORT_SYMBOL gbe_kernel_get_sampler_size_cb > *gbe_kernel_get_sampler_size = NULL; > GBE_EXPORT_SYMBOL gbe_kernel_get_sampler_data_cb > *gbe_kernel_get_sampler_data = NULL; > GBE_EXPORT_SYMBOL gbe_kernel_get_image_size_cb *gbe_kernel_get_image_size = > NULL; > @@ -810,6 +817,7 @@ namespace gbe > gbe_kernel_set_const_buffer_size = gbe::kernelSetConstBufSize; > gbe_kernel_get_required_work_group_size = > gbe::kernelGetRequiredWorkGroupSize; > gbe_kernel_use_slm = gbe::kernelUseSLM; > + gbe_kernel_get_slm_size = gbe::kernelGetSLMSize; > gbe_kernel_get_sampler_size = gbe::kernelGetSamplerSize; > gbe_kernel_get_sampler_data = gbe::kernelGetSamplerData; > gbe_kernel_get_image_size = gbe::kernelGetImageSize; > diff --git a/backend/src/backend/program.h b/backend/src/backend/program.h > index 8774344..10fcc49 100644 > --- a/backend/src/backend/program.h > +++ b/backend/src/backend/program.h > @@ -218,6 +218,9 @@ extern gbe_kernel_get_required_work_group_size_cb > *gbe_kernel_get_required_work_ > /*! Says if SLM is used. Required to reconfigure the L3 complex */ > typedef int32_t (gbe_kernel_use_slm_cb)(gbe_kernel); > extern gbe_kernel_use_slm_cb *gbe_kernel_use_slm; > +/*! Get slm size needed for kernel local variables */ > +typedef int32_t (gbe_kernel_get_slm_size_cb)(gbe_kernel); > +extern gbe_kernel_get_slm_size_cb *gbe_kernel_get_slm_size; > > #ifdef __cplusplus > } > diff --git a/backend/src/backend/program.hpp b/backend/src/backend/program.hpp > index 28a792d..895cd01 100644 > --- a/backend/src/backend/program.hpp > +++ b/backend/src/backend/program.hpp > @@ -104,6 +104,8 @@ namespace gbe { > INLINE uint32_t getSIMDWidth(void) const { return this->simdWidth; } > /*! Says if SLM is needed for it */ > INLINE bool getUseSLM(void) const { return this->useSLM; } > + /*! get slm size for kernel local variable */ > + INLINE uint32_t getSLMSize(void) const { return this->slmSize; } > /*! set constant buffer size and return the cb curbe offset */ > int32_t setConstBufSize(uint32_t argID, size_t sz) { > if(argID >= argNum) return -1; > @@ -169,6 +171,7 @@ namespace gbe { > uint32_t stackSize; //!< Stack size (may be 0 if unused) > uint32_t scratchSize; //!< Scratch memory size (may be 0 if unused) > bool useSLM; //!< SLM requires a special HW config > + uint32_t slmSize; //!< slm size for kernel variable > Context *ctx; //!< Save context after compiler to alloc > constant buffer curbe > ir::SamplerSet *samplerSet;//!< Copy from the corresponding function. > ir::ImageSet *imageSet; //!< Copy from the corresponding function. > diff --git a/backend/src/ir/function.cpp b/backend/src/ir/function.cpp > index 88aae08..c15c292 100644 > --- a/backend/src/ir/function.cpp > +++ b/backend/src/ir/function.cpp > @@ -43,7 +43,7 @@ namespace ir { > /////////////////////////////////////////////////////////////////////////// > > Function::Function(const std::string &name, const Unit &unit, Profile > profile) : > - name(name), unit(unit), profile(profile), simdWidth(0), useSLM(false) > + name(name), unit(unit), profile(profile), simdWidth(0), useSLM(false), > slmSize(0) > { > initProfile(*this); > samplerSet = GBE_NEW(SamplerSet); > diff --git a/backend/src/ir/function.hpp b/backend/src/ir/function.hpp > index 6e712cd..3d4733d 100644 > --- a/backend/src/ir/function.hpp > +++ b/backend/src/ir/function.hpp > @@ -301,6 +301,10 @@ namespace ir { > INLINE bool getUseSLM(void) const { return this->useSLM; } > /*! Change the SLM config for the function */ > INLINE bool setUseSLM(bool useSLM) { return this->useSLM = useSLM; } > + /*! get SLM size needed for local variable inside kernel function */ > + INLINE uint32_t getSLMSize(void) const { return this->slmSize; } > + /*! set slm size needed for local variable inside kernel function */ > + INLINE void setSLMSize(uint32_t size) { this->slmSize = size; } > /*! Get sampler set in this function */ > SamplerSet* getSamplerSet(void) const {return samplerSet; } > /*! Get image set in this function */ > @@ -320,6 +324,7 @@ namespace ir { > LocationMap locationMap; //!< Pushed function arguments (loc->reg) > uint32_t simdWidth; //!< 8 or 16 if forced, 0 otherwise > bool useSLM; //!< Is SLM required? > + uint32_t slmSize; //!< local variable size inside kernel > function > SamplerSet *samplerSet; //!< samplers used in this function. > ImageSet* imageSet; //!< Image set in this function's > arguments.. > GBE_CLASS(Function); //!< Use custom allocator > diff --git a/backend/src/llvm/llvm_gen_backend.cpp > b/backend/src/llvm/llvm_gen_backend.cpp > index 5b6857d..7af5bb8 100644 > --- a/backend/src/llvm/llvm_gen_backend.cpp > +++ b/backend/src/llvm/llvm_gen_backend.cpp > @@ -321,7 +321,9 @@ namespace gbe > /*! Allocate a new scalar register */ > ir::Register newScalar(Value *value, Value *key = NULL, uint32_t index = > 0u) > { > - GBE_ASSERT(dyn_cast<Constant>(value) == NULL); > + // we don't allow normal constant, but GlobalValue is a special case, > + // it needs a register to store its address > + GBE_ASSERT(! (isa<Constant>(value) && !isa<GlobalValue>(value))); > Type *type = value->getType(); > auto typeID = type->getTypeID(); > switch (typeID) { > @@ -477,7 +479,8 @@ namespace gbe > } > > virtual bool doFinalization(Module &M) { return false; } > - > + /*! handle global variable register allocation (local, constant space) */ > + void allocateGlobalVariableRegister(Function &F); > /*! Emit the complete function code and declaration */ > void emitFunction(Function &F); > /*! Handle input and output function parameters */ > @@ -488,6 +491,8 @@ namespace gbe > void emitMovForPHI(BasicBlock *curr, BasicBlock *succ); > /*! Alocate one or several registers (if vector) for the value */ > INLINE void newRegister(Value *value, Value *key = NULL); > + /*! get the register for a llvm::Constant */ > + ir::Register getConstantRegister(Constant *c, uint32_t index = 0); > /*! Return a valid register from an operand (can use LOADI to make one) > */ > INLINE ir::Register getRegister(Value *value, uint32_t index = 0); > /*! Create a new immediate from a constant */ > @@ -838,40 +843,46 @@ namespace gbe > }; > } > > - ir::Register GenWriter::getRegister(Value *value, uint32_t elemID) { > - //the real value may be constant, so get real value before constant check > - regTranslator.getRealValue(value, elemID); > + ir::Register GenWriter::getConstantRegister(Constant *c, uint32_t elemID) { > + GBE_ASSERT(c != NULL); > > - 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); > - } > + if(isa<GlobalValue>(c)) { > + return regTranslator.getScalar(c, elemID); > } > - Constant *CPV = dyn_cast<Constant>(value); > - if (CPV) { > - if (isa<GlobalValue>(CPV)) { > - auto name = CPV->getName().str(); > - uint16_t reg = unit.getConstantSet().getConstant(name).getReg(); > - return ir::Register(reg); > - } > - if (isa<ConstantExpr>(CPV)) { > + > + if(isa<ConstantExpr>(c)) { > + ConstantExpr * ce = dyn_cast<ConstantExpr>(c); > + > + if(ce->isCast()) { > + Value* op = ce->getOperand(0); > + ir::Register pointer_reg; > + if(isa<ConstantExpr>(op)) { > + // try to get the real pointer register, for case like: > + // store i64 ptrtoint (i8 addrspace(3)* getelementptr inbounds ... > + // in which ptrtoint and getelementptr are ConstantExpr. > + pointer_reg = getConstantRegister(dyn_cast<Constant>(op), elemID); > + } else { > + pointer_reg = regTranslator.getScalar(op, elemID); > + } > + // if ptrToInt request another type other than 32bit, convert as > requested > + ir::Type dstType = getType(ctx, ce->getType()); > + if(ce->getOpcode() == Instruction::PtrToInt && ir::TYPE_S32 != > dstType) { > + ir::Register tmp = ctx.reg(getFamily(dstType)); > + ctx.CVT(dstType, ir::TYPE_S32, tmp, pointer_reg); > + return tmp; > + } > + return pointer_reg; > + } else { > uint32_t TypeIndex; > uint32_t constantOffset = 0; > uint32_t offset = 0; > - ConstantExpr *CE = dyn_cast<ConstantExpr>(CPV); > > // currently only GetElementPtr is handled > - GBE_ASSERT(CE->getOpcode() == Instruction::GetElementPtr); > - Value *pointer = CE->getOperand(0); > + GBE_ASSERT(ce->getOpcode() == Instruction::GetElementPtr); > + Value *pointer = ce->getOperand(0); > CompositeType* CompTy = cast<CompositeType>(pointer->getType()); > - for(uint32_t op=1; op<CE->getNumOperands(); ++op) { > - ConstantInt* ConstOP = dyn_cast<ConstantInt>(CE->getOperand(op)); > + for(uint32_t op=1; op<ce->getNumOperands(); ++op) { > + ConstantInt* ConstOP = dyn_cast<ConstantInt>(ce->getOperand(op)); > GBE_ASSERT(ConstOP); > TypeIndex = ConstOP->getZExtValue(); > for(uint32_t ty_i=0; ty_i<TypeIndex; ty_i++) > @@ -889,21 +900,30 @@ namespace gbe > CompTy = > dyn_cast<CompositeType>(CompTy->getTypeAtIndex(TypeIndex)); > } > > - const std::string &pointer_name = pointer->getName().str(); > - ir::Register pointer_reg = > ir::Register(unit.getConstantSet().getConstant(pointer_name).getReg()); > + ir::Register pointer_reg; > + pointer_reg = regTranslator.getScalar(pointer, elemID); > ir::Register offset_reg = ctx.reg(ir::RegisterFamily::FAMILY_DWORD); > ctx.LOADI(ir::Type::TYPE_S32, offset_reg, > ctx.newIntegerImmediate(constantOffset, ir::Type::TYPE_S32)); > ir::Register reg = ctx.reg(ir::RegisterFamily::FAMILY_DWORD); > ctx.ADD(ir::Type::TYPE_S32, reg, pointer_reg, offset_reg); > return reg; > } > - const ir::ImmediateIndex immIndex = this->newImmediate(CPV, elemID); > - const ir::Immediate imm = ctx.getImmediate(immIndex); > - const ir::Register reg = ctx.reg(getFamily(imm.type)); > - ctx.LOADI(imm.type, reg, immIndex); > - return reg; > } > - else > + > + const ir::ImmediateIndex immIndex = this->newImmediate(c, elemID); > + const ir::Immediate imm = ctx.getImmediate(immIndex); > + const ir::Register reg = ctx.reg(getFamily(imm.type)); > + ctx.LOADI(imm.type, reg, immIndex); > + return reg; > + } > + > + ir::Register GenWriter::getRegister(Value *value, uint32_t elemID) { > + //the real value may be constant, so get real value before constant check > + regTranslator.getRealValue(value, elemID); > + if(isa<Constant>(value)) { > + Constant *c = dyn_cast<Constant>(value); > + return getConstantRegister(c, elemID); > + } else > return regTranslator.getScalar(value, elemID); > } > > @@ -1273,6 +1293,55 @@ namespace gbe > BVAR(OCL_OPTIMIZE_PHI_MOVES, true); > BVAR(OCL_OPTIMIZE_LOADI, true); > > + void GenWriter::allocateGlobalVariableRegister(Function &F) > + { > + // Allocate a address register for each global variable > + const Module::GlobalListType &globalList = TheModule->getGlobalList(); > + size_t j = 0; > + for(auto i = globalList.begin(); i != globalList.end(); i ++) { > + const GlobalVariable &v = *i; > + if(!v.isConstantUsed()) continue; > + > + ir::AddressSpace addrSpace = > addressSpaceLLVMToGen(v.getType()->getAddressSpace()); > + if(addrSpace == ir::MEM_LOCAL) { > + ir::Function &f = ctx.getFunction(); > + f.setUseSLM(true); > + const Constant *c = v.getInitializer(); > + Type *ty = c->getType(); > + uint32_t oldSlm = f.getSLMSize(); > + uint32_t align = 8 * getAlignmentByte(unit, ty); > + uint32_t padding = getPadding(oldSlm*8, align); > + > + f.setSLMSize(oldSlm + padding/8 + getTypeByteSize(unit, ty)); > + const Value * parent = cast<Value>(&v); > + // local variable can only be used in one kernel function. so, don't > need to check its all uses. > + // loop through the Constant to find the instruction that use the > global variable > + do { > + Value::const_use_iterator it = parent->use_begin(); > + parent = cast<Value>(*it); > + } while(isa<Constant>(parent)); > + > + const Instruction * insn = cast<Instruction>(parent); > + const BasicBlock * bb = insn->getParent(); > + const Function * func = bb->getParent(); > + if(func != &F) continue; > + > + this->newRegister(const_cast<GlobalVariable*>(&v)); > + ir::Register reg = > regTranslator.getScalar(const_cast<GlobalVariable*>(&v), 0); > + ctx.LOADI(ir::TYPE_S32, reg, ctx.newIntegerImmediate(oldSlm + > padding/8, ir::TYPE_S32)); > + } else if(addrSpace == ir::MEM_CONSTANT) { > + GBE_ASSERT(v.hasInitializer()); > + this->newRegister(const_cast<GlobalVariable*>(&v)); > + ir::Register reg = > regTranslator.getScalar(const_cast<GlobalVariable*>(&v), 0); > + ir::Constant &con = unit.getConstantSet().getConstant(j ++); > + ctx.LOADI(ir::TYPE_S32, reg, > ctx.newIntegerImmediate(con.getOffset(), ir::TYPE_S32)); > + } else { > + GBE_ASSERT(0); > + } > + } > + > + } > + > void GenWriter::emitFunction(Function &F) > { > switch (F.getCallingConv()) { > @@ -1293,21 +1362,7 @@ 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()); > - ir::Register reg = ctx.reg(ir::RegisterFamily::FAMILY_DWORD); > - ir::Constant &con = unit.getConstantSet().getConstant(j ++); > - con.setReg(reg.value()); > - ctx.LOADI(ir::TYPE_S32, reg, ctx.newIntegerImmediate(con.getOffset(), > ir::TYPE_S32)); > - } > - > + this->allocateGlobalVariableRegister(F); > // 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/kernels/compiler_local_slm.cl b/kernels/compiler_local_slm.cl > index 1a4b175..52c078c 100644 > --- a/kernels/compiler_local_slm.cl > +++ b/kernels/compiler_local_slm.cl > @@ -1,10 +1,24 @@ > -#if 0 > -__kernel void compiler_local_slm(__global int *dst, __local int *hop) { > -#else > +struct Test{ > + char t0; > + int t1; > +}; > + > +constant int two= 2; > + > __kernel void compiler_local_slm(__global int *dst) { > - __local int hop[10]; > -#endif > - hop[get_global_id(0)] = get_local_id(1); > - dst[get_global_id(0)] = hop[get_local_id(0)]; > + __local int hop[16]; > + __local char a; > + __local struct Test c; > + > + c.t1 = get_group_id(0); > + a = two;// seems clang currently has a bug if I write 'a=2;' so currently > workaroud it. > + hop[get_local_id(0)] = get_local_id(0); > + barrier(CLK_LOCAL_MEM_FENCE); > + dst[get_global_id(0)] = hop[get_local_id(0)] + (int)a + hop[1] + c.t1; > } > > +__kernel void compiler_local_slm1(__global ulong *dst) { > + __local int hop[16]; > + dst[1] = (ulong)&hop[1]; > + dst[0] = (ulong)&hop[0]; > +} > diff --git a/src/cl_command_queue_gen7.c b/src/cl_command_queue_gen7.c > index b85c0cd..be7bcef 100644 > --- a/src/cl_command_queue_gen7.c > +++ b/src/cl_command_queue_gen7.c > @@ -200,7 +200,8 @@ cl_curbe_fill(cl_kernel ker, > } > /* Handle the various offsets to SLM */ > const int32_t arg_n = gbe_kernel_get_arg_num(ker->opaque); > - int32_t arg, slm_offset = 0; > + /* align so that we kernel argument get good alignment */ > + int32_t arg, slm_offset = ALIGN(gbe_kernel_get_slm_size(ker->opaque), 32); > for (arg = 0; arg < arg_n; ++arg) { > const enum gbe_arg_type type = gbe_kernel_get_arg_type(ker->opaque, arg); > if (type != GBE_ARG_LOCAL_PTR) > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt > index a24c490..daa4d6f 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -96,6 +96,7 @@ set (utests_sources > compiler_local_memory_barrier.cpp > compiler_local_memory_barrier_wg64.cpp > compiler_local_memory_barrier_2.cpp > + compiler_local_slm.cpp > compiler_movforphi_undef.cpp > compiler_volatile.cpp > compiler_copy_image1.cpp > diff --git a/utests/compiler_local_slm.cpp b/utests/compiler_local_slm.cpp > index aa9a2fe..48a072f 100644 > --- a/utests/compiler_local_slm.cpp > +++ b/utests/compiler_local_slm.cpp > @@ -2,9 +2,33 @@ > > void compiler_local_slm(void) > { > - // Setup kernel and buffers > - OCL_CREATE_KERNEL("compiler_local_slm"); > + const size_t n = 32; > + OCL_CREATE_KERNEL_FROM_FILE("compiler_local_slm", "compiler_local_slm"); > + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint32_t), NULL); > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + globals[0] = n; > + locals[0] = 16; > + OCL_NDRANGE(1); > + OCL_MAP_BUFFER(0); > + for (uint32_t i = 0; i < n; ++i) > +// std::cout << ((int32_t*)buf_data[0])[i] << std::endl; > + OCL_ASSERT(((int32_t*)buf_data[0])[i] == (i%16 + 2 + 1+ i/16)); > + OCL_UNMAP_BUFFER(0); > } > > +void compiler_local_slm1(void) > +{ > + const size_t n = 2; > + OCL_CREATE_KERNEL_FROM_FILE("compiler_local_slm", "compiler_local_slm1"); > + OCL_CREATE_BUFFER(buf[0], 0, n * sizeof(uint64_t), NULL); > + OCL_SET_ARG(0, sizeof(cl_mem), &buf[0]); > + globals[0] = 1; > + locals[0] = 1; > + OCL_NDRANGE(1); > + OCL_MAP_BUFFER(0); > + uint64_t * ptr = (uint64_t*)buf_data[0]; > + OCL_ASSERT((ptr[1] -ptr[0]) == 4); > + OCL_UNMAP_BUFFER(0); > +} > MAKE_UTEST_FROM_FUNCTION(compiler_local_slm); > - > +MAKE_UTEST_FROM_FUNCTION(compiler_local_slm1); > -- > 1.7.9.5 > > _______________________________________________ > Beignet mailing list > Beignet@lists.freedesktop.org > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list Beignet@lists.freedesktop.org http://lists.freedesktop.org/mailman/listinfo/beignet