Commit Graph

3804 Commits

Author SHA1 Message Date
Kenneth Moreland
2dbc45ac08 Merge topic 'fix-cuda-warnings'
6d24343c5 Add exec to ArrayPortalFromIterators constructors
91df12305 Remove VTKM_EXEC modifiers from CPU devices

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1318
2018-07-12 13:18:25 -04:00
Robert Maynard
288719045c Merge topic 'colortable_properly_support_explicit_device_execution'
6dc06423d ColorTable can provide vtkm::exec::Colortable to a specific device

Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1316
2018-07-12 11:06:07 -04:00
Robert Maynard
6dc06423d8 ColorTable can provide vtkm::exec::Colortable to a specific device
Previously it wasn't possible to get a color table transfered
to a specific device.
2018-07-12 10:28:18 -04:00
Kenneth Moreland
6d24343c51 Add exec to ArrayPortalFromIterators constructors
There is no real reason why you cannot construct an
ArrayPortalFromIterators on a device, so go ahead and let that happen.
(This removes some CUDA warnings about calling __host__ from
__device__.)
2018-07-12 08:09:22 -06:00
Kenneth Moreland
91df123055 Remove VTKM_EXEC modifiers from CPU devices
Having VTKM_EXEC on algorithms for CPU devices was problematic because
the algorithms were specific to the CPU, but during a CUDA compile it
would try to compile device code (for no reasons since it was never
called on a device).

Remove these identifiers for the idea that a device implementation knows
specifically what function modifiers to use and does not need the VTK-m
defined catch-alls.
2018-07-11 16:45:30 -06:00
Haocheng LIU
69aafb1d46 Merge topic 'Create-warp-by-scale-workletAndFilter'
56993bc9e Add a warpScalar worklet and filter

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1276
2018-07-11 14:50:32 -04:00
Kenneth Moreland
abfc946f84 Merge topic 'exec-objects-as-alg-sort-compare'
f14021dd8 Shorten code for PrepareArgForExec
3b828608a Support ExecArg behavior in vtkm::cont::Algorithm methods

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1308
2018-07-10 19:01:40 -04:00
Kenneth Moreland
f14021dd84 Shorten code for PrepareArgForExec
By making is_base_of part of PrepareArgForExec, we can shorten not only
the C++ code but also the code that is generated by it.

Also, return && instead of by value when passing through the argument.

Changes thanks to Robert Maynard.
2018-07-10 13:48:20 -06:00
Haocheng LIU
2f57470a13 Merge branch 'upstream-diy' into update-diy
* upstream-diy:
  diy 2018-07-09 (38da7b11)
2018-07-10 13:52:20 -04:00
Haocheng LIU
56993bc9e2 Add a warpScalar worklet and filter
This commit adds a worklet as well as a filter that modify point coordinates by moving points
along point normals by the scalar amount times the scalar factor.
It's a simpified version of the vtkWarpScalar class in VTK. Additionally the filter doesn't
modify the point coordinates, but creates a new point coordinates that have been warped.
2018-07-09 15:33:25 -04:00
Haocheng LIU
a4c4c1931c Add a warpVector worklet and filter
This commit adds a worklet as well as a filter that modify point coordinates by moving points
along a vector by a certain scalar amount.
It's a simpified version of the vtkWarpVector in VTK.
It doesn't modify the point coordinates, but creates a new point coordinates that have been warped.
2018-07-09 14:22:12 -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
Robert Maynard
c631dccfcb Invocation parameters are now non const and can be 'modified'
The invocation parameters need to be non const as we want to
be able to call non-const methods like `PrepareForOutput` on them
from a transport function.

The original implementation abused the fact that everything
could be copied by value and have that work properly. But
when we start introducing virtual classes copying by value of
a base type can cause type slicing.
2018-07-06 14:27:36 -04:00
Kenneth Moreland
3b828608a4 Support ExecArg behavior in vtkm::cont::Algorithm methods
Most of the arguments given to device adapter algorithms are actually
control-side arguments that get converted to execution objects internally
(usually a `vtkm::cont::ArrayHandle`). However, some of the algorithms,
take an argument that is passed directly to the execution environment, such
as the predicate argument of `Sort`. If the argument is a plain-old-data
(POD) type, which is common enough, then you can just pass the object
straight through. However, if the object has any special elements that have
to be transferred to the execution environment, such as internal arrays,
passing this to the `vtkm::cont::Algorithm` functions becomes
problematic.

