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.
5960b8abf Suppress nvlink warnings about virtual methods not used
0af017b03 Move virtual methods of other CellLocators to vtkm_cont
e87864b0e Put CellLocatorBoundingIntervalHierarchy in vtkm_cont library
Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1594
These changes caused some warnings in clang to show up based on virtual
methods in other cell locators. Hence, the rest of the cell locators
have also had some of their code moved to vtkm_cont.
All of the methods in CellLocatorBoundingIntervalHierarchy were listed in
header files. This is sometimes problematic with virtual methods. Since
everything implemented in it can just be embedded in a library, move the
code into the vtkm_cont library.
When benchmarking the VTK-m algorithms on Summit I discovered
that our scheduling choices aren't optimal for the hardware.
This is a short term fix where we select good numbers for
Summit, and in the future make the defaults controllable
by the calling programming and/or environment variables.
Performance numbers can be found at:
https://gitlab.kitware.com/snippets/755