11145093866

Committed 02 Oct 2024 01:49PM UTC coverage: 59.477% (-0.01%) from 59.488%

Build # 11145093866

Build Type

Pull #2082

github

Committed by

web-flow

Commit Message

Merge 4e93ab4e9 into 9b37b0209

Pull Request Pull Request #2082: Implement a workaround to gemm issue in OneMKL

Run Details

3757 of 9438 branches covered (39.81%)

Branch coverage included in aggregate %.

18 of 33 new or added lines in 3 files covered. (54.55%)

1 existing line in 1 file now uncovered.

13772 of 20034 relevant lines covered (68.74%)

16823.87 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

69.51

/dpnp/backend/extensions/blas/gemm.cpp

//*****************************************************************************
// Copyright (c) 2024, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
// - Redistributions of source code must retain the above copyright notice,
//   this list of conditions and the following disclaimer.
// - Redistributions in binary form must reproduce the above copyright notice,
//   this list of conditions and the following disclaimer in the documentation
//   and/or other materials provided with the distribution.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
// THE POSSIBILITY OF SUCH DAMAGE.
//*****************************************************************************

#include <pybind11/pybind11.h>

// dpctl tensor headers
#include "utils/memory_overlap.hpp"
#include "utils/output_validation.hpp"
#include "utils/type_utils.hpp"

#include "gemm.hpp"
#include "types_matrix.hpp"

#include "dpnp_utils.hpp"

