/* ******************************************************************************
*
*
* This program and the accompanying materials are made available under the
* terms of the Apache License, Version 2.0 which is available at
* https://www.apache.org/licenses/LICENSE-2.0.
*
*  See the NOTICE file distributed with this work for additional
*  information regarding copyright ownership.
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS, WITHOUT
* WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the
* License for the specific language governing permissions and limitations
* under the License.
*
* SPDX-License-Identifier: Apache-2.0
******************************************************************************/

#pragma once
#ifndef OPS_H_
#define OPS_H_

#include <array/DataTypeUtils.h>
#include <helpers/shape.h>
#include <loops/ReduceType.h>
#include <loops/summarystatsreduce.h>
#include <math/templatemath.h>
#include <system/Environment.h>
#include <system/common.h>
#include <system/op_boilerplate.h>

#include <codecvt>
#include <vector>

#include "helpers/unicode.h"
#include "op_macros_meta.h"

// =============================================================================
// CONSTANTS
// =============================================================================

#define SELU_ALPHA 1.6732632423543772848170429916717
#define SELU_LAMBDA 1.0507009873554804934193349852946
#define SD_STRING_ASSIGN_TEMP_BUFFER_BYTES 256




namespace simdOps {


// =============================================================================
// BINARY ARITHMETIC OPERATIONS
// =============================================================================

DECLARE_BINARY_MATH_OP(Add, sd_add)
DECLARE_BINARY_MATH_OP(Subtract, sd_subtract)
DECLARE_BINARY_MATH_OP(Multiply, sd_multiply)
DECLARE_BINARY_MATH_OP(Divide, sd_divide)


// Reverse operations
DECLARE_REVERSE_BINARY_MATH_OP(ReverseSubtract, sd_subtract, 0.f)
DECLARE_REVERSE_BINARY_MATH_OP(ReverseDivide, sd_divide, 1)

// HardTanh - uses complex conditional macro
DECLARE_UNARY_COMPLEX_CONDITIONAL_OP(HardTanh,
                                     d1 < static_cast<X>(-1), static_cast<X>(-1),
                                     d1 > static_cast<X>(1), static_cast<X>(1),
                                     d1)

DECLARE_UNARY_CONDITIONAL_OP(RectifiedTanhDerivative,
                             d1 > static_cast<X>(0.f),
                             sd::math::sd_tanhderivative<X COMMA X>(d1),
                             static_cast<X>(0.f))



DECLARE_BINARY_MATH_OP_XZ(Atan2, sd_atan2)


DECLARE_BINARY_MATH_OP_WITH_STARTING(PowDerivative,
                                     static_cast<Z>(d2) * sd::math::sd_pow<X COMMA Z COMMA Z>(d1 COMMA static_cast<Z>(d2) - static_cast<Z>(1.f)),
                                     static_cast<Z>(d1),
                                     params[0] * sd::math::sd_pow<X COMMA Z COMMA Z>(d1 COMMA static_cast<Z>(params[0]) - static_cast<Z>(1.f)),
                                     static_cast<X>(0)
)

DECLARE_BINARY_COPY_OP(AMaxPairwise,
                       sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d1)) > sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d2)) ? static_cast<Z>(d1) : static_cast<Z>(d2),
                       sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d1)) > sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d2)) ? static_cast<Z>(d1) : static_cast<Z>(d2),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)

DECLARE_BINARY_COPY_OP(AMinPairwise,
                       sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d1)) < sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d2)) ? static_cast<Z>(d1) : static_cast<Z>(d2),
                       sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d1)) < sd::math::sd_abs<Z COMMA Z>(static_cast<Z>(d2)) ? static_cast<Z>(d1) : static_cast<Z>(d2),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)

DECLARE_BINARY_COPY_OP(MaxPairwise,
                       sd::math::sd_max<Z>(static_cast<Z>(d1) COMMA static_cast<Z>(d2)),
                       sd::math::sd_max<Z>(static_cast<Z>(d1) COMMA static_cast<Z>(d2)),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)

DECLARE_BINARY_COPY_OP(MinPairwise,
                       sd::math::sd_min<X COMMA Y COMMA Z>(d1 COMMA d2),
                       sd::math::sd_min<X COMMA Y COMMA Z>(d1 COMMA d2),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)

DECLARE_REDUCE3_OP_WITH_BOOL_SUPPORT(CosineSimilarity,
// Bool logic (for boolean input types)
                                     extraParamsRef[0] += static_cast<Y>(static_cast<int>(d1) * static_cast<int>(d1)); \
                                     extraParamsRef[1] += static_cast<Y>(static_cast<int>(d2) * static_cast<int>(d2)); \
                                     return static_cast<Y>(static_cast<int>(d1) * static_cast<int>(d2));,

// Normal logic (for non-boolean types)
                                     extraParamsRef[0] += static_cast<Y>(d1 * d1); \
                                     extraParamsRef[1] += static_cast<Y>(d2 * d2); \
                                     return static_cast<Y>(d1 * d2);,

                                     2,  // extraParamsLen
                                     0.0f,  // starting value

// Post process - calculate cosine similarity from dot product and norms
                                     reduction / (sd::math::sd_sqrt<Y COMMA Y>(extraParamsRef[0]) * sd::math::sd_sqrt<Y COMMA Y>(extraParamsRef[1]))
)



DECLARE_BINARY_MATH_OP(IGamma, sd_igamma)


DECLARE_BINARY_MATH_OP(IGammac, sd_igammac)



DECLARE_BINARY_COPY_OP(LogX,
                       sd::math::sd_log<X COMMA Z>(d1) / sd::math::sd_log<Y COMMA Z>(d2),
                       sd::math::sd_log<X COMMA Z>(d1) / sd::math::sd_log<Y COMMA Z>(d2),
                       static_cast<Z>(d1),
                       sd::math::sd_log<X COMMA Z>(d1) / sd::math::sd_log<Y COMMA Z>(params[0])
)




// ASinhDerivative - uses complex math expression macro
DECLARE_UNARY_COMPLEX_MATH_OP(ASinhDerivative,
                              static_cast<X>(1.f) / (sd::math::sd_sqrt<X, X>(sd::math::sd_pow<X, X, X>(d1, static_cast<X>(2.f)) + static_cast<X>(1.f))))

// ACoshDerivative - uses complex math expression macro
DECLARE_UNARY_COMPLEX_MATH_OP(ACoshDerivative,
                              static_cast<X>(1.f) / (sd::math::sd_sqrt<X, X>(d1 - static_cast<X>(1.f)) * sd::math::sd_sqrt<X, X>(d1 + static_cast<X>(1.f))))

// Power operations


DECLARE_POWER_OP(Pow, sd_pow)
DECLARE_REVERSE_BINARY_MATH_OP(ReversePow, sd_pow, 1)


DECLARE_UNARY_SIMPLE_OP(TanDerivative, static_cast<X>(1.f) / sd::math::sd_pow<X COMMA X COMMA X>(sd::math::sd_cos<X COMMA X>(d1) COMMA static_cast<X>(2.0f)))




DECLARE_SQUARED_SUBTRACT_OP(SquaredSubtract, sd_subtract)
DECLARE_SQUARED_REVERSE_SUBTRACT_OP(SquaredReverseSubtract, sd_subtract)

// =============================================================================
// COMPARISON OPERATIONS
// =============================================================================

DECLARE_COMPARISON_OP(EqualTo, ==)
DECLARE_COMPARISON_OP(NotEqualTo, !=)
DECLARE_COMPARISON_OP(GreaterThan, >)
DECLARE_COMPARISON_OP(GreaterThanOrEqual, >=)
DECLARE_COMPARISON_OP(LessThan, <)
DECLARE_COMPARISON_OP(LessThanOrEqual, <=)

// =============================================================================
// UNARY MATH OPERATIONS
// =============================================================================

DECLARE_UNARY_MATH_OP(Abs, sd_abs)
DECLARE_UNARY_MATH_OP(Ceiling, sd_ceil)
DECLARE_UNARY_MATH_OP(Cosine, sd_cos)
DECLARE_UNARY_MATH_OP(Exp, sd_exp)
DECLARE_UNARY_MATH_OP(Floor, sd_floor)
DECLARE_UNARY_MATH_OP(Log, sd_log)
DECLARE_UNARY_MATH_OP(Sin, sd_sin)
DECLARE_UNARY_MATH_OP(Tanh, sd_tanh)
DECLARE_UNARY_MATH_OP(Sigmoid, sd_sigmoid)

DECLARE_UNARY_SIMPLE_OP(Neg, -d1)
DECLARE_UNARY_SIMPLE_OP(Square, d1 * d1)
DECLARE_UNARY_SIMPLE_OP(Cube, d1 * d1 * d1)
DECLARE_UNARY_SIMPLE_OP(Identity, d1)
DECLARE_UNARY_SIMPLE_OP(OneMinus, static_cast<X>(1) - d1)
DECLARE_UNARY_SIMPLE_OP(Reciprocal, static_cast<X>(1) / d1)




// =============================================================================
// CONDITIONAL OPERATIONS
// =============================================================================


DECLARE_UNARY_CONDITIONAL_OP(Sign,
                             (d1 > static_cast<X>(0)) - (d1 < static_cast<X>(0)),
                             static_cast<X>(1), static_cast<X>(-1))


DECLARE_UNARY_CONDITIONAL_OP(HardTanhDerivative,
                             ((d1 >= static_cast<X>(-1.f) && d1 <= static_cast<X>(1.f)) ? static_cast<X>(1.f) : static_cast<X>(0.f)), d1, d1)

DECLARE_UNARY_CONDITIONAL_OP(HardSigmoidDerivative,
                             d1 < static_cast<X>(-2.5f) || d1 > static_cast<X>(2.5f) ? static_cast<X>(0.f) : static_cast<X>(0.2f), d1, d1)

DECLARE_BINARY_MATH_OP(Remainder, sd_remainder)
DECLARE_BINARY_MATH_OP(FMod, sd_fmod)


DECLARE_SAFE_DIVISION_OP(DivideNoNan, d2 == static_cast<Y>(0))
DECLARE_SAFE_DIVISION_OP(SafeDivide, d2 == static_cast<Y>(0))

// Floor division:
DECLARE_FLOOR_DIVISION_OP(FloorDiv, sd_floor)



DECLARE_BINARY_MATH_OP_WITH_STARTING(TruncateDiv,
                                     static_cast<Z>(sd::math::sd_divide<int,int,int>(static_cast<int>(d1), static_cast<int>(d2))),
                                     static_cast<Z>(d1),
                                     static_cast<Z>(sd::math::sd_divide<int,int,int>(static_cast<int>(d1), static_cast<int>(params[0]))),
                                     static_cast<X>(1)
)

DECLARE_BINARY_MATH_OP_WITH_STARTING(TruncateMod,
                                     static_cast<Z>(static_cast<int>(d1) % static_cast<int>(d2)),
                                     static_cast<Z>(d1),
                                     static_cast<Z>(static_cast<int>(d1) % static_cast<int>(params[0])),
                                     static_cast<X>(0)
)




DECLARE_UNARY_IDENTITY_OP(Copy)

template <typename X, typename Y, typename Z>
class FloorMod {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y d2) {
    Z m = sd::math::sd_fmod<X,Y,Z>(d1, d2);
    return (d1 < static_cast<X>(0)) == (d2 < static_cast<Y>(0))
           ? m
           : sd::math::sd_fmod<Z,Y,Z>(m + static_cast<Z>(d2), d2);
  }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y d2, Z *params) { return op_logic(d1, d2); }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1) { return static_cast<Z>(d1); }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y *params) {
    Z m = sd::math::sd_fmod<X,Y,Z>(d1, params[0]);
    return (d1 < static_cast<X>(0)) == (params[0] < static_cast<Y>(0))
           ? m
           : sd::math::sd_fmod<Z,Y,Z>(m + static_cast<Z>(params[0]), params[0]);
  }
  static SD_HOST_DEVICE SD_INLINE Z op_simd(X d1, Y d2) { return op_logic(d1, d2); }
  static SD_HOST_DEVICE SD_INLINE Z op_simd(X d1, Y d2, Z *params) { return op_logic(d1, d2, params); }
  static SD_HOST_DEVICE SD_INLINE Z op_simd(X d1) { return op_logic(d1); }
  static SD_HOST_DEVICE SD_INLINE Z op_simd(X d1, Y *params) { return op_logic(d1, params); }

 public:
  static SD_HOST_DEVICE SD_INLINE Z op(X d1, Y d2) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, d2);
    else return op_simd(d1, d2);
  }
  static SD_HOST_DEVICE SD_INLINE Z op(X d1, Y d2, Z *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, d2, params);
    else return op_simd(d1, d2, params);
  }
  static SD_HOST_DEVICE SD_INLINE Z op(X d1) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1);
    else return op_simd(d1);
  }
  static SD_HOST_DEVICE SD_INLINE Z op(X d1, Y *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, params);
    else return op_simd(d1, params);
  }
  SD_HOST_DEVICE SD_INLINE static X startingValue() { return static_cast<X>(0); }
};

DECLARE_BINARY_COPY_OP(CopyPws,
                       static_cast<Z>(d2),
                       static_cast<Z>(d2),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)

DECLARE_BINARY_COPY_OP(Copy2,
                       static_cast<Z>(d2),
                       static_cast<Z>(d2),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)



DECLARE_BINARY_COPY_OP(Axpy,
                       static_cast<Z>(d2 + d1),
                       params[0] * static_cast<Z>(d1) + static_cast<Z>(d2),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)



DECLARE_BINARY_PARAM_OP(LstmClip,
                        [&]() {
                          X _v = static_cast<X>(d2);
                          if (d1 > _v) return static_cast<Z>(_v);
                          else if (d1 < -_v) return static_cast<Z>(-_v);
                          else return static_cast<Z>(d1);
                        }(),
                        no_op_exec_special no_op_exec_special_cuda
)

DECLARE_BINARY_PARAM_OP(Step,
                        (d1 > static_cast<X>(d2) ? static_cast<Z>(1) : static_cast<Z>(0)),
                        no_op_exec_special_same no_op_exec_special_same_cuda
)

