DenseVector/Kernels/Morpheus_WAXPBY_Impl.hpp

 
#ifndef MORPHEUS_DENSEVECTOR_KERNELS_WAXPBY_IMPL_HPP
#define MORPHEUS_DENSEVECTOR_KERNELS_WAXPBY_IMPL_HPP
 
#if defined(MORPHEUS_ENABLE_HIP)
#include <impl/Morpheus_HIPUtils.hpp>
#elif defined(MORPHEUS_ENABLE_CUDA)
#include <impl/Morpheus_CudaUtils.hpp>
#endif
 
namespace Morpheus {
namespace Impl {
namespace Kernels {
 
template <typename SizeType, typename ValueType1, typename ValueType2,
          typename ValueType3>
__global__ void waxpby_kernel(SizeType n, ValueType1 alpha, const ValueType1* x,
                              ValueType2 beta, const ValueType2* y,
                              ValueType3* w) {
  const SizeType tid = blockDim.x * blockIdx.x + threadIdx.x;
  if (tid > n) return;
 
  if (alpha == 1.0) {
    w[tid] = x[tid] + beta * y[tid];
  } else if (beta == 1.0) {
    w[tid] = alpha * x[tid] + y[tid];
  } else {
    w[tid] = alpha * x[tid] + beta * y[tid];
  }
}
}  // namespace Kernels
}  // namespace Impl
}  // namespace Morpheus
 
#endif  // MORPHEUS_DENSEVECTOR_KERNELS_WAXPBY_IMPL_HPP