Replace tuple-style cast with pointer/stride access

It was noted in #1 (closed) that the get operator will not fully vectorize due to the tuple-style cast hidden in getStructMember. The pointers upfront are now cached in attempts to alleviate this. This changes the manner of data access in the AoSoA for arrays:

return getStructMember<I>(_raw_data[idx.s()]);

to:

return static_cast<struct_member_array_type<I> >( data<I>() + idx.s() * _strides[I] );

If this allows vectorization, then it may be well worth the extra arithmetic needed to index into the data block. Furthermore, if loop-invariant hoisting were done well by the compiler, idx.s() * _strides[I] would be hoisted above the inner loop over arrays as a further optimization.

Cacheing pointers and strides also had a nice side effect for future C interfaces - template-free pointer and stride access! A user can now get strides and (un-typed void*) pointers directly through the following interface:

size_t PosX = 0;
double* pos_x = (double*) aosoa.pointer(PosX);
size_t pos_x_stride = aosoa.stride(PosX);

core/example/cuda_uvm_example.cu

-#include "Cabana_AoSoA.hpp"
-#include "Cabana_Cuda.hpp"
+#include <Cabana_AoSoA.hpp>
+#include <Cabana_MemberDataTypes.hpp>
+#include <Cabana_Cuda.hpp>
 
 #include <cstdlib>
 #include <iostream>
@@ -60,7 +61,7 @@ __global__ void initializeParticles( ParticleList particles )
 {
     // Create a particle index. The Cuda block id is the struct id and the
     // thread id in the block is the array id in the struct.
-    ParticleList::Index idx( blockIdx.x, threadIdx.x );
+    Cabana::Index idx( array_size, blockIdx.x, threadIdx.x );
 
     // Only do the operation if we have a valid particle.
     if ( idx < particles.end() )