DECLARE_BINARY_PARAM_OP(SXELogitsSmoother,
                        static_cast<Z>(d1 * (static_cast<X>(1.f) - static_cast<X>(d2)) + static_cast<X>(0.5f) * static_cast<X>(d2)),
// no special boilerplate needed
)



DECLARE_UNARY_MATH_OP(Round, sd_round)
DECLARE_UNARY_MATH_OP(Rint, sd_rint)
DECLARE_UNARY_MATH_OP(Erf, sd_erf)
DECLARE_UNARY_MATH_OP(Erfc, sd_erfc)
DECLARE_UNARY_MATH_OP(ASin, sd_asin)
DECLARE_UNARY_MATH_OP(ACos, sd_acos)
DECLARE_UNARY_MATH_OP(ATan, sd_atan)
DECLARE_UNARY_MATH_OP(ATanh, sd_atanh)
DECLARE_UNARY_MATH_OP(ASinh, sd_asinh)
DECLARE_UNARY_MATH_OP(ACosh, sd_acosh)
DECLARE_UNARY_MATH_OP(Sinh, sd_sinh)
DECLARE_UNARY_MATH_OP(Cosh, sd_cosh)
DECLARE_UNARY_MATH_OP(Tan, sd_tan)
DECLARE_UNARY_MATH_OP(SoftSign, sd_softsign)

// Log1p - simple math expression
DECLARE_UNARY_SIMPLE_OP(Log1p, sd::math::sd_log<X COMMA X>(1 + d1))

// Expm1 - simple math expression
DECLARE_UNARY_SIMPLE_OP(Expm1, sd::math::sd_exp<X COMMA X>(d1) - static_cast<X>(1))

// StabilizeFP16 - conditional operation
DECLARE_UNARY_CONDITIONAL_OP(StabilizeFP16,
                             d1 <= static_cast<X>(0),
                             static_cast<X>(sd::DataTypeUtils::min<float16>()),
                             d1)
// StabilizeX - conditional operation
DECLARE_UNARY_CONDITIONAL_OP(StabilizeX,
                             d1 <= static_cast<X>(0),
                             sd::DataTypeUtils::min<X>(),
                             d1)

// SoftPlus - simple math function
DECLARE_UNARY_MATH_OP(SoftPlus, sd_softplus)

// SoftMax - complex math expression using params
DECLARE_UNARY_COMPLEX_MATH_OP(SoftMax,
                              sd::math::sd_exp<X COMMA X>(d1 - params[0]) / params[1])

// LogSoftMax - simple math expression using params
DECLARE_UNARY_SIMPLE_OP(LogSoftMax, (d1 - params[0]) - params[1])

// Sech - simple math expression (reciprocal of cosh)
DECLARE_UNARY_SIMPLE_OP(Sech, static_cast<X>(1) / sd::math::sd_cosh<X COMMA X>(d1))

// Csch - simple math expression (reciprocal of sinh)
DECLARE_UNARY_SIMPLE_OP(Csch, static_cast<X>(1) / sd::math::sd_sinh<X COMMA X>(d1))

// Coth - simple math expression (cosh/sinh)
DECLARE_UNARY_SIMPLE_OP(Coth, sd::math::sd_cosh<X COMMA X>(d1) / sd::math::sd_sinh<X COMMA X>(d1))



DECLARE_UNARY_COMPLEX_CONDITIONAL_OP(ClipByValue,
                                     d1 > params[1], params[1],
                                     d1 < params[0], params[0],
                                     d1)
template <typename X>
class LGamma {
 public:
  no_op_exec_special_same no_op_exec_special_same_cuda
      SD_HOST_DEVICE SD_INLINE static X op(X d1, X *params) {
    return sd::math::sd_lgamma<X, X>(d1);
  }
};



template <typename X>
class SetRange {
 private:
  static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X *params) {
    auto min = params[0];
    auto max = params[1];
    if (static_cast<X>(d1) >= min && static_cast<X>(d1) <= max) return d1;
    if (min == static_cast<X>(0) && max == static_cast<X>(1)) {
      auto val = static_cast<X>(1) / (static_cast<X>(1) + sd::math::sd_exp<X, X>(-d1));
      return (sd::math::sd_floor<X, X>(val * (max - min)) + min);
    }
    return (sd::math::sd_floor<X, X>(d1 * (max - min)) + min);
  }
  static SD_HOST_DEVICE SD_INLINE X op_simd(X d1, X *params) { return op_logic(d1, params); }

 public:
  no_op_exec_special_same no_op_exec_special_same_cuda;
  static SD_HOST_DEVICE SD_INLINE X op(X d1, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)
      return op_logic(d1, params);
    else
      return op_simd(d1, params);
  }
};


DECLARE_UNARY_SIMPLE_OP(Affine, params[0] * d1 + params[1])




template <typename X>
class Stabilize {
 private:
  static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X *params) {
    X k = params[0];
    if (d1 * k > static_cast<X>(SD_MAX_CUTFOFF))
      return static_cast<X>(SD_MAX_CUTFOFF) / k;
    else if (d1 * k < static_cast<X>(SD_MIN_CUTFOFF))
      return static_cast<X>(SD_MIN_CUTFOFF) / k;
    return d1;
  }
  static X op_simd(X d1, X *params) { return op_logic(d1, params); }

 public:
  no_op_exec_special_same no_op_exec_special_same_cuda;
  static SD_HOST_DEVICE SD_INLINE X op(X d1, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)
      return op_logic(d1, params);
    else
      return op_simd(d1, params);
  }
};

DECLARE_UNARY_SIMPLE_OP(Ones, static_cast<X>(1.0f))



DECLARE_REDUCE3_OP_WITH_BOOL_SUPPORT(JaccardDistance,
                                     Y num_val = static_cast<Y>(static_cast<int>(d1) & static_cast<int>(d2)); \
                                     Y denom_val = static_cast<Y>(static_cast<int>(d1) | static_cast<int>(d2)); \
                                     extraParamsRef[0] += num_val; \
                                     extraParamsRef[1] += denom_val; \
                                     return static_cast<Y>(0.0f);,
                                     Y num_val = static_cast<Y>(sd::math::sd_min<X>(d1 COMMA d2)); \
                                     Y denom_val = static_cast<Y>(sd::math::sd_max<X>(d1 COMMA d2)); \
                                     extraParamsRef[0] += num_val; \
                                     extraParamsRef[1] += denom_val; \
                                     return static_cast<Y>(0.0f);,
                                     2,
                                     0.0f,
                                     (static_cast<Y>(1.0f)) - (extraParamsRef[0] / extraParamsRef[1])
)


DECLARE_HAMMING_DISTANCE_OP_WITH_BOOL_SUPPORT(SimpleHammingDistance,
                                              (static_cast<int>(d1) == static_cast<int>(d2)) ? 0.0f : 1.0f,
                                              (d1 == d2) ? 0.0f : 1.0f,
                                              0.0f
)


DECLARE_REDUCE3_OP_WITH_BOOL_SUPPORT(CosineDistance,
                                     extraParamsRef[0] += static_cast<Y>(static_cast<int>(d1) * static_cast<int>(d1)); \
                                     extraParamsRef[1] += static_cast<Y>(static_cast<int>(d2) * static_cast<int>(d2)); \
                                     return static_cast<Y>(static_cast<int>(d1) * static_cast<int>(d2));,
                                     extraParamsRef[0] += static_cast<Y>(sd::math::sd_abs<X COMMA X>(d1) * sd::math::sd_abs<X COMMA X>(d1)); \
                                     extraParamsRef[1] += static_cast<Y>(sd::math::sd_abs<X COMMA X>(d2) * sd::math::sd_abs<X COMMA X>(d2)); \
                                     return static_cast<Y>(d1 * d2);,
                                     2,
                                     0.0f,
                                     (static_cast<Y>(1.0f)) - (reduction / (sd::math::sd_sqrt<Y COMMA Y>(extraParamsRef[0]) * sd::math::sd_sqrt<Y COMMA Y>(extraParamsRef[1])))
)

DECLARE_DISTANCE_OP_WITH_BOOL_SUPPORT(Dot,
                                      static_cast<int>(d1) * static_cast<int>(d2),
                                      d1 * d2,
                                      0.0f
)
DECLARE_BOOLEAN_OP_WITH_TYPE_SAFETY(EqualsWithEps,
                                    sd::math::sd_eq<X COMMA X>(d1 COMMA d2 COMMA eps),
                                    1.0f
)

DECLARE_DISTANCE_OP_WITH_BOOL_SUPPORT(EuclideanDistance,
                                      static_cast<int>(d1) != static_cast<int>(d2) ? 1 : 0,
                                      (d1 - d2) * (d1 - d2),
                                      0.0f
)

DECLARE_DISTANCE_OP_WITH_BOOL_SUPPORT(ManhattanDistance,
                                      static_cast<int>(d1) != static_cast<int>(d2) ? 1 : 0,
                                      sd::math::sd_abs<X COMMA X>(d1 - d2),
                                      0.0f
)

template <typename X>
class DropOut {
 public:
  no_op_exec_special_same no_op_exec_special_same_cuda
      SD_HOST_DEVICE SD_INLINE static X op(X d1, X *params) {
    X prob = params[0];
#ifdef __CUDACC__
    X length = params[1];
   X tid = blockIdx.x * blockDim.x + threadIdx.x;
   X rnd = sd::math::sd_abs<X,X>(sd::math::sd_cos<X>( static_cast<X>(tid) +
                                                      static_cast<X>(length) * static_cast<X>(tid)));
#else
    X rnd = static_cast<X>(rand() / RAND_MAX);
#endif
    return rnd >= prob ? static_cast<X>(0.0f) : d1;
  }
};

template <typename X, typename Y, typename Z>
class DropOutInverted {
 public:
  no_op_exec_special no_op_exec_special_cuda
#ifdef __CUDACC__
  SD_DEVICE
#endif
          SD_HOST_DEVICE SD_INLINE Z op(X d1, Y d2, Z *params) {
    Y prob = d2;
#ifdef __CUDACC__
    X length = params[1];
   X tid = blockIdx.x * blockDim.x + threadIdx.x;
   X rnd = sd::math::sd_abs<X,X>(sd::math::sd_cos<X>( static_cast<X>(tid) +
                                                      static_cast<X>(length) * static_cast<X>(tid)));
#else
    X rnd = static_cast<X>(rand() / RAND_MAX);
#endif
    return rnd >= static_cast<X>(prob) ? static_cast<Z>(0.0f) : reinterpret_cast<Z>(d1 / static_cast<X>(prob));
  }
};

DECLARE_BINARY_COPY_OP(ReplaceNans,
                       sd::math::sd_isnan(d1) ? static_cast<Z>(d2) : static_cast<Z>(d1),
                       sd::math::sd_isnan(d1) ? static_cast<Z>(d2) : static_cast<Z>(d1),
                       static_cast<Z>(d1),
                       static_cast<Z>(d1)
)




template <typename X, typename Y, typename Z>
class CompareAndReplace {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y d2, Z *params) {
    // Type guard for non-arithmetic types
    if constexpr (any_my_string_v<X, Y, Z>) {
      return static_cast<Z>(d1);  // For non-arithmetic types, just return cast of first value
    } else {
      // All the arithmetic logic is now inside the else block
      // This code is only compiled when none of X, Y, Z are string types
      auto zd1 = static_cast<Z>(d1);
      auto zd2 = static_cast<Z>(d2);
      auto compare = params[0];
      auto eps = params[2];
      int mode = (int)params[3];  // This line is now safe - only compiled for non-string types

      if (mode == 0)  // equals
        if (sd::math::sd_abs<Z,Z>(zd1 - compare) <= eps)
          return zd2;
        else
          return zd1;
      else if (mode == 1)  // not equals eps
        if (sd::math::sd_abs<Z,Z>(zd1 - compare) > eps)
          return zd2;
        else
          return zd1;
      else if (mode == 2)  // less_than eps
        if (zd1 < compare)
          return zd2;
        else
          return zd1;
      else if (mode == 3)  // greater_than
        if (zd1 > compare)
          return zd2;
        else
          return zd1;
      else if (mode == 4)  // less_or_equals_than
        if (zd1 <= compare)
          return zd2;
        else
          return zd1;
      else if (mode == 5)  // greater_or_equals_than
        if (zd1 >= compare)
          return zd2;
        else
          return zd1;
      else if (mode == 6)  // abs_less_than
        if (sd::math::sd_abs<Z,Z>(zd1) < compare)
          return zd2;
        else
          return zd1;
      else if (mode == 7)  // abs_greater_than
        if (sd::math::sd_abs<Z,Z>(zd1) > compare)
          return zd2;
        else
          return zd1;
      else if (mode == 8 || mode == 15)  // is inf
        if constexpr (std::is_arithmetic<X>::value) {
          if (sd::math::sd_isinf<X>(d1))  // Use original d1, not cast zd1
            return zd2;
          else
            return zd1;
        } else {
          return zd1;
        }
      else if (mode == 9)  // is nan
        if constexpr (std::is_arithmetic<X>::value) {
          if (sd::math::sd_isnan<X>(d1))  // Use original d1, not cast zd1
            return zd2;
          else
            return zd1;
        } else {
          return zd1;
        }
      else if (mode == 10)
        if (zd1 == compare)
          return zd2;
        else
          return zd1;
      else if (mode == 11)
        if (zd1 != compare)
          return zd2;
        else
          return zd1;
      else if (mode == 12)  // abs_greater_or_equals_than
        if (sd::math::sd_abs<Z,Z>(zd1) >= compare)
          return zd2;
        else
          return zd1;
      else if (mode == 13) {  // abs_less_or_equals_than
        if (sd::math::sd_abs<Z,Z>(zd1) <= compare)
          return zd2;
        else
          return zd1;
      }
      else if (mode == 14) {  // is_finite (not inf)
        if constexpr (std::is_arithmetic<X>::value) {
          if (!sd::math::sd_isinf<X>(d1))  // Use original d1, not cast zd1
            return zd2;
          else
            return zd1;
        } else {
          return zd1;
        }
      }
      else
        sd_printf("Undefined boolean operation: [%i]\n", mode);
      return zd1;
    }
  }
  
  static SD_HOST_DEVICE SD_INLINE Z op_simd(X d1, Y d2, Z *params) { 
    return op_logic(d1, d2, params); 
  }

 public:
  static SD_HOST_DEVICE SD_INLINE Z op(X d1, Y d2, Z *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, d2, params);
    else
      return op_simd(d1, d2, params);
  }
};

