tensor_contract_sycl_bench.cc File Reference

#include <CL/sycl.hpp>
#include <fstream>
#include <iostream>
#include <chrono>
#include <ctime>
#include <unsupported/Eigen/CXX11/Tensor>

Macros
#define	EIGEN_BENCH_CONTRACT_SYCL
#define	EIGEN_TEST_NO_LONGDOUBLE
#define	EIGEN_TEST_NO_COMPLEX
#define	EIGEN_DEFAULT_DENSE_INDEX_TYPE   int64_t

Functions
std::ofstream	out ("Result.txt")
std::chrono::time_point< std::chrono::system_clock >	get_time ()
template<typename Start , typename End , typename TensorIndex >
void	finalizeBenchmark (Start start, End end, TensorIndex m_, TensorIndex k_, TensorIndex n_, TensorIndex num_iters, std::string name)
template<typename T , typename Device , typename TensorIndex >
void	contraction (const Device &device_, TensorIndex num_iters, TensorIndex m_, TensorIndex k_, TensorIndex n_)
template<typename T , typename Device , typename TensorIndex >
void	contractionRowMajor (const Device &device_, TensorIndex num_iters, TensorIndex m_, TensorIndex k_, TensorIndex n_)
template<typename T , typename Device , typename TensorIndex >
void	contractionAT (const Device &device_, TensorIndex num_iters, TensorIndex m_, TensorIndex k_, TensorIndex n_)
template<typename T , typename Device , typename TensorIndex >
void	contractionBT (const Device &device_, TensorIndex num_iters, TensorIndex m_, TensorIndex k_, TensorIndex n_)
template<typename T , typename Device , typename TensorIndex >
void	contractionABT (const Device &device_, TensorIndex num_iters, TensorIndex m_, TensorIndex k_, TensorIndex n_)
int	main ()

Macro Definition Documentation

EIGEN_BENCH_CONTRACT_SYCL

#define EIGEN_BENCH_CONTRACT_SYCL

EIGEN_DEFAULT_DENSE_INDEX_TYPE

#define EIGEN_DEFAULT_DENSE_INDEX_TYPE   int64_t

EIGEN_TEST_NO_COMPLEX

#define EIGEN_TEST_NO_COMPLEX

EIGEN_TEST_NO_LONGDOUBLE

#define EIGEN_TEST_NO_LONGDOUBLE

Function Documentation

contraction()

template<typename T , typename Device , typename TensorIndex >

