From e72cfc3d8aeaf04a9dd46ef4e85b693076a5e5fc Mon Sep 17 00:00:00 2001 From: Li-Ta Lo Date: Fri, 29 May 2020 15:42:37 -0600 Subject: [PATCH 01/19] add particle density --- vtkm/filter/CMakeLists.txt | 2 + vtkm/filter/ParticleDensityNGP.h | 48 +++++++++ vtkm/filter/ParticleDensityNGP.hxx | 102 ++++++++++++++++++ .../testing/UnitTestParticleDensity.cxx | 29 +++++ 4 files changed, 181 insertions(+) create mode 100644 vtkm/filter/ParticleDensityNGP.h create mode 100644 vtkm/filter/ParticleDensityNGP.hxx create mode 100644 vtkm/filter/testing/UnitTestParticleDensity.cxx diff --git a/vtkm/filter/CMakeLists.txt b/vtkm/filter/CMakeLists.txt index 631314d61..6e3cd526e 100644 --- a/vtkm/filter/CMakeLists.txt +++ b/vtkm/filter/CMakeLists.txt @@ -50,6 +50,7 @@ set(headers MeshQuality.h NDEntropy.h NDHistogram.h + ParticleDensityNGP.h Pathline.h PointAverage.h PointElevation.h @@ -116,6 +117,7 @@ set(header_template_sources MeshQuality.hxx NDEntropy.hxx NDHistogram.hxx + ParticleDensityNGP.hxx Pathline.hxx PointAverage.hxx PointElevation.hxx diff --git a/vtkm/filter/ParticleDensityNGP.h b/vtkm/filter/ParticleDensityNGP.h new file mode 100644 index 000000000..9c1d162a3 --- /dev/null +++ b/vtkm/filter/ParticleDensityNGP.h @@ -0,0 +1,48 @@ +//============================================================================ +// Copyright (c) Kitware, Inc. +// All rights reserved. +// See LICENSE.txt for details. +// +// This software is distributed WITHOUT ANY WARRANTY; without even +// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR +// PURPOSE. See the above copyright notice for more information. +//============================================================================ + +#ifndef vtk_m_filter_particle_density_ngp_h +#define vtk_m_filter_particle_density_ngp_h + +#include + +namespace vtkm +{ +namespace filter +{ +/// \brief Estimate the density of particles using the Nearest Grid Point method + +// We only need the CoordinateSystem of the input dataset thus a FilterField +class ParticleDensityNGP : public vtkm::filter::FilterField +{ +public: + // ParticleDensity only support turning 2D/3D particle positions into density + using SupportedTypes = vtkm::TypeListFloatVec; + + // + ParticleDensityNGP(vtkm::Id3& dimension, vtkm::Vec3f& origin, vtkm::Vec3f& spacing); + + template + VTKM_CONT vtkm::cont::DataSet DoExecute(const vtkm::cont::DataSet& input, + const vtkm::cont::ArrayHandle& field, + const vtkm::filter::FieldMetadata& fieldMeta, + vtkm::filter::PolicyBase policy); + +private: + vtkm::Id3 Dimension; + vtkm::Vec3f Origin; + vtkm::Vec3f Spacing; +}; +} +} + +#include + +#endif //vtk_m_filter_particle_density_ngp_h diff --git a/vtkm/filter/ParticleDensityNGP.hxx b/vtkm/filter/ParticleDensityNGP.hxx new file mode 100644 index 000000000..91d8aff17 --- /dev/null +++ b/vtkm/filter/ParticleDensityNGP.hxx @@ -0,0 +1,102 @@ +//============================================================================ +// Copyright (c) Kitware, Inc. +// All rights reserved. +// See LICENSE.txt for details. +// +// This software is distributed WITHOUT ANY WARRANTY; without even +// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR +// PURPOSE. See the above copyright notice for more information. +//============================================================================ + +#ifndef vtk_m_filter_particle_density_ngp_hxx +#define vtk_m_filter_particle_density_ngp_hxx + +#include "ParticleDensityNGP.h" +#include +#include +#include +#include + +namespace vtkm +{ +namespace worklet +{ +class NGPWorklet : public vtkm::worklet::WorkletMapField +{ +public: + using ControlSignature = void(FieldIn coords, ExecObject locator, AtomicArrayInOut density); + using ExecutionSignature = void(_1, _2, _3); + + template + VTKM_EXEC void operator()(const Point& point, + const CellLocatorExecObj& locator, + AtomicArray& density) + { + vtkm::Id cellId; + vtkm::Vec3f parametric; + + // Find the cell containing the point + locator.FindCell(point, cellId, parametric, *this); + + // increment density + density.Add(cellId, 1); + } +}; //NGPWorklet +} //worklet +} //vtkm + + +namespace vtkm +{ +namespace filter +{ +inline VTKM_CONT ParticleDensityNGP::ParticleDensityNGP(vtkm::Id3& dimension, + vtkm::Vec3f& origin, + vtkm::Vec3f& spacing) + : Dimension(dimension) + , Origin(origin) + , Spacing(spacing) +{ +} + +template +inline VTKM_CONT vtkm::cont::DataSet ParticleDensityNGP::DoExecute( + const vtkm::cont::DataSet& input, + const vtkm::cont::ArrayHandle& field, + const vtkm::filter::FieldMetadata& fieldMeta, + vtkm::filter::PolicyBase policy) +{ + // TODO: it really doesn't need to be a UniformGrid, any CellSet with CellLocator will work. + // Make it another input rather an output generated. + + // We want to stores density as PointField which conforms to VTK/VTKm's idea of ImageDataset + // and works with the ImageConnectivity for segmentation purpose. We thus create a surrogate + // uniform dataset that has the cell dimension as the point dimension of the output. We use + // this dataset only for point in cell lookup purpose. This is a somewhat convolved way of + // doing some simple math. + auto surrogate = vtkm::cont::DataSetBuilderUniform::Create( + this->Dimension - vtkm::Id3{ 1, 1, 1 }, this->Origin, this->Spacing); + + // Create a CellSetLocator + vtkm::cont::CellLocatorUniformGrid locator; + locator.SetCellSet(surrogate.GetCellSet()); + locator.SetCoordinates(surrogate.GetCoordinateSystem()); + locator.Update(); + + // We still use an ArrayHandle and pass it to the Worklet as AtomicArrayInOut + // it will be type converted automatically. However the ArrayHandle needs to be + // allocated and initialized. The easily way to do it is to copy from an + // ArrayHandleConstant + vtkm::cont::ArrayHandle density; + vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleConstant(0, surrogate.GetNumberOfCells()), + density); + + this->Invoke(input, locator, density); + + surrogate.AddField(vtkm::cont::make_FieldCell("density", density)); + + return surrogate; +} +} +} +#endif //vtk_m_filter_particle_density_ngp_hxx diff --git a/vtkm/filter/testing/UnitTestParticleDensity.cxx b/vtkm/filter/testing/UnitTestParticleDensity.cxx new file mode 100644 index 000000000..b2d1f19c3 --- /dev/null +++ b/vtkm/filter/testing/UnitTestParticleDensity.cxx @@ -0,0 +1,29 @@ +//============================================================================ +// Copyright (c) Kitware, Inc. +// All rights reserved. +// See LICENSE.txt for details. +// +// This software is distributed WITHOUT ANY WARRANTY; without even +// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR +// PURPOSE. See the above copyright notice for more information. +//============================================================================ + +#include +#include + +#include + +void TestNGP() +{ + std::vector positions = { { 0.5, 0.5 } }; +} + +void TestParticleDensity() +{ + TestNGP(); +} + +int UnitTestParticleDensity(int argc, char* argv[]) +{ + return vtkm::cont::testing::Testing::Run(TestParticleDensity, argc, argv); +} \ No newline at end of file From 9c4129bfb1ea2b38f1aaf74708f9eb54d04a9a15 Mon Sep 17 00:00:00 2001 From: Li-Ta Lo Date: Mon, 19 Oct 2020 15:29:14 -0600 Subject: [PATCH 02/19] Neareast Grid Point particle density estimate --- vtkm/filter/ParticleDensityNGP.h | 8 ++- vtkm/filter/ParticleDensityNGP.hxx | 66 ++++++++++--------- vtkm/filter/testing/CMakeLists.txt | 1 + .../testing/UnitTestParticleDensity.cxx | 34 +++++++++- 4 files changed, 73 insertions(+), 36 deletions(-) diff --git a/vtkm/filter/ParticleDensityNGP.h b/vtkm/filter/ParticleDensityNGP.h index 9c1d162a3..9fb99612a 100644 --- a/vtkm/filter/ParticleDensityNGP.h +++ b/vtkm/filter/ParticleDensityNGP.h @@ -24,10 +24,12 @@ class ParticleDensityNGP : public vtkm::filter::FilterField { public: // ParticleDensity only support turning 2D/3D particle positions into density - using SupportedTypes = vtkm::TypeListFloatVec; + using SupportedTypes = vtkm::TypeListFieldVec3; // - ParticleDensityNGP(vtkm::Id3& dimension, vtkm::Vec3f& origin, vtkm::Vec3f& spacing); + ParticleDensityNGP(const vtkm::Id3& dimension, + const vtkm::Vec3f& origin, + const vtkm::Vec3f& spacing); template VTKM_CONT vtkm::cont::DataSet DoExecute(const vtkm::cont::DataSet& input, @@ -36,7 +38,7 @@ public: vtkm::filter::PolicyBase policy); private: - vtkm::Id3 Dimension; + vtkm::Id3 Dimension; // Cell dimension vtkm::Vec3f Origin; vtkm::Vec3f Spacing; }; diff --git a/vtkm/filter/ParticleDensityNGP.hxx b/vtkm/filter/ParticleDensityNGP.hxx index 9c8c50b33..6a6341249 100644 --- a/vtkm/filter/ParticleDensityNGP.hxx +++ b/vtkm/filter/ParticleDensityNGP.hxx @@ -30,16 +30,24 @@ public: template VTKM_EXEC void operator()(const Point& point, const CellLocatorExecObj& locator, - AtomicArray& density) + AtomicArray& density) const { - vtkm::Id cellId; + vtkm::Id cellId{}; vtkm::Vec3f parametric; // Find the cell containing the point - locator.FindCell(point, cellId, parametric, *this); + if (locator->FindCell(point, cellId, parametric) == vtkm::ErrorCode::Success) + { + // increment density + density.Add(cellId, 1); + } + else + { - // increment density - density.Add(cellId, 1); + // FIXME: what does mean when it is not found? + // We simply ignore that particular particle. + std::cout << "WTF: " << point << std::endl; + } } }; //NGPWorklet } //worklet @@ -50,9 +58,9 @@ namespace vtkm { namespace filter { -inline VTKM_CONT ParticleDensityNGP::ParticleDensityNGP(vtkm::Id3& dimension, - vtkm::Vec3f& origin, - vtkm::Vec3f& spacing) +inline VTKM_CONT ParticleDensityNGP::ParticleDensityNGP(const vtkm::Id3& dimension, + const vtkm::Vec3f& origin, + const vtkm::Vec3f& spacing) : Dimension(dimension) , Origin(origin) , Spacing(spacing) @@ -61,41 +69,39 @@ inline VTKM_CONT ParticleDensityNGP::ParticleDensityNGP(vtkm::Id3& dimension, template inline VTKM_CONT vtkm::cont::DataSet ParticleDensityNGP::DoExecute( - const vtkm::cont::DataSet& input, + const vtkm::cont::DataSet&, const vtkm::cont::ArrayHandle& field, - const vtkm::filter::FieldMetadata& fieldMeta, - vtkm::filter::PolicyBase policy) + const vtkm::filter::FieldMetadata&, + vtkm::filter::PolicyBase) { // TODO: it really doesn't need to be a UniformGrid, any CellSet with CellLocator will work. - // Make it another input rather an output generated. + // Make it another input rather an output generated. - // We want to stores density as PointField which conforms to VTK/VTKm's idea of ImageDataset - // and works with the ImageConnectivity for segmentation purpose. We thus create a surrogate - // uniform dataset that has the cell dimension as the point dimension of the output. We use - // this dataset only for point in cell lookup purpose. This is a somewhat convolved way of - // doing some simple math. - auto surrogate = vtkm::cont::DataSetBuilderUniform::Create( - this->Dimension - vtkm::Id3{ 1, 1, 1 }, this->Origin, this->Spacing); + // We stores density as CellField which conforms to physicists' idea of particle density + // better. However, VTK/VTKm's idea of "Image" Dataset and the ImageConnectivity filter + // expect a PointField. For better separation of concerns, we create a uniform dataset + // that has the cell dimension as expected and later convert the dataset to its dual. + auto uniform = vtkm::cont::DataSetBuilderUniform::Create( + this->Dimension + vtkm::Id3{ 1, 1, 1 }, this->Origin, this->Spacing); - // Create a CellSetLocator + // Create a CellLocator vtkm::cont::CellLocatorUniformGrid locator; - locator.SetCellSet(surrogate.GetCellSet()); - locator.SetCoordinates(surrogate.GetCoordinateSystem()); + locator.SetCellSet(uniform.GetCellSet()); + locator.SetCoordinates(uniform.GetCoordinateSystem()); locator.Update(); - // We still use an ArrayHandle and pass it to the Worklet as AtomicArrayInOut - // it will be type converted automatically. However the ArrayHandle needs to be - // allocated and initialized. The easily way to do it is to copy from an - // ArrayHandleConstant + // We create an ArrayHandle and pass it to the Worklet as AtomicArrayInOut. + // However the ArrayHandle needs to be allocated and initialized first. The + // easily way to do it is to copy from an ArrayHandleConstant vtkm::cont::ArrayHandle density; - vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleConstant(0, surrogate.GetNumberOfCells()), + vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleConstant(0, uniform.GetNumberOfCells()), density); - this->Invoke(input, locator, density); + this->Invoke(vtkm::worklet::NGPWorklet{}, field, locator, density); - surrogate.AddField(vtkm::cont::make_FieldCell("density", density)); + uniform.AddField(vtkm::cont::make_FieldCell("density", density)); - return surrogate; + return uniform; } } diff --git a/vtkm/filter/testing/CMakeLists.txt b/vtkm/filter/testing/CMakeLists.txt index 7a60a2b08..3b4e94072 100644 --- a/vtkm/filter/testing/CMakeLists.txt +++ b/vtkm/filter/testing/CMakeLists.txt @@ -47,6 +47,7 @@ set(unit_tests UnitTestNDEntropyFilter.cxx UnitTestNDHistogramFilter.cxx UnitTestParticleAdvectionFilter.cxx + UnitTestParticleDensity.cxx UnitTestPartitionedDataSetFilters.cxx UnitTestPartitionedDataSetHistogramFilter.cxx UnitTestPointAverageFilter.cxx diff --git a/vtkm/filter/testing/UnitTestParticleDensity.cxx b/vtkm/filter/testing/UnitTestParticleDensity.cxx index b2d1f19c3..d7e0bca87 100644 --- a/vtkm/filter/testing/UnitTestParticleDensity.cxx +++ b/vtkm/filter/testing/UnitTestParticleDensity.cxx @@ -8,14 +8,42 @@ // PURPOSE. See the above copyright notice for more information. //============================================================================ +#include +#include #include #include - -#include +#include void TestNGP() { - std::vector positions = { { 0.5, 0.5 } }; + const vtkm::Id N = 1000000; + auto composite = vtkm::cont::make_ArrayHandleCompositeVector( + vtkm::cont::ArrayHandleRandomUniformReal(N, 0xceed), + vtkm::cont::ArrayHandleRandomUniformReal(N, 0xdeed), + vtkm::cont::ArrayHandleRandomUniformReal(N, 0xabba)); + vtkm::cont::ArrayHandle positions; + vtkm::cont::ArrayCopy(composite, positions); + + vtkm::cont::ArrayHandle connectivity; + vtkm::cont::ArrayCopy(vtkm::cont::make_ArrayHandleIndex(N), connectivity); + + auto dataSet = vtkm::cont::DataSetBuilderExplicit::Create( + positions, vtkm::CellShapeTagVertex{}, 1, connectivity); + + auto cellDims = vtkm::Id3{ 3, 3, 3 }; + vtkm::filter::ParticleDensityNGP filter{ cellDims, + { 0.f, 0.f, 0.f }, + vtkm::Vec3f{ 1.f / 3.f, 1.f / 3.f, 1.f / 3.f } }; + filter.SetUseCoordinateSystemAsField(true); + auto density = filter.Execute(dataSet); + + vtkm::cont::ArrayHandle field; + density.GetCellField("density").GetData().AsArrayHandle(field); + auto field_f = vtkm::cont::make_ArrayHandleCast(field); + + auto result = vtkm::worklet::DescriptiveStatistics::Run(field_f); + VTKM_TEST_ASSERT(test_equal(result.Sum(), N)); + VTKM_TEST_ASSERT(test_equal(result.Mean(), N / density.GetNumberOfCells())); } void TestParticleDensity() From e6e0506e1f314c2577b665c00bf349794966b42f Mon Sep 17 00:00:00 2001 From: Li-Ta Lo Date: Mon, 19 Oct 2020 15:46:08 -0600 Subject: [PATCH 03/19] add a newline --- vtkm/filter/testing/UnitTestParticleDensity.cxx | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/vtkm/filter/testing/UnitTestParticleDensity.cxx b/vtkm/filter/testing/UnitTestParticleDensity.cxx index d7e0bca87..f5dab0594 100644 --- a/vtkm/filter/testing/UnitTestParticleDensity.cxx +++ b/vtkm/filter/testing/UnitTestParticleDensity.cxx @@ -54,4 +54,4 @@ void TestParticleDensity() int UnitTestParticleDensity(int argc, char* argv[]) { return vtkm::cont::testing::Testing::Run(TestParticleDensity, argc, argv); -} \ No newline at end of file +} From f1b27c2b8bc30dfef6d2872d191fe45bb3e15819 Mon Sep 17 00:00:00 2001 From: Li-Ta Lo Date: Mon, 19 Oct 2020 17:49:55 -0600 Subject: [PATCH 04/19] removed bad stuff --- vtkm/filter/ParticleDensityNGP.hxx | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/vtkm/filter/ParticleDensityNGP.hxx b/vtkm/filter/ParticleDensityNGP.hxx index 6a6341249..0fea15ea5 100644 --- a/vtkm/filter/ParticleDensityNGP.hxx +++ b/vtkm/filter/ParticleDensityNGP.hxx @@ -41,13 +41,9 @@ public: // increment density density.Add(cellId, 1); } - else - { - // FIXME: what does mean when it is not found? - // We simply ignore that particular particle. - std::cout << "WTF: " << point << std::endl; - } + // FIXME: what does mean when it is not found? + // We simply ignore that particular particle. } }; //NGPWorklet } //worklet From bcc840eaf6c3e122266f5259ab865d4155dd7891 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Wed, 21 Oct 2020 10:15:52 -0400 Subject: [PATCH 05/19] Add support for the Ampere line of NVIDIA GPU's --- CMake/VTKmDeviceAdapters.cmake | 16 +++++++++++++--- 1 file changed, 13 insertions(+), 3 deletions(-) diff --git a/CMake/VTKmDeviceAdapters.cmake b/CMake/VTKmDeviceAdapters.cmake index eb98d5c03..050ff3190 100644 --- a/CMake/VTKmDeviceAdapters.cmake +++ b/CMake/VTKmDeviceAdapters.cmake @@ -167,7 +167,10 @@ if(VTKm_ENABLE_CUDA) # 6 - volta # - Uses: --generate-code=arch=compute_70,code=sm_70 # 7 - turing - # - Uses: --generate-code=arch=compute_75code=sm_75 + # - Uses: --generate-code=arch=compute_75,code=sm_75 + # 8 - ampere + # - Uses: --generate-code=arch=compute_80,code=sm_80 + # - Uses: --generate-code=arch=compute_86,code=sm_86 # 8 - all # - Uses: --generate-code=arch=compute_30,code=sm_30 # - Uses: --generate-code=arch=compute_35,code=sm_35 @@ -175,12 +178,14 @@ if(VTKm_ENABLE_CUDA) # - Uses: --generate-code=arch=compute_60,code=sm_60 # - Uses: --generate-code=arch=compute_70,code=sm_70 # - Uses: --generate-code=arch=compute_75,code=sm_75 + # - Uses: --generate-code=arch=compute_80,code=sm_80 + # - Uses: --generate-code=arch=compute_86,code=sm_86 # 8 - none # #specify the property set(VTKm_CUDA_Architecture "native" CACHE STRING "Which GPU Architecture(s) to compile for") - set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta turing all none) + set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta turing ampere all none) #detect what the property is set too if(VTKm_CUDA_Architecture STREQUAL "native") @@ -234,13 +239,18 @@ if(VTKm_ENABLE_CUDA) set(arch_flags --generate-code=arch=compute_70,code=sm_70) elseif(VTKm_CUDA_Architecture STREQUAL "turing") set(arch_flags --generate-code=arch=compute_75,code=sm_75) + elseif(VTKm_CUDA_Architecture STREQUAL "ampere") + set(arch_flags --generate-code=arch=compute_80,code=sm_80) + set(arch_flags --generate-code=arch=compute_86,code=sm_86) elseif(VTKm_CUDA_Architecture STREQUAL "all") set(arch_flags --generate-code=arch=compute_30,code=sm_30 --generate-code=arch=compute_35,code=sm_35 --generate-code=arch=compute_50,code=sm_50 --generate-code=arch=compute_60,code=sm_60 --generate-code=arch=compute_70,code=sm_70 - --generate-code=arch=compute_75,code=sm_75) + --generate-code=arch=compute_75,code=sm_75 + --generate-code=arch=compute_80,code=sm_80 + --generate-code=arch=compute_86,code=sm_86) endif() string(REPLACE ";" " " arch_flags "${arch_flags}") From f66dc780c5857d07eabaaa949fe43a824f5cbebf Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Wed, 26 Aug 2020 07:58:59 -0600 Subject: [PATCH 06/19] Update ArrayHandleImplicit to new array style with Buffer Also updated the dependent `ArrayHandle`s that use `ArrayHandleImplicit`'s storage. --- vtkm/cont/ArrayHandleConstant.h | 18 +- vtkm/cont/ArrayHandleCounting.h | 37 +--- vtkm/cont/ArrayHandleImplicit.cxx | 75 +++++++ vtkm/cont/ArrayHandleImplicit.h | 192 ++++++++++++++++-- vtkm/cont/ArrayHandleIndex.h | 35 ++-- .../cont/ArrayHandleUniformPointCoordinates.h | 20 +- vtkm/cont/BitField.cxx | 20 +- vtkm/cont/BitField.h | 6 +- vtkm/cont/CMakeLists.txt | 2 +- vtkm/cont/StorageImplicit.h | 177 ---------------- vtkm/cont/testing/CMakeLists.txt | 1 - .../testing/UnitTestArrayHandleCounting.cxx | 55 +---- vtkm/cont/testing/UnitTestStorageImplicit.cxx | 134 ------------ .../ArrayPortalUniformPointCoordinates.h | 19 -- 14 files changed, 293 insertions(+), 498 deletions(-) create mode 100644 vtkm/cont/ArrayHandleImplicit.cxx delete mode 100644 vtkm/cont/StorageImplicit.h delete mode 100644 vtkm/cont/testing/UnitTestStorageImplicit.cxx diff --git a/vtkm/cont/ArrayHandleConstant.h b/vtkm/cont/ArrayHandleConstant.h index 51dfd097f..f5cc36460 100644 --- a/vtkm/cont/ArrayHandleConstant.h +++ b/vtkm/cont/ArrayHandleConstant.h @@ -47,20 +47,13 @@ using StorageTagConstantSuperclass = template struct Storage : Storage> { - using Superclass = Storage>; - using Superclass::Superclass; -}; - -template -struct ArrayTransfer - : ArrayTransfer, Device> -{ - using Superclass = ArrayTransfer, Device>; - using Superclass::Superclass; }; } // namespace internal +template +VTKM_ARRAY_HANDLE_NEW_STYLE(T, vtkm::cont::StorageTagConstant); + /// \brief An array handle with a constant value. /// /// ArrayHandleConstant is an implicit array handle with a constant value. A @@ -79,9 +72,8 @@ public: VTKM_CONT ArrayHandleConstant(T value, vtkm::Id numberOfValues = 0) - : Superclass(typename internal::Storage::PortalConstType( - internal::ConstantFunctor(value), - numberOfValues)) + : Superclass(internal::FunctorToArrayHandleImplicitBuffers(internal::ConstantFunctor(value), + numberOfValues)) { } }; diff --git a/vtkm/cont/ArrayHandleCounting.h b/vtkm/cont/ArrayHandleCounting.h index 19533ca0e..01658e386 100644 --- a/vtkm/cont/ArrayHandleCounting.h +++ b/vtkm/cont/ArrayHandleCounting.h @@ -10,8 +10,7 @@ #ifndef vtk_m_cont_ArrayHandleCounting_h #define vtk_m_cont_ArrayHandleCounting_h -#include -#include +#include #include @@ -52,24 +51,6 @@ public: { } - template - VTKM_EXEC_CONT ArrayPortalCounting(const ArrayPortalCounting& src) - : Start(src.Start) - , Step(src.Step) - , NumberOfValues(src.NumberOfValues) - { - } - - template - VTKM_EXEC_CONT ArrayPortalCounting& operator=( - const ArrayPortalCounting& src) - { - this->Start = src.Start; - this->Step = src.Step; - this->NumberOfValues = src.NumberOfValues; - return *this; - } - VTKM_EXEC_CONT ValueType GetStart() const { return this->Start; } @@ -98,20 +79,13 @@ using StorageTagCountingSuperclass = template struct Storage : Storage> { - using Superclass = Storage>; - using Superclass::Superclass; -}; - -template -struct ArrayTransfer - : ArrayTransfer, Device> -{ - using Superclass = ArrayTransfer, Device>; - using Superclass::Superclass; }; } // namespace internal +template +VTKM_ARRAY_HANDLE_NEW_STYLE(T, vtkm::cont::StorageTagCounting); + /// ArrayHandleCounting is a specialization of ArrayHandle. By default it /// contains a increment value, that is increment for each step between zero /// and the passed in length @@ -126,7 +100,8 @@ public: VTKM_CONT ArrayHandleCounting(CountingValueType start, CountingValueType step, vtkm::Id length) - : Superclass(internal::ArrayPortalCounting(start, step, length)) + : Superclass(internal::PortalToArrayHandleImplicitBuffers( + internal::ArrayPortalCounting(start, step, length))) { } }; diff --git a/vtkm/cont/ArrayHandleImplicit.cxx b/vtkm/cont/ArrayHandleImplicit.cxx new file mode 100644 index 000000000..92cc4cc6c --- /dev/null +++ b/vtkm/cont/ArrayHandleImplicit.cxx @@ -0,0 +1,75 @@ +//============================================================================ +// Copyright (c) Kitware, Inc. +// All rights reserved. +// See LICENSE.txt for details. +// +// This software is distributed WITHOUT ANY WARRANTY; without even +// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR +// PURPOSE. See the above copyright notice for more information. +//============================================================================ + +#include + +namespace vtkm +{ +namespace cont +{ +namespace internal +{ + +BufferMetaDataImplicit::BufferMetaDataImplicit(const BufferMetaDataImplicit& src) + : Deleter(src.Deleter) + , Copier(src.Copier) +{ + if (src.Portal) + { + VTKM_ASSERT(src.Deleter); + VTKM_ASSERT(src.Copier); + + this->Portal = src.Copier(src.Portal); + } + else + { + this->Portal = nullptr; + } +} + +BufferMetaDataImplicit::~BufferMetaDataImplicit() +{ + if (this->Portal) + { + VTKM_ASSERT(this->Deleter); + this->Deleter(this->Portal); + this->Portal = nullptr; + } +} + +std::unique_ptr BufferMetaDataImplicit::DeepCopy() const +{ + return std::unique_ptr(new BufferMetaDataImplicit(*this)); +} + +namespace detail +{ + +vtkm::cont::internal::BufferMetaDataImplicit* GetImplicitMetaData( + const vtkm::cont::internal::Buffer& buffer) +{ + vtkm::cont::internal::BufferMetaDataImplicit* metadata = + dynamic_cast(buffer.GetMetaData()); + VTKM_ASSERT(metadata && "Buffer for implicit array does not have correct metadata."); + return metadata; +} + +} // namespace detail + +} +} +} // namespace vtkm::cont::internal + +namespace vtkm +{ +namespace cont +{ +} +} // namespace vtkm::cont::detail diff --git a/vtkm/cont/ArrayHandleImplicit.h b/vtkm/cont/ArrayHandleImplicit.h index 058fc51bb..b9a304ff9 100644 --- a/vtkm/cont/ArrayHandleImplicit.h +++ b/vtkm/cont/ArrayHandleImplicit.h @@ -11,30 +11,15 @@ #define vtk_m_cont_ArrayHandleImplicit_h #include -#include + +#include namespace vtkm { -namespace cont -{ -namespace detail +namespace internal { -template -class VTKM_ALWAYS_EXPORT ArrayPortalImplicit; - -/// A convenience class that provides a typedef to the appropriate tag for -/// a implicit array container. -template -struct ArrayHandleImplicitTraits -{ - using ValueType = decltype(FunctorType{}(vtkm::Id{})); - using StorageTag = vtkm::cont::StorageTagImplicit>; - using Superclass = vtkm::cont::ArrayHandle; - using StorageType = vtkm::cont::internal::Storage; -}; - /// \brief An array portal that returns the result of a functor /// /// This array portal is similar to an implicit array i.e an array that is @@ -48,8 +33,8 @@ template class VTKM_ALWAYS_EXPORT ArrayPortalImplicit { public: - using ValueType = typename ArrayHandleImplicitTraits::ValueType; using FunctorType = FunctorType_; + using ValueType = decltype(FunctorType{}(vtkm::Id{})); VTKM_SUPPRESS_EXEC_WARNINGS VTKM_EXEC_CONT @@ -82,8 +67,170 @@ private: vtkm::Id NumberOfValues; }; +} // namespace internal + +namespace cont +{ + +/// \brief An implementation for read-only implicit arrays. +/// +/// It is sometimes the case that you want VTK-m to operate on an array of +/// implicit values. That is, rather than store the data in an actual array, it +/// is gerenated on the fly by a function. This is handled in VTK-m by creating +/// an ArrayHandle in VTK-m with a StorageTagImplicit type of \c Storage. This +/// tag itself is templated to specify an ArrayPortal that generates the +/// desired values. An ArrayHandle created with this tag will raise an error on +/// any operation that tries to modify it. +/// +template +struct VTKM_ALWAYS_EXPORT StorageTagImplicit +{ + using PortalType = ArrayPortalType; +}; + +namespace internal +{ + +struct VTKM_CONT_EXPORT BufferMetaDataImplicit : vtkm::cont::internal::BufferMetaData +{ + void* Portal; + + using DeleterType = void(void*); + DeleterType* Deleter; + + using CopierType = void*(void*); + CopierType* Copier; + + template + BufferMetaDataImplicit(const PortalType& portal) + : Portal(new PortalType(portal)) + , Deleter([](void* p) { delete reinterpret_cast(p); }) + , Copier([](void* p) -> void* { return new PortalType(*reinterpret_cast(p)); }) + { + } + + VTKM_CONT BufferMetaDataImplicit(const BufferMetaDataImplicit& src); + + BufferMetaDataImplicit& operator=(const BufferMetaDataImplicit&) = delete; + + VTKM_CONT ~BufferMetaDataImplicit() override; + + VTKM_CONT std::unique_ptr DeepCopy() const override; +}; + +namespace detail +{ + +VTKM_CONT_EXPORT vtkm::cont::internal::BufferMetaDataImplicit* GetImplicitMetaData( + const vtkm::cont::internal::Buffer& buffer); + } // namespace detail +template +struct VTKM_ALWAYS_EXPORT + Storage> +{ + VTKM_IS_TRIVIALLY_COPYABLE(ArrayPortalType); + + using ReadPortalType = ArrayPortalType; + + // Note that this portal is almost certainly read-only, so you will probably get + // an error if you try to write to it. + using WritePortalType = ArrayPortalType; + + // Implicit array has one buffer that should be empty (NumberOfBytes = 0), but holds + // the metadata for the array. + VTKM_CONT static vtkm::IdComponent GetNumberOfBuffers() { return 1; } + + VTKM_CONT static vtkm::Id GetNumberOfValues(const vtkm::cont::internal::Buffer* buffers) + { + vtkm::cont::internal::BufferMetaDataImplicit* metadata = + detail::GetImplicitMetaData(buffers[0]); + VTKM_ASSERT(metadata->Portal); + return reinterpret_cast(metadata->Portal)->GetNumberOfValues(); + } + + VTKM_CONT static void ResizeBuffers(vtkm::Id numValues, + vtkm::cont::internal::Buffer* buffers, + vtkm::CopyFlag, + vtkm::cont::Token&) + { + if (numValues == GetNumberOfValues(buffers)) + { + // In general, we don't allow resizing of the array, but if it was "allocated" to the + // correct size, we will allow that. + } + else + { + throw vtkm::cont::ErrorBadAllocation("Cannot allocate/resize implicit arrays."); + } + } + + VTKM_CONT static ReadPortalType CreateReadPortal(const vtkm::cont::internal::Buffer* buffers, + vtkm::cont::DeviceAdapterId, + vtkm::cont::Token&) + { + vtkm::cont::internal::BufferMetaDataImplicit* metadata = + detail::GetImplicitMetaData(buffers[0]); + VTKM_ASSERT(metadata->Portal); + return *reinterpret_cast(metadata->Portal); + } + + VTKM_CONT static WritePortalType CreateWritePortal(const vtkm::cont::internal::Buffer*, + vtkm::cont::DeviceAdapterId, + vtkm::cont::Token&) + { + throw vtkm::cont::ErrorBadAllocation("Cannot write to implicit arrays."); + } +}; + +/// Given an array portal, returns the buffers for the `ArrayHandle` with a storage that +/// is (or is compatible with) a storage tag of `StorageTagImplicit`. +template +VTKM_CONT inline std::vector PortalToArrayHandleImplicitBuffers( + const PortalType& portal) +{ + std::vector buffers(1); + buffers[0].SetMetaData(std::unique_ptr( + new vtkm::cont::internal::BufferMetaDataImplicit(portal))); + return buffers; +} + +/// Given a functor and the number of values, returns the buffers for the `ArrayHandleImplicit` +/// for the given functor. +template +VTKM_CONT inline std::vector FunctorToArrayHandleImplicitBuffers( + const FunctorType& functor, + vtkm::Id numValues) +{ + return PortalToArrayHandleImplicitBuffers( + vtkm::internal::ArrayPortalImplicit(functor, numValues)); +} + +} // namespace internal + +namespace detail +{ + +/// A convenience class that provides a typedef to the appropriate tag for +/// a implicit array container. +template +struct ArrayHandleImplicitTraits +{ + using ValueType = decltype(FunctorType{}(vtkm::Id{})); + using PortalType = vtkm::internal::ArrayPortalImplicit; + using StorageTag = vtkm::cont::StorageTagImplicit; + using Superclass = vtkm::cont::ArrayHandle; + using StorageType = vtkm::cont::internal::Storage; +}; + +} // namespace detail + +// This can go away once ArrayHandle is replaced with ArrayHandleNewStyle +template +VTKM_ARRAY_HANDLE_NEW_STYLE(typename PortalType::ValueType, + vtkm::cont::StorageTagImplicit); + /// \brief An \c ArrayHandle that computes values on the fly. /// /// \c ArrayHandleImplicit is a specialization of ArrayHandle. @@ -97,6 +244,7 @@ class VTKM_ALWAYS_EXPORT ArrayHandleImplicit { private: using ArrayTraits = typename detail::ArrayHandleImplicitTraits; + using PortalType = typename ArrayTraits::PortalType; public: VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleImplicit, @@ -105,7 +253,7 @@ public: VTKM_CONT ArrayHandleImplicit(FunctorType functor, vtkm::Id length) - : Superclass(typename ArrayTraits::StorageType::PortalType(functor, length)) + : Superclass(internal::PortalToArrayHandleImplicitBuffers(PortalType(functor, length))) { } }; @@ -144,7 +292,7 @@ struct SerializableTypeString> template struct SerializableTypeString::ValueType, - vtkm::cont::StorageTagImplicit>>> + vtkm::cont::StorageTagImplicit>>> : SerializableTypeString> { }; @@ -183,7 +331,7 @@ public: template struct Serialization::ValueType, - vtkm::cont::StorageTagImplicit>>> + vtkm::cont::StorageTagImplicit>>> : Serialization> { }; diff --git a/vtkm/cont/ArrayHandleIndex.h b/vtkm/cont/ArrayHandleIndex.h index f99cb94a1..4a152b77e 100644 --- a/vtkm/cont/ArrayHandleIndex.h +++ b/vtkm/cont/ArrayHandleIndex.h @@ -14,6 +14,17 @@ namespace vtkm { + +namespace internal +{ + +struct VTKM_ALWAYS_EXPORT IndexFunctor +{ + VTKM_EXEC_CONT vtkm::Id operator()(vtkm::Id index) const { return index; } +}; + +} // namespace internal + namespace cont { @@ -24,32 +35,19 @@ struct VTKM_ALWAYS_EXPORT StorageTagIndex namespace internal { -struct VTKM_ALWAYS_EXPORT IndexFunctor -{ - VTKM_EXEC_CONT - vtkm::Id operator()(vtkm::Id index) const { return index; } -}; - using StorageTagIndexSuperclass = - typename vtkm::cont::ArrayHandleImplicit::StorageTag; + typename vtkm::cont::ArrayHandleImplicit::StorageTag; template <> struct Storage : Storage { - using Superclass = Storage; - using Superclass::Superclass; -}; - -template -struct ArrayTransfer - : ArrayTransfer -{ - using Superclass = ArrayTransfer; - using Superclass::Superclass; }; } // namespace internal +template <> +VTKM_ARRAY_HANDLE_NEW_STYLE(vtkm::Id, vtkm::cont::StorageTagIndex); + /// \brief An implicit array handle containing the its own indices. /// /// \c ArrayHandleIndex is an implicit array handle containing the values @@ -65,8 +63,7 @@ public: VTKM_CONT ArrayHandleIndex(vtkm::Id length) : Superclass( - typename internal::Storage::PortalType(internal::IndexFunctor{}, - length)) + internal::FunctorToArrayHandleImplicitBuffers(vtkm::internal::IndexFunctor{}, length)) { } }; diff --git a/vtkm/cont/ArrayHandleUniformPointCoordinates.h b/vtkm/cont/ArrayHandleUniformPointCoordinates.h index cf3fa2e15..118062a20 100644 --- a/vtkm/cont/ArrayHandleUniformPointCoordinates.h +++ b/vtkm/cont/ArrayHandleUniformPointCoordinates.h @@ -10,8 +10,7 @@ #ifndef vtk_m_cont_ArrayHandleUniformPointCoordinates_h #define vtk_m_cont_ArrayHandleUniformPointCoordinates_h -#include -#include +#include #include namespace vtkm @@ -33,22 +32,13 @@ template <> struct Storage : Storage { - using Superclass = Storage; - - using Superclass::Superclass; -}; - -template -struct ArrayTransfer - : ArrayTransfer -{ - using Superclass = ArrayTransfer; - - using Superclass::Superclass; }; } // namespace internal +template <> +VTKM_ARRAY_HANDLE_NEW_STYLE(vtkm::Vec3f, vtkm::cont::StorageTagUniformPoints); + /// ArrayHandleUniformPointCoordinates is a specialization of ArrayHandle. It /// contains the information necessary to compute the point coordinates in a /// uniform orthogonal grid (extent, origin, and spacing) and implicitly @@ -70,7 +60,7 @@ public: ArrayHandleUniformPointCoordinates(vtkm::Id3 dimensions, ValueType origin = ValueType(0.0f, 0.0f, 0.0f), ValueType spacing = ValueType(1.0f, 1.0f, 1.0f)) - : Superclass(StorageType( + : Superclass(internal::PortalToArrayHandleImplicitBuffers( vtkm::internal::ArrayPortalUniformPointCoordinates(dimensions, origin, spacing))) { } diff --git a/vtkm/cont/BitField.cxx b/vtkm/cont/BitField.cxx index 399b544c0..37f073797 100644 --- a/vtkm/cont/BitField.cxx +++ b/vtkm/cont/BitField.cxx @@ -41,7 +41,7 @@ namespace cont namespace detail { -vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData( +vtkm::cont::internal::BufferMetaDataBitField* GetBitFieldMetaData( const vtkm::cont::internal::Buffer& buffer) { vtkm::cont::internal::BufferMetaData* generalMetaData = buffer.GetMetaData(); @@ -49,22 +49,22 @@ vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData( { VTKM_LOG_F(vtkm::cont::LogLevel::Warn, "BitField has buffer with no metadata. Setting."); const_cast(buffer).SetMetaData( - vtkm::cont::internal::BitFieldMetaData{}); + vtkm::cont::internal::BufferMetaDataBitField{}); generalMetaData = buffer.GetMetaData(); VTKM_ASSERT(generalMetaData != nullptr); } - vtkm::cont::internal::BitFieldMetaData* metadata = - dynamic_cast(generalMetaData); + vtkm::cont::internal::BufferMetaDataBitField* metadata = + dynamic_cast(generalMetaData); if (metadata == nullptr) { VTKM_LOG_F(vtkm::cont::LogLevel::Error, "BitField has a buffer with metadata of the wrong type. " "Replacing, but this will likely cause problems."); const_cast(buffer).SetMetaData( - vtkm::cont::internal::BitFieldMetaData{}); + vtkm::cont::internal::BufferMetaDataBitField{}); generalMetaData = buffer.GetMetaData(); - metadata = dynamic_cast(generalMetaData); + metadata = dynamic_cast(generalMetaData); VTKM_ASSERT(metadata != nullptr); } @@ -79,18 +79,18 @@ vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData( namespace internal { -BitFieldMetaData::~BitFieldMetaData() {} +BufferMetaDataBitField::~BufferMetaDataBitField() {} -std::unique_ptr BitFieldMetaData::DeepCopy() const +std::unique_ptr BufferMetaDataBitField::DeepCopy() const { - return std::unique_ptr(new BitFieldMetaData(*this)); + return std::unique_ptr(new BufferMetaDataBitField(*this)); } } // namespace internal BitField::BitField() { - this->Buffer.SetMetaData(internal::BitFieldMetaData{}); + this->Buffer.SetMetaData(internal::BufferMetaDataBitField{}); } vtkm::Id BitField::GetNumberOfBits() const diff --git a/vtkm/cont/BitField.h b/vtkm/cont/BitField.h index e7786515c..c08683eec 100644 --- a/vtkm/cont/BitField.h +++ b/vtkm/cont/BitField.h @@ -36,11 +36,11 @@ namespace internal struct StorageTagBitField; -struct VTKM_CONT_EXPORT BitFieldMetaData : vtkm::cont::internal::BufferMetaData +struct VTKM_CONT_EXPORT BufferMetaDataBitField : vtkm::cont::internal::BufferMetaData { vtkm::Id NumberOfBits = 0; - VTKM_CONT ~BitFieldMetaData() override; + VTKM_CONT ~BufferMetaDataBitField() override; VTKM_CONT std::unique_ptr DeepCopy() const override; }; @@ -99,7 +99,7 @@ struct BitCoordinate vtkm::Int32 BitOffset; // [0, bitsInWord) }; -VTKM_CONT_EXPORT vtkm::cont::internal::BitFieldMetaData* GetBitFieldMetaData( +VTKM_CONT_EXPORT vtkm::cont::internal::BufferMetaDataBitField* GetBitFieldMetaData( const vtkm::cont::internal::Buffer& buffer); /// Portal for performing bit or word operations on a BitField. diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index 2d6002e1d..dc3c3b8eb 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -109,7 +109,6 @@ set(headers Serialization.h Storage.h StorageExtrude.h - StorageImplicit.h StorageList.h StorageListTag.h Timer.h @@ -137,6 +136,7 @@ set(template_sources set(sources ArrayHandle.cxx ArrayHandleBasic.cxx + ArrayHandleImplicit.cxx ArrayHandleSOA.cxx BitField.cxx ColorTablePresets.cxx diff --git a/vtkm/cont/StorageImplicit.h b/vtkm/cont/StorageImplicit.h deleted file mode 100644 index 29f53f1a8..000000000 --- a/vtkm/cont/StorageImplicit.h +++ /dev/null @@ -1,177 +0,0 @@ -//============================================================================ -// Copyright (c) Kitware, Inc. -// All rights reserved. -// See LICENSE.txt for details. -// -// This software is distributed WITHOUT ANY WARRANTY; without even -// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR -// PURPOSE. See the above copyright notice for more information. -//============================================================================ -#ifndef vtk_m_cont_StorageImplicit -#define vtk_m_cont_StorageImplicit - -#include - -#include -#include -#include -#include - -#include - -namespace vtkm -{ -namespace cont -{ - -/// \brief An implementation for read-only implicit arrays. -/// -/// It is sometimes the case that you want VTK-m to operate on an array of -/// implicit values. That is, rather than store the data in an actual array, it -/// is gerenated on the fly by a function. This is handled in VTK-m by creating -/// an ArrayHandle in VTK-m with a StorageTagImplicit type of \c Storage. This -/// tag itself is templated to specify an ArrayPortal that generates the -/// desired values. An ArrayHandle created with this tag will raise an error on -/// any operation that tries to modify it. -/// -template -struct VTKM_ALWAYS_EXPORT StorageTagImplicit -{ - using PortalType = ArrayPortalType; -}; - -namespace internal -{ - -template -class VTKM_ALWAYS_EXPORT - Storage> -{ - using ClassType = - Storage>; - -public: - using ValueType = typename ArrayPortalType::ValueType; - using PortalConstType = ArrayPortalType; - - // Note that this portal is likely to be read-only, so you will probably get an error - // if you try to write to it. - using PortalType = ArrayPortalType; - - VTKM_CONT - Storage(const PortalConstType& portal = PortalConstType()) - : Portal(portal) - , NumberOfValues(portal.GetNumberOfValues()) - { - } - - VTKM_CONT Storage(const ClassType&) = default; - VTKM_CONT Storage(ClassType&&) = default; - VTKM_CONT ClassType& operator=(const ClassType&) = default; - VTKM_CONT ClassType& operator=(ClassType&&) = default; - - // All these methods do nothing but raise errors. - VTKM_CONT - PortalType GetPortal() { throw vtkm::cont::ErrorBadValue("Implicit arrays are read-only."); } - VTKM_CONT - PortalConstType GetPortalConst() const { return this->Portal; } - VTKM_CONT - vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; } - VTKM_CONT - void Allocate(vtkm::Id numberOfValues) - { - VTKM_ASSERT(numberOfValues <= this->Portal.GetNumberOfValues()); - this->NumberOfValues = numberOfValues; - } - VTKM_CONT - void Shrink(vtkm::Id numberOfValues) - { - VTKM_ASSERT(numberOfValues <= this->Portal.GetNumberOfValues()); - this->NumberOfValues = numberOfValues; - } - VTKM_CONT - void ReleaseResources() {} - -private: - PortalConstType Portal; - vtkm::Id NumberOfValues; -}; - -template -class ArrayTransfer, DeviceAdapterTag> -{ -public: - using StorageTag = StorageTagImplicit; - using StorageType = vtkm::cont::internal::Storage; - - using ValueType = T; - - using PortalControl = typename StorageType::PortalType; - using PortalConstControl = typename StorageType::PortalConstType; - using PortalExecution = PortalControl; - using PortalConstExecution = PortalConstControl; - - VTKM_CONT - ArrayTransfer(StorageType* storage) - : Storage(storage) - { - } - - VTKM_CONT - vtkm::Id GetNumberOfValues() const { return this->Storage->GetNumberOfValues(); } - - VTKM_CONT - PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData), vtkm::cont::Token&) - { - return this->Storage->GetPortalConst(); - } - -#if defined(VTKM_GCC) && defined(VTKM_ENABLE_OPENMP) && (__GNUC__ == 6 && __GNUC_MINOR__ == 1) -// When using GCC 6.1 with OpenMP enabled we cause a compiler ICE that is -// an identified compiler regression (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=71210) -// The easiest way to work around this is to make sure we aren't building with >= O2 -#define NO_OPTIMIZE_FUNC_ATTRIBUTE __attribute__((optimize(1))) -#else // gcc 6.1 openmp compiler ICE workaround -#define NO_OPTIMIZE_FUNC_ATTRIBUTE -#endif - - VTKM_CONT - NO_OPTIMIZE_FUNC_ATTRIBUTE - PortalExecution PrepareForInPlace(bool vtkmNotUsed(updateData), vtkm::cont::Token&) - { - throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be used for output or in place."); - } - - VTKM_CONT - NO_OPTIMIZE_FUNC_ATTRIBUTE - PortalExecution PrepareForOutput(vtkm::Id vtkmNotUsed(numberOfValues), vtkm::cont::Token&) - { - throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be used for output."); - } - -#undef NO_OPTIMIZE_FUNC_ATTRIBUTE - - VTKM_CONT - void RetrieveOutputData(StorageType* vtkmNotUsed(controlArray)) const - { - throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be used for output."); - } - - VTKM_CONT - void Shrink(vtkm::Id vtkmNotUsed(numberOfValues)) - { - throw vtkm::cont::ErrorBadValue("Implicit arrays cannot be resized."); - } - - VTKM_CONT - void ReleaseResources() {} - -private: - StorageType* Storage; -}; - -} // namespace internal -} -} // namespace vtkm::cont - -#endif //vtk_m_cont_StorageImplicit diff --git a/vtkm/cont/testing/CMakeLists.txt b/vtkm/cont/testing/CMakeLists.txt index 6ed6c21a7..fdd33cfc3 100644 --- a/vtkm/cont/testing/CMakeLists.txt +++ b/vtkm/cont/testing/CMakeLists.txt @@ -81,7 +81,6 @@ set(unit_tests UnitTestRuntimeDeviceInformation.cxx UnitTestRuntimeDeviceNames.cxx UnitTestScopedRuntimeDeviceTracker.cxx - UnitTestStorageImplicit.cxx UnitTestStorageList.cxx UnitTestStorageListTag.cxx UnitTestTimer.cxx diff --git a/vtkm/cont/testing/UnitTestArrayHandleCounting.cxx b/vtkm/cont/testing/UnitTestArrayHandleCounting.cxx index b5e975fa3..1b1ff7da5 100644 --- a/vtkm/cont/testing/UnitTestArrayHandleCounting.cxx +++ b/vtkm/cont/testing/UnitTestArrayHandleCounting.cxx @@ -18,57 +18,8 @@ namespace UnitTestArrayHandleCountingNamespace const vtkm::Id ARRAY_SIZE = 10; -// An unusual data type that represents a number with a string of a -// particular length. This makes sure that the ArrayHandleCounting -// works correctly with type casts. -class StringInt -{ -public: - StringInt() {} - StringInt(vtkm::Id v) - { - VTKM_ASSERT(v >= 0); - for (vtkm::Id i = 0; i < v; i++) - { - ++(*this); - } - } - - operator vtkm::Id() const { return vtkm::Id(this->Value.size()); } - - StringInt operator+(const StringInt& rhs) const { return StringInt(this->Value + rhs.Value); } - - StringInt operator*(const StringInt& rhs) const - { - StringInt result; - for (vtkm::Id i = 0; i < rhs; i++) - { - result = result + *this; - } - return result; - } - - bool operator==(const StringInt& other) const { return this->Value.size() == other.Value.size(); } - - StringInt& operator++() - { - this->Value.append("."); - return *this; - } - -private: - StringInt(const std::string& v) - : Value(v) - { - } - - std::string Value; -}; - } // namespace UnitTestArrayHandleCountingNamespace -VTKM_BASIC_TYPE_VECTOR(UnitTestArrayHandleCountingNamespace::StringInt) - namespace UnitTestArrayHandleCountingNamespace { @@ -81,7 +32,7 @@ struct TemplatedTests using PortalType = typename vtkm::cont::internal::Storage::PortalConstType; + typename ArrayHandleType::StorageTag>::ReadPortalType; void operator()(const ValueType& startingValue, const ValueType& step) { @@ -90,7 +41,7 @@ struct TemplatedTests ArrayHandleType arrayMake = vtkm::cont::make_ArrayHandleCounting(startingValue, step, ARRAY_SIZE); - ArrayHandleType2 arrayHandle = ArrayHandleType2(PortalType(startingValue, step, ARRAY_SIZE)); + ArrayHandleType2 arrayHandle = ArrayHandleType(startingValue, step, ARRAY_SIZE); VTKM_TEST_ASSERT(arrayConst.GetNumberOfValues() == ARRAY_SIZE, "Counting array using constructor has wrong size."); @@ -127,8 +78,6 @@ void TestArrayHandleCounting() TemplatedTests()(3.0f, -0.5f); TemplatedTests()(0.0, 1.0); TemplatedTests()(-3.0, 2.0); - TemplatedTests()(StringInt(0), StringInt(1)); - TemplatedTests()(StringInt(10), StringInt(2)); } diff --git a/vtkm/cont/testing/UnitTestStorageImplicit.cxx b/vtkm/cont/testing/UnitTestStorageImplicit.cxx deleted file mode 100644 index 8f1ed88fe..000000000 --- a/vtkm/cont/testing/UnitTestStorageImplicit.cxx +++ /dev/null @@ -1,134 +0,0 @@ -//============================================================================ -// Copyright (c) Kitware, Inc. -// All rights reserved. -// See LICENSE.txt for details. -// -// This software is distributed WITHOUT ANY WARRANTY; without even -// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR -// PURPOSE. See the above copyright notice for more information. -//============================================================================ - -#include -#include - -#include -#include - -#include - -#include - -#if defined(VTKM_STORAGE) -#undef VTKM_STORAGE -#endif - -#define VTKM_STORAGE VTKM_STORAGE_ERROR - -namespace -{ - -const vtkm::Id ARRAY_SIZE = 10; - -template -struct TestImplicitStorage -{ - using ValueType = T; - ValueType Temp; - - VTKM_EXEC_CONT - TestImplicitStorage() - : Temp(1) - { - } - - VTKM_EXEC_CONT - vtkm::Id GetNumberOfValues() const { return ARRAY_SIZE; } - - VTKM_EXEC_CONT - ValueType Get(vtkm::Id vtkmNotUsed(index)) const { return Temp; } -}; - -template -struct TemplatedTests -{ - using StorageTagType = vtkm::cont::StorageTagImplicit>; - using StorageType = vtkm::cont::internal::Storage; - - using ValueType = typename StorageType::ValueType; - using PortalType = typename StorageType::PortalType; - - void BasicAllocation() - { - StorageType arrayStorage; - - // The implicit portal defined for this test always returns ARRAY_SIZE for the - // number of values. We should get that. - VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE, - "Implicit Storage GetNumberOfValues returned wrong size."); - - // Make sure you can allocate and shrink to any value <= the reported portal size. - arrayStorage.Allocate(ARRAY_SIZE / 2); - VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE / 2, - "Cannot re-Allocate array to half size."); - - arrayStorage.Allocate(0); - VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == 0, "Cannot re-Allocate array to zero."); - - arrayStorage.Allocate(ARRAY_SIZE); - VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE, - "Cannot re-Allocate array to original size."); - - arrayStorage.Shrink(ARRAY_SIZE / 2); - VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE / 2, - "Cannot Shrink array to half size."); - - arrayStorage.Shrink(0); - VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == 0, "Cannot Shrink array to zero."); - - arrayStorage.Shrink(ARRAY_SIZE); - VTKM_TEST_ASSERT(arrayStorage.GetNumberOfValues() == ARRAY_SIZE, - "Cannot Shrink array to original size."); - - //verify that calling ReleaseResources doesn't throw an exception - arrayStorage.ReleaseResources(); - - //verify that you can allocate after releasing resources. - arrayStorage.Allocate(ARRAY_SIZE); - } - - void BasicAccess() - { - TestImplicitStorage portal; - vtkm::cont::ArrayHandle implictHandle(portal); - VTKM_TEST_ASSERT(implictHandle.GetNumberOfValues() == ARRAY_SIZE, "handle has wrong size"); - VTKM_TEST_ASSERT(implictHandle.ReadPortal().Get(0) == T(1), "portals first values should be 1"); - } - - void operator()() - { - BasicAllocation(); - BasicAccess(); - } -}; - -struct TestFunctor -{ - template - void operator()(T) const - { - TemplatedTests tests; - tests(); - } -}; - -void TestStorageBasic() -{ - vtkm::testing::Testing::TryTypes(TestFunctor()); -} - -} // Anonymous namespace - -int UnitTestStorageImplicit(int argc, char* argv[]) -{ - return vtkm::cont::testing::Testing::Run(TestStorageBasic, argc, argv); -} diff --git a/vtkm/internal/ArrayPortalUniformPointCoordinates.h b/vtkm/internal/ArrayPortalUniformPointCoordinates.h index ef4f2010d..aaeda6f5f 100644 --- a/vtkm/internal/ArrayPortalUniformPointCoordinates.h +++ b/vtkm/internal/ArrayPortalUniformPointCoordinates.h @@ -43,25 +43,6 @@ public: { } - VTKM_EXEC_CONT - ArrayPortalUniformPointCoordinates(const ArrayPortalUniformPointCoordinates& src) - : Dimensions(src.Dimensions) - , NumberOfValues(src.NumberOfValues) - , Origin(src.Origin) - , Spacing(src.Spacing) - { - } - - VTKM_EXEC_CONT - ArrayPortalUniformPointCoordinates& operator=(const ArrayPortalUniformPointCoordinates& src) - { - this->Dimensions = src.Dimensions; - this->NumberOfValues = src.NumberOfValues; - this->Origin = src.Origin; - this->Spacing = src.Spacing; - return *this; - } - VTKM_EXEC_CONT vtkm::Id GetNumberOfValues() const { return this->NumberOfValues; } From 77e1aec50af3dc6052a6b784a87387bb40cfb299 Mon Sep 17 00:00:00 2001 From: Li-Ta Lo Date: Wed, 21 Oct 2020 16:08:06 -0600 Subject: [PATCH 07/19] minor change based on code review --- vtkm/filter/CMakeLists.txt | 4 ++-- ...GP.h => ParticleDensityNearestGridPoint.h} | 13 ++++++++----- ...xx => ParticleDensityNearestGridPoint.hxx} | 19 ++++++++++++++----- .../testing/UnitTestParticleDensity.cxx | 8 ++++---- 4 files changed, 28 insertions(+), 16 deletions(-) rename vtkm/filter/{ParticleDensityNGP.h => ParticleDensityNearestGridPoint.h} (76%) rename vtkm/filter/{ParticleDensityNGP.hxx => ParticleDensityNearestGridPoint.hxx} (83%) diff --git a/vtkm/filter/CMakeLists.txt b/vtkm/filter/CMakeLists.txt index 20a17a9e4..7480b72e2 100644 --- a/vtkm/filter/CMakeLists.txt +++ b/vtkm/filter/CMakeLists.txt @@ -81,7 +81,7 @@ set(extra_headers MeshQuality.h NDEntropy.h NDHistogram.h - ParticleDensityNGP.h + ParticleDensityNearestGridPoint.h ParticleAdvection.h Pathline.h PointAverage.h @@ -130,7 +130,7 @@ set(extra_header_template_sources MeshQuality.hxx NDEntropy.hxx NDHistogram.hxx - ParticleDensityNGP.hxx + ParticleDensityNearestGridPoint.hxx ParticleAdvection.hxx Pathline.hxx PointAverage.hxx diff --git a/vtkm/filter/ParticleDensityNGP.h b/vtkm/filter/ParticleDensityNearestGridPoint.h similarity index 76% rename from vtkm/filter/ParticleDensityNGP.h rename to vtkm/filter/ParticleDensityNearestGridPoint.h index 9fb99612a..092e6d2cf 100644 --- a/vtkm/filter/ParticleDensityNGP.h +++ b/vtkm/filter/ParticleDensityNearestGridPoint.h @@ -20,16 +20,19 @@ namespace filter /// \brief Estimate the density of particles using the Nearest Grid Point method // We only need the CoordinateSystem of the input dataset thus a FilterField -class ParticleDensityNGP : public vtkm::filter::FilterField +class ParticleDensityNearestGridPoint + : public vtkm::filter::FilterField { public: // ParticleDensity only support turning 2D/3D particle positions into density using SupportedTypes = vtkm::TypeListFieldVec3; // - ParticleDensityNGP(const vtkm::Id3& dimension, - const vtkm::Vec3f& origin, - const vtkm::Vec3f& spacing); + ParticleDensityNearestGridPoint(const vtkm::Id3& dimension, + const vtkm::Vec3f& origin, + const vtkm::Vec3f& spacing); + + ParticleDensityNearestGridPoint(const vtkm::Id3& dimension, const vtkm::Bounds& bounds); template VTKM_CONT vtkm::cont::DataSet DoExecute(const vtkm::cont::DataSet& input, @@ -45,6 +48,6 @@ private: } } -#include +#include #endif //vtk_m_filter_particle_density_ngp_h diff --git a/vtkm/filter/ParticleDensityNGP.hxx b/vtkm/filter/ParticleDensityNearestGridPoint.hxx similarity index 83% rename from vtkm/filter/ParticleDensityNGP.hxx rename to vtkm/filter/ParticleDensityNearestGridPoint.hxx index 0fea15ea5..702438e26 100644 --- a/vtkm/filter/ParticleDensityNGP.hxx +++ b/vtkm/filter/ParticleDensityNearestGridPoint.hxx @@ -11,7 +11,7 @@ #ifndef vtk_m_filter_particle_density_ngp_hxx #define vtk_m_filter_particle_density_ngp_hxx -#include "ParticleDensityNGP.h" +#include "ParticleDensityNearestGridPoint.h" #include #include #include @@ -54,17 +54,26 @@ namespace vtkm { namespace filter { -inline VTKM_CONT ParticleDensityNGP::ParticleDensityNGP(const vtkm::Id3& dimension, - const vtkm::Vec3f& origin, - const vtkm::Vec3f& spacing) +inline VTKM_CONT ParticleDensityNearestGridPoint::ParticleDensityNearestGridPoint( + const vtkm::Id3& dimension, + const vtkm::Vec3f& origin, + const vtkm::Vec3f& spacing) : Dimension(dimension) , Origin(origin) , Spacing(spacing) { } +ParticleDensityNearestGridPoint::ParticleDensityNearestGridPoint(const Id3& dimension, + const vtkm::Bounds& bounds) + : Dimension(dimension) + , Origin({ bounds.X.Min, bounds.Y.Min, bounds.Y.Min }) + , Spacing(vtkm::Vec3f{ bounds.X.Length(), bounds.Y.Length(), bounds.Z.Length() } / Dimension) +{ +} + template -inline VTKM_CONT vtkm::cont::DataSet ParticleDensityNGP::DoExecute( +inline VTKM_CONT vtkm::cont::DataSet ParticleDensityNearestGridPoint::DoExecute( const vtkm::cont::DataSet&, const vtkm::cont::ArrayHandle& field, const vtkm::filter::FieldMetadata&, diff --git a/vtkm/filter/testing/UnitTestParticleDensity.cxx b/vtkm/filter/testing/UnitTestParticleDensity.cxx index f5dab0594..c0d605dcb 100644 --- a/vtkm/filter/testing/UnitTestParticleDensity.cxx +++ b/vtkm/filter/testing/UnitTestParticleDensity.cxx @@ -11,7 +11,7 @@ #include #include #include -#include +#include #include void TestNGP() @@ -31,9 +31,9 @@ void TestNGP() positions, vtkm::CellShapeTagVertex{}, 1, connectivity); auto cellDims = vtkm::Id3{ 3, 3, 3 }; - vtkm::filter::ParticleDensityNGP filter{ cellDims, - { 0.f, 0.f, 0.f }, - vtkm::Vec3f{ 1.f / 3.f, 1.f / 3.f, 1.f / 3.f } }; + vtkm::filter::ParticleDensityNearestGridPoint filter{ + cellDims, { 0.f, 0.f, 0.f }, vtkm::Vec3f{ 1.f / 3.f, 1.f / 3.f, 1.f / 3.f } + }; filter.SetUseCoordinateSystemAsField(true); auto density = filter.Execute(dataSet); From 8b6ae4f78c2ac20ff4597e331dc4edecd0056b4e Mon Sep 17 00:00:00 2001 From: Li-Ta Lo Date: Wed, 21 Oct 2020 16:32:04 -0600 Subject: [PATCH 08/19] fixed type conversion and typo --- vtkm/filter/ParticleDensityNearestGridPoint.hxx | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/vtkm/filter/ParticleDensityNearestGridPoint.hxx b/vtkm/filter/ParticleDensityNearestGridPoint.hxx index 702438e26..fc92735c3 100644 --- a/vtkm/filter/ParticleDensityNearestGridPoint.hxx +++ b/vtkm/filter/ParticleDensityNearestGridPoint.hxx @@ -67,8 +67,13 @@ inline VTKM_CONT ParticleDensityNearestGridPoint::ParticleDensityNearestGridPoin ParticleDensityNearestGridPoint::ParticleDensityNearestGridPoint(const Id3& dimension, const vtkm::Bounds& bounds) : Dimension(dimension) - , Origin({ bounds.X.Min, bounds.Y.Min, bounds.Y.Min }) - , Spacing(vtkm::Vec3f{ bounds.X.Length(), bounds.Y.Length(), bounds.Z.Length() } / Dimension) + , Origin({ static_cast(bounds.X.Min), + static_cast(bounds.Y.Min), + static_cast(bounds.Z.Min) }) + , Spacing(vtkm::Vec3f{ static_cast(bounds.X.Length()), + static_cast(bounds.Y.Length()), + static_cast(bounds.Z.Length()) } / + Dimension) { } From b22823a547b9b8fa0b167e20bddaec31d9eded9e Mon Sep 17 00:00:00 2001 From: Vicente Adolfo Bolea Sanchez Date: Tue, 13 Oct 2020 16:19:31 -0400 Subject: [PATCH 09/19] Contour: Zero-init arrays in Flying Edges - It fixes a lingering error triggered with BenchContour - It reenables BenchContour Signed-off-by: Vicente Adolfo Bolea Sanchez --- benchmarking/BenchmarkFilters.cxx | 2 +- vtkm/worklet/contour/FlyingEdgesPass1.h | 5 +++-- vtkm/worklet/contour/FlyingEdgesPass4Common.h | 2 +- 3 files changed, 5 insertions(+), 4 deletions(-) diff --git a/benchmarking/BenchmarkFilters.cxx b/benchmarking/BenchmarkFilters.cxx index 76411a93d..180d4355c 100644 --- a/benchmarking/BenchmarkFilters.cxx +++ b/benchmarking/BenchmarkFilters.cxx @@ -405,7 +405,7 @@ void BenchContourGenerator(::benchmark::internal::Benchmark* bm) } // :TODO: Disabled until SIGSEGV in Countour when passings field is resolved -//VTKM_BENCHMARK_APPLY(BenchContour, BenchContourGenerator); +VTKM_BENCHMARK_APPLY(BenchContour, BenchContourGenerator); void BenchExternalFaces(::benchmark::State& state) { diff --git a/vtkm/worklet/contour/FlyingEdgesPass1.h b/vtkm/worklet/contour/FlyingEdgesPass1.h index e683b2992..c1fa0d5e1 100644 --- a/vtkm/worklet/contour/FlyingEdgesPass1.h +++ b/vtkm/worklet/contour/FlyingEdgesPass1.h @@ -141,6 +141,7 @@ struct ComputePass1 : public vtkm::worklet::WorkletVisitPointsWithCells } } } + write_edge(device, startPos + (offset * end), edges, FlyingEdges3D::Below); } }; @@ -150,7 +151,7 @@ struct launchComputePass1 VTKM_CONT bool operator()(DeviceAdapterTag device, const ComputePass1& worklet, const vtkm::cont::ArrayHandle& inputField, - vtkm::cont::ArrayHandle edgeCases, + vtkm::cont::ArrayHandle& edgeCases, vtkm::cont::CellSetStructured<2>& metaDataMesh2D, Args&&... args) const { @@ -165,7 +166,7 @@ struct launchComputePass1 VTKM_CONT bool operator()(vtkm::cont::DeviceAdapterTagCuda device, const ComputePass1& worklet, const vtkm::cont::ArrayHandle& inputField, - vtkm::cont::ArrayHandle edgeCases, + vtkm::cont::ArrayHandle& edgeCases, vtkm::cont::CellSetStructured<2>& metaDataMesh2D, Args&&... args) const { diff --git a/vtkm/worklet/contour/FlyingEdgesPass4Common.h b/vtkm/worklet/contour/FlyingEdgesPass4Common.h index 268a48d4b..3c8490bfa 100644 --- a/vtkm/worklet/contour/FlyingEdgesPass4Common.h +++ b/vtkm/worklet/contour/FlyingEdgesPass4Common.h @@ -171,7 +171,7 @@ struct Pass4TrimState } if (ijk[AxisToSum::xindex] >= (pdims[AxisToSum::xindex] - 2)) { - boundaryStatus[AxisToSum::yindex] += FlyingEdges3D::MaxBoundary; + boundaryStatus[AxisToSum::xindex] += FlyingEdges3D::MaxBoundary; } if (ijk[AxisToSum::yindex] < 1) { From 1add731eff348733b4b1eb01aac666b2e4469763 Mon Sep 17 00:00:00 2001 From: Vicente Adolfo Bolea Sanchez Date: Tue, 13 Oct 2020 16:34:02 -0400 Subject: [PATCH 10/19] BenchmarkFilters: Add Unstructured bench for Contour --- benchmarking/BenchmarkFilters.cxx | 54 +++++++++++++++++-------------- 1 file changed, 29 insertions(+), 25 deletions(-) diff --git a/benchmarking/BenchmarkFilters.cxx b/benchmarking/BenchmarkFilters.cxx index 180d4355c..bf07837eb 100644 --- a/benchmarking/BenchmarkFilters.cxx +++ b/benchmarking/BenchmarkFilters.cxx @@ -347,10 +347,11 @@ void BenchContour(::benchmark::State& state) { const vtkm::cont::DeviceAdapterId device = Config.Device; - const vtkm::Id numIsoVals = static_cast(state.range(0)); - const bool mergePoints = static_cast(state.range(1)); - const bool normals = static_cast(state.range(2)); - const bool fastNormals = static_cast(state.range(3)); + const bool isStructured = static_cast(state.range(0)); + const vtkm::Id numIsoVals = static_cast(state.range(1)); + const bool mergePoints = static_cast(state.range(2)); + const bool normals = static_cast(state.range(3)); + const bool fastNormals = static_cast(state.range(4)); vtkm::filter::Contour filter; filter.SetActiveField(PointScalarsName, vtkm::cont::Field::Association::POINTS); @@ -376,11 +377,14 @@ void BenchContour(::benchmark::State& state) filter.SetComputeFastNormalsForUnstructured(fastNormals); vtkm::cont::Timer timer{ device }; + + vtkm::cont::DataSet input = isStructured ? InputDataSet : UnstructuredInputDataSet; + for (auto _ : state) { (void)_; timer.Start(); - auto result = filter.Execute(InputDataSet); + auto result = filter.Execute(input); ::benchmark::DoNotOptimize(result); timer.Stop(); @@ -390,13 +394,17 @@ void BenchContour(::benchmark::State& state) void BenchContourGenerator(::benchmark::internal::Benchmark* bm) { - bm->ArgNames({ "NIsoVals", "MergePts", "GenNormals", "FastNormals" }); + bm->ArgNames({ "IsStructuredDataSet", "NIsoVals", "MergePts", "GenNormals", "FastNormals" }); auto helper = [&](const vtkm::Id numIsoVals) { - bm->Args({ numIsoVals, 0, 0, 0 }); - bm->Args({ numIsoVals, 1, 0, 0 }); - bm->Args({ numIsoVals, 0, 1, 0 }); - bm->Args({ numIsoVals, 0, 1, 1 }); + bm->Args({ 0, numIsoVals, 0, 0, 0 }); + bm->Args({ 0, numIsoVals, 1, 0, 0 }); + bm->Args({ 0, numIsoVals, 0, 1, 0 }); + bm->Args({ 0, numIsoVals, 0, 1, 1 }); + bm->Args({ 1, numIsoVals, 0, 0, 0 }); + bm->Args({ 1, numIsoVals, 1, 0, 0 }); + bm->Args({ 1, numIsoVals, 0, 1, 0 }); + bm->Args({ 1, numIsoVals, 0, 1, 1 }); }; helper(1); @@ -993,26 +1001,22 @@ void InitDataSet(int& argc, char** argv) source.SetExtent({ 0 }, { waveletDim - 1 }); InputDataSet = source.Execute(); - - vtkm::cont::DataSet input = vtkm::cont::testing::MakeTestDataSet().Make2DUniformDataSet2(); - vtkm::filter::Triangulate triangulateFilter; - triangulateFilter.SetFieldsToPass( - vtkm::filter::FieldSelection(vtkm::filter::FieldSelection::MODE_ALL)); - UnstructuredInputDataSet = triangulateFilter.Execute(input); - } - - - if (tetra) - { - std::cerr << "[InitDataSet] Tetrahedralizing dataset...\n"; - vtkm::filter::Tetrahedralize tet; - tet.SetFieldsToPass(vtkm::filter::FieldSelection(vtkm::filter::FieldSelection::MODE_ALL)); - InputDataSet = tet.Execute(InputDataSet); } FindFields(); CreateMissingFields(); + std::cerr + << "[InitDataSet] Create UnstructuredInputDataSet from Tetrahedralized InputDataSet...\n"; + vtkm::filter::Tetrahedralize tet; + tet.SetFieldsToPass(vtkm::filter::FieldSelection(vtkm::filter::FieldSelection::MODE_ALL)); + UnstructuredInputDataSet = tet.Execute(InputDataSet); + + if (tetra) + { + InputDataSet = UnstructuredInputDataSet; + } + inputGenTimer.Stop(); std::cerr << "[InitDataSet] DataSet initialization took " << inputGenTimer.GetElapsedTime() From ff381bf8b4c58dc2ec2ffc1044f6d23070eeda7f Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 22 Oct 2020 16:48:30 +0000 Subject: [PATCH 11/19] vtkm/Swap works with hip --- vtkm/Swap.h | 18 ++++++++++++++++-- 1 file changed, 16 insertions(+), 2 deletions(-) diff --git a/vtkm/Swap.h b/vtkm/Swap.h index 3f4b1ddc0..f833a495c 100644 --- a/vtkm/Swap.h +++ b/vtkm/Swap.h @@ -13,7 +13,7 @@ #include -#ifdef __CUDACC__ +#ifdef VTKM_CUDA #include #else #include @@ -23,13 +23,27 @@ namespace vtkm { /// Performs a swap operation. Safe to call from cuda code. -#ifdef __CUDACC__ +#if defined(VTKM_CUDA) template VTKM_EXEC_CONT void Swap(T& a, T& b) { using namespace thrust; swap(a, b); } +#elif defined(VTKM_HIP) +template +__host__ void Swap(T& a, T& b) +{ + using namespace std; + swap(a, b); +} +template +__device__ void Swap(T& a, T& b) +{ + T temp = a; + a = b; + b = temp; +} #else template VTKM_EXEC_CONT void Swap(T& a, T& b) From 1ed5dfca0b8a93f46638d61c4804428166d9a445 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Thu, 22 Oct 2020 16:57:42 +0000 Subject: [PATCH 12/19] vtkm/Math frexp(float) version works now with HIP --- vtkm/Math.h | 3 ++- vtkm/Math.h.in | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/vtkm/Math.h b/vtkm/Math.h index fcdabfc94..d26e9a803 100644 --- a/vtkm/Math.h +++ b/vtkm/Math.h @@ -2562,7 +2562,8 @@ static inline VTKM_EXEC_CONT vtkm::Vec CopySign(const vtkm::Vec& x, inline VTKM_EXEC_CONT vtkm::Float32 Frexp(vtkm::Float32 x, vtkm::Int32 *exponent) { -#ifdef VTKM_CUDA + // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169 +#if defined(VTKM_CUDA) || defined(VTKM_HIP) return VTKM_CUDA_MATH_FUNCTION_32(frexp)(x, exponent); #else return std::frexp(x, exponent); diff --git a/vtkm/Math.h.in b/vtkm/Math.h.in index 4dd6dcbdc..494ef86f9 100644 --- a/vtkm/Math.h.in +++ b/vtkm/Math.h.in @@ -1164,7 +1164,8 @@ static inline VTKM_EXEC_CONT vtkm::Vec CopySign(const vtkm::Vec& x, inline VTKM_EXEC_CONT vtkm::Float32 Frexp(vtkm::Float32 x, vtkm::Int32 *exponent) { -#ifdef VTKM_CUDA + // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169 +#if defined(VTKM_CUDA) || defined(VTKM_HIP) return VTKM_CUDA_MATH_FUNCTION_32(frexp)(x, exponent); #else return std::frexp(x, exponent); From 0b2ddb83b33f70fd5d15098eeb880ad24a372b60 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Wed, 21 Oct 2020 16:55:01 -0400 Subject: [PATCH 13/19] UnitTestBounds custom compile flags expressed via generator expression --- vtkm/testing/CMakeLists.txt | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/vtkm/testing/CMakeLists.txt b/vtkm/testing/CMakeLists.txt index dbc547c66..ef4bd3493 100644 --- a/vtkm/testing/CMakeLists.txt +++ b/vtkm/testing/CMakeLists.txt @@ -55,11 +55,9 @@ set(unit_tests_device ) -#suppress gcc note: #variable tracking size limit exceeded with -fvar-tracking-assignments, retrying without -if (VTKM_COMPILER_IS_GNU) - set_source_files_properties(UnitTestBounds.cxx PROPERTIES COMPILE_OPTIONS "-fno-var-tracking-assignments") -endif() +set_source_files_properties(UnitTestBounds.cxx PROPERTIES + COMPILE_OPTIONS "$<$,$>:-fno-var-tracking-assignments>") vtkm_unit_tests(SOURCES ${unit_tests}) From b33c54bf613c928519343f3beeb60a9a7fbe5397 Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Sun, 25 Oct 2020 17:15:50 -0600 Subject: [PATCH 14/19] Add ScheduleTask to performance log When `DeviceAdapterAlgorithm::ScheduleTask` was called directly (i.e. not through `Schedule`), nothing was added to the log. Adding `VTKM_LOG_SCOPE` to these methods so that all scheduling is added to the performance log. --- vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h | 4 ++++ vtkm/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h | 4 ++++ vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.cxx | 4 ++++ vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx | 4 ++++ 4 files changed, 16 insertions(+) diff --git a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h index 8c6d8d254..4d62be8a5 100644 --- a/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h +++ b/vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h @@ -1639,6 +1639,8 @@ public: static void ScheduleTask(vtkm::exec::cuda::internal::TaskStrided1D& functor, vtkm::Id numInstances) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + VTKM_ASSERT(numInstances >= 0); if (numInstances < 1) { @@ -1671,6 +1673,8 @@ public: static void ScheduleTask(vtkm::exec::cuda::internal::TaskStrided3D& functor, vtkm::Id3 rangeMax) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + VTKM_ASSERT((rangeMax[0] >= 0) && (rangeMax[1] >= 0) && (rangeMax[2] >= 0)); if ((rangeMax[0] < 1) || (rangeMax[1] < 1) || (rangeMax[2] < 1)) { diff --git a/vtkm/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h b/vtkm/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h index be7c67f1e..fa49258f0 100644 --- a/vtkm/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h +++ b/vtkm/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h @@ -165,6 +165,8 @@ public: vtkm::exec::kokkos::internal::TaskBasic1D& functor, vtkm::Id numInstances) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + if (numInstances < 1) { // No instances means nothing to run. Just return. @@ -187,6 +189,8 @@ public: vtkm::exec::kokkos::internal::TaskBasic3D& functor, vtkm::Id3 rangeMax) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + if ((rangeMax[0] < 1) || (rangeMax[1] < 1) || (rangeMax[2] < 1)) { // No instances means nothing to run. Just return. diff --git a/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.cxx b/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.cxx index 905a54ce8..b2b0712d9 100644 --- a/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.cxx +++ b/vtkm/cont/openmp/internal/DeviceAdapterAlgorithmOpenMP.cxx @@ -24,6 +24,8 @@ void DeviceAdapterAlgorithm::ScheduleTask( vtkm::exec::openmp::internal::TaskTiling1D& functor, vtkm::Id size) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + static constexpr vtkm::Id MESSAGE_SIZE = 1024; char errorString[MESSAGE_SIZE]; errorString[0] = '\0'; @@ -66,6 +68,8 @@ void DeviceAdapterAlgorithm::ScheduleTask( vtkm::exec::openmp::internal::TaskTiling3D& functor, vtkm::Id3 size) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + static constexpr vtkm::Id MESSAGE_SIZE = 1024; char errorString[MESSAGE_SIZE]; errorString[0] = '\0'; diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx index fb88f3b7b..d5e64d3c0 100644 --- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx +++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.cxx @@ -19,6 +19,8 @@ void DeviceAdapterAlgorithm::ScheduleTask( vtkm::exec::tbb::internal::TaskTiling1D& functor, vtkm::Id size) { + VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf, "Schedule Task TBB 1D"); + const vtkm::Id MESSAGE_SIZE = 1024; char errorString[MESSAGE_SIZE]; errorString[0] = '\0'; @@ -40,6 +42,8 @@ void DeviceAdapterAlgorithm::ScheduleTask( vtkm::exec::tbb::internal::TaskTiling3D& functor, vtkm::Id3 size) { + VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf, "Schedule Task TBB 3D"); + static constexpr vtkm::UInt32 TBB_GRAIN_SIZE_3D[3] = { 1, 4, 256 }; const vtkm::Id MESSAGE_SIZE = 1024; char errorString[MESSAGE_SIZE]; From 0b1c48a3ee08528176636e64d4e88a3260d64299 Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Mon, 26 Oct 2020 09:00:09 -0600 Subject: [PATCH 15/19] Identify number of partitions in filter execution When you execute a filter, the default behavior is to do the execution on each partition of the data set independently. This code path is followed even for non-partitioned data; the `DataSet` is wrapped in a `PartitionedDataSet` of one partition. Make performance logging a bit more clear by only giving one scoped log for a basic `DataSet` and recording the number of partitions executed in the log. --- vtkm/filter/Filter.hxx | 12 ++++-------- 1 file changed, 4 insertions(+), 8 deletions(-) diff --git a/vtkm/filter/Filter.hxx b/vtkm/filter/Filter.hxx index 8cc86b767..852ebc310 100644 --- a/vtkm/filter/Filter.hxx +++ b/vtkm/filter/Filter.hxx @@ -249,9 +249,6 @@ inline VTKM_CONT Filter::~Filter() template inline VTKM_CONT vtkm::cont::DataSet Filter::Execute(const vtkm::cont::DataSet& input) { - VTKM_LOG_SCOPE( - vtkm::cont::LogLevel::Perf, "Filter: '%s'", vtkm::cont::TypeToString().c_str()); - Derived* self = static_cast(this); vtkm::cont::PartitionedDataSet output = self->Execute(vtkm::cont::PartitionedDataSet(input)); if (output.GetNumberOfPartitions() > 1) @@ -267,7 +264,8 @@ inline VTKM_CONT vtkm::cont::PartitionedDataSet Filter::Execute( const vtkm::cont::PartitionedDataSet& input) { VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf, - "Filter (PartitionedDataSet): '%s'", + "Filter (%d partitions): '%s'", + (int)input.GetNumberOfPartitions(), vtkm::cont::TypeToString().c_str()); Derived* self = static_cast(this); @@ -293,9 +291,6 @@ VTKM_CONT vtkm::cont::DataSet Filter::Execute( const vtkm::cont::DataSet& input, vtkm::filter::PolicyBase policy) { - VTKM_LOG_SCOPE( - vtkm::cont::LogLevel::Perf, "Filter: '%s'", vtkm::cont::TypeToString().c_str()); - Derived* self = static_cast(this); VTKM_DEPRECATED_SUPPRESS_BEGIN vtkm::cont::PartitionedDataSet output = @@ -316,7 +311,8 @@ VTKM_CONT vtkm::cont::PartitionedDataSet Filter::Execute( vtkm::filter::PolicyBase policy) { VTKM_LOG_SCOPE(vtkm::cont::LogLevel::Perf, - "Filter (PartitionedDataSet): '%s'", + "Filter (%d partitions): '%s'", + (int)input.GetNumberOfPartitions(), vtkm::cont::TypeToString().c_str()); Derived* self = static_cast(this); From 7c40254d396631fa62d9370f5a8f97bfc2f69cb9 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Mon, 26 Oct 2020 15:07:16 -0400 Subject: [PATCH 16/19] diy: update for corrections to allow HIP compilation --- vtkm/thirdparty/diy/update.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/vtkm/thirdparty/diy/update.sh b/vtkm/thirdparty/diy/update.sh index 29fdda655..1ee37c531 100755 --- a/vtkm/thirdparty/diy/update.sh +++ b/vtkm/thirdparty/diy/update.sh @@ -8,7 +8,7 @@ readonly name="diy" readonly ownership="Diy Upstream " readonly subtree="vtkm/thirdparty/$name/vtkm$name" readonly repo="https://gitlab.kitware.com/third-party/diy2.git" -readonly tag="for/vtk-m-20200826-master" +readonly tag="for/vtk-m-20201026-master" readonly paths=" cmake include From 25fed44bdfc182898eb421f84bef39b70635e4b2 Mon Sep 17 00:00:00 2001 From: Diy Upstream Date: Mon, 26 Oct 2020 17:18:41 -0400 Subject: [PATCH 17/19] diy 2020-10-26 (538957e7) Code extracted from: https://gitlab.kitware.com/third-party/diy2.git at commit 538957e7850b9e0007fe814bc37219152f533285 (for/vtk-m-20201026-master). --- include/vtkmdiy/master.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/vtkmdiy/master.hpp b/include/vtkmdiy/master.hpp index bff1cafa0..ccca263f9 100644 --- a/include/vtkmdiy/master.hpp +++ b/include/vtkmdiy/master.hpp @@ -166,7 +166,7 @@ namespace diy ExternalStorage* storage = 0, //!< storage object (path, method, etc.) for storing temporary blocks being shuffled in/out of core SaveBlock save = 0, //!< block save function; master manages saving if save != 0 LoadBlock load_ = 0, //!< block load function; master manages loading if load != 0 - QueuePolicy* q_policy = new QueueSizePolicy(4096)); //!< policy for managing message queues specifies maximum size of message queues to keep in memory + QueuePolicy* q_policy = nullptr); //!< policy for managing message queues specifies maximum size of message queues to keep in memory inline ~Master(); inline void clear(); @@ -370,7 +370,7 @@ Master(mpi::communicator comm, LoadBlock load_, QueuePolicy* q_policy): blocks_(create_, destroy_, storage, save, load_), - queue_policy_(q_policy), + queue_policy_( (q_policy==nullptr) ? new QueueSizePolicy(4096): q_policy), limit_(limit__), #if !defined(VTKMDIY_NO_THREADS) threads_(threads__ == -1 ? static_cast(thread::hardware_concurrency()) : threads__), From 5f0d04ad37e50b8f98654592c441e045bc3972ff Mon Sep 17 00:00:00 2001 From: Lightweight Cell Library Upstream Date: Wed, 28 Oct 2020 14:06:49 -0400 Subject: [PATCH 18/19] lcl 2020-10-28 (b7fe5404) Code extracted from: https://gitlab.kitware.com/vtk/lcl.git at commit b7fe5404f330425e44116f3e9ead66f3a30f83a7 (master). --- lcl/internal/Config.h | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/lcl/internal/Config.h b/lcl/internal/Config.h index a51ad5926..eb5aeb83c 100644 --- a/lcl/internal/Config.h +++ b/lcl/internal/Config.h @@ -13,7 +13,10 @@ #include #include -#ifdef __CUDACC__ +#if defined(__CUDACC__) +# define LCL_EXEC __device__ __host__ +#elif defined(__HIP__) +#include "hip/hip_runtime.h" //required for __device__ __host__ # define LCL_EXEC __device__ __host__ #else # define LCL_EXEC From 7a12c5e524cd27b669aad13b7f6a63cdeaf5381a Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Wed, 28 Oct 2020 18:10:38 -0600 Subject: [PATCH 19/19] Add missing logging for some TBB device algorithms Some of the algorithms listed near the end of DeviceAdapterAlgorithmTBB.h were missing the entry in the performance log letting us know it was called. --- vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h index d21dc1d0f..dc275d7ea 100644 --- a/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h +++ b/vtkm/cont/tbb/internal/DeviceAdapterAlgorithmTBB.h @@ -286,6 +286,8 @@ public: template VTKM_CONT static void Sort(vtkm::cont::ArrayHandle& values) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + //this is required to get sort to work with zip handles std::less lessOp; vtkm::cont::tbb::sort::parallel_sort(values, lessOp); @@ -295,6 +297,8 @@ public: VTKM_CONT static void Sort(vtkm::cont::ArrayHandle& values, BinaryCompare binary_compare) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::tbb::sort::parallel_sort(values, binary_compare); } @@ -302,6 +306,8 @@ public: VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle& keys, vtkm::cont::ArrayHandle& values) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::tbb::sort::parallel_sort_bykey(keys, values, std::less()); } @@ -310,6 +316,8 @@ public: vtkm::cont::ArrayHandle& values, BinaryCompare binary_compare) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::cont::tbb::sort::parallel_sort_bykey(keys, values, binary_compare); } @@ -323,6 +331,8 @@ public: VTKM_CONT static void Unique(vtkm::cont::ArrayHandle& values, BinaryCompare binary_compare) { + VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf); + vtkm::Id outputSize; { vtkm::cont::Token token;