Commit Graph

74 Commits

Author SHA1 Message Date
Robert Maynard
18a0cd35b0 vtkm::worklet::Invoker now supports scatter types
Fixes #297
2019-05-15 11:04:14 -04:00
Robert Maynard
77378993f8 DispatcherBase: Simplify remove_cvref and remove_pointer_and_decay.
designed easier to use remove_cvref and remove_pointer_and_decay
functions for the DispatcherBase.
2019-05-15 09:37:56 -04:00
nadavi
fbcea82e78 conslidate the license statement 2019-04-17 10:57:13 -06:00
Robert Maynard
6306ca66ce Merge topic 'refactor_cast_and_call_location'
72cb6343b Move vtkm::cont::CastAndCall into a header in vtkm/cont

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1622
2019-04-04 10:05:22 -04:00
Robert Maynard
72cb6343b8 Move vtkm::cont::CastAndCall into a header in vtkm/cont
Since CastAndCall is meant to be a user callable function,
we should make sure it is part of the vtkm cont headers
2019-04-03 14:15:26 -04:00
Robert Maynard
f1056affa7 Move select functions to host only to remove host/device suppressions
We previously had to mark the make_FunctionInterface function as
host/device as it could be used in either place. The ramifications
of this is that each time we launched a worklet any input parameter
had to either suppress cuda exceptions, or had to rely on the
DispatcherBase to suppress the warnings.

By making make_FunctionInterface only host callable ( as it is ),
we can remove all of our unneccesary suppression logic and better
expose real issues with code that is marked host/device but can't
be due to calling things such as std::abort
2019-04-03 12:48:33 -04:00
Kenneth Moreland
04254dbd25 Merge topic 'specialize-worklet-for-device'
4e34feecb Add ability to specialize worklet for device

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Sujin Philip <sujin.philip@kitware.com>
Merge-request: !1608
2019-04-02 12:08:01 -04:00
Kenneth Moreland
4e34feecb4 Add ability to specialize worklet for device
This adds an ExecutionSignature tag named Device that passes the
DeviceAdapterTag as an argument to the worklet's operator(). This allows
worklets to specialize their code based on the device.
2019-04-01 10:01:54 -06:00
Robert Maynard
18ff6681fb Less vtkm_cont cxx files bring in the cuda device
When you have CUDA enabled we need to make sure that all worklet
launches come from a cuda file otherwise we will generate ODR
violations.
2019-04-01 08:21:07 -04:00
Allison Vacanti
9014de0ebf Suppress new host/device warnings for cuda 10.1. 2019-03-26 14:53:42 -04:00
Allison Vacanti
d1db4ef8b3 Clarify intent of TypeString and TypeName functions.
TypeName is used for logging, and is now TypeToString.

TypeString is used for serialization, and is now SerializableTypeString.
2019-03-01 11:47:53 -05:00
Kenneth Moreland
191d6e5580 Add Mask capabilities to worklets
Mask objects allow you to specify which output values should be
generated when a worklet is run. That is, the Mask allows you to skip
the invocation of a worklet for any number of outputs.
2019-02-25 08:58:39 -07:00
Allison Vacanti
bd337854ec Initial implementation of general logging.
Addresses #291.
2018-10-02 11:37:55 -04:00
Robert Maynard
17a93921c3 Suppress CUDA 10 warnings about __host__ / __device__ 2018-09-20 14:02:16 -04:00
Kenneth Moreland
d879188de0 Make DispatcherBase invoke using a TryExecute
Rather than force all dispatchers to be templated on a device adapter,
instead use a TryExecute internally within the invoke to select a device
adapter.

Because this removes the need to declare a device when invoking a
worklet, this commit also removes the need to declare a device in
several other areas of the code.
2018-08-29 19:18:54 -07:00
haocheng liu
60363fff07 Provide a better error message when ExecutionSig != ControlSig
By facilitating brigand library, we can throw a meaningful compile time
error if the placeholder in execution signature exceeds the bound in
control signature.
2018-08-20 09:33:04 -04:00
Robert Maynard
64958b014b VTK-m now supports passing pointers when invoking worklets.
The original design of invoke and the transport infrastructure
relied on the implementation behavior of vtkm::cont types
such as ArrayHandle that used an internal shared_ptr to managed
state. This allowed passing by value instead of passing by
non-const ref when needing to transfer information to the device.

As VTK-m adds support for classes that use virtuals the ability
to pass by base pointer type allows for us to invoke worklets
using a base type without the risk of type slicing.

