Add new hypre_IntArray functions (#807)

The new functions are:

* Add hypre_IntArrayPrint
* Add hypre_IntArrayMigrate
* Add hypre_IntArrayRead
* Add hypre_IntArrayInverseMapping

src/utilities/_hypre_utilities.h

@@ -1950,18 +1950,23 @@ HYPRE_Int hypre_IntArrayCopy( hypre_IntArray *x, hypre_IntArray *y );
 hypre_IntArray* hypre_IntArrayCloneDeep_v2( hypre_IntArray *x,
                                             HYPRE_MemoryLocation memory_location );
 hypre_IntArray* hypre_IntArrayCloneDeep( hypre_IntArray *x );
+HYPRE_Int hypre_IntArrayMigrate( hypre_IntArray *v, HYPRE_MemoryLocation memory_location );
+HYPRE_Int hypre_IntArrayPrint( MPI_Comm comm, hypre_IntArray *array, const char *filename );
+HYPRE_Int hypre_IntArrayRead( MPI_Comm comm, const char *filename, hypre_IntArray **array_ptr );
 HYPRE_Int hypre_IntArraySetConstantValuesHost( hypre_IntArray *v, HYPRE_Int value );
 HYPRE_Int hypre_IntArraySetConstantValues( hypre_IntArray *v, HYPRE_Int value );
 HYPRE_Int hypre_IntArrayCountHost( hypre_IntArray *v, HYPRE_Int value,
                                    HYPRE_Int *num_values_ptr );
 HYPRE_Int hypre_IntArrayCount( hypre_IntArray *v, HYPRE_Int value,
                                HYPRE_Int *num_values_ptr );
+HYPRE_Int hypre_IntArrayInverseMapping( hypre_IntArray *v, hypre_IntArray **w_ptr );
 
 /* int_array_device.c */
 #if defined(HYPRE_USING_GPU)
 HYPRE_Int hypre_IntArraySetConstantValuesDevice( hypre_IntArray *v, HYPRE_Int value );
 HYPRE_Int hypre_IntArrayCountDevice ( hypre_IntArray *v, HYPRE_Int value,
                                       HYPRE_Int *num_values_ptr );
+HYPRE_Int hypre_IntArrayInverseMappingDevice( hypre_IntArray *v, hypre_IntArray *w );
 #endif
 
 /* memory_tracker.c */

src/utilities/int_array.c

@@ -138,6 +138,135 @@ hypre_IntArrayCloneDeep( hypre_IntArray *x )
    return hypre_IntArrayCloneDeep_v2(x, hypre_IntArrayMemoryLocation(x));
 }
 
