<commit_before>// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// <NAME>    Codeplay Software Ltd.
// <NAME>  Codeplay Software Ltd.
// <NAME>  Codeplay Software Ltd.
// Contact: <<EMAIL>>
// Copyright (C) 2016 <NAME> <<EMAIL>>

//
// 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/.

#if defined(EIGEN_USE_SYCL) && !defined(EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H)
#define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H

template <typename Scalar, size_t Align = EIGEN_MAX_ALIGN_BYTES, class Allocator = std::allocator<Scalar>>
struct SyclAllocator {
  typedef Scalar value_type;
  typedef typename std::allocator_traits<Allocator>::pointer pointer;
  typedef typename std::allocator_traits<Allocator>::size_type size_type;

  SyclAllocator( ){};
  Scalar* allocate(std::size_t elements) { return static_cast<Scalar*>(aligned_alloc(Align, elements)); }
  void deallocate(Scalar * p, std::size_t size) { EIGEN_UNUSED_VARIABLE(size); free(p); }
};

namespace Eigen {

  #define ConvertToActualTypeSycl(Scalar, buf_acc) reinterpret_cast<typename cl::sycl::global_ptr<Scalar>::pointer_t>((&(*buf_acc.get_pointer())))
  #define ConvertToActualSyclOffset(Scalar, offset) offset/sizeof(Scalar)


  template <typename Scalar, typename read_accessor, typename write_accessor> class MemCopyFunctor {
  public:
    MemCopyFunctor(read_accessor src_acc, write_accessor dst_acc, size_t rng, size_t i, size_t offset) : m_src_acc(src_acc), m_dst_acc(dst_acc), m_rng(rng), m_i(i), m_offset(offset) {}

    void operator()(cl::sycl::nd_item<1> itemID) {
      auto src_ptr = ConvertToActualTypeSycl(Scalar, m_src_acc);
      auto dst_ptr = ConvertToActualTypeSycl(Scalar, m_dst_acc);
      auto globalid = itemID.get_global_linear_id();
      if (globalid < m_rng) {
        dst_ptr[globalid + m_i] = src_ptr[globalid + m_offset];
      }
    }

  private:
    read_accessor m_src_acc;
    write_accessor m_dst_acc;
    size_t m_rng;
    size_t m_i;
    size_t m_offset;
  };

template<typename AccType>
  struct memsetkernelFunctor{
   AccType m_acc;
   const ptrdiff_t buff_offset;
   const size_t m_rng, m_c;
   memsetkernelFunctor(AccType acc, const ptrdiff_t buff_offset_, const size_t rng, const size_t c):m_acc(acc),  buff_offset(buff_offset_), m_rng(rng), m_c(c){}
   void operator()(cl::sycl::nd_item<1> itemID) {
     auto globalid=itemID.get_global_linear_id();
     if (globalid< m_rng) m_acc[globalid + buff_offset] = m_c;
   }

  };

struct memsetCghFunctor{
  cl::sycl::buffer<uint8_t, 1, SyclAllocator<uint8_t, EIGEN_MAX_ALIGN_BYTES> >& m_buf;
  const ptrdiff_t& buff_offset;
  const size_t& rng , GRange, tileSize;
  const int  &c;
  memsetCghFunctor(cl::sycl::buffer<uint8_t, 1, SyclAllocator<uint8_t, EIGEN_MAX_ALIGN_BYTES> >& buff,  const ptrdiff_t& buff_offset_, const size_t& rng_,  const size_t& GRange_,  const size_t& tileSize_, const int& c_)
  :m_buf(buff), buff_offset(buff_offset_), rng(rng_), GRange(GRange_), tileSize(tileSize_), c(c_){}

  void operator()(cl::sycl::handler &cgh) const {
    auto buf_acc = m_buf.template get_access<cl::sycl::access::mode::write, cl::sycl::access::target::global_buffer>(cgh);
    typedef decltype(buf_acc) AccType;
    cgh.parallel_for(cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), memsetkernelFunctor<AccType>(buf_acc, buff_offset, rng, c));
  }
};

 //get_devices returns all the available opencl devices. Either use device_selector or exclude devices that computecpp does not support (AMD OpenCL for CPU  and intel GPU)
