<|start_filename|>TensorSharp/Expression/SVar.cs<|end_filename|>
﻿using System;

namespace TensorSharp.Expression
{
    public class SVar
    {
        private readonly SExpression expression;


        public SVar(SExpression expression)
        {
            this.expression = expression;
        }


        public float Evaluate()
        {
            return expression.Evaluate();
        }

        public SExpression Expression => expression;


        public static implicit operator SVar(float value) { return new SVar(new ConstScalarExpression(value)); }

        public static SVar operator -(SVar src) { return new SVar(new UnaryScalarExpression(src.expression, val => -val)); }

        public static SVar operator +(SVar lhs, SVar rhs) { return new SVar(new BinaryScalarExpression(lhs.expression, rhs.expression, (l, r) => l + r)); }
        public static SVar operator -(SVar lhs, SVar rhs) { return new SVar(new BinaryScalarExpression(lhs.expression, rhs.expression, (l, r) => l - r)); }
        public static SVar operator *(SVar lhs, SVar rhs) { return new SVar(new BinaryScalarExpression(lhs.expression, rhs.expression, (l, r) => l * r)); }
        public static SVar operator /(SVar lhs, SVar rhs) { return new SVar(new BinaryScalarExpression(lhs.expression, rhs.expression, (l, r) => l / r)); }
        public static SVar operator %(SVar lhs, SVar rhs) { return new SVar(new BinaryScalarExpression(lhs.expression, rhs.expression, (l, r) => l % r)); }


        public SVar Abs() { return new SVar(new UnaryScalarExpression(expression, val => Math.Abs(val))); }
        public SVar Sign() { return new SVar(new UnaryScalarExpression(expression, val => Math.Sign(val))); }

        public SVar Sqrt() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Sqrt(val))); }
        public SVar Exp() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Exp(val))); }
        public SVar Log() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Log(val))); }
        public SVar Floor() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Floor(val))); }
        public SVar Ceil() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Ceiling(val))); }
        public SVar Round() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Round(val))); }
        public SVar Trunc() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Truncate(val))); }


        public SVar Sin() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Sin(val))); }
        public SVar Cos() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Cos(val))); }
        public SVar Tan() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Tan(val))); }

        public SVar Asin() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Asin(val))); }
        public SVar Acos() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Acos(val))); }
        public SVar Atan() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Atan(val))); }

        public SVar Sinh() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Sinh(val))); }
        public SVar Cosh() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Cosh(val))); }
        public SVar Tanh() { return new SVar(new UnaryScalarExpression(expression, val => (float)Math.Tanh(val))); }


        public SVar Pow(SVar y) { return new SVar(new BinaryScalarExpression(expression, y.expression, (xVal, yVal) => (float)Math.Pow(xVal, yVal))); }
        public SVar Clamp(SVar min, SVar max) { return new SVar(new DelegateScalarExpression(() => ClampFloat(expression.Evaluate(), min.expression.Evaluate(), max.expression.Evaluate()))); }

        // public TVar Pow(TVar y) { return new TVar(new BinaryScalarTensorExpression(this.Expression, y.Expression, Ops.Tpow)); }


        public static SVar Atan2(SVar y, SVar x) { return new SVar(new DelegateScalarExpression(() => (float)Math.Atan2(y.Evaluate(), x.Evaluate()))); }
        public static SVar Lerp(SVar a, SVar b, SVar weight) { return new SVar(new DelegateScalarExpression(() => LerpFloat(a.Evaluate(), b.Evaluate(), weight.Evaluate()))); }


        private static float LerpFloat(float a, float b, float weight)
        {
            return a + weight * (b - a);
        }

        private static float ClampFloat(float value, float min, float max)
        {
            if (value < min)
            {
                return min;
            }
            else if (value > max)
            {
                return max;
            }
            else
            {
                return value;
            }
        }
    }
}


<|start_filename|>TensorSharp.CUDA/DeviceCode/VarStdKernels.cs<|end_filename|>
﻿using ManagedCuda;
using ManagedCuda.BasicTypes;
using ManagedCuda.VectorTypes;
using System;
using TensorSharp.Core;

namespace TensorSharp.CUDA.DeviceCode
{
    [Precompile]
    public class VarStdKernels : CudaCode
    {
        private static readonly string Code = @"
// Given the sum of values and the sum of squares, compute the variance or standard deviation.
template<bool flag, bool apply_sqrt>
__forceinline__ __device__ float computeVar(float sum, float sum2, unsigned row_size) {
  if (flag) {
    sum /= row_size;
    sum2 /= row_size;
    sum2 -= sum * sum;
    sum2 = (sum2 < 0 ? 0 : sum2);
  }
  else {
    sum /= row_size;
    sum2 /= row_size - 1;
    sum2 -= ((float)row_size) / ((float)(row_size - 1)) * sum * sum;
    sum2 = (sum2 < 0 ? 0 : sum2);
  }
  if (apply_sqrt)
    return sqrt(sum2);
  else
    return sum2;
}

/* Compute the variance (or standard deviation) along an outer dimension of a tensor.
 *
 * - num_orows is the size of the flattened outer dimensions;
 * - num_irows is the size of the flattened inner dimensions;
 * - row_size is the size of the dimension along which to compute the variance;
 * - if flag is set, normalize by `row_size` instead of `row_size - 1`
 * - if apply_sqrt is set, compute the standard deviation instead of variance
 *
 * The dimensions to the outside and inside of the specified dimension are considered as flattened.
 * Thread blocks with the same blockIdx.y process an 'outer row' (i.e. an element of the flattened
 * outer dimensions, which contains several 'inner rows').
 * Each thread processes a single inner row at a time.
 */
template<bool flag, bool apply_sqrt>
__device__ void kernel_varOuterDim(float* tgt, float* src_, unsigned num_orows, unsigned num_irows, unsigned row_size)
{
    for (unsigned orow = blockIdx.x; orow < num_orows; orow += gridDim.x)
    {
        for (unsigned irow = blockIdx.y * blockDim.x + threadIdx.x; irow < num_irows; irow += gridDim.y * blockDim.x)
        {
            float* src = src_ + orow * row_size * num_irows + irow;
            float sum = 0, sum2 = 0;

            for (unsigned col = 0; col < row_size; ++col)
            {
                float val = *src;
                sum += val;
                sum2 += val * val;

                src += num_irows;
            }

            tgt[orow * num_irows + irow] = computeVar<flag, apply_sqrt>(sum, sum2, row_size);
        }
    }
}





/* Compute the variance (or standard deviation) of the innermost dimension of a tensor.
 *
 * - num_rows is the size of the flattened outer dimensions;
 * - row_size is the size of the innermost dimension;
 * - if flag is set, normalize by `row_size` instead of `row_size - 1`
 * - if apply_sqrt is set, compute the standard deviation instead of variance
 *
 * The outer dimensions of the tensor are considered as a single dimension, i.e. the tensor is
 * considered as having 'num_rows' rows of size 'row_size'.
 * Each thread block processes one or more sets of contiguous rows (processing multiple rows
 * per thread block is quicker than processing a single row, especially for short rows).
 */
template<bool flag, bool apply_sqrt>
__device__ void kernel_varInnermostDim(float *tgt, float *src_, unsigned num_rows, unsigned row_size)
{
  __shared__ float ssum[32][16];
  __shared__ float ssum2[32][16];

  for (unsigned block_row = blockIdx.x * blockDim.y; block_row < num_rows; block_row += blockDim.y * gridDim.x) {
    unsigned row = block_row + threadIdx.y;
    float sum = 0, sum2 = 0;
    if (row < num_rows) {
      float *src = src_ + row * row_size;
      // Sequential reduction within a thread.
      for (unsigned col = threadIdx.x; col < row_size; col += blockDim.x) {
        float val = src[col];
        sum += val;
        sum2 += val * val;
      }
    }
    ssum[threadIdx.y][threadIdx.x] = sum;
    ssum2[threadIdx.y][threadIdx.x] = sum2;
    __syncthreads();

    // Reduce intermediate values to single value.
    for (unsigned s = 8; s > 1; s >>= 1) {
      if (row < num_rows && threadIdx.x < s) {
        ssum[threadIdx.y][threadIdx.x] += ssum[threadIdx.y][threadIdx.x + s];
        ssum2[threadIdx.y][threadIdx.x] += ssum2[threadIdx.y][threadIdx.x + s];
      }
      __syncthreads();
    }

    if (row < num_rows && threadIdx.x == 0) {
      sum = ssum[threadIdx.y][0] + ssum[threadIdx.y][1];
      sum2 = ssum2[threadIdx.y][0] + ssum2[threadIdx.y][1];
      tgt[row] = computeVar<flag, apply_sqrt>(sum, sum2, row_size);
    }
    __syncthreads();
  }
}

#define DECLARE_OUTERMOST(FLAG_VALUE, SQRT_VALUE) \
    __global__ void kernel_varOuterDim_##FLAG_VALUE##_##SQRT_VALUE(float* tgt, float* src_, unsigned num_orows, unsigned num_irows, unsigned row_size)\
    {\
        kernel_varOuterDim<FLAG_VALUE, SQRT_VALUE>(tgt, src_, num_orows, num_irows, row_size);\
    }

#define DECLARE_INNERMOST(FLAG_VALUE, SQRT_VALUE) \
    __global__ void kernel_varInnermostDim_##FLAG_VALUE##_##SQRT_VALUE(float *tgt, float *src_, unsigned num_rows, unsigned row_size)\
    {\
        kernel_varInnermostDim<FLAG_VALUE, SQRT_VALUE>(tgt, src_, num_rows, row_size);\
    }

extern ""C"" {
    DECLARE_OUTERMOST(true, true)
    DECLARE_OUTERMOST(true, false)
    DECLARE_OUTERMOST(false, true)
    DECLARE_OUTERMOST(false, false)

    DECLARE_INNERMOST(true, true)
    DECLARE_INNERMOST(true, false)
    DECLARE_INNERMOST(false, true)
    DECLARE_INNERMOST(false, false)
}

";

