title: Efficiency of CUDA vector types (float2, float3, float4)
body: <p>I'm trying to understand the <code>integrate_functor</code> in <code>particles_kernel.cu</code> from CUDA examples:</p>
  
  <pre><code>struct integrate_functor
  {
      float deltaTime;    
      //constructor for functor
      //...
  
      template &lt;typename Tuple&gt;
      __device__
      void operator()(Tuple t)
      {
          volatile float4 posData = thrust::get&lt;2&gt;(t);
          volatile float4 velData = thrust::get&lt;3&gt;(t);
  
          float3 pos = make_float3(posData.x, posData.y, posData.z);
          float3 vel = make_float3(velData.x, velData.y, velData.z);
  
          // update position and velocity
          // ...
  
          // store new position and velocity
          thrust::get&lt;0&gt;(t) = make_float4(pos, posData.w);
          thrust::get&lt;1&gt;(t) = make_float4(vel, velData.w);
      }
  };
  </code></pre>
  
  <p>We call <code>make_float4(pos, age)</code> but <code>make_float4</code> is defined in <code>vector_functions.h</code> as</p>
  
  <pre><code>static __inline__ __host__ __device__ float4 make_float4(float x, float y, float z, float w)
  {
      float4 t; t.x = x; t.y = y; t.z = z; t.w = w; return t;
  } 
  </code></pre>
  
  <p>Are CUDA vector types (<code>float3</code> and <code>float4</code>) more efficient for the GPU and how does the compiler know how to overload the function <code>make_float4</code>?</p>