From 98ab3445b7f2542718aa8dfffb3120b78d1b3328 Mon Sep 17 00:00:00 2001 From: Wayne Mitchell Date: Tue, 2 Aug 2022 18:53:35 -0700 Subject: [PATCH] Sycl build fix (#707) Move fill functions in device utils back into the general functions block for use with sycl backend. --- src/utilities/device_utils.c | 178 +++++++++++++++++------------------ 1 file changed, 89 insertions(+), 89 deletions(-) diff --git a/src/utilities/device_utils.c b/src/utilities/device_utils.c index bad4bf7c3..e5899ca07 100644 --- a/src/utilities/device_utils.c +++ b/src/utilities/device_utils.c @@ -1713,6 +1713,95 @@ hypre_ResetDeviceRandGenerator( hypre_ulonglongint seed, #endif /* #if defined(HYPRE_USING_CURAND) || defined(HYPRE_USING_ROCRAND) || defined(HYPRE_USING_ONEMKLRAND) */ +/*-------------------------------------------------------------------- + * hypreGPUKernel_filln + *--------------------------------------------------------------------*/ + +template +__global__ void +hypreGPUKernel_filln(hypre_DeviceItem &item, T *x, size_t n, T v) +{ + HYPRE_Int i = hypre_gpu_get_grid_thread_id<1, 1>(item); + + if (i < n) + { + x[i] = v; + } +} + +/*-------------------------------------------------------------------- + * hypreDevice_Filln + *--------------------------------------------------------------------*/ + +template +HYPRE_Int +hypreDevice_Filln(T *d_x, size_t n, T v) +{ +#if 0 + HYPRE_THRUST_CALL( fill_n, d_x, n, v); +#else + if (n <= 0) + { + return hypre_error_flag; + } + + dim3 bDim = hypre_GetDefaultDeviceBlockDimension(); + dim3 gDim = hypre_GetDefaultDeviceGridDimension(n, "thread", bDim); + + HYPRE_GPU_LAUNCH( hypreGPUKernel_filln, gDim, bDim, d_x, n, v ); +#endif + + return hypre_error_flag; +} + +/*-------------------------------------------------------------------- + * hypreDevice_ComplexFilln + *--------------------------------------------------------------------*/ + +HYPRE_Int +hypreDevice_ComplexFilln( HYPRE_Complex *d_x, + size_t n, + HYPRE_Complex v ) +{ + return hypreDevice_Filln(d_x, n, v); +} + +/*-------------------------------------------------------------------- + * hypreDevice_CharFilln + *--------------------------------------------------------------------*/ + +HYPRE_Int +hypreDevice_CharFilln( char *d_x, + size_t n, + char v ) +{ + return hypreDevice_Filln(d_x, n, v); +} + +/*-------------------------------------------------------------------- + * hypreDevice_IntFilln + *--------------------------------------------------------------------*/ + +HYPRE_Int +hypreDevice_IntFilln( HYPRE_Int *d_x, + size_t n, + HYPRE_Int v ) +{ + return hypreDevice_Filln(d_x, n, v); +} + +/*-------------------------------------------------------------------- + * hypreDevice_BigIntFilln + *--------------------------------------------------------------------*/ + +HYPRE_Int +hypreDevice_BigIntFilln( HYPRE_BigInt *d_x, + size_t n, + HYPRE_BigInt v) +{ + return hypreDevice_Filln(d_x, n, v); +} + #endif // #if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL) /* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -1882,95 +1971,6 @@ hypreDevice_ComplexScalen( HYPRE_Complex *d_x, return hypreDevice_Scalen(d_x, n, d_y, v); } -/*-------------------------------------------------------------------- - * hypreGPUKernel_filln - *--------------------------------------------------------------------*/ - -template -__global__ void -hypreGPUKernel_filln(hypre_DeviceItem &item, T *x, size_t n, T v) -{ - HYPRE_Int i = hypre_gpu_get_grid_thread_id<1, 1>(item); - - if (i < n) - { - x[i] = v; - } -} - -/*-------------------------------------------------------------------- - * hypreDevice_Filln - *--------------------------------------------------------------------*/ - -template -HYPRE_Int -hypreDevice_Filln(T *d_x, size_t n, T v) -{ -#if 0 - HYPRE_THRUST_CALL( fill_n, d_x, n, v); -#else - if (n <= 0) - { - return hypre_error_flag; - } - - dim3 bDim = hypre_GetDefaultDeviceBlockDimension(); - dim3 gDim = hypre_GetDefaultDeviceGridDimension(n, "thread", bDim); - - HYPRE_GPU_LAUNCH( hypreGPUKernel_filln, gDim, bDim, d_x, n, v ); -#endif - - return hypre_error_flag; -} - -/*-------------------------------------------------------------------- - * hypreDevice_ComplexFilln - *--------------------------------------------------------------------*/ - -HYPRE_Int -hypreDevice_ComplexFilln( HYPRE_Complex *d_x, - size_t n, - HYPRE_Complex v ) -{ - return hypreDevice_Filln(d_x, n, v); -} - -/*-------------------------------------------------------------------- - * hypreDevice_CharFilln - *--------------------------------------------------------------------*/ - -HYPRE_Int -hypreDevice_CharFilln( char *d_x, - size_t n, - char v ) -{ - return hypreDevice_Filln(d_x, n, v); -} - -/*-------------------------------------------------------------------- - * hypreDevice_IntFilln - *--------------------------------------------------------------------*/ - -HYPRE_Int -hypreDevice_IntFilln( HYPRE_Int *d_x, - size_t n, - HYPRE_Int v ) -{ - return hypreDevice_Filln(d_x, n, v); -} - -/*-------------------------------------------------------------------- - * hypreDevice_BigIntFilln - *--------------------------------------------------------------------*/ - -HYPRE_Int -hypreDevice_BigIntFilln( HYPRE_BigInt *d_x, - size_t n, - HYPRE_BigInt v) -{ - return hypreDevice_Filln(d_x, n, v); -} - /*-------------------------------------------------------------------- * hypreGPUKernel_StridedCopy *--------------------------------------------------------------------*/