Commit Graph

112 Commits

Author SHA1 Message Date
Robert Maynard
59618b3a50 Merge branch 'sort_by_key' 2015-05-11 08:04:09 -04:00
Robert Maynard
1bafa94194 Verify the permutation handle can be used as output on all devices 2015-05-08 15:32:06 -04:00
Kenneth Moreland
dd897c0565 Remove template parameter for IncrementBy2
THe IncrementBy2 test type previously allowed any subtype including
floating point numbers. The meaning of this is actually a little unclear
and the feature was causing implicit type conversion warnings that were
hard to template out. The utility of of templating this class is dubious
in the first place, so class is now a fixed type.

I'm a little unsure whether we should keep this test class at all. It's
math operations are ad hoc and it could be difficult to determine if a
problem is caused by an actual problem or just bad math operators.
2015-05-08 10:53:28 -06:00
Kenneth Moreland
209053b79e Fix issues with implicitly converting ints to floats
MSVC complains that converting a 32-bit int to a 32-bit float could
cause loss of precision.
2015-05-08 10:33:59 -06:00
Kenneth Moreland
66aa10aab7 Fix issues with warnings about data loss from type conversion
Fix compile warnings that come up with the flags

  -Wconversion -Wno-sign-conversion

This catches several instances (mostly in the testing framework) where
types are implicitly converted. I expect these changes to fix some of
the warnings we are seeing in MSVC.

I was going to add these flags to the list of extra warning flags, but
unfortunately the Thrust library has several warnings of these types,
and I don't know a good way to turn on the warnings for our code but
turn them off for Thrust.
2015-05-07 16:34:06 -06:00
Kenneth Moreland
229579c337 Compare with tolerance in UnitTestArrayHandlePermutation
The UnitTestArrayHandlePermutation test was failing when compiled with
icc. I believe the issue is that the icc optimization takes some
liberties when computing literal floating point values versus computing
them at run time that makes the two slightly different. I changed all
the applicable comparisons in this test from using the == operator to
using the test_equal function, which adds a tolerance to the comparison.
I expect this to fix the test failure.
2015-05-07 13:22:10 -06:00
Kenneth Moreland
c9f6628c5b Make VTKM_ASSERT_CONT gracefully ignore CUDA devices
There are often instances where one wants to make an assert check in a
method that can run in either the control or execution environment. This
is rather difficult in general for the execution environment, but with
this change you can place the VTKM_ASSERT_CONT macro in such a method,
and it should compile even under CUDA. It works by removing the macro if
compiling for a CUDA device.
2015-05-07 13:13:50 -06:00
Robert Maynard
6564d7af1c Enable SortByKey Test on the Device Adapter. 2015-05-06 09:27:59 -04:00
Robert Maynard
cdac2ca29c Make sure we include as little form windows.h as possible.
Using NOMINMAX and WIN32_LEAN_AND_MEAN allows us to reduce the number
of macro's and functions we bring in from windows.
2015-05-06 09:24:49 -04:00
Robert Maynard
c2f2c166c0 Test all Fancy ArrayHandles and Array Portals on every device adapter.
Fancy array handles being Counting, Implicit, Permutation, Transform, and Zip.
2015-05-05 14:25:48 -04:00
Robert Maynard
1d170b8fdd Add ArrayHandleZip.
Allows both reading and writing to a vtkm::pair of values.
2015-05-05 14:03:40 -04:00
Robert Maynard
1348aa0ec1 Don't try to include vtkm/exec/cuda unless cuda is enabled. 2015-05-05 07:59:44 -04:00
Kenneth Moreland
fdcbb9cd21 Merge branch 'simplify-array-handle' 2015-05-04 17:40:51 -06:00
Kenneth Moreland
d15074153d Remove comments about textures from ArrayManagerExecutionThrustDevice
Per a discussion with Robert Maynard, we no longer have plans to bind
textures or use texture objects. Instead, we are now using the __ldg
command to load from global memory inside the execution portal, which
gives us most of the benefits of textures, but doesn't incur any
bookkeeping.
2015-05-04 14:18:00 -06:00
Robert Maynard
19ee2be2b6 Allow ArrayHandleTransform Control Portal to be accessed when using cuda.
Because of the device/control markups previously we couldn't access the
control portal when compiling with cuda enabled.
2015-05-04 10:58:07 -04:00
Robert Maynard
d101d41c8d ArrayManagerExecution needed to update the RetrieveOutputData signature. 2015-05-04 10:58:07 -04:00
Kenneth Moreland
1c733d4e32 Update the ArrayTransfer classes to new interface merged in
For what it's worth, the implementation of these classes simplified
quite a bit.
2015-04-30 22:01:02 -06:00
Kenneth Moreland
d9f45f09a5 Merge branch 'simplify-array-handle' into fancy-array-portals 2015-04-30 21:37:16 -06:00
Kenneth Moreland
e14fc427f0 Change storage references to storage pointers.
Previously ArrayTransfer and ArrayManagerExecution received a reference
to a Storage class in their constructor and held the reference as an
ivar. In retrospect, this is just asking for trouble. First, it is way
too easy to pass by value when you mean to pass by reference. Second, if
there ever is a bug where the Storage goes out of scope before the
classes holding a reference, it is that much harder to debug.
2015-04-30 21:29:58 -06:00
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
Kenneth Moreland
fe444e1267 Add ArrayHandlePermutation
Also found a problem with ArrayHandle that manifests itself with derived
types when you first do a PrepareForInput and then a PrepareForInPlace.
The ArrayHandle assumes the data is already moved to the device and
skips the in place call to the array transfer. However, this means the
transfer of the derived array handle does not have a chance to set up
for in place.

I think the appropriate solution may be to move the appropriate logic
from ArrayHandle to ArrayTransfer. I will look into that next.
2015-04-23 20:45:24 -04:00
Kenneth Moreland
04a3be81c1 Fix GetNumberOfValues in transformed array handle portal
The number of values in the array handle portal was screwy and the
GetNumberOfValues method was flat out wrong (thanks to Rob Maynard for
pointing that out). This is fixed.

Also fixed a subtle but nasty typing problem in the Storage's
GetPortalConst method.
2015-04-23 11:07:26 -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
Kenneth Moreland
f7daaf36a8 Add ArrayHandleTransform. 2015-04-13 14:43:18 -06:00
Kenneth Moreland
e6d9e120eb Add ArrayHandleImplicit. 2015-04-13 12:15:28 -06: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