update documentation, remove half.hpp, and put half func in gko

Co-authored-by: Marcel Koch <[email protected]>
Co-authored-by: Thomas Grützmacher <[email protected]>

CMakeLists.txt

@@ -33,7 +33,7 @@ option(GINKGO_BUILD_DOC "Generate documentation" OFF)
 option(GINKGO_FAST_TESTS "Reduces the input size for a few tests known to be time-intensive" OFF)
 option(GINKGO_TEST_NONDEFAULT_STREAM "Uses non-default streams in CUDA and HIP tests" OFF)
 option(GINKGO_MIXED_PRECISION "Instantiate true mixed-precision kernels (otherwise they will be conversion-based using implicit temporary storage)" OFF)
-option(GINKGO_ENABLE_HALF "Enable the half operation" ON)
+option(GINKGO_ENABLE_HALF "Enable the use of half precision" ON)
 option(GINKGO_SKIP_DEPENDENCY_UPDATE
     "Do not update dependencies each time the project is rebuilt" ON)
 option(GINKGO_WITH_CLANG_TIDY "Make Ginkgo call `clang-tidy` to find programming issues." OFF)

accessor/cuda_helper.hpp

@@ -21,6 +21,11 @@ struct __half;
 
 
 namespace gko {
+
+
+class half;
+
+
 namespace acc {
 namespace detail {
 

accessor/hip_helper.hpp

@@ -21,6 +21,11 @@ struct __half;
 
 
 namespace gko {
+
+
+class half;
+
+
 namespace acc {
 namespace detail {
 

benchmark/run_all_benchmarks.sh

@@ -114,7 +114,7 @@ elif [ "${BENCHMARK_PRECISION}" == "half" ]; then
     BENCH_SUFFIX="_half"
 else
     echo "BENCHMARK_PRECISION is set to the not supported \"${BENCHMARK_PRECISION}\"." 1>&2
-    echo "Currently supported values: \"double\", \"single\", \"dcomplex\" and \"scomplex\"" 1>&2
+    echo "Currently supported values: \"double\", \"single\", \"half\", \"dcomplex\" and \"scomplex\"" 1>&2
     exit 1
 fi
 

common/cuda_hip/components/reduction.hpp.inc

@@ -47,7 +47,7 @@ __device__ __forceinline__ int choose_pivot(const Group& group,
                                             bool is_pivoted)
 {
     using real = remove_complex<ValueType>;
-    real lmag = real(is_pivoted ? -one<real>() : abs(local_data));
+    auto lmag = static_cast<real>(is_pivoted ? -one<real>() : abs(local_data));
     const auto pivot =
         reduce(group, group.thread_rank(), [&](int lidx, int ridx) {
             const auto rmag = group.shfl(lmag, ridx);

common/cuda_hip/matrix/csr_kernels.hpp.inc

@@ -194,7 +194,7 @@ __global__ __launch_bounds__(spmv_block_size) void abstract_spmv(
 {
     using arithmetic_type = typename output_accessor::arithmetic_type;
     using output_type = typename output_accessor::storage_type;
-    const arithmetic_type scale_factor = static_cast<arithmetic_type>(alpha[0]);
+    const auto scale_factor = static_cast<arithmetic_type>(alpha[0]);
     spmv_kernel(nwarps, num_rows, val, col_idxs, row_ptrs, srow, b, c,
                 [&scale_factor](const arithmetic_type& x) {
                     return static_cast<output_type>(scale_factor * x);
@@ -402,7 +402,7 @@ __global__ __launch_bounds__(spmv_block_size) void abstract_reduce(
     const IndexType* __restrict__ last_row,
     const MatrixValueType* __restrict__ alpha, acc::range<output_accessor> c)
 {
-    const arithmetic_type alpha_val = static_cast<arithmetic_type>(alpha[0]);
+    const auto alpha_val = static_cast<arithmetic_type>(alpha[0]);
     merge_path_reduce(
         nwarps, last_val, last_row, c,
         [&alpha_val](const arithmetic_type& x) { return alpha_val * x; });

core/base/device_matrix_data_kernels.hpp

@@ -10,7 +10,6 @@
 
 #include <ginkgo/core/base/device_matrix_data.hpp>
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 
 #include "core/base/kernel_declaration.hpp"

core/base/utils.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 #include <type_traits>
 
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/polymorphic_object.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/base/utils.hpp>

core/components/absolute_array_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 

core/components/fill_array_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 
 #include "core/base/kernel_declaration.hpp"

core/components/format_conversion_kernels.hpp

@@ -10,7 +10,6 @@
 
 #include <ginkgo/core/base/device_matrix_data.hpp>
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 
 #include "core/base/kernel_declaration.hpp"

core/components/precision_conversion_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 

core/components/prefix_sum_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 
 #include "core/base/kernel_declaration.hpp"

core/components/reduce_array_kernels.hpp

@@ -10,7 +10,6 @@
 
 #include <ginkgo/core/base/array.hpp>
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 
 #include "core/base/kernel_declaration.hpp"

core/distributed/matrix_kernels.hpp

@@ -8,7 +8,6 @@
 
 #include <ginkgo/core/base/array.hpp>
 #include <ginkgo/core/base/device_matrix_data.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/matrix_data.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/distributed/partition.hpp>

core/factorization/cholesky_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>
 

core/factorization/factorization_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>
 

core/factorization/ilu_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/factorization/ilu.hpp>
 #include <ginkgo/core/matrix/csr.hpp>

core/factorization/lu_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>
 

core/factorization/par_ic_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/factorization/par_ic.hpp>
 #include <ginkgo/core/matrix/coo.hpp>

core/factorization/par_ict_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/factorization/par_ict.hpp>
 #include <ginkgo/core/matrix/coo.hpp>

core/factorization/par_ilu_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/factorization/par_ilu.hpp>
 #include <ginkgo/core/matrix/csr.hpp>

core/factorization/par_ilut_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/factorization/par_ilut.hpp>
 #include <ginkgo/core/matrix/coo.hpp>

core/matrix/coo_kernels.hpp

@@ -6,7 +6,6 @@
 #define GKO_CORE_MATRIX_COO_KERNELS_HPP_
 
 
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/coo.hpp>
 #include <ginkgo/core/matrix/csr.hpp>

core/matrix/csr_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/index_set.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/coo.hpp>

core/matrix/csr_lookup.hpp

@@ -8,7 +8,6 @@
 
 #include <type_traits>
 
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/intrinsics.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>

core/matrix/dense_kernels.hpp

@@ -8,7 +8,6 @@
 
 #include <memory>
 
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/matrix/diagonal_kernels.hpp

@@ -6,7 +6,6 @@
 #define GKO_CORE_MATRIX_DIAGONAL_KERNELS_HPP_
 
 
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/matrix/fbcsr_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/matrix/fft_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>
 

core/matrix/sparsity_csr_kernels.hpp

@@ -6,7 +6,6 @@
 #define GKO_CORE_MATRIX_SPARSITY_CSR_KERNELS_HPP_
 
 
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>
 #include <ginkgo/core/matrix/sparsity_csr.hpp>

core/preconditioner/jacobi_utils.hpp

@@ -6,7 +6,6 @@
 #define GKO_CORE_PRECONDITIONER_JACOBI_UTILS_HPP_
 
 
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 

core/reorder/rcm_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/lin_op.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/csr.hpp>

core/solver/bicg_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/bicgstab_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/cb_gmres_accessor.hpp

@@ -15,7 +15,6 @@
 #include <ginkgo/core/base/array.hpp>
 #include <ginkgo/core/base/dim.hpp>
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/std_extensions.hpp>
 #include <ginkgo/core/base/types.hpp>

core/solver/cb_gmres_kernels.hpp

@@ -8,7 +8,6 @@
 
 #include <ginkgo/core/base/array.hpp>
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/range.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/cg_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/cgs_kernels.hpp

@@ -9,7 +9,6 @@
 #include <memory>
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/common_gmres_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/gmres_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/idr_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/solver/ir_kernels.hpp

@@ -10,7 +10,6 @@
 
 #include <ginkgo/core/base/array.hpp>
 #include <ginkgo/core/base/executor.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/stop/stopping_status.hpp>
 

core/solver/multigrid_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/stop/criterion_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/stop/stopping_status.hpp>
 

core/stop/residual_norm_kernels.hpp

@@ -7,7 +7,6 @@
 
 
 #include <ginkgo/core/base/array.hpp>
-#include <ginkgo/core/base/half.hpp>
 #include <ginkgo/core/base/math.hpp>
 #include <ginkgo/core/base/types.hpp>
 #include <ginkgo/core/matrix/dense.hpp>

core/test/utils/assertions.hpp

@@ -259,7 +259,7 @@ void save_matrices_to_disk(Ostream& os, const MatrixData1& first,
 template <typename MatrixData1, typename MatrixData2>
 double get_relative_error(const MatrixData1& first, const MatrixData2& second)
 {
-    using std::abs;
+    using gko::abs;
     using vt = typename detail::biggest_valuetype<
         typename MatrixData1::value_type,
         typename MatrixData2::value_type>::type;