EIGEN_STRONG_INLINE auto get_sycl_supported_devices()->decltype(cl::sycl::device::get_devices()){
std::vector<cl::sycl::device> supported_devices;
auto plafrom_list =cl::sycl::platform::get_platforms();
for(const auto& platform : plafrom_list){
  auto device_list = platform.get_devices();
  auto platform_name =platform.template get_info<cl::sycl::info::platform::name>();
  std::transform(platform_name.begin(), platform_name.end(), platform_name.begin(), ::tolower);
  for(const auto& device : device_list){
    auto vendor = device.template get_info<cl::sycl::info::device::vendor>();
    std::transform(vendor.begin(), vendor.end(), vendor.begin(), ::tolower);
    bool unsuported_condition = (device.is_cpu() && platform_name.find("amd")!=std::string::npos && vendor.find("apu") == std::string::npos) ||
    (device.is_gpu() && platform_name.find("intel")!=std::string::npos);
    if(!unsuported_condition){<commit_msg>Restore misplaced comment<commit_after>// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// <NAME>    Codeplay Software Ltd.
// <NAME>  Codeplay Software Ltd.
// <NAME>  Codeplay Software Ltd.
// Contact: <<EMAIL>>
// Copyright (C) 2016 <NAME> <<EMAIL>>

//
// 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/.

#if defined(EIGEN_USE_SYCL) && !defined(EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H)
#define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H

template <typename Scalar, size_t Align = EIGEN_MAX_ALIGN_BYTES, class Allocator = std::allocator<Scalar>>
struct SyclAllocator {
  typedef Scalar value_type;
  typedef typename std::allocator_traits<Allocator>::pointer pointer;
  typedef typename std::allocator_traits<Allocator>::size_type size_type;

  SyclAllocator( ){};
  Scalar* allocate(std::size_t elements) { return static_cast<Scalar*>(aligned_alloc(Align, elements)); }
  void deallocate(Scalar * p, std::size_t size) { EIGEN_UNUSED_VARIABLE(size); free(p); }
};

namespace Eigen {

  #define ConvertToActualTypeSycl(Scalar, buf_acc) reinterpret_cast<typename cl::sycl::global_ptr<Scalar>::pointer_t>((&(*buf_acc.get_pointer())))
  #define ConvertToActualSyclOffset(Scalar, offset) offset/sizeof(Scalar)


  template <typename Scalar, typename read_accessor, typename write_accessor> class MemCopyFunctor {
  public:
    MemCopyFunctor(read_accessor src_acc, write_accessor dst_acc, size_t rng, size_t i, size_t offset) : m_src_acc(src_acc), m_dst_acc(dst_acc), m_rng(rng), m_i(i), m_offset(offset) {}

    void operator()(cl::sycl::nd_item<1> itemID) {
      auto src_ptr = ConvertToActualTypeSycl(Scalar, m_src_acc);
      auto dst_ptr = ConvertToActualTypeSycl(Scalar, m_dst_acc);
      auto globalid = itemID.get_global_linear_id();
      if (globalid < m_rng) {
        dst_ptr[globalid + m_i] = src_ptr[globalid + m_offset];
      }
    }

  private:
    read_accessor m_src_acc;
    write_accessor m_dst_acc;
    size_t m_rng;
    size_t m_i;
    size_t m_offset;
  };

template<typename AccType>
  struct memsetkernelFunctor{
   AccType m_acc;
   const ptrdiff_t buff_offset;
   const size_t m_rng, m_c;
   memsetkernelFunctor(AccType acc, const ptrdiff_t buff_offset_, const size_t rng, const size_t c):m_acc(acc),  buff_offset(buff_offset_), m_rng(rng), m_c(c){}
   void operator()(cl::sycl::nd_item<1> itemID) {
     auto globalid=itemID.get_global_linear_id();
     if (globalid< m_rng) m_acc[globalid + buff_offset] = m_c;
   }

  };

struct memsetCghFunctor{
  cl::sycl::buffer<uint8_t, 1, SyclAllocator<uint8_t, EIGEN_MAX_ALIGN_BYTES> >& m_buf;
  const ptrdiff_t& buff_offset;
  const size_t& rng , GRange, tileSize;
  const int  &c;
  memsetCghFunctor(cl::sycl::buffer<uint8_t, 1, SyclAllocator<uint8_t, EIGEN_MAX_ALIGN_BYTES> >& buff,  const ptrdiff_t& buff_offset_, const size_t& rng_,  const size_t& GRange_,  const size_t& tileSize_, const int& c_)
  :m_buf(buff), buff_offset(buff_offset_), rng(rng_), GRange(GRange_), tileSize(tileSize_), c(c_){}

  void operator()(cl::sycl::handler &cgh) const {
    auto buf_acc = m_buf.template get_access<cl::sycl::access::mode::write, cl::sycl::access::target::global_buffer>(cgh);
    typedef decltype(buf_acc) AccType;
    cgh.parallel_for(cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), memsetkernelFunctor<AccType>(buf_acc, buff_offset, rng, c));
  }
};

//get_devices returns all the available opencl devices. Either use device_selector or exclude devices that computecpp does not support (AMD OpenCL for CPU  and intel GPU)
EIGEN_STRONG_INLINE auto get_sycl_supported_devices()->decltype(cl::sycl::device::get_devices()){
std::vector<cl::sycl::device> supported_devices;
auto plafrom_list =cl::sycl::platform::get_platforms();
for(const auto& platform : plafrom_list){
  auto device_list = platform.get_devices();
  auto platform_name =platform.template get_info<cl::sycl::info::platform::name>();
  std::transform(platform_name.begin(), platform_name.end(), platform_name.begin(), ::tolower);
  for(const auto& device : device_list){
    auto vendor = device.template get_info<cl::sycl::info::device::vendor>();
    std::transform(vendor.begin(), vendor.end(), vendor.begin(), ::tolower);
    bool unsuported_condition = (device.is_cpu() && platform_name.find("amd")!=std::string::npos && vendor.find("apu") == std::string::npos) ||
    (device.is_gpu() && platform_name.find("intel")!=std::string::npos);
    if(!unsuported_condition){
{'fraction_non_alphanumeric': 0.13468814182250569, 'fraction_numerical': 0.00497243541238785, 'mean_word_length': 5.823008849557522, 'pattern_counts': {'":': 0, '<': 69, '<?xml version=': 0, '>': 67, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '4622771', 'n_tokens_mistral': 3364, 'n_tokens_neox': 3270, 'n_words': 693}