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

Merge branch 'sycl_par_matmat_domp' of github.com:hypre-space/hypre into parspgemm

Browse Source
This commit is contained in:
Ruipeng Li 2022-03-30 09:19:52 -07:00
commit 23c7777045
3 changed files with 17 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
src/utilities/_hypre_utilities.hpp
View File

@ -236,10 +236,9 @@ using dim3 = sycl::range<1>;
{ \
if ( gridsize[0] == 0 || blocksize[0] == 0 ) \
{ \
hypre_printf("Error %s %d: Invalid SYCL 1D launch parameters grid/block (%d) (%d)\n", \
/* hypre_printf("Warning %s %d: Zero SYCL 1D launch parameters grid/block (%d) (%d)\n", \
__FILE__, __LINE__, \
gridsize[0], blocksize[0]); \
assert(0); exit(1); \
gridsize[0], blocksize[0]); */ \
} \
else \
{ \
@ -643,10 +642,10 @@ struct hypre_GpuMatData
#endif //#if defined(HYPRE_USING_GPU)
/* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
* WM: generic device functions (cuda/hip/sycl)
* WM: Q - what about device openmp?
* generic device functions (cuda/hip/sycl)
* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL)
template <typename T>
static __device__ __forceinline__
T read_only_load( const T *ptr )
@ -657,9 +656,10 @@ T read_only_load( const T *ptr )
return *ptr;
#endif
}
#endif
/* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
* WM: cuda/hip functions
* cuda/hip functions
* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
@ -1234,6 +1234,7 @@ struct TupleComp3
* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */
/* device_utils.c */
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL)
dim3 hypre_GetDefaultDeviceBlockDimension();
dim3 hypre_GetDefaultDeviceGridDimension( HYPRE_Int n, const char *granularity, dim3 bDim );
@ -1247,6 +1248,7 @@ hypreDevice_StableSortTupleByTupleKey(HYPRE_Int N, T1 *keys1, T2 *keys2, T3 *val
template <typename T1, typename T2, typename T3> HYPRE_Int hypreDevice_ReduceByTupleKey(HYPRE_Int N,
T1 *keys1_in, T2 *keys2_in, T3 *vals_in, T1 *keys1_out, T2 *keys2_out, T3 *vals_out);
#endif
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)

3
src/utilities/device_utils.c
View File

@ -14,6 +14,8 @@
* generic device functions (cuda/hip/sycl)
* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL)
dim3
hypre_GetDefaultDeviceBlockDimension()
{
@ -628,6 +630,7 @@ hypreDevice_ReduceByTupleKey(HYPRE_Int N, T1 *keys1_in, T2 *keys2_in, T3 *vals
template HYPRE_Int hypreDevice_ReduceByTupleKey(HYPRE_Int N, HYPRE_Int *keys1_in,
HYPRE_Int *keys2_in, HYPRE_Complex *vals_in, HYPRE_Int *keys1_out, HYPRE_Int *keys2_out,
HYPRE_Complex *vals_out);
#endif
/* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
* cuda/hip functions

9
src/utilities/device_utils.h
View File

@ -585,10 +585,10 @@ struct hypre_GpuMatData
#endif //#if defined(HYPRE_USING_GPU)
/* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
* WM: generic device functions (cuda/hip/sycl)
* WM: Q - what about device openmp?
* generic device functions (cuda/hip/sycl)
* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL)
template <typename T>
static __device__ __forceinline__
T read_only_load( const T *ptr )
@ -599,9 +599,10 @@ T read_only_load( const T *ptr )
return *ptr;
#endif
}
#endif
/* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
* WM: cuda/hip functions
* cuda/hip functions
* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
@ -1176,6 +1177,7 @@ struct TupleComp3
* - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - */
/* device_utils.c */
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP) || defined(HYPRE_USING_SYCL)
dim3 hypre_GetDefaultDeviceBlockDimension();
dim3 hypre_GetDefaultDeviceGridDimension( HYPRE_Int n, const char *granularity, dim3 bDim );
@ -1189,6 +1191,7 @@ hypreDevice_StableSortTupleByTupleKey(HYPRE_Int N, T1 *keys1, T2 *keys2, T3 *val
template <typename T1, typename T2, typename T3> HYPRE_Int hypreDevice_ReduceByTupleKey(HYPRE_Int N,
T1 *keys1_in, T2 *keys2_in, T3 *vals_in, T1 *keys1_out, T2 *keys2_out, T3 *vals_out);
#endif
#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)