CFD/hypre
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Sycl build fix (#707)

Browse Source 
Move fill functions in device utils back into the general functions block for use with sycl backend.
This commit is contained in:
Wayne Mitchell 2022-08-02 18:53:35 -07:00 committed by GitHub
parent 6425099995
commit 98ab3445b7
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 89 additions and 89 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

178
src/utilities/device_utils.c
View File

@ -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<typename T>
__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<typename T>
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<typename T>
__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<typename T>
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
*--------------------------------------------------------------------*/