template <typename X, typename Y, typename Z>
class CompareAndSet {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X dX, Y dY, Z *params) {
    // Type guard for non-arithmetic types - early return
    if constexpr (any_my_string_v<X, Y, Z>) {
      return static_cast<Z>(dX);  // For non-arithmetic types, just return cast of first value
    } else {
      // All arithmetic logic inside else block - only compiled for non-string types
      auto d1 = static_cast<Z>(dX);
      auto d2 = static_cast<Z>(dY);
      auto compare = params[0];
      auto eps = params[2];
      auto mode = static_cast<int>(params[3]);  // Safe - only compiled for non-string types

      if (mode == 0)  // equals
        if (sd::math::sd_abs<Z,Z>(d2 - compare) <= eps)
          return d2;
        else
          return d1;
      else if (mode == 1)  // not equals
        if (sd::math::sd_abs<Z,Z>(d2 - compare) > eps)
          return d2;
        else
          return d1;
      else if (mode == 2)  // less_than
        if (d2 < compare)
          return d2;
        else
          return d1;
      else if (mode == 3)  // greater_than
        if (d2 > compare)
          return d2;
        else
          return d1;
      else if (mode == 4)  // less_or_equals_than
        if (d2 <= compare)
          return d2;
        else
          return d1;
      else if (mode == 5)  // greater_or_equals_than
        if (d2 >= compare)
          return d2;
        else
          return d1;
      else if (mode == 6)  // abs_less_than
        if (sd::math::sd_abs<Z,Z>(d2) < compare)
          return d2;
        else
          return d1;
      else if (mode == 7)  // abs_greater_than
        if (sd::math::sd_abs<Z,Z>(d2) > compare)
          return d2;
        else
          return d1;
      else if (mode == 8 || mode == 15) {  // is inf
        if constexpr (std::is_arithmetic<Y>::value) {
          if (sd::math::sd_isinf<Y>(dY))  // Use original dY, not cast d2
            return d2;
          else
            return d1;
        } else {
          return d1;
        }
      }
      else if (mode == 9) {  // is nan
        if constexpr (std::is_arithmetic<Y>::value) {
          if (sd::math::sd_isnan<Y>(dY))  // Use original dY, not cast d2
            return d2;
          else
            return d1;
        } else {
          return d1;
        }
      }
      else if (mode == 10)
        if (d2 == compare)
          return d2;
        else
          return d1;
      else if (mode == 11)
        if (d2 != compare)
          return d2;
        else
          return d1;
      else if (mode == 12)  // abs_greater_or_equals_than
        if (sd::math::sd_abs<Z,Z>(d1) >= compare)
          return d2;
        else
          return d1;
      else if (mode == 13)  // abs_less_or_equals_than
        if (sd::math::sd_abs<Z,Z>(d1) <= compare)
          return d2;
        else
          return d1;
      else if (mode == 14) {  // is_finite (not inf)
        if constexpr (std::is_arithmetic<X>::value) {
          if (!sd::math::sd_isinf<X>(dX))  // Use original dX, not cast d1
            return d2;
          else
            return d1;
        } else {
          return d1;
        }
      }
      else
        sd_printf("Undefined boolean operation: [%i]\n", mode);
      return d1;
    }
  }
  
  static SD_HOST_DEVICE SD_INLINE Z op_simd(X dX, Y dY, Z *params) { 
    return op_logic(dX, dY, params); 
  }

 public:
  static SD_HOST_DEVICE SD_INLINE Z op(X dX, Y dY, Z *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(dX, dY, params);
    else
      return op_simd(dX, dY, params);
  }
};

template <typename X>
class CompareAndSetTransform {
 private:
  static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X *params) {
    // Type guard for non-arithmetic types - early return
    if constexpr (!std::is_arithmetic<X>::value) {
      return d1;  // For non-arithmetic types, just return the input value
    } else {
      // All arithmetic logic inside else block - only compiled for arithmetic types
      auto compare = params[0];
      auto set = params[1];
      auto eps = params[2];
      int mode = (int)params[3];  // Safe - only compiled for arithmetic types

      if (mode == 0)  // equals
        if (sd::math::sd_abs<X,X>(d1 - compare) <= eps)
          return set;
        else
          return d1;
      else if (mode == 1)  // not equals
        if (sd::math::sd_abs<X,X>(d1 - compare) > eps)
          return set;
        else
          return d1;
      else if (mode == 2)  // less_than
        if (d1 < compare)
          return set;
        else
          return d1;
      else if (mode == 3)  // greater_than
        if (d1 > compare)
          return set;
        else
          return d1;
      else if (mode == 4)  // less_or_equals_than
        if (d1 <= compare)
          return set;
        else
          return d1;
      else if (mode == 5)  // greater_or_equals_than
        if (d1 >= compare)
          return set;
        else
          return d1;
      else if (mode == 6)  // abs_less_than
        if (sd::math::sd_abs<X,X>(d1) < compare)
          return set;
        else
          return d1;
      else if (mode == 7)  // abs_greater_than
        if (sd::math::sd_abs<X,X>(d1) > compare)
          return set;
        else
          return d1;
      else if (mode == 8)  // is inf
        if (sd::math::sd_isinf<X>(d1))
          return set;
        else
          return d1;
      else if (mode == 9)  // is nan
        if (sd::math::sd_isnan<X>(d1))
          return set;
        else
          return d1;
      else if (mode == 10)
        if (d1 == compare)
          return set;
        else
          return d1;
      else if (mode == 11)
        if (d1 != compare)
          return set;
        else
          return d1;
      else if (mode == 12)  // abs_greater_or_equals_than
        if (sd::math::sd_abs<X,X>(d1) >= compare)
          return set;
        else
          return d1;
      else if (mode == 13)  // abs_less_or_equals_than
        if (sd::math::sd_abs<X,X>(d1) <= compare)
          return set;
        else
          return d1;
      else if (mode == 14) {  // is_finite (not inf)
        if (!sd::math::sd_isinf<X>(d1))
          return compare;  // Note: original code returns compare, not set
        else
          return d1;
      }
      else
        sd_printf("Undefined boolean operation: [%i]\n", mode);
      return d1;
    }
  }
  
  static SD_HOST_DEVICE SD_INLINE X op_simd(X d1, X *params) { 
    return op_logic(d1, params); 
  }

 public:
  no_op_exec_special_same no_op_exec_special_same_cuda;
  
  static SD_HOST_DEVICE SD_INLINE X op(X d1, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)
      return op_logic(d1, params);
    else
      return op_simd(d1, params);
  }
};












DECLARE_UNARY_SIMD_SAFE_OP(SELUDerivative,
                           return d1 > static_cast<X>(0.f)
                                  ? static_cast<X>(SELU_LAMBDA)
                                  : static_cast<X>(SELU_ALPHA) * static_cast<X>(SELU_LAMBDA) * sd::math::sd_exp<X COMMA X>(d1);
)


DECLARE_UNARY_SIMD_SAFE_OP(HardSigmoid,
                           return sd::math::sd_min<X>(
                               static_cast<X>(1) COMMA sd::math::sd_max<X>(static_cast<X>(0) COMMA (static_cast<X>(0.2f)) * d1 + static_cast<X>(0.5f)));
)

DECLARE_UNARY_SIMD_SAFE_OP(SELU,
                           return d1 > static_cast<X>(0.0f)
                                  ? static_cast<X>(SELU_LAMBDA) * static_cast<X>(d1)
                                  : static_cast<X>(SELU_LAMBDA) * (static_cast<X>(SELU_ALPHA) * sd::math::sd_exp<X COMMA X>(d1) - static_cast<X>(SELU_ALPHA));
)





DECLARE_UNARY_SIMD_SAFE_OP(Swish,
                           return d1 * sd::math::sd_sigmoid<X COMMA X>(d1);
)


DECLARE_UNARY_SIMD_SAFE_OP(SwishDerivative,
                           X ex = sd::math::sd_pow<X COMMA X COMMA X>(static_cast<X>(M_E) COMMA d1);
                               return (ex * (d1 + ex + static_cast<X>(1.f))) / sd::math::sd_pow<X COMMA X COMMA X>((ex + static_cast<X>(1.f)) COMMA static_cast<X>(2.f));
)


DECLARE_UNARY_SIMD_SAFE_OP(Mish,
                           return d1 * sd::math::sd_tanh<X COMMA X>(sd::math::sd_softplus<X COMMA X>(d1));
)

DECLARE_UNARY_SIMD_SAFE_OP(MishDerivative,
                           auto ex = sd::math::sd_exp<X COMMA X>(d1);
                               auto e2x = ex * ex;
                               auto e3x = ex * ex * ex;
                               return (ex * (4 * (d1 + 1) + 4 * e2x + e3x + ex * (4 * d1 + 6))) / sd::math::sd_pow<X COMMA X COMMA X>((2 * ex + e2x + 2) COMMA (X)2.f);
)

DECLARE_UNARY_SIMD_SAFE_OP(GELU,
                           return d1 * sd::math::sd_sigmoid<X COMMA X>(static_cast<X>(1.702f) * d1);
)


DECLARE_UNARY_SIMD_SAFE_OP(PreciseGELU,
                           auto sp = sd::math::sd_sqrt<X COMMA X>(static_cast<X>(2) / static_cast<X>(M_PI));
                               auto xp = d1 + sd::math::sd_pow<X COMMA X COMMA X>(static_cast<X>(0.044715) * d1 COMMA static_cast<X>(3));
                               return (d1 / static_cast<X>(2)) * (static_cast<X>(1) + sd::math::sd_tanh<X COMMA X>(sp * xp));
)
DECLARE_UNARY_SIMD_SAFE_OP(GELUDerivative,
                           auto x17 = static_cast<X>(1.702f) * d1;
                               auto ep = sd::math::sd_exp<X COMMA X>(x17);
                               auto one_plus_ep = static_cast<X>(1.f) + ep;
                               return (ep * (one_plus_ep + x17)) / (one_plus_ep * one_plus_ep);

)
DECLARE_UNARY_SIMD_SAFE_OP(PreciseGELUDerivative,
                           auto x79 = static_cast<X>(0.797885) * d1;
                               auto temp1 = static_cast<X>(0.0356774) * d1;
                               auto x03 = temp1 * temp1 * temp1;  // cube without sd_pow
                               auto x39 = static_cast<X>(0.398942) * d1;
                               auto temp2 = static_cast<X>(0.0535161) * d1;
                               auto x05 = temp2 * temp2 * temp2;  // cube without sd_pow
                               auto scz = sd::math::sd_sech<X COMMA X>(x79 + x03);
                               return static_cast<X>(0.5) + (x39 + x05) * (scz * scz) + static_cast<X>(0.5) * sd::math::sd_tanh<X COMMA X>(x79 + x03);
)

DECLARE_UNARY_SIMD_SAFE_OP(LogSigmoid,
                           return sd::math::sd_log<X COMMA X>(sd::math::sd_sigmoid<X COMMA X>(d1));
)

DECLARE_UNARY_SIMD_SAFE_OP(LogSigmoidDerivative,
                           X ex = sd::math::sd_exp<X COMMA X>(d1);
                               return static_cast<X>(1.f) / (ex + static_cast<X>(1.f));
)
DECLARE_UNARY_MATH_OP(SigmoidDerivative, sd_sigmoidderivative)
DECLARE_UNARY_MATH_OP(TanhDerivative, sd_tanhderivative)
DECLARE_UNARY_MATH_OP(SinhDerivative, sd_cosh)  // sinh derivative is cosh
DECLARE_UNARY_MATH_OP(SoftSignDerivative, sd_softsignderivative)

DECLARE_MULTI_OP_SIMD_SAFE(And,
                           return d2 + d1;,
                           if (params != nullptr) {
                             auto comp = params[0];
                             return d1 != comp && d2 != comp ? static_cast<Z>(1) : static_cast<Z>(0);
                           } else {
                             auto b1 = static_cast<bool>(d1);
                             auto b2 = static_cast<bool>(d2);
                             return (b1 && b2) ? static_cast<Z>(1) : static_cast<Z>(0);
                           },
                           return d1;,
                           return static_cast<Z>(119);
)

DECLARE_UNARY_SIMD_SAFE_OP(RationalTanh,
                           auto dis = (static_cast<X>(2) / static_cast<X>(3)) * d1;
                               auto tanh = sd::math::sd_sgn<X COMMA X>(dis) *
                                           (static_cast<X>(1) -
                                            (static_cast<X>(1) / (static_cast<X>(1) + static_cast<X>(sd::math::sd_abs<X COMMA X>(dis)) +
                               sd::math::sd_pow<X COMMA X COMMA X>(dis COMMA static_cast<X>(2)) +
                               static_cast<X>(1.41645f) * sd::math::sd_pow<X COMMA X COMMA X>(dis COMMA static_cast<X>(4)))));
                               return static_cast<X>(1.7159f) * tanh;
)


DECLARE_UNARY_SIMD_SAFE_OP(RationalTanhDerivative,
                           auto dis = (static_cast<X>(2.f) / static_cast<X>(3.f)) * d1;
                               auto a = static_cast<X>(1.f) + sd::math::sd_abs<X COMMA X>(dis) + sd::math::sd_pow<X COMMA X COMMA X>(dis COMMA static_cast<X>(2.f)) +
                               static_cast<X>(1.41645f) * sd::math::sd_pow<X COMMA X COMMA X>(dis COMMA static_cast<X>(4));
                               auto tDeriv =
                           (static_cast<X>(1.f) + sd::math::sd_sign<X COMMA X>(dis) * (static_cast<X>(2.f) * dis +
                               static_cast<X>(4.f) * static_cast<X>(1.41645f) *
                               sd::math::sd_pow<X COMMA X COMMA X>(dis COMMA static_cast<X>(3)))) /
                           (a * a);
                               return static_cast<X>(1.7159f) * (static_cast<X>(2.f) / static_cast<X>(3.f)) * tDeriv;
)

