Method operator()

src/backend/oneapi/kernel/assign.hpp:54–103

        , nBBS1_(nBBS1) {}

    void operator()(sycl::nd_item<2> it) const {
        // retrive booleans that tell us which index to use
        const bool s0 = p_.isSeq[0];
        const bool s1 = p_.isSeq[1];
        const bool s2 = p_.isSeq[2];
        const bool s3 = p_.isSeq[3];

        sycl::group g = it.get_group();
        const int gz  = g.get_group_id(0) / nBBS0_;
        const int gw  = g.get_group_id(1) / nBBS1_;
        const int gx =
            g.get_local_range(0) * (g.get_group_id(0) - gz * nBBS0_) +
            it.get_local_id(0);
        const int gy =
            g.get_local_range(1) * (g.get_group_id(1) - gw * nBBS1_) +
            it.get_local_id(1);

        size_t idims0 = iInfo_.dims[0];
        size_t idims1 = iInfo_.dims[1];
        size_t idims2 = iInfo_.dims[2];
        size_t idims3 = iInfo_.dims[3];

        if (gx < idims0 && gy < idims1 && gz < idims2 && gw < idims3) {
            // calculate pointer offsets for input
            int i =
                p_.strds[0] *
                trimIndex(s0 ? gx + p_.offs[0] : p_.ptr[0][gx], oInfo_.dims[0]);
            int j =
                p_.strds[1] *
                trimIndex(s1 ? gy + p_.offs[1] : p_.ptr[1][gy], oInfo_.dims[1]);
            int k =
                p_.strds[2] *
                trimIndex(s2 ? gz + p_.offs[2] : p_.ptr[2][gz], oInfo_.dims[2]);
            int l =
                p_.strds[3] *
                trimIndex(s3 ? gw + p_.offs[3] : p_.ptr[3][gw], oInfo_.dims[3]);

            const T* iptr = in_.get_pointer();
            // offset input and output pointers
            const T* src =
                iptr + (gx * iInfo_.strides[0] + gy * iInfo_.strides[1] +
                        gz * iInfo_.strides[2] + gw * iInfo_.strides[3] +
                        iInfo_.offset);

            T* optr = out_.get_pointer();
            T* dst  = optr + (i + j + k + l) + oInfo_.offset;
            // set the output
            dst[0] = src[0];
        }
    }

   protected:
    write_accessor<T> out_;