Commit Graph

89 Commits

Author SHA1 Message Date
Kenneth Moreland
ec0adf8b16 Change interface of ArrayTransfer to be more like ArrayHandle.
This includes changing methods like LoadDataForInput to PrepareForInput.
It also changed the interface a bit to save a reference to the storage
object. (Maybe it would be better to save a pointer?) These changes also
extend up to the ArrayManagerExecution class, so it can effect device
adapter implementations.
2015-04-30 21:07:36 -06:00
Kenneth Moreland
c5b831b726 Remove LoadDataForInput that takes portals.
This API change effects both ArrayTransfer and ArrayManagerExecution.
This is in preparation for a future change to make the API more
consistent with ArrayHandle.
2015-04-28 11:30:30 -04:00
Kenneth Moreland
a2c280993c Remove UserPortal from ArrayHandle.
The UserPortal in ArrayHandle was used to copy a pointer the user
created into an ArrayHandle to use in VTK-m algorithms. However, this is
only really valid for a basic storage, so the functionality has been
moved there, and you have to construct an ArrayHandle with a storage
instead of an array portal.
2015-04-28 10:49:46 -04:00
Robert Maynard
30c5ee9cf4 Remove SetThrustForCuda as we explicitly use the thrust cuda system.
The only reason SetThrustForCuda was ever needed was that we used
::thrust::device_vector instead of ::thrust::system::cuda::vector.
2015-04-23 09:23:41 -04:00
Robert Maynard
b6f6056db6 Include the Cuda ArrayManagerExecution when it is the default backend. 2015-04-23 09:22:38 -04:00
Robert Maynard
dca1f3b487 Enable Algorithm testing on all the DeviceAdapters. 2015-04-16 10:25:44 -04:00
Robert Maynard
c806348c94 Remove debug code from cuda schedule algorithm. 2015-04-16 08:56:08 -04:00
Robert Maynard
3a10a63e8c make sure the 3d block size makes sense. 2015-04-08 15:57:35 -04:00
Robert Maynard
fdac208acb use cuda scheduling versus thrust scheduling. 2015-04-03 10:18:05 -04:00
Robert Maynard
d340a573d9 don't clear ::thrust::vectors as that is not our intended behavior. 2015-03-23 09:01:42 -04:00
Robert Maynard
86dc8f1d38 Move back to thrust::cuda::vector to properly handle allocating uint8's
Our approach of using the underlying allocator inside thrust was a bad approach,
for some reason it fails to properly allocate uint8's or int8's on the correct
boundaries. I expect that this logic is somewhere else in the code and
instead we should use thrust::system::cuda::vector which does this properly.
2015-03-10 10:10:03 -04:00
Robert Maynard
63b1f03187 Simplify the implementation of loading through textures.
We don't need this super complicated system for texture loading.
2015-03-09 16:37:45 -04:00
Robert Maynard
9b49973621 Use __ldg instead of texture object. 2015-03-05 18:31:44 -05:00
Kenneth Moreland
6141e83be4 Expose allocation in ArrayHandle.
Add an Allocate method in ArrayHandle that basically forwards the
alllocate request to the storage object. This allows some measure of
control of the array from the control side. You can allocate the array
and set values (by getting the control array portal) if you so desire.
2015-02-10 15:58:41 -07:00
Kenneth Moreland
2f781dfe7a Preserve data in ArrayHandle::ReleaseResourcesExecution
Previously when ReleaseResourcesExecution was called, the method blindly
deleted the execution array regardless of whether there was a valid copy
in the control environment. This could potentially lose data. What if
someone wanted to conserve memory on the device by clearing the array of
an output array?

There is also now an internal method that blindly deletes the array.
This is good for internal functions that are doing something to
invalidate the execution data anyway.
2015-02-10 15:49:55 -07:00
Kenneth Moreland
c224c2b98a Make ArrayHandle work better when uninitialized
Fixed a problem where ArrayHandle would cause a crash if you tried to
get the control portal on an uninitialized array (it was supposed to
throw an exception).