+/*--------------------------------------------------------------------------
+ * hypre_IntArrayMigrate
+ *--------------------------------------------------------------------------*/
+
+HYPRE_Int
+hypre_IntArrayMigrate( hypre_IntArray      *v,
+                       HYPRE_MemoryLocation memory_location )
+{
+   HYPRE_Int            size                = hypre_IntArraySize(v);
+   HYPRE_Int           *v_data              = hypre_IntArrayData(v);
+   HYPRE_MemoryLocation old_memory_location = hypre_IntArrayMemoryLocation(v);
+
+   HYPRE_Int           *w_data;
+
+   /* Update v's memory location */
+   hypre_IntArrayMemoryLocation(v) = memory_location;
+
+   if ( hypre_GetActualMemLocation(memory_location) !=
+        hypre_GetActualMemLocation(old_memory_location) )
+   {
+      w_data = hypre_TAlloc(HYPRE_Int, size, memory_location);
+      hypre_TMemcpy(w_data, v_data, HYPRE_Int, size,
+                    memory_location, old_memory_location);
+      hypre_TFree(v_data, old_memory_location);
+      hypre_IntArrayData(v) = w_data;
+   }
+
+   return hypre_error_flag;
+}
+
+/*--------------------------------------------------------------------------
+ * hypre_IntArrayPrint
+ *--------------------------------------------------------------------------*/
+
+HYPRE_Int
+hypre_IntArrayPrint( MPI_Comm        comm,
+                     hypre_IntArray *array,
+                     const char     *filename )
+{
+   HYPRE_Int             size            = hypre_IntArraySize(array);
+   HYPRE_MemoryLocation  memory_location = hypre_IntArrayMemoryLocation(array);
+
+   hypre_IntArray       *h_array;
+   HYPRE_Int            *data;
+
+   FILE                 *file;
+   HYPRE_Int             i, myid;
+   char                  new_filename[1024];
+
+   hypre_MPI_Comm_rank(comm, &myid);
+
+   /* Move data to host if needed*/
+   h_array = (hypre_GetActualMemLocation(memory_location) == hypre_MEMORY_DEVICE) ?
+             hypre_IntArrayCloneDeep_v2(array, HYPRE_MEMORY_HOST) : array;
+   data = hypre_IntArrayData(h_array);
+
+   /* Open file */
+   hypre_sprintf(new_filename, "%s.%05d", filename, myid);
+   if ((file = fopen(new_filename, "w")) == NULL)
+   {
+      hypre_error_w_msg(HYPRE_ERROR_GENERIC, "Error: can't open output file\n");
+      return hypre_error_flag;
+   }
+
+   /* Print to file */
+   hypre_fprintf(file, "%d\n", size);
+   for (i = 0; i < size; i++)
+   {
+      hypre_fprintf(file, "%d\n", data[i]);
+   }
+   fclose(file);
+
+   /* Free memory */
+   if (h_array != array)
+   {
+      hypre_IntArrayDestroy(h_array);
+   }
+
+   return hypre_error_flag;
+}
+
+/*--------------------------------------------------------------------------
+ * hypre_IntArrayRead
+ *--------------------------------------------------------------------------*/
+
+HYPRE_Int
+hypre_IntArrayRead( MPI_Comm         comm,
+                    const char      *filename,
+                    hypre_IntArray **array_ptr )
+{
+   hypre_IntArray       *array;
+   HYPRE_Int             size;
+   FILE                 *file;
+   HYPRE_Int             i, myid;
+   char                  new_filename[1024];
+
+   hypre_MPI_Comm_rank(comm, &myid);
+
+   /* Open file */
+   hypre_sprintf(new_filename, "%s.%05d", filename, myid);
+   if ((file = fopen(new_filename, "r")) == NULL)
+   {
+      hypre_error_w_msg(HYPRE_ERROR_GENERIC, "Error: can't open input file\n");
+      return hypre_error_flag;
+   }
+
+   /* Read array size from file */
+   hypre_fscanf(file, "%d\n", &size);
+
+   /* Create IntArray on the host */
+   array = hypre_IntArrayCreate(size);
+   hypre_IntArrayInitialize_v2(array, HYPRE_MEMORY_HOST);
+
+   /* Read array values from file */
+   for (i = 0; i < size; i++)
+   {
+      hypre_fscanf(file, "%d\n", &hypre_IntArrayData(array)[i]);
+   }
+   fclose(file);
+
+   /* Migrate to final memory location */
+   hypre_IntArrayMigrate(array, hypre_HandleMemoryLocation(hypre_handle()));
+
+   /* Set output pointer */
+   *array_ptr = array;
+
+   return hypre_error_flag;
+}
+
 /*--------------------------------------------------------------------------
  * hypre_IntArraySetConstantValuesHost
  *--------------------------------------------------------------------------*/
@@ -160,6 +289,7 @@ hypre_IntArraySetConstantValuesHost( hypre_IntArray *v,
 
    return hypre_error_flag;
 }
+
 /*--------------------------------------------------------------------------
  * hypre_IntArraySetConstantValues
  *--------------------------------------------------------------------------*/
@@ -246,3 +376,73 @@ hypre_IntArrayCount( hypre_IntArray *v,
 
    return hypre_error_flag;
 }
+
+/*--------------------------------------------------------------------------
+ * hypre_IntArrayInverseMappingHost
+ *--------------------------------------------------------------------------*/
+
+HYPRE_Int
+hypre_IntArrayInverseMappingHost( hypre_IntArray  *v,
+                                  hypre_IntArray  *w )
+{
+   HYPRE_Int   size    = hypre_IntArraySize(v);
+   HYPRE_Int  *v_data  = hypre_IntArrayData(v);
+   HYPRE_Int  *w_data  = hypre_IntArrayData(w);
+
+   HYPRE_Int   i;
+
+#if defined(HYPRE_USING_OPENMP)
+   #pragma omp parallel for private(i) HYPRE_SMP_SCHEDULE
+#endif
+   for (i = 0; i < size; i++)
+   {
+      w_data[v_data[i]] = i;
+   }
+
+   return hypre_error_flag;
+}
+
+/*--------------------------------------------------------------------------
+ * hypre_IntArrayInverseMapping
+ *
+ * Compute the reverse mapping (w) given an input array (v)
+ *--------------------------------------------------------------------------*/
+
+HYPRE_Int
+hypre_IntArrayInverseMapping( hypre_IntArray  *v,
+                              hypre_IntArray **w_ptr )
+{
+   HYPRE_Int             size = hypre_IntArraySize(v);
+   HYPRE_MemoryLocation  memory_location = hypre_IntArrayMemoryLocation(v);
+   hypre_IntArray       *w;
+
+   /* Create and initialize output array */
+   w = hypre_IntArrayCreate(size);
+   hypre_IntArrayInitialize_v2(w, memory_location);
+
+   /* Exit if array has no elements */
+   if (hypre_IntArraySize(w) <= 0)
+   {
+      *w_ptr = w;
+
+      return hypre_error_flag;
+   }
+
+#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL)
+   HYPRE_ExecutionPolicy exec = hypre_GetExecPolicy1(memory_location);
+
+   if (exec == HYPRE_EXEC_DEVICE)
+   {
+      hypre_IntArrayInverseMappingDevice(v, w);
+   }
+   else
+#endif
+   {
+      hypre_IntArrayInverseMappingHost(v, w);
+   }
+
+   /* Set output pointer */
+   *w_ptr = w;
+
+   return hypre_error_flag;
+}