DECLARE_UNARY_SIMD_SAFE_OP(ScaledTanh,
                           return params[0] * sd::math::sd_tanh<X COMMA X>(params[1] * d1);
)

// RectifiedTanh operation
DECLARE_UNARY_SIMD_SAFE_OP(RectifiedTanh,
                           return sd::math::sd_max<X>(static_cast<X>(0) COMMA sd::math::sd_tanh<X COMMA X>(d1));
)

// ELU operation
DECLARE_BINARY_SIMD_SAFE_OP(ELU,
                            return sd::math::sd_elu<X COMMA Z>(d1 COMMA static_cast<X>(d2));
)

// ELUDerivative operation
DECLARE_BINARY_SIMD_SAFE_OP(ELUDerivative,
                            return sd::math::sd_eluderivative<X COMMA Z>(d1 COMMA static_cast<X>(d2));
)

// RELU operation
DECLARE_BINARY_SIMD_SAFE_OP(RELU,
                            auto xt = static_cast<Z>(d1);
                                auto xf = static_cast<Z>(d2);
                                return xt < xf ? xf : xt;
)

// RELUDerivative operation
DECLARE_BINARY_SIMD_SAFE_OP(RELUDerivative,
                            auto xt = static_cast<Z>(d1);
                                auto xf = static_cast<Z>(d2);
                                return xt > xf ? static_cast<Z>(1.f) : static_cast<Z>(0.f);
)


DECLARE_BINARY_SIMD_SAFE_OP(RELU6,
                            auto relu = RELU<X COMMA Y COMMA Z>::op(d1 COMMA d2 COMMA params);
                                return relu < static_cast<Z>(6) ? relu : static_cast<Z>(6);
)

DECLARE_BINARY_SIMD_SAFE_OP(LeakyRELU,
                            auto val = static_cast<Z>(d1);
                                auto alpha = static_cast<Z>(d2);
                                return val < 0.0f ? alpha * val : val;
)

DECLARE_BINARY_SIMD_SAFE_OP(LeakyRELUDerivative,
                            if (d1 >= static_cast<X>(0))
                              return static_cast<Z>(1);
                            else
                              return static_cast<Z>(d2);
)

DECLARE_REDUCE_SIMD_SAFE_OP(IsNan,
                            return sd::math::sd_isnan<X>(d1) ? static_cast<Z>(1) : static_cast<Z>(0);
)
DECLARE_REDUCE_SIMD_SAFE_OP(IsPositive,
                            return d1 > (X)0.f;
)

DECLARE_REDUCE_SIMD_SAFE_OP(IsNegative,
                            return d1 < (X)0.f;
)

DECLARE_REDUCE_SIMD_SAFE_OP(IsInf,
                            return sd::math::sd_isinf<X>(d1) ? static_cast<Z>(1) : static_cast<Z>(0);
)

// IsInfOrNan operation
DECLARE_REDUCE_SIMD_SAFE_OP(IsFinite,
                            return sd::math::sd_isfin<X>(d1) ? static_cast<Z>(1) : static_cast<Z>(0);
)

DECLARE_REDUCE_SIMD_SAFE_OP(IsInfOrNan,
                            return sd::math::sd_isfin<X>(d1) ? static_cast<Z>(0) : static_cast<Z>(1);
)


DECLARE_UNARY_SIMPLE_OP(TimesOneMinus, d1 * (static_cast<X>(1) - d1))
DECLARE_UNARY_SIMPLE_OP(CubeDerivative, static_cast<X>(3) * d1 * d1)
DECLARE_UNARY_SIMPLE_OP(SpecialDerivative, d1 * (static_cast<X>(1.f) - d1))


DECLARE_COMPLEX_ACCUMULATION_SIMD_SAFE_OP(ASum,
                                          return sd::math::sd_abs<X COMMA X>(d1);,
                                          return sd::math::sd_abs<X COMMA X>(d1) + sd::math::sd_abs<X COMMA X>(d2);,
                                          return sd::math::sd_abs<X COMMA X>(d1) + sd::math::sd_abs<X COMMA X>(d2);,
                                          ASUM, static_cast<X>(0),
                                          sd::math::sd_abs<X COMMA X>(opOutput) + sd::math::sd_abs<X COMMA X>(old),
                                          sd::math::sd_abs<X COMMA X>(opOutput) + sd::math::sd_abs<X COMMA X>(old),
                                          sd::math::sd_abs<X COMMA X>(reduction)
)

DECLARE_SIMPLE_REDUCTION_OP(
    CountNonZero,
    ASUM,
    static_cast<Z>(0),
    (d1 == static_cast<X>(0.0f) ? static_cast<InterType>(0.0f) : static_cast<InterType>(1.0f)),
    (opOutput + old),
    (opOutput + old),
    static_cast<Z>(reduction)
)

DECLARE_SIMPLE_REDUCTION_OP(
    CountZero,
    SUM,
    static_cast<Z>(0.0f),
    (d1 == static_cast<X>(0) ? static_cast<InterType>(1) : static_cast<InterType>(0)),
    (opOutput + old),
    (opOutput + old),
    static_cast<Z>(reduction)
)

DECLARE_MIXED_ACCUMULATION_SIMD_SAFE_OP(Any,
                                        return d1;,
                                        SUM, static_cast<X>(0.0f),
                                        opOutput + old,
                                        opOutput + old,
                                        reduction > static_cast<Z>(0) ? static_cast<Z>(1) : static_cast<Z>(0)
)

DECLARE_MIXED_ACCUMULATION_SIMD_SAFE_OP(All,
                                        return d1;,
                                        SUM, static_cast<X>(1),
                                        static_cast<Z>(static_cast<bool>(opOutput) && static_cast<bool>(old) ? 1 : 0),
                                        static_cast<Z>(static_cast<bool>(opOutput) && static_cast<bool>(old) ? 1 : 0),
                                        reduction > static_cast<Z>(0) ? static_cast<Z>(1) : static_cast<Z>(0)
)

// AMean operation
DECLARE_ACCUMULATION_SIMD_SAFE_OP(AMean,
                                  return static_cast<InterType>(sd::math::sd_abs<X COMMA X>(d1));,
                                  ASUM, static_cast<X>(0),
                                  sd::math::sd_abs<InterType COMMA InterType>(opOutput) + sd::math::sd_abs<InterType COMMA InterType>(old),
                                  opOutput + old,
                                  static_cast<InterType>(reduction / static_cast<InterType>(n))
)

DECLARE_COMPLEX_ACCUMULATION_SIMD_SAFE_OP(AMax,
                                          return sd::math::sd_abs<X COMMA X>(d1);,
                                          return sd::math::sd_max<X>(sd::math::sd_abs<X COMMA X>(d1) COMMA sd::math::sd_abs<X COMMA X>(d2));,
                                          return sd::math::sd_abs<X COMMA X>(d1) > sd::math::sd_abs<X COMMA X>(d2) ? d1 : d2;,
                                          AMAX, input[0],
                                          sd::math::sd_max<X>(sd::math::sd_abs<X COMMA X>(old) COMMA sd::math::sd_abs<X COMMA X>(opOutput)),
                                          sd::math::sd_max<X>(sd::math::sd_abs<X COMMA X>(opOutput) COMMA sd::math::sd_abs<X COMMA X>(old)),
                                          sd::math::sd_abs<X COMMA X>(reduction)
)

DECLARE_COMPLEX_ACCUMULATION_SIMD_SAFE_OP(AMin,
                                          return sd::math::sd_abs<X COMMA X>(d1);,
                                          return sd::math::sd_min<X>(sd::math::sd_abs<X COMMA X>(d1) COMMA sd::math::sd_abs<X COMMA X>(d2));,
                                          return sd::math::sd_min<X>(sd::math::sd_abs<X COMMA X>(d1) COMMA sd::math::sd_abs<X COMMA X>(d2));,
                                          AMIN, input[0],
                                          sd::math::sd_min<X>(sd::math::sd_abs<X COMMA X>(old) COMMA sd::math::sd_abs<X COMMA X>(opOutput)),
                                          sd::math::sd_min<X>(sd::math::sd_abs<X COMMA X>(opOutput) COMMA sd::math::sd_abs<X COMMA X>(old)),
                                          sd::math::sd_abs<X COMMA X>(reduction)
)

DECLARE_ACCUMULATION_SIMD_SAFE_OP(Norm1,
                                  return static_cast<InterType>(sd::math::sd_abs<X COMMA X>(d1));,
                                  SUM, static_cast<X>(0),
                                  opOutput + old,
                                  opOutput + old,
                                  reduction
)

DECLARE_ACCUMULATION_SIMD_SAFE_OP(Norm2,
                                  auto v = static_cast<InterType>(d1);
                                  return v * v;,
                                               SUM, static_cast<X>(0),
                                               opOutput + old,
                                               opOutput + old,
                                               sd::math::sd_sqrt<InterType COMMA Z>(reduction)
)

DECLARE_ACCUMULATION_SIMD_SAFE_OP(SquaredNorm,
                                  auto v = static_cast<InterType>(d1);
                                  return v * v;,
                                               SUM, static_cast<X>(0),
                                               opOutput + old,
                                               opOutput + old,
                                               reduction
)

DECLARE_ACCUMULATION_SIMD_SAFE_OP(NormFrobenius,
                                  auto v = static_cast<InterType>(sd::math::sd_abs<X COMMA X>(d1));
                                  return v * v;,
                                               SUM, static_cast<X>(0),
                                               opOutput + old,
                                               opOutput + old,
                                               sd::math::sd_sqrt<InterType COMMA Z>(reduction)
)

DECLARE_MIXED_ACCUMULATION_SIMD_SAFE_OP(NormP,
                                  return sd::math::sd_pow<InterType COMMA Z COMMA InterType>(static_cast<InterType>(sd::math::sd_abs<X COMMA X>(d1)) COMMA static_cast<InterType>(params[0]));,
                                  SUM, static_cast<X>(0),
                                  opOutput + old,
                                  opOutput + old,
                                  sd::math::sd_pow<InterType COMMA Z COMMA Z>(reduction COMMA static_cast<Z>(1.0f) / extraParams[0])
)




DECLARE_MIXED_ACCUMULATION_SIMD_SAFE_OP(NormMax,
                                        return static_cast<Z>(d1);,
                                        SUM, static_cast<X>(0),
                                        opOutput + old,
                                        sd::math::sd_max<Z>(sd::math::sd_abs<Z COMMA Z>(old) COMMA sd::math::sd_abs<Z COMMA Z>(opOutput)),
                                        sd::math::sd_max<Z>(sd::math::sd_abs<Z COMMA Z>(reduction) COMMA sd::math::sd_abs<Z COMMA Z>(reduction))
)

// --- Generic Assign Template ---
template <typename X, typename Z>
class Assign {
 private:
  static SD_HOST_DEVICE SD_INLINE  Z op_logic(X d1, X *params) {
    if constexpr (std::is_same_v<X, Z>) {
      return d1; // No conversion needed
    } else if constexpr (std::is_convertible_v<X, Z>) {
      return static_cast<Z>(d1); // Use static_cast for direct convertibility
    } else {
      // This will trigger a compile error for unsupported types,
      // requiring a specialization like the ones below.
      return static_cast<Z>(d1);
    }
  }

  static SD_HOST_DEVICE SD_INLINE Z op_simd(X d1, X *params) { return op_logic(d1, params); }

 public:
  static SD_HOST_DEVICE SD_INLINE Z op(X d1, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value)
      return op_logic(d1, params);
    else
      return op_simd(d1, params);
  }
};

