From: R0CKSTAR <redacted>
Date: Wed, 11 Sep 2024 01:46:55 +0000 (+0800)
Subject: musa: remove Clang builtins mapping (llama/9421)
X-Git-Tag: upstream/0.0.1642~376
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=5b3d2e65b015486ff7cbbcc7dd2ff2dad14ef904;p=pkg%2Fggml%2Fsources%2Fggml

musa: remove Clang builtins mapping (llama/9421)

Signed-off-by: Xiaodong Ye <redacted>
---

diff --git a/src/ggml-cuda/vendors/musa.h b/src/ggml-cuda/vendors/musa.h
index e50a103a..8df57114 100644
--- a/src/ggml-cuda/vendors/musa.h
+++ b/src/ggml-cuda/vendors/musa.h
@@ -130,42 +130,3 @@
 #define cudaKernelNodeParams musaKernelNodeParams
 #define cudaStreamCaptureModeRelaxed musaStreamCaptureModeRelaxed
 #define cudaStreamEndCapture musaStreamEndCapture
-
-// XXX: Clang builtins mapping
-#define __vsub4   __vsub4_musa
-#define __vcmpeq4 __vcmpeq4_musa
-#define __vcmpne4 __vcmpne4_musa
-
-#ifndef __has_builtin
-    #define __has_builtin(x) 0
-#endif
-
-typedef uint8_t uint8x4_t __attribute__((ext_vector_type(4)));
-
-static __device__ __forceinline__ int __vsub4_musa(const int a, const int b) {
-    return __vsubss4(a, b);
-}
-
-static __device__ __forceinline__ unsigned int __vcmpeq4_musa(unsigned int a, unsigned int b) {
-    const uint8x4_t& va = reinterpret_cast<const uint8x4_t&>(a);
-    const uint8x4_t& vb = reinterpret_cast<const uint8x4_t&>(b);
-    unsigned int c;
-    uint8x4_t& vc = reinterpret_cast<uint8x4_t&>(c);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) {
-        vc[i] = va[i] == vb[i] ? 0xff : 0x00;
-    }
-    return c;
-}
-
-static __device__ __forceinline__ unsigned int __vcmpne4_musa(unsigned int a, unsigned int b) {
-    const uint8x4_t& va = reinterpret_cast<const uint8x4_t&>(a);
-    const uint8x4_t& vb = reinterpret_cast<const uint8x4_t&>(b);
-    unsigned int c;
-    uint8x4_t& vc = reinterpret_cast<uint8x4_t&>(c);
-#pragma unroll
-    for (int i = 0; i < 4; ++i) {
-        vc[i] = va[i] == vb[i] ? 0x00 : 0xff;
-    }
-    return c;
-}