I don't know why, but under some circumstances when compiling with CUDA
and Visual Studio, the compiler was giving syntax errors when the
IteratorFromArrayPortal was using the typename keyword to reference a
type in its superclass. Get around the problem by looking directly in
the superclass.
The build automatically sets some macros when building CUDA files. Some
of the CUDA sources were setting the same macros, which was causing
warnings. Change the code to be more careful about setting preprocessor
macros.
There are several tests in the cont directory that are in header files so
that they can be recompiled for different devices. Make sure that the
tests are exclusively using the device being tested by making the error
device adapter the default.
If any part of the test tries to use the default device (which could be
different than the one being tested), a compile error will occur. Several
of these compile errors are fixed in this commit.
There were a few places in the source code where
std::numeric_limits::min and max were used. There is an issue with these
methods on windows because the standard libraries there define macros
with the same name. Get around this problem by either places parentheses
so that they do not look like macros or use the vtkm::Infinity methods
instead.
When compiling under VisualStudio we need to first determine if checked
iterators are enabled ( _ITERATOR_DEBUG_LEVEL ). We don't want to use the
NDEBUG key, as we could be inside a project that is in Debug mode with
disabled checked iterators.
Secondly if they are enabled we need to handle the use case of NULL iterators
that get advanced by length zero. This last case is valid, but isn't supported
by the checked iterators so we need to work around it
First, be more explicit when we mean a range of values in a field or a
spacial bounds. Use the Range and Bounds structs in Field and
CoordinateSystem to make all of this more clear (and reduce a bit of
code as well).
The previous version of the bounds code required once less pass across
the data, but significantly increase the size of the resulting library:
Data for VTK-VTK-m interop:
- 7k more symbol table entries
- 1.5MB larger library
Because of the significant savings we need to use a less efficient
implementation that minimized the code size.
Certain algorithms like VertexClustering only work on certain cellset types.
It is pointless to generate the code paths for structured data cell sets,
when the algorithm can't operate on those dataset types.
Template instantiation is useful because when you are creating object files, as
uninstantiated template definitions are not are not added. Fully supporting
explicit instantiation like ITK does will require more code changes, but
this is a very minor step towards that goal.
VTK-m style guidelines dictate that all classes part of the API should
be in their own header file. Thus, we have moved the result classes out
of the filter header files and into their own headers.
We have also renamed these clases to ResultField and ResultDataSet to
better match the class naming conventions of VTK-m.
Since we are moving around these result classes, we have also revamped
their structure. They are now in a hierarchy with a common ResultBase
class. Additionally, all result classes have a reference to a complete
DataSet that can be considered the output of the filter. This will make
it easier for users to consistently deal with filter results.
These changes cover issues #60 and #62.
If Error.h was loaded before windows.h on windows compilers, then it
happily defined GetMessage and then GetMessageA and GetMessageB. However,
if the load order was reversed, then GetMessage is a macro that changes
it to GetMessageA or GetMessageB and you get an error about the same
method being declared twice.
Get around the problem by not defining GetMessage if it is already
declared as a macro.
These asserts are consolidated into the unified Assert.h. Also made some
minor edits to add asserts where appropriate and a little bit of
reconfiguring as found.
Previously the make_ArrayHandleCompositeVector template was declared in
such a way that if you passed it a fancy array handle, it would change
the type to a base ArrayHandle with a fancy storage. This was
inconsistent with types generated from ArrayHandleCompositeVectorType
when using the same fancy array handle types. This change makes that
more consistent.
This are just some minor things I ran into while documenting these classes.
One change is just an error in a comment. The other change allows you
to specify the type in make_ArrayHandleCast as a template argument instead
of creating an instance.
76b6179a Add missing copyright headers.
6a2a47d3 Remove order from new entries in MakeTestDataSet
d7c295f4 added rectilinear to test.
a8c7f92d Move mesa package to right place. Also, add in code for rectlinear mesh.
294b73e0 Make the rendering part of the build optional. Add the mesa libs to the link line.
719b715b Support for explicit cell sets.
98837735 Corrected error in projection matrix that was causing mismatches with ray tracing image and depth buffers.
f9c6f9cd cleanup of code.
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !383
removing all references to field order.
This removes all references to Field::order, since it is currently unused. If we decide to re-add it, it will be done in a more flexible and robust way.
See merge request !380
We have decided that we do not need the concept of an Extent in our data
model. Instead, we simply use dimensions, which can be represented in a
vtkm::Id3.
af7bba06 Refactor liner interpolation to use (1-w)*v0+w*v1
635f8c79 Add a CTestCustom file to to filter out warnings that cant be eliminated
f74c0d3c Remove type conversion related warnings for GCC
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !351
f81e8c64 Fix example that was using the old interface to VTKDataSetWriter
6aaf85ba Fix compile error with ofstream.
e863ee99 Change interface of VTKDataSetWriter
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !364
Make it match VTKDataSetReader better. It's also a bit easier to use
because you just have to give it a filename rather than open your own file
stream.
Previously we hinted to the compiler that it should vectorized operations
where the input and output are the same array. This obviously caused problems,
and these hints had to be removed.
In the future we need to first check for aliased arrays, and go from there.
e4237c3a Fix warnings found by the dashboard machines.
179b48e0 Reduce compile time for MarchingCubes by passing less info by Invoke.
8e4a47ef Update IsosurfaceUniformGrid to use the marching cubes filter.
d370155e MarchingCubes filter now generates coordinates when point merging is enabled.
c00fb53b Marching Cubes now generates vertices when merge duplicates is enabled.
f699c986 Renamed the ```Convert``` method to ```ApplyPolicy```
8e72ec8e Switch filter::threshold over to have a lower and upper bounds.
bcee8270 First draft of vtkm::filter design.
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !311
Even when running with the serial backend, the compiler might enable SIMD
vectorization when optimizations are turned on. When this occurs, we need
to use properly atomic Add's and CAE's.
A previous change used the Copy method of DeviceAdapterAlgorithm to
perform the actual copy in the CopyInto method. This works fine, except
that it uses the default device adapter, and the default device adapter
may need to copy the data to the device just to copy it to another
array. Instead, use the serial device adapter, which is guaranteed to
only perform one copy of the data.
I ran into a few minor issues with the constructors to the Field class.
The big change I made was that I removed the Field constructors that
take an example type and create an empty field of that type. The problem
was that the example type was easily confused with some other type that
was supposed to describe an array. This lead to some odd behavior in the
compiler and resulted in errors in unexpected places.
The use case for this constructor is dubious. There were several tests
in the code that would create an empty field, add it to a data set, then
get it back out to pass to the worklet. The code is much simpler if you
just make an ArrayHandle of the right type and use that in the worklet
invoke directly. It is also faster to compile with smaller code because
the type is known statically (whereas it is lost the other way).
The other change was to declare references to ArrayHandle and
DynamicArrayHandle as const. There is nothing in the behavior that
invalidates the const, and it accepts arrays constructed in the
parameter.
This makes is slightly easier to use as you do not actually have to
construct the DataSetBuilder object but just call its static method.
The DataSetBuilderExplicitIterative methods are not static because they
use state of the builder object to create the data.
Add GetCellDimensions to CellSetStructured
There was not an easy way to get the dimensions of cells in a structured
data set. Add a method to get that.
See merge request !323
4153c2c7 Found a few more places where we don't need to return by value.
dd85fc13 Document why we certain classes member variables need to be const ref.
6fb86da8 DynamicArrayHandle Casting methods now holds by const * const.
c1560e2d Perform less unnecessary copies when deducing a worklets parameters.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !320
Change Regular to Uniform
There was an inconsistency in naming classes where axes-aligned grids
with even spacing were sometimes called "uniform" and sometimes called
"regular". Maintain consistency by always calling them uniform.
See merge request !322
There was an inconsistency in naming classes where axes-aligned grids
with even spacing were sometimes called "uniform" and sometimes called
"regular". Maintain consistency by always calling them uniform.
Previously, DynamicArrayHandle and DynamicCellSet had slightly different
interfaces to their CastTo feature. It was a bit confusing and not all
that easy to use.
This change simplifies and unifies them by making each class have a single
CopyTo method that takes a reference to a cast object (an ArrayHandle or
CellSet, respectively) and fills that object with the data contained if
the cast is successfull. This interface gets around having to declare
strange types.
Each object also has a Cast method that has to have a template parameter
specified and returns a reference of that type (if possible).
In addition, the old behavior is preserved for DynamicArrayHandle (but
not DynamicCellSet). To avoid confusion, the name of that cast method is
CastToTypeStorage. However, the method was chaned to not take parameters
to make it consistent with the other Cast method.
Also, the IsType methods have been modified to reflect changes in
cast/copy. IsType now no longer takes arguments. However, an alternate
IsSameType does the same thing but does take an argument.
Previously in DynamicArrayHandle the dynamic_cast was contained in a
method that was reimplemented for every different instance of
DynamicArrayHandleBase. Change that to put the dynamic_cast in a
function outside of the class so that there is only one implementation
created per ArrayHandle type.
Similarly, DynamicCellSet had its dynamic_cast in a method plus there
was a second version in the functors used for the CastAndCall method.
Consolidate all of these into a function outside of either much like
DynamicArrayHandle.
Mainly issue dealing with dimensionality of cell sets and what that represents.
Have added in code to allow user to specify a custom dimensionality so that
tests continue to work properly.
Synchronize the CUDA timer on both the start and end events
Previously, the timer for CUDA devices only called cudaEventSynchronize
at the end event when asking for the elapsed time. This, however, could
allow time to pass from when the timer was reset to when the start event
happened that was not recorded in the timer. This added synchronization
should make sure that all time spent in CUDA is recorded.
See merge request !291
Reduce the number of conditions UnitTestArrayHandleCartesianProduct
tries so that it is faster. All the conditions are pretty similar, so it
should be OK to reduce some.
Previously UnitTestDataSetBuilderRectilinear and
UnitTestDataSetBuilderRegular did an exhaustive test of all possible
grid sizes within a certain amount of dimensions and with some number of
ways to build point arrays. This created hundreds of thousands of cases
that were checked, and it was running a long time. At best it wasted
time and at worst ctest reported the test as timed out.
It is not really necessary to perform an exhaustive test. The tests are
changed to do 10 trials using random values for dimensions and fill
methods.
f86382f0 Fix support for CoordinateSystems using ArrayHandleCartesianProduct.
d6a2a142 Add toleranced compare for values. Add tests for vtkm::Float32,Float64,Id typed arrays.
5d438353 Add toleranced comparisions for bounds validation. Also, add vtkm::Float32 and vtkm::Float64 to the testing for rectilinear and regular datasets.
b225ae97 Rectilinear coordinates (created with DataSetBuilderRectilinear) are now converted to vtkm::FloatDefault. This reduces the number of types to consider when casting inside CoordinateSystem, and was felt by all to be a reasonable restriction.
d755e43d Use ArrayHandleCompositeVector to represent separated point arrays for DataSetBuilderExplicit.h.
c7b0ffb8 Add tests for DataSetBuilderExplicit. Added cont/testing/ExplicitTestData.h which includes several explicit datasets. These datasets come from VTK data generated in VisIt. The new unit tests build datasets in several different ways and do some basic validation.
b4d04fff Add specialization of printSummary_ArrayHandle for UInt8. It prints them as characters, which are a little hard to understand to this computer scientist.
bd929c20 Fix compiler warnings.
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !262
It was previously using vtkm::cont::CellSetListTagCommon which would
break the ability for people to specify a custom value for
VTKM_DEFAULT_CELL_SET_LIST_TAG.
Add lots of checks to CUDA calls in the timer to try to identify any
problems that might be showing up on the dashboard.
Also adding some print statements around the sleep function in the
device adapter testing. For some reason the problem just went away with
them.
Previously, the timer for CUDA devices only called cudaEventSynchronize
at the end event when asking for the elapsed time. This, however, could
allow time to pass from when the timer was reset to when the start event
happened that was not recorded in the timer. This added synchronization
should make sure that all time spent in CUDA is recorded.
b70a000a Allow resetting the Type and Storage of a DynamicArrayHandle in a single call.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !302
vtkm::filter has the use case where we need to reset both the type and
storage of an array handle, by doing both at the same time we can reduce
the number of temporary objects, and invalid conversions of arrays.
Added cont/testing/ExplicitTestData.h which includes several explicit datasets. These datasets come from VTK data generated in VisIt. The new unit tests build datasets in several different ways and do some basic validation.
Add some new methods for DataSetFieldAdd class to improve usability.
c7756c78 TBB SortByKey works now when the key is a ArrayHandleZip.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !301
Previously, the ArrayHandleCompositeVector had a separate implementation
of ArrayPortal for the control and execution environments. Because I was
lazy when I implemented it, the control version did not support Get.
Since originally implementing this class, VTK-m now allows defining
methods that are declared as working in both control and execution
environments (VTKM_EXEC_CONT_EXPORT) but only work in one or the other
depending on methods of templated subclasses they call. Thus, solve this
problem by simply removing the control version of the portal and use the
same portal for both.
1. Additional ASSERT calls to validate arguments in: DataSetBuilderRegular
2. Fix some untested compile errors in DataSetBuilderRectilinear
3. Added a new unit test, cont/testing/UnitTestDataSetBuilderRectilinear.cxx
4. Provided additional tests for UnitTestDataSetBuilderRegular.cxx.
The new tests in (4) were also included in (3), and provide a much more robust way of validating datasets created. It has nested for loops to do an all-all test on various ways to specify the X,Y, and Z coordinates. It computes the bounds on the coordinate system and make sure they are correct.
Note: The GetBounds() call for Rectilinear is not working, and is an item for future discussion. It is disabled for now.
Previously each device adapter only had a unique string name. This was
not the best when it came to developing data structures to track the status
of a given device at runtime.
This adds in a unique numeric identifier to each device adapter. This will
allow classes to easily create bitmasks / lookup tables for the validity of
devices.
a7127f0f Adding vtkm::cont::RuntimeDeviceInformation.
7d249e89 Move DeviceAdapterTraits into vtkm::cont as they are user API.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !287
The RuntimeDeviceInformation class allows developers to check if a given
device is supported on a machine at runtime. This allows developers to properly
check for CUDA support before running any worklets.
When writing multiple backend code users of vtkm need to use the
DeviceAdapterTraits classes, so therefore we should move them to vtkm::cont
to signify this.
The WholeArrayIn, WholeArrayInOut, and WholeArrayOut ControlSignature
tags behave similarly to using an ExecObject tag with an
ExecutionWholeArray or ExecutionWholeArrayConst object. However, the
WholeArray* tags can simplify some implementations in two ways. First,
it allows you to specify more precisely what data is passed in. You have
to pass in an ArrayHandle or else an error will occur (as opposed to be
able to pass in any type of execution object). Second, this allows you
to easily pass in arrays stored in DynamicArrayHandle objects. The
Invoke mechanism will automatically find the appropriate static class.
This cannot be done easily with ExecutionWholeArray.
Scatter in worklets
Add the functionality to perform a scatter operation from input to output in a worklet invocation. This allows you to, for example, specify a variable amount of outputs generated for each input.
See merge request !221
Array transforms can now be created with an inverse functor, allowing for
casts back into the native array type. As a result, array transforms with
both a functor and inverse functor defined can perform read and write
operations. As an example, ArrayHandleCast now supports this operation. The
original implementation of ArrayHandleCast (i.e. read only) has been renamed
'ArrayHandleCastForInput'.
The original workaround for inclusive_scan bugs in thrust 1.8 only solved the
issue for basic arithmetic types such as int, float, double. Now we go one
step further and fix the problem for all types.
The solution is to provide a proper implementation of destructive_accumulate_n
and make sure it exists before any includes of thrust occur.
Previously, there was a declaration ConstArrayPortalFromThrust<const T>
in ArrayManagerExecutionThrustDevice. This proved problematic because
values read from the array in the worklet were typed as const T rather
than simply T. Any Vec or Matrix built from that type would then fail
because they are not meant to work with a const value (which means they
have to be set on construction and never changed.
Instead, declare ConstArrayPortalFromThrust<T> and internally set all
the Thrust pointers to have type const T. Also declare other thrust
pointers used as method parameters to have const T rather than T. This
should work as conversion from T to const T should be fine, but not the
other way around.
Now that ScatterCounting is implemented, we can use that to implement a
good part of the triangle generation in the isosurface algorithm. This
changes the worklet from a basic map to a topology map, which also
reduces a lot of code.
The parallel implementation in CellSetExplicit that builds cell-to-point
connectivity from point-to-cell connectivity uses a parallel sort-by-
key. The sort-by-key in the device adapter is not guaranteed to be
stable, so values associated with a particular key can be in any order.
The test for the result was expecting the connectivity array to be in a
particular order. Change the test to allow any connectivity ordering
that is still valid.
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.
adding VTK file exporter and test cases
This adds a legacy VTK file exporter which supports unstructured, explicit, and point meshes. (Single Cell Type cell sets are also supported.)
See merge request !247
CUDA default constructors, destructors, and assignment operators
Several classes exclusively work in the control environment. However, CUDA likes to add __device__ to constructors, destructors, and assignment operators it automatically creates. 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.
See merge request !245
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
Change Fetches to use ThreadIndices instead of Invocation.
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.
See merge request !233
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.
9a8809f9 Add CellSetPermutation which allows custom iteration over a cell set.
66f6db5a IsWriteableArrayHandle now can tell if an array handle can be written too
20f3fb50 Update VertexClustering to use vtkm::cont::CellSetSingleType.
154896b7 Extend the test for DataSetSingleType.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !228
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.
8bc40880 Add a test for CellSetExplicit::GetIndices
fa81d1de CellSetExplicit always calls methods using "this->"
9e496306 Allow incremental construction of CellSetSingleType.
3e307879 Mark CellSetExplicit::Fill as host side only.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !227
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.