        public VarStdKernels()
            : base(GetFullCode(), "General", "ReduceApplyUtils", "ReduceBlock", "Reduce", "ReduceMacros", "Math")
        {
        }

        private static string GetFullCode()
        {
            return Code;
        }

        private static string GetMangledNameSuffix(bool normByN, bool applySqrt)
        {
            return string.Format("_{0}_{1}", normByN, applySqrt).ToLower();
        }


        private void VarOuterDim(TSCudaContext context, Tensor result, Tensor src, int dimension, bool normByN, bool applySqrt)
        {
            CudaContext cudaContext = context.CudaContextForTensor(src);

            int ndim = src.DimensionCount;
            long num_orows = 1;
            for (int dim = 0; dim < dimension; dim++)
            {
                num_orows *= src.Sizes[dim];
            }
            long row_size = src.Sizes[dimension];
            // Treat all inner dimensions (i.e. dim > dimension) as one.
            long num_irows = 1;
            for (int dim = dimension + 1; dim < ndim; dim++)
            {
                num_irows *= src.Sizes[dim];
            }

            dim3 threads = new dim3((uint)Math.Min(512, num_irows));
            int maxGridDim = 1024;
            dim3 grid = new dim3((uint)Math.Min(maxGridDim, num_orows), (uint)Math.Min(maxGridDim, ApplyUtils.CeilDiv(num_irows, threads.x)));

            CUdeviceptr resultPtr = CudaHelpers.GetBufferStart(result);
            CUdeviceptr srcPtr = CudaHelpers.GetBufferStart(src);
            string kernelName = "kernel_varOuterDim" + GetMangledNameSuffix(normByN, applySqrt);

            Invoke(context, cudaContext, kernelName, grid, threads, 0, CUstream.NullStream, resultPtr, srcPtr, num_orows, num_irows, row_size);
        }


        private void VarInnermostDim(TSCudaContext context, Tensor result, Tensor src, bool normByN, bool applySqrt)
        {
            CudaContext cudaContext = context.CudaContextForTensor(src);

            int ndim = src.DimensionCount;
            long num_rows = 1;
            for (int dim = 0; dim < ndim - 1; dim++)
            {
                num_rows *= src.Sizes[dim];
            }
            long row_size = src.Sizes[ndim - 1];

            // (Comment from cuTorch source): From limited testing, 16x32 seemed a good compromise for handling both long and short dimensions.
            dim3 threads = new dim3(16, 32);
            dim3 grid = new dim3((uint)Math.Min(1024, ApplyUtils.CeilDiv(num_rows, threads.y)));

            CUdeviceptr resultPtr = CudaHelpers.GetBufferStart(result);
            CUdeviceptr srcPtr = CudaHelpers.GetBufferStart(src);
            string kernelName = "kernel_varInnermostDim" + GetMangledNameSuffix(normByN, applySqrt);

            Invoke(context, cudaContext, kernelName, grid, threads, 0, CUstream.NullStream, resultPtr, srcPtr, num_rows, row_size);
        }


        private Tensor RunVarOp(Tensor result, Tensor src, int dimension, bool normByN, bool applySqrt)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            long[] requiredOutputSize = (long[])src.Sizes.Clone();
            requiredOutputSize[dimension] = 1;
            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, true, requiredOutputSize);

            if (dimension == src.DimensionCount - 1)
            {
                VarInnermostDim(context, writeTarget, src, normByN, applySqrt);
            }
            else
            {
                VarOuterDim(context, writeTarget, src, dimension, normByN, applySqrt);
            }

            return writeTarget;
        }

        public Tensor Var(Tensor result, Tensor src, int dimension, bool normByN)
        {
            return RunVarOp(result, src, dimension, normByN, false);
        }

        public Tensor Std(Tensor result, Tensor src, int dimension, bool normByN)
        {
            return RunVarOp(result, src, dimension, normByN, true);
        }

        private void Invoke(TSCudaContext context, CudaContext cudaContext, string kernelName, dim3 grid, dim3 block, uint smemSize, CUstream stream, params object[] args)
        {
            byte[] ptx = GetPtx(context.Compiler);
            CudaKernel kernel = context.KernelCache.Get(cudaContext, ptx, kernelName);
            kernel.GridDimensions = grid;
            kernel.BlockDimensions = block;
            kernel.DynamicSharedMemory = smemSize;
            kernel.RunAsync(stream, args);
        }
    }
}


<|start_filename|>TensorSharp/OpRegistryAttributes.cs<|end_filename|>
﻿using System;
using System.Collections.Generic;
using System.Reflection;

namespace TensorSharp
{
    [AttributeUsage(AttributeTargets.Class)]
    public class OpsClassAttribute : Attribute
    {
    }

    [AttributeUsage(AttributeTargets.Method)]
    public abstract class RegisterOp : Attribute
    {
        public string OpName { get; private set; }

        public RegisterOp(string opName)
        {
            OpName = opName;
        }

        public abstract void DoRegister(object instance, MethodInfo method, IEnumerable<OpConstraint> paramConstraints);
    }

    /// <summary>
    /// Register a method where the only constraint is that the argument counts match.
    /// </summary>
    [AttributeUsage(AttributeTargets.Method)]
    public class RegisterOpArgCount : RegisterOp
    {
        public RegisterOpArgCount(string opName) : base(opName)
        {
        }

        public override void DoRegister(object instance, MethodInfo method, IEnumerable<OpConstraint> paramConstraints)
        {
            List<OpConstraint> constraints = new List<OpConstraint>();
            constraints.AddRange(paramConstraints);
            constraints.Add(new ArgCountConstraint(method.GetParameters().Length));

            OpRegistry.Register(OpName, args => method.Invoke(instance, args), constraints);
        }
    }


    [AttributeUsage(AttributeTargets.Method)]
    public class RegisterOpStorageType : RegisterOp
    {
        private readonly Type storageType;

        public RegisterOpStorageType(string opName, Type storageType) : base(opName)
        {
            this.storageType = storageType;
        }

        public override void DoRegister(object instance, MethodInfo method, IEnumerable<OpConstraint> paramConstraints)
        {
            List<OpConstraint> constraints = new List<OpConstraint>();
            constraints.AddRange(paramConstraints);
            constraints.Add(new ArgCountConstraint(method.GetParameters().Length));

            ParameterInfo[] methodParams = method.GetParameters();
            for (int i = 0; i < methodParams.Length; ++i)
            {
                if (methodParams[i].ParameterType == typeof(Tensor))
                {
                    constraints.Add(new ArgStorageTypeConstraint(i, storageType));
                }
            }

            OpRegistry.Register(OpName, args => method.Invoke(instance, args), constraints);
        }
    }




    [AttributeUsage(AttributeTargets.Parameter)]
    public abstract class ArgConstraintAttribute : Attribute
    {
        public ArgConstraintAttribute()
        {
        }

        public abstract IEnumerable<OpConstraint> GetConstraints(ParameterInfo parameter, object instance);
    }

    [AttributeUsage(AttributeTargets.Parameter)]
    public class OpArgStorageType : ArgConstraintAttribute
    {
        private readonly Type storageType;

        public OpArgStorageType(Type storageType)
        {
            this.storageType = storageType;
        }

        public override IEnumerable<OpConstraint> GetConstraints(ParameterInfo parameter, object instance)
        {
            yield return new ArgStorageTypeConstraint(parameter.Position, storageType);
        }
    }
}


<|start_filename|>TensorSharp/Cpu/CpuRandom.cs<|end_filename|>
﻿using System;
using System.Reflection;

namespace TensorSharp.Cpu
{
    [OpsClass]
    public class CpuRandom
    {
        private static readonly Random seedGen = new Random();


        public CpuRandom()
        {
        }


        // allArgs should start with a null placeholder for the RNG object
        private static void InvokeWithRng(int? seed, MethodInfo method, params object[] allArgs)
        {
            if (!seed.HasValue)
            {
                seed = seedGen.Next();
            }

            NativeWrapper.CheckResult(CpuOpsNative.TS_NewRNG(out IntPtr rng));
            NativeWrapper.CheckResult(CpuOpsNative.TS_SetRNGSeed(rng, seed.Value));
            allArgs[0] = rng;
            NativeWrapper.InvokeTypeMatch(method, allArgs);
            NativeWrapper.CheckResult(CpuOpsNative.TS_DeleteRNG(rng));
        }

        private readonly MethodInfo uniform_func = NativeWrapper.GetMethod("TS_RandomUniform");
        [RegisterOpStorageType("random_uniform", typeof(CpuStorage))]
        public void Uniform(Tensor result, int? seed, float min, float max) { InvokeWithRng(seed, uniform_func, null, result, min, max); }

        private readonly MethodInfo normal_func = NativeWrapper.GetMethod("TS_RandomNormal");
        [RegisterOpStorageType("random_normal", typeof(CpuStorage))]
        public void Normal(Tensor result, int? seed, float mean, float stdv) { InvokeWithRng(seed, normal_func, null, result, mean, stdv); }

        private readonly MethodInfo exponential_func = NativeWrapper.GetMethod("TS_RandomExponential");
        [RegisterOpStorageType("random_exponential", typeof(CpuStorage))]
        public void Exponential(Tensor result, int? seed, float lambda) { InvokeWithRng(seed, exponential_func, null, result, lambda); }

