[Beignet] [PATCH V2 2/6] Build the constant base address propagation registers.
Yang, Rong R
rong.r.yang at intel.com
Sun Apr 21 22:12:17 PDT 2013
Hi, Zhigang,
Yes, you are right. My previous implement is too complex. I have send out new patch as your proposal, It is much more clever. Thanks.
Yang Rong
-----Original Message-----
From: Zhigang Gong [mailto:zhigang.gong at linux.intel.com]
Sent: Friday, April 19, 2013 9:40 AM
To: Yang, Rong R; beignet at lists.freedesktop.org
Subject: RE: [Beignet] [PATCH V2 2/6] Build the constant base address propagation registers.
I have one major comment for the following patches. After review all of the code, I have a feeling that it's too hacky for me. You need to modify the binary instruction after the compilation.
Take the following cl kernel as an example, it seems that you ignore the value of the c0 itself, and you pass in the actual register offset by modifying the binary instruction on the fly. Right?
Why not just not ignore the value of c0/c1, and each time we seg argument, we may still need to reallocate curbe, but we don't need to modify the instruction, and just need to put the new offset to the c0 and c1. We may need to add one new gbe API to get offset of a constant region. And by this way, it will be much easier than you have done in this patchset. You need to analyze all the instructions and do the constant propagation and something really very complex for this function. And as you modify the instruction on the fly, we even can't get a valid ASM output for debugging.
__kernel void foo(__constant int *c0, __constant int *c1) {
}
Could you change to take my proposal and reimplement it? Or do I miss anything and that proposal has some big gap? Any thoughts?
> -----Original Message-----
> From:
> beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.org
> [mailto:beignet-bounces+zhigang.gong=linux.intel.com at lists.freedesktop.
> org] On Behalf Of Yang Rong
> Sent: Thursday, April 18, 2013 3:18 PM
> To: beignet at lists.freedesktop.org
> Cc: Yang Rong
> Subject: [Beignet] [PATCH V2 2/6] Build the constant base address
> propagation registers.
>
>
> Signed-off-by: Yang Rong <rong.r.yang at intel.com>
> ---
> backend/src/backend/context.cpp | 29
> +++++++++++++++++++++++++++++
> backend/src/backend/context.hpp | 8 ++++++++
> 2 files changed, 37 insertions(+)
>
> diff --git a/backend/src/backend/context.cpp
> b/backend/src/backend/context.cpp index c3ddb59..0a26b3d 100644
> --- a/backend/src/backend/context.cpp
> +++ b/backend/src/backend/context.cpp
> @@ -295,6 +295,7 @@ namespace gbe
> this->buildJIPs();
> this->buildStack();
> this->handleSLM();
> + this->buildConstBufs();
> if (this->emitCode() == false) {
> GBE_DELETE(this->kernel);
> this->kernel = NULL;
> @@ -548,6 +549,34 @@ namespace gbe
> kernel->useSLM = useSLM;
> }
>
> + void Context::buildConstBufs(void) {
> + uint32_t argNum = fn.argNum();
> + for (uint32_t argID = 0; argID < argNum; ++argID) {
> + const auto &arg = fn.getArg(argID);
> +
> + if(arg.type != ir::FunctionArgument::CONSTANT_POINTER)
> + continue;
> +
> + ir::Register srcReg = arg.reg;
> + constBaseProps.insert(std::make_pair(srcReg, srcReg));
> +
> + //build the const Propagation registers
> + fn.foreachInstruction([this](const ir::Instruction &insn) {
> + using namespace ir;
> + if(insn.getOpcode() == OP_LOAD) return;
> + const uint32_t srcNum = insn.getSrcNum();
> + for (uint32_t srcID = 0; srcID < srcNum; ++srcID) {
> + const ir::Register reg = insn.getSrc(srcID);
> + if((constBaseProps.contains(reg) != false)) {
> + const ir::Register dstReg = insn.getDst();
> + constBaseProps.insert(std::make_pair(dstReg,
> constBaseProps.find(reg)->second));
> + break;
> + }
> + }
> + });
> + }
> + }
> +
> bool Context::isScalarReg(const ir::Register ®) const {
> GBE_ASSERT(fn.getProfile() == ir::Profile::PROFILE_OCL);
> if (fn.getArg(reg) != NULL) return true; diff --git
> a/backend/src/backend/context.hpp b/backend/src/backend/context.hpp
> index 55a63a7..3b6e635 100644
> --- a/backend/src/backend/context.hpp
> +++ b/backend/src/backend/context.hpp
> @@ -82,6 +82,10 @@ namespace gbe
> INLINE bool hasJIP(const ir::Instruction *insn) const {
> return JIPs.find(insn) != JIPs.end();
> }
> + INLINE ir::Register getConstBaseReg(const ir::Register reg) const {
> + GBE_ASSERT(constBaseProps.find(reg) != constBaseProps.end());
> + return constBaseProps.find(reg)->second;
> + }
> /*! Allocate some memory in the register file */
> int16_t allocate(int16_t size, int16_t alignment);
> /*! Deallocate previously allocated memory */ @@ -105,12 +109,15
> @@ namespace gbe
> void buildJIPs(void);
> /*! Configure SLM use if needed */
> void handleSLM(void);
> + /* Build constant buffer map */
> + void buildConstBufs(void);
> /*! Insert a new entry with the given size in the Curbe. Return
> the offset
> * of the entry
> */
> void newCurbeEntry(gbe_curbe_type value, uint32_t subValue,
> uint32_t size, uint32_t alignment = 0);
> /*! Provide for each branch and label the label index target */
> typedef map<const ir::Instruction*, ir::LabelIndex> JIPMap;
> + typedef map<const ir::Register, ir::Register> ConstBasePropMap;
> const ir::Unit &unit; //!< Unit that contains the
> kernel
> const ir::Function &fn; //!< Function to compile
> std::string name; //!< Name of the kernel to
> compile
> @@ -120,6 +127,7 @@ namespace gbe
> RegisterFilePartitioner *partitioner; //!< Handle register file
> partionning
> set<ir::LabelIndex> usedLabels; //!< Set of all used labels
> JIPMap JIPs; //!< Where to jump all
> labels/branches
> + ConstBasePropMap constBaseProps; //!< const buffer reg
> propagation regs
> uint32_t simdWidth; //!< Number of lanes per
> HW threads
> GBE_CLASS(Context); //!< Use custom allocators
> };
> --
> 1.7.9.5
>
> _______________________________________________
> Beignet mailing list
> Beignet at lists.freedesktop.org
> http://lists.freedesktop.org/mailman/listinfo/beignet
More information about the Beignet
mailing list