[HIPIFY][6.3.0][BLAS] Sync with hipBLAS - Step 13

+ `hipblas(S|D|C|Z)trsm(_v2)?_64` support + Updated synthetic tests, the regenerated `hipify-perl`, and `BLAS` `CUDA2HIP` documentation
emankov · Nov 7, 2024 · 23e196b · 23e196b
1 parent a01cf56
commit 23e196b
Show file tree

Hide file tree

Showing 5 changed files with 64 additions and 32 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -4543,7 +4543,9 @@ sub simpleSubstitutions {
     subst("cublasCtrmv_v2_64", "hipblasCtrmv_v2_64", "library");
     subst("cublasCtrsm", "hipblasCtrsm_v2", "library");
     subst("cublasCtrsmBatched", "hipblasCtrsmBatched_v2", "library");
+    subst("cublasCtrsm_64", "hipblasCtrsm_v2_64", "library");
     subst("cublasCtrsm_v2", "hipblasCtrsm_v2", "library");
+    subst("cublasCtrsm_v2_64", "hipblasCtrsm_v2_64", "library");
     subst("cublasCtrsv", "hipblasCtrsv_v2", "library");
     subst("cublasCtrsv_64", "hipblasCtrsv_v2_64", "library");
     subst("cublasCtrsv_v2", "hipblasCtrsv_v2", "library");
@@ -4693,7 +4695,9 @@ sub simpleSubstitutions {
     subst("cublasDtrmv_v2_64", "hipblasDtrmv_64", "library");
     subst("cublasDtrsm", "hipblasDtrsm", "library");
     subst("cublasDtrsmBatched", "hipblasDtrsmBatched", "library");
+    subst("cublasDtrsm_64", "hipblasDtrsm_64", "library");
     subst("cublasDtrsm_v2", "hipblasDtrsm", "library");
+    subst("cublasDtrsm_v2_64", "hipblasDtrsm_64", "library");
     subst("cublasDtrsv", "hipblasDtrsv", "library");
     subst("cublasDtrsv_64", "hipblasDtrsv_64", "library");
     subst("cublasDtrsv_v2", "hipblasDtrsv", "library");
@@ -4940,7 +4944,9 @@ sub simpleSubstitutions {
     subst("cublasStrmv_v2_64", "hipblasStrmv_64", "library");
     subst("cublasStrsm", "hipblasStrsm", "library");
     subst("cublasStrsmBatched", "hipblasStrsmBatched", "library");
+    subst("cublasStrsm_64", "hipblasStrsm_64", "library");
     subst("cublasStrsm_v2", "hipblasStrsm", "library");
+    subst("cublasStrsm_v2_64", "hipblasStrsm_64", "library");
     subst("cublasStrsv", "hipblasStrsv", "library");
     subst("cublasStrsv_64", "hipblasStrsv_64", "library");
     subst("cublasStrsv_v2", "hipblasStrsv", "library");
@@ -5112,7 +5118,9 @@ sub simpleSubstitutions {
     subst("cublasZtrmv_v2_64", "hipblasZtrmv_v2_64", "library");
     subst("cublasZtrsm", "hipblasZtrsm_v2", "library");
     subst("cublasZtrsmBatched", "hipblasZtrsmBatched_v2", "library");
+    subst("cublasZtrsm_64", "hipblasZtrsm_v2_64", "library");
     subst("cublasZtrsm_v2", "hipblasZtrsm_v2", "library");
+    subst("cublasZtrsm_v2_64", "hipblasZtrsm_v2_64", "library");
     subst("cublasZtrsv", "hipblasZtrsv_v2", "library");
     subst("cublasZtrsv_64", "hipblasZtrsv_v2_64", "library");
     subst("cublasZtrsv_v2", "hipblasZtrsv_v2", "library");
@@ -11645,8 +11653,6 @@ sub warnHipOnlyUnsupportedFunctions {
         "cudnnAdvInferVersionCheck",
         "cudnnActivationStruct",
         "cublasZtrttp",
-        "cublasZtrsm_v2_64",
-        "cublasZtrsm_64",
         "cublasZtrsmBatched_64",
         "cublasZtpttr",
         "cublasZmatinvBatched",
@@ -11666,8 +11672,6 @@ sub warnHipOnlyUnsupportedFunctions {
         "cublasSwapEx_64",
         "cublasSwapEx",
         "cublasStrttp",
-        "cublasStrsm_v2_64",
-        "cublasStrsm_64",
         "cublasStrsmBatched_64",
         "cublasStpttr",
         "cublasSmatinvBatched",
@@ -11760,17 +11764,13 @@ sub warnHipOnlyUnsupportedFunctions {
         "cublasGemmBatchedEx_64",
         "cublasFree",
         "cublasDtrttp",
-        "cublasDtrsm_v2_64",
-        "cublasDtrsm_64",
         "cublasDtrsmBatched_64",
         "cublasDtpttr",
         "cublasDmatinvBatched",
         "cublasDgemmGroupedBatched_64",
         "cublasDgemmGroupedBatched",
         "cublasDdgmm_64",
         "cublasCtrttp",
-        "cublasCtrsm_v2_64",
-        "cublasCtrsm_64",
         "cublasCtrsmBatched_64",
         "cublasCtpttr",
         "cublasCsyrkEx_64",

diff --git a/docs/tables/CUBLAS_API_supported_by_HIP.md b/docs/tables/CUBLAS_API_supported_by_HIP.md
@@ -1569,9 +1569,9 @@
 |`cublasCtrmm_v2`| | | | |`hipblasCtrmm_v2`|6.0.0| | | | |
 |`cublasCtrmm_v2_64`|12.0| | | |`hipblasCtrmm_v2_64`|6.3.0| | | |6.3.0|
 |`cublasCtrsm`| | | | |`hipblasCtrsm_v2`|6.0.0| | | | |
-|`cublasCtrsm_64`|12.0| | | | | | | | | |
+|`cublasCtrsm_64`|12.0| | | |`hipblasCtrsm_v2_64`|6.3.0| | | |6.3.0|
 |`cublasCtrsm_v2`| | | | |`hipblasCtrsm_v2`|6.0.0| | | | |
-|`cublasCtrsm_v2_64`|12.0| | | | | | | | | |
+|`cublasCtrsm_v2_64`|12.0| | | |`hipblasCtrsm_v2_64`|6.3.0| | | |6.3.0|
 |`cublasDgemm`| | | | |`hipblasDgemm`|1.8.2| | | | |
 |`cublasDgemmBatched`| | | | |`hipblasDgemmBatched`|1.8.2| | | | |
 |`cublasDgemmBatched_64`|12.0| | | |`hipblasDgemmBatched_64`|6.3.0| | | |6.3.0|
@@ -1605,9 +1605,9 @@
 |`cublasDtrmm_v2`| | | | |`hipblasDtrmm`|3.2.0| |6.0.0| | |
 |`cublasDtrmm_v2_64`|12.0| | | |`hipblasDtrmm_64`|6.3.0| | | |6.3.0|
 |`cublasDtrsm`| | | | |`hipblasDtrsm`|1.8.2| | | | |
-|`cublasDtrsm_64`|12.0| | | | | | | | | |
+|`cublasDtrsm_64`|12.0| | | |`hipblasDtrsm_64`|6.3.0| | | |6.3.0|
 |`cublasDtrsm_v2`| | | | |`hipblasDtrsm`|1.8.2| | | | |
-|`cublasDtrsm_v2_64`|12.0| | | | | | | | | |
+|`cublasDtrsm_v2_64`|12.0| | | |`hipblasDtrsm_64`|6.3.0| | | |6.3.0|
 |`cublasGemmGroupedBatchedEx`|12.5| | | | | | | | | |
 |`cublasGemmGroupedBatchedEx_64`|12.5| | | | | | | | | |
 |`cublasHSHgemvBatched`|11.6| | | | | | | | | |
@@ -1657,9 +1657,9 @@
 |`cublasStrmm_v2`| | | | |`hipblasStrmm`|3.2.0| |6.0.0| | |
 |`cublasStrmm_v2_64`|12.0| | | |`hipblasStrmm_64`|6.3.0| | | |6.3.0|
 |`cublasStrsm`| | | | |`hipblasStrsm`|1.8.2| | | | |
-|`cublasStrsm_64`|12.0| | | | | | | | | |
+|`cublasStrsm_64`|12.0| | | |`hipblasStrsm_64`|6.3.0| | | |6.3.0|
 |`cublasStrsm_v2`| | | | |`hipblasStrsm`|1.8.2| | | | |
-|`cublasStrsm_v2_64`|12.0| | | | | | | | | |
+|`cublasStrsm_v2_64`|12.0| | | |`hipblasStrsm_64`|6.3.0| | | |6.3.0|
 |`cublasTSSgemvBatched`|11.6| | | | | | | | | |
 |`cublasTSSgemvBatched_64`|12.0| | | | | | | | | |
 |`cublasTSSgemvStridedBatched`|11.6| | | | | | | | | |
@@ -1715,9 +1715,9 @@
 |`cublasZtrmm_v2`| | | | |`hipblasZtrmm_v2`|6.0.0| | | | |
 |`cublasZtrmm_v2_64`|12.0| | | |`hipblasZtrmm_v2_64`|6.3.0| | | |6.3.0|
 |`cublasZtrsm`| | | | |`hipblasZtrsm_v2`|6.0.0| | | | |
-|`cublasZtrsm_64`|12.0| | | | | | | | | |
+|`cublasZtrsm_64`|12.0| | | |`hipblasZtrsm_v2_64`|6.3.0| | | |6.3.0|
 |`cublasZtrsm_v2`| | | | |`hipblasZtrsm_v2`|6.0.0| | | | |
-|`cublasZtrsm_v2_64`|12.0| | | | | | | | | |
+|`cublasZtrsm_v2_64`|12.0| | | |`hipblasZtrsm_v2_64`|6.3.0| | | |6.3.0|
 
 ## **8. BLAS-like Extension**
 

diff --git a/docs/tables/CUBLAS_API_supported_by_HIP_and_ROC.md b/docs/tables/CUBLAS_API_supported_by_HIP_and_ROC.md
@@ -1569,9 +1569,9 @@
 |`cublasCtrmm_v2`| | | | |`hipblasCtrmm_v2`|6.0.0| | | | |`rocblas_ctrmm`|3.5.0| |6.0.0| | |
 |`cublasCtrmm_v2_64`|12.0| | | |`hipblasCtrmm_v2_64`|6.3.0| | | |6.3.0|`rocblas_ctrmm_64`|6.3.0| | | |6.3.0|
 |`cublasCtrsm`| | | | |`hipblasCtrsm_v2`|6.0.0| | | | |`rocblas_ctrsm`|3.5.0| | | | |
-|`cublasCtrsm_64`|12.0| | | | | | | | | |`rocblas_ctrsm_64`|6.2.0| | | | |
+|`cublasCtrsm_64`|12.0| | | |`hipblasCtrsm_v2_64`|6.3.0| | | |6.3.0|`rocblas_ctrsm_64`|6.2.0| | | | |
 |`cublasCtrsm_v2`| | | | |`hipblasCtrsm_v2`|6.0.0| | | | |`rocblas_ctrsm`|3.5.0| | | | |
-|`cublasCtrsm_v2_64`|12.0| | | | | | | | | |`rocblas_ctrsm_64`|6.2.0| | | | |
+|`cublasCtrsm_v2_64`|12.0| | | |`hipblasCtrsm_v2_64`|6.3.0| | | |6.3.0|`rocblas_ctrsm_64`|6.2.0| | | | |
 |`cublasDgemm`| | | | |`hipblasDgemm`|1.8.2| | | | |`rocblas_dgemm`|1.5.0| | | | |
 |`cublasDgemmBatched`| | | | |`hipblasDgemmBatched`|1.8.2| | | | |`rocblas_dgemm_batched`|3.5.0| | | | |
 |`cublasDgemmBatched_64`|12.0| | | |`hipblasDgemmBatched_64`|6.3.0| | | |6.3.0|`rocblas_dgemm_batched_64`|6.3.0| | | |6.3.0|
@@ -1605,9 +1605,9 @@
 |`cublasDtrmm_v2`| | | | |`hipblasDtrmm`|3.2.0| |6.0.0| | |`rocblas_dtrmm`|3.5.0| |6.0.0| | |
 |`cublasDtrmm_v2_64`|12.0| | | |`hipblasDtrmm_64`|6.3.0| | | |6.3.0|`rocblas_dtrmm_64`|6.3.0| | | |6.3.0|
 |`cublasDtrsm`| | | | |`hipblasDtrsm`|1.8.2| | | | |`rocblas_dtrsm`|1.5.0| | | | |
-|`cublasDtrsm_64`|12.0| | | | | | | | | |`rocblas_dtrsm_64`|6.2.0| | | | |
+|`cublasDtrsm_64`|12.0| | | |`hipblasDtrsm_64`|6.3.0| | | |6.3.0|`rocblas_dtrsm_64`|6.2.0| | | | |
 |`cublasDtrsm_v2`| | | | |`hipblasDtrsm`|1.8.2| | | | |`rocblas_dtrsm`|1.5.0| | | | |
-|`cublasDtrsm_v2_64`|12.0| | | | | | | | | |`rocblas_dtrsm_64`|6.2.0| | | | |
+|`cublasDtrsm_v2_64`|12.0| | | |`hipblasDtrsm_64`|6.3.0| | | |6.3.0|`rocblas_dtrsm_64`|6.2.0| | | | |
 |`cublasGemmGroupedBatchedEx`|12.5| | | | | | | | | | | | | | | |
 |`cublasGemmGroupedBatchedEx_64`|12.5| | | | | | | | | | | | | | | |
 |`cublasHSHgemvBatched`|11.6| | | | | | | | | |`rocblas_hshgemv_batched`|6.0.0| | | | |
@@ -1657,9 +1657,9 @@
 |`cublasStrmm_v2`| | | | |`hipblasStrmm`|3.2.0| |6.0.0| | |`rocblas_strmm`|3.5.0| |6.0.0| | |
 |`cublasStrmm_v2_64`|12.0| | | |`hipblasStrmm_64`|6.3.0| | | |6.3.0|`rocblas_strmm_64`|6.3.0| | | |6.3.0|
 |`cublasStrsm`| | | | |`hipblasStrsm`|1.8.2| | | | |`rocblas_strsm`|1.5.0| | | | |
-|`cublasStrsm_64`|12.0| | | | | | | | | |`rocblas_strsm_64`|6.2.0| | | | |
+|`cublasStrsm_64`|12.0| | | |`hipblasStrsm_64`|6.3.0| | | |6.3.0|`rocblas_strsm_64`|6.2.0| | | | |
 |`cublasStrsm_v2`| | | | |`hipblasStrsm`|1.8.2| | | | |`rocblas_strsm`|1.5.0| | | | |
-|`cublasStrsm_v2_64`|12.0| | | | | | | | | |`rocblas_strsm_64`|6.2.0| | | | |
+|`cublasStrsm_v2_64`|12.0| | | |`hipblasStrsm_64`|6.3.0| | | |6.3.0|`rocblas_strsm_64`|6.2.0| | | | |
 |`cublasTSSgemvBatched`|11.6| | | | | | | | | |`rocblas_tssgemv_batched`|6.0.0| | | | |
 |`cublasTSSgemvBatched_64`|12.0| | | | | | | | | |`rocblas_tssgemv_batched_64`|6.2.0| | | | |
 |`cublasTSSgemvStridedBatched`|11.6| | | | | | | | | |`rocblas_tssgemv_strided_batched`|6.0.0| | | | |
@@ -1715,9 +1715,9 @@
 |`cublasZtrmm_v2`| | | | |`hipblasZtrmm_v2`|6.0.0| | | | |`rocblas_ztrmm`|3.5.0| |6.0.0| | |
 |`cublasZtrmm_v2_64`|12.0| | | |`hipblasZtrmm_v2_64`|6.3.0| | | |6.3.0|`rocblas_ztrmm_64`|6.3.0| | | |6.3.0|
 |`cublasZtrsm`| | | | |`hipblasZtrsm_v2`|6.0.0| | | | |`rocblas_ztrsm`|3.5.0| | | | |
-|`cublasZtrsm_64`|12.0| | | | | | | | | |`rocblas_ztrsm_64`|6.2.0| | | | |
+|`cublasZtrsm_64`|12.0| | | |`hipblasZtrsm_v2_64`|6.3.0| | | |6.3.0|`rocblas_ztrsm_64`|6.2.0| | | | |
 |`cublasZtrsm_v2`| | | | |`hipblasZtrsm_v2`|6.0.0| | | | |`rocblas_ztrsm`|3.5.0| | | | |
-|`cublasZtrsm_v2_64`|12.0| | | | | | | | | |`rocblas_ztrsm_64`|6.2.0| | | | |
+|`cublasZtrsm_v2_64`|12.0| | | |`hipblasZtrsm_v2_64`|6.3.0| | | |6.3.0|`rocblas_ztrsm_64`|6.2.0| | | | |
 
 ## **8. BLAS-like Extension**
 

diff --git a/src/CUDA2HIP_BLAS_API_functions.cpp b/src/CUDA2HIP_BLAS_API_functions.cpp
@@ -541,13 +541,13 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP {
 
   // TRSM
   {"cublasStrsm",                                          {"hipblasStrsm",                                              "rocblas_strsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_SUPPORTED_V2_ONLY}},
-  {"cublasStrsm_64",                                       {"hipblasStrsm_64",                                           "rocblas_strsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasStrsm_64",                                       {"hipblasStrsm_64",                                           "rocblas_strsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
   {"cublasDtrsm",                                          {"hipblasDtrsm",                                              "rocblas_dtrsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_SUPPORTED_V2_ONLY}},
-  {"cublasDtrsm_64",                                       {"hipblasDtrsm_64",                                           "rocblas_dtrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasDtrsm_64",                                       {"hipblasDtrsm_64",                                           "rocblas_dtrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
   {"cublasCtrsm",                                          {"hipblasCtrsm_v2",                                           "rocblas_ctrsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_SUPPORTED_V2_ONLY}},
-  {"cublasCtrsm_64",                                       {"hipblasCtrsm_64",                                           "rocblas_ctrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasCtrsm_64",                                       {"hipblasCtrsm_v2_64",                                        "rocblas_ctrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
   {"cublasZtrsm",                                          {"hipblasZtrsm_v2",                                           "rocblas_ztrsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_SUPPORTED_V2_ONLY}},
-  {"cublasZtrsm_64",                                       {"hipblasZtrsm_64",                                           "rocblas_ztrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasZtrsm_64",                                       {"hipblasZtrsm_v2_64",                                        "rocblas_ztrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
 
   // TRMM
   {"cublasStrmm",                                          {"hipblasStrmm",                                              "rocblas_strmm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_SUPPORTED_V2_ONLY}},
@@ -908,13 +908,13 @@ const std::map<llvm::StringRef, hipCounter> CUDA_BLAS_FUNCTION_MAP {
 
   // TRSM
   {"cublasStrsm_v2",                                       {"hipblasStrsm",                                              "rocblas_strsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
-  {"cublasStrsm_v2_64",                                    {"hipblasStrsm_64",                                           "rocblas_strsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasStrsm_v2_64",                                    {"hipblasStrsm_64",                                           "rocblas_strsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
   {"cublasDtrsm_v2",                                       {"hipblasDtrsm",                                              "rocblas_dtrsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
-  {"cublasDtrsm_v2_64",                                    {"hipblasDtrsm_64",                                           "rocblas_dtrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasDtrsm_v2_64",                                    {"hipblasDtrsm_64",                                           "rocblas_dtrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
   {"cublasCtrsm_v2",                                       {"hipblasCtrsm_v2",                                           "rocblas_ctrsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
-  {"cublasCtrsm_v2_64",                                    {"hipblasCtrsm_64",                                           "rocblas_ctrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasCtrsm_v2_64",                                    {"hipblasCtrsm_v2_64",                                        "rocblas_ctrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
   {"cublasZtrsm_v2",                                       {"hipblasZtrsm_v2",                                           "rocblas_ztrsm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
-  {"cublasZtrsm_v2_64",                                    {"hipblasZtrsm_64",                                           "rocblas_ztrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3, HIP_UNSUPPORTED}},
+  {"cublasZtrsm_v2_64",                                    {"hipblasZtrsm_v2_64",                                        "rocblas_ztrsm_64",                                   CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
 
   // TRMM
   {"cublasStrmm_v2",                                       {"hipblasStrmm",                                              "rocblas_strmm",                                      CONV_LIB_FUNC, API_BLAS, SEC::BLAS_LEVEL_3}},
@@ -2070,6 +2070,10 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_BLAS_FUNCTION_VER_MAP {
   {"hipblasDtrmm_64",                                      {HIP_6030, HIP_0,    HIP_0,  HIP_LATEST}},
   {"hipblasCtrmm_v2_64",                                   {HIP_6030, HIP_0,    HIP_0,  HIP_LATEST}},
   {"hipblasZtrmm_v2_64",                                   {HIP_6030, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hipblasStrsm_64",                                      {HIP_6030, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hipblasDtrsm_64",                                      {HIP_6030, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hipblasCtrsm_v2_64",                                   {HIP_6030, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hipblasZtrsm_v2_64",                                   {HIP_6030, HIP_0,    HIP_0,  HIP_LATEST}},
 
   {"rocblas_status_to_string",                             {HIP_3050, HIP_0,    HIP_0   }},
   {"rocblas_sscal",                                        {HIP_1050, HIP_0,    HIP_0   }},