Additional by moving over to a non-const ref Invocation we
can update all transports that have 'output' to now be
by ref and therefore support types that can't be copied while
being 'more' correct.
2018-07-06 14:27:36 -04:00
Allison Vacanti
04d9786688 Suppress NVCC warning 2885.
This warning is emitted when constructing a tuple that contains
an object with `__host__`-only constructors.

warning #2885-D: calling a __host__ function(...) from a __host__ __device__ function(...) is not allowed
2018-06-04 14:41:15 -04:00
Kenneth Moreland
edc4c85fd9 Move Scatter from Worklet to Dispatcher
Previously, when a Worklet needed a scatter, the scatter object was
stored in the Worklet object. That was problematic because that means
the Scatter, which is a control object, was shoved into the execution
environment.

To prevent that, move the Scatter into the Dispatcher object. The
worklet still declares a ScatterType alias, but no longer has a
GetScatter method. Instead, the Dispatcher now takes a Scatter object in
its constructor. If using the default scatter (ScatterIdentity), the
default constructor is used. If using another type of Scatter that
requires data to set up its state, then the caller of the worklet needs
to provide that to the dispatcher. For convenience, worklets are
encouraged to have a MakeScatter method to help construct a proper
scatter object.
2018-04-27 00:43:51 -04:00
Robert Maynard
73c05c10c7 Suppressions of __host__ warnings shouldn't trigger internal compile errors 2018-04-12 12:04:53 -04:00
Robert Maynard
5100e559c7 Suppress warnings about calling __host__ functions from CUDA 9.0 2018-04-12 10:31:19 -04:00
Robert Maynard
282c515ea8 Suppress warnings about calling __host__ functions from CUDA 8.0 2018-04-12 10:06:57 -04:00
Robert Maynard
8472b64e60 Suppress more CUDA host function from a host/device function warnings 2018-04-05 10:55:43 -04:00
Robert Maynard
8808b41fbd Merge branch 'master' into vtk-m-cmake_refactor 2018-03-29 22:51:26 -04:00
Robert Maynard
7c54125b66 Switch over from static const to static constexpr where possible. 2018-03-10 11:39:58 -05:00
Robert Maynard
a60e378003 CUDA 7.5 doesn't like suppressing warning codes that don't exist
Newer versions of CUDA have no problem suppressing warning codes that don't
exist.
2018-02-26 16:37:57 -05:00
Robert Maynard
ee69c7a4b7 Remove VS2013 workarounds from VTK-m. 2018-02-23 15:39:39 -05:00
Robert Maynard
e630ac5aa4 Merge branch 'master' into vtk-m-cmake_refactor 2018-02-23 14:52:00 -05:00
Robert Maynard
182f4707e7 vtkm prefers 'using' over typedef. 2018-02-23 10:47:20 -05:00
Robert Maynard
022c987182 Correct warnings and errors found with MSVC2017+CUDA9 2018-01-31 15:58:45 -05:00
luz.paz
80b11afa24 Misc. typos
Found via `codespell -q 3` via downstream VTK
2018-01-30 06:51:47 -05:00
Robert Maynard
5e0d992057 Suppress host/device warnings when building with cuda 9 2018-01-16 15:42:36 -05:00
Kenneth Moreland
fe9594c1d1 Add hint to read source code for help
In trying to give error diagnostics with template definitions of invalid
types, the user encounters some pretty confusing error messages at
first. There is no way to get the compiler to give exactly the
diagnostics we want in a nice readable error message, so we are putting
some verbose instructions as comments in the code. However, a user might
not know to look at the source code since the error happens deep in an
unfamiliar (and complicated) class. Thus, add (yet another) error at the
front that gives a (hopefully) clear indication to look at the source
code for help in understanding the errors.
2018-01-08 10:53:44 -07:00
Kenneth Moreland
c3889a1a67 Give better error reporting when an Invoke parameter is wrong
When one of the parameters to DispatcherBase::Invoke is incorrect, you
get an error in a strange place. It is deep in a call stack using some
heavily templated types rather than where the Invoke is actually called.

