issue/1081 - fix hygon nv_bf16

Author: wooway777

src/infiniop/ops/addcmul/cuda/kernel.cuh

@@ -22,7 +22,7 @@ public:
             float f_t2 = __half2float(t2);
             return __float2half(f_input + v * f_t1 * f_t2);
 
-        } else if constexpr (std::is_same_v<T, nv_bfloat16>) {
+        } else if constexpr (std::is_same_v<T, cuda_bfloat16>) {
             float f_input = __bfloat162float(input);
             float f_t1 = __bfloat162float(t1);
             float f_t2 = __bfloat162float(t2);

src/infiniop/ops/addcmul/nvidia/addcmul_nvidia.cu

@@ -166,7 +166,7 @@ infiniStatus_t Descriptor::calculate(
     case INFINI_DTYPE_F16:
         return launch_addcmul_kernel<half>(this, output, inputs, stream);
     case INFINI_DTYPE_BF16:
-        return launch_addcmul_kernel<nv_bfloat16>(this, output, inputs, stream);
+        return launch_addcmul_kernel<cuda_bfloat16>(this, output, inputs, stream);
     case INFINI_DTYPE_F32:
         return launch_addcmul_kernel<float>(this, output, inputs, stream);
     case INFINI_DTYPE_F64:

src/infiniop/ops/atanh/cuda/kernel.cuh

@@ -21,7 +21,7 @@ public:
         } else if constexpr (std::is_same_v<T, half>) {
             // half 类型先转为 float 计算再转回
             return __float2half(atanhf(__half2float(a)));
-        } else if constexpr (std::is_same_v<T, nv_bfloat16>) {
+        } else if constexpr (std::is_same_v<T, cuda_bfloat16>) {
             // bfloat16 类型处理同上
             return __float2bfloat16(atanhf(__bfloat162float(a)));
         } else if constexpr (std::is_same_v<T, float>) {

src/infiniop/ops/atanh/nvidia/atanh_nvidia.cu

@@ -44,7 +44,7 @@ infiniStatus_t Descriptor::calculate(
     case INFINI_DTYPE_F16:
         return _device_info->calculate<256, cuda::AtanhOp, half>(_info, workspace, output, inputs, stream);
     case INFINI_DTYPE_BF16:
-        return _device_info->calculate<256, cuda::AtanhOp, nv_bfloat16>(_info, workspace, output, inputs, stream);
+        return _device_info->calculate<256, cuda::AtanhOp, cuda_bfloat16>(_info, workspace, output, inputs, stream);
     case INFINI_DTYPE_F32:
         return _device_info->calculate<256, cuda::AtanhOp, float>(_info, workspace, output, inputs, stream);
     case INFINI_DTYPE_F64:

src/infiniop/ops/avg_pool1d/cuda/kernel.cuh

@@ -46,7 +46,7 @@ __device__ void avgPool1dKernel(
             }
         }
 
-#if defined(ENABLE_ILUVATAR_API)
+#if defined(ENABLE_ILUVATAR_API) || defined(ENABLE_HYGON_API)
         // Iluvatar __half doesn't accept size_t directly.
         y[y_offset] = sum / static_cast<T>(static_cast<double>(kernel_size));
 #else

src/infiniop/ops/var_mean/nvidia/var_mean_nvidia.cu

@@ -44,7 +44,6 @@ infiniStatus_t launchKernel(
     bool unbiased, bool keepdim,
     cudaStream_t stream, void *workspace, size_t workspace_size) {
     size_t input_ndim = info.permuted_input_shape.size();
-    size_t output_ndim = info.output_shape.size();
     size_t input_size = info.input_size;
     size_t output_size = info.output_size;
     size_t reduce_num = info.reduce_num;