Commit Graph

114 Commits

Author SHA1 Message Date
Vicente Adolfo Bolea Sanchez
d348b11183 Enable shared CUDA builds when not compiling virtuals
The reason why we did not support shared libraries when CUDA compiles
were on is that virtual methods require a special linking step to pull
together all virtual methods that might be called. I other words, you
cannot call a virtual CUDA method defined inside a library. This
requirement goes away when virtuals are removed.

Also removed the necessity of using seprable compilation with cuda.
Again, this is only needed when a CUDA function is defined in one
translation unit and used in another. Now we can enforce that all
translation units define their own CUDA functions.

Also, suppress warnings in cuda/internal/ExecutionPolicy.h

This is where we call the thrust algorithms. There must be some loop
where it, on some code path, calls back a host function. This must be in
an execution path that never happens. The thrust version has its own
suppress, but that does not seem to actually suppress the warning (it
just means that the warning does not tell you where the actual call is).

Get around the problem by suppressing the warnings in VTK-m.

Co-authored-by: Kenneth Moreland <morelandkd@ornl.gov>
Co-authored-by: Vicente Adolfo Bolea Sanchez <vicente.bolea@kitware.com>

Signed-off-by: Vicente Adolfo Bolea Sanchez <vicente.bolea@kitware.com>
2021-08-24 13:14:58 -04:00
Ben Boeckel
4c7fe13a98 cmake: avoid adding testing directories if testing is disabled
Some testing directories have side effects such as installing headers or
compiling code that ultimately doesn't end up getting used.
2021-06-01 18:40:40 -04:00
Kenneth Moreland
c55d15f397 Deprecate ArrayHandle::ExecutionTypes
The newer version of `ArrayHandle` no longer supports different types of
portals for different devices. Thus, the `ReadPortalType` and
`WritePortalType` are sufficient for all types of portals across all
devices.

This significantly simplifies supporting execution objects on devices,
and thus this change also includes many changes to various execution
objects to remove their dependence on the device adapter tag.
2021-02-08 12:17:37 -07:00
Kitware Robot
cf0cdcf7d1 clang-format: reformat the repository with clang-format-9 2020-08-24 14:01:08 -04:00
Kenneth Moreland
56bec1dd7b Replace basic ArrayHandle implementation to use Buffers
This encapsulates a lot of the required memory management into the
Buffer object and related code.

Many now unneeded classes were deleted.
2020-06-25 14:02:26 -06:00
Kenneth Moreland
a44392027a Fix compile error in UnitTestTaskStrided.cu from changes in Fetch
A recent change removed the thread indices parameters from the arguments
to the `Fetch` template. Somehow, an instance of using the old template
in the CUDA task strided tests snuck through the dashboard tests.
Correct that.
2020-04-27 17:47:37 -06:00
Vicente Adolfo Bolea Sanchez
d0396e2a40 relaxes ThreadIndicesType across multiple worklets fetchs
This change is needed for being able to use different thread indices types
without changing Fetchs. Basically decoupling those two areas.

1. This commit removes concrete specialization instantiations of
   ThreadIndicesTypes in all of the Fetch's specializations.

2. It also moves the ThreadIndicesType template parameter from the Fetch
   struct to a template parameter in their methods Load/Store.

Signed-off-by: Vicente Adolfo Bolea Sanchez <vicente.bolea@kitware.com>
2020-04-24 17:39:31 -04:00
Kenneth Moreland
4f9fa08fa1 Remove ArrayHandleStreaming capabilities
The `ArrayHandleStreaming` class stems from an old research project
experimenting with bringing data from an `ArrayHandle` in parts and
overlapping device transfer and execution. It works, but only in very
limited contexts. Thus, it is not actually used today. Plus, the feature
requires global indexing to be permutated throughout the worklet
dispatching classes of VTK-m for no further reason.