        private readonly MethodInfo cauchy_func = NativeWrapper.GetMethod("TS_RandomCauchy");
        [RegisterOpStorageType("random_cauchy", typeof(CpuStorage))]
        public void Cauchy(Tensor result, int? seed, float median, float sigma) { InvokeWithRng(seed, cauchy_func, null, result, median, sigma); }

        private readonly MethodInfo log_normal_func = NativeWrapper.GetMethod("TS_RandomLogNormal");
        [RegisterOpStorageType("random_lognormal", typeof(CpuStorage))]
        public void LogNormal(Tensor result, int? seed, float mean, float stdv) { InvokeWithRng(seed, log_normal_func, null, result, mean, stdv); }

        private readonly MethodInfo geometric_func = NativeWrapper.GetMethod("TS_RandomGeometric");
        [RegisterOpStorageType("random_geometric", typeof(CpuStorage))]
        public void Geometric(Tensor result, int? seed, float p) { InvokeWithRng(seed, geometric_func, null, result, p); }

        private readonly MethodInfo bernoulli_func = NativeWrapper.GetMethod("TS_RandomBernoulli");
        [RegisterOpStorageType("random_bernoulli", typeof(CpuStorage))]
        public void Bernoulli(Tensor result, int? seed, float p) { InvokeWithRng(seed, bernoulli_func, null, result, p); }
    }
}


<|start_filename|>TensorSharp.CUDA/PrecompileAttribute.cs<|end_filename|>
﻿using System;
using System.Reflection;
using TensorSharp.CUDA.RuntimeCompiler;

namespace TensorSharp.CUDA
{
    [AttributeUsage(AttributeTargets.Class, AllowMultiple = false, Inherited = false)]
    public class PrecompileAttribute : Attribute
    {
        public PrecompileAttribute()
        {
        }
    }

    public interface IPrecompilable
    {
        void Precompile(CudaCompiler compiler);
    }

    public static class PrecompileHelper
    {
        public static void PrecompileAllFields(object instance, CudaCompiler compiler)
        {
            Type type = instance.GetType();

            foreach (FieldInfo field in type.GetFields())
            {
                if (typeof(IPrecompilable).IsAssignableFrom(field.FieldType))
                {
                    IPrecompilable precompilableField = (IPrecompilable)field.GetValue(instance);
                    Console.WriteLine("Compiling field " + field.Name);
                    precompilableField.Precompile(compiler);
                }
            }
        }
    }
}


<|start_filename|>TensorSharp/Expression/SExpression.cs<|end_filename|>
﻿using System;

namespace TensorSharp.Expression
{
    public abstract class SExpression
    {
        public abstract float Evaluate();
    }


    public class ConstScalarExpression : SExpression
    {
        private readonly float value;

        public ConstScalarExpression(float value)
        {
            this.value = value;
        }

        public override float Evaluate()
        {
            return value;
        }
    }

    public class DelegateScalarExpression : SExpression
    {
        private readonly Func<float> evaluate;

        public DelegateScalarExpression(Func<float> evaluate)
        {
            this.evaluate = evaluate;
        }

        public override float Evaluate()
        {
            return evaluate();
        }
    }

    public class UnaryScalarExpression : SExpression
    {
        private readonly SExpression src;
        private readonly Func<float, float> evaluate;


        public UnaryScalarExpression(SExpression src, Func<float, float> evaluate)
        {
            this.src = src;
            this.evaluate = evaluate;
        }

        public override float Evaluate()
        {
            return evaluate(src.Evaluate());
        }
    }

    public class BinaryScalarExpression : SExpression
    {
        private readonly SExpression left;
        private readonly SExpression right;
        private readonly Func<float, float, float> evaluate;


        public BinaryScalarExpression(SExpression left, SExpression right, Func<float, float, float> evaluate)
        {
            this.left = left;
            this.right = right;
            this.evaluate = evaluate;
        }

        public override float Evaluate()
        {
            return evaluate(left.Evaluate(), right.Evaluate());
        }
    }
}


<|start_filename|>TensorSharp/OpRegistry.cs<|end_filename|>
﻿using System;
using System.Collections.Generic;
using System.Linq;
using System.Reflection;

namespace TensorSharp
{
    public delegate object OpHandler(object[] args);

    public static class OpRegistry
    {
        private class OpInstance
        {
            public OpHandler handler;
            public IEnumerable<OpConstraint> constraints;
        }

        private static readonly Dictionary<string, List<OpInstance>> opInstances = new Dictionary<string, List<OpInstance>>();
        // Remember which assemblies have been registered to avoid accidental double-registering
        private static readonly HashSet<Assembly> registeredAssemblies = new HashSet<Assembly>();

        static OpRegistry()
        {
            // Register CPU ops from this assembly
            RegisterAssembly(Assembly.GetExecutingAssembly());
        }

        public static void Register(string opName, OpHandler handler, IEnumerable<OpConstraint> constraints)
        {
            OpInstance newInstance = new OpInstance() { handler = handler, constraints = constraints };

            if (opInstances.TryGetValue(opName, out List<OpInstance> instanceList))
            {
                instanceList.Add(newInstance);
            }
            else
            {
                instanceList = new List<OpInstance>
                {
                    newInstance
                };
                opInstances.Add(opName, instanceList);
            }
        }

        public static object Invoke(string opName, params object[] args)
        {
            if (opInstances.TryGetValue(opName, out List<OpInstance> instanceList))
            {
                foreach (OpInstance instance in instanceList)
                {
                    if (instance.constraints.All(x => x.SatisfiedFor(args)))
                    {
                        return instance.handler.Invoke(args);
                    }
                }

                throw new ApplicationException("None of the registered handlers match the arguments for " + opName);
            }
            else
            {
                throw new ApplicationException("No handlers have been registered for op " + opName);
            }
        }

        public static void RegisterAssembly(Assembly assembly)
        {
            if (!registeredAssemblies.Contains(assembly))
            {
                registeredAssemblies.Add(assembly);

                IEnumerable<Type> types = assembly.TypesWithAttribute<OpsClassAttribute>(false)
                    .Select(x => x.Item1);

                foreach (Type type in types)
                {
                    object instance = Activator.CreateInstance(type);

                    IEnumerable<Tuple<MethodInfo, IEnumerable<RegisterOp>>> methods = type.MethodsWithAttribute<RegisterOp>(false);
                    foreach (Tuple<MethodInfo, IEnumerable<RegisterOp>> method in methods)
                    {
                        IEnumerable<OpConstraint> paramConstraints = GetParameterConstraints(method.Item1, instance);
                        foreach (RegisterOp attribute in method.Item2)
                        {
                            attribute.DoRegister(instance, method.Item1, paramConstraints);
                        }
                    }
                }
            }
        }

        private static IEnumerable<OpConstraint> GetParameterConstraints(MethodInfo method, object instance)
        {
            IEnumerable<OpConstraint> result = Enumerable.Empty<OpConstraint>();
            foreach (Tuple<ParameterInfo, IEnumerable<ArgConstraintAttribute>> parameter in method.ParametersWithAttribute<ArgConstraintAttribute>(false))
            {
                foreach (ArgConstraintAttribute attribute in parameter.Item2)
                {
                    result = Enumerable.Concat(result, attribute.GetConstraints(parameter.Item1, instance));
                }
            }

            return result;
        }
    }
}


<|start_filename|>Seq2SeqSharp/LearningRate/DecayLearningRate.cs<|end_filename|>
﻿using AdvUtils;
using System;

namespace Seq2SeqSharp
{
    public class DecayLearningRate : ILearningRate
    {
        private readonly float m_startLearningRate = 0.001f;
        private int m_weightsUpdateCount = 0;
        private readonly int m_warmupSteps = 8000;

        public DecayLearningRate(float startLearningRate, int warmupSteps, int weightsUpdatesCount)
        {
            Logger.WriteLine($"Creating decay learning rate. StartLearningRate = '{startLearningRate}', WarmupSteps = '{warmupSteps}', WeightsUpdatesCount = '{weightsUpdatesCount}'");
            m_startLearningRate = startLearningRate;
            m_warmupSteps = warmupSteps;
            m_weightsUpdateCount = weightsUpdatesCount;
        }

        public float GetCurrentLearningRate()
        {
            m_weightsUpdateCount++;
            float lr = m_startLearningRate * (float)(Math.Min(Math.Pow(m_weightsUpdateCount, -0.5), Math.Pow(m_warmupSteps, -1.5) * m_weightsUpdateCount) / Math.Pow(m_warmupSteps, -0.5));
            return lr;
        }
    }
}


<|start_filename|>Seq2SeqSharp/Utils/BeamSearch.cs<|end_filename|>
﻿using AdvUtils;
using Seq2SeqSharp.Tools;
using System.Collections.Generic;
using System.Linq;

namespace Seq2SeqSharp
{
    public class BeamSearchStatus
    {
        public List<int> OutputIds;
        public float Score;

        public List<IWeightTensor> HTs;
        public List<IWeightTensor> CTs;

        public BeamSearchStatus()
        {
            OutputIds = new List<int>();
            HTs = new List<IWeightTensor>();
            CTs = new List<IWeightTensor>();

            Score = 1.0f;
        }
    }

    public class BeamSearch
    {
        public static List<BeamSearchStatus> GetTopNBSS(List<BeamSearchStatus> bssList, int topN)
        {
            FixedSizePriorityQueue<ComparableItem<BeamSearchStatus>> q = new FixedSizePriorityQueue<ComparableItem<BeamSearchStatus>>(topN, new ComparableItemComparer<BeamSearchStatus>(false));

            for (int i = 0; i < bssList.Count; i++)
            {
                q.Enqueue(new ComparableItem<BeamSearchStatus>(bssList[i].Score, bssList[i]));
            }

            return q.Select(x => x.Value).ToList();
        }
    }
}