// --- Specialization: std::basic_string<char16_t> (UTF-16) -> std::basic_string<char> (UTF-8) ---
template <>
class Assign<std::basic_string<char16_t>, std::basic_string<char>> {
 public:
  static const bool requiresSpecial = false;
  static void execSpecial(const std::basic_string<char16_t> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char16_t> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char16_t> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char16_t> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char> *reductionPointer,
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  SD_HOST_DEVICE static std::basic_string<char>
  op(const std::basic_string<char16_t>& d1, std::basic_string<char16_t> * /*params*/) {
    char temp_output_buffer[SD_STRING_ASSIGN_TEMP_BUFFER_BYTES];
    const char16_t* input_data = d1.data();
    const uint32_t input_length_char16_units = static_cast<uint32_t>(d1.length());
    sd::LongType required_bytes = sd::unicode::offsetUtf16StringInUtf8(input_data, input_length_char16_units);

    if (required_bytes > 0 && static_cast<size_t>(required_bytes) <= SD_STRING_ASSIGN_TEMP_BUFFER_BYTES) {
      void* end_ptr = sd::unicode::utf16to8Ptr(input_data, input_data + input_length_char16_units, temp_output_buffer);
      size_t bytes_written = static_cast<char*>(end_ptr) - temp_output_buffer;
      if (bytes_written == static_cast<size_t>(required_bytes)) {
        return std::basic_string<char>(temp_output_buffer, bytes_written);
      }
    }
    return std::basic_string<char>();
  }
};

// --- Specialization: std::basic_string<char> (UTF-8) -> std::basic_string<char16_t> (UTF-16) ---
template <>
class Assign<std::basic_string<char>, std::basic_string<char16_t>> {
 public:
  static const bool requiresSpecial = false;
  static void execSpecial(const std::basic_string<char> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char16_t> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char16_t> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char16_t> *reductionPointer,
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  SD_HOST_DEVICE static std::basic_string<char16_t>
  op(const std::basic_string<char>& d1, std::basic_string<char> * /*params*/) {
    char16_t temp_output_buffer[SD_STRING_ASSIGN_TEMP_BUFFER_BYTES / sizeof(char16_t) + 1];
    const char* input_data = d1.data();
    const uint32_t input_length_bytes = static_cast<uint32_t>(d1.length());
    sd::LongType required_bytes_for_utf16 = sd::unicode::offsetUtf8StringInUtf16(input_data, input_length_bytes);

    if (required_bytes_for_utf16 > 0 && static_cast<size_t>(required_bytes_for_utf16) < sizeof(temp_output_buffer) ) {
      void* end_ptr = sd::unicode::utf8to16Ptr(input_data, input_data + input_length_bytes, temp_output_buffer);
      size_t char16_units_written = static_cast<char16_t*>(end_ptr) - temp_output_buffer;
      if (char16_units_written * sizeof(char16_t) == static_cast<size_t>(required_bytes_for_utf16)) {
        return std::basic_string<char16_t>(temp_output_buffer, char16_units_written);
      }
    }
    return std::basic_string<char16_t>();
  }
};

// --- Specialization: std::basic_string<char32_t> (UTF-32) -> std::basic_string<char> (UTF-8) ---
template <>
class Assign<std::basic_string<char32_t>, std::basic_string<char>> {
 public:
  static const bool requiresSpecial = false;
  static void execSpecial(const std::basic_string<char32_t> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char32_t> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char32_t> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char32_t> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char> *reductionPointer, // Z is std::basic_string<char>
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  SD_HOST_DEVICE static std::basic_string<char>
  op(const std::basic_string<char32_t>& d1, std::basic_string<char32_t> * /*params*/) {
    char temp_output_buffer[SD_STRING_ASSIGN_TEMP_BUFFER_BYTES];
    const char32_t* input_data = d1.data();
    const uint32_t input_length_char32_units = static_cast<uint32_t>(d1.length());
    sd::LongType required_bytes = sd::unicode::offsetUtf32StringInUtf8(input_data, input_length_char32_units);

    if (required_bytes > 0 && static_cast<size_t>(required_bytes) <= SD_STRING_ASSIGN_TEMP_BUFFER_BYTES) {
      void* end_ptr = sd::unicode::utf32to8Ptr(input_data, input_data + input_length_char32_units, temp_output_buffer);
      size_t bytes_written = static_cast<char*>(end_ptr) - temp_output_buffer;
      if (bytes_written == static_cast<size_t>(required_bytes)) {
        return std::basic_string<char>(temp_output_buffer, bytes_written);
      }
    }
    return std::basic_string<char>();
  }
};

// --- Specialization: std::basic_string<char> (UTF-8) -> std::basic_string<char32_t> (UTF-32) ---
template <>
class Assign<std::basic_string<char>, std::basic_string<char32_t>> {
 public:
  static const bool requiresSpecial = false;
  static void execSpecial(const std::basic_string<char> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char32_t> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char32_t> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char32_t> *reductionPointer,
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  SD_HOST_DEVICE static std::basic_string<char32_t>
  op(const std::basic_string<char>& d1, std::basic_string<char> * /*params*/) {
    char32_t temp_output_buffer[SD_STRING_ASSIGN_TEMP_BUFFER_BYTES / sizeof(char32_t) + 1];
    const char* input_data = d1.data();
    const uint32_t input_length_bytes = static_cast<uint32_t>(d1.length());
    sd::LongType required_bytes_for_utf32_data = sd::unicode::offsetUtf8StringInUtf32(input_data, input_length_bytes);

    if (required_bytes_for_utf32_data > 0 && static_cast<size_t>(required_bytes_for_utf32_data) < sizeof(temp_output_buffer) ) {
      void* end_ptr = sd::unicode::utf8to32Ptr(input_data, input_data + input_length_bytes, temp_output_buffer);
      size_t char32_units_written = static_cast<char32_t*>(end_ptr) - temp_output_buffer;
      if (char32_units_written * sizeof(char32_t) == static_cast<size_t>(required_bytes_for_utf32_data)) {
        return std::basic_string<char32_t>(temp_output_buffer, char32_units_written);
      }
    }
    return std::basic_string<char32_t>();
  }
};

// --- Identity Specializations ---
template <>
class Assign<std::basic_string<char>, std::basic_string<char>> {
 public:
  static const bool requiresSpecial = false;
  static void execSpecial(const std::basic_string<char> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char> *reductionPointer,
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif
  SD_HOST_DEVICE static std::basic_string<char>
  op(const std::basic_string<char>& d1, std::basic_string<char> * /*params*/) {
    return d1;
  }
};

template <typename X, typename Y, typename Z>
class LogPoissonLossFull {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X z, Y c) {
    auto zz = static_cast<Z>(z);
    auto zc = static_cast<Z>(c);
    return (sd::math::sd_exp<Y, Z>(c) - zz * zc +
            (zz * sd::math::sd_log<X, Z>(z) - zz +
             static_cast<Z>(0.5f) * sd::math::sd_log<Z, Z>(static_cast<Z>(SD_DOUBLE_PI_X) * zz)));
  }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X z, Y c, Z *params) { return op_logic(z, c); }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X z) {
    auto zz = static_cast<Z>(z);
    return (zz * sd::math::sd_log<Y, Z>(z) - zz +
            static_cast<Z>(0.5f) * sd::math::sd_log<Z, Z>(static_cast<Z>(SD_DOUBLE_PI_X) * zz));
  }
  static SD_HOST_DEVICE SD_INLINE X op_logic(X z, Y *params) {
    return (sd::math::sd_exp<X, X>(params[0]) - z * params[0] +
            (z * sd::math::sd_log<X, Z>(z) - z + static_cast<X>(0.5f) * sd::math::sd_log<X, Z>(SD_DOUBLE_PI_X * z)));
  }
  static SD_HOST_DEVICE Z op_simd(X z, Y c) { return op_logic(z, c); }
  static SD_HOST_DEVICE Z op_simd(X z, Y c, Z *params) { return op_logic(z, c, params); }
  static SD_HOST_DEVICE Z op_simd(X z) { return op_logic(z); }
  static SD_HOST_DEVICE X op_simd(X z, Y *params) { return op_logic(z, params); }

 public:
  static SD_HOST_DEVICE Z op(X z, Y c) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(z, c);
    else return op_simd(z, c);
  }
  static SD_HOST_DEVICE Z op(X z, Y c, Z *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(z, c, params);
    else return op_simd(z, c, params);
  }
  static SD_HOST_DEVICE Z op(X z) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(z);
    else return op_simd(z);
  }
  static SD_HOST_DEVICE X op(X z, Y *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(z, params);
    else return op_simd(z, params);
  }
};



template <typename X>
class Celu {
 private:
  static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X *params) {
    X alpha = params[0];
    return sd::math::sd_max<X>(static_cast<X>(0), d1) +
           sd::math::sd_min<X>(static_cast<X>(0), alpha * (sd::math::sd_exp<X, X>(d1/alpha) - static_cast<X>(1)));
  }
  static X op_simd(X d1, X *params) { return op_logic(d1, params); }

 public:
  no_op_exec_special_same no_op_exec_special_same_cuda;
  static SD_HOST_DEVICE X op(X d1, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)
      return op_logic(d1, params);
    else
      return op_simd(d1, params);
  }
};




// ELUAlpha - simple conditional with parameter
DECLARE_UNARY_CLIPPING_OP(ELUAlpha,
                          X alpha = params[0];
                              return d1 > static_cast<X>(0) ? d1 : alpha * (sd::math::sd_exp<X COMMA X>(d1) - static_cast<X>(1));
)

// PReLU - simple conditional with parameter
DECLARE_UNARY_CLIPPING_OP(PReLU,
                          X alpha = params[0];
                              return d1 > static_cast<X>(0) ? d1 : alpha * d1;
)

// ThresholdedReLU - simple threshold operation
DECLARE_UNARY_CLIPPING_OP(ThresholdedReLU,
                          X theta = params[0];
                              return d1 > theta ? d1 : static_cast<X>(0);
)

#define DECLARE_INTEGER_ONLY_BINARY_OP(OP_NAME, OPERATION)                                                \
  template <typename X>                                                                                    \
  class OP_NAME {                                                                                         \
   private:                                                                                               \
    static SD_INLINE X op_logic(X d1, X d2) {                                                            \
      if constexpr (std::is_integral<X>::value) {                                                        \
        return OPERATION;                                                                                \
      } else {                                                                                           \
        /* For floating point, interpret bits as integer */                                              \
        using IntType = typename std::conditional<sizeof(X) == 4, uint32_t,                              \
                        typename std::conditional<sizeof(X) == 8, uint64_t,                              \
                        typename std::conditional<sizeof(X) == 2, uint16_t, uint8_t>::type>::type>::type;\
        IntType i1, i2, iresult;                                                                         \
        std::memcpy(&i1, &d1, sizeof(X));                                                                \
        std::memcpy(&i2, &d2, sizeof(X));                                                                \
        /* Create a lambda to evaluate OPERATION with integer values */                                  \
        iresult = [&]() {                                                                                \
          auto d1 = i1;                                                                                  \
          auto d2 = i2;                                                                                  \
          return OPERATION;                                                                             \
        }();                                                                                             \
        X result;                                                                                        \
        std::memcpy(&result, &iresult, sizeof(X));                                                       \
        return result;                                                                                   \
      }                                                                                                  \
    }                                                                                                     \
    static SD_INLINE X op_logic(X d1, X d2, X *params) { return op_logic(d1, d2); }                       \
    static SD_INLINE SD_HOST_DEVICE X op_simd(X d1, X d2) { return op_logic(d1, d2); }                    \
    static SD_INLINE SD_HOST_DEVICE X op_simd(X d1, X d2, X *params) { return op_logic(d1, d2, params); } \
                                                                                                          \
   public:                                                                                                \
    static SD_INLINE SD_HOST_DEVICE X op(X d1, X d2) {                                                    \
      if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)                                   \
        return op_logic(d1, d2);                                                                          \
      else                                                                                                \
        return op_simd(d1, d2);                                                                           \
    }                                                                                                     \
    static SD_INLINE SD_HOST_DEVICE X op(X d1, X d2, X *params) {                                         \
      if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)                                   \
        return op_logic(d1, d2, params);                                                                  \
      else                                                                                                \
        return op_simd(d1, d2, params);                                                                   \
    }                                                                                                     \
  };
// Then use it exactly as before:
DECLARE_INTEGER_ONLY_BINARY_OP(ShiftLeft, d1 << d2)
DECLARE_INTEGER_ONLY_BINARY_OP(ShiftRight, d1 >> d2)
DECLARE_INTEGER_ONLY_BINARY_OP(IntOr, d2 | d1)
DECLARE_INTEGER_ONLY_BINARY_OP(IntAnd, d2 & d1)
DECLARE_INTEGER_ONLY_BINARY_OP(IntXor, d2 ^ d1)

#define DECLARE_INTEGER_ONLY_BINARY_TEMPLATE_OP(OP_NAME, OPERATION)                                       \
  template <typename X>                                                                                   \
  class OP_NAME {                                                                                         \
   private:                                                                                               \
    static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X d2) {                                             \
      if constexpr (std::is_integral<X>::value) {                                                        \
        return OPERATION;                                                                                \
      } else {                                                                                           \
        /* For floating point types, the template operation might not be valid */                        \
        /* This requires special handling based on the specific operation */                             \
        return static_cast<X>(0); /* Default fallback - you may need to specialize per operation */     \
      }                                                                                                  \
    }                                                                                                     \
    static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X d2, X *params) { return op_logic(d1, d2); }        \
    static SD_HOST_DEVICE SD_INLINE X op_simd(X d1, X d2) { return op_logic(d1, d2); }                    \
    static SD_HOST_DEVICE SD_INLINE X op_simd(X d1, X d2, X *params) { return op_logic(d1, d2, params); } \
                                                                                                          \
   public:                                                                                                \
    static SD_HOST_DEVICE SD_INLINE X op(X d1, X d2) {                                                    \
      if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)                                   \
        return op_logic(d1, d2);                                                                          \
      else                                                                                                \
        return op_simd(d1, d2);                                                                           \
    }                                                                                                     \
    static SD_HOST_DEVICE SD_INLINE X op(X d1, X d2, X *params) {                                         \
      if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)                                   \
        return op_logic(d1, d2, params);                                                                  \
      else                                                                                                \
        return op_simd(d1, d2, params);                                                                   \
    }                                                                                                     \
  };

// Then use it for cyclic shift operations:
DECLARE_INTEGER_ONLY_BINARY_TEMPLATE_OP(CyclicShiftLeft, sd::math::sd_rotl<X>(d1, d2))
DECLARE_INTEGER_ONLY_BINARY_TEMPLATE_OP(CyclicShiftRight, sd::math::sd_rotr<X>(d1, d2))


// Undefine before redefining with improved SFINAE-based constraint
#undef DECLARE_INTEGER_ONLY_BINARY_TEMPLATE_OP

#define DECLARE_INTEGER_ONLY_BINARY_TEMPLATE_OP(OP_NAME, OPERATION)                                       \
  template <typename X, typename = typename std::enable_if<std::is_integral<X>::value>::type>             \
  class OP_NAME {                                                                                         \
   private:                                                                                               \
    static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X d2) { return OPERATION; }                          \
    static SD_HOST_DEVICE SD_INLINE X op_logic(X d1, X d2, X *params) { return op_logic(d1, d2); }        \
    static SD_HOST_DEVICE SD_INLINE X op_simd(X d1, X d2) { return op_logic(d1, d2); }                    \
    static SD_HOST_DEVICE SD_INLINE X op_simd(X d1, X d2, X *params) { return op_logic(d1, d2, params); } \
                                                                                                          \
   public:                                                                                                \
    static SD_HOST_DEVICE SD_INLINE X op(X d1, X d2) {                                                    \
      if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)                                   \
        return op_logic(d1, d2);                                                                          \
      else                                                                                                \
        return op_simd(d1, d2);                                                                           \
    }                                                                                                     \
    static SD_HOST_DEVICE SD_INLINE X op(X d1, X d2, X *params) {                                         \
      if constexpr (simdOps::is_simd_unsupported_return_type<X>::value)                                   \
        return op_logic(d1, d2, params);                                                                  \
      else                                                                                                \
        return op_simd(d1, d2, params);                                                                   \
    }                                                                                                     \
  };