void contraction	(	const Device &	device_,
		TensorIndex	num_iters,
		TensorIndex	m_,
		TensorIndex	k_,
		TensorIndex	n_
	)

                                                                                                               {
  T* a_;
  T* b_;
  T* c_;
  a_ = (T*)device_.allocate(m_ * k_ * sizeof(T));
  b_ = (T*)device_.allocate(k_ * n_ * sizeof(T));
  c_ = (T*)device_.allocate(m_ * n_ * sizeof(T));
 
  // Initialize the content of the memory pools to prevent asan from
  // complaining.
  device_.fill(a_, a_ + (m_ * k_), T(12));
  device_.fill(b_, b_ + (k_ * n_), T(23));
  device_.fill(c_, c_ + (m_ * n_), T(31));
 
  Eigen::array<TensorIndex, 2> sizeA;
  sizeA[0] = m_;
  sizeA[1] = k_;
  Eigen::array<TensorIndex, 2> sizeB;
  sizeB[0] = k_;
  sizeB[1] = n_;
  Eigen::array<TensorIndex, 2> sizeC;
  sizeC[0] = m_;
  sizeC[1] = n_;
 
  const TensorMap<Tensor<T, 2>, Eigen::Aligned> A(a_, sizeA);
  const TensorMap<Tensor<T, 2>, Eigen::Aligned> B(b_, sizeB);
  TensorMap<Tensor<T, 2>, Eigen::Aligned> C(c_, sizeC);
 
  typedef typename Tensor<T, 2>::DimensionPair DimPair;
  Eigen::array<DimPair, 1> dims;
  dims[0] = DimPair(1, 0);
#ifdef EIGEN_USE_SYCL  // warmup for sycl
  for (int iter = 0; iter < 10; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
#endif
  auto start = get_time();
  for (int iter = 0; iter < num_iters; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
  auto end = get_time();
  // Record the number of FLOPs executed per second (size_ multiplications and
  // additions for each value in the resulting tensor)
  finalizeBenchmark(start, end, m_, k_, n_, num_iters, "contraction");
  device_.deallocate(a_);
  device_.deallocate(b_);
  device_.deallocate(c_);
  device_.synchronize();
}

References Eigen::Aligned, Eigen::placeholders::end, finalizeBenchmark(), get_time(), and oomph::CumulativeTimings::start().

contractionABT()

template<typename T , typename Device , typename TensorIndex >

void contractionABT	(	const Device &	device_,
		TensorIndex	num_iters,
		TensorIndex	m_,
		TensorIndex	k_,
		TensorIndex	n_
	)

                                                                                                                  {
  T* a_;
  T* b_;
  T* c_;
  a_ = (T*)device_.allocate(m_ * k_ * sizeof(T));
  b_ = (T*)device_.allocate(k_ * n_ * sizeof(T));
  c_ = (T*)device_.allocate(m_ * n_ * sizeof(T));
 
  // Initialize the content of the memory pools to prevent asan from
  // complaining.
  device_.memset(a_, 12, m_ * k_ * sizeof(T));
  device_.memset(b_, 23, k_ * n_ * sizeof(T));
  device_.memset(c_, 31, m_ * n_ * sizeof(T));
 
  Eigen::array<TensorIndex, 2> sizeA;
  sizeA[0] = k_;
  sizeA[1] = m_;
  Eigen::array<TensorIndex, 2> sizeB;
  sizeB[0] = n_;
  sizeB[1] = k_;
  Eigen::array<TensorIndex, 2> sizeC;
  sizeC[0] = m_;
  sizeC[1] = n_;
 
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> A(a_, sizeA);
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> B(b_, sizeB);
  TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> C(c_, sizeC);
 
  typedef typename Tensor<T, 2>::DimensionPair DimPair;
  Eigen::array<DimPair, 1> dims;
  dims[0] = DimPair(0, 1);
#ifdef EIGEN_USE_SYCL  // warmup for sycl
  for (int iter = 0; iter < 10; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
#endif
  auto start = get_time();
  for (int iter = 0; iter < num_iters; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
  auto end = get_time();
  // Record the number of FLOPs executed per second (size_ multiplications and
  // additions for each value in the resulting tensor)
  finalizeBenchmark(start, end, m_, k_, n_, num_iters, "contractionABT");
  device_.deallocate(a_);
  device_.deallocate(b_);
  device_.deallocate(c_);
  device_.synchronize();
}

References Eigen::Aligned, Eigen::placeholders::end, finalizeBenchmark(), get_time(), and oomph::CumulativeTimings::start().

contractionAT()

template<typename T , typename Device , typename TensorIndex >

void contractionAT	(	const Device &	device_,
		TensorIndex	num_iters,
		TensorIndex	m_,
		TensorIndex	k_,
		TensorIndex	n_
	)

                                                                                                                 {
  T* a_;
  T* b_;
  T* c_;
  a_ = (T*)device_.allocate(m_ * k_ * sizeof(T));
  b_ = (T*)device_.allocate(k_ * n_ * sizeof(T));
  c_ = (T*)device_.allocate(m_ * n_ * sizeof(T));
 
  // Initialize the content of the memory pools to prevent asan from
  // complaining.
  device_.memset(a_, 12, m_ * k_ * sizeof(T));
  device_.memset(b_, 23, k_ * n_ * sizeof(T));
  device_.memset(c_, 31, m_ * n_ * sizeof(T));
  Eigen::array<TensorIndex, 2> sizeA;
  sizeA[0] = k_;
  sizeA[1] = m_;
  Eigen::array<TensorIndex, 2> sizeB;
  sizeB[0] = k_;
  sizeB[1] = n_;
  Eigen::array<TensorIndex, 2> sizeC;
  sizeC[0] = m_;
  sizeC[1] = n_;
 
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> A(a_, sizeA);
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> B(b_, sizeB);
  TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> C(c_, sizeC);
 
  typedef typename Tensor<T, 2>::DimensionPair DimPair;
  Eigen::array<DimPair, 1> dims;
  dims[0] = DimPair(0, 0);
#ifdef EIGEN_USE_SYCL  // warmup for sycl
  for (int iter = 0; iter < 10; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
#endif
  auto start = get_time();
  for (int iter = 0; iter < num_iters; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
  auto end = get_time();
  // Record the number of FLOPs executed per second (size_ multiplications and
  // additions for each value in the resulting tensor)
  finalizeBenchmark(start, end, m_, k_, n_, num_iters, "contractionAT");
  device_.deallocate(a_);
  device_.deallocate(b_);
  device_.deallocate(c_);
  device_.synchronize();
}

References Eigen::Aligned, Eigen::placeholders::end, finalizeBenchmark(), get_time(), and oomph::CumulativeTimings::start().

contractionBT()

template<typename T , typename Device , typename TensorIndex >

void contractionBT	(	const Device &	device_,
		TensorIndex	num_iters,
		TensorIndex	m_,
		TensorIndex	k_,
		TensorIndex	n_
	)

                                                                                                                 {
  T* a_;
  T* b_;
  T* c_;
  a_ = (T*)device_.allocate(m_ * k_ * sizeof(T));
  b_ = (T*)device_.allocate(k_ * n_ * sizeof(T));
  c_ = (T*)device_.allocate(m_ * n_ * sizeof(T));
 
  // Initialize the content of the memory pools to prevent asan from
  // complaining.
  device_.memset(a_, 12, m_ * k_ * sizeof(T));
  device_.memset(b_, 23, k_ * n_ * sizeof(T));
  device_.memset(c_, 31, m_ * n_ * sizeof(T));
 
  Eigen::array<TensorIndex, 2> sizeA;
  sizeA[0] = m_;
  sizeA[1] = k_;
  Eigen::array<TensorIndex, 2> sizeB;
  sizeB[0] = n_;
  sizeB[1] = k_;
  Eigen::array<TensorIndex, 2> sizeC;
  sizeC[0] = m_;
  sizeC[1] = n_;
 
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> A(a_, sizeA);
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> B(b_, sizeB);
  TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> C(c_, sizeC);
 
  typedef typename Tensor<T, 2>::DimensionPair DimPair;
  Eigen::array<DimPair, 1> dims;
  dims[0] = DimPair(1, 1);
#ifdef EIGEN_USE_SYCL  // warmup for sycl
  for (int iter = 0; iter < 10; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
#endif
  auto start = get_time();
  for (int iter = 0; iter < num_iters; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
  auto end = get_time();
  // Record the number of FLOPs executed per second (size_ multiplications and
  // additions for each value in the resulting tensor)
  finalizeBenchmark(start, end, m_, k_, n_, num_iters, "contractionBT");
  device_.deallocate(a_);
  device_.deallocate(b_);
  device_.deallocate(c_);
  device_.synchronize();
}

References Eigen::Aligned, Eigen::placeholders::end, finalizeBenchmark(), get_time(), and oomph::CumulativeTimings::start().

contractionRowMajor()

template<typename T , typename Device , typename TensorIndex >

void contractionRowMajor	(	const Device &	device_,
		TensorIndex	num_iters,
		TensorIndex	m_,
		TensorIndex	k_,
		TensorIndex	n_
	)

                                                                                                                       {
  T* a_;
  T* b_;
  T* c_;
  a_ = (T*)device_.allocate(m_ * k_ * sizeof(T));
  b_ = (T*)device_.allocate(k_ * n_ * sizeof(T));
  c_ = (T*)device_.allocate(m_ * n_ * sizeof(T));
 
  // Initialize the content of the memory pools to prevent asan from
  // complaining.
  device_.memset(a_, T(12), T(m_ * k_ * sizeof(T)));
  device_.memset(b_, T(23), T(k_ * n_ * sizeof(T)));
  device_.memset(c_, T(31), T(m_ * n_ * sizeof(T)));
 
  Eigen::array<TensorIndex, 2> sizeA;
  sizeA[0] = m_;
  sizeA[1] = k_;
  Eigen::array<TensorIndex, 2> sizeB;
  sizeB[0] = k_;
  sizeB[1] = n_;
  Eigen::array<TensorIndex, 2> sizeC;
  sizeC[0] = m_;
  sizeC[1] = n_;
 
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> A(a_, sizeA);
  const TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> B(b_, sizeB);
  TensorMap<Tensor<T, 2, Eigen::RowMajor>, Eigen::Aligned> C(c_, sizeC);
 
  typedef typename Tensor<T, 2>::DimensionPair DimPair;
  Eigen::array<DimPair, 1> dims;
  dims[0] = DimPair(1, 0);
#ifdef EIGEN_USE_SYCL  // warmup for sycl
  for (int iter = 0; iter < 10; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
#endif
  auto start = get_time();
  for (int iter = 0; iter < num_iters; ++iter) {
    C.device(device_) = A.contract(B, dims);
  }
  auto end = get_time();
  // Record the number of FLOPs executed per second (size_ multiplications and
  // additions for each value in the resulting tensor)
  finalizeBenchmark(start, end, m_, k_, n_, num_iters, "contractionRowMajor");
  device_.deallocate(a_);
  device_.deallocate(b_);
  device_.deallocate(c_);
  device_.synchronize();
}

References Eigen::Aligned, Eigen::placeholders::end, finalizeBenchmark(), get_time(), and oomph::CumulativeTimings::start().

finalizeBenchmark()

template<typename Start , typename End , typename TensorIndex >

void finalizeBenchmark	(	Start	start,
		End	end,
		TensorIndex	m_,
		TensorIndex	k_,
		TensorIndex	n_,
		TensorIndex	num_iters,
		std::string	name
	)

                                       {
  std::chrono::duration<double> elapsed_seconds = end - start;
  std::cout << "Kernel Name : " << name << ", M : " << m_ << ",  N : " << n_ << ", K : " << k_ << " GFLOP/s : "
            << static_cast<float>((static_cast<int64_t>(2) * m_ * n_ * k_ * num_iters) / elapsed_seconds.count()) * 1e-9
            << "\n";
  out << "Kernel Name : " << name << ", M : " << m_ << ",  N : " << n_ << ", K : " << k_ << " GFLOP/s : "
      << static_cast<float>((static_cast<int64_t>(2) * m_ * n_ * k_ * num_iters) / elapsed_seconds.count()) * 1e-9
      << "\n";
}

References e(), Eigen::placeholders::end, plotDoE::name, out(), and oomph::CumulativeTimings::start().

Referenced by contraction(), contractionABT(), contractionAT(), contractionBT(), and contractionRowMajor().

get_time()

std::chrono::time_point<std::chrono::system_clock> get_time

(

)

                                                      {
  std::chrono::time_point<std::chrono::system_clock> start, end;
  return std::chrono::system_clock::now();
}

References Eigen::placeholders::end, and oomph::CumulativeTimings::start().

Referenced by contraction(), contractionABT(), contractionAT(), contractionBT(), and contractionRowMajor().

main()

int main

(

)

           {
  cl::sycl::gpu_selector selector;
  Eigen::QueueInterface queue(selector);
  Eigen::SyclDevice device(&queue);
  int64_t num_iters = 20;
  for (int64_t m = 32; m <= 4096; m *= 2)
    for (int64_t k = 32; k <= 4096; k *= 2)
      for (int64_t n = 32; n <= 4096; n *= 2) {
        (contraction<float>(device, num_iters, m, k, n));
        (contractionRowMajor<float>(device, num_iters, m, k, n));
        (contractionAT<float>(device, num_iters, m, k, n));
        (contractionBT<float>(device, num_iters, m, k, n));
        (contractionABT<float>(device, num_iters, m, k, n));
      }
  return 0;
}

References k, m, and n.

out()

std::ofstream out

(

"Result.txt"

)

Referenced by ChangingTOIParticle::actionsAfterTimeStep(), statistics_while_running< T >::appendToFile(), Eigen::internal::OP< MatrixSolver, MatrixType, Scalar, true >::applyOP(), Eigen::TensorEvaluator< const TensorPaddingOp< PaddingDimensions, ArgType >, Device >::block(), CFile::CFile(), check_sparse_solving(), Eigen::ArpackGeneralizedSelfAdjointEigenSolver< MatrixType, MatrixSolver, BisSPD >::compute(), CFile::copy_data(), CFile::copy_fstat(), create_fluid_and_solid_surface_mesh_from_fluid_xda_mesh(), cwise_ref(), oomph::GeneralisedElement::describe_dofs(), oomph::Mesh::describe_dofs(), oomph::Data::describe_dofs(), oomph::SolidNode::describe_dofs(), oomph::Problem::describe_dofs(), oomph::ElementWithExternalElement::describe_local_dofs(), oomph::HeatedLinearSurfaceContactElement< ELEMENT >::describe_local_dofs(), oomph::ElementWithMovingNodes::describe_local_dofs(), oomph::ElementWithSpecificMovingNodes< ELEMENT, NODE_TYPE >::describe_local_dofs(), oomph::GeneralisedElement::describe_local_dofs(), oomph::FiniteElement::describe_local_dofs(), oomph::SolidFiniteElement::describe_local_dofs(), oomph::FaceElementAsGeomObject< ELEMENT >::describe_local_dofs(), oomph::FSIWallElement::describe_local_dofs(), oomph::Mesh::describe_local_dofs(), oomph::ProjectableElement< ELEMENT >::describe_local_dofs(), oomph::PseudoSolidNodeUpdateElement< BASIC, SOLID >::describe_local_dofs(), oomph::RefineablePseudoSolidNodeUpdateElement< BASIC, SOLID >::describe_local_dofs(), oomph::SpectralElement::describe_local_dofs(), oomph::RefineableQSpectralPoissonElement< DIM, NNODE_1D >::describe_local_dofs(), oomph::FSIImposeDisplacementByLagrangeMultiplierElement< ELEMENT >::describe_local_dofs(), oomph::FiniteElement::describe_nodal_local_dofs(), oomph::SolidFiniteElement::describe_solid_local_dofs(), oomph::SpineMesh::describe_spine_dofs(), oomph::MyProblem::doc_boundaries(), oomph::MyProblem::doc_solution(), EIGEN_DECLARE_TEST(), Detail::VTKPointDescriptorEntryImpl< T, V >::emit(), Detail::emitProxy(), File::File(), finalizeBenchmark(), format(), helpers::getLineFromStringStream(), Eigen::PastixLU< MatrixType_, IsStrSym >::grabMatrix(), Eigen::PastixLLT< MatrixType_, UpLo_ >::grabMatrix(), Eigen::PastixLDLT< MatrixType_, UpLo_ >::grabMatrix(), HstopCurve(), DPMBase::incrementRunNumberInFile(), oomph::MyProblem::initial_doc(), DPMBase::initialiseSolve(), mag2(), main(), BaseCluster::makeAmatFile(), BaseCluster::makeCdatFile(), BaseCluster::makeGnuplotFile(), BaseCluster::makeIntenalStructureFile(), BaseCluster::makeOverlFile(), BaseVTKWriter< H >::makeVTKFileWithHeader(), VTKData::makeVTKFileWithHeader(), oomph::StreamfunctionProblem::my_output(), File::openWrite(), File::openWriteNoAppend(), Eigen::internal::pfrexp(), Eigen::internal::pldexp< Packet4d >(), Eigen::internal::pldexp< Packet8d >(), Eigen::internal::pldexp_generic(), printBenchStyle(), printStatheader(), Eigen::internal::ptranspose(), Eigen::internal::putDenseElt(), Eigen::internal::PutMatrixElt(), helpers::readArrayFromCommandLine(), helpers::readVectorFromCommandLine(), REF_FREXP(), run_and_compare_to_gpu(), run_and_verify(), run_on_cpu(), run_on_gpu(), run_on_gpu_meta_kernel(), Membrane::saveAsOFF(), Membrane::saveAsSTL(), Eigen::saveMarket(), Eigen::saveMarketDense(), strcicmp(), test_async_multithread_chip(), test_async_multithread_elementwise(), test_async_multithread_volume_patch(), test_binary_builtins_fixed_arg2(), test_binary_builtins_func(), test_broadcast_sycl(), test_broadcast_sycl_fixed(), test_coeff_wise(), test_complex_operators(), test_complex_sqrt(), test_custom_binary_op_sycl(), test_custom_unary_op_sycl(), test_device_exceptions(), test_diagonal(), test_eigenvalues_direct(), test_forced_eval_sycl(), test_innermost_first_dims(), test_innermost_last_dims(), test_matrix_inverse(), test_multithread_chip(), test_multithread_compound_assignment(), test_multithread_elementwise(), test_multithread_volume_patch(), test_numeric_limits(), test_product(), test_reduce_middle_dims(), test_redux(), test_replicate(), test_sigmoid_sycl(), test_static_dims(), test_sycl_cast(), test_sycl_computations(), test_sycl_mem_sync(), test_sycl_mem_sync_offsets(), test_sycl_memset_offsets(), test_sycl_random_normal(), test_sycl_random_uniform(), test_tanh_sycl(), test_unary_builtins_for_scalar(), test_unary_builtins_return_bool(), To32BitDims(), to_string_padded(), Eigen::TensorEvaluator< const TensorPaddingOp< PaddingDimensions, ArgType >, Device >::updateCostPerDimension(), CFile::writeP4P(), Calibration::writePSDToFile(), CFile::writeRestart(), helpers::writeToFile(), and statistics_while_running< T >::writeToFile().