<|start_filename|>TensorSharp/TensorSerialization.cs<|end_filename|>
﻿using System;
using System.IO;
using System.Runtime.InteropServices;

namespace TensorSharp
{
    public static class TensorSerialization
    {
        public static void Serialize(Tensor tensor, Stream stream)
        {
            using (Tensor src = Ops.AsContiguous(tensor))
            {
                // Note: don't dispose writer - it does not own the stream's lifetime
                BinaryWriter writer = new System.IO.BinaryWriter(stream);

                // Can infer strides - src is contiguous
                writer.Write(tensor.DimensionCount); // int32
                writer.Write((int)tensor.ElementType);
                for (int i = 0; i < tensor.DimensionCount; ++i)
                {
                    writer.Write(tensor.Sizes[i]);
                }

                long byteCount = src.ElementType.Size() * tensor.ElementCount();
                writer.Write(byteCount);
                WriteBytes(writer, src.Storage, src.StorageOffset, byteCount);

                writer.Flush();
            }
        }

        public static Tensor Deserialize(IAllocator allocator, Stream stream)
        {
            // Note: don't dispose reader - it does not own the stream's lifetime
            BinaryReader reader = new BinaryReader(stream);

            int dimCount = reader.ReadInt32();
            DType elementType = (DType)reader.ReadInt32();
            long[] sizes = new long[dimCount];
            for (int i = 0; i < dimCount; ++i)
            {
                sizes[i] = reader.ReadInt64();
            }

            long byteCount = reader.ReadInt64();
            Tensor result = new Tensor(allocator, elementType, sizes);

            ReadBytes(reader, result.Storage, result.StorageOffset, byteCount);

            return result;
        }

        private static void WriteBytes(BinaryWriter writer, Storage storage, long startIndex, long byteCount)
        {
            byte[] buffer = new byte[4096];
            GCHandle bufferHandle = GCHandle.Alloc(buffer, GCHandleType.Pinned);
            try
            {
                long curStart = startIndex;
                long afterLastByte = startIndex + byteCount;
                while (curStart < afterLastByte)
                {
                    int length = (int)Math.Min(buffer.Length, afterLastByte - curStart);
                    storage.CopyFromStorage(bufferHandle.AddrOfPinnedObject(), curStart, length);
                    writer.Write(buffer, 0, length);
                    curStart += length;
                }
            }
            finally
            {
                bufferHandle.Free();
            }
        }

        private static void ReadBytes(BinaryReader reader, Storage storage, long startIndex, long byteCount)
        {
            byte[] buffer = new byte[4096];
            GCHandle bufferHandle = GCHandle.Alloc(buffer, GCHandleType.Pinned);
            try
            {
                long curStart = startIndex;
                long afterLastByte = startIndex + byteCount;
                while (curStart < afterLastByte)
                {
                    int length = (int)Math.Min(buffer.Length, afterLastByte - curStart);
                    reader.Read(buffer, 0, length);
                    storage.CopyToStorage(curStart, bufferHandle.AddrOfPinnedObject(), length);
                    curStart += length;
                }
            }
            finally
            {
                bufferHandle.Free();
            }
        }
    }
}


<|start_filename|>TensorSharp.CUDA/KernelOps/ReductionOp.cs<|end_filename|>
﻿using ManagedCuda;
using ManagedCuda.VectorTypes;
using System;
using TensorSharp.Core;
using TensorSharp.CUDA.DeviceCode;
using TensorSharp.CUDA.RuntimeCompiler;

namespace TensorSharp.CUDA.KernelOps
{
    public static class ReductionOp
    {
        public static Tensor Invoke(CudaReduceKernels reduceKernels, string kernelName, float init, ReduceInitType initType, Tensor result, Tensor src, int dim, object extraArg = null)
        {
            if (src.DimensionCount == 0)
            {
                return result;
            }

            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            long[] requiredOutputSize = (long[])src.Sizes.Clone();
            requiredOutputSize[dim] = 1;
            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, requiredOutputSize);
            ThrowIfAnyTensorInvalid(writeTarget, src);

            long inElements = src.ElementCount();
            long reductionSize = src.Sizes[dim];
            long reductionStride = src.Strides[dim];
            long outElements = inElements / reductionSize;
            bool contigReduction = reductionStride == 1;


            // We must make sure that when the tensor is passed to the kernel, src.Sizes[dim] is set to 1
            // This includes for the purposes of determining which tensor specializations to use (changing
            // the dimension size to 1 may make the tensor non-contiguous
            long[] newSizes = (long[])src.Sizes.Clone();
            newSizes[dim] = 1;
            Tensor srcSlim = new Tensor(newSizes, src.Strides, src.Storage, src.StorageOffset);

            ApplySpecialization config = new ApplySpecialization(writeTarget, srcSlim);
            object totalSlices = config.Use32BitIndices ? (uint)outElements : (ulong)outElements;
            object reductionSizeTyped = config.Use32BitIndices ? (uint)reductionSize : (ulong)reductionSize;
            object reductionStrideTyped = config.Use32BitIndices ? (uint)reductionStride : (ulong)reductionStride;
            object initValueTyped = ReduceInitConverter.GetInitValue(init, initType, src.ElementType);

            byte[] ptx = reduceKernels.GetPtx(context.Compiler);

            if (contigReduction)
            {
                dim3 block = GetContigReduceBlock(cudaContext, outElements, reductionSize);
                dim3 grid = GetContigReduceGrid(outElements);
                uint smemSize = (uint)src.ElementType.Size() * block.x;

                string fullName = "contig_" + PermutationGenerator.GetMangledName(kernelName, config);
                if (extraArg == null)
                {
                    InvokeReduce(context, cudaContext, ptx, fullName, grid, block, smemSize, config, writeTarget, srcSlim, reductionSizeTyped, totalSlices, initValueTyped);
                }
                else
                {
                    InvokeReduce(context, cudaContext, ptx, fullName, grid, block, smemSize, config, writeTarget, srcSlim, reductionSizeTyped, totalSlices, initValueTyped, extraArg);
                }
            }
            else
            {
                CudaDeviceProperties deviceProps = context.DeviceInfoForContext(cudaContext);
                dim3 block = GetNonContigReduceBlock(deviceProps);
                dim3 grid = GetNoncontigReduceGrid(deviceProps, outElements);
                uint smemSize = 0;

                string fullName = "noncontig_" + PermutationGenerator.GetMangledName(kernelName, config);
                if (extraArg == null)
                {
                    InvokeReduce(context, cudaContext, ptx, fullName, grid, block, smemSize, config, writeTarget, srcSlim, reductionStrideTyped, reductionSizeTyped, totalSlices, initValueTyped);
                }
                else
                {
                    InvokeReduce(context, cudaContext, ptx, fullName, grid, block, smemSize, config, writeTarget, srcSlim, reductionStrideTyped, reductionSizeTyped, totalSlices, initValueTyped, extraArg);
                }
            }

            return writeTarget;
        }

        public static void InvokeReduce(TSCudaContext context, CudaContext cudaContext, byte[] ptx, string kernelName, dim3 grid, dim3 block, uint smemSize, ApplySpecialization spec, params object[] args)
        {
            ConvertTensorArgs.Convert(cudaContext, spec.Use32BitIndices, args);

            CudaKernel kernel = context.KernelCache.Get(cudaContext, ptx, kernelName);

            kernel.GridDimensions = grid;
            kernel.BlockDimensions = block;
            kernel.DynamicSharedMemory = smemSize;

            kernel.Run(args);

        }

        private static int GetNonContigReduceBlockSize(CudaDeviceProperties deviceProps)
        {
            return Math.Min(DeviceCode.Headers.Reduce.NonContigReduceBlockSize, (int)deviceProps.MaxBlockDim.x);
        }

        private static dim3 GetNonContigReduceBlock(CudaDeviceProperties deviceProps)
        {
            return new dim3(GetNonContigReduceBlockSize(deviceProps));
        }

        private static dim3 GetContigReduceBlock(CudaContext cudaContext, long numSlices, long reductionSize)
        {
            // If the number of slices is low but the reduction dimension size
            // is high, then we should increase block size for greater parallelism.
            // Aim for at least 32 warps per SM (assume 15 SMs; don't bother
            // inquiring the real number for now).
            int smCount = 15;
            int maxWarps = 4; // better occupancy if many blocks are around
                              // For numSlices > smCount * 8, there are > 32 warps active per SM.
            if (numSlices < smCount * 8)
            {
                maxWarps = 8;
                if (numSlices < smCount * 4)
                {
                    maxWarps = 16;
                    if (numSlices < smCount * 2)
                    {
                        maxWarps = 32;
                    }
                }
            }

            // Scale up block size based on the reduction dimension size
            long warpsInReductionSize = ApplyUtils.CeilDiv(reductionSize, 32);
            int numWarps =
              warpsInReductionSize > maxWarps ? maxWarps : (int)warpsInReductionSize;

            int targetSize = numWarps * 32;

            targetSize = Math.Min(targetSize, (int)cudaContext.GetDeviceInfo().MaxBlockDim.x);
            return new dim3(targetSize);
        }

        private static dim3 GetNoncontigReduceGrid(CudaDeviceProperties deviceProps, long elements)
        {
            // One output point per thread
            return GridFromTiles(ApplyUtils.CeilDiv(elements, GetNonContigReduceBlockSize(deviceProps)));
        }

        private static dim3 GetContigReduceGrid(long elements)
        {
            // One output point per block
            return GridFromTiles(elements);
        }


        private const long MaxGridSize = 65535;

