From b2dd8faa225a3fc56ef2aa3a130ea4d15cb5f49a Mon Sep 17 00:00:00 2001
From: yhmtsai <yhmtsai@gmail.com>
Date: Wed, 15 Apr 2020 14:25:59 +0200
Subject: [PATCH] assign a variable to avoid hip use too many VGPRs

---
 common/components/warp_blas.hpp.inc | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

diff --git a/common/components/warp_blas.hpp.inc b/common/components/warp_blas.hpp.inc
index 424200cc52e..5cf7f901b52 100644
--- a/common/components/warp_blas.hpp.inc
+++ b/common/components/warp_blas.hpp.inc
@@ -215,17 +215,16 @@ __device__ __forceinline__ void copy_matrix(
     size_type stride)
 {
     GKO_ASSERT(problem_size <= max_problem_size);
-#if !defined(__NVCC__)
-#pragma unroll 1
-#else
 #pragma unroll
-#endif  // !defined(__NVCC__)
     for (int32 i = 0; i < max_problem_size; ++i) {
         if (i < problem_size) {
             const auto idx = group.shfl(col_perm, i);
             if (group.thread_rank() < problem_size) {
+                // Need to assign a variable for the source_row, or hip
+                // will use a lot of VGPRs in unroll. This might lead a problem.
+                const auto val = source_row[i * increment];
                 destination[get_row_major_index<mod>(idx, row_perm, stride)] =
-                    static_cast<ResultValueType>(source_row[i * increment]);
+                    static_cast<ResultValueType>(val);
             }
         }
     }