Method operator()

src/backend/oneapi/kernel/lookup.hpp:55–88

        , nBBS1_(nBBS1)
        , DIM_(DIM) {}
    void operator()(sycl::nd_item<2> it) const {
        sycl::group g = it.get_group();

        int lx = it.get_local_id(0);
        int ly = it.get_local_id(1);

        int gz = g.get_group_id(0) / nBBS0_;
        int gw = g.get_group_id(1) / nBBS1_;

        int gx = g.get_local_range(0) * (g.get_group_id(0) - gz * nBBS0_) + lx;
        int gy = g.get_local_range(1) * (g.get_group_id(1) - gw * nBBS1_) + ly;

        const idx_t *idxPtr = indices_.get_pointer() + idxInfo_.offset;

        int i = iInfo_.strides[0] *
                (DIM_ == 0 ? trimIndex((int)idxPtr[gx], iInfo_.dims[0]) : gx);
        int j = iInfo_.strides[1] *
                (DIM_ == 1 ? trimIndex((int)idxPtr[gy], iInfo_.dims[1]) : gy);
        int k = iInfo_.strides[2] *
                (DIM_ == 2 ? trimIndex((int)idxPtr[gz], iInfo_.dims[2]) : gz);
        int l = iInfo_.strides[3] *
                (DIM_ == 3 ? trimIndex((int)idxPtr[gw], iInfo_.dims[3]) : gw);

        const in_t *inPtr = in_.get_pointer() + (i + j + k + l) + iInfo_.offset;
        in_t *outPtr =
            out_.get_pointer() +
            (gx * oInfo_.strides[0] + gy * oInfo_.strides[1] +
             gz * oInfo_.strides[2] + gw * oInfo_.strides[3] + oInfo_.offset);

        if (gx < oInfo_.dims[0] && gy < oInfo_.dims[1] && gz < oInfo_.dims[2] &&
            gw < oInfo_.dims[3]) {
            outPtr[0] = inPtr[0];
        }
    }

   private:
    write_accessor<in_t> out_;