Because it is not really used, there are other more promising approaches
on the horizon, and it makes further scheduling improvements difficult,
we are removing this functionality.
2020-03-24 15:01:56 -06:00
Robert Maynard
8377806778 Merge topic 'introduce_mapfield_3d_scheduling'
1f1688483 Initial infrastructure to allow WorkletMapField to have 3D scheduling

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1938
2020-02-27 08:02:52 -05:00
Kenneth Moreland
ec34cb56c4 Use new ways to get array portal in control environment
Also fix deadlocks that occur when portals are not destroyed
in time.
2020-02-26 13:10:46 -07:00
Robert Maynard
1f1688483e Initial infrastructure to allow WorkletMapField to have 3D scheduling 2020-02-25 15:23:41 -05:00
Kenneth Moreland
3671cbe168 Fix token issues with CUDA 2020-02-25 09:39:30 -07:00
Allison Vacanti
813f5a422f Fixup custom portal iterator logic.
The convenience functions `ArrayPortalToIteratorBegin()` and
`ArrayPortalToIteratorEnd()` wouldn't detect specializations of
`ArrayPortalToIterators<PortalType>` since the specializations aren't
visible when the `Begin`/`End` functions are declared.

Since the CUDA iterators rely on a specialization, the convenience
functions would not compile on CUDA.

Now, instead of specializing `ArrayPortalToIterators` to provide custom
iterators for a particular portal, the portal may advertise custom
iterators by defining `IteratorType`, `GetIteratorBegin()`, and
`GetIteratorEnd()`. `ArrayPortalToIterators` will detect such portals
and automatically switch to using the specialized portals.

This eliminates the need for the specializations to be visible to the
convenience functions and allows them to be usable on CUDA.
2019-12-17 15:39:51 -05:00
Robert Maynard
5c56ff945f Label tests which exercise a given Device Adapter
This allows developers an easy way to run all OpenMP tests
2019-09-13 15:52:40 -04:00
Robert Maynard
bc78126876 CUDA 10.1 didn't fix stateless_resource_allocator warnings
The patch to fix this warning from thrust has not be approved
by upstream thrust. So keep suppressing it in CUDA 10.1 update 1
2019-09-12 17:36:12 -04:00
Allison Vacanti
0510c6f053 Patch thrust to be happy with aligned_r_cast<long> and friends. 2019-08-12 15:21:11 -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
Robert Maynard
6775685c72 Update ThrustPatches to be aware of issues fixed in thrust 1.9.6 2019-05-31 10:08:47 -04:00
Robert Maynard
9937f51fe1 Merge topic 'thrust_patch_order'
63c931e63 Correct location of ThrustPatches which clang formatter moved

Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !1653
2019-04-23 17:37:05 -04:00
Robert Maynard
6fafcf01fd correct compile issues caused by clang formatting.
A couple of tests require vtkm/testing/Testing to be the first include
2019-04-23 17:16:26 -04:00
Robert Maynard
63c931e639 Correct location of ThrustPatches which clang formatter moved 2019-04-23 15:02:58 -04:00
Robert Maynard
ff687016ee For VTK-m libs all includes of DeviceAdapterTagCuda happen from cuda files
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.
2019-04-22 10:39:54 -04:00
nadavi
fbcea82e78 conslidate the license statement 2019-04-17 10:57:13 -06:00
Robert Maynard
1d980ed147 Thrust Patches tried to apply CUDA 10.1 patches to CUDA 10.0 2019-04-10 14:46:56 -04:00
Robert Maynard
f05940aaca ThrustPatches now only patches fixes for relevant cuda versions
Rather than always patch Thrust functions, we now only patch
based on the Thrust version.
2019-04-10 11:18:42 -04:00
Robert Maynard
20d6201a98 Suppress thrust::mr::stateless_resource_allocator host/device warnings 2019-04-10 09:46:34 -04:00
Robert Maynard
58884a7299 Fix warnings found when cuda is the cuda host compiler 2019-04-05 15:40:03 -04:00
Robert Maynard
4f2156dfaf Thrust detail::aligned_reinterpret_cast doesn't warn now
We specialize aligned_reinterpret_cast inside vtk-m to fix the
issues related to missing __host__ __device__ markups on the
function.
2019-04-03 12:48:33 -04:00
Robert Maynard
f4840618cf Make sure ThrustPatches is included before thrust. 2019-04-03 08:51:05 -04:00
Kenneth Moreland
191d6e5580 Add Mask capabilities to worklets
Mask objects allow you to specify which output values should be
generated when a worklet is run. That is, the Mask allows you to skip
the invocation of a worklet for any number of outputs.
2019-02-25 08:58:39 -07:00
Kenneth Moreland
1ca55ac319 Add specialized operators for ArrayPortalValueReference
The ArrayPortalValueReference is supposed to behave just like the value
it encapsulates and does so by automatically converting to the base type
when necessary. However, when it is possible to convert that to
something else, it is possible to get errors about ambiguous overloads.
To avoid these, add specialized versions of the operators to specify
which ones should be used.