template <typename X, typename Y, typename Z>
class Mod {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y d2) {
    auto dx = static_cast<X>(d2);
    auto f = sd::math::sd_floor<X, X>(d1 / dx);
    auto r = f * dx;
    return static_cast<Z>(d1 - r);
  }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y d2, Z *params) { return op_logic(d1, d2); }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y *params) { return op_logic(d1, params[0]); }
  static SD_HOST_DEVICE Z op_simd(X d1, Y d2) { return op_logic(d1, d2); }
  static SD_HOST_DEVICE Z op_simd(X d1, Y d2, Z *params) { return op_logic(d1, d2, params); }
  static SD_HOST_DEVICE Z op_simd(X d1, Y *params) { return op_logic(d1, params); }

 public:
  static SD_HOST_DEVICE Z op(X d1, Y d2) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, d2);
    else return op_simd(d1, d2);
  }
  static SD_HOST_DEVICE Z op(X d1, Y d2, Z *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, d2, params);
    else return op_simd(d1, d2, params);
  }
  static SD_HOST_DEVICE Z op(X d1, Y *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, params);
    else return op_simd(d1, params);
  }
};


DECLARE_BINARY_COPY_OP(ReverseMod,
                       static_cast<Z>(static_cast<int>(d2) % static_cast<int>(d1)),
                       static_cast<Z>(static_cast<int>(d2) % static_cast<int>(d1)),
                       static_cast<Z>(d1),
                       static_cast<Z>(static_cast<int>(params[0]) % static_cast<int>(d1))
)

template <typename X, typename Z>
class Epsilon {
 private:
  static SD_HOST_DEVICE SD_INLINE  Z op_logic(X d1, X d2) {
    X diff = d1 - d2;
    X absDiff = sd::math::sd_abs<X,X>(diff);
    if (absDiff <= static_cast<X>(SD_MIN_V)) return static_cast<Z>(1);
    return static_cast<Z>(0);
  }
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, X d2, X *params) {
    X diff = d1 - d2;
    X absDiff = sd::math::sd_abs<X,X>(diff);
    if(params != nullptr && absDiff <= static_cast<X>(params[0])) {
      return static_cast<Z>(1);
    } else if(absDiff <= static_cast<X>(1e-5)) {
      return static_cast<Z>(1);
    }
    return static_cast<Z>(0);
  }
  static SD_HOST_DEVICE SD_INLINE  Z op_logic(X d1, X *params) { return static_cast<Z>(d1); }
  static SD_HOST_DEVICE SD_INLINE  Z op_simd(X d1, X d2) { return op_logic(d1, d2); }
  static SD_HOST_DEVICE SD_INLINE  Z op_simd(X d1, X d2, X *params) { return op_logic(d1, d2, params); }
  static SD_HOST_DEVICE SD_INLINE  Z op_simd(X d1, X *params) { return op_logic(d1, params); }

 public:
  static SD_HOST_DEVICE SD_INLINE  Z op(X d1, X d2) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1, d2);
    else return op_simd(d1, d2);
  }
  static SD_HOST_DEVICE SD_INLINE  Z op(X d1, X d2, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1, d2, params);
    else return op_simd(d1, d2, params);
  }
  static SD_HOST_DEVICE SD_INLINE  Z op(X d1, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1, params);
    else return op_simd(d1, params);
  }
};

template <typename X, typename Z>
class MatchConditionBool {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, X *extraParams) {
    X compare = extraParams[0];
    X eps = extraParams[1];
    auto mode = static_cast<int>(extraParams[2]);
    sd_debug("value: %f; comp: %f; eps: %f; mode: %i;\n", d1, compare, eps, mode);

    switch (mode) {
      case 0:  // equals
        return sd::math::sd_abs<X,X>(d1 - compare) <= eps ? static_cast<Z>(1) : static_cast<Z>(0);
      case 1:  // not equals
        return sd::math::sd_abs<X,X>(d1 - compare) > eps ? static_cast<Z>(1) : static_cast<Z>(0);
      case 2:  // less_than
        return d1 < compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 3:  // greater_than
        return d1 > compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 4:  // less_or_equals_than
        return d1 <= compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 5:  // greater_or_equals_than
        return d1 >= compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 6:  // abs_less_than
        return sd::math::sd_abs<X,X>(d1) < compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 7:  // abs_greater_than
        return sd::math::sd_abs<X,X>(d1) > compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 8:  // is inf
        return sd::math::sd_isinf<X>(d1) ? static_cast<Z>(1) : static_cast<Z>(0);
      case 9:  // is nan
        return sd::math::sd_isnan<X>(d1) ? static_cast<Z>(1) : static_cast<Z>(0);
      case 10:
        return (d1 == compare) ? static_cast<Z>(1) : static_cast<Z>(0);
      case 11:
        return (d1 != compare) ? static_cast<Z>(1) : static_cast<Z>(0);
      case 12:  // abs_greater_or_equals_than
        return sd::math::sd_abs<X,X>(d1) >= compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 13:  // abs_less_or_equals_than
        return sd::math::sd_abs<X,X>(d1) <= compare ? static_cast<Z>(1) : static_cast<Z>(0);
      case 14:
        // isFinite
        return !(sd::math::sd_isinf<X>(d1) || sd::math::sd_isnan<X>(d1)) ? static_cast<Z>(1) : static_cast<Z>(0);
      case 15:
        // isInfinite
        return (sd::math::sd_isinf<X>(d1) || sd::math::sd_isnan<X>(d1)) ? static_cast<Z>(1) : static_cast<Z>(0);
      default:
        sd_debug("Undefined match condition: [%i]\n", mode);
    }
    return static_cast<Z>(d1);
  }

  // Remove SD_OP_DEF to avoid SIMD issues with float16/bfloat16
  static Z op_simd(X d1, X *extraParams) { return op_logic(d1, extraParams); }

 public:
  // Fix: Use explicit declarations instead of problematic macros
  no_op_exec_special no_op_exec_special_cuda;

  // Special handling for bool operations with type compatibility
  static const bool requiresSpecialAccumulation = false;

  // Primary execSpecial function with Z_TYPE* extraParams (for boolean case)
  static void execSpecial(const X *x, const sd::LongType *xShapeInfo, Z *extraParams, Z *result,
                          const sd::LongType *resultShapeInfoBuffer, sd::LongType *dimension, sd::LongType dimensionLength,
                          const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset) {}

  // Template overload to handle type conversions (for cases where extraParams is sd::LongType*)
  template<typename ExtraParamsType>
  static void execSpecial(const X *x, const sd::LongType *xShapeInfo, ExtraParamsType *extraParams, Z *result,
                          const sd::LongType *resultShapeInfoBuffer, sd::LongType *dimension, sd::LongType dimensionLength,
                          const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset) {
    // Handle type conversion if needed - this handles the sd::LongType* to Z* conversion
    // For most cases, this will be empty since we don't actually implement special accumulation
  }

#ifdef __CUDACC__
  static SD_INLINE SD_DEVICE void execSpecialCuda(
      const X *dx, const sd::LongType *xShapeInfo, Z *extraParams, Z *result,
      const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength,
      Z *reductionBuffer, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets) {}

  template<typename ExtraParamsType>
  static SD_INLINE SD_DEVICE void execSpecialCuda(
      const X *dx, const sd::LongType *xShapeInfo, ExtraParamsType *extraParams, Z *result,
      const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength,
      Z *reductionBuffer, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  static SD_HOST_DEVICE SD_INLINE Z op(X d1, X *extraParams) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1, extraParams);
    else
      return op_simd(d1, extraParams);
  }
};


DECLARE_MULTI_OP_SIMD_SAFE(Or,
                           return d2 + d1;,
                           if (params != nullptr) {
                             auto comp = params[0];
                             return d1 != comp || d2 != comp ? static_cast<Z>(1) : static_cast<Z>(0);
                           } else {
                             auto b1 = static_cast<bool>(d1);
                             auto b2 = static_cast<bool>(d2);
                             return b1 || b2 ? static_cast<Z>(1) : static_cast<Z>(0);
                           },
                           return d1;,
                           return static_cast<Z>(119);
)


DECLARE_XOR_SIMD_SAFE(Xor,
                      return d2 + d1;,
                      if (params != nullptr) {
                        auto comp = params[0];
                        return ((d1 == comp && d2 != comp) || (d1 != comp && d2 == comp)) ? static_cast<Z>(1) : static_cast<Z>(0);
                      } else {
                        auto b1 = static_cast<bool>(d1);
                        auto b2 = static_cast<bool>(d2);
                        return (!b1 && b2) || (b1 && !b2) ? static_cast<Z>(1) : static_cast<Z>(0);
                      },
                      return d1;
)




DECLARE_NOT_SIMD_SAFE(Not,
                      return static_cast<Z>(0);,
                      return d1 != d2 ? static_cast<Z>(1) : static_cast<Z>(0);,
                      auto b1 = static_cast<bool>(d1);
                          return !b1;
)


DECLARE_ACCUMULATION_SIMD_SAFE_OP(Variance,
                                  X mean = static_cast<InterType>(params[0]);
                                      X ret = d1 - mean;
                                      return ret * ret;,
                                  SUM, static_cast<X>(0.0f),
                                  old + opOutput,
                                  old + opOutput,
                                  static_cast<Z>(reduction / static_cast<InterType>(n - 1))
)

DECLARE_ACCUMULATION_SIMD_SAFE_OP(StandardDeviation,
                                  InterType mean = static_cast<InterType>(params[0]);
                                      InterType ret = d1 - mean;
                                      return ret * ret;,
                                  SUM, static_cast<X>(0.0f),
                                  old + opOutput,
                                  old + opOutput,
                                  sd::math::sd_sqrt<InterType COMMA Z>(static_cast<InterType>(reduction / static_cast<InterType>(n - 1)))
)


DECLARE_ACCUMULATION_SIMD_SAFE_OP(ShannonEntropy,
                                  auto p = d1;
                                      return static_cast<Z>(p) * sd::math::sd_log2<X COMMA Z>(p);,
                                  SUM, static_cast<X>(0),
                                  opOutput + old,
                                  opOutput + old,
                                  -reduction
)


DECLARE_ACCUMULATION_SIMD_SAFE_OP(LogEntropy,
                                  return static_cast<InterType>(d1) * sd::math::sd_log<X COMMA InterType>(d1);,
                                  SUM, static_cast<X>(0),
                                  opOutput + old,
                                  opOutput + old,
                                  sd::math::sd_log<InterType COMMA Z>(-reduction)
)



