3c85f7f40 Correct undefined behavior from missing export on ErrorMessageBuffer
e298e05bc Test will now fail if programs such as UBSAN output "runtime error"
feb66f970 Make VTK-m use consistent symbol visibility for benchmarks
b43d61533 Make VTK-m use consistent symbol visibility for tests.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1699
a9e6a919d Correct a subtle difference in the SplitSharp multiple pass algorithm
0f61808da SplitSharpEdges worklet does more algorithm execution on the device.
20fa341a2 SplitSharpEdgesFilter test now checks against a structured dataset
99865f47d Update Wavelet worklet to support a runtime device adapter id
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Allison Vacanti <allison.vacanti@kitware.com>
Merge-request: !1693
Fixes#373
The SplitSharp algorithm has two passes that have the same algorithm,
but does different work with the final computed state. The issue was
that in the original implementation the two algorithms differed slightly.
This refactors the code to have a single implementation of the algorithm,
to avoid this form of issue going forward.
fbbc18590 Use logging to output status of VTKDataSetReaderBase
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1689
2e0f4dd37 Fix floating point error in test
b766d9a92 Improve support of testing recursive types
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1691
The original implementation predates the logging and so just outputted
error and warning messages to stderr. Changed these to the more
appropriate logging mechanism.
The UnitTestCudaArrayHandleFancy was failing because it was
was calling TestValue in a worklet and comparing that to
a value that was generated by calling TestValue on the host.
Because these were called on different architecture, it is
possible that they will not be bit-wise exact. Add a little
bit of tolerance to the check to avoid false failures for
this reason.
HasVecTraits<T> is a template that will resolve to std::true_type if
VecTraits<T> is properly defined and std::false_type otherwise. This is
helpful when defining classes and methods that use VecTraits but you
still want them to compile if VecTraits does not exist.
All C++11 compilers support _Pragma to insert pragmas inline into code
and within macros. All compilers, that is, except for visual studio
because Microsoft has to be contrarian and make life miserable for all
programmers. Instead, you have to use __pragma with visual studio.
401b12bd6 Merge branch 'master' of https://gitlab.kitware.com/vtk/vtk-m into add_polyLine
fea18190f Specialized cases for cell-edge functions on polylines.
d310ec3aa return type is void. Call vertex/line methods, then just return.
d6898b805 Fix cell deriv for polylines and remove print statements.
9157004ac Merge branch 'master' of https://gitlab.kitware.com/vtk/vtk-m into add_polyLine
d6e2e9588 Remove debugging print statements.
b9d109ab3 Fix for CellEdgeFace test. Case is identical to polygon.
d7e793861 Fix compiler warnings. Comment out std::cout usage for testing with cuda.
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1677
42a3602ed Had missed a few unused variables when fixing rebase issues.
63a4e696c vim retabed and removed trailing whitespaces form Unit Test
e0000f35f Should be working now
63165904c Merged with the pull (reword)
5cc1de297 Clearing out unused variables
c7bd79e78 Extract Structured with Offset
56598ebfa Maybe closer
5feadd166 Pushing for Abhishek (Don't merge)
...
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1644
Did a bad rebase to fix old commit message. Propagated merge conflicts
through so things got really confusing. Code looks to be back to normal
now though.
bcaf7d9be ScopedRuntimeDeviceTracker have better controls of setting devices.
4212d0c04 RuntimeDeviceInformation now says the AnyTag exists.
fa03dc664 ScopedRuntimeDeviceTracker requires a device to execute on when constructed.
4020f5198 RuntimeDeviceTracker can't be copied and is only accessible via reference.
e9482018e ScopedRuntimeDeviceTracker has the same API as RuntimeDeviceTracker
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1676
These methods need to be marked inline as they are used by multiple
TU's. When they aren't marked as inline we see random failures in
downstream VTK-m users generally when CUDA is enabled due to ODR
violations.
The ScopedRuntimeDeviceTracker now can force, enable, or disable
devices. Additionally the ScopedRuntimeDeviceTracker and the
RuntimeDeviceTracker handle the DeviceAdapterTagAny robustly
across all methods.
As the RuntimeDeviceTracker is a per thread construct we now make
it explicit that you can only get a reference to the per-thread
version and can't copy it.
63fe0f096 DispatcherReduceByKey uses the scheduling_range like all other dispatchers
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1674
The 8x8x8 is a better launch strategy for most VTK-m kernels.
The current problem is that a couple of VTK-m kernels use a
high number of registers and this number of threads combines to
require too many registers.
What we should do in the longer run is have more controls over
kernel launches on a per kernel basis. This will require VTK-m
to extract the number of registers being used by each kernel
41b8236a2 For GCC 4.8.4 'half' shadows a global variable with that name
770912f99 Correct compiler issues found with GCC 4.8.5 + CUDA 9.2 on summit
b248b2c93 Correct unused-parameter warnings from defaulted methods.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Allison Vacanti <allison.vacanti@kitware.com>
Merge-request: !1666
The consistent API for control to execution memory transfers is
the ArrayHandle class. Previously the tests would verify memory
transfer by calling the ArrayManagerExecution class directly. This
is problematic as the class isn't used by ArrayHandle<T, StorageBasic>.
When TransferInfo is given memory from VirtualObjectTransferShareWithControl
it doesn't have a bound function ptr for the destruction. In those cases
we need to make sure the HostCopyOfDevice is properly deleted, otherwise
we will cause a memory leak.
bdabfbe11 Make sure ArrayPortalUniformPointCoordinates constructor is explicit
b3d951b50 vtkm::Range Include function now requires half as many min/max calls
ddaa0df26 ArrayHandleVirtualCoordinates now calls the proper parent constructor
61e800379 Make sure all execution side CellLocator objects have explicit destructors
307898ff6 Cleanup the CellLocatorBoundingIntervalHierarchy.cxx style.
0f31c69f3 Remove unnecessary constructor from ParameterContainer
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1657
Previously it was calling the ArrayHandle<T,StorageTagVirtual>
constructor and not the ArrayHandleVirtual constructor which
generated a warning with some compilers
There is a small section in the code generated from Math.h.in that is
subject to clang formatting. A recent change reformatted that bit of
Math.h, so we need to update Math.h.in accordingly.
ff687016e For VTK-m libs all includes of DeviceAdapterTagCuda happen from cuda files
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1648
It is very easy to cause ODR violations with DeviceAdapterTagCuda.
If you include that header from a C++ file and a CUDA file inside
the same program we an ODR violation. The reasons is that the C++
versions will say the tag is invalid, and the CUDA will say the
tag is valid.
The solution to this is that any compilation unit that includes
DeviceAdapterTagCuda from a version of VTK-m that has CUDA enabled
must be invoked by the cuda compiler.
d8cc067ca Remove DeviceAdapterError as it isn't needed any more.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1649
Fixes#277
DeviceAdapterError existed to make sure that the default device adapter
template was being handled properly. Since the default device adapter doesn't
exist, and nothing is templated over it we can now remove DeviceAdapterError.
9c2920072 UnitTestBoundingIntervalHierarchy handles systems under load better
671c1df5c Timer logs the proper device name when called with an invalid device
d3d66a331 GameOfLife example always uses the proper device adapter
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1645
The UnitTestBoundingIntervalHierarchy has historically had problems
when the machine is already under-load when the algorithm is executed.
By limiting the number of openMP threads the test uses we can
reduce the amount of CPU time slicing that this test causes.
VTK-m now offers a more GPU aware set of defaults for kernel scheduling.
When VTK-m first launches a kernel we do system introspection and determine
what GPU's are on the machine and than match this information to a preset
table of values. The implementation is designed in a way that allows for
VTK-m to offer both specific presets for a given GPU ( V100 ) or for
an entire generation of cards ( Pascal ).
Currently VTK-m offers preset tables for the following GPU's:
- Tesla V100
- Tesla P100
If the hardware doesn't match a specific GPU card we than try to find the
nearest know hardware generation and use those defaults. Currently we offer
defaults for
- Older than Pascal Hardware
- Pascal Hardware
- Volta+ Hardware
Some users have workloads that don't align with the defaults provided by
VTK-m. When that is the cause, it is possible to override the defaults
by binding a custom function to `vtkm::cont::cuda::InitScheduleParameters`.
As shown below:
```cpp
ScheduleParameters CustomScheduleValues(char const* name,
int major,
int minor,
int multiProcessorCount,
int maxThreadsPerMultiProcessor,
int maxThreadsPerBlock)
{
ScheduleParameters params {
64 * multiProcessorCount, //1d blocks
64, //1d threads per block
64 * multiProcessorCount, //2d blocks
{ 8, 8, 1 }, //2d threads per block
64 * multiProcessorCount, //3d blocks
{ 4, 4, 4 } }; //3d threads per block
return params;
}
vtkm::cont::cuda::InitScheduleParameters(&CustomScheduleValues);
```
661fb64de AtomicInterfaceControl functions are marked with VTKM_SUPPRESS_EXEC_WARNINGS
0c70f9b9a Add BitFieldIn/Out/InOut worklet signature tags.
a66510e81 Add ArrayHandleBitField, a boolean-valued AH backed by a BitField.
56cc5c3d3 Add support for BitFields.
d01b97382 Allow VTKM_SUPPRESS_EXEC_WARNINGS to be used inside macros.
2f2ca9370 Add bit operations FindFirstSetBit and CountSetBits to Math.h.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1629
BitFields are:
- Stored in memory using a contiguous buffer of bits.
- Accessible via portals, a la ArrayHandle.
- Portals operate on individual bits or words.
- Operations may be atomic for safe use from concurrent kernels.
The new BitFieldToUnorderedSet device algorithm produces an ArrayHandle
containing the indices of all set bits, in no particular order.
The new AtomicInterface classes provide an abstraction into bitwise
atomic operations across control and execution environments and are used
to implement the BitPortals.
When reducing an input type that differs from the output type
you need to write a custom binary operator that also implements
how to do the unary transformation.
0130088b8 Suppress more self-assign-overloaded warnings found by clang
fa5455854 UnitTestArrayPortalValueReference doesn't warn when compiled with appleclang
5cc0d03f6 Merge branch 'upstream-diy' into clang_warnings
dbd3781d5 diy 2019-04-09 (f7a68da4)
5c2f2ebce suppress warnings found by Wself-assign-overloaded
77426f044 Correct casting long to long long warning from clang.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1635
After the change to `ArrayHandleVirtual` where it became a subclass of
`vtkm::cont::ArrayHandle`, a few extra changes are required.
1. Functions with `ArrayHandleVirtual` as parameters will not be callable
with the superclass `ArrayHandle`. This is fixed by changing the argument
types to the superclass.
2. Add Serialization classes specializations for the superclass, as "is-a"
relation is not considered for class template parameters.
18f3b0dcc Document all the free functions for VariantArrayHandle
5180d6a73 DynamicCellSet has the same free function support as VariantArrayHandle
8f545662d CellSet uses NewInstance to be consistent with the rest of VTK-m
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Sujin Philip <sujin.philip@kitware.com>
Merge-request: !1620
f1056affa Move select functions to host only to remove host/device suppressions
4f2156dfa Thrust detail::aligned_reinterpret_cast doesn't warn now
f4840618c Make sure ThrustPatches is included before thrust.
b2bbd66e6 Merge branch 'upstream-taotuple' into update_taoo
4ec6fc812 taotuple 2019-04-03 (8e70fa8a)
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Allison Vacanti <allison.vacanti@kitware.com>
Merge-request: !1607
566e220ea Suppress dashboard warnings
f20d7e788 Document the changes that are part of this MR.
f78e763be Add CellLocatorGeneral
c6bead838 Rename CellLocatorTwoLevelUniformGrid to CellLocatorUniformBins
ee838b829 Stylistic changes to CellLocators to match VTK-m
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1615
The CastAndCall for ArrayHandleVirtual was always a bad idea.
The ArrayHandleVirtual is not considered a dynamic type by the
dispatch engine so it would never have tried to use this specialization.
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
A general purpose `CellLocator` that should work with any type of dataset.
Internally, it tries to chose one of the existing cell locators that would be
optimal for the input data.
Also removes `CellLocatorHelper`.
18ff6681f Less vtkm_cont cxx files bring in the cuda device
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1612
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.
The connected components filters had some simple issues that were
inconsistent with the operation of other filters.
* Set the output to the OutputFieldName defined in the filter rather
than hardcoding it to "component". The default value is still
"component", but now you can change it.
* Make sure the output field association is correct. The field
association for CellSetConnectivity is always cell and the field
association for ImageConnectivity is always point.
* Check that the field association for the input field of
ImageConnectivity is point. The filter will fail if it is not.
* Make both filters inherit from FilterCell instead of FilterField.
The superclasses are similar but the latter allows the user to set the
active cell set correctly.
* Properly get the cell set specified by GetActiveCellSetIndex.
Previously it was set to the active coordinates, which is wrong.
CellInterpolationHelper.h(43): warning #2913-D: __device__ annotation is
ignored on a function("~CellInterpolationHelper") that is explicitly
defaulted on its first declaration.
0581b368f Fix issues with importing installed vtkm::cuda after refactor.
ed4374b05 diy 2019-03-13 (e8e68c7c)
b85dcf229 Add more vtkm_cont compilation units to device_sources.
9014de0eb Suppress new host/device warnings for cuda 10.1.
04a464cc4 Update cuda detection to create vtkm::cuda as an interface library.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1582
Previously we just took the optionparser.h file and stuck it right in
our source code. That was problematic for a variety of reasons.
1. It incorrectly assigned our license to external code.
2. It made lots of unnecessary changes to the original source (like
reformatting).
3. It made it near impossible to track patches we make and updates to
the original software.
Instead, use the third-party system to track changes to optionparser.h
in a different repository and then pull that into ours.
When a library requires reading some command line arguments through a
function like Initialize, it is typical that it will parse through
arguments it supports and then remove those arguments from argc and argv
so that the remaining arguments can be parsed by the calling program.
VTK-m's initialize did not do that, so add that functionality.
The RuntimeDeviceTracker had grown organically to handle multiple
different roles inside VTK-m. Now that we have device tags
that can be passed around at runtime, large portions of
the RuntimeDeviceTracker API aren't needed.
Additionally the RuntimeDeviceTracker had a dependency on knowing
the names of each device, and this wasn't possible
as that information was part of its self. Now we have moved that
information into RuntimeDeviceInformation and have broken
the recursion.
e1f5c4dd9 Modify VariantAH::AsVirtual to cast to new ValueType if needed.
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1585
525317249 Merge branch 'master' of https://gitlab.kitware.com/vtk/vtk-m into temporal_particle_advect
2c1d3ee51 Updating example with new interfaces
cb8ea7dff Adding missing APIs
a7510a3c0 Fixing out of seq init warning
f3f2469b8 Adding direct constructor and vtkm::Lerp
2daf18956 Updating copyright for Unit Test
eca618e10 Fixing typos
be80dcc1a Adding changes for temporal advection, and adding test
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Dave Pugmire <dpugmire@gmail.com>
Merge-request: !1590
E.g:
```
ArrayHandle<Float64> doubleArray;
VariantArrayHandle varHandle{doubleArray};
ArrayHandleVirtual<Float32> = varHandle.AsVirtual<Float32>();
```
If there is a loss in range and/or precision, a warning is logged. If
the ValueTypes are Vecs with mismatched widths, an ErrorBadType is thrown.
Internally, an ArrayHandleCast is used between the VariantArrayHandle's
stored array and the ArrayHandleVirtual.
b81b2f817 Make CellClassification more clear
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Dave Pugmire <dpugmire@gmail.com>
Merge-request: !1597
a09bb9eca Add a new test for CellSet API
8868fb989 Improve CellSet API
Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1589
Previously, the CellClassification enum was stored in a header file
named GhostCell.h, which made it hard to find and obscured its purpose.
Moved it to an appropriately named file. Also, renamed the DUPLICATE
field to GHOST to make its intention more clear.