Also consolidated the CUDA version of an ArrayPortalValueReference to the
standard one. The two implementations were equivalent and we would like
changes to apply to both.
2019-02-20 13:33:55 -07:00
Robert Maynard
5508d17c31 Merge topic 'correct_broken_install'
24e71d251 VTK-m yet again has properly installed headers.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Kenneth Moreland <kmorel@sandia.gov>
Merge-request: !1525
2019-01-24 14:59:41 -05:00
Robert Maynard
24e71d251b VTK-m yet again has properly installed headers.
Fixes the install issues mentioned in #342
2019-01-24 14:26:40 -05:00
Allison Vacanti
03fc7b66d0 Add VTKM_CUDA_DEVICE_PASS preprocessing definition.
This is only set while compiling device code, and is useful
for code that needs different implementations on devices (e.g.
they call CUDA device intrinsics, etc).
2019-01-24 11:23:45 -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
Robert Maynard
4ec5bae02d Remove VTK-m TestBuild infrastructure
The purpose of the TestBuild infrastructure was to confirm that
VTK-m didn't have any lexical issues when it was a pure header
only project. As we now move to have more compiled components
the need for this form of testing is mitigated. Combined
with the issue of TestBuilds causing MSVC issues, we should
just remove this infrastructure.
2019-01-16 10:04:33 -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
Shreeraj Jadhav
947496550e constexpr construction for Vec classes
Vec class objects can now be constructed during compile-time
as constant expressions by calling Vec( T, ... ) constructors
or through brace-initialization.
Constant expression using fill constructor and nested vectors
of sizes greater than 4 are not supported yet.
Changes made to WrappedOperators.h for resolving overload
ambiguities in Vec construction and typecasting.
Appropriate test cases were added to UnitTestTypes.cxx.
Addresses issue #199.
2018-06-20 14:15:29 -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
Sujin Philip
2f92e40503 Suppress false positive cuda warnings
Suppresses "declared but never referenced" warnings.
2018-05-09 13:17:37 -04:00
Robert Maynard
b56894dd09 Move VTK-m Cuda backend over to a grid-stride iteration pattern.
This allows for easier host side logic when determining grid and block
sizes, and allows for a smaller library side by moving some logic
into compiled in functions.
2018-04-30 17:29:26 -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
505e7aa1ec VTK-m now has defines for the CUDA version even when not using nvcc.
This is needed so so that the CUDA and C++ compiler generate the same code
when scanning a shared header but generating different translation units
2018-02-26 16:38:26 -05:00
Robert Maynard
41d968f68d Cuda ExecutionPolicy when using CUDA 7.5 is aware we use raw pointer now
The failure was caused by not updating CUDA 7.5 code paths when we
removed the usage of ::thrust::cuda::pointer.
2018-02-26 16:37:57 -05:00
Robert Maynard
e630ac5aa4 Merge branch 'master' into vtk-m-cmake_refactor 2018-02-23 14:52:00 -05: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
Robert Maynard
0660c67fef Merge branch 'master' into vtk-m-cmake_refactor 2018-01-16 15:42:28 -05:00
Robert Maynard
004bfe7b12 Prefer using existence of targets when looking for TBB/CUDA support 2018-01-08 14:00:57 -05:00
Robert Maynard
93bc0198fe Suppress false positive warnings about calling host device functions. 2018-01-02 10:40:49 -05:00