template <typename X, typename Z>
class IndexAbsoluteMax {
 public:
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> val,
                                                                        X *extraParams) {
    return sd::math::sd_abs<X,X>(val);
  }
  static SD_HOST_DEVICE SD_INLINE   functions::indexreduce::IndexValue<X> update(
      functions::indexreduce::IndexValue<X> &old, functions::indexreduce::IndexValue<X> &opOutput, X *extraParams) {
    opOutput.value = sd::math::sd_abs<X,X>(opOutput.value);
    old.value = sd::math::sd_abs<X,X>(old.value);
    if (opOutput.value > old.value) return opOutput;
#ifdef __CUDACC__
    else if (opOutput.value == old.value && opOutput.index < old.index)
     return opOutput;
#endif
    return old;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> merge(functions::indexreduce::IndexValue<X> f1,
                                                                           functions::indexreduce::IndexValue<X> f2,
                                                                           X *extraParams) {
    if (sd::math::sd_abs<X,X>(f1.value) > sd::math::sd_abs<X,X>(f2.value)) return f2;
    return f1;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> postProcess(
      functions::indexreduce::IndexValue<X> reduction, int n, int xOffset, X *dx, int incx, X *extraParams, X *result) {
    return reduction;
  }
  static SD_HOST_DEVICE SD_INLINE  X startingValue(const X *input) { return static_cast<X>(0); }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> startingIndexValue(const X *input) {
    functions::indexreduce::IndexValue<X> local;
    local.value = startingValue(input);
    local.index = 0;
    return local;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> d1,
                                                                        functions::indexreduce::IndexValue<X> d2,
                                                                        X *extraParams) {
    return d1;
  }
};

template <typename X, typename Z>
class IndexAbsoluteMin {
 public:
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> val,
                                                                        X *extraParams) {
    return val;
  }
  static SD_HOST_DEVICE SD_INLINE  X startingValue(const X *input) { return sd::DataTypeUtils::infOrMax<X>(); }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> startingIndexValue(const X *input) {
    functions::indexreduce::IndexValue<X> local;
    local.value = startingValue(input);
    local.index = 0;
    return local;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> update(
      functions::indexreduce::IndexValue<X> &old, functions::indexreduce::IndexValue<X> &opOutput, X *extraParams) {
    opOutput.value = sd::math::sd_abs<X,X>(opOutput.value);
    old.value = sd::math::sd_abs<X,X>(old.value);
    if (opOutput.value < old.value) return opOutput;
#ifdef __CUDACC__
    else if (opOutput.value == old.value && opOutput.index < old.index)
     return opOutput;
#endif
    return old;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> merge(functions::indexreduce::IndexValue<X> f1,
                                                                           functions::indexreduce::IndexValue<X> f2,
                                                                           X *extraParams) {
    if (sd::math::sd_abs<X,X>(f1.value) < sd::math::sd_abs<X,X>(f2.value)) return f2;
    return f1;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> postProcess(
      functions::indexreduce::IndexValue<X> reduction, int n, int xOffset, X *dx, int incx, X *extraParams, X *result) {
    return reduction;
  }
  static  SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> d1,
                                                                        functions::indexreduce::IndexValue<X> d2,
                                                                        X *extraParams) {
    return d1;
  }
};

template <typename X, typename Z>
class FirstIndex {
 public:
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> val,
                                                                        X *extraParams) {
    return val;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> update(functions::indexreduce::IndexValue<X> &old,
                                                                     functions::indexreduce::IndexValue<X> &opOutput,
                                                                     X *extraParams) {
#ifdef __CUDACC__
    if (opOutput.index < 0) return old;
#endif
    auto res = MatchConditionBool<X, X>::op(opOutput.value, extraParams);
    if (res == static_cast<X>(0)) return old;
    if (old.index < 0) return opOutput;
    if (old.index > opOutput.index) return opOutput;
    return old;
  }
  static SD_HOST_DEVICE SD_INLINE  X startingValue(const X *input) { return -sd::DataTypeUtils::infOrMax<X>(); }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> startingIndexValue(const X *input) {
    functions::indexreduce::IndexValue<X> local;
    local.value = startingValue(input);
    local.index = -1;
    return local;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> d1,
                                                                        functions::indexreduce::IndexValue<X> d2,
                                                                        X *extraParams) {
    return d1;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> merge(functions::indexreduce::IndexValue<X> f1,
                                                                           functions::indexreduce::IndexValue<X> f2,
                                                                           X *extraParams) {
    if (f1.index > f2.index) return f2;
    return f1;
  }
  static SD_HOST_DEVICE SD_INLINE functions::indexreduce::IndexValue<X> postProcess(
      functions::indexreduce::IndexValue<X> reduction, int n, int xOffset, X *dx, int incx, X *extraParams, X *result) {
    return reduction;
  }
};


template <typename X, typename Z>
class LastIndex {
 public:
  static SD_HOST_DEVICE SD_INLINE functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> val,
                                                                        X *extraParams) {
    return val;
  }
  static SD_HOST_DEVICE SD_INLINE  functions::indexreduce::IndexValue<X> update(functions::indexreduce::IndexValue<X> &old,
                                                                     functions::indexreduce::IndexValue<X> &opOutput,
                                                                     X *extraParams) {
#ifdef __CUDACC__
    if (opOutput.index < 0) return old;
#endif
    auto res = MatchConditionBool<X, X>::op(opOutput.value, extraParams);
    if (res == static_cast<X>(0)) return old;
    if (old.index < 0) return opOutput;
    if (old.index < opOutput.index) return opOutput;
    return old;
  }
  static SD_HOST_DEVICE SD_INLINE X startingValue(const X *input) { return -sd::DataTypeUtils::infOrMax<X>(); }
  static SD_HOST_DEVICE SD_INLINE functions::indexreduce::IndexValue<X> startingIndexValue(const X *input) {
    functions::indexreduce::IndexValue<X> local;
    local.value = startingValue(input);
    local.index = -1;
    return local;
  }
  static SD_HOST_DEVICE SD_INLINE functions::indexreduce::IndexValue<X> op(functions::indexreduce::IndexValue<X> d1,
                                                                        functions::indexreduce::IndexValue<X> d2,
                                                                        X *extraParams) {
    return d1;
  }
  static SD_HOST_DEVICE SD_INLINE functions::indexreduce::IndexValue<X> merge(functions::indexreduce::IndexValue<X> f1,
                                                                           functions::indexreduce::IndexValue<X> f2,
                                                                           X *extraParams) {
    if (f1.index < f2.index) return f2;
    return f1;
  }
  static SD_HOST_DEVICE SD_INLINE functions::indexreduce::IndexValue<X> postProcess(
      functions::indexreduce::IndexValue<X> reduction, int n, int xOffset, X *dx, int incx, X *extraParams, X *result) {
    return reduction;
  }
};

DECLARE_ACCUMULATION_SIMD_SAFE_OP(Entropy,
                                  return static_cast<InterType>(d1) * sd::math::sd_log<X COMMA InterType>(d1);,
                                  SUM, static_cast<X>(0),
                                  opOutput + old,
                                  opOutput + old,
                                  static_cast<Z>(-reduction)
)


template <typename X, typename Z>
class SummaryStatsVariance {
 public:
  static SD_HOST_DEVICE SD_INLINE Z getValue(const bool biasCorrected, functions::summarystats::SummaryStatsData<X> val) {
    if (biasCorrected) {
      Z ret = static_cast<Z>(val.varianceBiasCorrected());
      if (ret < static_cast<Z>(0.0f)) return static_cast<Z>(val.variance());
      return ret;
    }
    return static_cast<Z>(val.variance());
  }
  static SD_HOST_DEVICE SD_INLINE functions::summarystats::SummaryStatsData<X> op(
      functions::summarystats::SummaryStatsData<X> d1, Z *extraParams) {
    return d1;
  }
};

template <typename X, typename Z>
class SummaryStatsStandardDeviation {
 public:
  static SD_HOST_DEVICE SD_INLINE Z getValue(const bool biasCorrected, functions::summarystats::SummaryStatsData<X> val) {
    if (biasCorrected) {
      auto ret = static_cast<Z>(val.varianceBiasCorrected());
      if (ret < static_cast<Z>(0.0f))
        return sd::math::sd_sqrt<double, Z>(val.variance());
      else
        return sd::math::sd_sqrt<double, Z>(ret);
    }
    return sd::math::sd_sqrt<double, Z>(val.variance());
  }
  static SD_HOST_DEVICE SD_INLINE functions::summarystats::SummaryStatsData<X> op(
      functions::summarystats::SummaryStatsData<X> d1, Z *extraParams) {
    return d1;
  }
};

// =============================================================================
// MISSING ADDITIONAL UNARY OPERATIONS
// =============================================================================

DECLARE_UNARY_SIMD_SAFE_OP(Sqr,
                           return sd::math::sd_pow<X COMMA X COMMA X>(d1 COMMA static_cast<X>(2));
)

DECLARE_UNARY_MATH_OP_XZ(Sqrt, sd_sqrt)

// For RSqrt, use the complex math version:
DECLARE_UNARY_COMPLEX_MATH_OP_XZ(RSqrt,
                                 static_cast<Z>(1.0) / static_cast<Z>(sd::math::sd_sqrt<X, Z>(d1)))



DECLARE_BINARY_COPY_OP(RelativeError,
                       static_cast<Z>(sd::math::sd_re<X>(d1, static_cast<X>(d2))),
                       static_cast<Z>(sd::math::sd_re<X>(d1, static_cast<X>(d2))),
                       static_cast<Z>(0),
                       static_cast<Z>(0)
)



template <typename X, typename Y, typename Z>
class BinaryRelativeError {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y d2, Z *params) {
    // Type guard for non-arithmetic types
    if constexpr (any_my_string_v<X, Y, Z>) {
      return static_cast<Z>(d1);  // For non-arithmetic types, just return cast of first value
    } else {
      // All arithmetic operations inside else block
      X threshold = static_cast<X>(params[0]);
      return sd::math::sd_re<X>(d1, static_cast<X>(d2)) > threshold ? static_cast<Z>(1) : static_cast<Z>(0);
    }
  }
  
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1) { 
    // Type guard for non-arithmetic types
    if constexpr (any_my_string_v<X, Z>) {
      return static_cast<Z>(d1);  // For non-arithmetic types, just return cast of first value
    } else {
      return static_cast<Z>(0);
    }
  }
  
  SD_HOST_DEVICE SD_INLINE static Z op_simd(X d1, Y d2, Z *params) { 
    return op_logic(d1, d2, params); 
  }
  
  SD_HOST_DEVICE SD_INLINE static Z op_simd(X d1) { 
    return op_logic(d1); 
  }

 public:
  no_op_exec_special no_op_exec_special_cuda
  
  static SD_HOST_DEVICE Z op(X d1, Y d2, Z *params) {
    if constexpr (any_my_string_v<X, Y, Z> ||
                  simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, d2, params);
    else 
      return op_simd(d1, d2, params);
  }
  
  static SD_HOST_DEVICE Z op(X d1) {
    if constexpr (any_my_string_v<X, Z> ||
                  simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1);
    else 
      return op_simd(d1);
  }
};


// BinaryRelativeError - Custom conditional logic, manual implementation needed
template <typename X, typename Y, typename Z>
class BinaryMinimumAbsoluteRelativeError {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, X *params) {
    // Type guard for non-arithmetic types
    if constexpr (any_my_string_v<X, Z>) {
      return static_cast<Z>(d1);  // For non-arithmetic types, just return cast of first value
    } else {
      // All arithmetic operations inside else block
      X d2 = params[0];
      X thresholdRelative = params[1];
      X thresholdAbsolute = params[2];
      return sd::math::sd_re<X>(d1, d2) > thresholdRelative
             ? (sd::math::sd_abs<X,X>(d1 - static_cast<X>(d2)) < thresholdAbsolute ? static_cast<Z>(0)
                                                                                   : static_cast<Z>(1))
             : static_cast<Z>(0);
    }
  }
  
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, Y d2, Z *params) {
    // Type guard for non-arithmetic types
    if constexpr (any_my_string_v<X, Y, Z>) {
      return static_cast<Z>(d1);  // For non-arithmetic types, just return cast of first value
    } else {
      // All arithmetic operations inside else block
      X thresholdRelative = static_cast<X>(params[0]);
      X thresholdAbsolute = static_cast<X>(params[1]);
      return sd::math::sd_re<X>(d1, static_cast<X>(d2)) > thresholdRelative
             ? (sd::math::sd_abs<X,X>(d1 - static_cast<X>(d2)) < thresholdAbsolute ? static_cast<Z>(0)
                                                                                   : static_cast<Z>(1))
             : static_cast<Z>(0);
    }
  }
  
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1) { 
    // Type guard for non-arithmetic types
    if constexpr (any_my_string_v<X, Z>) {
      return static_cast<Z>(d1);  // For non-arithmetic types, just return cast of first value
    } else {
      return static_cast<Z>(0);
    }
  }
  
  SD_HOST_DEVICE SD_INLINE static Z op_simd(X d1, X *params) { 
    return op_logic(d1, params); 
  }
  
  SD_HOST_DEVICE SD_INLINE static Z op_simd(X d1, Y d2, Z *params) { 
    return op_logic(d1, d2, params); 
  }
  
  SD_HOST_DEVICE SD_INLINE static Z op_simd(X d1) { 
    return op_logic(d1); 
  }
  
 public:
  no_op_exec_special no_op_exec_special_cuda
  
  static SD_HOST_DEVICE SD_INLINE Z op(X d1, X *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1, params);
    else 
      return op_simd(d1, params);
  }
  
  static SD_HOST_DEVICE SD_INLINE Z op(X d1, Y d2, Z *params) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value ||
                  simdOps::is_simd_unsupported_argument_type<Y>::value)
      return op_logic(d1, d2, params);
    else 
      return op_simd(d1, d2, params);
  }
  
  static SD_HOST_DEVICE SD_INLINE Z op(X d1) {
    if constexpr (simdOps::is_simd_unsupported_return_type<Z>::value ||
                  simdOps::is_simd_unsupported_argument_type<X>::value)
      return op_logic(d1);
    else 
      return op_simd(d1);
  }
};



// =============================================================================
// REDUCE OPERATIONS
// =============================================================================

DECLARE_REDUCE_OP(Sum, SUM, static_cast<X>(0.0f), opOutput + old, opOutput + old, reduction)
DECLARE_REDUCE_OP(Prod, PRODUCT, static_cast<X>(1), opOutput * old, opOutput * old, reduction)
DECLARE_REDUCE_OP(Max, MAX, -sd::DataTypeUtils::infOrMax<X>(),
                  sd::math::sd_max<X>(old, opOutput), sd::math::sd_max<X>(opOutput, old), reduction)
DECLARE_REDUCE_OP(Min, MIN, sd::DataTypeUtils::infOrMax<X>(),
                  sd::math::sd_min<X>(old, opOutput), sd::math::sd_min<X>(opOutput, old), reduction)

DECLARE_ACCUMULATION_SIMD_SAFE_OP(Mean,
                                  return static_cast<InterType>(d1);,
                                  SUM,
                                  static_cast<X>(0),
                                  old + opOutput,
                                  old + opOutput,
                                  reduction / static_cast<InterType>(n)
)

// =============================================================================
// INDEX REDUCE OPERATIONS
// =============================================================================

DECLARE_INDEX_REDUCE_OP(IndexMax, -sd::DataTypeUtils::infOrMax<X>(),
                        opOutput.value > old.value, f1.value > f2.value)
DECLARE_INDEX_REDUCE_OP(IndexMin, sd::DataTypeUtils::infOrMax<X>(),
                        opOutput.value < old.value, f1.value < f2.value)


DECLARE_BINARY_COPY_OP(LogPoissonLoss,
                       sd::math::sd_exp<Y COMMA Z>(d2) - static_cast<Z>(d1) * static_cast<Z>(d2),
                       sd::math::sd_exp<Y COMMA Z>(d2) - static_cast<Z>(d1) * static_cast<Z>(d2),
                       static_cast<Z>(d1),
                       sd::math::sd_exp<Y COMMA Z>(params[0]) - static_cast<Z>(d1) * static_cast<Z>(params[0])
)



// LogicalNot - using existing binary copy pattern
DECLARE_BINARY_COPY_OP(LogicalNot,
                       static_cast<Z>(!((int)d1 && (int)d2)),
                       static_cast<Z>(!(static_cast<int>(d1) && static_cast<int>(d2))),
                       static_cast<Z>(d1),
                       static_cast<Z>(119)
)

// LogicalXor - bitwise XOR logic
DECLARE_BINARY_COPY_OP(LogicalXor,
                       static_cast<Z>((static_cast<int>(d1) | static_cast<int>(d2)) & ~(static_cast<int>(d1) & static_cast<int>(d2))),
                       static_cast<Z>((static_cast<int>(d1) | static_cast<int>(d2)) & ~(static_cast<int>(d1) & static_cast<int>(d2))),
                       static_cast<Z>(d1),
                       static_cast<Z>(119)
)

// LogicalAnd - bitwise AND logic
DECLARE_BINARY_COPY_OP(LogicalAnd,
                       static_cast<Z>(static_cast<int>(d1) & static_cast<int>(d2)),
                       static_cast<Z>(static_cast<int>(d1) & static_cast<int>(d2)),
                       static_cast<Z>(d1),
                       static_cast<Z>(119)
)

