Commit Graph

154 Commits

Author SHA1 Message Date
Allison Vacanti
afe1bd12dd Add ScanExtended device algorithm.
This behaves just like `ScanExclusive`, but rather than returning the
total sum, it is appended to the end of the output array.

This is in preparation for the CellSetExplicit refactoring described in
issue #408.
2019-09-03 15:02:41 -04:00
Allison Vacanti
884616788a Simplify and extend AtomicArray implementation.
- Use AtomicInterface to implement device-specific atomic operations.
- Remove DeviceAdapterAtomicArrayImplementations.
- Extend supported atomic types to include unsigned 32/64-bit ints.
- Add a static_assert to check that AtomicArray type is supported.
- Add documentation for AtomicArrayExecutionObject, including a CAS
  example.
- Add a `T Get(idx)` method to AtomicArrayExecutionObject that does
  an atomic load, and update existing CAS usage to use this instead
  of `Add(idx, 0)`.
2019-08-23 15:40:37 -04:00
Allison Vacanti
112024dae2 Fix CUDA shfl usage.
There was a bug in the implementations of CountSetBits and
BitFieldToUnorderedSet.
2019-08-01 10:57:57 -04:00
Kenneth Moreland
0be50c119d Update VTK-m code to use new Vec aliases
Should make the code easier to read.
2019-07-31 12:55:40 -06:00
Allison Vacanti
f370857c15 Add CountSetBits and Fill device algorithms. 2019-06-25 11:30:39 -04:00
Robert Maynard
4020f51988 RuntimeDeviceTracker can't be copied and is only accessible via reference.
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.
2019-05-20 11:43:05 -04:00
Robert Maynard
065d117838 Testing Device Adapter now uses ArrayHandle for all device transfers
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>.
2019-04-30 13:50:08 -04:00
Robert Maynard
d8cc067caa Remove DeviceAdapterError as it isn't needed any more.
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.
2019-04-18 15:09:57 -04:00
nadavi
fbcea82e78 conslidate the license statement 2019-04-17 10:57:13 -06:00
Robert Maynard
a5dbe1ece3 Merge topic 'bitfields'
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
2019-04-11 12:32:03 -04:00
Allison Vacanti
56cc5c3d3a Add support for BitFields.
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.
2019-04-11 08:27:17 -04:00
Robert Maynard
89ec4aae2f Reduction on CUDA handles different input and output types better
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.
2019-04-10 14:44:44 -04:00
Robert Maynard
ae11e115a0 RuntimeDeviceTracker: Remove Global from names 2019-03-22 08:53:26 -07:00
Haocheng LIU
415252c662 Introduce asynchronous and device independent timer
The timer class now is asynchronous and device independent. it's using an
similiar API as vtkOpenGLRenderTimer with Start(), Stop(), Reset(), Ready(),
and GetElapsedTime() function. For convenience and backward compability, Each
Start() function call will call Reset() internally and each GetElapsedTime()
function call will call Stop() function if it hasn't been called yet for keeping
backward compatibility purpose.

Bascially it can be used in two modes:

* Create a Timer without any device info. vtkm::cont::Timer time;

  * It would enable timers for all enabled devices on the machine. Users can get a
specific elapsed time by passing a device id into the GetElapsedtime function.
If no device is provided, it would pick the maximum of all timer results - the
logic behind this decision is that if cuda is disabled, openmp, serial and tbb
roughly give the same results; if cuda is enabled it's safe to return the
maximum elapsed time since users are more interested in the device execution
time rather than the kernal launch time. The Ready function can be handy here
to query the status of the timer.

* Create a Timer with a device id. vtkm::cont::Timer time((vtkm::cont::DeviceAdapterTagCuda()));

  * It works as the old timer that times for a specific device id.
2019-02-05 12:01:56 -05:00
Robert Maynard
d6f66d17a3 Testing run methods now take argc/argv to init logging/runtime device
`vtkm::cont::testing` now initializes with logging enabled and support
for device being passed on the command line, `vtkm::testing` only
enables logging.
2019-01-17 13:16:27 -06:00
Haocheng LIU
bb06717803 Make RuntimeDeviceInformation class template independent
By making RuntimeDeviceInformation class template independent, vtkm is
able to detect
device info at runtime with a runtime specified deviceId. In the past
it's impossible
because the CRTP pattern does not allow function overloading(compiler
would complain
that DeviceAdapterRuntimeDetector does not have Exists() function
defined).
2018-11-01 14:55:33 -04:00
Robert Maynard
45422478cf Refactor VirtualObjectHandle to support new virtual design 2018-10-15 17:38:54 -04:00
luz.paz
d5beb69ec1 Misc. typos
Found via `codespell`
2018-10-04 10:30:33 -04:00
Robert Maynard
554bc3d369 At runtime TryExecute supports a specific deviceId to execute on.
Instead of always using the first enabled device, now TryExecute
can be told which device at runtime to use.
2018-08-07 17:22:18 -04:00
Robert Maynard
3533975694 Remove usages of std::vector from OpenMP reduction algorithm
The OpenMP Device Reduction algorithm previously used a std::vector<T>
to store the reduction results of each thread. This caused problems
when T=bool as the types became a proxy type which isn't usable
with vtkm BinaryOperators.

