devtools/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2026-04-10 11:34:33 +08:00
Code Issues Packages Projects Releases Wiki Activity
Files
d368998120aec845ff5bdffb4a14a03fcdb09e1a
eigen/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

1654 lines
86 KiB
C
Raw Normal View History

[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// This file is part of Eigen, a lightweight C++ template library for linear algebra.
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
//
// Mehdi Goli Codeplay Software Ltd.
// Ralph Potter Codeplay Software Ltd.
// Luke Iwanski Codeplay Software Ltd.
// Contact: <eigen@codeplay.com>
//
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// This Source Code Form is subject to the terms of the Mozilla Public License v. 2.0. If a copy of the MPL was not
// distributed with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
/*****************************************************************
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
* TensorContractionSycl.h
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
*
* \brief:
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
* TensorContractionSycl.h, provides various tensor contraction kernel for SYCL backend
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
*
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
*****************************************************************/
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
#ifndef EIGEN_CXX11_TENSOR_TENSOR_CONTRACTION_SYCL_H
#define EIGEN_CXX11_TENSOR_TENSOR_CONTRACTION_SYCL_H
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
Add IWYU private pragmas to internal headers.
2023-08-21 16:25:22 +00:00
// IWYU pragma: private
Issue an error in case of direct inclusion of internal headers.
2021-09-10 19:12:26 +00:00
#include "./InternalHeaderCheck.h"
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
namespace Eigen {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
namespace TensorSycl {
namespace internal {
#ifndef EIGEN_SYCL_DISABLE_GEMV
/*!
* \brief TVPanelSize, a template class used for setting the panel size required for launching General TensorVector
* contraction kernel on various hardware devices.
*
* \tparam Scalar: determines the element type of the tensor/vector
*
* \tparam StorageIndex determines the Index type.
*
* \tparam NCWindow: determines the number of non-contracting element to be process by each work-group
*
* \tparam CFactor: determines the number of contracting element to be process by each thread
*
* \tparam NCFactor: determines the number of non-contracting element to be process by each thread
*/
template <typename Scalar, typename StorageIndex, StorageIndex NCWindow, StorageIndex CFactor, StorageIndex NCFactor>
struct TVPanelSize {
// LocalThreadSizeC: determines total number of thread per workgroup for the contracting dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex LocalThreadSizeC = EIGEN_SYCL_LOCAL_THREAD_DIM0;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// LocalThreadSizeNC: determines total number of thread per workgroup for the non-contracting dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constxpr StorageIndex LocalThreadSizeNC = EIGEN_SYCL_LOCAL_THREAD_DIM1;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// TileSizeDimNC: determines the tile size for the non-contracting dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex TileSizeDimNC = NCWindow / NCFactor;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// TileSizeDimC: determines the tile size for the contracting dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex TileSizeDimC = CFactor * LocalThreadSizeNC * LocalThreadSizeC;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// WorkLoadPerThreadNC : determines workload per thread for loading the non-contracting dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex WorkLoadPerThreadNC = TileSizeDimNC / LocalThreadSizeNC;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// WorkLoadPerThreadC: determines workload per thread for loading the non-contracting dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex WorkLoadPerThreadC = TileSizeDimC / LocalThreadSizeC;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// BC : determines if supporting bank conflict is required
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool BC = false;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
};
#endif
/*!
* \brief TTPanelSize, a template class used for setting the panel size required for launching General Tensor Tensor
contraction kernel on various hardware devices.
*
* \tparam Scalar: determines the element type of the tensor
*
* \tparam StorageIndex: determines the Index type.
*
* \tparam REG_SIZE_M: determines workload per thread for loading the M dimension This can be varied based on the
available register on a chosen device(can be controlled by EIGEN_SYCL_REG_M macro).
*
* \tparam REG_SIZE_N: determines workload per thread for loading the N dimension This can be varied based on the
available register on a chosen device(can be controlled by EIGEN_SYCL_REG_N macro).
*
* \tparam TSDK: determines Tile size for dimension K. The packet size is assumed to be considered
*/
template <typename Scalar, typename StorageIndex, StorageIndex REG_SIZE_M, StorageIndex REG_SIZE_N, StorageIndex TSDK>
struct TTPanelSize {
// TileSizeDimK: determines Tile size for dimension K. The packet size is assumed to be considered
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex TileSizeDimK = TSDK;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// WorkLoadPerThreadM : determines workload per thread for loading the M dimension This can be varied based on the
// available register on a chosen device(can be controlled by EIGEN_SYCL_REG_M macro//
#ifndef EIGEN_SYCL_REG_M
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex WorkLoadPerThreadM = REG_SIZE_M;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#else
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex WorkLoadPerThreadM = EIGEN_SYCL_REG_M;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#endif
// WorkLoadPerThreadN : determines workload per thread for loading the N dimension This can be varied based on the
// available register on a chosen device(can be controlled by EIGEN_SYCL_REG_N macro
#ifndef EIGEN_SYCL_REG_N
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex WorkLoadPerThreadN = REG_SIZE_N;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#else
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex WorkLoadPerThreadN = EIGEN_SYCL_REG_N;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#endif
// LocalThreadSizeM: determines total number of thread per workgroup for the m dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex LocalThreadSizeM = EIGEN_SYCL_LOCAL_THREAD_DIM0;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// LocalThreadSizeN: determines total number of thread per workgroup for the n dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex LocalThreadSizeN = EIGEN_SYCL_LOCAL_THREAD_DIM1;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// TileSizeDimM: determines the tile size for the m dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex TileSizeDimM = LocalThreadSizeM * WorkLoadPerThreadM;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// TileSizeDimN: determines the tile size for the n dimension
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex TileSizeDimN = LocalThreadSizeN * WorkLoadPerThreadN;
Fix typos
2024-08-02 00:06:24 +00:00
// LoadPerThreadLhs: determines workload per thread for loading Lhs Tensor. This must be divisible by packetsize
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex LoadPerThreadLhs =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
((TileSizeDimK * WorkLoadPerThreadM * WorkLoadPerThreadN) / (TileSizeDimN));
Fix typos
2024-08-02 00:06:24 +00:00
// LoadPerThreadRhs: determines workload per thread for loading Rhs Tensor. This must be divisible by packetsize
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex LoadPerThreadRhs =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
((TileSizeDimK * WorkLoadPerThreadM * WorkLoadPerThreadN) / (TileSizeDimM));
// BC : determines if supporting bank conflict is required
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool BC = true;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// DoubleBuffer: determines if double buffering technique should be used (This can be disabled by
Fix some typos found
2021-09-23 15:22:00 +00:00
// EIGEN_SYCL_DISABLE_DOUBLE_BUFFER macro when the device does not have sufficient local memory)
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool DoubleBuffer =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#ifdef EIGEN_SYCL_DISABLE_DOUBLE_BUFFER
false;
#else
true;
#endif
};
/* !
* \brief contraction_type: an enum class representing the Tensor Contraction implementation algorithm. This is used to
* specialize the contraction algorithm based on device support for dedicated local memory.
*/
enum class contraction_type { local, no_local };
/* !
* \brief data_source an enum class determining the location of the data in a memory hierarchy (global, local, private).
*/
enum class data_source { global_mem, local_mem, private_mem };
/*!
* \brief read, a template function used for loading the data from global
memory. This function is used to guarantee coalesced and vectorized load whenever possible
*
* \tparam PacketLoad: determines if the each element of this tensor block should be loaded in a packet mode
*
* \param is_coalesced_layout: determines whether or not the Tensor data in a memory can be access coalesced and
vectorized when possible. Coalesced memory access is a key factor in Kernel performance. When a tensor is 2d and the
contracting dimension is 1, it is always possible to accessed tensor data coalesced and vectorized. This is the case
when RHS(right hand side) Tensor is transposed or when LHS(left hand side) Tensor is not transposed.
*
* \tparam PacketType: determines the type of packet
*
* \tparam TensorMapper: determines the input tensor mapper type
*
* \tparam StorageIndex: determines the Index type
* \param tensorMapper: is the input tensor
*
* \param NCIndex: is the non-contracting dim index
*
* \param CIndex is the contracting dim index
*
* \param ld: is the leading dimension of the flattened tensor
*/
template <bool PacketLoad, bool is_coalesced_layout, bool, typename PacketType, typename TensorMapper,
typename StorageIndex>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<PacketLoad, PacketType> read(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const TensorMapper &tensorMapper, const StorageIndex &NCIndex, const StorageIndex &CIndex, const StorageIndex &ld) {
const StorageIndex row = (is_coalesced_layout) ? NCIndex : CIndex;
const StorageIndex col = (is_coalesced_layout) ? CIndex : NCIndex;
return tensorMapper.get_tensor().template packet<Unaligned>(row + (col * ld));
}
/*!
* \brief read, special overload of read function, when the read access is not vectorized
*
* \tparam PacketLoad: determines if the each element of this tensor block should be loaded in a packet mode
*
* \param is_coalesced_layout: determines whether or not the Tensor data in a memory can be access coalesced and
vectorized when possible. Coalesced memory access is a key factor in Kernel performance. When a tensor is 2d and the
contracting dimension is 1, it is always possible to accessed tensor data coalesced and vectorized. This is the case
when RHS(right hand side) Tensor is transposed or when LHS(left hand side) Tensor is not transposed.
*
* \tparam PacketType: determines the type of packet
*
* \tparam TensorMapper: determines the input tensor mapper type
*
* \tparam StorageIndex: determines the Index type
* \param tensorMapper: is the input tensor
*
* \param NCIndex: is the non-contracting dim index
*
* \param CIndex: is the contracting dim index
*/
template <bool PacketLoad, bool, bool IsRhs, typename PacketType, typename TensorMapper, typename StorageIndex>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<!PacketLoad, PacketType> read(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const TensorMapper &tensorMapper, const StorageIndex &NCIndex, const StorageIndex &CIndex, const StorageIndex &) {
const StorageIndex row = (IsRhs) ? CIndex : NCIndex;
const StorageIndex col = (IsRhs) ? NCIndex : CIndex;
return tensorMapper(row, col);
}
/*!
* \brief write, a template function used for storing the data to local memory. This function is used to guarantee
* coalesced and vectorized store whenever possible.
*
* \tparam StorageIndex: determines the Index type
*
* \param ld is the leading dimension of the local memory. ld is a compile time value for the local memory
*
* \tparam data_source: an enum value representing if the location of the data in a memory hierarchy.
*
* \tparam PacketType: determines the type of packet
*
* \tparam DataScalar: determines the output data type
*
* \param packet_data: the data to be written in the local memory
*
* \param ptr: a pointer to the local memory
*
* \param CIndex is the contracting dim index
*/
template <typename StorageIndex, StorageIndex ld, data_source dt, typename PacketType, typename DataScalar>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<dt != data_source::global_mem, void> write(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
PacketType &packet_data, DataScalar ptr) {
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr int PacketSize = Eigen::internal::unpacket_traits<PacketType>::size;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
EIGEN_UNROLL_LOOP
for (int i = 0; i < PacketSize; i++) {
*ptr = PacketWrapper<PacketType, PacketSize>::scalarize(i, packet_data);
ptr += ld;
}
}
/*!
* \brief Overloading the write function for storing the data to global memory, when vectorization enabled This function
* is used to guarantee coalesced and vectorized store whenever possible.
*
* \tparam data_source: an enum value representing if the location of the data in a memory hierarchy.
*
* \tparam PacketType: determines the type of packet
*
* \tparam DataScalar: determines the output data type
*
* \param packet_data: the data to be written in the local memory
*
* \param ptr: a pointer to the local memory
*/
template <data_source dt, typename PacketType, typename DataScalar>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
typename std::enable_if_t<Eigen::internal::unpacket_traits<PacketType>::size != 1 && dt == data_source::global_mem,
void>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
write(PacketType &packet_data, DataScalar *ptr) {
::Eigen::internal::pstoreu<DataScalar, PacketType>(ptr, packet_data);
}
/*!
* \brief Overloading the write function for storing the data to global memory, when vectorization is disabled.
*
* \tparam data_source: an enum value representing if the location of the data in a memory hierarchy.
*
* \tparam PacketType: determines the type of packet
*
* \tparam DataScalar: determines the output data type
*
* \param packet_data: the data to be written in the local memory
*
* \param ptr: a pointer to the local memory
*/
template <data_source dt, typename PacketType, typename DataScalar>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
typename std::enable_if_t<Eigen::internal::unpacket_traits<PacketType>::size == 1 && dt == data_source::global_mem,
void>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
write(PacketType &packet_data, DataScalar *ptr) {
*ptr = packet_data;
}
/*!
* \brief check_boundary: is used to check the edge condition for non-internal blocks.
*
* \tparam is_internal: determines if the block is internal
*/
template <bool is_internal>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool check_boundary(bool) {
return true;
}
/*!
* \brief check_boundary: specialization of the check_boundary for non-internal blocks.
*
* \param cond: true when the data is in range. Otherwise false
*/
template <>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE bool check_boundary<false>(bool cond) {
return cond;
}
/*!
* \brief BlockProperties is a template class that provides different characteristic of a block of each Tensor processed
* by each workgroup.
*
* \tparam is_transposed: iff true, determines whether or not the block of the Tensor is transposed
*
* \tparam packet_load_: determines if the each element of this tensor block should be loaded in a packet mode
*
* \tparam PacketType: determines the type of packet
*
* \tparam OutType: determines the type of each element for this block of tensor. If packet load is true, it will be
* packetType; Otherwise it will be scalar Type
*
* \param elements_per_access determines the size of each element based on OutType
*
* \param is_coalesced_layout determines whether or not the Tensor data in a memory can be access coalesced and
* vectorized when possible. Coalesced memory access is a key factor in Kernel performance. When a tensor is 2d and the
* contracting dimension is 1, it is always possible to accessed tensor data coalesced and vectorized. This is the case
* when RHS(right hand side) Tensor is transposed or when LHS(left hand side) Tensor is not transposed.
*
* \param nc_stride determines the stride of non-contracting dimension to access the next adjustment element within the
* Tensor Block for each workgroup
*
* \param c_stride determines the stride of contracting dimension to access the next adjustment element within the
* Tensor Block for each workgroup
*/
template <bool is_transposed, bool is_rhs_, bool packet_load_, typename PacketType>
struct BlockProperties {
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool packet_load = packet_load_;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef typename Eigen::internal::unpacket_traits<PacketType>::type OutScalar;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool is_rhs = is_rhs_;
Fix usages of wrong namespace
2022-03-21 15:07:53 +00:00
typedef std::conditional_t<packet_load, PacketType, OutScalar> OutType;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr int elements_per_access = Eigen::internal::unpacket_traits<OutType>::size;
static constexpr bool is_coalesced_layout = !(is_transposed ^ is_rhs);
static constexpr int nc_stride = (is_coalesced_layout ? elements_per_access : 1);
static constexpr int c_stride = (is_coalesced_layout ? 1 : elements_per_access);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
};
/*!
* \brief ThreadProperties is a template class that provides each thread's properties within a workgroup. Please see
* the sycl-1.2.1 specification (https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf) for the workgroup,
* work-items
*
* \tparam StorageIndex: determines the StorageIndex Type
*
* \param linearLocalThreadId: determines the linearized location of a thread within a work-group
*
* \param kGroupId: determines the logical group id in a k dimension of the flattened tensor. It will be > 1 when
* tall/skinny algorithm is used
*
* \param mGroupOffset: determines the logical start position of all thread within a workgroup for the m dimension of
* the flattened tensor.
*
* \param kGroupOffset determines the logical start position of all thread within a workgroup for the k dimension of the
* flattened tensor. It will be > 1 when tall/skinny algorithm is used.
*
* \param mLocalOffset: determines the logical start position of each thread within a workgroup for the m dimension of a
* flattened tensor. The position determines the distance of each thread within the workgroup from each other
* independent from their global position.
*
* \param nLocalOffset: determines the logical start position of each thread within a workgroup for the n dimension of a
* flattened tensor. The position determines the distance of each thread within the workgroup from each other
* independent from their global position.
*
* \param mGlobalOffset: determines the logical start position of each thread a thread for the m dimension on a
* flattened tensor
*
* \param nGlobalOffset: determines the logical start position of each thread a thread for the n dimension on a
* flattened tensor
*
* \param kSize : determine the number of the k elements of the flattened Tensor to be processed by each thread for the
* given tensor block. This is !=K dimension of Flattened Tensor when Tall/Skinny matrix is used.
*
* \param is_internal : this will determined if the thread within the work-group computes an internal block of tensor or
* the edge blocks. When it is internal, there is no need to check the boundaries and all the if stantement can be
* resolve by compiler.
*/
template <typename StorageIndex>
struct ThreadProperties {
const StorageIndex linearLocalThreadId;
const StorageIndex kGroupId;
const StorageIndex mGroupOffset;
const StorageIndex nGroupOffset;
const StorageIndex kGroupOffset;
const StorageIndex mLocalOffset;
const StorageIndex nLocalOffset;
const StorageIndex mGlobalOffset;
const StorageIndex nGlobalOffset;
StorageIndex kSize;
const bool is_internal;
// this is used to adjust the last block
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE ThreadProperties(
const StorageIndex linearLocalThreadId_, const StorageIndex kGroupId_, const StorageIndex mGroupOffset_,
const StorageIndex nGroupOffset_, const StorageIndex kGroupOffset_, const StorageIndex mLocalOffset_,
const StorageIndex nLocalOffset_, const StorageIndex mGlobalOffset_, const StorageIndex nGlobalOffset_,
StorageIndex kSize_, const bool is_internal_)
: linearLocalThreadId(linearLocalThreadId_),
kGroupId(kGroupId_),
mGroupOffset(mGroupOffset_),
nGroupOffset(nGroupOffset_),
kGroupOffset(kGroupOffset_),
mLocalOffset(mLocalOffset_),
nLocalOffset(nLocalOffset_),
mGlobalOffset(mGlobalOffset_),
nGlobalOffset(nGlobalOffset_),
kSize(kSize_),
is_internal(is_internal_) {}
};
/*!
* \brief TensorContractionKernel is a template class that provides Tensor -Tensor contraction operation.
*
* \tparam OutScalar: determines the output scalar type
*
* \tparam LhsScalar: determines the left-hand-side scalar type
*
* \tparam RhsScalar: determines the right-hand-side scalar type
*
* \tparam OutAccessor: determines the sycl accessor type for out put (please see the sycl-1.2.1 specification
(https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf) for accessor definition)
*
* \tparam LhsMapper determines the tensor contraction mapper type for left-hand-side matrix
*
* \tparam RhsMapper determines the tensor contraction mapper type for right-hand-side matrix
*
* \tparam StorageIndex: determines the StorageIndex Type
*
* \tparam Properties: determines the Contraction Panel properties
*
* \tparam TripleDim: determines the M, K, N dimensions for the flatten tensors in order to treat them as a matrix
*
* \tparam Vectorizable: determines whether or not the vectorization is enabled for the Eigen expression.
*
* \tparam input_mapper_properties : determine if the input tensors are matrix. If they are matrix, special memory
access is used to guarantee that always the memory access are coalesced.
*
Fix all the doxygen warnings.
2025-02-01 00:00:31 +00:00
* \tparam IsFinal : determine if this is the final kernel. If so, the result will be written in a final output.
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
Otherwise, the result of contraction will be written iin a temporary buffer. This is the case when Tall/Skinny
contraction is used. So in this case, a final reduction step is required to compute final output.
Fix some typos found
2021-09-23 15:22:00 +00:00
* \tparam contraction_tp: it is an enum value representing whether the local memory/no local memory implementation of
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
the algorithm to be used
*
* \param scratch: local memory containing tiles of LHS and RHS tensors for each work-group
*
* \param lhs: determines the left-hand-side flattened tensor (tensor mapper)
*
* \param rhs: determines the right-hand-side flattened tensor (tensor mapper)
*
* \param out_res: determines the output tensor containing the contraction result
*
* \param groupSizeM: a logical number determining the number of work-group for m dimension
*
* \param groupSizeN: a logical number determining the number of work-group for n dimension
*
* \param numTiles: determines total number of tiles on the k dimension
*
* \param TripleDim: determines the M, K, N dimensions for the flatten tensors in order to treat them as a matrix
*/
template <typename OutScalar, typename LhsScalar, typename RhsScalar, typename OutAccessor, typename LhsMapper,
typename RhsMapper, typename StorageIndex, typename Properties, typename TripleDim, bool Vectorizable,
typename input_mapper_properties, bool IsFinal, contraction_type contraction_tp>
class TensorContractionKernel {
public:
typedef typename Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketReturnType
PacketReturnType;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr int PacketSize =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketSize;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool is_lhs_transposed =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
!::Eigen::internal::TensorContractionInputMapperTrait<LhsMapper>::inner_dim_contiguous;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool is_rhs_transposed =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
!::Eigen::internal::TensorContractionInputMapperTrait<RhsMapper>::inner_dim_contiguous;
typedef BlockProperties<is_lhs_transposed, false, input_mapper_properties::is_lhs_matrix && Vectorizable,
PacketReturnType>
LHSBlockProperties;
typedef BlockProperties<is_rhs_transposed, true, input_mapper_properties::is_rhs_matrix && Vectorizable,
PacketReturnType>
RHSBlockProperties;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex NStride =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
contraction_tp == contraction_type::local ? Properties::WorkLoadPerThreadN : RHSBlockProperties::nc_stride;
typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Scratch;
typedef cl::sycl::multi_ptr<OutScalar, cl::sycl::access::address_space::local_space> local_ptr;
typedef OutScalar * /*cl::sycl::multi_ptr<OutScalar, cl::sycl::access::address_space::private_space>*/ private_ptr;
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
typedef std::conditional_t<contraction_tp == contraction_type::local, local_ptr, private_ptr> tile_ptr;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex LSDL = contraction_tp == contraction_type::local
? Properties::TileSizeDimM + Properties::BC
: Properties::WorkLoadPerThreadM;
static constexpr StorageIndex LSDR = contraction_tp == contraction_type::local
? Properties::TileSizeDimN + Properties::BC
: Properties::WorkLoadPerThreadN;
static constexpr StorageIndex LocalOffset = Properties::LocalThreadSizeM * Properties::LocalThreadSizeN;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
/**
* \brief MemHolder this is a place holder struct for creating memory hierarchy in SYCL. Inside SYCL kernel it is not
* allowed to have dynamic memory allocation. While the local memory is created outside of the kernel and passed to
* the kernel as an accessor, the private memory can only allowed to be allocated statically. Since we are abstracting
* the TiledMemory for both local and private memory, the MemHolder structs is used as a helper to abstract out
* different type of memory needed when local/no_local memory computation is called.
*
Fix some typos found
2021-09-23 15:22:00 +00:00
* \tparam contraction_type: it is an enum value representing whether the local memory/no local memory implementation
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
of the algorithm to be used
* \tparam the private memory size
* \param ptr the tile memory pointer type
*/
template <contraction_type, StorageIndex>
struct MemHolder {
tile_ptr ptr;
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE MemHolder(local_ptr block_start_ptr) : ptr(block_start_ptr) {}
};
/**
* \brief specialization of memHolder class when no local memory kernel is used.
*/
template <StorageIndex MemSize>
struct MemHolder<contraction_type::no_local, MemSize> {
OutScalar ptr[MemSize] = {OutScalar{0}};
};
/**
* \brief TiledMemory: contains required memory pointer for loading each tile of the TensorContraction panel from
* global memory to local/private memory when local/no_local algorithm used.
*
* \param lhs_scratch_extract : determines the LHS tile memory. It is either private or local memory based on the
* selected contraction_type.
*
* \param rhs_scratch_extract : determines the RHS tile memory. It is either private or local memory based on the
* selected contraction_type.
*
fix various typos
2021-09-22 16:15:06 +00:00
* \param lhs_extract_index: determines the position of each thread on a local memory for lhs input. When private
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
* memory is used this is set to zero as this is not applicable in case of private memory.
*
fix various typos
2021-09-22 16:15:06 +00:00
* \param rhs_extract_index: determines the position of each thread on a local memory for rhs input. When private
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
* memory is used this is set to zero as this is not applicable in case of private memory.
*
* \param lhs_scratch_compute : determines the location to load for computation for lhs_local memory. This is the
* same as lhs_scratch_extract for private memory.
*
* \param rhs_scratch_compute : determines the location to load for computation for rhs_local memory. This is the
* same as rhs_scratch_extract for private memory.
*/
struct TiledMemory {
MemHolder<contraction_tp, Properties::WorkLoadPerThreadM * Properties::TileSizeDimK> lhs_scratch_extract;
MemHolder<contraction_tp, Properties::WorkLoadPerThreadN * Properties::TileSizeDimK> rhs_scratch_extract;
tile_ptr lhs_scratch_ptr_compute;
tile_ptr rhs_scratch_ptr_compute;
const std::pair<StorageIndex, StorageIndex> lhs_extract_index;
const std::pair<StorageIndex, StorageIndex> rhs_extract_index;
template <contraction_type tp = contraction_tp>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TiledMemory(const ThreadProperties<StorageIndex> &, local_ptr,
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
std::enable_if_t<tp == contraction_type::no_local> * = 0)
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
: lhs_scratch_extract{},
rhs_scratch_extract{},
lhs_scratch_ptr_compute(lhs_scratch_extract.ptr),
rhs_scratch_ptr_compute(rhs_scratch_extract.ptr),
lhs_extract_index(std::pair<StorageIndex, StorageIndex>(StorageIndex{0}, StorageIndex{0})),
rhs_extract_index(std::pair<StorageIndex, StorageIndex>(StorageIndex{0}, StorageIndex{0})) {}
template <contraction_type tp = contraction_tp>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TiledMemory(const ThreadProperties<StorageIndex> &thread_properties,
local_ptr block_start_ptr,
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
std::enable_if_t<tp == contraction_type::local> * = 0)
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
: lhs_scratch_extract{block_start_ptr},
rhs_scratch_extract{lhs_scratch_extract.ptr +
((Properties::DoubleBuffer + 1) * LSDL * Properties::TileSizeDimK)},
lhs_scratch_ptr_compute(lhs_scratch_extract.ptr + thread_properties.mLocalOffset),
rhs_scratch_ptr_compute(rhs_scratch_extract.ptr + thread_properties.nLocalOffset),
lhs_extract_index(
local_id_extract<LHSBlockProperties, Properties::TileSizeDimM>(thread_properties.linearLocalThreadId)),
rhs_extract_index(
local_id_extract<RHSBlockProperties, Properties::TileSizeDimN>(thread_properties.linearLocalThreadId)) {}
};
Scratch scratch;
const LhsMapper lhs;
const RhsMapper rhs;
OutAccessor out_res;
const StorageIndex groupSizeM;
const StorageIndex groupSizeN;
const StorageIndex numTiles;
const TripleDim triple_dim;
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorContractionKernel(Scratch scratch_, const LhsMapper lhs_,
const RhsMapper rhs_, OutAccessor out_res_,
const StorageIndex groupSizeM_,
const StorageIndex groupSizeN_,
const StorageIndex numTiles_,
const TripleDim triple_dim_)
: scratch(scratch_),
lhs(lhs_),
rhs(rhs_),
out_res(out_res_),
groupSizeM(groupSizeM_),
groupSizeN(groupSizeN_),
numTiles(numTiles_),
triple_dim(triple_dim_) {}
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorContractionKernel(Scratch scratch_, const LhsMapper lhs_,
const RhsMapper rhs_, OutAccessor out_res_,
const StorageIndex groupSizeM_,
const StorageIndex numTiles_,
const TripleDim triple_dim_)
: TensorContractionKernel(scratch_, lhs_, rhs_, out_res_, groupSizeM_, 1, numTiles_, triple_dim_) {}
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void operator()(cl::sycl::nd_item<1> itemID) const {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const StorageIndex linearLocalThreadId = itemID.get_local_id(0);
const StorageIndex nLocalThreadId = linearLocalThreadId / Properties::LocalThreadSizeM;
const StorageIndex mLocalThreadId = linearLocalThreadId % Properties::LocalThreadSizeM;
const StorageIndex mGroupId = itemID.get_group(0) % groupSizeM;
const StorageIndex tmp = itemID.get_group(0) / groupSizeM;
const StorageIndex nGroupId = IsFinal ? tmp : tmp % groupSizeN;
const StorageIndex kGroupId = IsFinal ? 0 : tmp / groupSizeN;
const StorageIndex mGroupOffset = mGroupId * Properties::TileSizeDimM;
const StorageIndex nGroupOffset = nGroupId * Properties::TileSizeDimN;
const StorageIndex mLocalOffset = PacketSize * mLocalThreadId;
const StorageIndex nLocalOffset = NStride * nLocalThreadId;
const StorageIndex mGlobalOffset = mGroupOffset + mLocalOffset;
const StorageIndex nGlobalOffset = nGroupOffset + nLocalOffset;
const StorageIndex kSizePerWG = IsFinal ? triple_dim.K : numTiles * Properties::TileSizeDimK;
StorageIndex kGroupOffset = kGroupId * kSizePerWG;
const bool is_internal = triple_dim.M - mGroupOffset >= Properties::TileSizeDimM &&
triple_dim.N - nGroupOffset >= Properties::TileSizeDimN &&
triple_dim.K - kGroupOffset >= kSizePerWG;
// this is used to adjust the last block
StorageIndex kSize = IsFinal ? triple_dim.K : std::min(kSizePerWG, triple_dim.K - kGroupOffset);
// This is used to find out the lats K offset so that kGroupOffset -kSize can compute the coffset for loading to
// tile
kGroupOffset += kSize;
auto thread_properties =
ThreadProperties<StorageIndex>(linearLocalThreadId, kGroupId, mGroupOffset, nGroupOffset, kGroupOffset,
mLocalOffset, nLocalOffset, mGlobalOffset, nGlobalOffset, kSize, is_internal);
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
auto out_ptr = out_res + (IsFinal ? 0 : thread_properties.kGroupId * triple_dim.M * triple_dim.N);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
(thread_properties.is_internal) ? compute_panel<true>(itemID, thread_properties, out_ptr)
: compute_panel<false>(itemID, thread_properties, out_ptr);
}
// The compute block computes the contraction operation private block for each thread and store the resutl in the
// privateRes memory of Each computation the compute block function is independent of local and no local concepts as
// it only compute the block on each thread's private memory space
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void compute_block_per_tile(OutScalar *lhs_block_ptr, OutScalar *rhs_block_ptr,
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
PacketReturnType *privateRes) const {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
StorageIndex idx = 0;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex lhs_stride =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
contraction_tp == contraction_type::local ? (PacketSize * Properties::LocalThreadSizeM) : 1;
EIGEN_UNROLL_LOOP
for (StorageIndex wLPTN = 0; wLPTN < Properties::WorkLoadPerThreadN; wLPTN++) {
auto rhsPacket = PacketReturnType{*(rhs_block_ptr + wLPTN)};
StorageIndex lhs_index = 0;
EIGEN_UNROLL_LOOP
for (StorageIndex wLPTM = 0; wLPTM < Properties::WorkLoadPerThreadM / PacketSize; wLPTM++) {
PacketReturnType lhsPack{};
Eigen::TensorSycl::internal::PacketWrapper<PacketReturnType, PacketSize>::set_packet(lhsPack,
lhs_block_ptr + lhs_index);
privateRes[idx] = ::Eigen::internal::pmadd(lhsPack, rhsPacket, privateRes[idx]);
lhs_index += lhs_stride;
idx++;
}
}
}
// The store function write the computed contraction operation in the private memory of each thread to the global
// memory. The store function is independent of local and no local concepts s that it can be abstract out in the base
// class.
template <bool is_internal_block, StorageIndex PrivateNStride, typename OutPtr>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void store(OutPtr *out_ptr, PacketReturnType *privateRes,
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
StorageIndex mGlobalOffset, StorageIndex nGlobalOffset) const {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
auto chk_bound = [&](const StorageIndex &mIndex, const StorageIndex &nIndex) EIGEN_DEVICE_FUNC {
return (mIndex + PacketSize - 1 < triple_dim.M && nGlobalOffset + nIndex < triple_dim.N);
};
// when local memory is not used M and N are both accessed in a coalesced way. However, when local memory is
// available the k*N is transposed in the local to N*K therefore, each blocks operates on blockId*
// WorkLoadPerThreadN slice of N
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex GlobalNStride = contraction_tp == contraction_type::local ? 1 : Properties::LocalThreadSizeN;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
EIGEN_UNROLL_LOOP
for (StorageIndex wLPTN = 0; wLPTN < Properties::WorkLoadPerThreadN / PrivateNStride; wLPTN++) {
// output leading dimension
StorageIndex outputLD = 0;
Fix typos
2024-08-02 00:06:24 +00:00
// When local memory is used the PrivateNstride is always 1 because the coalesced access on N is loaded into Local
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// memory and extracting from local to global is the same as no transposed version. However, when local memory is
// not used and RHS is transposed we packetize the load for RHS.
EIGEN_UNROLL_LOOP
for (StorageIndex nId = 0; nId < PrivateNStride; nId++) {
StorageIndex globalRow = mGlobalOffset;
EIGEN_UNROLL_LOOP
for (StorageIndex wLPTM = 0; wLPTM < Properties::WorkLoadPerThreadM / PacketSize; wLPTM++) {
PacketReturnType privetOut = privateRes[wLPTM];
if (check_boundary<is_internal_block>(chk_bound(globalRow, nId))) {
// Store the final results in C. The C matrix has always M as a first StorageIndex and N as a second
// StorageIndex Therefore it is always coalesced layout
write<data_source::global_mem>(privetOut, out_ptr + outputLD + globalRow);
} else {
EIGEN_UNROLL_LOOP
for (StorageIndex mId = 0; mId < PacketSize; mId++) {
StorageIndex mOffset = globalRow + mId;
if (mOffset < triple_dim.M && (nGlobalOffset + nId < triple_dim.N)) {
out_ptr[mOffset + outputLD] =
Eigen::TensorSycl::internal::PacketWrapper<PacketReturnType, PacketSize>::scalarize(mId, privetOut);
}
}
}
globalRow += (PacketSize * Properties::LocalThreadSizeM);
}
outputLD += triple_dim.M;
privateRes += Properties::WorkLoadPerThreadM / PacketSize;
}
out_ptr += (GlobalNStride * outputLD);
nGlobalOffset += (PrivateNStride * GlobalNStride);
}
}
// when no local memory is used the following extract_block will be enabled
template <typename InputBlockProperties, bool is_internal_block, typename Input, typename PrivateReg,
contraction_type contract_tp = contraction_tp>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<contract_tp == contraction_type::no_local> extract_block(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const Input &inpt, PrivateReg private_ptr, const std::pair<StorageIndex, StorageIndex> &,
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
const StorageIndex &ncOffset, const StorageIndex cOffset) const {
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex LocalThreadSizeNC =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
InputBlockProperties::is_rhs ? Properties::LocalThreadSizeN : Properties::LocalThreadSizeM;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex WorkLoadPerThreadNC =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
InputBlockProperties::is_rhs ? Properties::WorkLoadPerThreadN : Properties::WorkLoadPerThreadM;
const StorageIndex &NC = InputBlockProperties::is_rhs ? triple_dim.N : triple_dim.M;
auto chk_bound = [&](const StorageIndex &CIndex, const StorageIndex &NCIndex) EIGEN_DEVICE_FUNC {
return ((CIndex + InputBlockProperties::c_stride - 1 < triple_dim.K) &&
(NCIndex + InputBlockProperties::nc_stride - 1 < NC));
};
const StorageIndex ld = InputBlockProperties::is_coalesced_layout ? NC : triple_dim.K;
StorageIndex cIndex = cOffset;
EIGEN_UNROLL_LOOP
for (StorageIndex cId = 0; cId < Properties::TileSizeDimK / InputBlockProperties::c_stride; cId++) {
StorageIndex ncIndex = ncOffset;
EIGEN_UNROLL_LOOP
for (StorageIndex ncId = 0; ncId < WorkLoadPerThreadNC / InputBlockProperties::nc_stride; ncId++) {
if (check_boundary<is_internal_block>(chk_bound(cIndex, ncIndex))) {
auto val =
read<InputBlockProperties::packet_load, InputBlockProperties::is_coalesced_layout,
InputBlockProperties::is_rhs, typename InputBlockProperties::OutType>(inpt, ncIndex, cIndex, ld);
write<StorageIndex, (InputBlockProperties::is_coalesced_layout ? 1 : WorkLoadPerThreadNC),
data_source::private_mem>(val, private_ptr);
} else {
EIGEN_UNROLL_LOOP
for (StorageIndex i = 0; i < InputBlockProperties::elements_per_access; i++) {
const StorageIndex ncInd = ncIndex + (InputBlockProperties::is_coalesced_layout ? i : 0);
const StorageIndex cInd = cIndex + (InputBlockProperties::is_coalesced_layout ? 0 : i);
OutScalar val =
(ncInd < NC && cInd < triple_dim.K)
? read<false, InputBlockProperties::is_coalesced_layout, InputBlockProperties::is_rhs, OutScalar>(
inpt, ncInd, cInd, ld)
: OutScalar(0);
write<StorageIndex, (InputBlockProperties::is_coalesced_layout ? 1 : WorkLoadPerThreadNC),
data_source::private_mem>(
val, private_ptr + (InputBlockProperties::is_coalesced_layout ? i : 0) +
((InputBlockProperties::is_coalesced_layout ? 0 : i) * WorkLoadPerThreadNC));
}
}
// if it is lhs we have to load it packetised when the packet size is > 1, because the output is coalesced. So
// even if M is not accessed in a coalesced mode, we have to load packet_size number of m per thread.
ncIndex = (!InputBlockProperties::is_rhs && InputBlockProperties::nc_stride == 1 && PacketSize != 1)
? ncOffset + (ncId + 1) % PacketSize + ((ncId + 1) / PacketSize) * LocalThreadSizeNC
: (ncIndex + InputBlockProperties::nc_stride * LocalThreadSizeNC);
private_ptr += InputBlockProperties::nc_stride;
}
// the previous for loop ( private_ptr += (ncId * nc_stride)) has already moved ptr with one WorkLoadPerThreadNC
private_ptr += (InputBlockProperties::c_stride - 1) * WorkLoadPerThreadNC;
cIndex += InputBlockProperties::c_stride;
}
}
template <typename InputBlockProperties, StorageIndex TileSizeDimNC>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::pair<StorageIndex, StorageIndex> local_id_extract(
const StorageIndex &linearLocalThreadId) {
const StorageIndex localThreadNC =
(InputBlockProperties::is_coalesced_layout)
? linearLocalThreadId % (TileSizeDimNC / InputBlockProperties::nc_stride)
: linearLocalThreadId / (Properties::TileSizeDimK / InputBlockProperties::c_stride);
const StorageIndex localThreadC =
(InputBlockProperties::is_coalesced_layout)
? linearLocalThreadId / (TileSizeDimNC / InputBlockProperties::nc_stride)
: linearLocalThreadId % (Properties::TileSizeDimK / InputBlockProperties::c_stride);
return std::pair<StorageIndex, StorageIndex>(localThreadNC, localThreadC);
}
template <bool db = Properties::DoubleBuffer, contraction_type ctp = contraction_tp>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<db && ctp == contraction_type::local> sync_mem(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const cl::sycl::nd_item<1> &, bool &db_offset) noexcept {
db_offset = !db_offset;
}
template <bool db = Properties::DoubleBuffer, contraction_type ctp = contraction_tp>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<!db && ctp == contraction_type::local> sync_mem(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const cl::sycl::nd_item<1> &itemID, bool &) noexcept {
itemID.barrier(cl::sycl::access::fence_space::local_space);
}
template <contraction_type ctp = contraction_tp>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<ctp == contraction_type::no_local> sync_mem(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const cl::sycl::nd_item<1> &, bool &) noexcept {
return;
}
template <bool need_sync, contraction_type ctp = contraction_tp>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<need_sync && ctp == contraction_type::no_local>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
sync_thread(const cl::sycl::nd_item<1> &
#ifdef EIGEN_SYCL_ARM_GPU_CACHE_OPTIMISATION
itemID
#endif
) noexcept {
#ifdef EIGEN_SYCL_ARM_GPU_CACHE_OPTIMISATION
itemID.barrier(cl::sycl::access::fence_spacce::local_space);
#else
return;
#endif
}
template <bool need_sync, contraction_type ctp = contraction_tp>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<need_sync && ctp == contraction_type::local>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
sync_thread(const cl::sycl::nd_item<1> &itemID) {
itemID.barrier(cl::sycl::access::fence_space::local_space);
}
template <bool need_sync>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<!need_sync> sync_thread(const cl::sycl::nd_item<1> &) {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
return;
}
Fix TensorContractionOp evaluators for GPU and SYCL
2018-07-17 14:09:37 -07:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
template <bool is_internal_block>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void compute_tile_per_panel(const cl::sycl::nd_item<1> &itemID,
ThreadProperties<StorageIndex> &thread_properties,
TiledMemory &tiled_input_block,
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
PacketReturnType *privateRes,
bool &db_offset) const {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// Tiling the Rhs block from global to local memory
extract_block<RHSBlockProperties, is_internal_block>(
rhs, tiled_input_block.rhs_scratch_extract.ptr + (db_offset * Properties::TileSizeDimK * LSDR),
tiled_input_block.rhs_extract_index,
contraction_tp == contraction_type::local ? thread_properties.nGroupOffset : thread_properties.nGlobalOffset,
thread_properties.kGroupOffset - thread_properties.kSize);
sync_thread<contraction_tp == contraction_type::no_local>(itemID);
// Tiling the Lhs block from global to local memory
extract_block<LHSBlockProperties, is_internal_block>(
lhs, tiled_input_block.lhs_scratch_extract.ptr + (db_offset * LSDL * Properties::TileSizeDimK),
tiled_input_block.lhs_extract_index,
contraction_tp == contraction_type::local ? thread_properties.mGroupOffset : thread_properties.mGlobalOffset,
thread_properties.kGroupOffset - thread_properties.kSize);
// itemID.barrier(cl::sycl::access::fence_space::local_space);
sync_thread<contraction_tp == contraction_type::local>(itemID);
// switch to compute mede
StorageIndex lhs_offset = (db_offset * LSDL * Properties::TileSizeDimK);
StorageIndex rhs_offset = (db_offset * Properties::TileSizeDimK * LSDR);
// Loop over the values of a single tile
for (StorageIndex k = 0; k < Properties::TileSizeDimK; k++) {
compute_block_per_tile(tiled_input_block.lhs_scratch_ptr_compute + lhs_offset,
tiled_input_block.rhs_scratch_ptr_compute + rhs_offset, privateRes);
lhs_offset += LSDL;
rhs_offset += LSDR;
}
// computing the K index for the next tile
thread_properties.kSize -= Properties::TileSizeDimK;
sync_mem(itemID, db_offset);
}
// when local memory is available the following compute_panel will be enabled
template <bool is_internal_block, typename OutPtr>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void compute_panel(const cl::sycl::nd_item<1> &itemID,
ThreadProperties<StorageIndex> &thread_properties,
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
OutPtr out_ptr) const {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
auto tiled_input_block = TiledMemory{thread_properties, scratch.get_pointer()};
// Allocate register space
PacketReturnType privateRes[Properties::WorkLoadPerThreadM * Properties::WorkLoadPerThreadN / PacketSize] = {
PacketReturnType{0}};
bool db_offset = 0;
while (thread_properties.kSize >= Properties::TileSizeDimK) {
compute_tile_per_panel<is_internal_block>(itemID, thread_properties, tiled_input_block, privateRes, db_offset);
}
if (thread_properties.kSize > 0) {
compute_tile_per_panel<false>(itemID, thread_properties, tiled_input_block, privateRes, db_offset);
}
// Storing the final results in the output
store<is_internal_block,
contraction_tp == contraction_type::local ? static_cast<StorageIndex>(1) : RHSBlockProperties::nc_stride>(
out_ptr + thread_properties.nGlobalOffset * triple_dim.M, privateRes, thread_properties.mGlobalOffset,
thread_properties.nGlobalOffset);
}
// When local memory is available the following extract_block will be enabled
template <typename InputBlockProperties, bool is_internal_block, typename Input, typename Local,
contraction_type contract_tp = contraction_tp>
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::enable_if_t<contract_tp == contraction_type::local> extract_block(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const Input &inpt, Local local_ptr, const std::pair<StorageIndex, StorageIndex> &local_index,
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
const StorageIndex &ncOffset, const StorageIndex cOffset) const {
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex TileSizeDimNC =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
InputBlockProperties::is_rhs ? Properties::TileSizeDimN : Properties::TileSizeDimM;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex LoadPerThread =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
InputBlockProperties::is_rhs ? Properties::LoadPerThreadRhs : Properties::LoadPerThreadLhs;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex LSD = InputBlockProperties::is_rhs ? LSDR : LSDL;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
static_assert(((LocalOffset % (TileSizeDimNC / InputBlockProperties::nc_stride) == 0) &&
(LocalOffset % (Properties::TileSizeDimK / InputBlockProperties::c_stride) == 0)),
Fix typos
2024-08-02 00:06:24 +00:00
" LocalOffset must be divisible by stride");
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const StorageIndex &NC = InputBlockProperties::is_rhs ? triple_dim.N : triple_dim.M;
StorageIndex localThreadNC = local_index.first;
StorageIndex localThreadC = local_index.second;
auto chk_bound = [&](const StorageIndex &CIndex, const StorageIndex &NCIndex) EIGEN_DEVICE_FUNC {
return ((CIndex + InputBlockProperties::c_stride - 1 < triple_dim.K) &&
(NCIndex + InputBlockProperties::nc_stride - 1 < NC));
};
EIGEN_UNROLL_LOOP
for (StorageIndex lPT = 0; lPT < LoadPerThread / InputBlockProperties::elements_per_access; lPT++) {
const StorageIndex CIndex = cOffset + (InputBlockProperties::c_stride * localThreadC);
const StorageIndex NCIndex = ncOffset + (InputBlockProperties::nc_stride * localThreadNC);
const StorageIndex ld = InputBlockProperties::is_coalesced_layout ? NC : triple_dim.K;
if (check_boundary<is_internal_block>(chk_bound(CIndex, NCIndex))) {
auto val =
read<InputBlockProperties::packet_load, InputBlockProperties::is_coalesced_layout,
InputBlockProperties::is_rhs, typename InputBlockProperties::OutType>(inpt, NCIndex, CIndex, ld);
write<StorageIndex, (InputBlockProperties::is_coalesced_layout ? 1 : LSD), data_source::local_mem>(
val, local_ptr + (InputBlockProperties::nc_stride * localThreadNC) +
(InputBlockProperties::c_stride * localThreadC * LSD));
} else {
EIGEN_UNROLL_LOOP
for (StorageIndex i = 0; i < InputBlockProperties::elements_per_access; i++) {
const StorageIndex nCInd = NCIndex + (InputBlockProperties::is_coalesced_layout ? i : 0);
const StorageIndex cInd = CIndex + (InputBlockProperties::is_coalesced_layout ? 0 : i);
OutScalar val =
(nCInd < NC && cInd < triple_dim.K)
? read<false, InputBlockProperties::is_coalesced_layout, InputBlockProperties::is_rhs, OutScalar>(
inpt, nCInd, cInd, ld)
: OutScalar(0);
write<StorageIndex, (InputBlockProperties::is_coalesced_layout ? 1 : LSD), data_source::local_mem>(
val, local_ptr + (InputBlockProperties::nc_stride * localThreadNC) +
(InputBlockProperties::is_coalesced_layout ? i : 0) +
((InputBlockProperties::c_stride * localThreadC +
(InputBlockProperties::is_coalesced_layout ? 0 : i)) *
LSD));
}
}
localThreadNC += (InputBlockProperties::is_coalesced_layout)
? LocalOffset % (TileSizeDimNC / InputBlockProperties::nc_stride)
: LocalOffset / (Properties::TileSizeDimK / InputBlockProperties::c_stride);
localThreadC += (InputBlockProperties::is_coalesced_layout)
? LocalOffset / (TileSizeDimNC / InputBlockProperties::nc_stride)
: LocalOffset % (Properties::TileSizeDimK / InputBlockProperties::c_stride);
}
}
};
#ifndef EIGEN_SYCL_DISABLE_GEMV
/*!
* \brief GeneralVectorTensor is a template class that provides Tensor -vector contraction operation, which is a special
* case of Tensor Tensor contraction.
*
* \tparam OutScalar: determines the output scalar type
*
* \tparam OutAccessor: determines the sycl accessor type for out put (please see the sycl-1.2.1 specification
* (https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf) for accessor definition)
*
* \tparam VectorMapper: determines the tensor contraction mapper for the vector input (can be lhs or rhs)
*
* \tparam TensorMapper: determines the tensor contraction mapper for the tensor input (can be lhs or rhs)
*
* \tparam StorageIndex: determines the StorageIndex Type
*
* \tparam Properties: determines the Contraction Panel properties
*
* \tparam KFactor: determines the number of elements in K dimension in a Tile
*
* \tparam Vectorizable: determines whether or not the vectorization is enabled for the Eigen expression.
*
* \tparam is_lhs_vec: determines whether lhs is a vector or rhs is a vector
*
* \tparam IsFinal: determine if this is the final kernel. If so, the result will be written in a final output.
* Otherwise, the result of contraction will be written iin a temporary buffer.
*
* \param scratch: determines the local memory containing the vector block for each work-group
*
* \param vec: determines the vector input (tensor mapper)
*
* \param mat: determines the tensor input (tensor mapper)
*
* \param out_res: determines the output vector containing the contraction result
*
* \param nonContractGroupSize: a logical number determining the number of work-group for non-contracting dimension
*
* \param nonContractDim: determines the size of non contracting dimension for the flattened tensor
*
* \param contractDim: determines the size of non contracting dimension for the flattened tensor
*
*/
template <typename OutScalar, typename OutAccessor, typename VectorMapper, typename TensorMapper, typename StorageIndex,
typename Properties, StorageIndex KFactor, bool Vectorizable, bool is_lhs_vec, bool IsFinal>
struct GeneralVectorTensor {
typedef typename Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketReturnType
PacketReturnType;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr int PacketSize =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
Eigen::TensorSycl::internal::Vectorise<OutScalar, Eigen::SyclDevice, Vectorizable>::PacketSize;
typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Scratch;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr StorageIndex OutScratchOffset =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
KFactor * Properties::LocalThreadSizeC * Properties::LocalThreadSizeNC;
// Since the access layout for a vector can always be coalesced, when LHS is a vector, we pass false and false to make
// sure that the !^ is true When RHS is a vector, we pass true and true to make sure that the !^ is true.
typedef BlockProperties<is_lhs_vec ? false : true, is_lhs_vec ? false : true, Vectorizable, PacketReturnType>
VecBlockProperties;
Scratch scratch;
const VectorMapper vec;
const TensorMapper mat;
OutAccessor out_res;
const StorageIndex nonContractGroupSize;
const StorageIndex nonContractDim;
const StorageIndex contractDim;
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE GeneralVectorTensor(Scratch scratch_, const VectorMapper vec_,
const TensorMapper mat_, OutAccessor out_res_,
const StorageIndex nonContractGroupSize_,
const StorageIndex nonContractDim_,
const StorageIndex contractDim_)
: scratch(scratch_),
vec(vec_),
mat(mat_),
out_res(out_res_),
nonContractGroupSize(nonContractGroupSize_),
nonContractDim(nonContractDim_),
contractDim(contractDim_) {}
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void operator()(cl::sycl::nd_item<1> itemID) const {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
auto scratch_ptr = scratch.get_pointer();
const StorageIndex linearLocalThreadId = itemID.get_local_id(0);
StorageIndex nonContractId = is_lhs_vec ? linearLocalThreadId / Properties::LocalThreadSizeC
: linearLocalThreadId % Properties::LocalThreadSizeNC;
StorageIndex contractId = is_lhs_vec ? linearLocalThreadId % Properties::LocalThreadSizeC
: linearLocalThreadId / Properties::LocalThreadSizeNC;
const StorageIndex cGroupSize = itemID.get_group_range(0) / nonContractGroupSize;
const StorageIndex nonContractGroupId =
is_lhs_vec ? itemID.get_group(0) / cGroupSize : itemID.get_group(0) % nonContractGroupSize;
const StorageIndex contractGroupId =
is_lhs_vec ? itemID.get_group(0) % cGroupSize : itemID.get_group(0) / nonContractGroupSize;
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
auto out_ptr = out_res + (IsFinal ? 0 : contractGroupId * nonContractDim);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const StorageIndex nonContractGroupOffset = nonContractGroupId * Properties::TileSizeDimNC;
const StorageIndex contractGroupOffset = contractGroupId * Properties::TileSizeDimC;
auto outScratchIndex = nonContractId + contractId * Properties::LocalThreadSizeNC;
const StorageIndex globalNonContractDimOffset = nonContractGroupOffset + nonContractId;
const StorageIndex globalContractDimOffset = contractGroupOffset + contractId;
auto local_output = scratch_ptr + OutScratchOffset;
const bool is_internal = nonContractDim - nonContractGroupOffset >= Properties::TileSizeDimNC &&
contractDim - contractGroupOffset >= Properties::TileSizeDimC;
is_internal
? compute_panel<true>(itemID, vec, mat, local_output, out_ptr,
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_ON
scratch_ptr, contractGroupOffset,
#endif
nonContractGroupOffset, linearLocalThreadId, contractDim, nonContractDim, contractId,
nonContractId, globalContractDimOffset, globalNonContractDimOffset, outScratchIndex)
: compute_panel<false>(itemID, vec, mat, local_output, out_ptr,
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_ON
scratch_ptr, contractGroupOffset,
#endif
nonContractGroupOffset, linearLocalThreadId, contractDim, nonContractDim, contractId,
nonContractId, globalContractDimOffset, globalNonContractDimOffset, outScratchIndex);
}
template <bool is_internal_block, typename OutPtr>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void compute_panel(
const cl::sycl::nd_item<1> &itemID, const VectorMapper &vec, const TensorMapper &mat, OutScalar *local_output,
OutPtr out_ptr,
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_ON
OutScalar *scratch_ptr, const StorageIndex contractGroupOffset,
#endif
const StorageIndex nonContractGroupOffset, const StorageIndex linearLocalThreadId, StorageIndex contractDim,
StorageIndex nonContractDim, StorageIndex contractId, StorageIndex nonContractId,
StorageIndex globalContractDimOffset, StorageIndex globalNonContractDimOffset, StorageIndex outScratchIndex) {
OutScalar outScalar[Properties::WorkLoadPerThreadNC] = {OutScalar(0)};
// Reading the vector
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_ON
const StorageIndex vectorOffset = contractGroupOffset + linearLocalThreadId;
extract_block<VecBlockProperties, is_internal_block, KFactor,
Properties::LocalThreadSizeNC * Properties::LocalThreadSizeC>(vec, scratch_ptr, linearLocalThreadId,
vectorOffset, contractDim);
itemID.barrier(cl::sycl::access::fence_space::local_space);
auto in_scratch_ptr = scratch_ptr + contractId;
#endif
StorageIndex privateOffsetC = 0;
EIGEN_UNROLL_LOOP
for (StorageIndex i = 0; i < Properties::WorkLoadPerThreadC; i++) {
StorageIndex privateOffsetNC = 0;
bool contract_conds = ((globalContractDimOffset + privateOffsetC) < contractDim);
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_ON
auto vecScalar = *in_scratch_ptr;
#else
auto vecScalar = (check_boundary<is_internal_block>(contract_conds))
? vec(is_lhs_vec ? StorageIndex(0) : globalContractDimOffset + privateOffsetC,
is_lhs_vec ? globalContractDimOffset + privateOffsetC : StorageIndex(0))
: OutScalar(0);
#endif
EIGEN_UNROLL_LOOP
for (StorageIndex j = 0; j < Properties::WorkLoadPerThreadNC; j++) {
auto matScalar = (check_boundary<is_internal_block>(
contract_conds && ((globalNonContractDimOffset + privateOffsetNC) < nonContractDim)))
? mat(is_lhs_vec ? globalContractDimOffset + privateOffsetC
: globalNonContractDimOffset + privateOffsetNC,
is_lhs_vec ? globalNonContractDimOffset + privateOffsetNC
: globalContractDimOffset + privateOffsetC)
: OutScalar(0);
[SYCL-2020] Enabling half precision support for SYCL.
2023-04-27 14:01:11 +01:00
outScalar[j] = ::Eigen::internal::pmadd(matScalar, vecScalar, outScalar[j]);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
privateOffsetNC += Properties::LocalThreadSizeNC;
}
privateOffsetC += Properties::LocalThreadSizeC;
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_ON
in_scratch_ptr += Properties::LocalThreadSizeC;
#endif
}
auto out_scratch_ptr = local_output + outScratchIndex;
// Each block of 16*16 element in shared memory should reduce to 16*1
EIGEN_UNROLL_LOOP
for (StorageIndex j = 0; j < Properties::WorkLoadPerThreadNC; j++) {
*out_scratch_ptr = outScalar[j];
out_scratch_ptr += (Properties::LocalThreadSizeNC * Properties::LocalThreadSizeC);
}
if (is_lhs_vec) {
nonContractId = linearLocalThreadId % Properties::LocalThreadSizeNC;
contractId = linearLocalThreadId / Properties::LocalThreadSizeNC;
outScratchIndex = nonContractId + contractId * Properties::LocalThreadSizeNC;
}
out_scratch_ptr = local_output + outScratchIndex;
EIGEN_UNROLL_LOOP
for (StorageIndex j = 0; j < Properties::WorkLoadPerThreadNC; j++) {
EIGEN_UNROLL_LOOP
for (StorageIndex offset = Properties::LocalThreadSizeC >> 1; offset > 0; offset >>= 1) {
itemID.barrier(cl::sycl::access::fence_space::local_space);
if (contractId < offset) {
StorageIndex myNeigbourId = (Properties::LocalThreadSizeNC * offset);
*out_scratch_ptr += out_scratch_ptr[myNeigbourId];
}
}
// moving to next 16 by 16 block
out_scratch_ptr += (Properties::LocalThreadSizeNC * Properties::LocalThreadSizeC);
}
if (contractId == 0) {
out_scratch_ptr = local_output + nonContractId;
StorageIndex global_final_offset = nonContractGroupOffset + nonContractId;
out_ptr += global_final_offset;
EIGEN_UNROLL_LOOP
for (StorageIndex j = 0; j < Properties::WorkLoadPerThreadNC; j++) {
if (check_boundary<is_internal_block>(global_final_offset < nonContractDim)) {
auto res = *out_scratch_ptr;
*out_ptr = res;
out_ptr += Properties::LocalThreadSizeNC;
}
// moving to next 16 by 16 block to ge the next 16 reduced elements
out_scratch_ptr += (Properties::LocalThreadSizeNC * Properties::LocalThreadSizeC);
if (!(is_internal_block)) global_final_offset += Properties::LocalThreadSizeNC;
}
}
}
template <typename InputBlockProperties, bool is_internal_block, int CFactor, int GroupSize, typename Input,
typename Local>
static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void extract_block(const Input &inpt, Local *local_ptr,
const StorageIndex &linearLocalThreadId,
const StorageIndex &cOffset, const StorageIndex &C) {
local_ptr += InputBlockProperties::c_stride * linearLocalThreadId;
StorageIndex cIndex = cOffset;
for (StorageIndex cId = 0; cId < CFactor / InputBlockProperties::c_stride; cId++) {
if (check_boundary<is_internal_block>(cIndex + InputBlockProperties::c_stride - 1 < C)) {
auto val = read<InputBlockProperties::packet_load, InputBlockProperties::is_coalesced_layout,
InputBlockProperties::is_rhs, typename InputBlockProperties::OutType>(inpt, StorageIndex(0),
cIndex, StorageIndex(1));
write<StorageIndex, 1, data_source::local_mem>(val, local_ptr);
} else {
EIGEN_UNROLL_LOOP
for (StorageIndex i = 0; i < InputBlockProperties::elements_per_access; i++) {
OutScalar val =
(cIndex + i < C)
? read<false, InputBlockProperties::is_coalesced_layout, InputBlockProperties::is_rhs, OutScalar>(
inpt, StorageIndex(0), cIndex + i, StorageIndex(1))
: OutScalar(0);
write<StorageIndex, 1, data_source::local_mem>(val, local_ptr + i);
}
}
local_ptr += InputBlockProperties::c_stride * GroupSize;
cIndex += InputBlockProperties::c_stride * GroupSize;
}
}
};
#endif
#ifndef EIGEN_SYCL_DISABLE_SCALAR
/*!
* \brief GeneralScalarContraction is a template class that provides the scalar value of Tensor -Tensor contraction
* operation, when all the dimensions are contracting dimensions. This Kernel reduces two tensors to an scalar
*
* \tparam OutScalar: determines the output scalar type
*
* \tparam LhsScalar: determines the left-hand-side scalar type
*
* \tparam RhsScalar: determines the right-hand-side scalar type
*
* \tparam OutAccessor: determines the sycl accessor type for out put (please see the sycl-1.2.1 specification
* (https://www.khronos.org/registry/SYCL/specs/sycl-1.2.1.pdf) for accessor definition)
*
* \tparam LhsMapper: determines the tensor contraction mapper type for left-hand-side matrix
*
* \tparam RhsMapper: determines the tensor contraction mapper type for right-hand-side matrix
*
* \tparam StorageIndex: determines the StorageIndex Type
*
* \tparam Vectorizable: determines whether or not the vectorization is enabled for the Eigen expression.
*
* \param scratch: local memory containing tiles of LHS and RHS tensors for each work-group
*
* \param lhs: determines the left-hand-side flattened tensor (tensor mapper)
*
* \param rhs: determines the right-hand-side flattened tensor (tensor mapper)
*
* \param out_res: determines the output tensor containing the contraction result
*
fix various typos
2021-09-22 16:15:06 +00:00
* \param rng: determines the total input data size
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
*/
template <typename OutScalar, typename LhsScalar, typename RhsScalar, typename OutAccessor, typename LhsMapper,
typename RhsMapper, typename StorageIndex, bool Vectorizable>
struct GeneralScalarContraction {
typedef cl::sycl::accessor<OutScalar, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::local> Scratch;
Scratch scratch;
const LhsMapper lhs;
const RhsMapper rhs;
OutAccessor out_res;
const StorageIndex rng;
EIGEN_DEVICE_FUNC GeneralScalarContraction(Scratch scratch_, const LhsMapper lhs_, const RhsMapper rhs_,
OutAccessor out_res_, const StorageIndex rng_)
: scratch(scratch_), lhs(lhs_), rhs(rhs_), out_res(out_res_), rng(rng_) {}
[SYCL-2020 Support] Enabling Intel DPCPP Compiler support to Eigen
2023-01-16 07:04:08 +00:00
EIGEN_DEVICE_FUNC void operator()(cl::sycl::nd_item<1> itemID) const {
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
auto out_ptr = out_res;
OutScalar *scratch_ptr = scratch.get_pointer();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
StorageIndex globalid = itemID.get_global_id(0);
StorageIndex localid = itemID.get_local_id(0);
OutScalar accumulator = OutScalar(0);
for (StorageIndex i = globalid; i < rng; i += itemID.get_global_range(0)) {
[SYCL-2020] Enabling half precision support for SYCL.
2023-04-27 14:01:11 +01:00
accumulator = Eigen::internal::pmadd(lhs(0, i), rhs(i, 0), accumulator);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
}
auto out_scratch_ptr = scratch_ptr + localid;
*out_scratch_ptr = accumulator;
for (StorageIndex offset = itemID.get_local_range(0) >> 1; offset > 0; offset >>= 1) {
itemID.barrier(cl::sycl::access::fence_space::local_space);
if (localid < offset) {
*out_scratch_ptr = (accumulator += out_scratch_ptr[offset]);
}
}
if (localid == 0) {
out_ptr[itemID.get_group(0)] = accumulator;
}
}
};
#endif
} // namespace internal
} // namespace TensorSycl
template <typename Indices, typename LeftArgType, typename RightArgType, typename OutputKernelType>
struct TensorEvaluator<const TensorContractionOp<Indices, LeftArgType, RightArgType, OutputKernelType>,
Eigen::SyclDevice>
: public TensorContractionEvaluatorBase<TensorEvaluator<
const TensorContractionOp<Indices, LeftArgType, RightArgType, OutputKernelType>, Eigen::SyclDevice>> {
Fix TensorContractionOp evaluators for GPU and SYCL
2018-07-17 14:09:37 -07:00
static_assert(std::is_same<OutputKernelType, const NoOpOutputKernel>::value,
"SYCL tensor contraction does not support output kernels.");
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef Eigen::SyclDevice Device;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
Fix TensorContractionOp evaluators for GPU and SYCL
2018-07-17 14:09:37 -07:00
typedef TensorEvaluator<const TensorContractionOp<Indices, LeftArgType, RightArgType, OutputKernelType>, Device> Self;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
typedef TensorContractionEvaluatorBase<Self> Base;
Fix TensorContractionOp evaluators for GPU and SYCL
2018-07-17 14:09:37 -07:00
typedef TensorContractionOp<Indices, LeftArgType, RightArgType, OutputKernelType> XprType;
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
typedef std::remove_const_t<typename XprType::Scalar> Scalar;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef typename XprType::Index StorageIndex;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
typedef typename XprType::CoeffReturnType CoeffReturnType;
typedef typename PacketType<CoeffReturnType, Device>::type PacketReturnType;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef typename Base::Storage Storage;
typedef typename Base::EvaluatorPointerType EvaluatorPointerType;
struct TripleDim {
const StorageIndex M;
const StorageIndex N;
const StorageIndex K;
TripleDim(const StorageIndex M_, const StorageIndex N_, const StorageIndex K_) : M(M_), N(N_), K(K_) {}
};
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
enum {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
PacketAccess = (PacketType<CoeffReturnType, Device>::size > 1),
BlockAccess = false,
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
};
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
static constexpr int Layout = TensorEvaluator<LeftArgType, Device>::Layout;
static const class members turned into constexpr
2022-04-04 17:33:33 +00:00
static constexpr int LDims = Base::LDims;
static constexpr int RDims = Base::RDims;
static constexpr int ContractDims = Base::ContractDims;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef array<StorageIndex, LDims> left_dim_mapper_t;
typedef array<StorageIndex, RDims> right_dim_mapper_t;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef array<StorageIndex, ContractDims> contract_t;
typedef array<StorageIndex, LDims - ContractDims> left_nocontract_t;
typedef array<StorageIndex, RDims - ContractDims> right_nocontract_t;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
static const class members turned into constexpr
2022-04-04 17:33:33 +00:00
static constexpr int NumDims = LDims + RDims - 2 * ContractDims;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef DSizes<StorageIndex, NumDims> Dimensions;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef TensorEvaluator<typename Base::EvalLeftArgType, Device> LeftEvaluator;
typedef TensorEvaluator<typename Base::EvalRightArgType, Device> RightEvaluator;
Replace Eigen type metaprogramming with corresponding std types and make use of alias templates
2022-03-16 16:43:40 +00:00
typedef std::remove_const_t<typename LeftEvaluator::CoeffReturnType> LhsScalar;
typedef std::remove_const_t<typename RightEvaluator::CoeffReturnType> RhsScalar;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
typedef typename LeftEvaluator::Dimensions LeftDimensions;
typedef typename RightEvaluator::Dimensions RightDimensions;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
template <bool lhs_inner_dim_contiguous, bool rhs_inner_dim_contiguous, bool rhs_inner_dim_reordered>
struct input_mapper_propertis {
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
static constexpr bool is_lhs_matrix = (LDims == 2 && ContractDims == 1) || lhs_inner_dim_contiguous;
static constexpr bool is_rhs_matrix =
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
(RDims == 2 && ContractDims == 1) || (rhs_inner_dim_contiguous && !rhs_inner_dim_reordered);
};
Fix calls to device functions from host code
2021-05-11 22:47:49 +00:00
TensorEvaluator(const XprType &op, const Device &device) : Base(op, device) {}
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
// We need to redefine this method to make nvcc happy
Fix calls to device functions from host code
2021-05-11 22:47:49 +00:00
EIGEN_STRONG_INLINE bool evalSubExprsIfNeeded(typename Base::EvaluatorPointerType data) {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
this->m_leftImpl.evalSubExprsIfNeeded(NULL);
this->m_rightImpl.evalSubExprsIfNeeded(NULL);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
if (!data) {
this->m_result = this->m_device.get(
static_cast<Scalar *>(this->m_device.allocate_temp(this->dimensions().TotalSize() * sizeof(Scalar))));
data = this->m_result;
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
evalToSycl(data);
return (this->m_result != NULL);
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
const Eigen::SyclDevice &device() const { return this->m_device; }
void evalToSycl(typename Base::EvaluatorPointerType buffer) const {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
if (this->m_lhs_inner_dim_contiguous) {
if (this->m_rhs_inner_dim_contiguous) {
if (this->m_rhs_inner_dim_reordered) {
evalTyped<true, true, true, Unaligned>(buffer);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
evalTyped<true, true, false, Unaligned>(buffer);
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
if (this->m_rhs_inner_dim_reordered) {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
evalTyped<true, false, true, Unaligned>(buffer);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
evalTyped<true, false, false, Unaligned>(buffer);
}
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
if (this->m_rhs_inner_dim_contiguous) {
if (this->m_rhs_inner_dim_reordered) {
evalTyped<false, true, true, Unaligned>(buffer);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
evalTyped<false, true, false, Unaligned>(buffer);
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
if (this->m_rhs_inner_dim_reordered) {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
evalTyped<false, false, true, Unaligned>(buffer);
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
evalTyped<false, false, false, Unaligned>(buffer);
}
}
}
}
template <bool lhs_inner_dim_contiguous, bool rhs_inner_dim_contiguous, bool rhs_inner_dim_reordered, int Alignment>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
void evalTyped(typename Base::EvaluatorPointerType buffer) const {
const auto triple_dim = TripleDim{this->m_i_size, this->m_j_size, this->m_k_size};
typedef internal::TensorContractionInputMapper<
LhsScalar, StorageIndex, internal::Lhs, LeftEvaluator, left_nocontract_t, contract_t,
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
PacketType<CoeffReturnType, Device>::size, lhs_inner_dim_contiguous, false, Unaligned, MakePointer>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
LhsMapper;
typedef internal::TensorContractionInputMapper<RhsScalar, StorageIndex, internal::Rhs, RightEvaluator,
right_nocontract_t, contract_t,
PacketType<CoeffReturnType, Device>::size, rhs_inner_dim_contiguous,
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
rhs_inner_dim_reordered, Unaligned, MakePointer>
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
RhsMapper;
// initialize data mappers
LhsMapper lhs(this->m_leftImpl, this->m_left_nocontract_strides, this->m_i_strides,
this->m_left_contracting_strides, this->m_k_strides);
RhsMapper rhs(this->m_rightImpl, this->m_right_nocontract_strides, this->m_j_strides,
this->m_right_contracting_strides, this->m_k_strides);
#ifndef EIGEN_SYCL_DISABLE_SCALAR
if (triple_dim.M == 1 && triple_dim.N == 1) {
launchSC(buffer, lhs, rhs, triple_dim.K);
} else
#endif
#ifndef EIGEN_SYCL_DISABLE_GEMV
if (triple_dim.M != 1 && triple_dim.N == 1) {
LaunchVT<false>(buffer, rhs, lhs, triple_dim.M, triple_dim.K);
} else if (triple_dim.M == 1 && triple_dim.N != 1) {
LaunchVT<true>(buffer, lhs, rhs, triple_dim.N, triple_dim.K);
} else // This is equivalent of if (m!=1 && n!=1)
#endif
{
typedef input_mapper_propertis<lhs_inner_dim_contiguous, rhs_inner_dim_contiguous, rhs_inner_dim_reordered>
inpt_mapper_properties;
#ifndef EIGEN_SYCL_DISABLE_SKINNY
bool skinny = false;
auto platform_name = this->device().getPlatformName();
// This is based on empirical calculation for AMD r9-nano and Fiji
if (platform_name.find("AMD") == 0) {
skinny = (triple_dim.M < triple_dim.K || triple_dim.N < triple_dim.K) &&
((triple_dim.M < 1024 && triple_dim.N < 1024) ||
(uint64_t(triple_dim.M * triple_dim.N) < uint64_t(triple_dim.K)));
} else {
skinny = (((std::max(triple_dim.K, triple_dim.N) / std::min(triple_dim.K, triple_dim.N)) > 100) ||
((std::max(triple_dim.K, triple_dim.M) / std::min(triple_dim.K, triple_dim.M)) > 100) ||
((std::max(triple_dim.N, triple_dim.M) / std::min(triple_dim.N, triple_dim.M)) > 100));
}
if (skinny)
adjustTT<true, inpt_mapper_properties>(buffer, lhs, rhs, triple_dim);
else
#endif // EIGEN_SYCL_DISABLE_SKINNY
adjustTT<false, inpt_mapper_properties>(buffer, lhs, rhs, triple_dim);
}
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
template <bool skinny, typename input_mapper_properties, typename LhsMapper, typename RhsMapper>
void EIGEN_ALWAYS_INLINE adjustTT(EvaluatorPointerType buffer, const LhsMapper &lhs, const RhsMapper &rhs,
const TripleDim &triple_dim) const {
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_ON
if (device().has_local_memory()) {
typedef TensorSycl::internal::TTPanelSize<CoeffReturnType, StorageIndex, 4, 4, 16> PanelParameters;
launchTT<TensorSycl::internal::contraction_type::local, skinny, input_mapper_properties, PanelParameters>(
buffer, lhs, rhs, triple_dim);
}
#endif
#ifdef EIGEN_SYCL_LOCAL_MEM_UNSET_OR_OFF
if (!(device().has_local_memory())) {
typedef TensorSycl::internal::TTPanelSize<CoeffReturnType, StorageIndex, 4, 4, 4> PanelParameters;
launchTT<TensorSycl::internal::contraction_type::no_local, skinny, input_mapper_properties, PanelParameters>(
buffer, lhs, rhs, triple_dim);
}
#endif
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
template <TensorSycl::internal::contraction_type ct, bool skinny, typename input_mapper_properties,
typename Properties, typename LhsMapper, typename RhsMapper>
void launchTT(EvaluatorPointerType buffer, const LhsMapper &lhs, const RhsMapper &rhs,
const TripleDim &triple_dim) const {
const StorageIndex roundUpM = Eigen::TensorSycl::internal::roundUp(triple_dim.M, Properties::TileSizeDimM);
const StorageIndex roundUpN = Eigen::TensorSycl::internal::roundUp(triple_dim.N, Properties::TileSizeDimN);
const StorageIndex groupSizeM = roundUpM / Properties::TileSizeDimM;
const StorageIndex groupSizeN = roundUpN / Properties::TileSizeDimN;
const StorageIndex roundUpK = Eigen::TensorSycl::internal::roundUp(triple_dim.K, Properties::TileSizeDimK);
StorageIndex totalTilesK = roundUpK / Properties::TileSizeDimK;
StorageIndex groupSizeK =
skinny
? std::max(std::min(totalTilesK,
(StorageIndex)(device().getPowerOfTwo(device().getNumSyclMultiProcessors(), true) * 4) /
(groupSizeM * groupSizeN)),
StorageIndex(1))
: StorageIndex(1);
const StorageIndex numTilesPerGroup = Eigen::TensorSycl::internal::roundUp(totalTilesK, groupSizeK) / groupSizeK;
const StorageIndex totalGroupSize = groupSizeM * groupSizeN * groupSizeK;
const StorageIndex localRange = Properties::LocalThreadSizeM * Properties::LocalThreadSizeN;
const StorageIndex globalRange = totalGroupSize * localRange;
const StorageIndex scratchSize = (ct == TensorSycl::internal::contraction_type::local)
? ((Properties::DoubleBuffer + 1) *
(Properties::TileSizeDimM + Properties::BC) * (Properties::TileSizeDimK)) +
((Properties::DoubleBuffer + 1) * (Properties::TileSizeDimK) *
(Properties::TileSizeDimN + Properties::BC))
: StorageIndex(1);
auto thread_range = cl::sycl::nd_range<1>(cl::sycl::range<1>(globalRange), cl::sycl::range<1>(localRange));
if (groupSizeK == 1) {
typedef TensorSycl::internal::TensorContractionKernel<CoeffReturnType, LhsScalar, RhsScalar, EvaluatorPointerType,
LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim,
PacketAccess, input_mapper_properties, true, ct>
ContractKernelName;
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
device()
.template binary_kernel_launcher<CoeffReturnType, ContractKernelName>(
lhs, rhs, buffer, thread_range, scratchSize, groupSizeM, groupSizeN, numTilesPerGroup, triple_dim)
.wait();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
} else {
typedef TensorSycl::internal::TensorContractionKernel<CoeffReturnType, LhsScalar, RhsScalar, EvaluatorPointerType,
LhsMapper, RhsMapper, StorageIndex, Properties, TripleDim,
PacketAccess, input_mapper_properties, false, ct>
ContractKernelName;
CoeffReturnType *temp_pointer = static_cast<CoeffReturnType *>(
device().allocate_temp(triple_dim.M * triple_dim.N * groupSizeK * sizeof(CoeffReturnType)));
EvaluatorPointerType tmp_global_accessor = device().get(temp_pointer);
device()
.template binary_kernel_launcher<CoeffReturnType, ContractKernelName>(
lhs, rhs, tmp_global_accessor, thread_range, scratchSize, groupSizeM, groupSizeN, numTilesPerGroup,
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
triple_dim)
.wait();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef Eigen::internal::SumReducer<CoeffReturnType> Op;
auto op = Op();
typedef TensorSycl::internal::SecondStepPartialReduction<CoeffReturnType, StorageIndex, EvaluatorPointerType,
EvaluatorPointerType, Op>
ReductionKernel;
device()
.template unary_kernel_launcher<CoeffReturnType, ReductionKernel>(
tmp_global_accessor, buffer,
cl::sycl::nd_range<1>(cl::sycl::range<1>(StorageIndex(
Eigen::TensorSycl::internal::roundUp(triple_dim.M * triple_dim.N, localRange))),
cl::sycl::range<1>(localRange)),
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
StorageIndex(1), op, StorageIndex(triple_dim.M * triple_dim.N), groupSizeK)
.wait();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
device().deallocate_temp(temp_pointer);
}
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#ifndef EIGEN_SYCL_DISABLE_GEMV
template <bool is_lhs_vec, typename VectorMapper, typename TensorMapper, typename StorageIndex>
void EIGEN_ALWAYS_INLINE LaunchVT(EvaluatorPointerType buffer, const VectorMapper &vec, const TensorMapper &mat,
StorageIndex NC, StorageIndex C) const {
const StorageIndex nonContractDim = NC;
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex NCFactor = 1;
constexpr StorageIndex CFactor = 1;
constexpr StorageIndex NCWindow = 16;
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef Eigen::TensorSycl::internal::TVPanelSize<CoeffReturnType, StorageIndex, NCWindow, CFactor, NCFactor>
Properties;
const StorageIndex roundUpC = Eigen::TensorSycl::internal::roundUp(C, Properties::TileSizeDimC);
const StorageIndex cNumGroups = roundUpC / (Properties::LocalThreadSizeC * Properties::WorkLoadPerThreadC);
const StorageIndex roundUpNC = Eigen::TensorSycl::internal::roundUp(nonContractDim, Properties::TileSizeDimNC);
const StorageIndex nCNumGroups = roundUpNC / (Properties::LocalThreadSizeNC * Properties::WorkLoadPerThreadNC);
const StorageIndex globalRange =
(roundUpNC / (Properties::WorkLoadPerThreadNC)) * (roundUpC / (Properties::WorkLoadPerThreadC));
const StorageIndex localRange = Properties::LocalThreadSizeNC * Properties::LocalThreadSizeC;
const StorageIndex scratchSize =
(Properties::WorkLoadPerThreadNC + CFactor) * Properties::LocalThreadSizeC * Properties::LocalThreadSizeNC;
auto thread_range = cl::sycl::nd_range<1>(cl::sycl::range<1>(globalRange), cl::sycl::range<1>(localRange));
if (cNumGroups > 1) {
typedef Eigen::TensorSycl::internal::GeneralVectorTensor<CoeffReturnType, EvaluatorPointerType, VectorMapper,
TensorMapper, StorageIndex, Properties, CFactor, false,
is_lhs_vec, false>
ContractKernelName;
CoeffReturnType *temp_pointer =
static_cast<CoeffReturnType *>(device().allocate_temp(nonContractDim * cNumGroups * sizeof(CoeffReturnType)));
EvaluatorPointerType tmp_global_accessor = device().get(temp_pointer);
Converting all parallel for lambda to functor in order to prevent kernel duplication name error; adding tensorConcatinationOp backend for sycl.
2016-12-16 19:46:45 +00:00
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
device()
.template binary_kernel_launcher<CoeffReturnType, ContractKernelName>(
vec, mat, tmp_global_accessor, thread_range, scratchSize, nCNumGroups, nonContractDim, C)
.wait();
Converting all parallel for lambda to functor in order to prevent kernel duplication name error; adding tensorConcatinationOp backend for sycl.
2016-12-16 19:46:45 +00:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef Eigen::internal::SumReducer<CoeffReturnType> Op;
typedef TensorSycl::internal::SecondStepPartialReduction<CoeffReturnType, StorageIndex, EvaluatorPointerType,
EvaluatorPointerType, Op>
ReductionKernel;
Converting all parallel for lambda to functor in order to prevent kernel duplication name error; adding tensorConcatinationOp backend for sycl.
2016-12-16 19:46:45 +00:00
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
device()
.template unary_kernel_launcher<CoeffReturnType, ReductionKernel>(
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
tmp_global_accessor, buffer,
cl::sycl::nd_range<1>(
cl::sycl::range<1>(Eigen::TensorSycl::internal::roundUp(nonContractDim, localRange)),
cl::sycl::range<1>(localRange)),
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
StorageIndex(1), Op(), nonContractDim, cNumGroups)
.wait();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
device().deallocate_temp(temp_pointer);
} else {
typedef Eigen::TensorSycl::internal::GeneralVectorTensor<CoeffReturnType, EvaluatorPointerType, VectorMapper,
TensorMapper, StorageIndex, Properties, CFactor, false,
is_lhs_vec, true>
ContractKernelName;
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
device()
.template binary_kernel_launcher<CoeffReturnType, ContractKernelName>(
vec, mat, buffer, thread_range, scratchSize, nCNumGroups, nonContractDim, C)
.wait();
Converting all parallel for lambda to functor in order to prevent kernel duplication name error; adding tensorConcatinationOp backend for sycl.
2016-12-16 19:46:45 +00:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
}
#endif
Converting all parallel for lambda to functor in order to prevent kernel duplication name error; adding tensorConcatinationOp backend for sycl.
2016-12-16 19:46:45 +00:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#ifndef EIGEN_SYCL_DISABLE_SCALAR
template <typename LhsMapper, typename RhsMapper>
EIGEN_ALWAYS_INLINE void launchSC(EvaluatorPointerType buffer, const LhsMapper &lhs, const RhsMapper &rhs,
StorageIndex K) const {
EIGEN_STATIC_ASSERT(!((EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1) &
(EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1 - 1)),
"The Local thread size must be a power of 2 for the reduction "
"operation");
Replace instances of EIGEN_CONSTEXPR macro
2024-11-15 21:11:01 -08:00
constexpr StorageIndex local_range = EIGEN_SYCL_LOCAL_THREAD_DIM0 * EIGEN_SYCL_LOCAL_THREAD_DIM1;
Avoid using #define since they can conflict with user code
2016-12-14 19:49:15 -08:00
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
// Here we force the code not to be more than 2-step reduction: Our empirical research shows that if each thread
// reduces at least 512 elementss individually, we get better performance.
const StorageIndex num_work_group = ((K + (512 * local_range - 1)) / (512 * local_range) > 1 ? local_range : 1);
const StorageIndex global_range = num_work_group * local_range;
typedef Eigen::TensorSycl::internal::GeneralScalarContraction<
CoeffReturnType, LhsScalar, RhsScalar, EvaluatorPointerType, LhsMapper, RhsMapper, StorageIndex, false>
ContractKernelName;
auto thread_range = cl::sycl::nd_range<1>(cl::sycl::range<1>(global_range), cl::sycl::range<1>(local_range));
if (num_work_group > 1) {
CoeffReturnType *temp_pointer =
static_cast<CoeffReturnType *>(device().allocate_temp(num_work_group * sizeof(CoeffReturnType)));
EvaluatorPointerType tmp_global_accessor = device().get(temp_pointer);
device()
.template binary_kernel_launcher<CoeffReturnType, ContractKernelName>(lhs, rhs, tmp_global_accessor,
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
thread_range, local_range, K)
.wait();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
typedef Eigen::internal::SumReducer<CoeffReturnType> Op;
typedef TensorSycl::internal::SecondStepFullReducer<CoeffReturnType, Op, EvaluatorPointerType,
EvaluatorPointerType, StorageIndex, local_range>
GenericRKernel;
device()
.template unary_kernel_launcher<CoeffReturnType, GenericRKernel>(
tmp_global_accessor, buffer,
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
cl::sycl::nd_range<1>(cl::sycl::range<1>(local_range), cl::sycl::range<1>(local_range)), local_range,
Op())
.wait();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
device().deallocate_temp(temp_pointer);
} else {
device()
.template binary_kernel_launcher<CoeffReturnType, ContractKernelName>(lhs, rhs, buffer, thread_range,
[SYCL-2020] Enabling USM support for SYCL. SYCL-1.2.1 did not have support for USM.
2023-05-05 17:30:36 +00:00
local_range, K)
.wait();
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
}
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
}
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
#endif
Adding tensor contraction operation backend for Sycl; adding test for contractionOp sycl backend; adding temporary solution to prevent memory leak in buffer; cleaning up cxx11_tensor_buildins_sycl.h
2016-12-14 15:30:37 +00:00
Fix calls to device functions from host code
2021-05-11 22:47:49 +00:00
EIGEN_STRONG_INLINE void cleanup() {
[SYCL] Rebasing the SYCL support branch on top of the Einge upstream master branch. * Unifying all loadLocalTile from lhs and rhs to an extract_block function. * Adding get_tensor operation which was missing in TensorContractionMapper. * Adding the -D method missing from cmake for Disable_Skinny Contraction operation. * Wrapping all the indices in TensorScanSycl into Scan parameter struct. * Fixing typo in Device SYCL * Unifying load to private register for tall/skinny no shared * Unifying load to vector tile for tensor-vector/vector-tensor operation * Removing all the LHS/RHS class for extracting data from global * Removing Outputfunction from TensorContractionSkinnyNoshared. * Combining the local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining the no-local memory version of tall/skinny and normal tensor contraction into one kernel. * Combining General Tensor-Vector and VectorTensor contraction into one kernel. * Making double buffering optional for Tensor contraction when local memory is version is used. * Modifying benchmark to accept custom Reduction Sizes * Disabling AVX optimization for SYCL backend on the host to allow SSE optimization to the host * Adding Test for SYCL * Modifying SYCL CMake
2019-11-28 10:08:54 +00:00
this->m_leftImpl.cleanup();
this->m_rightImpl.cleanup();
if (this->m_result) {
this->m_device.deallocate_temp(this->m_result);
this->m_result = NULL;
}
}
};
} // namespace Eigen
#endif // EIGEN_CXX11_TENSOR_TENSOR_CONTRACTION_SYCL_H
Reference in New Issue Copy Permalink