namespace dpnp::extensions::blas
{
namespace mkl_blas = oneapi::mkl::blas;
namespace py = pybind11;
namespace type_utils = dpctl::tensor::type_utils;

typedef sycl::event (*gemm_impl_fn_ptr_t)(sycl::queue &,
                                          oneapi::mkl::transpose,
                                          oneapi::mkl::transpose,
                                          const std::int64_t,
                                          const std::int64_t,
                                          const std::int64_t,
                                          const char *,
                                          const std::int64_t,
                                          const char *,
                                          const std::int64_t,
                                          char *,
                                          const std::int64_t,
                                          const bool,
                                          const std::vector<sycl::event> &);

static gemm_impl_fn_ptr_t gemm_dispatch_table[dpctl_td_ns::num_types]
                                             [dpctl_td_ns::num_types];

template <typename Tab, typename Tc>
static sycl::event gemm_impl(sycl::queue &exec_q,
                             oneapi::mkl::transpose transA,
                             oneapi::mkl::transpose transB,
                             const std::int64_t m,
                             const std::int64_t n,
                             const std::int64_t k,
                             const char *matrixA,
                             const std::int64_t lda,
                             const char *matrixB,
                             const std::int64_t ldb,
                             char *resultC,
                             const std::int64_t ldc,
                             const bool is_row_major,
                             const std::vector<sycl::event> &depends)
{
    type_utils::validate_type_for_device<Tab>(exec_q);
    type_utils::validate_type_for_device<Tc>(exec_q);

    const Tab *a = reinterpret_cast<const Tab *>(matrixA);
    const Tab *b = reinterpret_cast<const Tab *>(matrixB);
    Tc *res = reinterpret_cast<Tc *>(resultC);

    std::stringstream error_msg;
    bool is_exception_caught = false;

    sycl::event gemm_event;
    try {
        auto gemm_func =
            [&](sycl::queue &q, oneapi::mkl::transpose transA,
                oneapi::mkl::transpose transB, const std::int64_t m,
                const std::int64_t n, const std::int64_t k, Tab alpha,
                const Tab *a, const std::int64_t lda, const Tab *b,
                const std::int64_t ldb, Tab beta, Tc *c, const std::int64_t ldc,
                const std::vector<sycl::event> &deps) -> sycl::event {
#if defined(USE_ONEMKL_CUBLAS)
            return mkl_blas::column_major::gemm(q, transA, transB, m, n, k,
                                                alpha, a, lda, b, ldb, beta, c,
                                                ldc, deps);
#else
            if (is_row_major) {
                return mkl_blas::row_major::gemm(q, transA, transB, m, n, k,
                                                 alpha, a, lda, b, ldb, beta, c,
                                                 ldc, deps);
            }
            else {
                return mkl_blas::column_major::gemm(q, transA, transB, m, n, k,
                                                    alpha, a, lda, b, ldb, beta,
                                                    c, ldc, deps);
            }
#endif // USE_ONEMKL_CUBLAS
        };
        gemm_event = gemm_func(
            exec_q,
            transA, // Defines the transpose operation for matrix A:
                    // 'N' indicates no transpose, 'T' for transpose,
                    // or 'C' for a conjugate transpose.
            transB, // Same as transA but for matrix B.
            m,      // Number of rows in matrices A and C.
            n,      // Number of columns in matrices B and C.
            k,      // Number of columns in matrix A and rows in matrix B.
            Tab(1), // Scaling factor for the product of matrices A and B.
            a,      // Pointer to matrix A.
            lda,    // Leading dimension of matrix A, which is the
                    // stride between successive rows (for row major
                    // layout).
            b,      // Pointer to matrix B.
            ldb,    // Leading dimension of matrix B, similar to lda.
            Tab(0), // Scaling factor for matrix C.
            res,    // Pointer to matrix C, where the result is stored.
            ldc,    // Leading dimension of matrix C.
            depends);
    } catch (oneapi::mkl::exception const &e) {
        error_msg
            << "Unexpected MKL exception caught during gemm() call:\nreason: "
            << e.what();
        is_exception_caught = true;
    } catch (sycl::exception const &e) {
        error_msg << "Unexpected SYCL exception caught during gemm() call:\n"
                  << e.what();
        is_exception_caught = true;
    }

    if (is_exception_caught) // an unexpected error occurs
    {
        throw std::runtime_error(error_msg.str());
    }

    return gemm_event;
}

std::tuple<sycl::event, sycl::event, bool>
    gemm(sycl::queue &exec_q,
         const dpctl::tensor::usm_ndarray &matrixA,
         const dpctl::tensor::usm_ndarray &matrixB,
         const dpctl::tensor::usm_ndarray &resultC,
         const std::vector<sycl::event> &depends)
{
    const int matrixA_nd = matrixA.get_ndim();
    const int matrixB_nd = matrixB.get_ndim();
    const int resultC_nd = resultC.get_ndim();

    if ((matrixA_nd != 2) || (matrixB_nd != 2) || (resultC_nd != 2)) {
        throw py::value_error("Input matrices must be two-dimensional.");
    }

    auto const &overlap = dpctl::tensor::overlap::MemoryOverlap();
    if (overlap(matrixA, resultC)) {
        throw py::value_error(
            "The first input array and output array are overlapping "
            "segments of memory");
    }
    if (overlap(matrixB, resultC)) {
        throw py::value_error(
            "The second input array and output array are overlapping "
            "segments of memory");
    }

    if (!dpctl::utils::queues_are_compatible(
            exec_q,
            {matrixA.get_queue(), matrixB.get_queue(), resultC.get_queue()}))
    {
        throw py::value_error(
            "USM allocations are not compatible with the execution queue.");
    }

    const py::ssize_t *a_shape = matrixA.get_shape_raw();
    const py::ssize_t *b_shape = matrixB.get_shape_raw();
    const py::ssize_t *c_shape = resultC.get_shape_raw();
    const std::int64_t m = a_shape[0];
    const std::int64_t n = b_shape[1];
    const std::int64_t k = a_shape[1];
    if (a_shape[1] != b_shape[0]) {
        throw py::value_error("The number of columns in A must be equal to "
                              "the number of rows in B.");
    }
    if (a_shape[0] != c_shape[0]) {
        throw py::value_error("The number of rows in A must be equal to "
                              "the number of rows in result array.");
    }
    if (b_shape[1] != c_shape[1]) {
        throw py::value_error("The number of columns in B must be equal to "
                              "the number of columns in result array.");
    }

    const std::size_t src_nelems = m * n;
    dpctl::tensor::validation::CheckWritable::throw_if_not_writable(resultC);
    dpctl::tensor::validation::AmpleMemory::throw_if_not_ample(resultC,
                                                               src_nelems);

    const bool is_matrixA_f_contig = matrixA.is_f_contiguous();
    const bool is_matrixB_f_contig = matrixB.is_f_contiguous();
    const bool is_resultC_f_contig = resultC.is_f_contiguous();
    const bool is_matrixA_c_contig = matrixA.is_c_contiguous();
    const bool is_matrixB_c_contig = matrixB.is_c_contiguous();
    const bool is_resultC_c_contig = resultC.is_c_contiguous();

    if (!is_matrixA_f_contig and !is_matrixA_c_contig) {
        throw py::value_error(
            "The first input array is not c-contiguous nor f-contiguous.");
    }
    if (!is_matrixB_f_contig and !is_matrixB_c_contig) {
        throw py::value_error(
            "The second input array is not c-contiguous nor f-contiguous.");
    }
    if (!is_resultC_f_contig and !is_resultC_c_contig) {
        throw py::value_error(
            "Result array is not c-contiguous nor f-contiguous.");
    }

    oneapi::mkl::transpose transA;
    oneapi::mkl::transpose transB;
    std::int64_t lda;
    std::int64_t ldb;

#if defined(USE_ONEMKL_CUBLAS)
    const bool is_row_major = false;

    transA = is_matrixA_c_contig ? oneapi::mkl::transpose::T
                                 : oneapi::mkl::transpose::N;
    transB = is_matrixB_c_contig ? oneapi::mkl::transpose::T
                                 : oneapi::mkl::transpose::N;

    if (transA == oneapi::mkl::transpose::N) {
        lda = m;
    }
    else {
        lda = k;
    }
    if (transB == oneapi::mkl::transpose::N) {
        ldb = k;
    }
    else {
        ldb = n;
    }
#else
    bool is_row_major = true;
    if (is_matrixA_f_contig && is_matrixB_f_contig) {
        is_row_major = false;
    }

    if (is_row_major) {
        transA = is_matrixA_f_contig ? oneapi::mkl::transpose::T
                                     : oneapi::mkl::transpose::N;
        transB = is_matrixB_f_contig ? oneapi::mkl::transpose::T
                                     : oneapi::mkl::transpose::N;
        if (transA == oneapi::mkl::transpose::N) {
            lda = k;
        }
        else {
            lda = m;
        }
        if (transB == oneapi::mkl::transpose::N) {
            ldb = n;
        }
        else {
            ldb = k;
        }
    }
    else {
        transA = oneapi::mkl::transpose::N;
        transB = oneapi::mkl::transpose::N;
        lda = m;
        ldb = k;
    }
#endif // USE_ONEMKL_CUBLAS

    const std::int64_t ldc = is_row_major ? n : m;

    const int matrixA_typenum = matrixA.get_typenum();
    const int matrixB_typenum = matrixB.get_typenum();
    const int resultC_typenum = resultC.get_typenum();

    if (matrixA_typenum != matrixB_typenum) {
        throw py::value_error("matrixA and matrixB must be of the same type.");
    }

    auto array_types = dpctl_td_ns::usm_ndarray_types();
    const int matrixAB_type_id =
        array_types.typenum_to_lookup_id(matrixA_typenum);
    const int resultC_type_id =
        array_types.typenum_to_lookup_id(resultC_typenum);

    gemm_impl_fn_ptr_t gemm_fn =
        gemm_dispatch_table[matrixAB_type_id][resultC_type_id];
    if (gemm_fn == nullptr) {
        throw py::value_error(
            "Types of input matrices and result matrix are mismatched.");
    }

    const char *a_typeless_ptr = matrixA.get_data();
    const char *b_typeless_ptr = matrixB.get_data();
    char *r_typeless_ptr = resultC.get_data();

    sycl::event gemm_ev = gemm_fn(exec_q, transA, transB, m, n, k,
                                  a_typeless_ptr, lda, b_typeless_ptr, ldb,
                                  r_typeless_ptr, ldc, is_row_major, depends);

    sycl::event args_ev = dpctl::utils::keep_args_alive(
        exec_q, {matrixA, matrixB, resultC}, {gemm_ev});

    return std::make_tuple(args_ev, gemm_ev, is_row_major);
}

bool _is_lnl_arl_architecture(sycl::device &dev)
{
#if !defined(USE_ONEMKL_CUBLAS)
    if (dev.ext_oneapi_architecture_is(
            sycl::ext::oneapi::experimental::architecture::
                intel_gpu_lnl_m)) /* Lunar Lake */
    {
        std::cout << "_is_lnl_arl_architecture: running on Lunar Lake" << std::endl;
        return true;
    }
    if (dev.ext_oneapi_architecture_is(
            sycl::ext::oneapi::experimental::architecture::
                intel_gpu_bmg_g21)) /* Battlemage G21 */
    {
        std::cout << "_is_lnl_arl_architecture: running on Battlemage G21" << std::endl;
        return true;
    }
    // else if (dev.ext_oneapi_architecture_is(
    //              sycl::ext::oneapi::experimental::architecture::
    //                  intel_gpu_arl_h)) /* Arrow Lake H */
    // {
    //     std::cout << "_is_lnl_arl_architecture: running on Arrow Lake H" << std::endl;
    //     return true;
    // }
    // else if (dev.ext_oneapi_architecture_is(
    //              sycl::ext::oneapi::experimental::architecture::
    //                  intel_gpu_arl_u)) /* Arrow Lake U */
    // {
    //     std::cout << "_is_lnl_arl_architecture: running on Arrow Lake U" << std::endl;
    //     return true;
    // }
    // else if (dev.ext_oneapi_architecture_is(
    //              sycl::ext::oneapi::experimental::architecture::
    //                  intel_gpu_arl_s)) /* Arrow Lake S */
    // {
    //     std::cout << "_is_lnl_arl_architecture: running on Arrow Lake S" << std::endl;
    //     return true;
    // }
#endif // !defined(USE_ONEMKL_CUBLAS)
    return false;
}

template <typename fnT, typename Tab, typename Tc>
struct GemmContigFactory
{
    fnT get()
    {
        if constexpr (types::GemmTypePairSupportFactory<Tab, Tc>::is_defined) {
            return gemm_impl<Tab, Tc>;
        }
        else {
            return nullptr;
        }
    }
};

void init_gemm_dispatch_table(void)
{
    dpctl_td_ns::DispatchTableBuilder<gemm_impl_fn_ptr_t, GemmContigFactory,
                                      dpctl_td_ns::num_types>
        contig;
    contig.populate_dispatch_table(gemm_dispatch_table);
}
} // namespace dpnp::extensions::blas

