title: Converting uint16_t to float in CUDA kernel
body: <p><strong>Problem Statement:</strong></p>
  
  <p>I am currently processing image data in CUDA through a kernel. What I'll do in that kernel when it's complete is irrelevant for this conversation. Before I get to the fun stuff of developing the inner workings of the kernel, I'm writing a passthrough function that reads data from one buffer, converts it to float, and then pushes that data into another buffer.</p>
  
  <p>You can see from the function signature that I'm passing in a pointer to uint64_t data and passing out a pointer to float. Inside the kernel, I bit shift the uint64_t data into four chunks of uint16_t data. From here, in order to get to float, I run each uint16_t through two conversion functions: <code>__half2float(__ushort_as_half(x));</code></p>
  
  <ul>
  <li>Yes, CUDA recognizes the uint16_t and uint64_t data types.</li>
  <li>i can technically perform the function: <code>uint16_t val = 4; float fVal = val;</code> in CUDA</li>
  <li>but it yields a very different float value than <code>__half2float(__ushort_as_half(x));</code> does.</li>
  <li>i am including <code>stdint</code> because both host and device have no issue with the uint64_t or uint16_t variables.</li>
  </ul>
  
  <hr>
  
  <p><strong>Question:</strong></p>
  
  <p>Does anyone know if there's an easier way to convert from <code>uint16_t</code> to <code>float</code> in CUDA, while maintaining the same numeric results that <code>__half2float(__ushort_as_half(x));</code> yields?</p>
  
  <hr>
  
  <p><strong>CODE</strong></p>
  
  <pre><code>__global__ void _64bitPassThrough(const uint64_t *in, float *out, const int w, const int h, const int n)
  {
      int position = blockDim.x * blockIdx.x + threadIdx.x;
      int maxval = w * h;
  
      if (position &lt; maxval)
      {
          uint16_t a = (in[position] &gt;&gt; 48);
          uint16_t b = (in[position] &gt;&gt; 32) &amp; 0xffff;
          uint16_t g = (in[position] &gt;&gt; 16) &amp; 0xffff;
          uint16_t r = in[position] &amp; 0xffff;
          int new_position = n * position;
  
          out[new_position + 0] = __half2float(__ushort_as_half(r));
          out[new_position + 1] = __half2float(__ushort_as_half(g));
          out[new_position + 2] = __half2float(__ushort_as_half(b));
          out[new_position + 3] = __half2float(__ushort_as_half(a));
      }
  }
  </code></pre>
  
  <hr>
  
  <p><strong>Operating Parameters:</strong></p>
  
  <ul>
  <li><em>OS</em>: Windows 10</li>
  <li><em>GPU</em>: GeForce GTX 960M</li>
  <li><em>Compute Capability</em>: compute_50,sm_50</li>
  <li><em>CUDA Ver</em>: 9.0</li>
  <li><em>IDE &amp; Compiler</em>: Visual Studio 2015 Community &amp; MSVC</li>
  </ul>
  
  <hr>
  
  <p><strong>Sidebar:</strong></p>
  
  <p>I'm open to answering questions you may have about my approach. I will explain as best I can and take all relevant and useful advice as befits the project.</p>