Also changed some methods in ArrayHandle so that they work resonably
without error when used with an uninitialized array. Specifically, the
aforementioned behavior was changed to "allocate" an array of size 0
(that is, call Allocate(0) on the storage object) to return an empty
array portal rather than throw an error. Although this use of
ArrayHandle can be considered erroneous, it is convenient the get an
empty array portal when dealing with a potentially unallocated array
rather than create a special condition.
2015-02-09 14:54:04 -07:00
Robert Maynard
3c8ce36666 Properly deallocate cuda memory when we are done with it. 2015-01-28 15:54:45 -05:00
Robert Maynard
46df484ca7 Redesign the way we implement when to choose texture portals.
The logic to select texture portals is easier to understand now.
2015-01-20 13:58:28 -05:00
Robert Maynard
89f8f07806 Properly specify the iterator type for ArrayHandleCounting. 2015-01-20 13:58:28 -05:00
Robert Maynard
e26040282b Properly map TexturePortals to Iterators.
Also at the same time disable the Texture support by default.
2015-01-20 13:58:28 -05:00
Robert Maynard
04bc41cad3 Fix multiple issue in the cuda array handle unit tests.
The namespaces need to be different for each test, or else only the first
implementation of the function will be used for all tests that call that
function.

Also updated the test to verify that we can count starting from a non zero
number.
2015-01-20 09:03:24 -05:00
Robert Maynard
f8bb4214ad Implicit Storage containers don't return anything from GetPortal/Const()
This means you can't assign the results of the function call to a variable.
This was causing nvcc to crash.
2015-01-20 09:03:24 -05:00
Kenneth Moreland
51b5cc63c4 Merge branch 'no-function-interface-zip' into cuda_DeviceAdapterAlgorithm 2015-01-15 22:35:38 -07:00
Kenneth Moreland
694010b8d7 Fix annoying warning about converting signed to unsigned. 2015-01-15 22:32:04 -07:00
Kenneth Moreland
37dac92052 Add index tags to FunctionInterface features.
The functors in the ForEach, StaticTransform, and DynamicTransform
methods sometimes can use the index of the parameter that they are
operating on. This can be a helpful diagnostic in compile and run-time
errors. It is also helpful when linking parameters from one
FunctionInterface with those of another.

This new features are now replacing implementations using the Zip
functionality that was removed earlier. The implementation is actually
simplified a bit.
2015-01-15 22:13:47 -07:00
Robert Maynard
1b5c5a6ce5 Add in initial support for texture binding of input arrays. 2014-12-19 13:47:28 -05:00
Robert Maynard
a509dae909 Use pinned memory for error reporting on the cuda backend.
This reduces the amount of explicit cuda Device to Host memory copies.
2014-12-19 13:47:28 -05:00
Robert Maynard
d9270e408d Adding a cuda device adapter to vtkm.
Porting the dax device adapter over to vtkm. Unlike the dax version, doesn't
use the thrust::device_vector, but instead uses thrust::system calls so that
we can support multiple thrust based backends.

Also this has Texture Memory support for input array handles. Some more work
will need to be done to ArrayHandle so that everything works when using an
ArrayHandle inplace with texture memory bindings.
2014-12-19 13:47:28 -05:00
Robert Maynard
e2eb901be3 move DeviceAdapterAlgorithm to correct folder as it's namespace 2014-12-18 11:09:15 -05:00
Robert Maynard
735a2ff25a Update the documentation to use vtkm::FloatDefault instead of vtkm::Scalar. 2014-12-10 10:38:03 -05:00
Robert Maynard
519509bba1 Update the documentation to note have references to dax. 2014-12-10 10:37:25 -05:00
Kenneth Moreland
a6622aeea9 Documentation fixes.
Fix documentation comments in FunctionInterface.h and some declarations
that are only for documentation in DeviceAdapterAlgorithm.h.
2014-12-09 16:07:27 -07:00
Kenneth Moreland
4ab4dc9cf2 Use test_equal instead of == for comparing numbers in exec worklets.
A couple of tests were failing with the Intel compiler due to
imprecision in comparing floating point values.

Also snuck in some minor documentation fixes in a comment for
FunctionInterface.
2014-11-12 08:53:03 -07:00
Kenneth Moreland
72c62a3556 Merge branch 'dispatch'
Conflicts:
	vtkm/internal/testing/UnitTestFunctionInterface.cxx
