vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

iq3_xxs: hopefully fix ROCm

Browse source
This commit is contained in:
Iwan Kawrakow 2024-01-29 16:36:19 +02:00
parent 6efbc690dc
commit 62623434af
1 changed files with 4 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

4
ggml-cuda.cu
View file

@ -191,6 +191,10 @@ static __device__ __forceinline__ int __vsubss4(const int a, const int b) {
#endif // __has_builtin(__builtin_elementwise_sub_sat) #endif // __has_builtin(__builtin_elementwise_sub_sat)
} }
static __device__ __forceinline__ int __vsub4(const int a, const int b) {
return __vsubss4(a, b);
}
static __device__ __forceinline__ int __dp4a(const int a, const int b, int c) { static __device__ __forceinline__ int __dp4a(const int a, const int b, int c) {
#if defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx1030__) #if defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx1030__)
c = __builtin_amdgcn_sdot4(a, b, c, false); c = __builtin_amdgcn_sdot4(a, b, c, false);