        private static dim3 GridFromTiles(long gridTiles)
        {
            if (gridTiles > MaxGridSize * MaxGridSize * MaxGridSize)
            {
                throw new ArgumentException("gridTiles exceeds the maximum allowed tile count", "gridTiles");
            }

            long gridX = gridTiles > MaxGridSize ? MaxGridSize : gridTiles;
            long gridY = 1;
            long gridZ = 1;

            if (gridTiles > MaxGridSize)
            {
                gridTiles = ApplyUtils.CeilDiv(gridTiles, MaxGridSize);
                gridY = gridTiles > MaxGridSize ? MaxGridSize : gridTiles;

                if (gridTiles > MaxGridSize)
                {
                    gridTiles = ApplyUtils.CeilDiv(gridTiles, MaxGridSize);
                    gridZ = gridTiles > MaxGridSize ? MaxGridSize : gridTiles;
                }
            }

            return new dim3((uint)gridX, (uint)gridY, (uint)gridZ);
        }



        private static void ThrowIfAnyTensorInvalid(params Tensor[] args)
        {
            foreach (Tensor tensor in args)
            {
                if (tensor.DimensionCount > TSCudaContext.MaxDims)
                {
                    throw new InvalidOperationException("Tensors with dimension count > " + TSCudaContext.MaxDims + " are not supported");
                }
            }
        }
    }
}


<|start_filename|>TensorSharp.CUDA/RuntimeCompiler/CudaIncludeAttribute.cs<|end_filename|>
﻿using System;

namespace TensorSharp.CUDA.RuntimeCompiler
{
    [AttributeUsage(AttributeTargets.Class, AllowMultiple = true, Inherited = false)]
    public class CudaIncludeAttribute : Attribute
    {
        public string FieldName { get; private set; }
        public string IncludeName { get; private set; }

        public CudaIncludeAttribute(string fieldName, string includeName)
        {
            FieldName = fieldName;
            IncludeName = includeName;
        }
    }
}


<|start_filename|>TensorSharp.CUDA/Util/ObjectPool.cs<|end_filename|>
﻿using System;
using System.Collections.Generic;

namespace TensorSharp.CUDA.Util
{
    public class PooledObject<T> : IDisposable
    {
        private readonly Action<PooledObject<T>> onDispose;
        private readonly T value;

        private bool disposed = false;

        public PooledObject(T value, Action<PooledObject<T>> onDispose)
        {
            if (onDispose == null)
            {
                throw new ArgumentNullException("onDispose");
            }

            this.onDispose = onDispose;
            this.value = value;
        }

        public T Value
        {
            get
            {
                if (disposed)
                {
                    throw new ObjectDisposedException(ToString());
                }

                return value;
            }
        }

        public void Dispose()
        {
            if (!disposed)
            {
                onDispose(this);
                disposed = true;
            }
            else
            {
                throw new ObjectDisposedException(ToString());
            }
        }
    }

    public class ObjectPool<T> : IDisposable
    {
        private readonly Func<T> constructor;
        private readonly Action<T> destructor;
        private readonly Stack<T> freeList = new Stack<T>();
        private bool disposed = false;


        public ObjectPool(int initialSize, Func<T> constructor, Action<T> destructor)
        {
            if (constructor == null)
            {
                throw new ArgumentNullException("constructor");
            }

            if (destructor == null)
            {
                throw new ArgumentNullException("destructor");
            }

            this.constructor = constructor;
            this.destructor = destructor;

            for (int i = 0; i < initialSize; ++i)
            {
                freeList.Push(constructor());
            }
        }

        public void Dispose()
        {
            if (!disposed)
            {
                disposed = true;
                foreach (T item in freeList)
                {
                    destructor(item);
                }
                freeList.Clear();
            }
        }

        public PooledObject<T> Get()
        {
            T value = freeList.Count > 0 ? freeList.Pop() : constructor();
            return new PooledObject<T>(value, Release);
        }

        private void Release(PooledObject<T> handle)
        {
            freeList.Push(handle.Value);
        }
    }
}


<|start_filename|>TensorSharp.CUDA/DeviceCode/CudaCode.cs<|end_filename|>
﻿using TensorSharp.CUDA.RuntimeCompiler;

namespace TensorSharp.CUDA.DeviceCode
{
    public abstract class CudaCode : IPrecompilable
    {
        private readonly string code;
        private readonly string[] requiredHeaders;
        private byte[] ptx = null;

        protected CudaCode(string code, params string[] requiredHeaders)
        {
            this.code = code;
            this.requiredHeaders = requiredHeaders;
        }

        public byte[] GetPtx(CudaCompiler compiler)
        {
            if (ptx == null)
            {
                Precompile(compiler);
            }
            return ptx;
        }

        public void Precompile(CudaCompiler compiler)
        {
            ptx = compiler.CompileToPtx(code, requiredHeaders);
        }
    }

}


<|start_filename|>TensorSharp/ReflectionExtensions.cs<|end_filename|>
﻿using System;
using System.Collections.Generic;
using System.Linq;
using System.Reflection;

namespace TensorSharp
{
    public static class AssemblyExtensions
    {
        public static IEnumerable<Tuple<Type, IEnumerable<T>>> TypesWithAttribute<T>(this Assembly assembly, bool inherit)
        {
            foreach (Type type in assembly.GetTypes())
            {
                object[] attributes = type.GetCustomAttributes(typeof(T), inherit);
                if (attributes.Any())
                {
                    yield return Tuple.Create(type, attributes.Cast<T>());
                }
            }
        }
    }

    public static class TypeExtensions
    {
        public static IEnumerable<Tuple<MethodInfo, IEnumerable<T>>> MethodsWithAttribute<T>(this Type type, bool inherit)
        {
            foreach (MethodInfo method in type.GetMethods())
            {
                object[] attributes = method.GetCustomAttributes(typeof(T), inherit);
                if (attributes.Any())
                {
                    yield return Tuple.Create(method, attributes.Cast<T>());
                }
            }
        }
    }

    public static class MethodExtensions
    {
        public static IEnumerable<Tuple<ParameterInfo, IEnumerable<T>>> ParametersWithAttribute<T>(this MethodInfo method, bool inherit)
        {
            foreach (ParameterInfo paramter in method.GetParameters())
            {
                object[] attributes = paramter.GetCustomAttributes(typeof(T), inherit);
                if (attributes.Any())
                {
                    yield return Tuple.Create(paramter, attributes.Cast<T>());
                }
            }
        }
    }
}


<|start_filename|>TensorSharp.CUDA/RuntimeCompiler/KernelConfig.cs<|end_filename|>
﻿using System.Collections.Generic;
using System.Text;

namespace TensorSharp.CUDA.RuntimeCompiler
{
    public class KernelConfig
    {
        private readonly SortedDictionary<string, string> values = new SortedDictionary<string, string>();


        public KernelConfig()
        {
        }

        public IEnumerable<string> Keys => values.Keys;

        public IEnumerable<KeyValuePair<string, string>> AllValues()
        {
            return values;
        }

        public override bool Equals(object obj)
        {
            KernelConfig o = obj as KernelConfig;
            if (o == null)
            {
                return false;
            }

            if (values.Count != o.values.Count)
            {
                return false;
            }

            foreach (KeyValuePair<string, string> kvp in values)
            {
                if (values.TryGetValue(kvp.Key, out string oValue))
                {
                    if (!kvp.Value.Equals(oValue))
                    {
                        return false;
                    }
                }
                else
                {
                    return false;
                }
            }

            return true;
        }

        public override int GetHashCode()
        {
            int result = 0;
            foreach (KeyValuePair<string, string> kvp in values)
            {
                result ^= kvp.Key.GetHashCode();
                result ^= kvp.Value.GetHashCode();
            }
            return result;
        }

        public bool ContainsKey(string name)
        {
            return values.ContainsKey(name);
        }

        public void Set(string name, string value)
        {
            values[name] = value;
        }

        public string ApplyToTemplate(string templateCode)
        {
            StringBuilder fullCode = new StringBuilder();
            foreach (KeyValuePair<string, string> item in values)
            {
                fullCode.AppendFormat("#define {0} {1}\n", item.Key, item.Value);
            }
            fullCode.AppendLine(templateCode);
            return fullCode.ToString();
        }
    }
}


<|start_filename|>Seq2SeqSharp/Tools/ComparableItem.cs<|end_filename|>
﻿using System;
using System.Collections.Generic;

namespace Seq2SeqSharp.Tools
{

    public class ComparableItem<T>
    {
        public float Score { get; }
        public T Value { get; }

        public ComparableItem(float score, T value)
        {
            Score = score;
            Value = value;
        }
    }

    public class ComparableItemComparer<T> : IComparer<ComparableItem<T>>
    {
        public ComparableItemComparer(bool fAscending)
        {
            m_fAscending = fAscending;
        }

        public int Compare(ComparableItem<T> x, ComparableItem<T> y)
        {
            int iSign = Math.Sign(x.Score - y.Score);
            if (!m_fAscending)
            {
                iSign = -iSign;
            }

            return iSign;
        }

        protected bool m_fAscending;
    }
}


<|start_filename|>TensorSharp/Cpu/CpuNativeHelpers.cs<|end_filename|>
﻿using System;

namespace TensorSharp.Cpu
{
    public static class CpuNativeHelpers
    {
        public static IntPtr GetBufferStart(Tensor tensor)
        {
            IntPtr buffer = ((CpuStorage)tensor.Storage).buffer;
            return PtrAdd(buffer, tensor.StorageOffset * tensor.ElementType.Size());
        }

        private static IntPtr PtrAdd(IntPtr ptr, long offset)
        {
            return new IntPtr(ptr.ToInt64() + offset);
        }

    }
}


<|start_filename|>TensorSharp.CUDA/ApplyUtils.cs<|end_filename|>
﻿using ManagedCuda;
using ManagedCuda.BasicTypes;
using ManagedCuda.VectorTypes;
using System;
using System.Linq;

