Method operator()

src/backend/oneapi/kernel/scan_first.hpp:52–123

        , s_tmp_(s_tmp) {}

    void operator()(sycl::nd_item<2> it) const {
        sycl::group g   = it.get_group();
        const uint lidx = it.get_local_id(0);
        const uint lidy = it.get_local_id(1);

        const uint zid       = g.get_group_id(0) / groups_x_;
        const uint wid       = g.get_group_id(1) / groups_y_;
        const uint groupId_x = g.get_group_id(0) - (groups_x_)*zid;
        const uint groupId_y = g.get_group_id(1) - (groups_y_)*wid;
        const uint xid       = groupId_x * g.get_local_range(0) * lim_ + lidx;
        const uint yid       = groupId_y * g.get_local_range(1) + lidy;

        bool cond_yzw = (yid < oInfo_.dims[1]) && (zid < oInfo_.dims[2]) &&
                        (wid < oInfo_.dims[3]);

        // if (!cond_yzw) return;  // retire warps early TODO: move

        const Ti *iptr = in_acc_.get_pointer();
        To *optr       = out_acc_.get_pointer();
        To *tptr       = tmp_acc_.get_pointer();

        iptr += wid * iInfo_.strides[3] + zid * iInfo_.strides[2] +
                yid * iInfo_.strides[1] + iInfo_.offset;
        optr += wid * oInfo_.strides[3] + zid * oInfo_.strides[2] +
                yid * oInfo_.strides[1];
        tptr += wid * tInfo_.strides[3] + zid * tInfo_.strides[2] +
                yid * tInfo_.strides[1];

        To *sptr = s_val_.get_pointer() + lidy * (2 * DIMX_ + 1);

        common::Transform<Ti, To, op> transform;
        common::Binary<To, op> binop;

        const To init = common::Binary<To, op>::init();
        int id        = xid;
        To val        = init;

        const bool isLast = (lidx == (DIMX_ - 1));
        for (int k = 0; k < lim_; k++) {
            if (isLast) s_tmp_[lidy] = val;

            bool cond  = (id < oInfo_.dims[0]) && cond_yzw;
            val        = cond ? transform(iptr[id]) : init;
            sptr[lidx] = val;
            group_barrier(g);

            int start = 0;
            for (int off = 1; off < DIMX_; off *= 2) {
                if (lidx >= off) val = binop(val, sptr[(start - off) + lidx]);
                start              = DIMX_ - start;
                sptr[start + lidx] = val;

                group_barrier(g);
            }

            val = binop(val, s_tmp_[lidy]);

            if (inclusive_scan_) {