eficient conversion from and to HALF

hi

Is there any ‘more’ efficient float to half and half to float conversion than the code below ?

Any tips’n’tricks in this topic ?

Handling of NaN and +/-Inf is probably not required :)

__device__ unsigned short __float_as_half(float Value)

{

 //

 // Zero case

 //

 if (Value == 0.0f) 

  return 0;

//

 // Generic case

 //

 int i = __float_as_int(Value);

 int s =  (i >> 16) & 0x00008000;

 int e = ((i >> 23) & 0x000000ff) - (127 - 15);

 int m =   i        & 0x007fffff;

if (e <= 0)

 {

  //

  // Below half_min

  //

  if (e < -10)

   return 0;

    m = (m | 0x00800000) >> (1 - e);

    if (m &  0x00001000)

         m += 0x00002000;

    return s | (m >> 13);

 }

 else 

 if (e == 0xff - (127 - 15))

 {

  if (m == 0)

  {

         return s | 0x7c00;

  }

    m >>= 13;

     

  return s | 0x7c00 | m | (m == 0);

 }

if (m & 0x00001000)

 {

     m += 0x00002000;

    if (m & 0x00800000)

     {

   m =  0; 

         e += 1; 

     }

 }

if (e > 30)

     return s | 0x7c00;

return s | (e << 10) | (m >> 13);

}

__device__ float __half_as_float(unsigned int Value)

{

 //

 // Zero case

 //

 if (Value == 0)

  return 0.0f;

//

 // Generic case

 //

 unsigned int Mantissa;

 unsigned int Exponent;

 unsigned int Result;

Mantissa = (Value & 0x03FF);

if ((Value & 0x7C00) != 0)

 {

  Exponent = ((Value >> 10) & 0x1F);

 }

 else 

 if (Mantissa != 0)

 {

     Exponent = 1;

    do

     {

   Exponent--;

         Mantissa <<= 1;

     } while ((Mantissa & 0x0400) == 0);

    Mantissa &= 0x03FF;

 }

 else

 {

  Exponent = (unsigned int)-112;

 }

Result = ((Value & 0x8000) << 16) | // Sign

          ((Exponent + 112) << 23) | // Exponent

           (Mantissa << 13);         // Mantissa

return __int_as_float(Result);

}

One trick - you can do half to float conversion (or any function on halfs) using a 256 x 256 texture as a lookup table (since half floats only have 65536 possible values).