2014-11-10 12:27:31 -07:00
Kenneth Moreland
bc3e1ebd21 delete operator and std::allocator are not necessarily the same
The unit test for StorageBasic tested the StealArray feature and then
used the delete[] operator on the stolen array to deallocate it. For
many standard libraries the default implementation for delete[] is
the same as (or at least compatible with) std::allocator, but for
the PGI compiler they were not compatible and this resulted in a
run-time error. This change fixes the problem with the test by using
the same allocator as the StorageBasic test.
2014-10-23 15:39:05 -06:00
Kenneth Moreland
51e3b2bb1d Add checks for signature tags
It's easy to put accidently put something that is not a valid tag in a
ControlSignature or ExecutionSignature. Previously, when you did that
you got a weird error at the end of a very long template instantiation
chain that made it difficult to find the offending worklet.

This adds some type checks when the dispatcher is instantated to check
the signatures. It doesn't point directly to the signature or its
parameter, but it is much closer.
2014-10-23 08:28:50 -06:00
Kenneth Moreland
6b1db2cf04 Check value type in Invoke input arrays.
Instead of just checking that a dispatcher's Invoke input is an
ArrayHandle, also check that the ValueType of the ArrayHandle is
compatible with the types of the worklet operator. This is done by
adding a template argument to the ControlSignature tags that is a type
list tag that gets passed to the type check.
2014-10-22 16:51:42 -06:00
Kenneth Moreland
f4fb9f0ace Merge branch 'new-list-features' into distpach 2014-10-22 15:21:48 -06:00
Kenneth Moreland
421746d041 Add basic support for type checking of Invoke arguments 2014-10-22 15:21:18 -06:00
Kenneth Moreland
b0b0bee4f6 Add VTKM_IS_LIST_TAG checks to dynamic classes 2014-10-22 10:46:54 -06:00
Kenneth Moreland
b368830fc0 Add VTKM_IS_LIST_TAG macro
As part of this is a ListTagCheck class that can be used for other
template magic.
2014-10-22 07:24:04 -06:00
Kenneth Moreland
d4ad846e65 Merge branch 'any-scalar-width' 2014-10-21 16:53:50 -06:00
Kenneth Moreland
7cf25331e2 Fix MSVC compiler warnings
Generally, the MSVC compiler tends to be more picky about implicit type
conversions. It warns if there is any possibility of precision loss.
2014-10-21 16:52:24 -06:00
Kenneth Moreland
30558cf7ed Add Transport class
The Transport class is responsible for moving data from the control
environment to the execution environment. (Actually, it might be more
accurate to say it gets the execution environment associated with a
given control object.) The Transport class is templated with a tag that
controls the mechanism used for the transport.
2014-10-14 10:44:48 -06:00
Kenneth Moreland
10ba2efbf7 Unify test values.
Lots of tests have to move values in and out of arrays and check them
against expected values. It is also often the case that these tests are
run on lots of different types. There is some repeated code for
generating known values for particular indices. This change unifies some
of that. This can probably also encourage making more generic tests.
2014-10-14 10:00:34 -06:00
Kenneth Moreland
a7e6666037 Rename some type lists to be a bit more consistent. 2014-10-10 11:30:10 -06:00
Kenneth Moreland
7f94eafc9c Remove vtkm::Scalar and vtkm::Vector# types
Providing these types tends to "lock in" the precision of the algorithms
used in VTK-m. Since we are using templating anyway, our templates
should be generic enough to handle difference precision in the data.
Usually the appropriate type can be determined by the data provided. In
the case where there is no hint on the precision of data to use (for
example, in the code that provides coordinates for uniform data), there
is a vtkm::FloatDefault.
2014-10-09 08:54:56 -06:00
Kenneth Moreland
4881ed4ddb Change tests that try all base types to use different precision
Before we assumed that we would only use the basic types specified by
the widths of vtkm::Scalar and vtkm::Id. We want to expand this to make
sure the code works on whatever data precision we need.
2014-10-08 16:56:33 -06:00
Kenneth Moreland
0cc9d27e26 Expand the list of types to include multiple widths.
Since we want our code to generally handle data of different precision
(for example either float or double) expand the types in our list types
to include multiple precision.
2014-10-08 15:40:20 -06:00
Kenneth Moreland
b2298172ee Allow longer base lists
Previously we just hand coded base lists up to 4 entries, which was fine
for what we were using it for. However, now that we want to support base
types of different sizes, we are going to need much longer lists.
2014-10-08 12:53:01 -06:00