private static half __float2half(float a)
        {
            half r         = new half();
            uint sign      = 0;
            uint remainder = 0;

            r.x = __internal_float2half(a, ref sign, ref remainder);
            if ((remainder > 0x80000000U) || ((remainder == 0x80000000U) && ((r.x & 0x1U) != 0U)))
            {
                r.x++;
            }

            return(r);
        }
        public static extern CublasStatus cublasHgemmStridedBatched(CudaBlasHandle handle,
																 Operation transa,
																 Operation transb,
																 int m,
																 int n,
																 int k,
																 ref half alpha,  // host or device poi$
																 CUdeviceptr A,
																 int lda,
																 long strideA,   // purposely signed 
																 CUdeviceptr B,
																 int ldb,
																 long strideB,
																 ref half beta,   // host or device poi$
																 CUdeviceptr C,
																 int ldc,
																 long strideC,
																 int batchCount);
 public static extern CublasStatus cublasHgemm(CudaBlasHandle handle, 
                                               Operation transa,
                                               Operation transb, 
                                               int m,
                                               int n,
                                               int k,
                                               ref half alpha, /* host or device pointer */  
                                               CUdeviceptr A, 
                                               int lda,
                                               CUdeviceptr B,
                                               int ldb, 
                                               ref half beta, /* host or device pointer */  
                                               CUdeviceptr C,
                                               int ldc);
        private static half __double2half(double x)
        {
            // Perform rounding to 11 bits of precision, convert value
            // to float and call existing float to half conversion.
            // By pre-rounding to 11 bits we avoid additional rounding
            // in float to half conversion.
            ulong absx;

            ulong[]  ux = new ulong[1];
            double[] xa = new double[] { x };
            Buffer.BlockCopy(xa, 0, ux, 0, sizeof(double));

            absx = (ux[0] & 0x7fffffffffffffffUL);
            if ((absx >= 0x40f0000000000000UL) || (absx <= 0x3e60000000000000UL))
            {
                // |x| >= 2^16 or NaN or |x| <= 2^(-25)
                // double-rounding is not a problem
                return(__float2half((float)x));
            }

            // here 2^(-25) < |x| < 2^16
            // prepare shifter value such that x + shifter
            // done in double precision performs round-to-nearest-even
            // and (x + shifter) - shifter results in x rounded to
            // 11 bits of precision. Shifter needs to have exponent of
            // x plus 53 - 11 = 42 and a leading bit in mantissa to guard
            // against negative values.
            // So need to have |x| capped to avoid overflow in exponent.
            // For inputs that are smaller than half precision minnorm
            // we prepare fixed shifter exponent.
            ulong shifterBits = ux[0] & 0x7ff0000000000000UL;

            if (absx >= 0x3f10000000000000UL)
            {               // |x| >= 2^(-14)
                // add 42 to exponent bits
                shifterBits += 42ul << 52;
            }

            else
            {               // 2^(-25) < |x| < 2^(-14), potentially results in denormal
                // set exponent bits to 42 - 14 + bias
                shifterBits = ((42ul - 14 + 1023) << 52);
            }
            // set leading mantissa bit to protect against negative inputs
            shifterBits |= 1ul << 51;
            ulong[]  shifterBitsArr = new ulong[] { shifterBits };
            double[] shifter        = new double[1];

            Buffer.BlockCopy(shifterBitsArr, 0, shifter, 0, sizeof(double));

            double xShiftRound = x + shifter[0];

            double[] xShiftRoundArr = new double[] { xShiftRound };

            // Prevent the compiler from optimizing away x + shifter - shifter
            // by doing intermediate memcopy and harmless bitwize operation
            ulong[] xShiftRoundBits = new ulong[1];

            Buffer.BlockCopy(xShiftRoundArr, 0, xShiftRoundBits, 0, sizeof(double));

            // the value is positive, so this operation doesn't change anything
            xShiftRoundBits[0] &= 0x7ffffffffffffffful;

            Buffer.BlockCopy(xShiftRoundBits, 0, xShiftRoundArr, 0, sizeof(double));

            double xRounded = xShiftRound - shifter[0];
            float  xRndFlt  = (float)xRounded;
            half   res      = __float2half(xRndFlt);

            return(res);
        }
 /// <summary>
 ///
 /// </summary>
 public half(half h16)
 {
     x = h16.x;
 }