Formerly, the code was structured to give a very obfuscated error
message. Try to make this easier on users by first adding helpful hints
in the code around where the error is to explain why the error occured
and helpful advice on how to track down the problem. Second, reconfigure
the static_assert so that the compiler gives a readable error message.
Third, add some auxiliary error messages that provide additional
information about the types and classes involved in the error.
2018-01-08 10:53:43 -07:00
Robert Maynard
93bc0198fe Suppress false positive warnings about calling host device functions. 2018-01-02 10:40:49 -05:00
Robert Maynard
d51ddda9b6 Since sizeof returns std::size_t we should iterate using unsigned integers 2017-12-06 13:57:05 -05:00
Robert Maynard
f6e18ac452 Remove IntegerSequence.h as we don't need it in vtk-m anymore 2017-11-29 10:55:04 -05:00
Robert Maynard
7f76220427 Redesign the Dispatcher to not need FunctionInterface to convert dynamic types 2017-11-28 11:01:02 -05:00
Kenneth Moreland
c3a3184d51 Update copyright for Sandia
Sandia National Laboratories recently changed management from the
Sandia Corporation to the National Technology & Engineering Solutions
of Sandia, LLC (NTESS). The copyright statements need to be updated
accordingly.
2017-09-20 15:33:44 -06:00
Robert Maynard
ce80383238 Adds WorkletPointNeighborhood and DispatcherPointNeighborhood types.
VTK-m is now able to run algorithms on structured points that require the
local point neighbors in a highly efficient manner.
2017-08-23 16:42:00 -04:00
Robert Maynard
5dd346007b Respect VTK-m convention of parameters all or nothing on a line
clang-format BinPack settings have been disabled to make sure that the
VTK-m style guideline is obeyed.
2017-05-26 13:53:28 -04:00
Robert Maynard
60a405ef65 Add TaskTiling1D/3D which use faux virtuals to reduce binary size.
Redesigns the TBB and Serial backends and the vtkm::exec::Task concept so that
we can re-use the same launching logic for all Worklets, instead of generating
per worlet code. To keep the performance the same the TilingTask now is past
a range of indices to work on, rather than a single index.

Binary size reduction:
WorkletTests_SERIAL old - 19MB
WorkletTests_SERIAL new - 18MB

WorkletTests_TBB old - 39MB
WorkletTests_TBB new - 18MB

libvtkAcceleratorsVTKm old - 48MB
libvtkAcceleratorsVTKm new - 19MB
2017-05-25 11:00:01 -04:00
Kitware Robot
4ade5f5770 clang-format: apply to the entire tree 2017-05-25 07:51:37 -04:00
Robert Maynard
022c36fa4f Add vtkm::exec::TaskBase, and rename WorkletInvokeFunctor to TaskSingular
Previously WorkletInvokeFunctor inherited from vtkm::exec::FunctorBase,
which is also the base class for all users Worklets and for all functors
based to DeviceAdapter::Schedule.

This is done for a few reasons. The first is that we reduce the
minimum size of user worklets. Previously the users worklet would hold
a reference to the error message, and so would the wrapper class added
when calling DeviceAdapter::Schedule. Now we only have the users worklet
holding a reference.

Second, by refactoring to have two base classes we can better improve
the documentation on what responsibilities FunctorBase.h has, compared
to TaskBase.
2017-05-02 16:38:43 -04:00
Kenneth Moreland
0b12cab1dd Fix overshadowing of class member with template parameter 2017-03-24 15:36:58 -06:00
Kenneth Moreland
dc192b793d Add input range to arguments of transport
Previously the arguments to the operator of a vtkm::cont::arg::Transport
were the control object, the input domain object, and the output range.
If you wanted to, for example, check the size of an input array to make
sure it matched the input range, you would have to know the meaning of
the input domain object to query its range. This made it hard to create
generic transports, like TransportTagArrayIn, that accept data from
multiple different input domains but need to know the input range.
2017-03-24 15:14:18 -06:00
Sujin Philip
8c4bbc39ad Use C++11 =delete keyword 2017-02-24 09:39:22 -05:00
David C. Lonie
f601e38ba8 Simplify exception hierarchy.
Remove the ErrorControl class such that all subclasses now inherit from
error. Renamed all exception classes via s/ErrorControl/Error/.

See issue #57.
2017-02-07 15:42:38 -05:00
Kenneth Moreland
f23ff9fa49 Fix warnings about assignment operators not being generated
For some reason when VTK-m was being compiled as an accelerator in VTK,
Visual Studio 2013 gave a bunch of warnings about not being able to generate
assignment operators for many classes. This happened for classes with a
const ivar that could not be automatically set. (Automatic copy constructors
are fine on this count.) I'm not sure why these warnings did not happen
when just compiling VTK-m, nor am I sure why they were generated at all as
no code actually used the copy constructors.

This commit fixes the problems by adding a private declaration for assignment
operators that cannot be automatically created. No implementation is
provided, nor should any be needed.
2017-01-10 11:10:38 -07:00
Kenneth Moreland
9ea10317d9 Add passing input domain to transport operator
Previously, the operator for a Transport class took the object being
transported to the execution environment and the size of the output
domain. This change also passes in the control-side argument for the
input domain. This will help check input array sizes as well as make
other potential transformations based on the input domain.
2016-12-19 17:03:15 -07:00