namespace TensorSharp.CUDA
{
    public struct TensorInfo
    {
        public long[] sizes;
        public long[] strides;
        public CUdeviceptr buffer;
    }

    public static class ApplyUtils
    {
        public const int ApplyThreadsPerBlock = 32 * 16;


        public static dim3 GetApplyBlock()
        {
            return new dim3(ApplyThreadsPerBlock);
        }

        // returns Ceil(x / y)
        public static long CeilDiv(long x, long y)
        {
            return (x + y - 1) / y;
        }

        public static dim3 GetApplyGrid(CudaDeviceProperties deviceInfo, long totalElements)
        {
            int smCount = deviceInfo.MultiProcessorCount;

            // Rationale for grid size - from cuTorch source code:
            // 16 warps per block * 4 per SM gives 64 warps per SM at maximum,
            // which seems to be a good sweetspot for latency hiding
            int maxSize = 4 * smCount;
            long targetSize = CeilDiv(totalElements, ApplyThreadsPerBlock);
            return new dim3((uint)Math.Min(targetSize, maxSize));
        }



        public static bool CanUse32BitIndexMath(Tensor tensor)
        {
            long elements = tensor.ElementCount();
            if (elements >= uint.MaxValue)
            {
                return false;
            }

            long offset = 0;
            long linearId = elements - 1;

            for (int i = tensor.DimensionCount - 1; i >= 0; --i)
            {
                long curDimIndex = linearId % tensor.Sizes[i];
                long curDimOffset = curDimIndex * tensor.Strides[i];
                offset += curDimOffset;
                linearId /= tensor.Sizes[i];
            }

            if (offset >= uint.MaxValue)
            {
                return false;
            }

            return true;
        }

        private struct SizeAndStride : IComparable
        {
            public long size; public long stride;

            public int CompareTo(object obj)
            {
                if (!(obj is SizeAndStride))
                {
                    throw new InvalidOperationException();
                }

                SizeAndStride o = (SizeAndStride)obj;
                return -stride.CompareTo(o.stride); // negated because we require descending order
            }
        }

        public static bool HasOverlappingIndices(Tensor tensor)
        {
            // In this function, we don't care about permutations of the
            // size/stride arrays (transpositions).
            // We order the size/stride arrays by stride, skipping dimensions of
            // size 1. Strides of dimensions of size 1 don't matter, since there
            // is only one addressing point in them.
            // In this reordered view, the tensor is contiguous if
            // stride[dim] == size[dim + 1] * stride[dim + 1] for all `dim`.
            // The tensor has holes if
            // stride[dim] > size[dim + 1] * stride[dim + 1] for one or more
            // `dim`.
            // The tensor has overlaps if
            // stride[dim] < size[dim + 1] * stride[dim + 1] for one or more
            // `dim`, or the innermost stride is 0.

            // Extract size/stride arrays; only consider size >1 dims.
            SizeAndStride[] info = Enumerable.Range(0, tensor.DimensionCount)
                .Select(x => new SizeAndStride() { size = tensor.Sizes[x], stride = tensor.Strides[x] })
                .Where(x => x.size > 1)
                .ToArray();

            if (info.Length == 0)
            {
                return false; // no overlap
            }

            // Descending order (innermost dimension in sorted view is at last index)
            Array.Sort(info);

            // Base case: innermost dimension must have stride >= 1
            if (info[info.Length - 1].stride < 1)
            {
                return true;
            }

            // Subsequent dimensions, if any
            for (int i = info.Length - 2; i >= 0; --i)
            {
                if (info[i].stride < info[i + 1].size * info[i + 1].stride)
                {
                    // There are overlaps
                    return true;
                }
            }

            // Tensor has holes or is contiguous
            return false;
        }



        private static int GetInnermostNon1Dim(long[] sizes, int excludeDim)
        {
            for (int i = sizes.Length - 1; i >= 0; --i)
            {
                if (i == excludeDim)
                {
                    return i;
                }

                if (sizes[i] != 1)
                {
                    return i;
                }
            }

            return -1;
        }

        //TODO this is not actually used at the moment. It probably should be.
        // excludeDim may be -1 to not exclude any dimension
        public static int CollapseDims(Tensor tensor, int excludeDim, out TensorInfo info)
        {
            info.buffer = CudaHelpers.GetBufferStart(tensor);
            int firstNonOneDim = GetInnermostNon1Dim(tensor.Sizes, excludeDim);

            // If all dims are size 1 (ie. tensor contains 1 element)
            if (firstNonOneDim == -1)
            {
                if (excludeDim != -1)
                {
                    throw new ArgumentException("excludeDim must equal -1 if all dims are of size 1", "excludeDim");
                }

                info.sizes = new long[] { 1 };
                info.strides = new long[] { 1 };
                return 0;
            }


            // Count the number of successive dimensions that can be collapsed, from
            // innermost to outermost.
            int numCollapsed = 0;

            // Skip the leading size 1 dims
            numCollapsed += tensor.DimensionCount - 1 - firstNonOneDim;

            // We perform one pass through to determine how many dimensions we
            // can collapse, before calculating the actual size of the collapsed
            // dimensions.
            // size/strideInner are the size/strides of the previous inner
            // non-collapsible dim we encounter.
            long sizeInner = tensor.Sizes[firstNonOneDim];
            long strideInner = tensor.Strides[firstNonOneDim];

            for (int i = firstNonOneDim - 1; i >= 0; --i)
            {
                long sizeOuter = tensor.Sizes[i];
                long strideOuter = tensor.Strides[i];

                // Don't collapse this dimension if we want to exclude it from
                // collapsing.
                // Since this code is attempting to collapse a subsequent
                // dimension (i) with the preceding dimension (i + 1), we can only
                // perform collapsing if the preceding dimension can be collapsed
                // (i.e., not excludeDim)
                if ((excludeDim != i) && (excludeDim != i + 1))
                {
                    // The next outermost dimension can be skipped if size 1
                    if (sizeOuter == 1)
                    {
                        ++numCollapsed;
                        continue;
                    }

                    // If the next outermost dimension is contiguous with the
                    // previous non-collapsed one, collapse it
                    if (strideOuter == strideInner * sizeInner)
                    {
                        ++numCollapsed;

                        // This is the run of collapsed dimensions' size
                        sizeInner = sizeInner * sizeOuter;
                        continue;
                    }
                }

                // Otherwise, this new outer dimension at `i` cannot be collapsed
                // because it is excluded from collapsing, or it is not contiguous
                // with the previous inner dimension.
                sizeInner = sizeOuter;
                strideInner = strideOuter;
            }

            // This will be our new size/stride and dimension.
            long[] newSizes = new long[TSCudaContext.MaxDims];
            long[] newStrides = new long[TSCudaContext.MaxDims];

            int newDims = tensor.DimensionCount - numCollapsed;

            // We return the index of the excluded dimension that is excluded
            // from being collapsed here.
            int returnDim = -1;

            // We perform a second pass through the dimensions to actually
            // calculate the size of the collapsed dimensions.
            int collapsedIndex = tensor.DimensionCount - numCollapsed - 1;
            newSizes[collapsedIndex] = tensor.Sizes[firstNonOneDim];
            newStrides[collapsedIndex] = tensor.Strides[firstNonOneDim];

            if (firstNonOneDim == excludeDim)
            {
                returnDim = collapsedIndex;
            }

            for (int i = firstNonOneDim - 1; i >= 0; --i)
            {
                long sizeOuter = tensor.Sizes[i];
                long strideOuter = tensor.Strides[i];

                if ((excludeDim != i) && (excludeDim != i + 1))
                {
                    if (sizeOuter == 1)
                    {
                        // skip
                        continue;
                    }

                    if (strideOuter == newSizes[collapsedIndex] * newStrides[collapsedIndex])
                    {
                        // collapse
                        newSizes[collapsedIndex] *= sizeOuter;
                        continue;
                    }
                }

                // Otherwise, strides don't match, or dim `i` is excluded from
                // collapsing.
                --collapsedIndex;
                //assert(collapsedIndex >= 0);
                //assert(collapsedIndex < newDims);
                newSizes[collapsedIndex] = sizeOuter;
                newStrides[collapsedIndex] = strideOuter;

                if (excludeDim == i)
                {
                    returnDim = collapsedIndex;
                }

            }

            info.sizes = newSizes.Take(newDims).ToArray();
            info.strides = newStrides.Take(newDims).ToArray();
            return returnDim;
        }
    }
}


<|start_filename|>TensorSharp.CUDA/KernelOps/ElementwiseTOp.cs<|end_filename|>
﻿using ManagedCuda;
using ManagedCuda.BasicTypes;
using System;
using System.Collections.Generic;
using System.Linq;
using TensorSharp.Core;
using TensorSharp.CUDA.DeviceCode;
using TensorSharp.CUDA.RuntimeCompiler;

namespace TensorSharp.CUDA.KernelOps
{
    public static class ApplyOpInvoke
    {
        public static void Invoke(TSCudaContext context, CudaContext cudaContext, byte[] ptx, string baseName, params object[] args)
        {
            ThrowIfAnyTensorInvalid(args);

            cudaContext.SetCurrent();

            CudaDeviceProperties deviceInfo = context.DeviceInfoForContext(cudaContext);

            IEnumerable<Tensor> allTensors = args.OfType<Tensor>();
            Tensor firstTensor = allTensors.First();
            long elementCount = firstTensor.ElementCount();
            ApplySpecialization spec = new ApplySpecialization(allTensors.ToArray());

            ConvertTensorArgs.Convert(cudaContext, spec.Use32BitIndices, args);

            ManagedCuda.VectorTypes.dim3 block = ApplyUtils.GetApplyBlock();
            ManagedCuda.VectorTypes.dim3 grid = ApplyUtils.GetApplyGrid(deviceInfo, elementCount);

            string fullKernelName = PermutationGenerator.GetMangledName(baseName, spec);
            CudaKernel kernel = context.KernelCache.Get(cudaContext, ptx, fullKernelName);

            kernel.GridDimensions = grid;
            kernel.BlockDimensions = block;
            kernel.RunAsync(CUstream.NullStream, args);

        }


