A recent change to the DeviceAdapter header includes the TBB device if
available instead of the serial device. Thus, DeviceAdapterTagSerial was
not defined automatically in all cases for the build of
UnitTestDataSetPermutation. Add the header for that explicitly.

The DynamicArrayHandle and DynamicCellSet classes exclusively work in
the control environment. However, CUDA likes to add __device__ to
constructors, destructors, and assignment operators it automatically
adds. This in turn causes warnings about the __device__ function using
host-only classes (like boost::shared_ptr). Solve this problem by adding
explicit methods for all of these.

The CellSet classes all exclusively work in the control environment.
However, CUDA likes to add __device__ to constructors, destructors, and
assignment operators it automatically adds. This in turn causes warnings
about the __device__ function using host-only classes (like
boost::shared_ptr). Solve this problem by adding explicit methods for
all of these.

The ArrayHandle classes all exclusively work in the control environment.
However, CUDA likes to add __device__ to constructors, destructors, and
assignment operators it automatically adds. This in turn causes warnings
about the __device__ function using host-only classes (like
boost::shared_ptr). Solve this problem by adding explicit methods for
all of these.

Implemented this by wrapping up all these default objects in a macro.
This also solved the problem of other constructors that are necessary
for array handles such as a constructor that takes the base array
handle.

There is a strange nvcc warning in CUDA 7.5 that sometimes happens on MSVC
that causes it to emit a warning for an undefined method that is clearly
defined. The CUDA development team is aware of the problem and is going
to fix it, but these changes will work around the problem for now.

Thanks to Tom Fogal from NVIDIA for these fixes.

Under CUDA, the default constructors and destructors created are exported
as __host__ and __device__, which causes problems because they used a boost
pointer that only works on the host. The explicit copy constructors and
destructors do the same thing as the default ones except declared to only
work on the host.

This now allows for even more efficient construction of uniform point
coordinates when running under the 3d scheduler, since we don't need to go
from 3d index to flat index to 3d index, instead we stay in 3d index

When ICC replaced copy_backwards with copy it causes a failure when the
input and output arrays are the same.

Array handles for cuda device pointers have been implemented. The data for
these handles exists solely on the exec side (info such as length can be
queried from the cont side).

Previously, all Fetch objects received an Invocation object in their
Load and Store methods. The point of this was that it allowed the Fetch
to get data from any of the execution objects. However, every Fetch
either just got data directly from its associated execution object or
else used a secondary execution object (the input domain) to get indices
into their own execution object.

This left two potential areas for improvement. First, pulling data out
of the Invocation object was unnecessarily complicated. It would be much
nicer to get data directly from the associated execution object. Second,
when getting index information from the input domain, it was often the
case that extra computations were necessary (particularly on structured
cell sets). There was no way to share the index information among
Fetches, and therefore the computations were replicated.

This change removes the Invocation from the Fetch Load and Store.
Instead, it passes the associated execution object and a new object type
called the ThreadIndices. The ThreadIndices are customized for the input
domain and therefore have all the information needed for a redirected
lookup. It is also a thread-local object so it can cache computed
indices and save on computation time.

When you create a CellSetPermutation you provide an array of the cell ids that
you want to iterate. This allows the user to do custom blanking of a data set,
or to do multi iteration over a set of cells.

Remove the const from the ValueType of the delegate portal in
ArrayPortalGroupVec. This was creating a Vec with a const type, which
was immutable, which was problematic when trying to create the Vec in
the first place.

The testing of ArrayHandleGroupVec was just using the == operator to
check values. Even though we are not doing any math, optimizers can
sometimes make float values slightly different anyway, so test_equal
should give the correct comparison.

This is a fancy array handle that groups consecutive entries in a
delegate array handle and presents them as a single vec.

Previously the CellSetExplicit presumed that the CellToPoint and PointToCell
Storage tags matched, which they don't when using custom storage tags.

Previously if you created a cell set explicit and didn't set the number of
points you would get a runtime error when you over-ran an array's bounds.
Now we account for this use case and properly generate the Cell To Point
Connectivity.

Even when using implicit index's the ConnectivityExplicit would generate
the code to compute the IndexOffsets, which would than fail to compile as
the ArrayHandle would only support read operations. This fixes that issue.