We need call PrepareForInput on the input argument before invoking a function.
The order of execution of parameters of a function is undefined, so we need to
make sure input is called before output, or else in-place use case breaks.
Instead of having a single specialization for sort and zip handles,
we know handle any fancy handles being passed to the cuda device adapter.
This was done by reworking how we represent fancy iterators inside thrust,
and instead of using a transform iterator + counting iterator we just use
a iterator_facade.
This is required so that we can use ArrayHandleZip with Sort/Reduce and
custom comparison operators. ArrayZip and PortalValue don't combine
well together, when used with DeviceAlgorithm that has a custom operator.
The custom operator is actually passed the PortalValue instead of
the real values, and by that point we can't fix anything since we
don't know what the original operator is.
We have increased the ARRAY_SIZE for the TestingDeviceAdapter to work
around issues with backends not using parallel algorithms for arrays of length
500 or less.
This is the first step to having better 3d scheduling for the serial backend.
While this doesn't improve the performance, it lays the foundation for
allowing I,J,K indexing when iterating Cells of a uniform grid.
The driving desire for this change was to make vtk-m ArrayHandle behave more
like the standard vector. The canonical example of what we want to support is:
```
std::vector<vtkm::Id> values;
values.resize(100);
vtkm::cont::ArrayHandle<vtkm::Id> array = vtkm::cont::make_ArrayHandle(values);
vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>::Sort(array);
```
And with-out the ability to modify user provided memory, this was impossible.
CellSetExplicit can now use different storage backends to allow
for cheaper representations of the data. For example a pool of triangles
can now implicit handles for shape and num indices.
This does raise the question if runtime polymorphism is the best approach
for CellSet, or if we should do some from of CRTP.
ExplicitConnectivity can now use different storage backends to allow
for cheaper representations of the data. For example a pool of triangles
can now implicit handles for shape and num indices.
The original implementation use the naive approach of the device adapter
to copy the memory to the dynamic array handle. The issue with this approach
is that it would do a upload to device, and than a copy on device. Instead
we should prefer a copy on host, and than an upload to device.
Previously we had accidentally bound the PrepareForInput method to upload
only to the default device adapter, which breaks the use case of allowing
multiple simultaneous backends at runtime.
The original implementation use the naive approach of the device adapter
to copy the memory to the dynamic array handle. The issue with this approach
is that it would do a upload to device, and than a copy on device. Instead
we should prefer a copy on host, and than an upload to device.
THe IncrementBy2 test type previously allowed any subtype including
floating point numbers. The meaning of this is actually a little unclear
and the feature was causing implicit type conversion warnings that were
hard to template out. The utility of of templating this class is dubious
in the first place, so class is now a fixed type.
I'm a little unsure whether we should keep this test class at all. It's
math operations are ad hoc and it could be difficult to determine if a
problem is caused by an actual problem or just bad math operators.
Fix compile warnings that come up with the flags
-Wconversion -Wno-sign-conversion
This catches several instances (mostly in the testing framework) where
types are implicitly converted. I expect these changes to fix some of
the warnings we are seeing in MSVC.
I was going to add these flags to the list of extra warning flags, but
unfortunately the Thrust library has several warnings of these types,
and I don't know a good way to turn on the warnings for our code but
turn them off for Thrust.
The UnitTestArrayHandlePermutation test was failing when compiled with
icc. I believe the issue is that the icc optimization takes some
liberties when computing literal floating point values versus computing
them at run time that makes the two slightly different. I changed all
the applicable comparisons in this test from using the == operator to
using the test_equal function, which adds a tolerance to the comparison.
I expect this to fix the test failure.
There are often instances where one wants to make an assert check in a
method that can run in either the control or execution environment. This
is rather difficult in general for the execution environment, but with
this change you can place the VTKM_ASSERT_CONT macro in such a method,
and it should compile even under CUDA. It works by removing the macro if
compiling for a CUDA device.
Per a discussion with Robert Maynard, we no longer have plans to bind
textures or use texture objects. Instead, we are now using the __ldg
command to load from global memory inside the execution portal, which
gives us most of the benefits of textures, but doesn't incur any
bookkeeping.
Previously ArrayTransfer and ArrayManagerExecution received a reference
to a Storage class in their constructor and held the reference as an
ivar. In retrospect, this is just asking for trouble. First, it is way
too easy to pass by value when you mean to pass by reference. Second, if
there ever is a bug where the Storage goes out of scope before the
classes holding a reference, it is that much harder to debug.
This includes changing methods like LoadDataForInput to PrepareForInput.
It also changed the interface a bit to save a reference to the storage
object. (Maybe it would be better to save a pointer?) These changes also
extend up to the ArrayManagerExecution class, so it can effect device
adapter implementations.
This API change effects both ArrayTransfer and ArrayManagerExecution.
This is in preparation for a future change to make the API more
consistent with ArrayHandle.
The UserPortal in ArrayHandle was used to copy a pointer the user
created into an ArrayHandle to use in VTK-m algorithms. However, this is
only really valid for a basic storage, so the functionality has been
moved there, and you have to construct an ArrayHandle with a storage
instead of an array portal.
Also found a problem with ArrayHandle that manifests itself with derived
types when you first do a PrepareForInput and then a PrepareForInPlace.
The ArrayHandle assumes the data is already moved to the device and
skips the in place call to the array transfer. However, this means the
transfer of the derived array handle does not have a chance to set up
for in place.
I think the appropriate solution may be to move the appropriate logic
from ArrayHandle to ArrayTransfer. I will look into that next.
The number of values in the array handle portal was screwy and the
GetNumberOfValues method was flat out wrong (thanks to Rob Maynard for
pointing that out). This is fixed.
Also fixed a subtle but nasty typing problem in the Storage's
GetPortalConst method.
Our approach of using the underlying allocator inside thrust was a bad approach,
for some reason it fails to properly allocate uint8's or int8's on the correct
boundaries. I expect that this logic is somewhere else in the code and
instead we should use thrust::system::cuda::vector which does this properly.
Add an Allocate method in ArrayHandle that basically forwards the
alllocate request to the storage object. This allows some measure of
control of the array from the control side. You can allocate the array
and set values (by getting the control array portal) if you so desire.
Previously when ReleaseResourcesExecution was called, the method blindly
deleted the execution array regardless of whether there was a valid copy
in the control environment. This could potentially lose data. What if
someone wanted to conserve memory on the device by clearing the array of
an output array?
There is also now an internal method that blindly deletes the array.
This is good for internal functions that are doing something to
invalidate the execution data anyway.
Fixed a problem where ArrayHandle would cause a crash if you tried to
get the control portal on an uninitialized array (it was supposed to
throw an exception).
Also changed some methods in ArrayHandle so that they work resonably
without error when used with an uninitialized array. Specifically, the
aforementioned behavior was changed to "allocate" an array of size 0
(that is, call Allocate(0) on the storage object) to return an empty
array portal rather than throw an error. Although this use of
ArrayHandle can be considered erroneous, it is convenient the get an
empty array portal when dealing with a potentially unallocated array
rather than create a special condition.
The namespaces need to be different for each test, or else only the first
implementation of the function will be used for all tests that call that
function.
Also updated the test to verify that we can count starting from a non zero
number.
The functors in the ForEach, StaticTransform, and DynamicTransform
methods sometimes can use the index of the parameter that they are
operating on. This can be a helpful diagnostic in compile and run-time
errors. It is also helpful when linking parameters from one
FunctionInterface with those of another.
This new features are now replacing implementations using the Zip
functionality that was removed earlier. The implementation is actually
simplified a bit.
Porting the dax device adapter over to vtkm. Unlike the dax version, doesn't
use the thrust::device_vector, but instead uses thrust::system calls so that
we can support multiple thrust based backends.
Also this has Texture Memory support for input array handles. Some more work
will need to be done to ArrayHandle so that everything works when using an
ArrayHandle inplace with texture memory bindings.
A couple of tests were failing with the Intel compiler due to
imprecision in comparing floating point values.
Also snuck in some minor documentation fixes in a comment for
FunctionInterface.
The unit test for StorageBasic tested the StealArray feature and then
used the delete[] operator on the stolen array to deallocate it. For
many standard libraries the default implementation for delete[] is
the same as (or at least compatible with) std::allocator, but for
the PGI compiler they were not compatible and this resulted in a
run-time error. This change fixes the problem with the test by using
the same allocator as the StorageBasic test.
It's easy to put accidently put something that is not a valid tag in a
ControlSignature or ExecutionSignature. Previously, when you did that
you got a weird error at the end of a very long template instantiation
chain that made it difficult to find the offending worklet.
This adds some type checks when the dispatcher is instantated to check
the signatures. It doesn't point directly to the signature or its
parameter, but it is much closer.
Instead of just checking that a dispatcher's Invoke input is an
ArrayHandle, also check that the ValueType of the ArrayHandle is
compatible with the types of the worklet operator. This is done by
adding a template argument to the ControlSignature tags that is a type
list tag that gets passed to the type check.
The Transport class is responsible for moving data from the control
environment to the execution environment. (Actually, it might be more
accurate to say it gets the execution environment associated with a
given control object.) The Transport class is templated with a tag that
controls the mechanism used for the transport.
Lots of tests have to move values in and out of arrays and check them
against expected values. It is also often the case that these tests are
run on lots of different types. There is some repeated code for
generating known values for particular indices. This change unifies some
of that. This can probably also encourage making more generic tests.
Providing these types tends to "lock in" the precision of the algorithms
used in VTK-m. Since we are using templating anyway, our templates
should be generic enough to handle difference precision in the data.
Usually the appropriate type can be determined by the data provided. In
the case where there is no hint on the precision of data to use (for
example, in the code that provides coordinates for uniform data), there
is a vtkm::FloatDefault.
Before we assumed that we would only use the basic types specified by
the widths of vtkm::Scalar and vtkm::Id. We want to expand this to make
sure the code works on whatever data precision we need.
Since we want our code to generally handle data of different precision
(for example either float or double) expand the types in our list types
to include multiple precision.
Previously we just hand coded base lists up to 4 entries, which was fine
for what we were using it for. However, now that we want to support base
types of different sizes, we are going to need much longer lists.
There are multiple reasons for this name change:
* The name Tuple conflicts with the boost::Tuple class, which as a
different interface and feature set. This gets confusing, especially
since VTK-m uses boost quite a bit.
* The use of this class is usually (although not always) as a
mathematical vector.
* The vtkm::Scalar and vtkm::Vector* classes are going to go away soon
to better support multiple base data type widths. Having this
abbriviated name will hopefully make the code a bit nicer when these
types have to be explicitly specified.
Also modified the implementation a bit to consolidate some of the code.
In preparation for supporting base types with more widths, add typedefs
for the base types with explicit widths (number of bits).
Also added a IdComponent type that should be used for indices for
components into tuples and vectors. There now should be no reason to use
"int" inside of VTK-m code (especially for indexing). This change cleans
up many of the int types that were used throughout.
When compiling with 32-bit Ids for a 64 bit machine (which is not
uncommon), it is possible that the distance between two iterators
is larger than the maximum value that can be stored in vtkm::Id.
If two such iterators were passed to ArrayPortalFromIterators, that
would cause problems.
This change checks for that condition and throws an out of memory
exception if it occurs. That would be a pretty darn big array and
is more likely to be the cause of an error somewhere else in the
code, but either way the check and error is good. This change also
fixes a warning we have been getting with MSVC.
For MSVC we use the non-portable wrapper stdext::checked_array_iterators
because the compiler insists on it for safety. When we check to make sure
our templates are giving us raw pointers, we have to check for this wrapper
instead of the raw pointer itself.
This moves the ability to get an iterator from an array portal out of
the portal itself. The next step is to move the GetIteratorBegin/End out
of ArrayPortal. This should make the implemenation a bit cleaner.
MSVC likes to warn about using raw pointers as iterators in generic
algorithms because they have been known to lead to problems. When
compiling with that compiler, wrap raw pointers in
stdext::checked_array_pointer to suppress the error and also add a bit
more checking.
I wanted to test ArrayHandleCounting with some non-standard data type.
I was using a class that looked like a number that counts by two, but
the operator behavior was not a proper group and that was causing issues.
Replaced that with a class that inefficiently represents an unsigned
integer as a string with that many characters. The inefficiency does not
matter because it is just a test.
It appears that when the Intel compiler is optimizing, constant floating
point values can be slightly different than the same value stored in memory
and never changed. This change uses the test_equal method to compare
these floating point values that might have a slight numeric error.
After a talk with Robert Maynard, we decided to change the name
ArrayContainerControl to Storage. There are several reasons for this
change.
1. The name ArrayContainerControl is unwieldy. It is long, hard for
humans to parse, and makes for long lines and wraparound. It is also
hard to distinguish from other names like ArrayHandleFoo and
ArrayExecutionManager.
2. The word container is getting overloaded. For example, there is a
SimplePolymorphicContainer. Container is being used for an object that
literally acts like a container for data. This class really manages
data.
3. The data does not necessarily have to be on the control side.
Implicit containers store the data nowhere. Derivative containers might
have all the real data on the execution side. It is possible in the
future to have storage on the execution environment instead of the
control (think interfacing with a simulator on the GPU).
Storage is not a perfect word (what does implicit storage really mean?),
but its the best English word we came up with.
It was originally put there to support CopyInto in ArrayHandle, but that
has already been removed. It really only makes sense for trivial
examples and testing code, and it sometimes causes complications with
coding.
There is a special version of the testing methods for use in the control
environment that handles execeptions that can be thrown there. There are
tests to make sure you are using the correct version of the testing
framework, but it was broken until the last commit. Now that it's fixed,
here are two places where the wrong testing method was used.
We made this change a while ago to help with completion in IDEs.
(Completion was matching a bunch of wrapper macros that were almost
never used anywhere.) Most of the changes are in comments, but there are
a few bad macro definitions.
Whenever creating a functor to be launched in the execution environment
using the device adapter Schedule algorithm, you had to also create a
couple of methods to handle error message buffers. For convenience, lots
of code started to just inherit from WorkletBase. Although this worked,
it was a misnomer (and might cause problems in the future if worklets
later require different things from its base). To get around this
problem, add a FunctorBase class that is intended to be used as the
superclass to functors called with Schedule.
Previously, VTKM_DEVICE_ADAPTER_UNDEFINED and
VTKM_ARRAY_CONTAINER_CONTROL_UNDEFINED were set to 0. The problem is
that if someone set VTKM_DEVICE_ADAPTER or VTKM_ARRAY_CONTAINER_CONTROL
to something invalid, that would test positive when compared to 0. Thus,
you get an error about not defining the default when in fact the problem
is setting an invalid flag.
This change makes the undefined constants -1 so that the comparison will
fail unless the macro is actually properly set.
Getting the type right for ArrayHandleCompositeVector can be a bit
tricky. It is expressed in a somewhat strange function signature format
and you have to extract the right component type for the return. This
adds an ArrayHandleCompositeVectorType that makes it easier (although no
less verbose).
Each type of point coordinates has its own class with the name
PointCoordinates*. Currently there is a PointCoordiantesArray that contains
an ArrayHandle holding the point coordinates and a PointCoordinatesUniform
that takes the standard extent, origin, and spacing for a uniform rectilinear
grid and defines point coordiantes for that. Creating new PointCoordinates
arrays is pretty easy, and we will almost definitely add more. For example,
we should have an elevation version that takes uniform coordinates for
a 2D grid and then an elevation in the third dimension. We can probably
also use a basic composite point coordinates that can build them from
other coordinates.
There is also a DynamicPointCoordinates class that polymorphically stores
an instance of a PointCoordinates class. It has a CastAndCall method that
behaves like DynamicArrayHandle; it can call a functor with an array handle
(possible implicit) that holds the point coordinates.
This derived array handle creates an array of vectors whose components come
from other arrays of vectors. In either case ArrayHandleCompositeVector
handles scalars as vectors of size 1.
This is used with the FunctionInterface::DynamicTransformCont method to
convert a call of arguments using dynamic array handles to a function
templated on concrete types.
Use this mechanism in the dynamic array handle to skip over trying
invalid array handle types (and thereby incurring a compiler error even
though we never intended to use these classes).
The dynamic array handle holds a reference to an array handle of an
unknown type. It contains the ability to try to cast it to an instance
of array handle or to try lists of types and containers.
There is currently an issue that is causing the test code not to
compile. It is the case that some combinations of types and containers
are not compatible. For example, an implict container is bound to a
certain type, and the container is undefined if they do not agree. There
needs to be a mechanism to detect these invalid combinations and skip
over them in the MTP for each.
Provies a list of types in a template like boost::mpl::vector and a
method to call a functor on each type. However, rather than explicitly
list each type, uses tags to identify the list. This provides the
following main advantages:
1. Can use these type lists without creating horrendously long class
names based on them, making compiler errors easier to read. For example,
you would have a typename like MyClass<TypeListTagVectors> instead of
MyClass<TypeList<Id3,Vector2,Vector3,Vector4> > (or worse if variadic
templates are not supported). This is the main motivation for this
implementation.
2. Do not require variadic templates and usually few constructions. That
should speed compile times.
There is one main disadvantage to this approach: It is difficult to get
a printed list of items in a list during an error. If necessary, it
probably would not be too hard to make a template to convert a tag to a
boost mpl vector.
