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

ldg only for sm >= 35 (#516)

Browse Source 
This PR fixes compile issues with CUDA sm_30. See #511
This commit is contained in:
Ruipeng Li 2022-01-12 08:55:58 -08:00 committed by GitHub
parent 436e09cba2
commit bcccb117ef
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
2 changed files with 8 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/utilities/_hypre_utilities.hpp
View File

@ -694,7 +694,11 @@ template <typename T>
static __device__ __forceinline__ static __device__ __forceinline__
T read_only_load( const T *ptr ) T read_only_load( const T *ptr )
{ {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
return __ldg( ptr ); return __ldg( ptr );
#else
return *ptr;
#endif
} }
/* exclusive prefix scan */ /* exclusive prefix scan */

4
src/utilities/device_utils.h
View File

@ -637,7 +637,11 @@ template <typename T>
static __device__ __forceinline__ static __device__ __forceinline__
T read_only_load( const T *ptr ) T read_only_load( const T *ptr )
{ {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350
return __ldg( ptr ); return __ldg( ptr );
#else
return *ptr;
#endif
} }
/* exclusive prefix scan */ /* exclusive prefix scan */