Zero imag portion of diagonal in syrk/herk (#1648) (#1651)

daineAMD · web-flow · commit f08d23e349e1 · 2025-06-16T13:54:00.000-04:00
* Zero imag portion of diagonal in syrk/herk (#1648) * Bump version to 4.4.1 and update changelog
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -3,6 +3,12 @@
 rocBLAS documentation is available at
 [https://rocm.docs.amd.com/projects/rocBLAS/en/latest/index.html](https://rocm.docs.amd.com/projects/rocBLAS/en/latest/index.html).
 
+## rocBLAS 4.4.1 for ROCm 6.4.2
+
+### Resolved issues
+
+* Zero imaginary portion of diagonal of C matrix for cherk/zherk for gfx90a/gfx942 with problem sizes `k > 500`
+
 ## rocBLAS 4.4.0 for ROCm 6.4
 
 ### Added
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -1,5 +1,5 @@
 # ########################################################################
-# Copyright (C) 2016-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (C) 2016-2025 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -100,7 +100,7 @@ get_os_id(OS_ID)
 message (STATUS "OS detected is ${OS_ID}")
 
 # Versioning via rocm-cmake
-set ( VERSION_STRING "4.4.0" )
+set ( VERSION_STRING "4.4.1" )
 rocm_setup_version( VERSION ${VERSION_STRING} )
 
 # Users may override HIP path by specifying their own in CMAKE_MODULE_PATH
diff --git a/clients/include/blas3/testing_herk.hpp b/clients/include/blas3/testing_herk.hpp
@@ -212,7 +212,7 @@ void testing_herk(const Arguments& arg)
     rocblas_init_matrix(
         hA, arg, rocblas_client_alpha_sets_nan, rocblas_client_general_matrix, true, true);
     rocblas_init_matrix(
-        hC, arg, rocblas_client_beta_sets_nan, rocblas_client_hermitian_matrix, false, true);
+        hC, arg, rocblas_client_beta_sets_nan, rocblas_client_general_matrix, false, true);
 
     hC_gold = hC;
 
diff --git a/clients/include/blas3/testing_herk_batched.hpp b/clients/include/blas3/testing_herk_batched.hpp
@@ -259,7 +259,7 @@ void testing_herk_batched(const Arguments& arg)
     rocblas_init_matrix(
         hA, arg, rocblas_client_alpha_sets_nan, rocblas_client_general_matrix, true, true);
     rocblas_init_matrix(
-        hC, arg, rocblas_client_beta_sets_nan, rocblas_client_hermitian_matrix, false, true);
+        hC, arg, rocblas_client_beta_sets_nan, rocblas_client_general_matrix, false, true);
 
     hC_gold.copy_from(hC);
 
diff --git a/clients/include/blas3/testing_herk_strided_batched.hpp b/clients/include/blas3/testing_herk_strided_batched.hpp
@@ -412,7 +412,7 @@ void testing_herk_strided_batched(const Arguments& arg)
     rocblas_init_matrix(
         hA, arg, rocblas_client_alpha_sets_nan, rocblas_client_general_matrix, true, true);
     rocblas_init_matrix(
-        hC, arg, rocblas_client_beta_sets_nan, rocblas_client_hermitian_matrix, false, true);
+        hC, arg, rocblas_client_beta_sets_nan, rocblas_client_general_matrix, false, true);
 
     hC_gold.copy_from(hC);
 
