text: /* ************************************************************************
   * Copyright 2016-2019 Advanced Micro Devices, Inc.
   * ************************************************************************ */
  #include "fetch_template.h"
  #include "handle.h"
  #include "logging.h"
  #include "reduction.h"
  #include "rocblas.h"
  #include "utility.h"
  
  namespace
  {
      template <class To>
      struct rocblas_fetch_asum
      {
          template <typename Ti>
          __forceinline__ __device__ To operator()(Ti x, ptrdiff_t)
          {
              return {fetch_asum(x)};
          }
      };
  
      template <typename>
      constexpr char rocblas_asum_name[] = "unknown";
      template <>
      constexpr char rocblas_asum_name<float>[] = "rocblas_sasum";
      template <>
      constexpr char rocblas_asum_name<double>[] = "rocblas_dasum";
      template <>
      constexpr char rocblas_asum_name<rocblas_float_complex>[] = "rocblas_scasum";
      template <>
      constexpr char rocblas_asum_name<rocblas_double_complex>[] = "rocblas_dzasum";
  
      // HIP support up to 1024 threads/work itmes per thread block/work group
      constexpr int NB = 512;
  
      // allocate workspace inside this API
      template <typename Ti, typename To>
      rocblas_status rocblas_asum(
          rocblas_handle handle, rocblas_int n, const Ti* x, rocblas_int incx, To* result)
      {
          if(!handle)
              return rocblas_status_invalid_handle;
  
          auto layer_mode = handle->layer_mode;
          if(layer_mode & rocblas_layer_mode_log_trace)
              log_trace(handle, rocblas_asum_name<Ti>, n, x, incx);
  
          if(layer_mode & rocblas_layer_mode_log_bench)
              log_bench(handle,
                        "./rocblas-bench -f asum -r",
                        rocblas_precision_string<Ti>,
                        "-n",
                        n,
                        "--incx",
                        incx);
  
          if(layer_mode & rocblas_layer_mode_log_profile)
              log_profile(handle, rocblas_asum_name<Ti>, "N", n, "incx", incx);
  
          if(!x || !result)
              return rocblas_status_invalid_pointer;
  
          // Quick return if possible.
          if(n <= 0 || incx <= 0)
          {
              if(handle->is_device_memory_size_query())
                  return rocblas_status_size_unchanged;
              else if(rocblas_pointer_mode_device == handle->pointer_mode)
                  RETURN_IF_HIP_ERROR(hipMemset(result, 0, sizeof(*result)));
              else
                  *result = 0;
              return rocblas_status_success;
          }
  
          auto blocks = (n - 1) / NB + 1;
          if(handle->is_device_memory_size_query())
              return handle->set_optimal_device_memory_size(sizeof(To) * blocks);
  
          auto mem = handle->device_malloc(sizeof(To) * blocks);
          if(!mem)
              return rocblas_status_memory_error;
  
          return rocblas_reduction_kernel<NB, rocblas_fetch_asum<To>>(
              handle, n, x, incx, result, (To*)mem, blocks);
      }
  
  } // namespace
  
  /*
   * ===========================================================================
   *    C wrapper
   * ===========================================================================
   */
  
  extern "C" {
  
  rocblas_status rocblas_sasum(
      rocblas_handle handle, rocblas_int n, const float* x, rocblas_int incx, float* result)
  {
      return rocblas_asum(handle, n, x, incx, result);
  }
  
  rocblas_status rocblas_dasum(
      rocblas_handle handle, rocblas_int n, const double* x, rocblas_int incx, double* result)
  {
      return rocblas_asum(handle, n, x, incx, result);
  }
  
  rocblas_status rocblas_scasum(rocblas_handle               handle,
                                rocblas_int                  n,
                                const rocblas_float_complex* x,
                                rocblas_int                  incx,
                                float*                       result)
  {
      return rocblas_asum(handle, n, x, incx, result);
  }
  
  rocblas_status rocblas_dzasum(rocblas_handle                handle,
                                rocblas_int                   n,
                                const rocblas_double_complex* x,
                                rocblas_int                   incx,
                                double*                       result)
  {
      return rocblas_asum(handle, n, x, incx, result);
  }
  
  } // extern "C"
  