        private static void ThrowIfAnyTensorInvalid(object[] args)
        {
            foreach (Tensor tensor in args.OfType<Tensor>())
            {
                if (tensor.DimensionCount > TSCudaContext.MaxDims)
                {
                    throw new InvalidOperationException("Tensors with dimension count > " + TSCudaContext.MaxDims + " are not supported");
                }
            }
        }


        public static void ApplyPrecompile(CudaCompiler compiler, DeviceKernelTemplate template, int tensorCount)
        {
            foreach (ApplySpecialization spec in ApplySpecialization.AllSpecializations(tensorCount))
            {
                template.PtxForConfig(compiler, spec.GetConfig());
            }
        }
    }

    public static class ElementwiseTTOp
    {
        public static Tensor Invoke(CudaCode kernels, string funcName, Tensor result, Tensor src)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, src.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == src)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, src, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTSOp
    {
        public static Tensor Invoke(CudaCode kernels, string funcName, Tensor result, Tensor src, float value)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, src.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == src)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, value, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, src, value, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTTSOp
    {
        public static Tensor Invoke(CudaCode kernels, string funcName, Tensor result, Tensor src, Tensor src2, float value)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, src.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == src)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, src2, value, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, src, src2, value, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTTTSOp
    {
        public static Tensor Invoke(ElementwiseKernels kernels, string funcName, Tensor result, Tensor src, Tensor src2, Tensor src3, float value)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, src.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == src)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, src2, src3, value, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, src, src2, src3, value, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTTTSSOp
    {
        public static Tensor Invoke(ElementwiseKernels kernels, string funcName, Tensor result, Tensor src, Tensor src2, Tensor src3, float step_size, float value)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, src.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == src)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, src2, src3, step_size, value, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, src, src2, src3, step_size, value, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTTSSOp
    {
        public static Tensor Invoke(ElementwiseKernels kernels, string funcName, Tensor result, Tensor src, Tensor src2, float step_size, float value)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, src.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == src)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, src2, step_size, value, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, src, src2, step_size, value, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTTOp
    {
        public static Tensor Invoke(CudaCode kernels, string funcName, Tensor result, Tensor lhs, Tensor rhs)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(lhs);
            CudaContext cudaContext = context.CudaContextForTensor(lhs);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, lhs, false, lhs.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == lhs)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, rhs, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, lhs, rhs, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTTTOp
    {
        public static Tensor Invoke(CudaCode kernels, string funcName, Tensor result, Tensor lhs, Tensor rhs, Tensor rhs2)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(lhs);
            CudaContext cudaContext = context.CudaContextForTensor(lhs);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, lhs, false, lhs.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == lhs)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, rhs, rhs2, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, lhs, rhs, rhs2, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ElementwiseTTTTTOp
    {
        public static Tensor Invoke(ElementwiseKernels kernels, string funcName, Tensor result, Tensor lhs, Tensor rhs, Tensor rhs2, Tensor rhs3)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(lhs);
            CudaContext cudaContext = context.CudaContextForTensor(lhs);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, lhs, false, lhs.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == lhs)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, rhs, rhs2, rhs3, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, lhs, rhs, rhs2, rhs3, elementCount);
            }

            return writeTarget;
        }
    }

    public static class ClampOp
    {
        public static Tensor Invoke(ElementwiseKernels kernels, Tensor result, Tensor src, float min, float max)
        {
            string funcName = "clamp";
            TSCudaContext context = CudaHelpers.TSContextForTensor(src);
            CudaContext cudaContext = context.CudaContextForTensor(src);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, src, false, src.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);

            if (result == src)
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t1_" + funcName, writeTarget, min, max, elementCount);
            }
            else
            {
                ApplyOpInvoke.Invoke(context, cudaContext, ptx, "t2_" + funcName, writeTarget, src, min, max, elementCount);
            }

            return writeTarget;
        }
    }

    public static class Atan2Op
    {
        public static Tensor Invoke(CudaCode kernels, Tensor result, Tensor srcY, Tensor srcX)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(srcY);
            CudaContext cudaContext = context.CudaContextForTensor(srcY);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, srcY, false, srcY.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);
            ApplyOpInvoke.Invoke(context, cudaContext, ptx, "atan2", writeTarget, srcY, srcX, elementCount);

            return writeTarget;
        }
    }

    public static class LerpOp
    {
        public static Tensor Invoke(ElementwiseKernels kernels, Tensor result, Tensor srcA, Tensor srcB, float weight)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(srcA);
            CudaContext cudaContext = context.CudaContextForTensor(srcA);

            cudaContext.SetCurrent();

            Tensor writeTarget = TensorResultBuilder.GetWriteTarget(result, srcA, false, srcA.Sizes);
            long elementCount = writeTarget.ElementCount();

            byte[] ptx = kernels.GetPtx(context.Compiler);
            ApplyOpInvoke.Invoke(context, cudaContext, ptx, "lerp", writeTarget, srcA, srcB, weight, elementCount);

            return writeTarget;
        }
    }

    public static class CopyOp
    {
        public static void Invoke(FillCopyKernels kernels, TSCudaContext context, CudaContext cudaContext, Tensor result, Tensor src)
        {
            //    cudaContext.SetCurrent();

            byte[] ptx = kernels.GetPtx(context.Compiler);
            long elementCount = result.ElementCount();
            ApplyOpInvoke.Invoke(context, cudaContext, ptx, "copy", result, src, elementCount);
        }
    }

    public static class FillOp
    {
        public static void Invoke(FillCopyKernels kernels, Tensor result, float value)
        {
            TSCudaContext context = CudaHelpers.TSContextForTensor(result);
            CudaContext cudaContext = context.CudaContextForTensor(result);

            cudaContext.SetCurrent();

            byte[] ptx = kernels.GetPtx(context.Compiler);
            long elementCount = result.ElementCount();
            ApplyOpInvoke.Invoke(context, cudaContext, ptx, "fill", result, value, elementCount);
        }
    }
}


<|start_filename|>TensorSharp/RefCounted.cs<|end_filename|>
﻿using System;
using System.Threading;

namespace TensorSharp
{
    /// <summary>
    /// Provides a thread safe reference counting implementation. Inheritors need only implement the Destroy() method,
    /// which will be called when the reference count reaches zero. The reference count automatically starts at 1.
    /// </summary>

    [Serializable]
    public abstract class RefCounted
    {
        private int refCount = 1;

        /// <summary>
        /// Construct a new reference counted object. The reference count automatically starts at 1.
        /// </summary>
        public RefCounted()
        {
        }

        ~RefCounted()
        {
            if (refCount > 0)
            {
                Destroy();
                refCount = 0;
            }
        }

        /// <summary>
        /// This method is called when the reference count reaches zero. It will be called at most once to allow subclasses to release resources.
        /// </summary>
        protected abstract void Destroy();

        /// <summary>
        /// Returns true if the object has already been destroyed; false otherwise.
        /// </summary>
        /// <returns>true if the object is destroyed; false otherwise.</returns>
        protected bool IsDestroyed()
        {
            return refCount == 0;
        }

        /// <summary>
        /// Throws an exception if the object has been destroyed, otherwise does nothing.
        /// </summary>
        protected void ThrowIfDestroyed()
        {
            if (IsDestroyed())
            {
                throw new InvalidOperationException("Reference counted object has been destroyed");
            }
        }

        protected int GetCurrentRefCount()
        {
            return refCount;
        }

        /// <summary>
        /// Increments the reference count. If the object has previously been destroyed, an exception is thrown.
        /// </summary>
        public void AddRef()
        {
            int curRefCount;
            int original;
            SpinWait spin = new SpinWait();
            while (true)
            {
                curRefCount = refCount;
                if (curRefCount == 0)
                {
                    throw new InvalidOperationException("Cannot AddRef - object has already been destroyed");
                }

                int desiredRefCount = curRefCount + 1;
                original = Interlocked.CompareExchange(ref refCount, desiredRefCount, curRefCount);
                if (original == curRefCount)
                {
                    break;
                }

                spin.SpinOnce();
            }
        }

        /// <summary>
        /// Decrements the reference count. If the reference count reaches zero, the object is destroyed.
        /// If the object has previously been destroyed, an exception is thrown.
        /// </summary>
        public void Release()
        {
            int original;
            int curRefCount;
            SpinWait spin = new SpinWait();
            while (true)
            {
                curRefCount = refCount;
                if (curRefCount == 0)
                {
                    throw new InvalidOperationException("Cannot release object - object has already been destroyed");
                }

                int desiredRefCount = refCount - 1;
                original = Interlocked.CompareExchange(ref refCount, desiredRefCount, curRefCount);
                if (original == curRefCount)
                {
                    break;
                }

                spin.SpinOnce();
            }

            if (refCount <= 0)
            {
                Destroy();
            }
        }
    }
}


<|start_filename|>Seq2SeqSharp/Tools/RandomGenerator.cs<|end_filename|>
﻿using System;

namespace Seq2SeqSharp
{
    [Serializable]
    public static class RandomGenerator
    {

        public static bool Return_V { get; set; }
        public static float V_Val { get; set; }

