title: CUDA __float_as_int in acosf implementation
body: <p>CUDA C's maths function implementation (<code>cuda/math_function.h</code>) of <code>acosf</code> contains the passage:</p>
  
  <pre><code>if (__float_as_int(a) &lt; 0) {
    t1 = CUDART_PI_F - t1;
  }
  </code></pre>
  
  <p>where <code>a</code> and <code>t1</code> are <code>floats</code> and <code>CUDART_PI_F</code> is a <code>float</code> previously set to a numerical value close to the mathematical constant Pi.
  I am trying to understand what the conditional (if-clause) is testing for and what would be the C equivalent of it or the function/macro <code>__float_as_int(a)</code>. I searched for the implementation of <code>__float_as_int()</code> but without success. It seems that <code>__float_as_int()</code> is a built-in macro or function to NVIDIA NVCC. Looking at the PTX that NVCC produces out of the above passage:</p>
  
  <pre><code>    .reg .u32 %r&lt;4&gt;;
      .reg .f32 %f&lt;46&gt;;
      .reg .pred %p&lt;4&gt;;
      // ...
      mov.b32         %r1, %f1;
      mov.s32         %r2, 0;
      setp.lt.s32     %p2, %r1, %r2;
      selp.f32        %f44, %f43, %f41, %p2;
  </code></pre>
  
  <p>it becomes clear that <code>__float_as_int()</code> is not a <code>float</code> to <code>int</code> rounding. (This would have yielded a <code>cvt.s32.f32</code>.) Instead it assigns the <code>float %f1</code> as a bit-copy (<code>b32</code>) to <code>%r1</code> (notice: <code>%r1</code> is of type <code>u32</code> (unsigned int)!!) and then compares <code>%r1</code> as if it was a <code>s32</code> (signed int, confusing!!) with <code>%r2</code> (who's value is <code>0</code>).</p>
  
  <p>To me this looks a little odd. But obviously it is correct.</p>
  
  <p>Can someone explain what's going on and especially explain what <code>__float_as_int()</code> is doing in the context of the if-clause testing for being negative (<code>&lt;0</code>)? .. and provide a C equivalent of the if-clause and/or <code>__float_as_int()</code> marco ?</p>