message: Test detection with verbose output
diff: @@ -46,8 +46,18 @@ SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  #include "cuda/components/uninitialized_array.hpp"
  #include "cuda/components/warp_blas.cuh"
  +#include <cstdio>
  namespace gko {
  +
  +
  +template <typename T>
  +__device__ __forceinline__ constexpr real(const thrust::complex<T> &x)
  +{
  + return x.real();
  +}
  +
  +
  namespace kernels {
  namespace cuda {
  @@ -70,13 +80,18 @@ __device__ __forceinline__ bool validate_precision_reduction_feasibility(
  {
  using gko::detail::float_traits;
  // save original data and reduce precision
  + if (group.thread_rank() < block_size) {
  #pragma unroll
  for (auto i = 0u; i < max_block_size; ++i) {
  if (i >= block_size) {
  break;
  }
  - work[i * stride + group.thread_rank()] = row[i];
  + auto tmp = work[i * stride + group.thread_rank()] = row[i];
  row[i] = static_cast<ValueType>(static_cast<ReducedType>(row[i]));
  + printf("%d %d %d (%d): %d >> %.16le -> %.16le\n", threadIdx.x,
  + threadIdx.y, blockIdx.x, int(sizeof(ReducedType)), int(i),
  + double(real(tmp)), double(real(row[i])));
  + }
  }
  // compute the condition number
  @@ -84,17 +99,29 @@ __device__ __forceinline__ bool validate_precision_reduction_feasibility(
  auto trans_perm = perm;
  auto block_cond = compute_infinity_norm<max_block_size>(group, block_size,
  block_size, row);
  + printf("%d %d %d (%d): ||A|| = %.16le\n", threadIdx.x, threadIdx.y,
  + blockIdx.x, int(sizeof(ReducedType)), double(block_cond));
  invert_block<max_block_size>(group, block_size, row, perm, trans_perm);
  - block_cond *= compute_infinity_norm<max_block_size>(group, block_size,
  + auto tmp = compute_infinity_norm<max_block_size>(group, block_size,
  block_size, row);
  -
  + printf("%d %d %d (%d): ||inv(A)|| = %.16le\n", threadIdx.x, threadIdx.y,
  + blockIdx.x, int(sizeof(ReducedType)), double(tmp));
  + block_cond *= tmp;
  + printf("%d %d %d (%d): k(A) = %.16le\n", threadIdx.x, threadIdx.y,
  + blockIdx.x, int(sizeof(ReducedType)), double(block_cond));
  // restore original data
  + if (group.thread_rank() < block_size) {
  #pragma unroll
  for (auto i = 0u; i < max_block_size; ++i) {
  if (i >= block_size) {
  break;
  }
  + auto tmp = row[i];
  row[i] = work[i * stride + group.thread_rank()];
  + printf("%d %d %d (%d): %d >> %.16le -> %.16le\n", threadIdx.x,
  + threadIdx.y, blockIdx.x, int(sizeof(ReducedType)), int(i),
  + double(real(tmp)), double(real(row[i])));
  + }
  }
  return block_cond > 0.0 &&
  