// LogicalOr - bitwise OR logic
DECLARE_BINARY_COPY_OP(LogicalOr,
                       static_cast<Z>(static_cast<int>(d1) | static_cast<int>(d2)),
                       static_cast<Z>(static_cast<int>(d1) | static_cast<int>(d2)),
                       static_cast<Z>(d1),
                       static_cast<Z>(119)
)
template <typename X, typename Z>
class MatchCondition {
 private:
  static SD_HOST_DEVICE SD_INLINE Z op_logic(X d1, X compare, X eps, int mode) {
    switch (mode) {
      case 0: return static_cast<Z>(sd::math::sd_abs<X,X>(d1 - compare) <= eps ? 1 : 0);
      case 1: return static_cast<Z>(sd::math::sd_abs<X,X>(d1 - compare) > eps ? 1 : 0);
      case 2: return static_cast<Z>(d1 < compare ? 1 : 0);
      case 3: return static_cast<Z>(d1 > compare ? 1 : 0);
      case 4: return static_cast<Z>(d1 <= compare ? 1 : 0);
      case 5: return static_cast<Z>(d1 >= compare ? 1 : 0);
      case 6: return static_cast<Z>(sd::math::sd_abs<X,X>(d1) < compare ? 1 : 0);
      case 7: return static_cast<Z>(sd::math::sd_abs<X,X>(d1) > compare ? 1 : 0);
      case 8: return static_cast<Z>(sd::math::sd_isinf(d1) ? 1 : 0);
      case 9: return static_cast<Z>(sd::math::sd_isnan(d1) ? 1 : 0);
      case 10: return static_cast<Z>((d1 == compare) ? 1 : 0);
      case 11: return static_cast<Z>((d1 != compare) ? 1 : 0);
      case 12: return static_cast<Z>(sd::math::sd_abs<X,X>(d1) >= compare ? 1 : 0);
      case 13: return static_cast<Z>(sd::math::sd_abs<X,X>(d1) <= compare ? 1 : 0);
      case 14: return static_cast<Z>(!(sd::math::sd_isinf<X>(d1) || sd::math::sd_isnan<X>(d1)) ? 1 : 0);
      case 15: return static_cast<Z>(sd::math::sd_isinf<X>(d1) || sd::math::sd_isnan<X>(d1) ? 1 : 0);
      default: sd_printf("Undefined match condition: [%i]\n", mode);
    }
    return static_cast<Z>(d1);
  }

 public:
  static const bool requiresSpecialAccumulation = false;

  using InterType = typename AggregateType<Z>::type;

  // execSpecial signatures - matches what reduce_long.hpp expects
  static void execSpecial(const X *x, const sd::LongType *xShapeInfo, sd::LongType *extraParams, Z *result,
                          const sd::LongType *resultShapeInfoBuffer, sd::LongType *dimension, sd::LongType dimensionLength,
                          const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset) {}

#ifdef __CUDACC__
  static SD_INLINE SD_DEVICE void execSpecialCuda(
      const X *dx, const sd::LongType *xShapeInfo, sd::LongType *extraParams, Z *result,
      const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength,
      Z *reductionBuffer, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  // Core reduction operation methods - these use X* parameters for reduce_long
  SD_HOST_DEVICE SD_INLINE static Z startingValue(const X* input) { return static_cast<Z>(0); }
  SD_HOST_DEVICE SD_INLINE static InterType merge(InterType old, InterType opOutput, X* extraParams) { return old + opOutput; }
  SD_HOST_DEVICE SD_INLINE static InterType update(InterType old, InterType opOutput, X* extraParams) { return old + opOutput; }
  SD_HOST_DEVICE SD_INLINE static Z postProcess(InterType reduction, sd::LongType n, X* extraParams) { return static_cast<Z>(reduction); }

  // Core op methods - these use X* parameters for reduce_long to preserve comparison value types
  static SD_HOST_DEVICE SD_INLINE InterType op(X d1, X* extraParams) {
    if (extraParams == nullptr) return static_cast<InterType>(0);
    X compare = extraParams[0];
    X eps = extraParams[1];
    auto mode = static_cast<int>(extraParams[2]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  static SD_HOST_DEVICE SD_INLINE InterType op(X d1, X d2, X* extraParams) {
    if (extraParams == nullptr) {
      // If no extraParams, use d2 as compare value, default eps=0, mode=0 (equals)
      return static_cast<InterType>(op_logic(d1, d2, static_cast<X>(0), 0));
    }

    // Use d2 as comparison value, extraParams for eps and mode
    X compare = d2;
    X eps = extraParams[0];
    auto mode = static_cast<int>(extraParams[1]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  static SD_HOST_DEVICE SD_INLINE InterType op(X d1, X d2) {
    // Default: compare d1 to d2 with eps=0 and mode=0 (equals)
    return static_cast<InterType>(op_logic(d1, d2, static_cast<X>(0), 0));
  }

  // *** TEMPLATE OVERLOADS FOR DIFFERENT PARAMETER TYPES ***

  // Template overloads for X* parameters - only when X != Z
  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, X> && !std::is_same_v<X, Z>, InterType>::type
      op(X d1, ParamType* extraParams) {
    if (extraParams == nullptr) return static_cast<InterType>(0);
    X compare = extraParams[0];
    X eps = extraParams[1];
    auto mode = static_cast<int>(extraParams[2]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      op(X d1, X d2, ParamType* extraParams) {
    if (extraParams == nullptr) {
      return static_cast<InterType>(op_logic(d1, d2, static_cast<X>(0), 0));
    }
    X compare = d2;
    X eps = extraParams[0];
    auto mode = static_cast<int>(extraParams[1]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      merge(InterType old, InterType opOutput, ParamType* extraParams) {
    return old + opOutput;
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      update(InterType old, InterType opOutput, ParamType* extraParams) {
    return old + opOutput;
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, Z>::type
      postProcess(InterType reduction, sd::LongType n, ParamType* extraParams) {
    return static_cast<Z>(reduction);
  }

  // Template overloads for sd::LongType* parameters
  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, sd::LongType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      op(X d1, ParamType* extraParams) {
    if (extraParams == nullptr) return static_cast<InterType>(0);
    X compare = static_cast<X>(extraParams[0]);
    X eps = static_cast<X>(extraParams[1]);
    auto mode = static_cast<int>(extraParams[2]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, sd::LongType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      op(X d1, X d2, ParamType* extraParams) {
    if (extraParams == nullptr) {
      return static_cast<InterType>(op_logic(d1, d2, static_cast<X>(0), 0));
    }
    X compare = d2;
    X eps = static_cast<X>(extraParams[0]);
    auto mode = static_cast<int>(extraParams[1]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, sd::LongType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      merge(InterType old, InterType opOutput, ParamType* extraParams) {
    return old + opOutput;
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, sd::LongType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      update(InterType old, InterType opOutput, ParamType* extraParams) {
    return old + opOutput;
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_same_v<ParamType, sd::LongType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, Z>::type
      postProcess(InterType reduction, sd::LongType n, ParamType* extraParams) {
    return static_cast<Z>(reduction);
  }

  // Template overloads for float* parameters (for cases like bfloat16/bfloat16 with float* extraParams)
  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_floating_point_v<ParamType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      op(X d1, ParamType* extraParams) {
    if (extraParams == nullptr) return static_cast<InterType>(0);
    X compare = static_cast<X>(extraParams[0]);
    X eps = static_cast<X>(extraParams[1]);
    auto mode = static_cast<int>(extraParams[2]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_floating_point_v<ParamType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      op(X d1, X d2, ParamType* extraParams) {
    if (extraParams == nullptr) {
      return static_cast<InterType>(op_logic(d1, d2, static_cast<X>(0), 0));
    }
    X compare = d2;
    X eps = static_cast<X>(extraParams[0]);
    auto mode = static_cast<int>(extraParams[1]);
    return static_cast<InterType>(op_logic(d1, compare, eps, mode));
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_floating_point_v<ParamType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      merge(InterType old, InterType opOutput, ParamType* extraParams) {
    return old + opOutput;
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_floating_point_v<ParamType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, InterType>::type
      update(InterType old, InterType opOutput, ParamType* extraParams) {
    return old + opOutput;
  }

  template<typename ParamType>
  static SD_HOST_DEVICE SD_INLINE
      typename std::enable_if<std::is_floating_point_v<ParamType> && !std::is_same_v<ParamType, X> && !std::is_same_v<ParamType, Z>, Z>::type
      postProcess(InterType reduction, sd::LongType n, ParamType* extraParams) {
    return static_cast<Z>(reduction);
  }
};


// --- Specialization: std::basic_string<char32_t> (UTF-32) -> std::basic_string<char16_t> (UTF-16) ---
template <>
class Assign<std::basic_string<char32_t>, std::basic_string<char16_t>> {
 public:
  static const bool requiresSpecial = false;
  static SD_HOST_DEVICE void execSpecial(const std::basic_string<char32_t> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char16_t> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char32_t> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char32_t> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char16_t> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char32_t> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char16_t> *reductionPointer, // Z is std::basic_string<char16_t>
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  SD_HOST_DEVICE SD_INLINE  static std::basic_string<char16_t>
  op(const std::basic_string<char32_t>& d1, std::basic_string<char32_t> * /*params*/) {
    char16_t temp_output_buffer[SD_STRING_ASSIGN_TEMP_BUFFER_BYTES / sizeof(char16_t) + 1];
    const char32_t* input_data = d1.data();
    const uint32_t input_length_char32_units = static_cast<uint32_t>(d1.length());

    sd::LongType required_bytes_for_utf16 = sd::unicode::offsetUtf32StringInUtf16(input_data, input_length_char32_units);

    if (required_bytes_for_utf16 > 0 && static_cast<size_t>(required_bytes_for_utf16) < sizeof(temp_output_buffer)) {
      void* end_ptr = sd::unicode::utf32to16Ptr(input_data, input_data + input_length_char32_units, temp_output_buffer);
      size_t char16_units_written = static_cast<char16_t*>(end_ptr) - temp_output_buffer;
      if (char16_units_written * sizeof(char16_t) == static_cast<size_t>(required_bytes_for_utf16)) {
        return std::basic_string<char16_t>(temp_output_buffer, char16_units_written);
      }
    }
    return std::basic_string<char16_t>();
  }
};

// --- Specialization: std::basic_string<char16_t> (UTF-16) -> std::basic_string<char32_t> (UTF-32) ---
template <>
class Assign<std::basic_string<char16_t>, std::basic_string<char32_t>> {
 public:
  static const bool requiresSpecial = false;
  static SD_HOST_DEVICE SD_INLINE  void execSpecial(const std::basic_string<char16_t> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char32_t> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char16_t> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char16_t> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char32_t> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char16_t> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char32_t> *reductionPointer, // Z is std::basic_string<char32_t>
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif

  SD_HOST_DEVICE SD_INLINE  static std::basic_string<char32_t>
  op(const std::basic_string<char16_t>& d1, std::basic_string<char16_t> * /*params*/) {
    char32_t temp_output_buffer[SD_STRING_ASSIGN_TEMP_BUFFER_BYTES / sizeof(char32_t) + 1];
    const char16_t* input_data = d1.data();
    const uint32_t input_length_char16_units = static_cast<uint32_t>(d1.length());

    sd::LongType required_bytes_for_utf32_data = sd::unicode::offsetUtf16StringInUtf32(input_data, input_length_char16_units);

    if (required_bytes_for_utf32_data > 0 && static_cast<size_t>(required_bytes_for_utf32_data) < sizeof(temp_output_buffer)) {
      void* end_ptr = sd::unicode::utf16to32Ptr(input_data, input_data + input_length_char16_units, temp_output_buffer);
      size_t char32_units_written = static_cast<char32_t*>(end_ptr) - temp_output_buffer;
      if (char32_units_written * sizeof(char32_t) == static_cast<size_t>(required_bytes_for_utf32_data)) {
        return std::basic_string<char32_t>(temp_output_buffer, char32_units_written);
      }
    }
    return std::basic_string<char32_t>();
  }
};


// --- Identity Specializations (Redundant in modernized version but included for completeness) ---
template <>
class Assign<std::basic_string<char16_t>, std::basic_string<char16_t>> {
 public:
  static const bool requiresSpecial = false;
  static SD_HOST_DEVICE SD_INLINE  void execSpecial(const std::basic_string<char16_t> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char16_t> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char16_t> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char16_t> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char16_t> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char16_t> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char16_t> *reductionPointer,
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif
  SD_HOST_DEVICE SD_INLINE   static std::basic_string<char16_t>
  op(const std::basic_string<char16_t>& d1, std::basic_string<char16_t> * /*params*/) {
    return d1;
  }
};

template <>
class Assign<std::basic_string<char32_t>, std::basic_string<char32_t>> {
 public:
  static const bool requiresSpecial = false;
  static  SD_HOST_DEVICE SD_INLINE  void execSpecial(const std::basic_string<char32_t> *dx, const sd::LongType *xShapeBuffer,
                          std::basic_string<char32_t> *result, const sd::LongType *resultShapeBuffer,
                          std::basic_string<char32_t> *extraParams, const sd::LongType *tadShapeInfo,
                          const sd::LongType *tadOffsets) {}
#ifdef __CUDACC__
  static SD_DEVICE void execSpecialCuda(const std::basic_string<char32_t> *dx, const sd::LongType *xShapeBuffer,
                                        std::basic_string<char32_t> *result, const sd::LongType *resultShapeBuffer,
                                        std::basic_string<char32_t> *extraParams,
                                        sd::LongType *allocationPointer,
                                        std::basic_string<char32_t> *reductionPointer,
                                        const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets) {}
#endif
  SD_HOST_DEVICE SD_INLINE  static std::basic_string<char32_t>
  op(const std::basic_string<char32_t>& d1, std::basic_string<char32_t> * /*params*/) {
    return d1;
  }
};

}  // namespace simdOps

#endif