[BACKEND] KUNLUNXIN xpu update to Commit 9855424(20251203) (#164)

Author: Ason93

third_party/xpu/device/trigonometric.xpu

@@ -65,7 +65,7 @@ __device__ inline void printBool(bool b, int customIdx) {
 __device__ inline void printFloat(float a, int customIdx) {
   int clusterId = cluster_id();
   int coreId = core_id();
-  if (clusterId == __CHECK_CLUSTER_ID && coreId == 0) {
+  if (clusterId == __CHECK_CLUSTER_ID && coreId < 8) {
     printf("[printFloat_%d] cluster_id = %d, core_id = %d, local_x[0] = %f \n",
            customIdx, clusterId, coreId, a);
   }
@@ -76,7 +76,7 @@ __device__ inline void printFloat_spec(float a, int clu_id, int cid,
   int clusterId = cluster_id();
   int coreId = core_id();
   if (clusterId == __CHECK_CLUSTER_ID && coreId == cid) {
-    printf("[printFloat_%d] cluster_id = %d, core_id = %d, local_x[0] = %f \n",
+    printf("[printFloat_spec_%d] cluster_id = %d, core_id = %d, local_x[0] = %f \n",
            customIdx, clusterId, coreId, a);
   }
 }
@@ -93,21 +93,41 @@ __device__ inline void printFloat_all(float a, int customIdx) {
 __device__ inline void printInt(int a, int customIdx) {
   int clusterId = cluster_id();
   int coreId = core_id();
-  if (coreId == 0) {
+  if (clusterId == __CHECK_CLUSTER_ID && coreId < 8) {
     printf("[printInt_%d] cluster_id = %d, core_id = %d, local_x[0] = %d \n",
            customIdx, clusterId, coreId, a);
   }
 }
 
+__device__ inline void printInt_spec(int a, int clu_id, int cid,
+                                       int customIdx) {
+  int clusterId = cluster_id();
+  int coreId = core_id();
+  if (clusterId == __CHECK_CLUSTER_ID && coreId == cid) {
+    printf("[printInt_spec_%d] cluster_id = %d, core_id = %d, local_x[0] = %d \n",
+           customIdx, clusterId, coreId, a);
+  }
+}
+
 __device__ inline void printInt64(int64_t a, int customIdx) {
   int clusterId = cluster_id();
   int coreId = core_id();
-  if (clusterId == __CHECK_CLUSTER_ID && coreId == 0) {
+  if (clusterId == __CHECK_CLUSTER_ID && coreId < 8) {
     printf("[printInt64_%d] cluster_id = %d, core_id = %d, local_x[0] = %ld \n",
            customIdx, clusterId, coreId, a);
   }
 }
 
+__device__ inline void printInt64_spec(int64_t a, int clu_id, int cid,
+                                       int customIdx) {
+  int clusterId = cluster_id();
+  int coreId = core_id();
+  if (clusterId == __CHECK_CLUSTER_ID && coreId == cid) {
+    printf("[printInt64_spec_%d] cluster_id = %d, core_id = %d, local_x[0] = %ld \n",
+           customIdx, clusterId, coreId, a);
+  }
+}
+
 __device__ inline void printMMaOp(int a, int b, int c, int isAcc) {
   int clusterId = cluster_id();
   int coreId = core_id();
@@ -287,6 +307,27 @@ __device__ inline void vstore2_lm(bfloat16* ptr, float32x16_t vl, float32x16_t v
     vstore_lm_float16x32(reinterpret_cast<float16*>(ptr), reinterpret_cast<float16x32_t>(vl));
 }
 
+
+__device__ inline void vstore2_lm_unordered(bfloat16* ptr, float32x16_t veven, float32x16_t vodd) {
+    int mask_even = sveq_uint32x16(0, svand_uint32x16(0x10000, reinterpret_cast<uint32x16_t>(veven)));
+    int mask_odd = sveq_uint32x16(0, svand_uint32x16(0x10000, reinterpret_cast<uint32x16_t>(vodd)));
+
+    veven = reinterpret_cast<float32x16_t>(
+            svadd_uint32x16_mh(0x8000, reinterpret_cast<uint32x16_t>(veven), reinterpret_cast<uint32x16_t>(veven), ~mask_even));
+    vodd = reinterpret_cast<float32x16_t>(
+            svadd_uint32x16_mh(0x8000, reinterpret_cast<uint32x16_t>(vodd), reinterpret_cast<uint32x16_t>(vodd), ~mask_odd));
+    veven = reinterpret_cast<float32x16_t>(
+            svadd_uint32x16_mh(0x7FFF, reinterpret_cast<uint32x16_t>(veven), reinterpret_cast<uint32x16_t>(veven), mask_even));
+    vodd = reinterpret_cast<float32x16_t>(
+            svadd_uint32x16_mh(0x7FFF, reinterpret_cast<uint32x16_t>(vodd), reinterpret_cast<uint32x16_t>(vodd), mask_odd));
+
+    constexpr int mask = 0xaaaaaaaa;    // 0b10101010101010101010101010101010
+    vstore_lm_int16x32_mh(ptr, reinterpret_cast<int16x32_t>(veven), mask);
+    constexpr int pose = 16;
+    __asm__ __volatile__("vsrl.p %0, %1, %2" : "=&v"(vodd) : "r"(pose), "v"(vodd));
+    vstore_lm_int16x32_mh(ptr, reinterpret_cast<int16x32_t>(vodd), (~mask));
+}
+
 static __device__ inline void taylor_sin(float *C1, float *C3, float *C5,
                                          float *C7, float *C9) {
   *C1 = 1;