src/utilities/int_array_device.c

@@ -39,6 +39,55 @@ hypre_IntArraySetConstantValuesDevice( hypre_IntArray *v,
    return hypre_error_flag;
 }
 
+/*--------------------------------------------------------------------------
+ * hypreGPUKernel_IntArrayInverseMapping
+ *--------------------------------------------------------------------------*/
+
+__global__ void
+hypreGPUKernel_IntArrayInverseMapping( hypre_DeviceItem  &item,
+                                       HYPRE_Int          size,
+                                       HYPRE_Int         *v_data,
+                                       HYPRE_Int         *w_data )
+{
+   HYPRE_Int i = hypre_gpu_get_grid_thread_id<1, 1>(item);
+
+   if (i < size)
+   {
+      w_data[v_data[i]] = i;
+   }
+}
+
+/*--------------------------------------------------------------------------
+ * hypre_IntArrayInverseMappingDevice
+ *--------------------------------------------------------------------------*/
+
+HYPRE_Int
+hypre_IntArrayInverseMappingDevice( hypre_IntArray  *v,
+                                    hypre_IntArray  *w )
+{
+   HYPRE_Int   size    = hypre_IntArraySize(v);
+   HYPRE_Int  *v_data  = hypre_IntArrayData(v);
+   HYPRE_Int  *w_data  = hypre_IntArrayData(w);
+
+#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL)
+   dim3 bDim = hypre_GetDefaultDeviceBlockDimension();
+   dim3 gDim = hypre_GetDefaultDeviceGridDimension(size, "thread", bDim);
+
+   HYPRE_GPU_LAUNCH( hypreGPUKernel_IntArrayInverseMapping, gDim, bDim, size, v_data, w_data );
+
+#elif defined(HYPRE_USING_DEVICE_OPENMP)
+   HYPRE_Int i;
+
+   #pragma omp target teams distribute parallel for private(i) is_device_ptr(v_data, w_data)
+   for (i = 0; i < size; i++)
+   {
+      w_data[v_data[i]] = i;
+   }
+#endif
+
+   return hypre_error_flag;
+}
+
 /*--------------------------------------------------------------------------
  * hypre_IntArrayCountDevice
  *--------------------------------------------------------------------------*/

src/utilities/protos.h

@@ -384,18 +384,23 @@ HYPRE_Int hypre_IntArrayCopy( hypre_IntArray *x, hypre_IntArray *y );
 hypre_IntArray* hypre_IntArrayCloneDeep_v2( hypre_IntArray *x,
                                             HYPRE_MemoryLocation memory_location );
 hypre_IntArray* hypre_IntArrayCloneDeep( hypre_IntArray *x );
+HYPRE_Int hypre_IntArrayMigrate( hypre_IntArray *v, HYPRE_MemoryLocation memory_location );
+HYPRE_Int hypre_IntArrayPrint( MPI_Comm comm, hypre_IntArray *array, const char *filename );
+HYPRE_Int hypre_IntArrayRead( MPI_Comm comm, const char *filename, hypre_IntArray **array_ptr );
 HYPRE_Int hypre_IntArraySetConstantValuesHost( hypre_IntArray *v, HYPRE_Int value );
 HYPRE_Int hypre_IntArraySetConstantValues( hypre_IntArray *v, HYPRE_Int value );
 HYPRE_Int hypre_IntArrayCountHost( hypre_IntArray *v, HYPRE_Int value,
                                    HYPRE_Int *num_values_ptr );
 HYPRE_Int hypre_IntArrayCount( hypre_IntArray *v, HYPRE_Int value,
                                HYPRE_Int *num_values_ptr );
+HYPRE_Int hypre_IntArrayInverseMapping( hypre_IntArray *v, hypre_IntArray **w_ptr );
 
 /* int_array_device.c */
 #if defined(HYPRE_USING_GPU)
 HYPRE_Int hypre_IntArraySetConstantValuesDevice( hypre_IntArray *v, HYPRE_Int value );
 HYPRE_Int hypre_IntArrayCountDevice ( hypre_IntArray *v, HYPRE_Int value,
                                       HYPRE_Int *num_values_ptr );
+HYPRE_Int hypre_IntArrayInverseMappingDevice( hypre_IntArray *v, hypre_IntArray *w );
 #endif
 
 /* memory_tracker.c */