group_blas_fill.hh

Go to the documentation of this file.

#pragma once
 
#include <batchlas/blas/device/detail/group_blas_common.hh>
 
namespace batchlas::device {
 
/// @addtogroup api_device_blas
/// @{
 
/// @brief Sets every element of `x` to `value`, cooperatively across `group`.
/// @tparam Group  `sycl::group`, `sycl::sub_group` or an `nd_item`; every work-item must call
/// @param group  executor whose work-items share the elements
/// @param x      single vector (`batch_size() == 1`), written in place
/// @param value  fill value
template <typename Group, typename T>
inline constexpr void fill(const Group& group,
                           const VectorView<T>& x,
                           const T& value) {
    detail::validate_vector_operand(x, "fill");
    const int local_id = detail::group_local_linear_id(group);
    const int local_size = detail::group_local_linear_range(group);
 
    for (int index = local_id; index < x.size(); index += local_size) {
        x(index) = value;
    }
}
 
/// @brief Sets every element of the dense matrix `a` to `value`; the padding rows between `a.rows()` and `a.ld()` are not touched.
/// @param group  executor whose work-items share the elements
/// @param a      single matrix, written in place
/// @param value  fill value
template <typename Group, typename T>
inline constexpr void fill(const Group& group,
                           const KernelMatrixView<T, MatrixFormat::Dense>& a,
                           const T& value) {
    detail::validate_single_problem(a, Transpose::NoTrans);
    const int local_id = detail::group_local_linear_id(group);
    const int local_size = detail::group_local_linear_range(group);
    const int extent = a.rows() * a.cols();
 
    for (int linear_index = local_id; linear_index < extent; linear_index += local_size) {
        const int row = linear_index % a.rows();
        const int col = linear_index / a.rows();
        a(row, col) = value;
    }
}
 
/// @}
 
} // namespace batchlas::device