DenseVector/HIP/Morpheus_WAXPBY_Impl.hpp

 
#ifndef MORPHEUS_DENSEVECTOR_HIP_WAXPBY_IMPL_HPP
#define MORPHEUS_DENSEVECTOR_HIP_WAXPBY_IMPL_HPP
 
#include <Morpheus_Macros.hpp>
#if defined(MORPHEUS_ENABLE_HIP)
 
#include <Morpheus_SpaceTraits.hpp>
#include <Morpheus_FormatTraits.hpp>
#include <Morpheus_FormatTags.hpp>
#include <Morpheus_Spaces.hpp>
 
#include <impl/Morpheus_HIPUtils.hpp>
#include <impl/DenseVector/Kernels/Morpheus_WAXPBY_Impl.hpp>
 
#include <cassert>
 
namespace Morpheus {
namespace Impl {
 
template <typename ExecSpace, typename SizeType, typename Vector1,
          typename Vector2, typename Vector3>
inline void waxpby(
    const SizeType n, const typename Vector1::value_type alpha,
    const Vector1& x, const typename Vector2::value_type beta, const Vector2& y,
    Vector3& w,
    typename std::enable_if_t<
        Morpheus::is_dense_vector_format_container_v<Vector1> &&
        Morpheus::is_dense_vector_format_container_v<Vector2> &&
        Morpheus::is_dense_vector_format_container_v<Vector3> &&
        Morpheus::has_custom_backend_v<ExecSpace> &&
        Morpheus::has_hip_execution_space_v<ExecSpace> &&
        Morpheus::has_access_v<ExecSpace, Vector1, Vector2, Vector3>>* =
        nullptr) {
  assert(x.size() >= n);
  assert(y.size() >= n);
  assert(w.size() >= n);
 
  const SizeType BLOCK_SIZE = 256;
  const SizeType NUM_BLOCKS = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;
 
  Kernels::waxpby_kernel<SizeType, typename Vector1::value_type,
                         typename Vector2::value_type,
                         typename Vector3::value_type>
      <<<NUM_BLOCKS, BLOCK_SIZE, 0>>>(n, alpha, x.data(), beta, y.data(),
                                      w.data());
#if defined(DEBUG) || defined(MORPHEUS_DEBUG)
  getLastHIPError("spmv_waxpby_kernel: Kernel execution failed");
#endif
}
 
}  // namespace Impl
}  // namespace Morpheus
 
#endif  // MORPHEUS_ENABLE_HIP
#endif  // MORPHEUS_DENSEVECTOR_HIP_WAXPBY_IMPL_HPP