From 152f33fb57756b494ba1c6518820ebf2283fe914 Mon Sep 17 00:00:00 2001
From: PanZezhong <panzezhong@qiyuanlab.com>
Date: Wed, 8 Jul 2026 10:39:21 +0800
Subject: [PATCH] issue/1376 deltanet for hygon

---
 .../ops/chunk_gated_delta_rule/operator.cc         | 14 +++++++++++++-
 .../ops/recurrent_gated_delta_rule/cuda/kernel.cuh |  4 ++--
 .../ops/recurrent_gated_delta_rule/operator.cc     | 14 +++++++++++++-
 3 files changed, 28 insertions(+), 4 deletions(-)

diff --git a/src/infiniop/ops/chunk_gated_delta_rule/operator.cc b/src/infiniop/ops/chunk_gated_delta_rule/operator.cc
index 631d9ca44..7e5bd04d0 100644
--- a/src/infiniop/ops/chunk_gated_delta_rule/operator.cc
+++ b/src/infiniop/ops/chunk_gated_delta_rule/operator.cc
@@ -4,7 +4,7 @@
 #include "../../handle.h"
 #include "infiniop/ops/chunk_gated_delta_rule.h"
 
-#if defined(ENABLE_NVIDIA_API)
+#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_HYGON_API)
 #include "nvidia/chunk_gated_delta_rule_nvidia.cuh"
 #endif
 
@@ -41,6 +41,9 @@ __INFINI_C infiniStatus_t infiniopCreateChunkGatedDeltaRuleDescriptor(
 #ifdef ENABLE_NVIDIA_API
         CREATE(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        CREATE(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -61,6 +64,9 @@ __INFINI_C infiniStatus_t infiniopGetChunkGatedDeltaRuleWorkspaceSize(
 #ifdef ENABLE_NVIDIA_API
         GET(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        GET(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -88,6 +94,9 @@ __INFINI_C infiniStatus_t infiniopChunkGatedDeltaRule(
 #ifdef ENABLE_NVIDIA_API
         CALCULATE(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        CALCULATE(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -107,6 +116,9 @@ __INFINI_C infiniStatus_t infiniopDestroyChunkGatedDeltaRuleDescriptor(
 #ifdef ENABLE_NVIDIA_API
         DESTROY(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        DESTROY(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
diff --git a/src/infiniop/ops/recurrent_gated_delta_rule/cuda/kernel.cuh b/src/infiniop/ops/recurrent_gated_delta_rule/cuda/kernel.cuh
index 211e508df..9e952eb5a 100644
--- a/src/infiniop/ops/recurrent_gated_delta_rule/cuda/kernel.cuh
+++ b/src/infiniop/ops/recurrent_gated_delta_rule/cuda/kernel.cuh
@@ -35,9 +35,9 @@ __device__ inline float loadAsFloat<cuda_bfloat16>(const cuda_bfloat16 *ptr, ptr
 __device__ inline float warpReduceSum(float value) {
 #pragma unroll
     for (int offset = 16; offset > 0; offset >>= 1) {
-        value += __shfl_down_sync(0xffffffff, value, offset);
+        value += __shfl_down_sync(0xffffffff, value, offset, 32);
     }
-    return __shfl_sync(0xffffffff, value, 0);
+    return __shfl_sync(0xffffffff, value, 0, 32);
 }
 template <typename Tdata, typename Tgate, typename Tcompute, size_t Dk, size_t Dv, size_t WARPS_PER_BLOCK>
 __device__ void recurrentGatedDeltaRuleIndexedPoolWarpKernel(
diff --git a/src/infiniop/ops/recurrent_gated_delta_rule/operator.cc b/src/infiniop/ops/recurrent_gated_delta_rule/operator.cc
index 465e8d6c9..84d061ed9 100644
--- a/src/infiniop/ops/recurrent_gated_delta_rule/operator.cc
+++ b/src/infiniop/ops/recurrent_gated_delta_rule/operator.cc
@@ -4,7 +4,7 @@
 #include "../../handle.h"
 #include "infiniop/ops/recurrent_gated_delta_rule.h"
 
-#if defined(ENABLE_NVIDIA_API)
+#if defined(ENABLE_NVIDIA_API) || defined(ENABLE_HYGON_API)
 #include "nvidia/recurrent_gated_delta_rule_nvidia.cuh"
 #endif
 
@@ -37,6 +37,9 @@ __INFINI_C infiniStatus_t infiniopCreateRecurrentGatedDeltaRuleDescriptor(
 #ifdef ENABLE_NVIDIA_API
         CREATE(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        CREATE(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -57,6 +60,9 @@ __INFINI_C infiniStatus_t infiniopGetRecurrentGatedDeltaRuleWorkspaceSize(
 #ifdef ENABLE_NVIDIA_API
         GET(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        GET(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -85,6 +91,9 @@ __INFINI_C infiniStatus_t infiniopRecurrentGatedDeltaRule(
 #ifdef ENABLE_NVIDIA_API
         CALCULATE(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        CALCULATE(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -104,6 +113,9 @@ __INFINI_C infiniStatus_t infiniopDestroyRecurrentGatedDeltaRuleDescriptor(
 #ifdef ENABLE_NVIDIA_API
         DESTROY(INFINI_DEVICE_NVIDIA, nvidia)
 #endif
+#ifdef ENABLE_HYGON_API
+        DESTROY(INFINI_DEVICE_HYGON, nvidia)
+#endif
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;