To cover this use case, all the `vtkm::cont::Algorithm` functions now
support automatically transferring objects that support the `ExecObject`
worklet convention. If any argument to any of the `vtkm::cont::Algorithm`
functions inherits from `vtkm::cont::ExecutionObjectBase`, then the
`PrepareForExecution` method is called with the device the algorithm is
running on, which allows these device-specific objects to be used without
the hassle of creating a `TryExecute`.
2018-07-06 18:57:54 +02:00
Robert Maynard
9238cedcab Merge topic 'ice_nvcc_on_renar'
5ced0da8f Try to ice the ubuntu 17.10 + cuda 9.1 compiler

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1305
2018-07-05 11:36:16 -04:00
Allison Vacanti
be31ce0124 Merge branch 'upstream-taotuple' into fix_msvc2017_tuple_get
* upstream-taotuple:
  taotuple 2018-07-05 (ee1b4af3)
2018-07-05 09:50:31 -04:00
Robert Maynard
5ced0da8f5 Try to ice the ubuntu 17.10 + cuda 9.1 compiler 2018-07-05 09:14:52 -04:00
Robert Maynard
e5090e1289 Make sure the PointLocatorUniform uses the correct runtime device 2018-07-03 17:42:57 -04:00
Robert Maynard
38e0e4c33e Mark PointLocatorUniformGrid constructors as host only 2018-07-03 17:41:56 -04:00
Allison Vacanti
8b0777c4b1 Merge topic 'cuda_32bit_id_fix'
1751b4932 Fix build error and host/device warnings on 32-bit builds with CUDA.
768086a0f Merge branch 'upstream-taotuple' into cuda_32bit_id_fix
d4b38a78a taotuple 2018-06-29 (4bd72b5e)
b4110cd5d Update update.sh for taotuple to reflect upstream changes.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1295
2018-07-03 10:47:55 -04:00
Haocheng LIU
6bc579e461 Merge topic 'Allow-NDHistogram-to-take-custom-type'
b47b1f9ae Allow NDHistogram to take custom type

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Sujin Philip <sujin.philip@kitware.com>
Merge-request: !1299
2018-07-02 17:10:57 -04:00
Haocheng LIU
b47b1f9ae4 Allow NDHistogram to take custom type 2018-07-02 15:25:49 -04:00
Allison Vacanti
1e6c30b7af Make pair construction rvalue-friendly. 2018-07-02 14:48:06 -04:00
Kenneth Moreland
71486f0864 Fix TestingPointLocatorUniformGrid for double precision
There was an error in TestingPointLocatorUniformGrid in which it was
creating arrays of type vtkm::Float32 and passing them to a worklet that
expected vtkm::FloatDefault. This is corrected.
2018-07-01 11:27:35 +02:00
ayenpure
e2dccee099 Merge branch 'master' of https://gitlab.kitware.com/vtk/vtk-m into spatialsearch 2018-06-30 11:56:33 -06:00
ayenpure
ebdad45d44 Removing VTKM_CONT_EXPORT for the CellLocator 2018-06-30 07:09:41 -06:00
Abhishek Yenpure
e42c270f89 fixing -Wmissing-field-initializers for BoundingIntevalHierarchyNode 2018-06-29 18:58:28 -07:00
Abhishek Yenpure
b5f2e7060e Moving changes of .cxx to .hxx
-For the BoundingIntervalHierarchy CUDA had failures with using
 .cxx file to implement the virtual methods
-Moving the contents to the .hxx file after discussing with Rob
 over email
-Need to still work on the .cxx implementation after merge
2018-06-29 16:11:16 -07:00
ayenpure
e0296a24b6 Fixes for resolving Rob's suggestions
- Releasing GPU resources when not needed.
- Adding VTK_EXEC to TreeNode and SplitProperties constructors
2018-06-29 15:55:23 -07:00
Abhishek Yenpure
dcbb444a74 Warning fixes for BoundingIntervalHierarchy
- Adding VTKM_CONT_EXPORT to CellLocator
- Removing whitespaces
- Fixing warnings of unused typedef
2018-06-29 14:51:50 -06:00
Allison Vacanti
1751b49327 Fix build error and host/device warnings on 32-bit builds with CUDA.
Error: Throwing an exception in CUDA code.
Fix: Change method throwing exception to VTKM_CONT.
New warning: host/device warning in taotuple.
Fix: Markup additional taotuple methods with suppressions.

This also updates our taotuple checkout to match upstream master.
2018-06-29 16:46:18 -04:00
Allison Vacanti
768086a0f3 Merge branch 'upstream-taotuple' into cuda_32bit_id_fix
* upstream-taotuple:
  taotuple 2018-06-29 (4bd72b5e)
