ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'develop' into amoeba-gpu

Browse Source
This commit is contained in:
Trung Nguyen
2022-10-01 00:38:24 -05:00
parent 9a1f23a079 f136bcbef7
commit 2ef6a59c0a
1805 changed files with 403466 additions and 63446 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
lib/gpu/lal_pre_cuda_hip.h
View File

@ -30,7 +30,7 @@
// -------------------------------------------------------------------------
#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
#define CONFIG_ID 303
#define SIMD_SIZE 64
#else
@ -112,7 +112,7 @@
// KERNEL MACROS - TEXTURES
// -------------------------------------------------------------------------
#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
#define _texture(name, type) __device__ type* name
#define _texture_2d(name, type) __device__ type* name
#else
@ -134,9 +134,12 @@
int2 qt = tex1Dfetch(q_tex,i); \
ans=__hiloint2double(qt.y, qt.x); \
}
#elif defined(__HIP_PLATFORM_SPIRV__)
#define fetch4(ans,i,pos_tex) tex1Dfetch(&ans, pos_tex, i);
#define fetch(ans,i,q_tex) tex1Dfetch(&ans, q_tex,i);
#else
#define fetch4(ans,i,pos_tex) ans=tex1Dfetch(pos_tex, i);
#define fetch(ans,i,q_tex) ans=tex1Dfetch(q_tex,i);
#define fetch4(ans,i,pos_tex) ans=tex1Dfetch(pos_tex, i);
#define fetch(ans,i,q_tex) ans=tex1Dfetch(q_tex,i);
#endif
#else
#define fetch4(ans,i,x) ans=x[i]
@ -152,7 +155,7 @@
#define mu_tex mu_
#endif
#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
#if defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
#undef fetch4
#undef fetch
@ -211,7 +214,7 @@
#endif
#endif
#if defined(CUDA_PRE_NINE) || defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)
#if defined(CUDA_PRE_NINE) || defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__) || defined(__HIP_PLATFORM_SPIRV__)
#ifdef _SINGLE_SINGLE
#define shfl_down __shfl_down