Additionally by fixing this issue in the FunctorsOpenMP we
can remove a workaround in FunctorsGeneral that caused
compile failures when using complex BinaryOperators
such as MinAndMax.
2018-08-06 13:08:33 -04:00
Sujin Philip
259d670ab5 Merge topic 'cuda-per-thread-streams-2'
06dee259f Minimize cuda synchronizations

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1288
2018-07-25 15:07:39 -04:00
Matthew Letter
f892695f16 fixing so wierd merging issue
somehow a rebase removed code that shouldnt have been removed so I added the code back in, gonna see if I can squash this commit later
2018-07-03 14:28:29 -06:00
Sujin Philip
06dee259f7 Minimize cuda synchronizations
1. Have a per-thread pinned array for cuda errors
2. Check for errors before scheduling new tasks and at explicit sync points
3. Remove explicit synchronizations from most places

Addresses part 2 of #168
2018-07-03 14:19:06 -04:00
Matthew Letter
db1c9bfeee Change the namespacing of atomic array
moved the execution object to vtkm::exec and AtomicArray to vtkm::cont
2018-07-02 16:25:48 -06:00
Matthew Letter
6b9cb536a8 added new lines at the end of bot recommended files 2018-07-02 16:21:00 -06:00
Matthew Letter
3c9e7f049f updated Atomic array to use new excution object paradigm 2018-07-02 16:21:00 -06:00
Allison Vacanti
d833d9285f Merge topic 'openmp'
183bcf109 Add initial version of an OpenMP backend.
7b5ad3e80 Expand device scheduler test to check for overlap.
e621b6ba3 Generalize the TBB radix sort implementation.
d60278434 Specialize swap for ArrayPortalValueReference types.
761f8986f Cache inputs to SSI::Unique benchmark.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1099
2018-05-31 16:49:15 -04:00
Shreeraj Jadhav
6e099b7c09 Fix issue #110 TestingDeviceAdapter uniform_int_distribution
This fix addresses issue #110. Test functions in TestingDeviceAdapter have
been updated to test large arrays through random indices using
std::uniform_int_distribution rather than testing every 100th or 50th value.
2018-05-30 11:04:02 -04:00
Allison Vacanti
7b5ad3e80c Expand device scheduler test to check for overlap.
Add a test that fails if DeviceAdapterAlgorithm::Schedule
executes the same instance twice. Useful for checking block
scheduling algorithms.
2018-05-29 14:49:44 -04:00
Allison Vacanti
93506d25e2 Change function signatures to use 'using' aliases.
Also cleaned up some lingering type typedefs.
2018-05-25 17:18:41 -04:00
Robert Maynard
22ea58335a iVTK-m CUDA backend doesn't use thrust::cuda::pointer any more.
This was removed as CUDA 9.0 on MSVC has issues where CUB/Thrust
would fail to compile when given these types.
2018-02-02 08:33:17 -05:00
Sujin Philip
8c242cef91 Switch from faux to true virtuals 2017-11-06 15:25:29 -05:00
Robert Maynard
27d1275249 Correct issues on windows with debug tests timing out.
The tests actually raised a std assert which was causing a timeout as it
required user intervention to click through.
2017-10-31 13:35:13 -04:00
Kenneth Moreland
e50ec6b667 Fix type error in ScanExclusiveByKey
The implementation of ScanExclusiveByKey in
DeviceAdapterAlgorithmGeneral by shifting values in the input values
array and then calling ScanInclusiveByKey. However, the temporary
shifted values array was created using the key type instead of the
values type. This caused a compile error when the keys and values had
different types.
2017-10-24 16:12:30 -06:00
Allison Vacanti
1018d981a0 Check for overlap in CopySubRange.
Some parallel copy implementations will not handle this sanely.
2017-10-11 16:52:32 -04: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
f68635941e Convert VTK-m over to use 'using' instead of 'typedef' 2017-08-17 10:47:25 -04:00
Robert Maynard
b85cdd9080 Convert VTK-m over to use 'using' instead of 'typedef' 2017-08-07 14:05:43 -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
Kitware Robot
4ade5f5770 clang-format: apply to the entire tree 2017-05-25 07:51:37 -04:00
Kitware Robot
efbde1d54b clang-format: sort include directives 2017-05-18 12:59:33 -04:00
Sujin Philip
e9898cc5cf Merge topic 'virtual-methods'
4049b5b2 Add ClipWithImplicitFunction Filter
82d02e46 Modify ImplicitFunctions to use Virtual Methods
968960c1 Add Virtual Methods Framework

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !750
2017-05-02 16:12:04 -04:00
Sujin Philip
968960c1a1 Add Virtual Methods Framework 2017-05-01 16:51:42 -04:00
Li-Ta Lo
5c735a385c this should resovle all the type conversion warnings 2017-04-25 10:12:11 -06:00
Li-Ta Lo
d8b02329c5 Merge branch 'scanbykey' of gitlab.kitware.com:ollielo/vtk-m into scanbykey 2017-04-20 15:13:05 -06:00
Li-Ta Lo
58ef7c8de0 one more attemp to get the data type right 2017-04-20 15:09:38 -06:00
Li-Ta Lo
22d0e35521 attempt to fix warning on type conversion 2017-04-20 10:22:43 -06:00
Li-Ta Lo
ded4583aca attempt to fix warning on type conversion 2017-04-20 08:54:09 -06:00
Li-Ta Lo
897b2f0f63 add tests for 1, 2 and ARRAY_SIZE elements for both ScanInclusiveByKey and ScanExclusiveByKey 2017-04-19 13:38:28 -06:00
Li-Ta Lo
c05a2c323e Merge branch 'scanbykey' of gitlab.kitware.com:ollielo/vtk-m into scanbykey 2017-04-18 16:13:51 -06:00