2018-06-29 16:45:56 -04:00
Allison Vacanti
ec00a0dced Fix doxygen warning in for exec/PointLocatorUniformGrid. 2018-06-29 15:41:33 -04:00
Allison Vacanti
b4110cd5dd Update update.sh for taotuple to reflect upstream changes. 2018-06-29 14:49:07 -04:00
Abhishek Yenpure
50963290dc Moving BoundingIntervalHierarchy.cxx to device sources
Rectifying print statements in the unit test
2018-06-29 10:46:31 -07:00
ayenpure
bf224dbb7c Adding VTKM_CONT_EXPORT to BoundingIntervalHierarchy 2018-06-29 10:57:25 -06:00
ayenpure
98c45403d3 Removing initialization for Nodes and CellIds for CUDA failure 2018-06-29 10:00:46 -06:00
ayenpure
2714da8f0e Adding changes suggested by Rob on GitLab
- Seperate BoundingIntervalHierarchy.h into
  - BoundingIntervalHierarchy.h
  - BoundingIntervalHierarchy.hxx
  - BoundingIntervalHierarchy.cxx
2018-06-29 09:09:24 -06:00
Abhishek Yenpure
ab81c6335b Changing Float64 to FloatDefault 2018-06-28 13:00:33 -07:00
Abhishek Yenpure
adde6491a1 SplitProperties and TreeNode -Wmissing-field-initializers fixes 2018-06-28 12:46:45 -07:00
Kenneth Moreland
4459ab9174 Merge branch 'master' into 'pointlocator-general-interface'
# Conflicts:
#   vtkm/cont/PointLocatorUniformGrid.h
2018-06-28 12:51:08 -04:00
Robert Maynard
9a1b44cc88 Merge topic 'update_osx_rtti_doc'
46d85146d Update the documentation on why VTKM_ALWAYS_EXPORT exists.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1293
2018-06-28 11:46:38 -04:00
Robert Maynard
46d85146d5 Update the documentation on why VTKM_ALWAYS_EXPORT exists. 2018-06-28 11:46:13 -04:00
Allison Vacanti
404939d6d4 Merge topic 'doxygen_fixes'
1e53b86ea Identifiers in doxygen @param blocks must appear in signature.
828bb3179 Silence expected doxygen warning.
6db40ae49 Newline needed after \file command.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1291
2018-06-28 11:35:43 -04:00
Allison Vacanti
0236595f94 Merge topic 'suppress_hd_warnings_for_atomics'
a8d8b3670 Suppress host/device warnings on CUDA atomics.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1286
2018-06-28 10:05:59 -04:00
Dave Pugmire
912d393b07 Merge topic 'coordSysFilter'
80b12e325 Merge branch 'coordSysFilter' of https://gitlab.kitware.com/dpugmire/vtk-m into coordSysFilter
ab5eeab18 Fixes for making the filters non-templated.
27dade145 Fixes for coordinate systems w/ help from Sujin.
db5ded3a6 Add files for coord sys transform filters.
17087a26a Filter for coordinate system transform.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1284
2018-06-28 05:23:21 -04:00
Abhishek Yenpure
72d4556f7d Adding changes for Ken's review from GitLab
- Reducing the stack allocation for CUDA for the BIH unit test
- Adding changes from Ken's review
- Suppress ptxas stack size warning for BoundingIntervalHierarchy
2018-06-27 22:40:13 -06:00
Kenneth Moreland
51fd4a117d Fix warning about __host__/__device__ on default constructor
The vtkm::exec::PointLocatorUniformGrid has a default constructor. It
was "helpfully" declared as VTKM_EXEC_CONT, but apparently that is the
wrong thing to do for constructors that are set to default.
2018-06-28 05:05:54 +02:00
Kenneth Moreland
6f75cd008b Fix crash in CUDA compiler
Previously when PointLocatorUniformGrid.h was compiled by the CUDA
compiler, the compiler would crash. Apparently during the ptxas
part of the compiler goes into a crazy recursion and runs out of
stack space. This appears to be a long-standing bug in CUDA
(been there for multiple releases) without a clear reason why it
sometimes rears its ugly head. (See for example
https://devtalk.nvidia.com/default/topic/1028825/cuda-programming-and-performance/-ptxas-died-with-status-0xc00000fd-stack_overflow-/)

The problem appears to be when having a doubly or triply nested
loop over a box of values to check in the uniform array. This
appears to fix the problem by converting that to a single for
loop with some index magic to convert that to 3D indices.
2018-06-27 23:52:48 +02:00
Allison Vacanti
1e53b86ea4 Identifiers in doxygen @param blocks must appear in signature. 2018-06-27 14:10:05 -04:00