diff --git a/library/src/blas3/herk_syrk_device.hpp b/library/src/blas3/herk_syrk_device.hpp
@@ -1,5 +1,5 @@
 /* ************************************************************************
- * Copyright (C) 2020-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2020-2025 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -1327,14 +1327,20 @@ rocblas_status rocblas_syr2k_her2k_dispatch(rocblas_handle    handle,
     return rocblas_status_success;
 }
 
-template <bool copy_from_C_to_W_C, bool is_upper, typename T, typename TPtr, int DIM_X, int DIM_Y>
+template <bool copy_from_C_to_W_C,
+          bool is_upper,
+          bool HERM,
+          typename T,
+          typename TPtr,
+          int DIM_X,
+          int DIM_Y>
 ROCBLAS_KERNEL(DIM_X* DIM_Y)
-rocblas_copy_triangular_excluding_diagonal_kernel(rocblas_int    n,
-                                                  TPtr           d_C,
-                                                  rocblas_int    ldc,
-                                                  rocblas_stride stride_C,
-                                                  T*             W_C,
-                                                  rocblas_int    batch_count)
+rocblas_copy_triangular_syrk_herk_kernel(rocblas_int    n,
+                                         TPtr           d_C,
+                                         rocblas_int    ldc,
+                                         rocblas_stride stride_C,
+                                         T*             W_C,
+                                         rocblas_int    batch_count)
 {
     uint32_t batch = blockIdx.z;
 
@@ -1379,19 +1385,24 @@ rocblas_copy_triangular_excluding_diagonal_kernel(rocblas_int    n,
             }
         }
 
+        // When copying back to C, we need to zero-out diagonal imaginary
+        if constexpr(HERM && !copy_from_C_to_W_C)
+            if(row == col && row < n)
+                C[row + row * int64_t(ldc)] = std::real(C[row + row * int64_t(ldc)]);
+
 #if DEVICE_GRID_YZ_16BIT
     }
 #endif
 }
 
-template <bool copy_from_C_to_W_C, bool is_upper, typename T, typename TPtr>
-rocblas_status rocblas_copy_triangular_excluding_diagonal(rocblas_handle handle,
-                                                          rocblas_int    n,
-                                                          TPtr           C,
-                                                          rocblas_int    ldc,
-                                                          rocblas_stride stride_C,
-                                                          T*             W_C,
-                                                          rocblas_int    batch_count)
+template <bool copy_from_C_to_W_C, bool is_upper, bool HERM, typename T, typename TPtr>
+rocblas_status rocblas_copy_triangular_syrk_herk(rocblas_handle handle,
+                                                 rocblas_int    n,
+                                                 TPtr           C,
+                                                 rocblas_int    ldc,
+                                                 rocblas_stride stride_C,
+                                                 T*             W_C,
+                                                 rocblas_int    batch_count)
 {
     hipStream_t rocblas_stream = handle->get_stream();
 
@@ -1405,12 +1416,13 @@ rocblas_status rocblas_copy_triangular_excluding_diagonal(rocblas_handle handle,
     dim3 gridDim((n - 1) / blockDim.x + 1, (n - 1) / blockDim.y + 1, batches);
 
     // Launch kernel
-    ROCBLAS_LAUNCH_KERNEL((rocblas_copy_triangular_excluding_diagonal_kernel<copy_from_C_to_W_C,
-                                                                             is_upper,
-                                                                             T,
-                                                                             TPtr,
-                                                                             DIM_X,
-                                                                             DIM_Y>),
+    ROCBLAS_LAUNCH_KERNEL((rocblas_copy_triangular_syrk_herk_kernel<copy_from_C_to_W_C,
+                                                                    is_upper,
+                                                                    HERM,
+                                                                    T,
+                                                                    TPtr,
+                                                                    DIM_X,
+                                                                    DIM_Y>),
                           gridDim,
                           blockDim,
                           0,
diff --git a/library/src/blas3/rocblas_syrk_herk_kernels.cpp b/library/src/blas3/rocblas_syrk_herk_kernels.cpp
@@ -1,5 +1,5 @@
 /* ************************************************************************
- * Copyright (C) 2020-2024 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (C) 2020-2025 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -104,10 +104,10 @@ rocblas_status rocblas_internal_syrk_herk_template(rocblas_handle    handle,
 
         // Launch kernel to copy the data from triangular matrix to the workspace memory
         if(rocblas_fill_upper == uplo)
-            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_excluding_diagonal<true, true>(
+            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_syrk_herk<true, true, HERM>(
                 handle, n, C, ldc, stride_C, (T*)w_mem, batch_count)));
         else
-            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_excluding_diagonal<true, false>(
+            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_syrk_herk<true, false, HERM>(
                 handle, n, C, ldc, stride_C, (T*)w_mem, batch_count)));
 
         RETURN_IF_ROCBLAS_ERROR((rocblas_internal_gemm_64<BATCHED>(handle,
@@ -134,10 +134,10 @@ rocblas_status rocblas_internal_syrk_herk_template(rocblas_handle    handle,
 
         // Launch kernel to copy the data from workspace memory back to triangular matrix
         if(rocblas_fill_upper == uplo)
-            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_excluding_diagonal<false, true>(
+            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_syrk_herk<false, true, HERM>(
                 handle, n, C, ldc, stride_C, (T*)w_mem, batch_count)));
         else
-            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_excluding_diagonal<false, false>(
+            RETURN_IF_ROCBLAS_ERROR((rocblas_copy_triangular_syrk_herk<false, false, HERM>(
                 handle, n, C, ldc, stride_C, (T*)w_mem, batch_count)));
 
         return rocblas_status_success;