        private static readonly Random random = new Random(DateTime.Now.Millisecond);
        public static float GaussRandom()
        {
            if (Return_V)
            {
                Return_V = false;
                return V_Val;
            }
            double u = 2 * random.NextDouble() - 1;
            double v = 2 * random.NextDouble() - 1;
            double r = (u * u) + (v * v);

            if (r == 0 || r > 1)
            {
                return GaussRandom();
            }

            double c = Math.Sqrt(-2 * Math.Log(r) / r);
            V_Val = (float)(v * c);
            Return_V = true;
            return (float)(u * c);
        }

        public static float floatRandom(float a, float b)
        {

            return (float)(random.NextDouble() * (b - a) + a);

        }

        public static float IntegarRandom(float a, float b)
        {
            return (float)(Math.Floor(random.NextDouble() * (b - a) + a));
        }
        ////public static float NormalRandom(float mu, float std)
        ////{
        ////    return mu + GaussRandom() * std;
        ////}

        public static float NormalRandom(float mu, float std)
        {
            return mu + floatRandom(-1.0f, 1.0f) * std;
        }


    }

}


<|start_filename|>TensorSharp/Core/DelegateDisposable.cs<|end_filename|>
﻿using System;

namespace TensorSharp.Core
{
    public class DelegateDisposable : IDisposable
    {
        private readonly Action action;

        public DelegateDisposable(Action action)
        {
            this.action = action;
        }

        public virtual void Dispose()
        {
            action();
        }
    }
}


<|start_filename|>TensorSharp/OpConstraint.cs<|end_filename|>
﻿using System;

namespace TensorSharp
{
    public abstract class OpConstraint
    {
        public abstract bool SatisfiedFor(object[] args);
    }

    public class ArgCountConstraint : OpConstraint
    {
        private readonly int argCount;

        public ArgCountConstraint(int argCount) { this.argCount = argCount; }

        public override bool SatisfiedFor(object[] args)
        {
            return args.Length == argCount;
        }
    }

    public class ArgTypeConstraint : OpConstraint
    {
        private readonly int argIndex;
        private readonly Type requiredType;

        public ArgTypeConstraint(int argIndex, Type requiredType)
        {
            this.argIndex = argIndex;
            this.requiredType = requiredType;
        }

        public override bool SatisfiedFor(object[] args)
        {
            return requiredType.IsAssignableFrom(args[argIndex].GetType());
        }
    }

    public class ArgStorageTypeConstraint : OpConstraint
    {
        private readonly int argIndex;
        private readonly Type requiredType;
        private readonly bool allowNull;

        public ArgStorageTypeConstraint(int argIndex, Type requiredType, bool allowNull = true)
        {
            this.argIndex = argIndex;
            this.requiredType = requiredType;
            this.allowNull = allowNull;
        }

        public override bool SatisfiedFor(object[] args)
        {
            if (allowNull && args[argIndex] == null)
            {
                return true;
            }
            else if (!allowNull && args[argIndex] == null)
            {
                return false;
            }

            Storage argStorage = ((Tensor)args[argIndex]).Storage;
            return requiredType.IsAssignableFrom(argStorage.GetType());
        }
    }
}


<|start_filename|>TensorSharp/Expression/TExpression.cs<|end_filename|>
﻿using System;

namespace TensorSharp.Expression
{
    public abstract class TExpression
    {
        public bool IsValidLvalue { get; private set; }

        public TExpression(bool isValidLvalue = false)
        {
            IsValidLvalue = isValidLvalue;
        }

        public abstract Tensor Evaluate(Tensor writeTarget);
    }

    public class ViewExpression : TExpression
    {
        private readonly TExpression src;
        private readonly Func<Tensor, Tensor> evaluate;

        public ViewExpression(TExpression src, Func<Tensor, Tensor> evaluate)
            : base(src.IsValidLvalue)
        {
            this.src = src;
            this.evaluate = evaluate;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            if (writeTarget != null)
            {
                throw new InvalidOperationException("Cannot Select directly into another tensor");
            }

            using (Tensor s = src.Evaluate(null))
            {
                return evaluate(s);
            }
        }
    }

    public class FromArrayExpression : TExpression
    {
        private readonly IAllocator allocator;
        private readonly Array array;

        public FromArrayExpression(IAllocator allocator, Array array)
            : base(false)
        {
            this.allocator = allocator;
            this.array = array;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            if (writeTarget != null)
            {
                writeTarget.CopyFrom(array);
                return writeTarget;
            }
            else
            {
                return Tensor.FromArray(allocator, array);
            }
        }
    }

    public class AsTypeExpression : TExpression
    {
        private readonly TExpression src;
        private readonly DType type;

        public AsTypeExpression(TExpression src, DType type)
        {
            this.src = src;
            this.type = type;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor srcVal = src.Evaluate(null))
            {
                if (writeTarget == null)
                {
                    writeTarget = new Tensor(srcVal.Allocator, type, srcVal.Sizes);
                }

                Ops.Copy(writeTarget, srcVal);
                return writeTarget;
            }
        }
    }

    public class ToDeviceExpression : TExpression
    {
        private readonly TExpression src;
        private readonly IAllocator allocator;

        public ToDeviceExpression(TExpression src, IAllocator allocator)
        {
            this.src = src;
            this.allocator = allocator;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor srcVal = src.Evaluate(null))
            {
                if (writeTarget == null)
                {
                    writeTarget = new Tensor(allocator, srcVal.ElementType, srcVal.Sizes);
                }

                Ops.Copy(writeTarget, srcVal);
                return writeTarget;
            }
        }
    }

    public class ScatterFillExpression : TExpression
    {
        private readonly TExpression src;
        private readonly TExpression indices;
        private readonly SVar value;
        private readonly int dimension;


        public ScatterFillExpression(TExpression src, SVar value, int dimension, TExpression indices)
        {
            this.src = src;
            this.value = value;
            this.dimension = dimension;
            this.indices = indices;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor s = src.Evaluate(null))
            using (Tensor i = indices.Evaluate(null))
            {
                if (!writeTarget.Equals(s))
                {
                    Ops.Copy(writeTarget, s);
                }
                Ops.ScatterFill(writeTarget, value.Evaluate(), dimension, i);
            }

            return writeTarget;
        }
    }

    public class FillExpression : TExpression
    {
        private readonly IAllocator allocator;
        private readonly DType elementType;
        private readonly long[] sizes;
        private readonly Action<Tensor> fillAction;


        public FillExpression(IAllocator allocator, DType elementType, long[] sizes, Action<Tensor> fillAction)
        {
            this.allocator = allocator;
            this.elementType = elementType;
            this.sizes = sizes;
            this.fillAction = fillAction;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            if (writeTarget == null)
            {
                writeTarget = new Tensor(allocator, elementType, sizes);
            }

            fillAction(writeTarget);

            return writeTarget;
        }
    }

    public class AddmmExpression : TExpression
    {
        private readonly TExpression src, m1, m2;
        private readonly float alpha, beta;

        public AddmmExpression(float beta, TExpression src, float alpha, TExpression m1, TExpression m2)
        {
            this.beta = beta;
            this.src = src;
            this.alpha = alpha;
            this.m1 = m1;
            this.m2 = m2;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor s = src.Evaluate(null))
            using (Tensor m1Val = m1.Evaluate(null))
            using (Tensor m2Val = m2.Evaluate(null))
            {
                return Ops.Addmm(writeTarget, beta, s, alpha, m1Val, m2Val);
            }
        }
    }


    public class TensorValueExpression : TExpression
    {
        private readonly Tensor value;

        public TensorValueExpression(Tensor value)
            : base(true)
        {
            this.value = value;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            if (writeTarget == null)
            {
                return value.CopyRef();
            }
            else
            {
                Ops.Copy(writeTarget, value);
                return writeTarget;
            }
        }
    }

    public class BinaryTensorTensorExpression : TExpression
    {
        private readonly TExpression left, right;
        private readonly Func<Tensor, Tensor, Tensor, Tensor> evaluate;

        public BinaryTensorTensorExpression(TExpression left, TExpression right, Func<Tensor, Tensor, Tensor, Tensor> evaluate)
        {
            this.left = left;
            this.right = right;
            this.evaluate = evaluate;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor lhs = left.Evaluate(null))
            using (Tensor rhs = right.Evaluate(null))
            {
                return evaluate(writeTarget, lhs, rhs);
            }
        }
    }

    public class UnaryTensorExpression : TExpression
    {
        private readonly TExpression src;
        private readonly Func<Tensor, Tensor, Tensor> evaluate;

        public UnaryTensorExpression(TExpression src, Func<Tensor, Tensor, Tensor> evaluate)
        {
            this.src = src;
            this.evaluate = evaluate;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor s = src.Evaluate(null))
            {
                return evaluate(writeTarget, s);
            }
        }
    }

    public class BinaryScalarTensorExpression : TExpression
    {
        public readonly SExpression left;
        public readonly TExpression right;
        public readonly Func<Tensor, float, Tensor, Tensor> evaluate;

        public BinaryScalarTensorExpression(SExpression left, TExpression right, Func<Tensor, float, Tensor, Tensor> evaluate)
        {
            this.left = left;
            this.right = right;
            this.evaluate = evaluate;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor rhs = right.Evaluate(null))
            {
                return evaluate(writeTarget, left.Evaluate(), rhs);
            }
        }
    }

    public class BinaryTensorScalarExpression : TExpression
    {
        public readonly TExpression left;
        public readonly SExpression right;
        public readonly Func<Tensor, Tensor, float, Tensor> evaluate;

        public BinaryTensorScalarExpression(TExpression left, SExpression right, Func<Tensor, Tensor, float, Tensor> evaluate)
        {
            this.left = left;
            this.right = right;
            this.evaluate = evaluate;
        }

        public override Tensor Evaluate(Tensor writeTarget)
        {
            using (Tensor lhs = left.Evaluate(null))
            {
                return evaluate(writeTarget, lhs, right.Evaluate());
            }
        }
    }
}
