Method operator()

src/backend/oneapi/kernel/transpose.hpp:66–124

        , IS32MULTIPLE_(IS32MULTIPLE)
        , shrdMem_(shrdMem) {}
    void operator()(sycl::nd_item<2> it) const {
        const int shrdStride = TILE_DIM + 1;

        const int oDim0 = out_.dims[0];
        const int oDim1 = out_.dims[1];
        const int iDim0 = in_.dims[0];
        const int iDim1 = in_.dims[1];

        // calculate strides
        const int oStride1 = out_.strides[1];
        const int iStride1 = in_.strides[1];

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

        // batch based block Id
        sycl::group g        = it.get_group();
        const int batchId_x  = g.get_group_id(0) / blocksPerMatX_;
        const int blockIdx_x = (g.get_group_id(0) - batchId_x * blocksPerMatX_);

        const int batchId_y  = g.get_group_id(1) / blocksPerMatY_;
        const int blockIdx_y = (g.get_group_id(1) - batchId_y * blocksPerMatY_);

        const int x0 = TILE_DIM * blockIdx_x;
        const int y0 = TILE_DIM * blockIdx_y;

        // calculate global in_dices
        int gx = lx + x0;
        int gy = ly + y0;

        // offset in_ and out_ based on batch id
        // also add the subBuffer offsets
        const T *iDataPtr = iData_.get_pointer();
        T *oDataPtr       = oData_.get_pointer();
        iDataPtr += batchId_x * in_.strides[2] + batchId_y * in_.strides[3] +
                    in_.offset;
        oDataPtr += batchId_x * out_.strides[2] + batchId_y * out_.strides[3] +
                    out_.offset;

        for (int repeat = 0; repeat < TILE_DIM; repeat += THREADS_Y) {
            int gy_ = gy + repeat;
            if (IS32MULTIPLE_ || (gx < iDim0 && gy_ < iDim1))
                shrdMem_[(ly + repeat) * shrdStride + lx] =
                    iDataPtr[gy_ * iStride1 + gx];
        }
        it.barrier();

        gx = lx + y0;
        gy = ly + x0;

        for (int repeat = 0; repeat < TILE_DIM; repeat += THREADS_Y) {
            int gy_ = gy + repeat;
            if (IS32MULTIPLE_ || (gx < oDim0 && gy_ < oDim1)) {
                const T val = shrdMem_[lx * shrdStride + ly + repeat];
                oDataPtr[gy_ * oStride1 + gx] =
                    conjugate_ ? getConjugate(val) : val;
            }
        }