From b9d3eccbe3d5c19da705681d44ba96d2db41500a Mon Sep 17 00:00:00 2001 From: "Victor A. Paludetto Magri" <50467563+victorapm@users.noreply.github.com> Date: Thu, 5 Jan 2023 19:02:11 -0500 Subject: [PATCH] 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 | 5 + src/utilities/int_array.c | 200 +++++++++++++++++++++++++++++++ src/utilities/int_array_device.c | 49 ++++++++ src/utilities/protos.h | 5 + 4 files changed, 259 insertions(+) diff --git a/src/utilities/_hypre_utilities.h b/src/utilities/_hypre_utilities.h index 3a589a0c5..32a1e1252 100644 --- a/src/utilities/_hypre_utilities.h +++ b/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 */ diff --git a/src/utilities/int_array.c b/src/utilities/int_array.c index f7aee530d..dc347bcb2 100644 --- a/src/utilities/int_array.c +++ b/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; +} diff --git a/src/utilities/int_array_device.c b/src/utilities/int_array_device.c index 05c3dff79..cea33c09c 100644 --- a/src/utilities/int_array_device.c +++ b/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 *--------------------------------------------------------------------------*/ diff --git a/src/utilities/protos.h b/src/utilities/protos.h index 5d774d02b..0b85efc33 100644 --- a/src/utilities/protos.h +++ b/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 */