Commit Graph

61 Commits

Author SHA1 Message Date
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
Kenneth Moreland
fdaccc22db Remove exports for header-only functions/methods
Change the VTKM_CONT_EXPORT to VTKM_CONT. (Likewise for EXEC and
EXEC_CONT.) Remove the inline from these macros so that they can be
applied to everything, including implementations in a library.

Because inline is not declared in these modifies, you have to add the
keyword to functions and methods where the implementation is not inlined
in the class.
2016-11-15 22:22:13 -07:00
Robert Maynard
a931b8e295 FunctionInterface and DispatcherBase don't require boost now. 2016-09-23 16:39:20 -04:00
Robert Maynard
aff8c0211f Remove DispatcherBaseDetailInvoke as it unnecessary with C++11. 2016-09-01 10:41:45 -04:00
Robert Maynard
12810165bb Switch over to c++11 type_traits. 2016-08-31 16:11:26 -04:00
Robert Maynard
40896e2bab Allocate the scatter arrays to be proper length.
The previous code would over allocate all the scatter arrays since
it was computing the output range, than using that as the input range.
2016-03-15 13:13:34 -04:00
Robert Maynard
c1560e2d3f Perform less unnecessary copies when deducing a worklets parameters.
One of the causes of the large library size and slow compile times has been
that vtkm has been creating unnecessary copies when not needed. When the
objects being copied use shared_ptr this causes a bloom in library size. I
presume this bloom is caused by the atomic increment/decrement that is
required by shared_ptr.

For testing I used the following example:
```
struct ExampleFieldWorklet : public vtkm::worklet::WorkletMapField
{
  typedef void ControlSignature( FieldIn<>, FieldIn<>, FieldIn<>,
                                 FieldOut<>, FieldOut<>, FieldOut<> );
  typedef void ExecutionSignature( _1, _2, _3, _4, _5, _6 );

  template<typename T, typename U, typename V>
  VTKM_EXEC_EXPORT
  void operator()( const vtkm::Vec< T, 3 > & vec,
                   const U & scalar1,
                   const V& scalar2,
                   vtkm::Vec<T, 3>& out_vec,
                   U& out_scalar1,
                   V& out_scalar2 ) const
  {
    out_vec = vec * scalar1;
    out_scalar1 = scalar1 + scalar2;
    out_scalar2 = scalar2;
  }

  template<typename T, typename U, typename V, typename W, typename X, typename Y>
  VTKM_EXEC_EXPORT
  void operator()( const T & vec,
                   const U & scalar1,
                   const V& scalar2,
                   W& out_vec,
                   X& out_scalar,
                   Y& ) const
  {
  //no-op
  }
};

int main(int argc, char** argv)
{
  std::vector< vtkm::Vec<vtkm::Float32, 3> > inputVec;
  std::vector< vtkm::Int32 > inputScalar1;
  std::vector< vtkm::Float64 > inputScalar2;

  vtkm::cont::ArrayHandle< vtkm::Vec<vtkm::Float32, 3> > handleV =
    vtkm::cont::make_ArrayHandle(inputVec);

  vtkm::cont::ArrayHandle< vtkm::Vec<vtkm::Float32, 3> > handleS1 =
    vtkm::cont::make_ArrayHandle(inputVec);

  vtkm::cont::ArrayHandle< vtkm::Vec<vtkm::Float32, 3> > handleS2 =
    vtkm::cont::make_ArrayHandle(inputVec);

  vtkm::cont::ArrayHandle< vtkm::Vec<vtkm::Float32, 3> > handleOV;
  vtkm::cont::ArrayHandle< vtkm::Vec<vtkm::Float32, 3> > handleOS1;
  vtkm::cont::ArrayHandle< vtkm::Vec<vtkm::Float32, 3> > handleOS2;

  std::cout << "Making 3 output DynamicArrayHandles " << std::endl;
  vtkm::cont::DynamicArrayHandle out1(handleOV), out2(handleOS1), out3(handleOS2);

  typedef vtkm::worklet::DispatcherMapField<ExampleFieldWorklet> DispatcherType;

  std::cout << "Invoking ExampleFieldWorklet" << std::endl;
  DispatcherType dispatcher;

  dispatcher.Invoke(handleV, handleS1, handleS2, out1, out2, out3);

}
```

Original vtkm would generate a binary of size 4684kb and would perform 91
ArrayHandle copies or assignments. With this branch the binary size is
reduced to 2392kb and will perform 36 copies or assignments.
2016-01-19 09:20:49 -05:00
Kenneth Moreland
e6a9c96c96 Adding ScatterCounting 2015-11-06 18:05:20 -07:00
Kenneth Moreland
b0c5a32611 Add Scatter parameters to Invocation.
We are passing in execution objects with the Invocation when the Worklet
is scheduled, but we are not using it yet.
2015-11-06 18:05:20 -07:00
Robert Maynard
9fdc0f09fc Improve the error message for Invoke type mismatch at compile time.
Now that we can skip generating the DynamicTransform code when all the
arguments are statically known, we need a way to produce nicer error messages.
2015-10-26 16:35:02 -04:00
Robert Maynard
54d25fae55 Only perform DynamicTransformCont if at least one parameter is dynamic.
Previously we always ran DynamicTransformCont even if we knew all the types.
By checking for Dynamic types first, we save roughly 3% on the binary size.

This also is a good starting point for a redesign of DynamicTransformCont
2015-10-26 13:48:25 -04:00
Robert Maynard
05d397cbf7 Remove unnecessary template parameters from DispatcherMapField
DispatcherMapField was templated on the device adapter but it
actually doesn't need to be, only BasicInvoke and subsequent
methods need to be templated on the device.
2015-10-22 12:22:59 -04:00
Kenneth Moreland
b15940c1e3 Declare new VTKM_STATIC_ASSERT
This is to be used in place of BOOST_STATIC_ASSERT so that we can
control its implementation.

The implementation is designed to fix the issue where the latest XCode
clang compiler gives a warning about a unused typedefs when the boost
static assert is used within a function. (This warning also happens when
using the C++11 static_assert keyword.) You can suppress this warning
with _Pragma commands, but _Pragma commands inside a block is not
supported in GCC. The implementation of VTKM_STATIC_ASSERT handles all
current cases.
2015-09-17 14:40:39 -06:00
Kenneth Moreland
2ff6576c65 Add third party wrappers around boost macros.
The boost assert macros seem to have an issue where they define an
unused typedef. This is causing the XCode 7 compiler to issue a warning.
Since the offending code is in a macro, the warning is identified with
the VTK-m header even though the code is in boost. To get around this,
wrap all uses of the boost assert that is causing the warning in the
third party pre/post macros to disable the warning.
2015-09-16 23:34:49 -06:00