1	//*****************************************************************************
2	// Copyright (c) 2024, Intel Corporation
3	// All rights reserved.
4	//
5	// Redistribution and use in source and binary forms, with or without
6	// modification, are permitted provided that the following conditions are met:
7	// - Redistributions of source code must retain the above copyright notice,
8	// this list of conditions and the following disclaimer.
9	// - Redistributions in binary form must reproduce the above copyright notice,
10	// this list of conditions and the following disclaimer in the documentation
11	// and/or other materials provided with the distribution.
12	//
13	// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14	// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15	// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16	// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17	// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18	// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19	// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20	// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21	// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22	// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23	// THE POSSIBILITY OF SUCH DAMAGE.
24	//*****************************************************************************
25
26	#include <pybind11/pybind11.h>
27
28	// dpctl tensor headers
29	#include "utils/memory_overlap.hpp"
30	#include "utils/output_validation.hpp"
31	#include "utils/type_utils.hpp"
32
33	#include "gemm.hpp"
34	#include "types_matrix.hpp"
35
36	#include "dpnp_utils.hpp"
37
38	namespace dpnp::extensions::blas
39	{
40	namespace mkl_blas = oneapi::mkl::blas;
41	namespace py = pybind11;
42	namespace type_utils = dpctl::tensor::type_utils;
43
44	typedef sycl::event (*gemm_impl_fn_ptr_t)(sycl::queue &,
45	oneapi::mkl::transpose,
46	oneapi::mkl::transpose,
47	const std::int64_t,
48	const std::int64_t,
49	const std::int64_t,
50	const char *,
51	const std::int64_t,
52	const char *,
53	const std::int64_t,
54	char *,
55	const std::int64_t,
56	const bool,
57	const std::vector<sycl::event> &);
58
59	static gemm_impl_fn_ptr_t gemm_dispatch_table[dpctl_td_ns::num_types]
60	[dpctl_td_ns::num_types];
61
62	template <typename Tab, typename Tc>
63	static sycl::event gemm_impl(sycl::queue &exec_q,
64	oneapi::mkl::transpose transA,
65	oneapi::mkl::transpose transB,
66	const std::int64_t m,
67	const std::int64_t n,
68	const std::int64_t k,
69	const char *matrixA,
70	const std::int64_t lda,
71	const char *matrixB,
72	const std::int64_t ldb,
73	char *resultC,
74	const std::int64_t ldc,
75	const bool is_row_major,
76	const std::vector<sycl::event> &depends)
77	{	5,559✔
78	type_utils::validate_type_for_device<Tab>(exec_q);	5,559✔
79	type_utils::validate_type_for_device<Tc>(exec_q);	5,559✔
80
81	const Tab a = reinterpret_cast<const Tab >(matrixA);	5,559✔
82	const Tab b = reinterpret_cast<const Tab >(matrixB);	5,559✔
83	Tc res = reinterpret_cast<Tc >(resultC);	5,559✔
84
85	std::stringstream error_msg;	5,559✔
86	bool is_exception_caught = false;	5,559✔
87
88	sycl::event gemm_event;	5,559✔
89	try {	5,559✔
90	auto gemm_func =	5,559✔
91	[&](sycl::queue &q, oneapi::mkl::transpose transA,	5,559✔
92	oneapi::mkl::transpose transB, const std::int64_t m,	5,559✔
93	const std::int64_t n, const std::int64_t k, Tab alpha,	5,559✔
94	const Tab a, const std::int64_t lda, const Tab b,	5,559✔
95	const std::int64_t ldb, Tab beta, Tc *c, const std::int64_t ldc,	5,559✔
96	const std::vector<sycl::event> &deps) -> sycl::event {	5,559✔
97	#if defined(USE_ONEMKL_CUBLAS)
98	return mkl_blas::column_major::gemm(q, transA, transB, m, n, k,
99	alpha, a, lda, b, ldb, beta, c,
100	ldc, deps);
101	#else
102	if (is_row_major) {	5,559!
103	return mkl_blas::row_major::gemm(q, transA, transB, m, n, k,	4,130✔
104	alpha, a, lda, b, ldb, beta, c,	4,130✔
105	ldc, deps);	4,130✔
106	}	4,130✔
107	else {	1,429✔
108	return mkl_blas::column_major::gemm(q, transA, transB, m, n, k,	1,429✔
109	alpha, a, lda, b, ldb, beta,	1,429✔
110	c, ldc, deps);	1,429✔
111	}	1,429✔
112	#endif // USE_ONEMKL_CUBLAS	5,559✔
113	};	5,559✔
114	gemm_event = gemm_func(	5,559✔
115	exec_q,	5,559✔
116	transA, // Defines the transpose operation for matrix A:	5,559✔
117	// 'N' indicates no transpose, 'T' for transpose,
118	// or 'C' for a conjugate transpose.
119	transB, // Same as transA but for matrix B.	5,559✔
120	m, // Number of rows in matrices A and C.	5,559✔
121	n, // Number of columns in matrices B and C.	5,559✔
122	k, // Number of columns in matrix A and rows in matrix B.	5,559✔
123	Tab(1), // Scaling factor for the product of matrices A and B.	5,559✔
124	a, // Pointer to matrix A.	5,559✔
125	lda, // Leading dimension of matrix A, which is the	5,559✔
126	// stride between successive rows (for row major
127	// layout).
128	b, // Pointer to matrix B.	5,559✔
129	ldb, // Leading dimension of matrix B, similar to lda.	5,559✔
130	Tab(0), // Scaling factor for matrix C.	5,559✔
131	res, // Pointer to matrix C, where the result is stored.	5,559✔
132	ldc, // Leading dimension of matrix C.	5,559✔
133	depends);	5,559✔
134	} catch (oneapi::mkl::exception const &e) {	5,559✔
135	error_msg	×
136	<< "Unexpected MKL exception caught during gemm() call:\nreason: "	×
137	<< e.what();	×
138	is_exception_caught = true;	×
139	} catch (sycl::exception const &e) {	×
140	error_msg << "Unexpected SYCL exception caught during gemm() call:\n"	×
141	<< e.what();	×
142	is_exception_caught = true;	×
143	}	×
144
145	if (is_exception_caught) // an unexpected error occurs	5,559!
146	{	×
147	throw std::runtime_error(error_msg.str());	×
148	}	×
149
150	return gemm_event;	5,559✔
151	}	5,559✔
152
153	std::tuple<sycl::event, sycl::event, bool>
154	gemm(sycl::queue &exec_q,
155	const dpctl::tensor::usm_ndarray &matrixA,
156	const dpctl::tensor::usm_ndarray &matrixB,
157	const dpctl::tensor::usm_ndarray &resultC,
158	const std::vector<sycl::event> &depends)
159	{	5,559✔
160	const int matrixA_nd = matrixA.get_ndim();	5,559✔
161	const int matrixB_nd = matrixB.get_ndim();	5,559✔
162	const int resultC_nd = resultC.get_ndim();	5,559✔
163
164	if ((matrixA_nd != 2) \|\| (matrixB_nd != 2) \|\| (resultC_nd != 2)) {	5,559!
165	throw py::value_error("Input matrices must be two-dimensional.");	×
166	}	×
167
168	auto const &overlap = dpctl::tensor::overlap::MemoryOverlap();	5,559✔
169	if (overlap(matrixA, resultC)) {	5,559!
170	throw py::value_error(	×
171	"The first input array and output array are overlapping "	×
172	"segments of memory");	×
173	}	×
174	if (overlap(matrixB, resultC)) {	5,559!
175	throw py::value_error(	×
176	"The second input array and output array are overlapping "	×
177	"segments of memory");	×
178	}	×
179
180	if (!dpctl::utils::queues_are_compatible(	5,559!
181	exec_q,	5,559✔
182	{matrixA.get_queue(), matrixB.get_queue(), resultC.get_queue()}))	5,559✔
183	{	×
184	throw py::value_error(	×
185	"USM allocations are not compatible with the execution queue.");	×
186	}	×
187
188	const py::ssize_t *a_shape = matrixA.get_shape_raw();	5,559✔
189	const py::ssize_t *b_shape = matrixB.get_shape_raw();	5,559✔
190	const py::ssize_t *c_shape = resultC.get_shape_raw();	5,559✔
191	const std::int64_t m = a_shape[0];	5,559✔
192	const std::int64_t n = b_shape[1];	5,559✔
193	const std::int64_t k = a_shape[1];	5,559✔
194	if (a_shape[1] != b_shape[0]) {	5,559!
195	throw py::value_error("The number of columns in A must be equal to "	×
196	"the number of rows in B.");	×
197	}	×
198	if (a_shape[0] != c_shape[0]) {	5,559!
199	throw py::value_error("The number of rows in A must be equal to "	×
200	"the number of rows in result array.");	×
201	}	×
202	if (b_shape[1] != c_shape[1]) {	5,559!
203	throw py::value_error("The number of columns in B must be equal to "	×
204	"the number of columns in result array.");	×
205	}	×
206
207	const std::size_t src_nelems = m * n;	5,559✔
208	dpctl::tensor::validation::CheckWritable::throw_if_not_writable(resultC);	5,559✔
209	dpctl::tensor::validation::AmpleMemory::throw_if_not_ample(resultC,	5,559✔
210	src_nelems);	5,559✔
211
212	const bool is_matrixA_f_contig = matrixA.is_f_contiguous();	5,559✔
213	const bool is_matrixB_f_contig = matrixB.is_f_contiguous();	5,559✔
214	const bool is_resultC_f_contig = resultC.is_f_contiguous();	5,559✔
215	const bool is_matrixA_c_contig = matrixA.is_c_contiguous();	5,559✔
216	const bool is_matrixB_c_contig = matrixB.is_c_contiguous();	5,559✔
217	const bool is_resultC_c_contig = resultC.is_c_contiguous();	5,559✔
218
219	if (!is_matrixA_f_contig and !is_matrixA_c_contig) {	5,559!
220	throw py::value_error(	×
221	"The first input array is not c-contiguous nor f-contiguous.");	×
222	}	×
223	if (!is_matrixB_f_contig and !is_matrixB_c_contig) {	5,559!
224	throw py::value_error(	×
225	"The second input array is not c-contiguous nor f-contiguous.");	×
226	}	×
227	if (!is_resultC_f_contig and !is_resultC_c_contig) {	5,559!
228	throw py::value_error(	×
229	"Result array is not c-contiguous nor f-contiguous.");	×
230	}	×
231
232	oneapi::mkl::transpose transA;	5,559✔
233	oneapi::mkl::transpose transB;	5,559✔
234	std::int64_t lda;	5,559✔
235	std::int64_t ldb;	5,559✔
236
237	#if defined(USE_ONEMKL_CUBLAS)
238	const bool is_row_major = false;
239
240	transA = is_matrixA_c_contig ? oneapi::mkl::transpose::T
241	: oneapi::mkl::transpose::N;
242	transB = is_matrixB_c_contig ? oneapi::mkl::transpose::T
243	: oneapi::mkl::transpose::N;
244
245	if (transA == oneapi::mkl::transpose::N) {
246	lda = m;
247	}
248	else {
249	lda = k;
250	}
251	if (transB == oneapi::mkl::transpose::N) {
252	ldb = k;
253	}
254	else {
255	ldb = n;
256	}
257	#else
258	bool is_row_major = true;	5,559✔
259	if (is_matrixA_f_contig && is_matrixB_f_contig) {	5,559✔
260	is_row_major = false;	1,429✔
261	}	1,429✔
262
263	if (is_row_major) {	5,559✔
264	transA = is_matrixA_f_contig ? oneapi::mkl::transpose::T	4,130✔
265	: oneapi::mkl::transpose::N;	4,130✔
266	transB = is_matrixB_f_contig ? oneapi::mkl::transpose::T	4,130✔
267	: oneapi::mkl::transpose::N;	4,130✔
268	if (transA == oneapi::mkl::transpose::N) {	4,130✔
269	lda = k;	3,405✔
270	}	3,405✔
271	else {	725✔
272	lda = m;	725✔
273	}	725✔
274	if (transB == oneapi::mkl::transpose::N) {	4,130✔
275	ldb = n;	3,659✔
276	}	3,659✔
277	else {	471✔
278	ldb = k;	471✔
279	}	471✔
280	}	4,130✔
281	else {	1,429✔
282	transA = oneapi::mkl::transpose::N;	1,429✔
283	transB = oneapi::mkl::transpose::N;	1,429✔
284	lda = m;	1,429✔
285	ldb = k;	1,429✔
286	}	1,429✔
287	#endif // USE_ONEMKL_CUBLAS	5,559✔
288
289	const std::int64_t ldc = is_row_major ? n : m;	5,559✔
290
291	const int matrixA_typenum = matrixA.get_typenum();	5,559✔
292	const int matrixB_typenum = matrixB.get_typenum();	5,559✔
293	const int resultC_typenum = resultC.get_typenum();	5,559✔
294
295	if (matrixA_typenum != matrixB_typenum) {	5,559!
296	throw py::value_error("matrixA and matrixB must be of the same type.");	×
297	}	×
298
299	auto array_types = dpctl_td_ns::usm_ndarray_types();	5,559✔
300	const int matrixAB_type_id =	5,559✔
301	array_types.typenum_to_lookup_id(matrixA_typenum);	5,559✔
302	const int resultC_type_id =	5,559✔
303	array_types.typenum_to_lookup_id(resultC_typenum);	5,559✔
304
305	gemm_impl_fn_ptr_t gemm_fn =	5,559✔
306	gemm_dispatch_table[matrixAB_type_id][resultC_type_id];	5,559✔
307	if (gemm_fn == nullptr) {	5,559!
308	throw py::value_error(	×
309	"Types of input matrices and result matrix are mismatched.");	×
310	}	×
311
312	const char *a_typeless_ptr = matrixA.get_data();	5,559✔
313	const char *b_typeless_ptr = matrixB.get_data();	5,559✔
314	char *r_typeless_ptr = resultC.get_data();	5,559✔
315
316	sycl::event gemm_ev = gemm_fn(exec_q, transA, transB, m, n, k,	5,559✔
317	a_typeless_ptr, lda, b_typeless_ptr, ldb,	5,559✔
318	r_typeless_ptr, ldc, is_row_major, depends);	5,559✔
319
320	sycl::event args_ev = dpctl::utils::keep_args_alive(	5,559✔
321	exec_q, {matrixA, matrixB, resultC}, {gemm_ev});	5,559✔
322
323	return std::make_tuple(args_ev, gemm_ev, is_row_major);	5,559✔
324	}	5,559✔
325
326	bool _is_lnl_arl_architecture(sycl::device &dev)
327	{	5,559✔
328	#if !defined(USE_ONEMKL_CUBLAS)	5,559✔
329	if (dev.ext_oneapi_architecture_is(	5,559!
330	sycl::ext::oneapi::experimental::architecture::	5,559✔
331	intel_gpu_lnl_m)) /* Lunar Lake */	5,559✔
NEW 332	{	×
NEW 333	std::cout << "_is_lnl_arl_architecture: running on Lunar Lake" << std::endl;	×
NEW 334	return true;	×
NEW 335	}	×
336	if (dev.ext_oneapi_architecture_is(	5,559!
337	sycl::ext::oneapi::experimental::architecture::	5,559✔
338	intel_gpu_bmg_g21)) /* Battlemage G21 */	5,559✔
NEW 339	{	×
NEW 340	std::cout << "_is_lnl_arl_architecture: running on Battlemage G21" << std::endl;	×
NEW 341	return true;	×
NEW UNCOV 342	}	×
343	// else if (dev.ext_oneapi_architecture_is(
344	// sycl::ext::oneapi::experimental::architecture::
345	// intel_gpu_arl_h)) /* Arrow Lake H */
346	// {
347	// std::cout << "_is_lnl_arl_architecture: running on Arrow Lake H" << std::endl;
348	// return true;
349	// }
350	// else if (dev.ext_oneapi_architecture_is(
351	// sycl::ext::oneapi::experimental::architecture::
352	// intel_gpu_arl_u)) /* Arrow Lake U */
353	// {
354	// std::cout << "_is_lnl_arl_architecture: running on Arrow Lake U" << std::endl;
355	// return true;
356	// }
357	// else if (dev.ext_oneapi_architecture_is(
358	// sycl::ext::oneapi::experimental::architecture::
359	// intel_gpu_arl_s)) /* Arrow Lake S */
360	// {
361	// std::cout << "_is_lnl_arl_architecture: running on Arrow Lake S" << std::endl;
362	// return true;
363	// }
364	#endif // !defined(USE_ONEMKL_CUBLAS)	5,559✔
365	return false;	5,559✔
366	}	5,559✔
367
368	template <typename fnT, typename Tab, typename Tc>
369	struct GemmContigFactory
370	{
371	fnT get()
372	{	196✔
373	if constexpr (types::GemmTypePairSupportFactory<Tab, Tc>::is_defined) {	196✔
374	return gemm_impl<Tab, Tc>;	8✔
375	}
376	else {	188✔
377	return nullptr;	188✔
378	}	188✔
379	}	196✔
380	};
381
382	void init_gemm_dispatch_table(void)
383	{	1✔
384	dpctl_td_ns::DispatchTableBuilder<gemm_impl_fn_ptr_t, GemmContigFactory,	1✔
385	dpctl_td_ns::num_types>	1✔
386	contig;	1✔
387	contig.populate_dispatch_table(gemm_dispatch_table);	1✔
388	}	1✔
389	} // namespace dpnp::extensions::blas

IntelPython / dpnp / 11145093866

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous