Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-16 17:22:55 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Merge branch 'master' of https://gitlab.kitware.com/vtk/vtk-m into zero_velocity_particle_advection_terminate

Browse Source
This commit is contained in:
Dave Pugmire 2022-08-29 15:23:10 -04:00
commit 8aa56b7aff
104 changed files with 1275 additions and 861 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
.gitlab-ci.yml
View File

@ -126,9 +126,9 @@
when: on_success
- when: never
.run_master: &run_master
.run_upstream_branches: &run_upstream_branches
rules:
- if: '$CI_PROJECT_PATH == "vtk/vtk-m" && $CI_COMMIT_BRANCH == "master"'
- if: '$CI_PROJECT_PATH == "vtk/vtk-m" && $CI_MERGE_REQUEST_ID == null'
when: on_success
- when: never

4
.gitlab/ci/ubuntu1604.yml
View File

@ -50,7 +50,7 @@ build:ubuntu1604_gcc5_2:
extends:
- .ubuntu1604_cuda
- .cmake_build_linux
- .run_master
- .run_upstream_branches
- .use_minimum_supported_cmake
variables:
CC: "gcc-5"
@ -69,7 +69,7 @@ test:ubuntu1804_test_ubuntu1604_gcc5_2:
extends:
- .ubuntu1804_cuda
- .cmake_test_linux
- .run_master
- .run_upstream_branches
variables:
CTEST_EXCLUSIONS: "built_against_test_install"
dependencies:

4
.gitlab/ci/ubuntu1804.yml
View File

@ -91,7 +91,7 @@ build:ubuntu1804_clang_cuda:
- .ubuntu1804_cuda
- .cmake_build_linux
- .run_automatically
# - .run_master
# - .run_upstream_branches
variables:
CC: "clang-8"
CXX: "clang++-8"
@ -110,7 +110,7 @@ test:ubuntu1804_clang_cuda:
- .ubuntu1804_cuda
- .cmake_test_linux
- .run_automatically
# - .run_master
# - .run_upstream_branches
dependencies:
- build:ubuntu1804_clang_cuda
needs:

4
.gitlab/ci/ubuntu2004.yml
View File

@ -69,7 +69,7 @@ build:ubuntu2004_hip_kokkos:
extends:
- .ubuntu2004_hip_kokkos
- .cmake_build_linux
- .run_scheduled
- .run_upstream_branches
variables:
CMAKE_BUILD_TYPE: RelWithDebInfo
VTKM_SETTINGS: "benchmarks+kokkos+hip+no_virtual+no_rendering+ccache"
@ -96,7 +96,7 @@ test:ubuntu2004_hip_kokkos:
extends:
- .ubuntu2004_hip_kokkos
- .cmake_test_linux
- .run_scheduled
- .run_upstream_branches
variables:
CTEST_TIMEOUT: "30"
dependencies:

2
CMake/testing/VTKmCheckSourceInInstall.cmake
View File

@ -114,6 +114,8 @@ function(do_verify root_dir prefix)
set(file_exceptions
thirdparty/diy/vtkmdiy/cmake/mpi_types.h
thirdparty/lodepng/vtkmlodepng/lodepng.h
thirdparty/loguru/vtkmloguru/loguru.hpp
# Ignore deprecated virtual classes (which are not installed if VTKm_NO_DEPRECATED_VIRTUAL
# is on). These exceptions can be removed when these files are completely removed.

3
CMake/testing/VTKmTestWrappers.cmake
View File

@ -144,7 +144,8 @@ function(vtkm_unit_tests)
if(VTKm_ENABLE_KOKKOS AND (enable_all_backends OR NOT per_device_suffix))
list(APPEND per_device_command_line_arguments --vtkm-device=kokkos)
list(APPEND per_device_suffix "KOKKOS")
list(APPEND per_device_timeout 600)
#may require more time because of kernel generation.
list(APPEND per_device_timeout 1500)
list(APPEND per_device_serial FALSE)
endif()
if(VTKm_ENABLE_TBB AND (enable_all_backends OR NOT per_device_suffix))

1
config/vtkm.pc.in
View File

@ -27,6 +27,5 @@ Libs: -L${libdir} \
-lvtkm_worklet-@VTKm_VERSION@ \
-lvtkm_source-@VTKm_VERSION@ \
-lvtkm_io-@VTKm_VERSION@ \
-lvtkm_lodepng-@VTKm_VERSION@ \
-lvtkm_cont-@VTKm_VERSION@ \
-lvtkmdiympi_nompi

1
config/vtkm_config.mk.in
View File

@ -34,6 +34,5 @@ VTKm_LIB_FLAGS = -L $(VTKm_DIR)/lib \
-lvtkm_worklet-$(VTKM_VERSION) \
-lvtkm_source-$(VTKM_VERSION) \
-lvtkm_io-$(VTKM_VERSION) \
-lvtkm_lodepng-$(VTKM_VERSION) \
-lvtkm_cont-$(VTKM_VERSION) \
-lvtkmdiympi_nompi

3
data/data/misc/warpXfields.vtk Normal file
View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:7a569aff0c6872611ef75a4dcb597e1320cb966b80f840a0dbd76a29b2760dbb
size 50335052

3
data/data/misc/warpXparticles.vtk Normal file
View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:671345bdb045aeadc8e9fa1060de51e53286c74929c6c8c60a529b318f02bbfc
size 3795

5
docs/changelog/variant-istype.md Normal file
View File

@ -0,0 +1,5 @@
# Add Variant::IsType
The `Variant` class was missing a way to check the type. You could do it
indirectly using `variant.GetIndex() == variant.GetIndexOf<T>()`, but
having this convenience function is more clear.

4
examples/clipping/CMakeLists.txt
View File

@ -15,7 +15,3 @@ find_package(VTKm REQUIRED QUIET)
add_executable(Clipping Clipping.cxx)
target_link_libraries(Clipping PRIVATE vtkm_filter vtkm_io)
vtkm_add_target_information(Clipping
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES Clipping.cxx)

6
examples/contour_tree/CMakeLists.txt
View File

@ -18,9 +18,3 @@ target_link_libraries(ContourTreeMesh2D vtkm_filter)
add_executable(ContourTreeMesh3D ContourTreeMesh3D.cxx)
target_link_libraries(ContourTreeMesh3D vtkm_filter)
vtkm_add_target_information(ContourTreeMesh2D ContourTreeMesh3D
DROP_UNUSED_SYMBOLS
MODIFY_CUDA_FLAGS
DEVICE_SOURCES
ContourTreeMesh2D.cxx ContourTreeMesh3D.cxx)

4
examples/histogram/CMakeLists.txt
View File

@ -16,6 +16,6 @@ if (VTKm_ENABLE_MPI)
add_executable(Histogram Histogram.cxx HistogramMPI.h HistogramMPI.hxx)
target_link_libraries(Histogram PRIVATE vtkm_filter MPI::MPI_CXX)
vtkm_add_target_information(Histogram
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES Histogram.cxx)
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES Histogram.cxx)
endif()

8
examples/mesh_quality/CMakeLists.txt
View File

@ -27,11 +27,3 @@ find_package(VTKm REQUIRED QUIET)
add_executable(MeshQuality MeshQuality.cxx)
target_link_libraries(MeshQuality PRIVATE vtkm_filter vtkm_io)
if(TARGET vtkm::tbb)
target_compile_definitions(MeshQuality PRIVATE BUILDING_TBB_VERSION)
endif()
vtkm_add_target_information(MeshQuality
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES MeshQuality.cxx)

3
examples/oscillator/CMakeLists.txt
View File

@ -15,6 +15,3 @@ find_package(VTKm REQUIRED QUIET)
add_executable(Oscillator Oscillator.cxx)
target_link_libraries(Oscillator PRIVATE vtkm_source)
vtkm_add_target_information(Oscillator
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES Oscillator.cxx)

6
examples/oscillator/Oscillator.cxx
View File

@ -9,19 +9,13 @@
//============================================================================
#include <algorithm>
#include <cctype>
#include <fstream>
#include <iostream>
#include <sstream>
#include <vtkm/Math.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/DataSetBuilderUniform.h>
#include <vtkm/cont/Initialize.h>
#include <vtkm/filter/FilterDataSet.h>
#include <vtkm/cont/TryExecute.h>
#include <vtkm/source/Oscillator.h>
#if !defined(_WIN32) || defined(__CYGWIN__)

6
examples/polyline_archimedean_helix/CMakeLists.txt
View File

@ -13,8 +13,8 @@ project(PolyLineArchimedeanHelix CXX)
find_package(VTKm REQUIRED QUIET)
if (VTKm_ENABLE_RENDERING)
add_executable(PolyLineArchimedeanHelix PolyLineArchimedeanHelix.cxx)
vtkm_add_target_information(PolyLineArchimedeanHelix
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES PolyLineArchimedeanHelix.cxx)
target_link_libraries(PolyLineArchimedeanHelix PRIVATE vtkm_filter vtkm_rendering)
vtkm_add_target_information(PolyLineArchimedeanHelix
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES PolyLineArchimedeanHelix.cxx)
endif()

6
examples/tetrahedra/CMakeLists.txt
View File

@ -18,9 +18,3 @@ target_link_libraries(Tetrahedralize PRIVATE vtkm_filter vtkm_io)
add_executable(Triangulate Triangulate.cxx)
target_link_libraries(Triangulate PRIVATE vtkm_filter vtkm_io)
vtkm_add_target_information(Tetrahedralize Triangulate
DROP_UNUSED_SYMBOLS
MODIFY_CUDA_FLAGS
DEVICE_SOURCES
Tetrahedralize.cxx Triangulate.cxx)

37
vtkm/Particle.h
View File

@ -176,9 +176,9 @@ public:
VTKM_EXEC_CONT
ChargedParticle(const vtkm::Vec3f& position,
const vtkm::Id& id,
const vtkm::FloatDefault& mass,
const vtkm::FloatDefault& charge,
const vtkm::FloatDefault& weighting,
const vtkm::Float64& mass,
const vtkm::Float64& charge,
const vtkm::Float64& weighting,
const vtkm::Vec3f& momentum,
const vtkm::Id& numSteps = 0,
const vtkm::ParticleStatus& status = vtkm::ParticleStatus(),
@ -196,16 +196,16 @@ public:
}
VTKM_EXEC_CONT
vtkm::FloatDefault Gamma(vtkm::Vec3f momentum, bool reciprocal = false) const
vtkm::Float64 Gamma(vtkm::Vec3f momentum, bool reciprocal = false) const
{
constexpr vtkm::FloatDefault c2 = SPEED_OF_LIGHT * SPEED_OF_LIGHT;
const auto fMom2 = vtkm::MagnitudeSquared(momentum);
const auto m2 = this->Mass * this->Mass;
const auto m2_c2_reci = 1.0 / (m2 * c2);
const vtkm::Float64 fMom2 = vtkm::MagnitudeSquared(momentum);
const vtkm::Float64 m2 = this->Mass * this->Mass;
const vtkm::Float64 m2_c2_reci = 1.0 / (m2 * c2);
if (reciprocal)
return static_cast<vtkm::FloatDefault>(vtkm::RSqrt(1.0 + fMom2 * m2_c2_reci));
return vtkm::RSqrt(1.0 + fMom2 * m2_c2_reci);
else
return static_cast<vtkm::FloatDefault>(vtkm::Sqrt(1.0 + fMom2 * m2_c2_reci));
return vtkm::Sqrt(1.0 + fMom2 * m2_c2_reci);
}
VTKM_EXEC_CONT
@ -220,11 +220,11 @@ public:
vtkm::Vec3f eField = vectors[0];
vtkm::Vec3f bField = vectors[1];
const vtkm::FloatDefault QoM = this->Charge / this->Mass;
const vtkm::Float64 QoM = this->Charge / this->Mass;
const vtkm::Vec3f mom_minus = this->Momentum + (0.5 * this->Charge * eField * length);
// Get reciprocal of Gamma
vtkm::Vec3f gamma_reci = this->Gamma(mom_minus, true);
vtkm::Vec3f gamma_reci = static_cast<vtkm::FloatDefault>(this->Gamma(mom_minus, true));
const vtkm::Vec3f t = 0.5 * QoM * length * bField * gamma_reci;
const vtkm::Vec3f s = 2.0f * t * (1.0 / (1.0 + vtkm::Magnitude(t)));
const vtkm::Vec3f mom_prime = mom_minus + vtkm::Cross(mom_minus, t);
@ -266,9 +266,9 @@ public:
vtkm::FloatDefault Time = 0;
private:
vtkm::FloatDefault Mass;
vtkm::FloatDefault Charge;
vtkm::FloatDefault Weighting;
vtkm::Float64 Mass;
vtkm::Float64 Charge;
vtkm::Float64 Weighting;
vtkm::Vec3f Momentum;
constexpr static vtkm::FloatDefault SPEED_OF_LIGHT =
static_cast<vtkm::FloatDefault>(2.99792458e8);
@ -283,11 +283,10 @@ public:
+ sizeof(vtkm::Id) // NumSteps
+ sizeof(vtkm::UInt8) // Status
+ sizeof(vtkm::FloatDefault) // Time
+ sizeof(vtkm::FloatDefault) //Mass
+ sizeof(vtkm::FloatDefault) //Charge
+ sizeof(vtkm::FloatDefault) //Weighting
+ sizeof(vtkm::Vec3f) //Momentum
+ sizeof(vtkm::FloatDefault); //Speed_of_light
+ sizeof(vtkm::Float64) //Mass
+ sizeof(vtkm::Float64) //Charge
+ sizeof(vtkm::Float64) //Weighting
+ sizeof(vtkm::Vec3f); //Momentum
return sz;
}

1
vtkm/cont/CellLocatorBoundingIntervalHierarchy.h
View File

@ -39,6 +39,7 @@ public:
vtkm::ListTransform<SupportedCellSets, vtkm::exec::CellLocatorBoundingIntervalHierarchy>;
using ExecObjType = vtkm::ListApply<CellLocatorExecList, vtkm::exec::CellLocatorMultiplexer>;
using LastCell = typename ExecObjType::LastCell;
VTKM_CONT
CellLocatorBoundingIntervalHierarchy(vtkm::IdComponent numPlanes = 4,

1
vtkm/cont/CellLocatorGeneral.h
View File

@ -54,6 +54,7 @@ public:
vtkm::cont::internal::ExecutionObjectType<vtkm::cont::CellLocatorTwoLevel>>;
using ExecObjType = vtkm::ListApply<ExecLocatorList, vtkm::exec::CellLocatorMultiplexer>;
using LastCell = typename ExecObjType::LastCell;
VTKM_CONT ExecObjType PrepareForExecution(vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token) const;

2
vtkm/cont/CellLocatorRectilinearGrid.h
View File

@ -36,6 +36,8 @@ public:
~CellLocatorRectilinearGrid() = default;
using LastCell = vtkm::exec::CellLocatorRectilinearGrid::LastCell;
VTKM_CONT vtkm::exec::CellLocatorRectilinearGrid PrepareForExecution(
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token) const;

1
vtkm/cont/CellLocatorTwoLevel.h
View File

@ -58,6 +58,7 @@ public:
vtkm::ListTransform<CellExecObjectList, vtkm::exec::CellLocatorTwoLevel>;
using ExecObjType = vtkm::ListApply<CellLocatorExecList, vtkm::exec::CellLocatorMultiplexer>;
using LastCell = typename ExecObjType::LastCell;
CellLocatorTwoLevel()
: DensityL1(32.0f)

2
vtkm/cont/CellLocatorUniformGrid.h
View File

@ -25,6 +25,8 @@ class VTKM_CONT_EXPORT CellLocatorUniformGrid
using Superclass = vtkm::cont::internal::CellLocatorBase<CellLocatorUniformGrid>;
public:
using LastCell = vtkm::exec::CellLocatorUniformGrid::LastCell;
VTKM_CONT vtkm::exec::CellLocatorUniformGrid PrepareForExecution(
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token) const;

83
vtkm/cont/RuntimeDeviceTracker.cxx
View File

@ -28,8 +28,42 @@ namespace detail
struct RuntimeDeviceTrackerInternals
{
RuntimeDeviceTrackerInternals() = default;
RuntimeDeviceTrackerInternals(const RuntimeDeviceTrackerInternals* v) { this->CopyFrom(v); }
RuntimeDeviceTrackerInternals& operator=(const RuntimeDeviceTrackerInternals* v)
{
this->CopyFrom(v);
return *this;
}
bool GetRuntimeAllowed(std::size_t deviceId) const { return this->RuntimeAllowed[deviceId]; }
void SetRuntimeAllowed(std::size_t deviceId, bool flag) { this->RuntimeAllowed[deviceId] = flag; }
bool GetThreadFriendlyMemAlloc() const { return this->ThreadFriendlyMemAlloc; }
void SetThreadFriendlyMemAlloc(bool flag) { this->ThreadFriendlyMemAlloc = flag; }
void ResetRuntimeAllowed()
{
std::fill_n(this->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, false);
}
void Reset() { this->ResetRuntimeAllowed(); }
private:
void CopyFrom(const RuntimeDeviceTrackerInternals* v)
{
std::copy(std::cbegin(v->RuntimeAllowed),
std::cend(v->RuntimeAllowed),
std::begin(this->RuntimeAllowed));
this->SetThreadFriendlyMemAlloc(v->GetThreadFriendlyMemAlloc());
}
bool RuntimeAllowed[VTKM_MAX_DEVICE_ADAPTER_ID];
bool ThreadFriendlyMemAlloc = false;
};
}
VTKM_CONT
@ -65,7 +99,7 @@ bool RuntimeDeviceTracker::CanRunOn(vtkm::cont::DeviceAdapterId deviceId) const
{ //If at least a single device is enabled, than any device is enabled
for (vtkm::Int8 i = 1; i < VTKM_MAX_DEVICE_ADAPTER_ID; ++i)
{
if (this->Internals->RuntimeAllowed[static_cast<std::size_t>(i)])
if (this->Internals->GetRuntimeAllowed(static_cast<std::size_t>(i)))
{
return true;
}
@ -75,18 +109,29 @@ bool RuntimeDeviceTracker::CanRunOn(vtkm::cont::DeviceAdapterId deviceId) const
else
{
this->CheckDevice(deviceId);
return this->Internals->RuntimeAllowed[deviceId.GetValue()];
return this->Internals->GetRuntimeAllowed(deviceId.GetValue());
}
}
VTKM_CONT
bool RuntimeDeviceTracker::GetThreadFriendlyMemAlloc() const
{
return this->Internals->GetThreadFriendlyMemAlloc();
}
VTKM_CONT
void RuntimeDeviceTracker::SetDeviceState(vtkm::cont::DeviceAdapterId deviceId, bool state)
{
this->CheckDevice(deviceId);
this->Internals->RuntimeAllowed[deviceId.GetValue()] = state;
this->Internals->SetRuntimeAllowed(deviceId.GetValue(), state);
}
VTKM_CONT
void RuntimeDeviceTracker::SetThreadFriendlyMemAlloc(bool state)
{
this->Internals->SetThreadFriendlyMemAlloc(state);
}
VTKM_CONT void RuntimeDeviceTracker::ResetDevice(vtkm::cont::DeviceAdapterId deviceId)
{
@ -106,7 +151,7 @@ VTKM_CONT void RuntimeDeviceTracker::ResetDevice(vtkm::cont::DeviceAdapterId dev
VTKM_CONT
void RuntimeDeviceTracker::Reset()
{
std::fill_n(this->Internals->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, false);
this->Internals->Reset();
// We use this instead of calling CheckDevice/SetDeviceState so that
// when we use logging we get better messages stating we are reseting
@ -118,7 +163,7 @@ void RuntimeDeviceTracker::Reset()
if (device.IsValueValid())
{
const bool state = runtimeDevice.Exists(device);
this->Internals->RuntimeAllowed[device.GetValue()] = state;
this->Internals->SetRuntimeAllowed(device.GetValue(), state);
}
}
this->LogEnabledDevices();
@ -128,7 +173,7 @@ VTKM_CONT void RuntimeDeviceTracker::DisableDevice(vtkm::cont::DeviceAdapterId d
{
if (deviceId == vtkm::cont::DeviceAdapterTagAny{})
{
std::fill_n(this->Internals->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, false);
this->Internals->ResetRuntimeAllowed();
}
else
{
@ -157,18 +202,15 @@ void RuntimeDeviceTracker::ForceDevice(DeviceAdapterId deviceId)
throw vtkm::cont::ErrorBadValue(message.str());
}
std::fill_n(this->Internals->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, false);
this->Internals->RuntimeAllowed[deviceId.GetValue()] = runtimeExists;
this->Internals->ResetRuntimeAllowed();
this->Internals->SetRuntimeAllowed(deviceId.GetValue(), runtimeExists);
this->LogEnabledDevices();
}
}
VTKM_CONT void RuntimeDeviceTracker::CopyStateFrom(const vtkm::cont::RuntimeDeviceTracker& tracker)
{
std::copy(std::cbegin(tracker.Internals->RuntimeAllowed),
std::cend(tracker.Internals->RuntimeAllowed),
std::begin(this->Internals->RuntimeAllowed));
*(this->Internals) = tracker.Internals;
}
VTKM_CONT
@ -208,11 +250,9 @@ VTKM_CONT
ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(vtkm::cont::DeviceAdapterId device,
RuntimeDeviceTrackerMode mode)
: RuntimeDeviceTracker(GetRuntimeDeviceTracker().Internals, false)
, SavedState(new detail::RuntimeDeviceTrackerInternals())
, SavedState(new detail::RuntimeDeviceTrackerInternals(this->Internals))
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Entering scoped runtime region");
std::copy_n(
this->Internals->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, this->SavedState->RuntimeAllowed);
if (mode == RuntimeDeviceTrackerMode::Force)
{
@ -234,11 +274,10 @@ ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(
RuntimeDeviceTrackerMode mode,
const vtkm::cont::RuntimeDeviceTracker& tracker)
: RuntimeDeviceTracker(tracker.Internals, false)
, SavedState(new detail::RuntimeDeviceTrackerInternals())
, SavedState(new detail::RuntimeDeviceTrackerInternals(this->Internals))
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Entering scoped runtime region");
std::copy_n(
this->Internals->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, this->SavedState->RuntimeAllowed);
if (mode == RuntimeDeviceTrackerMode::Force)
{
this->ForceDevice(device);
@ -257,19 +296,17 @@ VTKM_CONT
ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(
const vtkm::cont::RuntimeDeviceTracker& tracker)
: RuntimeDeviceTracker(tracker.Internals, false)
, SavedState(new detail::RuntimeDeviceTrackerInternals())
, SavedState(new detail::RuntimeDeviceTrackerInternals(this->Internals))
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Entering scoped runtime region");
std::copy_n(
this->Internals->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, this->SavedState->RuntimeAllowed);
}
VTKM_CONT
ScopedRuntimeDeviceTracker::~ScopedRuntimeDeviceTracker()
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Leaving scoped runtime region");
std::copy_n(
this->SavedState->RuntimeAllowed, VTKM_MAX_DEVICE_ADAPTER_ID, this->Internals->RuntimeAllowed);
*(this->Internals) = this->SavedState.get();
this->LogEnabledDevices();
}

8
vtkm/cont/RuntimeDeviceTracker.h
View File

@ -110,7 +110,13 @@ public:
///
VTKM_CONT void ForceDevice(DeviceAdapterId deviceId);
/// \brief Copyies the state from the given device.
/// \brief Get/Set use of thread-friendly memory allocation for a device.
///
///
VTKM_CONT bool GetThreadFriendlyMemAlloc() const;
VTKM_CONT void SetThreadFriendlyMemAlloc(bool state);
/// \brief Copies the state from the given device.
///
/// This is a convenient way to allow the `RuntimeDeviceTracker` on one thread
/// copy the behavior from another thread.

37
vtkm/cont/cuda/internal/CudaAllocator.cu
View File

@ -12,6 +12,7 @@
#include <mutex>
#include <vtkm/cont/Logging.h>
#include <vtkm/cont/RuntimeDeviceInformation.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
#include <vtkm/cont/cuda/ErrorCuda.h>
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
#include <vtkm/cont/cuda/internal/DeviceAdapterTagCuda.h>
@ -152,7 +153,15 @@ void* CudaAllocator::Allocate(std::size_t numBytes)
}
void* ptr = nullptr;
if (ManagedMemoryEnabled)
#if CUDART_VERSION >= 11030
const auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
if (tracker.GetThreadFriendlyMemAlloc())
{
VTKM_CUDA_CALL(cudaMallocAsync(&ptr, numBytes, cudaStreamPerThread));
}
else
#endif
if (ManagedMemoryEnabled)
{
VTKM_CUDA_CALL(cudaMallocManaged(&ptr, numBytes));
}
@ -174,7 +183,18 @@ void* CudaAllocator::Allocate(std::size_t numBytes)
void* CudaAllocator::AllocateUnManaged(std::size_t numBytes)
{
void* ptr = nullptr;
VTKM_CUDA_CALL(cudaMalloc(&ptr, numBytes));
#if CUDART_VERSION >= 11030
const auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
if (tracker.GetThreadFriendlyMemAlloc())
{
VTKM_CUDA_CALL(cudaMallocAsync(&ptr, numBytes, cudaStreamPerThread));
}
else
#endif
{
VTKM_CUDA_CALL(cudaMalloc(&ptr, numBytes));
}
{
VTKM_LOG_F(vtkm::cont::LogLevel::MemExec,
"Allocated CUDA array of %s at %p.",
@ -195,7 +215,18 @@ void CudaAllocator::Free(void* ptr)
}
VTKM_LOG_F(vtkm::cont::LogLevel::MemExec, "Freeing CUDA allocation at %p.", ptr);
VTKM_CUDA_CALL(cudaFree(ptr));
#if CUDART_VERSION >= 11030
const auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
if (tracker.GetThreadFriendlyMemAlloc())
{
VTKM_CUDA_CALL(cudaFreeAsync(ptr, cudaStreamPerThread));
}
else
#endif
{
VTKM_CUDA_CALL(cudaFree(ptr));
}
}
void CudaAllocator::FreeDeferred(void* ptr, std::size_t numBytes)

3
vtkm/cont/cuda/internal/CudaAllocator.h
View File

@ -40,6 +40,9 @@ struct VTKM_CONT_EXPORT CudaAllocator
/// VTK-m will ignore the request and continue to use unmanaged memory (aka cudaMalloc).
static VTKM_CONT void ForceManagedMemoryOn();
static VTKM_CONT void ForceSyncMemoryAllocator();
static VTKM_CONT void ForceAsyncMemoryAllocator();
/// Returns true if the pointer is accessible from a CUDA device.
static VTKM_CONT bool IsDevicePointer(const void* ptr);

1
vtkm/cont/kokkos/internal/DeviceAdapterAlgorithmKokkos.h
View File

@ -27,7 +27,6 @@
VTKM_THIRDPARTY_PRE_INCLUDE
#include <Kokkos_Core.hpp>
#include <Kokkos_DualView.hpp>
#include <Kokkos_Parallel_Reduce.hpp>
#include <Kokkos_Sort.hpp>
VTKM_THIRDPARTY_POST_INCLUDE

87
vtkm/cont/testing/UnitTestCellLocatorGeneral.cxx
View File

@ -12,10 +12,9 @@
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/DataSetBuilderRectilinear.h>
#include <vtkm/cont/DataSetBuilderUniform.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/exec/CellInterpolate.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterPermutation.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/WorkletMapTopology.h>
@ -140,10 +139,13 @@ void GenerateRandomInput(const vtkm::cont::DataSet& ds,
pcoordsPortal.Set(i, pc);
}
vtkm::worklet::DispatcherMapTopology<ParametricToWorldCoordinates> dispatcher(
ParametricToWorldCoordinates::MakeScatter(cellIds));
dispatcher.Invoke(
ds.GetCellSet(), ds.GetCoordinateSystem().GetDataAsMultiplexer(), pcoords, wcoords);
vtkm::cont::Invoker invoker;
invoker(ParametricToWorldCoordinates{},
ParametricToWorldCoordinates::MakeScatter(cellIds),
ds.GetCellSet(),
ds.GetCoordinateSystem().GetDataAsMultiplexer(),
pcoords,
wcoords);
}
//-----------------------------------------------------------------------------
@ -170,6 +172,56 @@ public:
}
};
class FindCellWorkletWithLastCell : public vtkm::worklet::WorkletMapField
{
public:
using ControlSignature = void(FieldIn points,
ExecObject locator,
FieldOut cellIds,
FieldOut pcoords,
FieldInOut lastCell);
using ExecutionSignature = void(_1, _2, _3, _4, _5);
template <typename LocatorType>
VTKM_EXEC void operator()(const vtkm::Vec3f& point,
const LocatorType& locator,
vtkm::Id& cellId,
vtkm::Vec3f& pcoords,
typename LocatorType::LastCell& lastCell) const
{
vtkm::ErrorCode status = locator.FindCell(point, cellId, pcoords, lastCell);
if (status != vtkm::ErrorCode::Success)
this->RaiseError(vtkm::ErrorString(status));
}
};
void TestLastCell(vtkm::cont::CellLocatorGeneral& locator,
vtkm::Id numPoints,
vtkm::cont::ArrayHandle<vtkm::cont::CellLocatorGeneral::LastCell>& lastCell,
const vtkm::cont::ArrayHandle<PointType>& points,
const vtkm::cont::ArrayHandle<vtkm::Id>& expCellIds,
const vtkm::cont::ArrayHandle<PointType>& expPCoords)
{
vtkm::cont::ArrayHandle<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<PointType> pcoords;
vtkm::cont::Invoker invoker;
invoker(FindCellWorkletWithLastCell{}, points, locator, cellIds, pcoords, lastCell);
auto cellIdPortal = cellIds.ReadPortal();
auto expCellIdsPortal = expCellIds.ReadPortal();
auto pcoordsPortal = pcoords.ReadPortal();
auto expPCoordsPortal = expPCoords.ReadPortal();
for (vtkm::Id i = 0; i < numPoints; ++i)
{
VTKM_TEST_ASSERT(cellIdPortal.Get(i) == expCellIdsPortal.Get(i), "Incorrect cell ids");
VTKM_TEST_ASSERT(test_equal(pcoordsPortal.Get(i), expPCoordsPortal.Get(i), 1e-3),
"Incorrect parameteric coordinates");
}
}
void TestWithDataSet(vtkm::cont::CellLocatorGeneral& locator, const vtkm::cont::DataSet& dataset)
{
locator.SetCellSet(dataset.GetCellSet());
@ -184,8 +236,8 @@ void TestWithDataSet(vtkm::cont::CellLocatorGeneral& locator, const vtkm::cont::
vtkm::cont::ArrayHandle<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<PointType> pcoords;
vtkm::worklet::DispatcherMapField<FindCellWorklet> dispatcher;
dispatcher.Invoke(points, locator, cellIds, pcoords);
vtkm::cont::Invoker invoker;
invoker(FindCellWorklet{}, points, locator, cellIds, pcoords);
auto cellIdPortal = cellIds.ReadPortal();
auto expCellIdsPortal = expCellIds.ReadPortal();
@ -197,6 +249,25 @@ void TestWithDataSet(vtkm::cont::CellLocatorGeneral& locator, const vtkm::cont::
VTKM_TEST_ASSERT(test_equal(pcoordsPortal.Get(i), expPCoordsPortal.Get(i), 1e-3),
"Incorrect parameteric coordinates");
}
//Test locator using lastCell
//Test it with initialized.
vtkm::cont::ArrayHandle<vtkm::cont::CellLocatorGeneral::LastCell> lastCell;
lastCell.AllocateAndFill(64, vtkm::cont::CellLocatorGeneral::LastCell{});
TestLastCell(locator, 64, lastCell, points, expCellIds, pcoords);
//Call it again using the lastCell just computed to validate.
TestLastCell(locator, 64, lastCell, points, expCellIds, pcoords);
//Test it with uninitialized array.
vtkm::cont::ArrayHandle<vtkm::cont::CellLocatorGeneral::LastCell> lastCell2;
lastCell2.Allocate(64);
TestLastCell(locator, 64, lastCell2, points, expCellIds, pcoords);
//Call it again using the lastCell just computed to validate.
TestLastCell(locator, 64, lastCell2, points, expCellIds, pcoords);
}
void TestCellLocatorGeneral()

86
vtkm/cont/testing/UnitTestCellLocatorTwoLevel.cxx
View File

@ -11,14 +11,13 @@
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/CellLocatorTwoLevel.h>
#include <vtkm/cont/DataSetBuilderUniform.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/exec/ParametricCoordinates.h>
#include <vtkm/filter/geometry_refinement/worklet/Tetrahedralize.h>
#include <vtkm/filter/geometry_refinement/worklet/Triangulate.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/DispatcherMapTopology.h>
#include <vtkm/worklet/ScatterPermutation.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/WorkletMapTopology.h>
@ -150,10 +149,13 @@ void GenerateRandomInput(const vtkm::cont::DataSet& ds,
pcoords.WritePortal().Set(i, pc);
}
vtkm::worklet::DispatcherMapTopology<ParametricToWorldCoordinates> dispatcher(
ParametricToWorldCoordinates::MakeScatter(cellIds));
dispatcher.Invoke(
ds.GetCellSet(), ds.GetCoordinateSystem().GetDataAsMultiplexer(), pcoords, wcoords);
vtkm::cont::Invoker invoker;
invoker(ParametricToWorldCoordinates{},
ParametricToWorldCoordinates::MakeScatter(cellIds),
ds.GetCellSet(),
ds.GetCoordinateSystem().GetDataAsMultiplexer(),
pcoords,
wcoords);
}
class FindCellWorklet : public vtkm::worklet::WorkletMapField
@ -179,6 +181,56 @@ public:
}
};
class FindCellWorkletWithLastCell : public vtkm::worklet::WorkletMapField
{
public:
using ControlSignature = void(FieldIn points,
ExecObject locator,
FieldOut cellIds,
FieldOut pcoords,
FieldInOut lastCell);
using ExecutionSignature = void(_1, _2, _3, _4, _5);
template <typename LocatorType>
VTKM_EXEC void operator()(const vtkm::Vec3f& point,
const LocatorType& locator,
vtkm::Id& cellId,
vtkm::Vec3f& pcoords,
typename LocatorType::LastCell& lastCell) const
{
vtkm::ErrorCode status = locator.FindCell(point, cellId, pcoords, lastCell);
if (status != vtkm::ErrorCode::Success)
this->RaiseError(vtkm::ErrorString(status));
}
};
void TestLastCell(vtkm::cont::CellLocatorTwoLevel& locator,
vtkm::Id numPoints,
vtkm::cont::ArrayHandle<vtkm::cont::CellLocatorTwoLevel::LastCell>& lastCell,
const vtkm::cont::ArrayHandle<PointType>& points,
const vtkm::cont::ArrayHandle<vtkm::Id>& expCellIds,
const vtkm::cont::ArrayHandle<PointType>& expPCoords)
{
vtkm::cont::ArrayHandle<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<PointType> pcoords;
vtkm::cont::Invoker invoker;
invoker(FindCellWorkletWithLastCell{}, points, locator, cellIds, pcoords, lastCell);
auto cellIdPortal = cellIds.ReadPortal();
auto expCellIdsPortal = expCellIds.ReadPortal();
auto pcoordsPortal = pcoords.ReadPortal();
auto expPCoordsPortal = expPCoords.ReadPortal();
for (vtkm::Id i = 0; i < numPoints; ++i)
{
VTKM_TEST_ASSERT(cellIdPortal.Get(i) == expCellIdsPortal.Get(i), "Incorrect cell ids");
VTKM_TEST_ASSERT(test_equal(pcoordsPortal.Get(i), expPCoordsPortal.Get(i), 1e-3),
"Incorrect parameteric coordinates");
}
}
template <vtkm::IdComponent DIMENSIONS>
void TestCellLocator(const vtkm::Vec<vtkm::Id, DIMENSIONS>& dim, vtkm::Id numberOfPoints)
{
@ -202,8 +254,8 @@ void TestCellLocator(const vtkm::Vec<vtkm::Id, DIMENSIONS>& dim, vtkm::Id number
vtkm::cont::ArrayHandle<vtkm::Id> cellIds;
vtkm::cont::ArrayHandle<PointType> pcoords;
vtkm::worklet::DispatcherMapField<FindCellWorklet> dispatcher;
dispatcher.Invoke(points, locator, cellIds, pcoords);
vtkm::cont::Invoker invoker;
invoker(FindCellWorklet{}, points, locator, cellIds, pcoords);
auto cellIdsPortal = cellIds.ReadPortal();
auto expCellIdsPortal = expCellIds.ReadPortal();
@ -215,6 +267,24 @@ void TestCellLocator(const vtkm::Vec<vtkm::Id, DIMENSIONS>& dim, vtkm::Id number
VTKM_TEST_ASSERT(test_equal(pcoordsPortal.Get(i), expPCoordsPortal.Get(i), 1e-3),
"Incorrect parameteric coordinates");
}
//Test locator using lastCell
//Test it with initialized.
vtkm::cont::ArrayHandle<vtkm::cont::CellLocatorTwoLevel::LastCell> lastCell;
lastCell.AllocateAndFill(numberOfPoints, vtkm::cont::CellLocatorTwoLevel::LastCell{});
TestLastCell(locator, numberOfPoints, lastCell, points, expCellIds, pcoords);
//Call it again using the lastCell just computed to validate.
TestLastCell(locator, numberOfPoints, lastCell, points, expCellIds, pcoords);
//Test it with uninitialized array.
vtkm::cont::ArrayHandle<vtkm::cont::CellLocatorTwoLevel::LastCell> lastCell2;
lastCell2.Allocate(numberOfPoints);
TestLastCell(locator, numberOfPoints, lastCell2, points, expCellIds, pcoords);
//Call it again using the lastCell2 just computed to validate.
TestLastCell(locator, numberOfPoints, lastCell2, points, expCellIds, pcoords);
}
void TestingCellLocatorTwoLevel()

94
vtkm/exec/CellLocatorBoundingIntervalHierarchy.h
View File

@ -83,11 +83,64 @@ public:
{
}
struct LastCell
{
vtkm::Id CellId = -1;
vtkm::Id NodeIdx = -1;
};
VTKM_EXEC
vtkm::ErrorCode FindCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric) const
{
LastCell lastCell;
return this->FindCellImpl(point, cellId, parametric, lastCell);
}
VTKM_EXEC
vtkm::ErrorCode FindCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric,
LastCell& lastCell) const
{
cellId = -1;
//Check the last cell.
if ((lastCell.CellId >= 0) && (lastCell.CellId < this->CellSet.GetNumberOfElements()))
{
if (this->PointInCell(point, lastCell.CellId, parametric) == vtkm::ErrorCode::Success)
{
cellId = lastCell.CellId;
return vtkm::ErrorCode::Success;
}
}
//Check the last leaf node.
if ((lastCell.NodeIdx >= 0) && (lastCell.NodeIdx < this->Nodes.GetNumberOfValues()))
{
const auto& node = this->Nodes.Get(lastCell.NodeIdx);
if (node.ChildIndex < 0)
{
VTKM_RETURN_ON_ERROR(this->FindInLeaf(point, parametric, node, cellId));
if (cellId != -1)
{
lastCell.CellId = cellId;
return vtkm::ErrorCode::Success;
}
}
}
//No fastpath. Do a full search.
return this->FindCellImpl(point, cellId, parametric, lastCell);
}
VTKM_EXEC
vtkm::ErrorCode FindCellImpl(const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric,
LastCell& lastCell) const
{
cellId = -1;
vtkm::Id nodeIndex = 0;
@ -98,7 +151,8 @@ public:
switch (state)
{
case FindCellState::EnterNode:
VTKM_RETURN_ON_ERROR(this->EnterNode(state, point, cellId, nodeIndex, parametric));
VTKM_RETURN_ON_ERROR(
this->EnterNode(state, point, cellId, nodeIndex, parametric, lastCell));
break;
case FindCellState::AscendFromNode:
this->AscendFromNode(state, nodeIndex);
@ -141,7 +195,8 @@ private:
const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Id nodeIndex,
vtkm::Vec3f& parametric) const
vtkm::Vec3f& parametric,
LastCell& lastCell) const
{
VTKM_ASSERT(state == FindCellState::EnterNode);
@ -152,6 +207,11 @@ private:
// In a leaf node. Look for a containing cell.
VTKM_RETURN_ON_ERROR(this->FindInLeaf(point, parametric, node, cellId));
state = FindCellState::AscendFromNode;
if (cellId != -1)
{
lastCell.CellId = cellId;
lastCell.NodeIdx = nodeIndex;
}
}
else
{
@ -230,26 +290,40 @@ private:
const vtkm::exec::CellLocatorBoundingIntervalHierarchyNode& node,
vtkm::Id& containingCellId) const
{
using IndicesType = typename CellSetPortal::IndicesType;
for (vtkm::Id i = node.Leaf.Start; i < node.Leaf.Start + node.Leaf.Size; ++i)
{
vtkm::Id cellId = this->CellIds.Get(i);
IndicesType cellPointIndices = this->CellSet.GetIndices(cellId);
vtkm::VecFromPortalPermute<IndicesType, CoordsPortal> cellPoints(&cellPointIndices,
this->Coords);
bool found;
VTKM_RETURN_ON_ERROR(this->IsPointInCell(
point, parametric, this->CellSet.GetCellShape(cellId), cellPoints, found));
if (found)
if (this->PointInCell(point, cellId, parametric) == vtkm::ErrorCode::Success)
{
containingCellId = cellId;
return vtkm::ErrorCode::Success;
}
}
containingCellId = -1;
return vtkm::ErrorCode::Success;
}
// template <typename CoordsType, typename CellShapeTag>
VTKM_EXEC vtkm::ErrorCode PointInCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric) const
{
using IndicesType = typename CellSetPortal::IndicesType;
IndicesType cellPointIndices = this->CellSet.GetIndices(cellId);
vtkm::VecFromPortalPermute<IndicesType, CoordsPortal> cellPoints(&cellPointIndices,
this->Coords);
auto cellShape = this->CellSet.GetCellShape(cellId);
bool isInside;
VTKM_RETURN_ON_ERROR(IsPointInCell(point, parametric, cellShape, cellPoints, isInside));
if (isInside && vtkm::exec::CellInside(parametric, cellShape))
return vtkm::ErrorCode::Success;
return vtkm::ErrorCode::CellNotFound;
}
template <typename CoordsType, typename CellShapeTag>
VTKM_EXEC static vtkm::ErrorCode IsPointInCell(const vtkm::Vec3f& point,
vtkm::Vec3f& parametric,

26
vtkm/exec/CellLocatorMultiplexer.h
View File

@ -33,6 +33,21 @@ struct FindCellFunctor
{
return locator.FindCell(point, cellId, parametric);
}
template <typename Locator, typename LastCell>
VTKM_EXEC vtkm::ErrorCode operator()(Locator&& locator,
const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric,
LastCell& lastCell) const
{
using ConcreteLastCell = typename std::decay_t<Locator>::LastCell;
if (!lastCell.template IsType<ConcreteLastCell>())
{
lastCell = ConcreteLastCell{};
}
return locator.FindCell(point, cellId, parametric, lastCell.template Get<ConcreteLastCell>());
}
};
} // namespace detail
@ -45,6 +60,8 @@ class VTKM_ALWAYS_EXPORT CellLocatorMultiplexer
public:
CellLocatorMultiplexer() = default;
using LastCell = vtkm::exec::internal::Variant<typename LocatorTypes::LastCell...>;
template <typename Locator>
VTKM_CONT CellLocatorMultiplexer(const Locator& locator)
: Locators(locator)
@ -58,6 +75,15 @@ public:
return this->Locators.CastAndCall(detail::FindCellFunctor{}, point, cellId, parametric);
}
VTKM_EXEC vtkm::ErrorCode FindCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric,
LastCell& lastCell) const
{
return this->Locators.CastAndCall(
detail::FindCellFunctor{}, point, cellId, parametric, lastCell);
}
VTKM_DEPRECATED(1.6, "Locators are no longer pointers. Use . operator.")
VTKM_EXEC CellLocatorMultiplexer* operator->() { return this; }
VTKM_DEPRECATED(1.6, "Locators are no longer pointers. Use . operator.")

13
vtkm/exec/CellLocatorRectilinearGrid.h
View File

@ -42,6 +42,10 @@ private:
VTKM_CONT static vtkm::Id3 ToId3(vtkm::Id&& src) { return vtkm::Id3(src, 1, 1); }
public:
struct LastCell
{
};
template <vtkm::IdComponent dimensions>
VTKM_CONT CellLocatorRectilinearGrid(const vtkm::Id planeSize,
const vtkm::Id rowSize,
@ -87,6 +91,15 @@ public:
return inside;
}
VTKM_EXEC
vtkm::ErrorCode FindCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric,
LastCell& vtkmNotUsed(lastCell)) const
{
return this->FindCell(point, cellId, parametric);
}
VTKM_EXEC
vtkm::ErrorCode FindCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,

126
vtkm/exec/CellLocatorTwoLevel.h
View File

@ -153,12 +153,110 @@ public:
{
}
struct LastCell
{
vtkm::Id CellId = -1;
vtkm::Id LeafIdx = -1;
};
VTKM_EXEC
vtkm::ErrorCode FindCell(const FloatVec3& point, vtkm::Id& cellId, FloatVec3& parametric) const
{
LastCell lastCell;
return this->FindCellImpl(point, cellId, parametric, lastCell);
}
VTKM_EXEC
vtkm::ErrorCode FindCell(const FloatVec3& point,
vtkm::Id& cellId,
FloatVec3& parametric,
LastCell& lastCell) const
{
vtkm::Vec3f pc;
//See if point is inside the last cell.
if ((lastCell.CellId >= 0) && (lastCell.CellId < this->CellSet.GetNumberOfElements()) &&
this->PointInCell(point, lastCell.CellId, pc) == vtkm::ErrorCode::Success)
{
parametric = pc;
cellId = lastCell.CellId;
return vtkm::ErrorCode::Success;
}
//See if it's in the last leaf.
if ((lastCell.LeafIdx >= 0) && (lastCell.LeafIdx < this->CellCount.GetNumberOfValues()) &&
this->PointInLeaf(point, lastCell.LeafIdx, cellId, pc) == vtkm::ErrorCode::Success)
{
parametric = pc;
lastCell.CellId = cellId;
return vtkm::ErrorCode::Success;
}
//Call the full point search.
return this->FindCellImpl(point, cellId, parametric, lastCell);
}
VTKM_DEPRECATED(1.6, "Locators are no longer pointers. Use . operator.")
VTKM_EXEC CellLocatorTwoLevel* operator->() { return this; }
VTKM_DEPRECATED(1.6, "Locators are no longer pointers. Use . operator.")
VTKM_EXEC const CellLocatorTwoLevel* operator->() const { return this; }
private:
VTKM_EXEC
vtkm::ErrorCode PointInCell(const vtkm::Vec3f& point,
const vtkm::Id& cid,
vtkm::Vec3f& parametric) const
{
auto indices = this->CellSet.GetIndices(cid);
auto pts = vtkm::make_VecFromPortalPermute(&indices, this->Coords);
vtkm::Vec3f pc;
bool inside;
auto status = PointInsideCell(point, this->CellSet.GetCellShape(cid), pts, pc, inside);
if (status == vtkm::ErrorCode::Success && inside)
{
parametric = pc;
return vtkm::ErrorCode::Success;
}
return vtkm::ErrorCode::CellNotFound;
}
VTKM_EXEC
vtkm::ErrorCode PointInLeaf(const FloatVec3& point,
const vtkm::Id& leafIdx,
vtkm::Id& cellId,
FloatVec3& parametric) const
{
vtkm::Id start = this->CellStartIndex.Get(leafIdx);
vtkm::Id end = start + this->CellCount.Get(leafIdx);
for (vtkm::Id i = start; i < end; ++i)
{
vtkm::Vec3f pc;
vtkm::Id cid = this->CellIds.Get(i);
if (this->PointInCell(point, cid, pc) == vtkm::ErrorCode::Success)
{
cellId = cid;
parametric = pc;
return vtkm::ErrorCode::Success;
}
}
return vtkm::ErrorCode::CellNotFound;
}
VTKM_EXEC
vtkm::ErrorCode FindCellImpl(const FloatVec3& point,
vtkm::Id& cellId,
FloatVec3& parametric,
LastCell& lastCell) const
{
using namespace vtkm::internal::cl_uniform_bins;
cellId = -1;
lastCell.CellId = -1;
lastCell.LeafIdx = -1;
DimVec3 binId3 = static_cast<DimVec3>((point - this->TopLevel.Origin) / this->TopLevel.BinSize);
if (binId3[0] >= 0 && binId3[0] < this->TopLevel.Dimensions[0] && binId3[1] >= 0 &&
@ -180,37 +278,19 @@ public:
leafId3 = vtkm::Max(DimVec3(0), vtkm::Min(ldim - DimVec3(1), leafId3));
vtkm::Id leafStart = this->LeafStartIndex.Get(binId);
vtkm::Id leafId = leafStart + ComputeFlatIndex(leafId3, leafGrid.Dimensions);
vtkm::Id leafIdx = leafStart + ComputeFlatIndex(leafId3, leafGrid.Dimensions);
vtkm::Id start = this->CellStartIndex.Get(leafId);
vtkm::Id end = start + this->CellCount.Get(leafId);
for (vtkm::Id i = start; i < end; ++i)
if (this->PointInLeaf(point, leafIdx, cellId, parametric) == vtkm::ErrorCode::Success)
{
vtkm::Id cid = this->CellIds.Get(i);
auto indices = this->CellSet.GetIndices(cid);
auto pts = vtkm::make_VecFromPortalPermute(&indices, this->Coords);
FloatVec3 pc;
bool inside;
VTKM_RETURN_ON_ERROR(
PointInsideCell(point, this->CellSet.GetCellShape(cid), pts, pc, inside));
if (inside)
{
cellId = cid;
parametric = pc;
return vtkm::ErrorCode::Success;
}
lastCell.CellId = cellId;
lastCell.LeafIdx = leafIdx;
return vtkm::ErrorCode::Success;
}
}
return vtkm::ErrorCode::CellNotFound;
}
VTKM_DEPRECATED(1.6, "Locators are no longer pointers. Use . operator.")
VTKM_EXEC CellLocatorTwoLevel* operator->() { return this; }
VTKM_DEPRECATED(1.6, "Locators are no longer pointers. Use . operator.")
VTKM_EXEC const CellLocatorTwoLevel* operator->() const { return this; }
private:
vtkm::internal::cl_uniform_bins::Grid TopLevel;
ReadPortal<DimVec3> LeafDimensions;

13
vtkm/exec/CellLocatorUniformGrid.h
View File

@ -55,6 +55,19 @@ public:
return inside;
}
struct LastCell
{
};
VTKM_EXEC
vtkm::ErrorCode FindCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,
vtkm::Vec3f& parametric,
LastCell& vtkmNotUsed(lastCell)) const
{
return this->FindCell(point, cellId, parametric);
}
VTKM_EXEC
vtkm::ErrorCode FindCell(const vtkm::Vec3f& point,
vtkm::Id& cellId,

2
vtkm/exec/internal/testing/UnitTestVariant.cxx
View File

@ -365,6 +365,8 @@ void TestGet()
VariantType variant = expectedValue;
VTKM_TEST_ASSERT(variant.GetIndex() == 2);
VTKM_TEST_ASSERT(variant.IsType<vtkm::Id>());
VTKM_TEST_ASSERT(!variant.IsType<vtkm::Float32>());
VTKM_TEST_ASSERT(variant.Get<2>() == expectedValue);

5
vtkm/filter/CMakeLists.txt
View File

@ -45,14 +45,19 @@ set(deprecated_headers
MIRFilter.h
NDEntropy.h
NDHistogram.h
ParticleAdvection.h
ParticleDensityCloudInCell.h
ParticleDensityNearestGridPoint.h
Pathline.h
PathParticle.h
PointAverage.h
PointElevation.h
PointTransform.h
Probe.h
Slice.h
SplitSharpEdges.h
Streamline.h
StreamSurface.h
SurfaceNormals.h
Tetrahedralize.h
Threshold.h

19
vtkm/filter/NewFilter.cxx
View File

@ -20,12 +20,17 @@ namespace vtkm
{
namespace filter
{
namespace
{
void RunFilter(NewFilter* self,
vtkm::filter::DataSetQueue& input,
vtkm::filter::DataSetQueue& output)
{
auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
bool prevVal = tracker.GetThreadFriendlyMemAlloc();
tracker.SetThreadFriendlyMemAlloc(true);
std::pair<vtkm::Id, vtkm::cont::DataSet> task;
while (input.GetTask(task))
{
@ -34,8 +39,10 @@ void RunFilter(NewFilter* self,
}
vtkm::cont::Algorithm::Synchronize();
tracker.SetThreadFriendlyMemAlloc(prevVal);
}
}
} // anonymous namespace
NewFilter::~NewFilter() = default;
@ -115,21 +122,17 @@ vtkm::Id NewFilter::DetermineNumberOfThreads(const vtkm::cont::PartitionedDataSe
{
vtkm::Id numDS = input.GetNumberOfPartitions();
//Aribitrary constants.
const vtkm::Id threadsPerGPU = 8;
const vtkm::Id threadsPerCPU = 4;
vtkm::Id availThreads = 1;
auto& tracker = vtkm::cont::GetRuntimeDeviceTracker();
if (tracker.CanRunOn(vtkm::cont::DeviceAdapterTagCuda{}))
availThreads = threadsPerGPU;
availThreads = this->NumThreadsPerGPU;
else if (tracker.CanRunOn(vtkm::cont::DeviceAdapterTagKokkos{}))
{
//Kokkos doesn't support threading on the CPU.
#ifdef VTKM_KOKKOS_CUDA
availThreads = threadsPerGPU;
availThreads = this->NumThreadsPerGPU;
#else
availThreads = 1;
#endif
@ -137,7 +140,7 @@ vtkm::Id NewFilter::DetermineNumberOfThreads(const vtkm::cont::PartitionedDataSe
else if (tracker.CanRunOn(vtkm::cont::DeviceAdapterTagSerial{}))
availThreads = 1;
else
availThreads = threadsPerCPU;
availThreads = this->NumThreadsPerCPU;
vtkm::Id numThreads = std::min<vtkm::Id>(numDS, availThreads);
return numThreads;

13
vtkm/filter/NewFilter.h
View File

@ -17,6 +17,7 @@
#include <vtkm/cont/PartitionedDataSet.h>
#include <vtkm/filter/FieldSelection.h>
#include <vtkm/filter/TaskQueue.h>
#include <vtkm/filter/vtkm_filter_core_export.h>
namespace vtkm
@ -223,6 +224,16 @@ public:
VTKM_CONT
virtual bool CanThread() const;
VTKM_CONT
void SetThreadsPerCPU(vtkm::Id numThreads) { this->NumThreadsPerCPU = numThreads; }
VTKM_CONT
void SetThreadsPerGPU(vtkm::Id numThreads) { this->NumThreadsPerGPU = numThreads; }
VTKM_CONT
vtkm::Id GetThreadsPerCPU() const { return this->NumThreadsPerCPU; }
VTKM_CONT
vtkm::Id GetThreadsPerGPU() const { return this->NumThreadsPerGPU; }
VTKM_CONT
bool GetRunMultiThreadedFilter() const
{
@ -441,6 +452,8 @@ private:
vtkm::filter::FieldSelection FieldsToPass = vtkm::filter::FieldSelection::Mode::All;
bool RunFilterWithMultipleThreads = false;
vtkm::Id NumThreadsPerCPU = 4;
vtkm::Id NumThreadsPerGPU = 8;
};
}
} // namespace vtkm::filter

34
vtkm/filter/ParticleAdvection.h Normal file
View File

@ -0,0 +1,34 @@
//============================================================================
// 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_ParticleAdvection_h
#define vtk_m_filter_ParticleAdvection_h
#include <vtkm/Deprecated.h>
#include <vtkm/filter/flow/ParticleAdvection.h>
namespace vtkm
{
namespace filter
{
VTKM_DEPRECATED(
1.8,
"Use vtkm/filter/flow/ParticleAdvection.h instead of vtkm/filter/ParticleAdvection.h")
inline void ParticleAdvection_deprecated() {}
inline void ParticleAdvection_deprecated_warning()
{
ParticleAdvection_deprecated();
}
}
}
#endif //vtk_m_filter_ParticleAdvection_h

32
vtkm/filter/PathParticle.h Normal file
View File

@ -0,0 +1,32 @@
//============================================================================
// 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_PathParticle_h
#define vtk_m_filter_PathParticle_h
#include <vtkm/Deprecated.h>
#include <vtkm/filter/flow/PathParticle.h>
namespace vtkm
{
namespace filter
{
VTKM_DEPRECATED(1.8, "Use vtkm/filter/flow/PathParticle.h instead of vtkm/filter/PathParticle.h")
inline void PathParticle_deprecated() {}
inline void PathParticle_deprecated_warning()
{
PathParticle_deprecated();
}
}
}
#endif //vtk_m_filter_PathParticle_h

32
vtkm/filter/Pathline.h Normal file
View File

@ -0,0 +1,32 @@
//============================================================================
// 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_Pathline_h
#define vtk_m_filter_Pathline_h
#include <vtkm/Deprecated.h>
#include <vtkm/filter/flow/Pathline.h>
namespace vtkm
{
namespace filter
{
VTKM_DEPRECATED(1.8, "Use vtkm/filter/flow/Pathline.h instead of vtkm/filter/Pathline.h")
inline void Pathline_deprecated() {}
inline void Pathline_deprecated_warning()
{
Pathline_deprecated();
}
}
}
#endif //vtk_m_filter_Pathline_h

32
vtkm/filter/StreamSurface.h Normal file
View File

@ -0,0 +1,32 @@
//============================================================================
// 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_StreamSurface_h
#define vtk_m_filter_StreamSurface_h
#include <vtkm/Deprecated.h>
#include <vtkm/filter/flow/StreamSurface.h>
namespace vtkm
{
namespace filter
{
VTKM_DEPRECATED(1.8, "Use vtkm/filter/flow/StreamSurface.h instead of vtkm/filter/StreamSurface.h")
inline void StreamSurface_deprecated() {}
inline void StreamSurface_deprecated_warning()
{
StreamSurface_deprecated();
}
}
}
#endif //vtk_m_filter_StreamSurface_h

32
vtkm/filter/Streamline.h Normal file
View File

@ -0,0 +1,32 @@
//============================================================================
// 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_Streamline_h
#define vtk_m_filter_Streamline_h
#include <vtkm/Deprecated.h>
#include <vtkm/filter/flow/Streamline.h>
namespace vtkm
{
namespace filter
{
VTKM_DEPRECATED(1.8, "Use vtkm/filter/flow/Streamline.h instead of vtkm/filter/Streamline.h")
inline void Streamline_deprecated() {}
inline void Streamline_deprecated_warning()
{
Streamline_deprecated();
}
}
}
#endif //vtk_m_filter_Streamline_h

1
vtkm/filter/TaskQueue.h
View File

@ -105,7 +105,6 @@ public:
while (this->GetTask(task))
{
dataSets[static_cast<std::size_t>(task.first)] = std::move(task.second);
std::cout << "****** Get: " << task.first << std::endl;
}
pds.AppendPartitions(dataSets);

13
vtkm/filter/clean_grid/worklet/RemoveDegenerateCells.h
View File

@ -133,22 +133,11 @@ struct RemoveDegenerateCells
return output;
}
struct CallWorklet
{
template <typename CellSetType>
void operator()(const CellSetType& cellSet,
RemoveDegenerateCells& self,
vtkm::cont::CellSetExplicit<>& output) const
{
output = self.Run(cellSet);
}
};
template <typename CellSetList>
vtkm::cont::CellSetExplicit<> Run(const vtkm::cont::UncertainCellSet<CellSetList>& cellSet)
{
vtkm::cont::CellSetExplicit<> output;
cellSet.CastAndCall(CallWorklet(), *this, output);
cellSet.CastAndCall([&](const auto& concrete) { output = this->Run(concrete); });
return output;
}

2
vtkm/filter/connected_components/CellSetConnectivity.cxx
View File

@ -21,7 +21,7 @@ VTKM_CONT vtkm::cont::DataSet CellSetConnectivity::DoExecute(const vtkm::cont::D
{
vtkm::cont::ArrayHandle<vtkm::Id> component;
vtkm::worklet::connectivity::CellSetConnectivity().Run(input.GetCellSet(), component);
vtkm::worklet::connectivity::CellSetConnectivity::Run(input.GetCellSet(), component);
return this->CreateResultFieldCell(input, this->GetOutputFieldName(), component);
}

2
vtkm/filter/connected_components/testing/UnitTestGraphConnectivityWorklet.cxx
View File

@ -93,7 +93,7 @@ public:
vtkm::cont::ArrayHandle<int> conns_h = vtkm::cont::make_ArrayHandle(conns, vtkm::CopyFlag::Off);
vtkm::cont::ArrayHandle<vtkm::Id> comps_h;
vtkm::worklet::connectivity::GraphConnectivity().Run(counts_h, offsets_h, conns_h, comps_h);
vtkm::worklet::connectivity::GraphConnectivity::Run(counts_h, offsets_h, conns_h, comps_h);
VTKM_TEST_ASSERT(vtkm::cont::Algorithm::Reduce(comps_h, vtkm::Id(0), vtkm::Maximum{}) ==
ncomps - 1,

8
vtkm/filter/connected_components/worklet/CellSetConnectivity.h
View File

@ -23,17 +23,17 @@ namespace connectivity
class CellSetConnectivity
{
public:
template <typename CellSetType>
void Run(const CellSetType& cellSet, vtkm::cont::ArrayHandle<vtkm::Id>& componentArray) const
static void Run(const vtkm::cont::UnknownCellSet& cellSet,
vtkm::cont::ArrayHandle<vtkm::Id>& componentArray)
{
vtkm::cont::ArrayHandle<vtkm::Id> numIndicesArray;
vtkm::cont::ArrayHandle<vtkm::Id> indexOffsetsArray;
vtkm::cont::ArrayHandle<vtkm::Id> connectivityArray;
// create cell to cell connectivity graph (dual graph)
CellSetDualGraph().Run(cellSet, numIndicesArray, indexOffsetsArray, connectivityArray);
CellSetDualGraph::Run(cellSet, numIndicesArray, indexOffsetsArray, connectivityArray);
// find the connected component of the dual graph
GraphConnectivity().Run(numIndicesArray, indexOffsetsArray, connectivityArray, componentArray);
GraphConnectivity::Run(numIndicesArray, indexOffsetsArray, connectivityArray, componentArray);
}
};
}

30
vtkm/filter/connected_components/worklet/CellSetDualGraph.h
View File

@ -99,19 +99,11 @@ struct CellToCellConnectivity : public vtkm::worklet::WorkletMapField
class CellSetDualGraph
{
public:
using Algorithm = vtkm::cont::Algorithm;
struct degree2
{
VTKM_EXEC
bool operator()(vtkm::Id degree) const { return degree >= 2; }
};
template <typename CellSet>
void EdgeToCellConnectivity(const CellSet& cellSet,
vtkm::cont::ArrayHandle<vtkm::Id>& cellIds,
vtkm::cont::ArrayHandle<vtkm::Id2>& cellEdges) const
static void EdgeToCellConnectivity(const vtkm::cont::UnknownCellSet& cellSet,
vtkm::cont::ArrayHandle<vtkm::Id>& cellIds,
vtkm::cont::ArrayHandle<vtkm::Id2>& cellEdges)
{
// Get number of edges for each cell and use it as scatter count.
vtkm::cont::ArrayHandle<vtkm::IdComponent> numEdgesPerCell;
@ -124,11 +116,17 @@ public:
edgeExtractDisp.Invoke(cellSet, cellIds, cellEdges);
}
template <typename CellSetType>
void Run(const CellSetType& cellSet,
vtkm::cont::ArrayHandle<vtkm::Id>& numIndicesArray,
vtkm::cont::ArrayHandle<vtkm::Id>& indexOffsetArray,
vtkm::cont::ArrayHandle<vtkm::Id>& connectivityArray) const
public:
struct degree2
{
VTKM_EXEC
bool operator()(vtkm::Id degree) const { return degree >= 2; }
};
static void Run(const vtkm::cont::UnknownCellSet& cellSet,
vtkm::cont::ArrayHandle<vtkm::Id>& numIndicesArray,
vtkm::cont::ArrayHandle<vtkm::Id>& indexOffsetArray,
vtkm::cont::ArrayHandle<vtkm::Id>& connectivityArray)
{
// calculate the uncompressed Edge to Cell connectivity from Point to Cell connectivity
// in the CellSet

8
vtkm/filter/connected_components/worklet/GraphConnectivity.h
View File

@ -69,10 +69,10 @@ class GraphConnectivity
{
public:
template <typename InputArrayType, typename OutputArrayType>
void Run(const InputArrayType& numIndicesArray,
const InputArrayType& indexOffsetsArray,
const InputArrayType& connectivityArray,
OutputArrayType& componentsOut) const
static void Run(const InputArrayType& numIndicesArray,
const InputArrayType& indexOffsetsArray,
const InputArrayType& connectivityArray,
OutputArrayType& componentsOut)
{
VTKM_IS_ARRAY_HANDLE(InputArrayType);
VTKM_IS_ARRAY_HANDLE(OutputArrayType);

40
vtkm/filter/connected_components/worklet/ImageConnectivity.h
View File

@ -90,13 +90,12 @@ public:
// Parallel and Distributed Computing. 2018.
class ImageConnectivity
{
public:
class RunImpl
{
public:
template <int Dimension, typename T, typename StorageT, typename OutputPortalType>
void operator()(const vtkm::cont::ArrayHandle<T, StorageT>& pixels,
const vtkm::cont::CellSetStructured<Dimension>& input,
void operator()(const vtkm::cont::CellSetStructured<Dimension>& input,
const vtkm::cont::ArrayHandle<T, StorageT>& pixels,
OutputPortalType& componentsOut) const
{
using Algorithm = vtkm::cont::Algorithm;
@ -115,43 +114,14 @@ public:
}
};
class ResolveUnknownCellSet
{
public:
template <int Dimension, typename T, typename StorageT, typename OutputPortalType>
void operator()(const vtkm::cont::CellSetStructured<Dimension>& input,
const vtkm::cont::ArrayHandle<T, StorageT>& pixels,
OutputPortalType& components) const
{
vtkm::cont::CastAndCall(pixels, RunImpl(), input, components);
}
};
template <int Dimension, typename OutputPortalType>
void Run(const vtkm::cont::CellSetStructured<Dimension>& input,
const vtkm::cont::UnknownArrayHandle& pixels,
OutputPortalType& componentsOut) const
{
using Types = vtkm::TypeListScalarAll;
using Storages = VTKM_DEFAULT_STORAGE_LIST;
vtkm::cont::CastAndCall(pixels.ResetTypes<Types, Storages>(), RunImpl(), input, componentsOut);
}
template <int Dimension, typename T, typename S, typename OutputPortalType>
void Run(const vtkm::cont::CellSetStructured<Dimension>& input,
const vtkm::cont::ArrayHandle<T, S>& pixels,
OutputPortalType& componentsOut) const
{
vtkm::cont::CastAndCall(pixels, RunImpl(), input, componentsOut);
}
public:
template <typename T, typename S, typename OutputPortalType>
void Run(const vtkm::cont::UnknownCellSet& input,
const vtkm::cont::ArrayHandle<T, S>& pixels,
OutputPortalType& componentsOut) const
{
input.ResetCellSetList(vtkm::cont::CellSetListStructured())
.CastAndCall(ResolveUnknownCellSet(), pixels, componentsOut);
input.template CastAndCallForTypes<vtkm::cont::CellSetListStructured>(
RunImpl(), pixels, componentsOut);
}
};
} // namespace connectivity

56
vtkm/filter/entity_extraction/ExtractGeometry.cxx
View File

@ -16,45 +16,6 @@
namespace
{
struct CallWorker
{
vtkm::cont::UnknownCellSet& Output;
vtkm::worklet::ExtractGeometry& Worklet;
const vtkm::cont::CoordinateSystem& Coords;
const vtkm::ImplicitFunctionGeneral& Function;
bool ExtractInside;
bool ExtractBoundaryCells;
bool ExtractOnlyBoundaryCells;
CallWorker(vtkm::cont::UnknownCellSet& output,
vtkm::worklet::ExtractGeometry& worklet,
const vtkm::cont::CoordinateSystem& coords,
const vtkm::ImplicitFunctionGeneral& function,
bool extractInside,
bool extractBoundaryCells,
bool extractOnlyBoundaryCells)
: Output(output)
, Worklet(worklet)
, Coords(coords)
, Function(function)
, ExtractInside(extractInside)
, ExtractBoundaryCells(extractBoundaryCells)
, ExtractOnlyBoundaryCells(extractOnlyBoundaryCells)
{
}
template <typename CellSetType>
void operator()(const CellSetType& cellSet) const
{
this->Output = this->Worklet.Run(cellSet,
this->Coords,
this->Function,
this->ExtractInside,
this->ExtractBoundaryCells,
this->ExtractOnlyBoundaryCells);
}
};
bool DoMapField(vtkm::cont::DataSet& result,
const vtkm::cont::Field& field,
const vtkm::worklet::ExtractGeometry& worklet)
@ -97,14 +58,15 @@ vtkm::cont::DataSet ExtractGeometry::DoExecute(const vtkm::cont::DataSet& input)
vtkm::worklet::ExtractGeometry worklet;
vtkm::cont::UnknownCellSet outCells;
CallWorker worker(outCells,
worklet,
coords,
this->Function,
this->ExtractInside,
this->ExtractBoundaryCells,
this->ExtractOnlyBoundaryCells);
cells.CastAndCallForTypes<VTKM_DEFAULT_CELL_SET_LIST>(worker);
cells.CastAndCallForTypes<VTKM_DEFAULT_CELL_SET_LIST>([&](const auto& concrete) {
outCells = worklet.Run(concrete,
coords,
this->Function,
this->ExtractInside,
this->ExtractBoundaryCells,
this->ExtractOnlyBoundaryCells);
});
// create the output dataset
auto mapper = [&](auto& result, const auto& f) { DoMapField(result, f, worklet); };

11
vtkm/filter/entity_extraction/GhostCellRemove.cxx
View File

@ -352,17 +352,12 @@ VTKM_CONT vtkm::cont::DataSet GhostCellRemove::DoExecute(const vtkm::cont::DataS
if (this->GetRemoveAllGhost())
{
cellOut = worklet.Run(cells.ResetCellSetList<VTKM_DEFAULT_CELL_SET_LIST>(),
fieldArray,
field.GetAssociation(),
RemoveAllGhosts());
cellOut = worklet.Run(cells, fieldArray, field.GetAssociation(), RemoveAllGhosts());
}
else if (this->GetRemoveByType())
{
cellOut = worklet.Run(cells.ResetCellSetList<VTKM_DEFAULT_CELL_SET_LIST>(),
fieldArray,
field.GetAssociation(),
RemoveGhostByType(this->GetRemoveType()));
cellOut = worklet.Run(
cells, fieldArray, field.GetAssociation(), RemoveGhostByType(this->GetRemoveType()));
}
else
{

24
vtkm/filter/entity_extraction/Mask.cxx
View File

@ -13,26 +13,6 @@
namespace
{
struct CallWorklet
{
vtkm::Id Stride;
vtkm::cont::UnknownCellSet& Output;
vtkm::worklet::Mask& Worklet;
CallWorklet(vtkm::Id stride, vtkm::cont::UnknownCellSet& output, vtkm::worklet::Mask& worklet)
: Stride(stride)
, Output(output)
, Worklet(worklet)
{
}
template <typename CellSetType>
void operator()(const CellSetType& cells) const
{
this->Output = this->Worklet.Run(cells, this->Stride);
}
};
VTKM_CONT bool DoMapField(vtkm::cont::DataSet& result,
const vtkm::cont::Field& field,
const vtkm::worklet::Mask& worklet)
@ -66,8 +46,8 @@ VTKM_CONT vtkm::cont::DataSet Mask::DoExecute(const vtkm::cont::DataSet& input)
vtkm::cont::UnknownCellSet cellOut;
vtkm::worklet::Mask worklet;
CallWorklet workletCaller(this->Stride, cellOut, worklet);
cells.CastAndCallForTypes<VTKM_DEFAULT_CELL_SET_LIST>(workletCaller);
cells.CastAndCallForTypes<VTKM_DEFAULT_CELL_SET_LIST>(
[&](const auto& concrete) { cellOut = worklet.Run(concrete, this->Stride); });
// create the output dataset
auto mapper = [&](auto& result, const auto& f) { DoMapField(result, f, worklet); };

11
vtkm/filter/entity_extraction/Threshold.cxx
View File

@ -83,15 +83,8 @@ vtkm::cont::DataSet Threshold::DoExecute(const vtkm::cont::DataSet& input)
vtkm::cont::UnknownCellSet cellOut;
auto resolveArrayType = [&](const auto& concrete) {
// Note: there are two overloads of .Run, the first one taking an UncertainCellSet, which is
// the desired entry point in the following call. The other is a function template on the input
// CellSet. Without the call to .ResetCellSetList to turn an UnknownCellSet to an UncertainCellSet,
// the compiler will pick the function template (i.e. wrong overload).
cellOut = worklet.Run(cells.ResetCellSetList<VTKM_DEFAULT_CELL_SET_LIST>(),
concrete,
field.GetAssociation(),
predicate,
this->GetAllInRange());
cellOut =
worklet.Run(cells, concrete, field.GetAssociation(), predicate, this->GetAllInRange());
};
field.GetData().CastAndCallForTypes<vtkm::TypeListScalarAll, VTKM_DEFAULT_STORAGE_LIST>(

69
vtkm/filter/entity_extraction/worklet/Threshold.h
View File

@ -30,12 +30,6 @@ namespace worklet
class Threshold
{
public:
enum class FieldType
{
Point,
Cell
};
template <typename UnaryPredicate>
class ThresholdByPointField : public vtkm::worklet::WorkletVisitCellsWithPoints
{
@ -84,22 +78,8 @@ public:
bool ReturnAllInRange;
};
struct ThresholdCopy : public vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn, FieldOut, WholeArrayIn);
template <typename ScalarType, typename WholeFieldIn>
VTKM_EXEC void operator()(vtkm::Id& index,
ScalarType& output,
const WholeFieldIn& inputField) const
{
output = inputField.Get(index);
}
};
template <typename CellSetType, typename ValueType, typename StorageType, typename UnaryPredicate>
vtkm::cont::CellSetPermutation<CellSetType> Run(
vtkm::cont::CellSetPermutation<CellSetType> RunImpl(
const CellSetType& cellSet,
const vtkm::cont::ArrayHandle<ValueType, StorageType>& field,
const vtkm::cont::Field::Association fieldType,
@ -141,53 +121,18 @@ public:
return OutputType(this->ValidCellIds, cellSet);
}
template <typename FieldArrayType, typename UnaryPredicate>
struct CallWorklet
{
vtkm::cont::UnknownCellSet& Output;
vtkm::worklet::Threshold& Worklet;
const FieldArrayType& Field;
const vtkm::cont::Field::Association FieldType;
const UnaryPredicate& Predicate;
const bool ReturnAllInRange;
CallWorklet(vtkm::cont::UnknownCellSet& output,
vtkm::worklet::Threshold& worklet,
const FieldArrayType& field,
const vtkm::cont::Field::Association fieldType,
const UnaryPredicate& predicate,
const bool returnAllInRange)
: Output(output)
, Worklet(worklet)
, Field(field)
, FieldType(fieldType)
, Predicate(predicate)
, ReturnAllInRange(returnAllInRange)
{
}
template <typename CellSetType>
void operator()(const CellSetType& cellSet) const
{
// Copy output to an explicit grid so that other units can guess what this is.
this->Output = vtkm::worklet::CellDeepCopy::Run(this->Worklet.Run(
cellSet, this->Field, this->FieldType, this->Predicate, this->ReturnAllInRange));
}
};
template <typename CellSetList, typename ValueType, typename StorageType, typename UnaryPredicate>
vtkm::cont::UnknownCellSet Run(const vtkm::cont::UncertainCellSet<CellSetList>& cellSet,
template <typename ValueType, typename StorageType, typename UnaryPredicate>
vtkm::cont::UnknownCellSet Run(const vtkm::cont::UnknownCellSet& cellSet,
const vtkm::cont::ArrayHandle<ValueType, StorageType>& field,
const vtkm::cont::Field::Association fieldType,
const UnaryPredicate& predicate,
const bool returnAllInRange = false)
{
using Worker = CallWorklet<vtkm::cont::ArrayHandle<ValueType, StorageType>, UnaryPredicate>;
vtkm::cont::UnknownCellSet output;
Worker worker(output, *this, field, fieldType, predicate, returnAllInRange);
cellSet.CastAndCall(worker);
CastAndCall(cellSet, [&](auto concrete) {
output = vtkm::worklet::CellDeepCopy::Run(
this->RunImpl(concrete, field, fieldType, predicate, returnAllInRange));
});
return output;
}

6
vtkm/filter/field_transform/CylindricalCoordinateTransform.cxx
View File

@ -26,8 +26,10 @@ vtkm::cont::DataSet CylindricalCoordinateTransform::DoExecute(const vtkm::cont::
// use std::decay to remove const ref from the decltype of concrete.
using T = typename std::decay_t<decltype(concrete)>::ValueType;
vtkm::cont::ArrayHandle<T> result;
vtkm::worklet::CylindricalCoordinateTransform worklet{ this->CartesianToCylindrical };
worklet.Run(concrete, result);
if (this->CartesianToCylindrical)
this->Invoke(vtkm::worklet::CarToCyl{}, concrete, result);
else
this->Invoke(vtkm::worklet::CylToCar{}, concrete, result);
outArray = result;
};
this->CastAndCallVecField<3>(this->GetFieldFromDataSet(inDataSet), resolveType);

6
vtkm/filter/field_transform/SphericalCoordinateTransform.cxx
View File

@ -25,8 +25,10 @@ vtkm::cont::DataSet SphericalCoordinateTransform::DoExecute(const vtkm::cont::Da
// use std::decay to remove const ref from the decltype of concrete.
using T = typename std::decay_t<decltype(concrete)>::ValueType;
vtkm::cont::ArrayHandle<T> result;
vtkm::worklet::SphericalCoordinateTransform worklet{ this->CartesianToSpherical };
worklet.Run(concrete, result);
if (this->CartesianToSpherical)
this->Invoke(vtkm::worklet::CarToSphere{}, concrete, result);
else
this->Invoke(vtkm::worklet::SphereToCar{}, concrete, result);
outArray = result;
};
this->CastAndCallVecField<3>(this->GetFieldFromDataSet(inDataSet), resolveType);

4
vtkm/filter/field_transform/WarpScalar.cxx
View File

@ -68,8 +68,8 @@ VTKM_CONT vtkm::cont::DataSet WarpScalar::DoExecute(const vtkm::cont::DataSet& i
using VecType = typename std::decay_t<decltype(concrete)>::ValueType;
vtkm::cont::ArrayHandle<VecType> result;
vtkm::worklet::WarpScalar worklet;
worklet.Run(concrete, normalArray, scaleFactorArray, this->ScaleAmount, result);
vtkm::worklet::WarpScalar worklet{ this->ScaleAmount };
this->Invoke(worklet, concrete, normalArray, scaleFactorArray, result);
outArray = result;
};
const auto& field = this->GetFieldFromDataSet(inDataSet);

4
vtkm/filter/field_transform/WarpVector.cxx
View File

@ -55,8 +55,8 @@ VTKM_CONT vtkm::cont::DataSet WarpVector::DoExecute(const vtkm::cont::DataSet& i
using VecType = typename std::decay_t<decltype(concrete)>::ValueType;
vtkm::cont::ArrayHandle<VecType> result;
vtkm::worklet::WarpVector worklet;
worklet.Run(concrete, vectorArray, this->Scale, result);
vtkm::worklet::WarpVector worklet{ this->Scale };
this->Invoke(worklet, concrete, vectorArray, result);
outArray = result;
};
const auto& field = this->GetFieldFromDataSet(inDataSet);

105
vtkm/filter/field_transform/worklet/CoordinateSystemTransform.h
View File

@ -12,25 +12,19 @@
#define vtk_m_worklet_CoordinateSystemTransform_h
#include <vtkm/Math.h>
#include <vtkm/Matrix.h>
#include <vtkm/Transform3D.h>
#include <vtkm/VectorAnalysis.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
namespace vtkm
{
namespace worklet
{
namespace detail
{
template <typename T>
struct CylToCar : public vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn, FieldOut);
using ExecutionSignature = _2(_1);
//Functor
template <typename T>
VTKM_EXEC vtkm::Vec<T, 3> operator()(const vtkm::Vec<T, 3>& vec) const
{
vtkm::Vec<T, 3> res(vec[0] * static_cast<T>(vtkm::Cos(vec[1])),
@ -40,13 +34,13 @@ struct CylToCar : public vtkm::worklet::WorkletMapField
}
};
template <typename T>
struct CarToCyl : public vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn, FieldOut);
using ExecutionSignature = _2(_1);
//Functor
template <typename T>
VTKM_EXEC vtkm::Vec<T, 3> operator()(const vtkm::Vec<T, 3>& vec) const
{
T R = vtkm::Sqrt(vec[0] * vec[0] + vec[1] * vec[1]);
@ -64,13 +58,13 @@ struct CarToCyl : public vtkm::worklet::WorkletMapField
}
};
template <typename T>
struct SphereToCar : public vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn, FieldOut);
using ExecutionSignature = _2(_1);
//Functor
template <typename T>
VTKM_EXEC vtkm::Vec<T, 3> operator()(const vtkm::Vec<T, 3>& vec) const
{
T R = vec[0];
@ -91,13 +85,13 @@ struct SphereToCar : public vtkm::worklet::WorkletMapField
}
};
template <typename T>
struct CarToSphere : public vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn, FieldOut);
using ExecutionSignature = _2(_1);
//Functor
template <typename T>
VTKM_EXEC vtkm::Vec<T, 3> operator()(const vtkm::Vec<T, 3>& vec) const
{
T R = vtkm::Sqrt(vtkm::Dot(vec, vec));
@ -112,97 +106,6 @@ struct CarToSphere : public vtkm::worklet::WorkletMapField
}
};
}
class CylindricalCoordinateTransform
{
public:
VTKM_CONT
explicit CylindricalCoordinateTransform(bool car2cyl)
: CartesianToCylindrical(car2cyl)
{
}
template <typename T, typename InStorageType, typename OutStorageType>
void Run(const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>, InStorageType>& inPoints,
vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>, OutStorageType>& outPoints) const
{
if (CartesianToCylindrical)
{
vtkm::worklet::DispatcherMapField<detail::CarToCyl<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
else
{
vtkm::worklet::DispatcherMapField<detail::CylToCar<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
}
template <typename T, typename CoordsStorageType>
void Run(const vtkm::cont::CoordinateSystem& inPoints,
vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>, CoordsStorageType>& outPoints) const
{
if (CartesianToCylindrical)
{
vtkm::worklet::DispatcherMapField<detail::CarToCyl<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
else
{
vtkm::worklet::DispatcherMapField<detail::CylToCar<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
}
private:
bool CartesianToCylindrical;
};
class SphericalCoordinateTransform
{
public:
VTKM_CONT
explicit SphericalCoordinateTransform(bool car2sph)
: CartesianToSpherical(car2sph)
{
}
template <typename T, typename InStorageType, typename OutStorageType>
void Run(const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>, InStorageType>& inPoints,
vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>, OutStorageType>& outPoints) const
{
if (CartesianToSpherical)
{
vtkm::worklet::DispatcherMapField<detail::CarToSphere<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
else
{
vtkm::worklet::DispatcherMapField<detail::SphereToCar<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
}
template <typename T, typename CoordsStorageType>
void Run(const vtkm::cont::CoordinateSystem& inPoints,
vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>, CoordsStorageType>& outPoints) const
{
if (CartesianToSpherical)
{
vtkm::worklet::DispatcherMapField<detail::CarToSphere<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
else
{
vtkm::worklet::DispatcherMapField<detail::SphereToCar<T>> dispatcher;
dispatcher.Invoke(inPoints, outPoints);
}
}
private:
bool CartesianToSpherical;
};
}
} // namespace vtkm::worklet
#endif // vtk_m_worklet_CoordinateSystemTransform_h

81
vtkm/filter/field_transform/worklet/WarpScalar.h
View File

@ -11,12 +11,8 @@
#ifndef vtk_m_worklet_WarpScalar_h
#define vtk_m_worklet_WarpScalar_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/VecTraits.h>
namespace vtkm
{
namespace worklet
@ -26,60 +22,37 @@ namespace worklet
// Useful for creating carpet or x-y-z plots.
// It doesn't modify the point coordinates, but creates a new point coordinates
// that have been warped.
class WarpScalar
class WarpScalar : public vtkm::worklet::WorkletMapField
{
public:
class WarpScalarImp : public vtkm::worklet::WorkletMapField
using ControlSignature = void(FieldIn, FieldIn, FieldIn, FieldOut);
using ExecutionSignature = void(_1, _2, _3, _4);
VTKM_CONT
explicit WarpScalar(vtkm::FloatDefault scaleAmount)
: ScaleAmount(scaleAmount)
{
public:
using ControlSignature = void(FieldIn, FieldIn, FieldIn, FieldOut);
using ExecutionSignature = void(_1, _2, _3, _4);
VTKM_CONT
explicit WarpScalarImp(vtkm::FloatDefault scaleAmount)
: ScaleAmount(scaleAmount)
{
}
VTKM_EXEC
void operator()(const vtkm::Vec3f& point,
const vtkm::Vec3f& normal,
const vtkm::FloatDefault& scaleFactor,
vtkm::Vec3f& result) const
{
result = point + this->ScaleAmount * scaleFactor * normal;
}
template <typename T1, typename T2, typename T3>
VTKM_EXEC void operator()(const vtkm::Vec<T1, 3>& point,
const vtkm::Vec<T2, 3>& normal,
const T3& scaleFactor,
vtkm::Vec<T1, 3>& result) const
{
result = point + static_cast<T1>(this->ScaleAmount * scaleFactor) * normal;
}
private:
vtkm::FloatDefault ScaleAmount;
};
// Execute the WarpScalar worklet given the points, vector and a scale factor.
// Scale factor can differs per point.
template <typename PointType,
typename NormalType,
typename ScaleFactorType,
typename ResultType,
typename ScaleAmountType>
void Run(PointType point,
NormalType normal,
ScaleFactorType scaleFactor,
ScaleAmountType scale,
ResultType warpedPoint)
{
WarpScalarImp warpScalarImp(scale);
vtkm::worklet::DispatcherMapField<WarpScalarImp> dispatcher(warpScalarImp);
dispatcher.Invoke(point, normal, scaleFactor, warpedPoint);
}
VTKM_EXEC
void operator()(const vtkm::Vec3f& point,
const vtkm::Vec3f& normal,
const vtkm::FloatDefault& scaleFactor,
vtkm::Vec3f& result) const
{
result = point + this->ScaleAmount * scaleFactor * normal;
}
template <typename T1, typename T2, typename T3>
VTKM_EXEC void operator()(const vtkm::Vec<T1, 3>& point,
const vtkm::Vec<T2, 3>& normal,
const T3& scaleFactor,
vtkm::Vec<T1, 3>& result) const
{
result = point + static_cast<T1>(this->ScaleAmount * scaleFactor) * normal;
}
private:
vtkm::FloatDefault ScaleAmount;
};
}
} // namespace vtkm::worklet

66
vtkm/filter/field_transform/worklet/WarpVector.h
View File

@ -11,13 +11,8 @@
#ifndef vtk_m_worklet_WarpVector_h
#define vtk_m_worklet_WarpVector_h
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/VecTraits.h>
namespace vtkm
{
namespace worklet
@ -27,48 +22,33 @@ namespace worklet
// Useful for showing flow profiles or mechanical deformation.
// This worklet does not modify the input points but generate new point coordinate
// instance that has been warped.
class WarpVector
class WarpVector : public vtkm::worklet::WorkletMapField
{
public:
class WarpVectorImp : public vtkm::worklet::WorkletMapField
using ControlSignature = void(FieldIn, FieldIn, FieldOut);
using ExecutionSignature = _3(_1, _2);
VTKM_CONT
explicit WarpVector(vtkm::FloatDefault scale)
: Scale(scale)
{
public:
using ControlSignature = void(FieldIn, FieldIn, FieldOut);
using ExecutionSignature = _3(_1, _2);
VTKM_CONT
explicit WarpVectorImp(vtkm::FloatDefault scale)
: Scale(scale)
{
}
VTKM_EXEC
vtkm::Vec3f operator()(const vtkm::Vec3f& point, const vtkm::Vec3f& vector) const
{
return point + this->Scale * vector;
}
template <typename T>
VTKM_EXEC vtkm::Vec<T, 3> operator()(const vtkm::Vec<T, 3>& point,
const vtkm::Vec<T, 3>& vector) const
{
return point + static_cast<T>(this->Scale) * vector;
}
private:
vtkm::FloatDefault Scale;
};
// Execute the WarpVector worklet given the points, vector and a scale factor.
// Returns:
// warped points
template <typename PointType, typename VectorType, typename ResultType>
void Run(PointType point, VectorType vector, vtkm::FloatDefault scale, ResultType warpedPoint)
{
WarpVectorImp warpVectorImp(scale);
vtkm::worklet::DispatcherMapField<WarpVectorImp> dispatcher(warpVectorImp);
dispatcher.Invoke(point, vector, warpedPoint);
}
VTKM_EXEC
vtkm::Vec3f operator()(const vtkm::Vec3f& point, const vtkm::Vec3f& vector) const
{
return point + this->Scale * vector;
}
template <typename T>
VTKM_EXEC vtkm::Vec<T, 3> operator()(const vtkm::Vec<T, 3>& point,
const vtkm::Vec<T, 3>& vector) const
{
return point + static_cast<T>(this->Scale) * vector;
}
private:
vtkm::FloatDefault Scale;
};
}
} // namespace vtkm::worklet

2
vtkm/filter/flow/NewFilterParticleAdvection.cxx
View File

@ -37,7 +37,7 @@ VTKM_CONT void NewFilterParticleAdvection::ValidateOptions() const
if (this->Seeds.GetNumberOfValues() == 0)
throw vtkm::cont::ErrorFilterExecution("No seeds provided.");
if (!this->Seeds.IsBaseComponentType<vtkm::Particle>() &&
this->Seeds.IsBaseComponentType<vtkm::ChargedParticle>())
!this->Seeds.IsBaseComponentType<vtkm::ChargedParticle>())
throw vtkm::cont::ErrorFilterExecution("Unsupported particle type in seed array.");
if (this->NumberOfSteps == 0)
throw vtkm::cont::ErrorFilterExecution("Number of steps not specified.");

22
vtkm/filter/flow/NewFilterParticleAdvection.h
View File

@ -11,6 +11,7 @@
#ifndef vtk_m_filter_flow_NewFilterParticleAdvection_h
#define vtk_m_filter_flow_NewFilterParticleAdvection_h
#include <vtkm/Particle.h>
#include <vtkm/filter/NewFilterField.h>
#include <vtkm/filter/flow/FlowTypes.h>
#include <vtkm/filter/flow/vtkm_filter_flow_export.h>
@ -30,6 +31,9 @@ namespace flow
class VTKM_FILTER_FLOW_EXPORT NewFilterParticleAdvection : public vtkm::filter::NewFilterField
{
public:
VTKM_CONT
bool CanThread() const override { return false; }
VTKM_CONT
void SetStepSize(vtkm::FloatDefault s) { this->StepSize = s; }
@ -51,12 +55,30 @@ public:
VTKM_CONT
void SetSolverRK4() { this->SolverType = vtkm::filter::flow::IntegrationSolverType::RK4_TYPE; }
VTKM_CONT
void SetSolverEuler()
{
this->SolverType = vtkm::filter::flow::IntegrationSolverType::EULER_TYPE;
}
VTKM_CONT
void SetVectorFieldType(vtkm::filter::flow::VectorFieldType vecFieldType)
{
this->VecFieldType = vecFieldType;
}
//@{
/// Choose the field to operate on. Note, if
/// `this->UseCoordinateSystemAsField` is true, then the active field is not used.
VTKM_CONT void SetEField(const std::string& name) { this->SetActiveField(0, name); }
VTKM_CONT void SetBField(const std::string& name) { this->SetActiveField(1, name); }
VTKM_CONT std::string GetEField() const { return this->GetActiveFieldName(0); }
VTKM_CONT std::string GetBField() const { return this->GetActiveFieldName(1); }
VTKM_CONT
bool GetUseThreadedAlgorithm() { return this->UseThreadedAlgorithm; }

59
vtkm/filter/flow/NewFilterParticleAdvectionSteadyState.cxx
View File

@ -24,12 +24,14 @@ namespace flow
namespace
{
VTKM_CONT std::vector<vtkm::filter::flow::internal::DataSetIntegratorSteadyState>
CreateDataSetIntegrators(const vtkm::cont::PartitionedDataSet& input,
const std::string& activeField,
const vtkm::filter::flow::internal::BoundsMap& boundsMap,
const vtkm::filter::flow::IntegrationSolverType solverType,
const vtkm::filter::flow::VectorFieldType vecFieldType,
const vtkm::filter::flow::FlowResultType resultType)
CreateDataSetIntegrators(
const vtkm::cont::PartitionedDataSet& input,
const vtkm::cont::internal::Variant<std::string, std::pair<std::string, std::string>>&
activeField,
const vtkm::filter::flow::internal::BoundsMap& boundsMap,
const vtkm::filter::flow::IntegrationSolverType solverType,
const vtkm::filter::flow::VectorFieldType vecFieldType,
const vtkm::filter::flow::FlowResultType resultType)
{
using DSIType = vtkm::filter::flow::internal::DataSetIntegratorSteadyState;
@ -38,14 +40,27 @@ CreateDataSetIntegrators(const vtkm::cont::PartitionedDataSet& input,
{
vtkm::Id blockId = boundsMap.GetLocalBlockId(i);
auto ds = input.GetPartition(i);
if (!ds.HasPointField(activeField) && !ds.HasCellField(activeField))
throw vtkm::cont::ErrorFilterExecution("Unsupported field assocation");
if (activeField.IsType<DSIType::VelocityFieldNameType>())
{
const auto& fieldNm = activeField.Get<DSIType::VelocityFieldNameType>();
if (!ds.HasPointField(fieldNm) && !ds.HasCellField(fieldNm))
throw vtkm::cont::ErrorFilterExecution("Unsupported field assocation");
}
else if (activeField.IsType<DSIType::ElectroMagneticFieldNameType>())
{
const auto& fieldNm = activeField.Get<DSIType::ElectroMagneticFieldNameType>();
const auto& electric = fieldNm.first;
const auto& magnetic = fieldNm.second;
if (!ds.HasPointField(electric) && !ds.HasCellField(electric))
throw vtkm::cont::ErrorFilterExecution("Unsupported field assocation");
if (!ds.HasPointField(magnetic) && !ds.HasCellField(magnetic))
throw vtkm::cont::ErrorFilterExecution("Unsupported field assocation");
}
dsi.emplace_back(ds, blockId, activeField, solverType, vecFieldType, resultType);
}
return dsi;
}
} //anonymous namespace
VTKM_CONT vtkm::cont::PartitionedDataSet NewFilterParticleAdvectionSteadyState::DoExecutePartitions(
@ -54,13 +69,25 @@ VTKM_CONT vtkm::cont::PartitionedDataSet NewFilterParticleAdvectionSteadyState::
using DSIType = vtkm::filter::flow::internal::DataSetIntegratorSteadyState;
this->ValidateOptions();
using VariantType =
vtkm::cont::internal::Variant<std::string, std::pair<std::string, std::string>>;
VariantType variant;
if (this->VecFieldType == vtkm::filter::flow::VectorFieldType::VELOCITY_FIELD_TYPE)
{
const auto& field = this->GetActiveFieldName();
variant.Emplace<DSIType::VelocityFieldNameType>(field);
}
else if (this->VecFieldType == vtkm::filter::flow::VectorFieldType::ELECTRO_MAGNETIC_FIELD_TYPE)
{
const auto& electric = this->GetEField();
const auto& magnetic = this->GetBField();
variant.Emplace<DSIType::ElectroMagneticFieldNameType>(electric, magnetic);
}
vtkm::filter::flow::internal::BoundsMap boundsMap(input);
auto dsi = CreateDataSetIntegrators(input,
this->GetActiveFieldName(),
boundsMap,
this->SolverType,
this->VecFieldType,
this->GetResultType());
auto dsi = CreateDataSetIntegrators(
input, variant, boundsMap, this->SolverType, this->VecFieldType, this->GetResultType());
vtkm::filter::flow::internal::ParticleAdvector<DSIType> pav(
boundsMap, dsi, this->UseThreadedAlgorithm, this->GetResultType());

5
vtkm/filter/flow/ParticleAdvection.h
View File

@ -36,6 +36,11 @@ private:
};
}
struct VTKM_DEPRECATED(1.8, "Use vtkm::filter::flow::ParticleAdvection.") ParticleAdvection
: vtkm::filter::flow::ParticleAdvection
{
using vtkm::filter::flow::ParticleAdvection::ParticleAdvection;
};
}
} // namespace vtkm::filter::flow

11
vtkm/filter/flow/PathParticle.h
View File

@ -8,8 +8,8 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_filter_PathParticle_h
#define vtk_m_filter_PathParticle_h
#ifndef vtk_m_filter_flow_PathParticle_h
#define vtk_m_filter_flow_PathParticle_h
#include <vtkm/filter/flow/FlowTypes.h>
#include <vtkm/filter/flow/NewFilterParticleAdvectionUnsteadyState.h>
@ -35,7 +35,12 @@ private:
};
}
struct VTKM_DEPRECATED(1.8, "Use vtkm::filter::flow::PathParticle.") PathParticle
: vtkm::filter::flow::PathParticle
{
using vtkm::filter::flow::PathParticle::PathParticle;
};
}
} // namespace vtkm::filter::flow
#endif // vtk_m_filter_PathParticle_h
#endif // vtk_m_filter_flow_PathParticle_h

11
vtkm/filter/flow/Pathline.h
View File

@ -8,8 +8,8 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_filter_Pathline_h
#define vtk_m_filter_Pathline_h
#ifndef vtk_m_filter_flow_Pathline_h
#define vtk_m_filter_flow_Pathline_h
#include <vtkm/filter/flow/FlowTypes.h>
#include <vtkm/filter/flow/NewFilterParticleAdvectionUnsteadyState.h>
@ -35,7 +35,12 @@ private:
};
}
struct VTKM_DEPRECATED(1.8, "Use vtkm::filter::flow::Pathline.") Pathline
: vtkm::filter::flow::Pathline
{
using vtkm::filter::flow::Pathline::Pathline;
};
}
} // namespace vtkm::filter::flow
#endif // vtk_m_filter_Pathline_h
#endif // vtk_m_filter_flow_Pathline_h

5
vtkm/filter/flow/StreamSurface.h
View File

@ -58,6 +58,11 @@ private:
};
}
struct VTKM_DEPRECATED(1.8, "Use vtkm::filter::flow::StreamSurface.") StreamSurface
: vtkm::filter::flow::StreamSurface
{
using vtkm::filter::flow::StreamSurface::StreamSurface;
};
}
} // namespace vtkm::filter::flow

5
vtkm/filter/flow/Streamline.h
View File

@ -35,6 +35,11 @@ private:
};
}
struct VTKM_DEPRECATED(1.8, "Use vtkm::filter::flow::Streamline.") Streamline
: vtkm::filter::flow::Streamline
{
using vtkm::filter::flow::Streamline::Streamline;
};
}
} // namespace vtkm::filter::flow

4
vtkm/filter/flow/internal/DataSetIntegrator.h
View File

@ -60,9 +60,11 @@ using DSIHelperInfoType = vtkm::cont::internal::Variant<DSIHelperInfo<vtkm::Part
template <typename Derived>
class DataSetIntegrator
{
protected:
public:
using VelocityFieldNameType = std::string;
using ElectroMagneticFieldNameType = std::pair<std::string, std::string>;
protected:
using FieldNameType =
vtkm::cont::internal::Variant<VelocityFieldNameType, ElectroMagneticFieldNameType>;

107
vtkm/filter/flow/internal/DataSetIntegratorSteadyState.h
View File

@ -55,7 +55,7 @@ protected:
VTKM_CONT void GetVelocityField(
vtkm::worklet::flow::VelocityField<ArrayType>& velocityField) const
{
if (this->FieldName.GetIndex() == this->FieldName.GetIndexOf<VelocityFieldNameType>())
if (this->FieldName.IsType<VelocityFieldNameType>())
{
const auto& fieldNm = this->FieldName.Get<VelocityFieldNameType>();
auto assoc = this->DataSet.GetField(fieldNm).GetAssociation();
@ -68,6 +68,31 @@ protected:
throw vtkm::cont::ErrorFilterExecution("Velocity field vector type not available");
}
template <typename ArrayType>
VTKM_CONT void GetElectroMagneticField(
vtkm::worklet::flow::ElectroMagneticField<ArrayType>& ebField) const
{
if (this->FieldName.IsType<ElectroMagneticFieldNameType>())
{
const auto& fieldNm = this->FieldName.Get<ElectroMagneticFieldNameType>();
const auto& electric = fieldNm.first;
const auto& magnetic = fieldNm.second;
auto eAssoc = this->DataSet.GetField(electric).GetAssociation();
auto bAssoc = this->DataSet.GetField(magnetic).GetAssociation();
if (eAssoc != bAssoc)
{
throw vtkm::cont::ErrorFilterExecution("E and B field need to have same association");
}
ArrayType eField, bField;
vtkm::cont::ArrayCopyShallowIfPossible(this->DataSet.GetField(electric).GetData(), eField);
vtkm::cont::ArrayCopyShallowIfPossible(this->DataSet.GetField(magnetic).GetData(), bField);
ebField = vtkm::worklet::flow::ElectroMagneticField<ArrayType>(eField, bField, eAssoc);
}
else
throw vtkm::cont::ErrorFilterExecution("Electromagnetic field vector type not available");
}
private:
vtkm::cont::DataSet DataSet;
};
@ -76,17 +101,23 @@ private:
namespace internal
{
using ArrayType = vtkm::cont::ArrayHandle<vtkm::Vec3f>;
using VelocityFieldType = vtkm::worklet::flow::VelocityField<ArrayType>;
using SteadyStateGridEvalType = vtkm::worklet::flow::GridEvaluator<VelocityFieldType>;
using VelocityEvalType = vtkm::worklet::flow::GridEvaluator<VelocityFieldType>;
template <typename GridEvalType, typename ParticleType>
class AdvectHelper;
using EBFieldType = vtkm::worklet::flow::ElectroMagneticField<ArrayType>;
using EBEvalType = vtkm::worklet::flow::GridEvaluator<EBFieldType>;
template <typename ParticleType>
class AdvectHelper<SteadyStateGridEvalType, ParticleType>
//template <typename GridEvalType, typename ParticleType>
//class AdvectHelper;
template <typename FieldType, typename ParticleType>
class AdvectHelper //<FieldType, ParticleType>
{
public:
static void Advect(const VelocityFieldType& velField,
using SteadyStateGridEvalType = vtkm::worklet::flow::GridEvaluator<FieldType>;
static void Advect(const FieldType& vecField,
const vtkm::cont::DataSet& ds,
vtkm::cont::ArrayHandle<ParticleType>& seedArray,
vtkm::FloatDefault stepSize,
@ -99,20 +130,20 @@ public:
DoAdvect<vtkm::worklet::flow::ParticleAdvection,
vtkm::worklet::flow::ParticleAdvectionResult,
vtkm::worklet::flow::RK4Integrator>(
velField, ds, seedArray, stepSize, maxSteps, result);
vecField, ds, seedArray, stepSize, maxSteps, result);
}
else if (solverType == IntegrationSolverType::EULER_TYPE)
{
DoAdvect<vtkm::worklet::flow::ParticleAdvection,
vtkm::worklet::flow::ParticleAdvectionResult,
vtkm::worklet::flow::EulerIntegrator>(
velField, ds, seedArray, stepSize, maxSteps, result);
vecField, ds, seedArray, stepSize, maxSteps, result);
}
else
throw vtkm::cont::ErrorFilterExecution("Unsupported Integrator type");
}
static void Advect(const VelocityFieldType& velField,
static void Advect(const FieldType& vecField,
const vtkm::cont::DataSet& ds,
vtkm::cont::ArrayHandle<ParticleType>& seedArray,
vtkm::FloatDefault stepSize,
@ -125,14 +156,14 @@ public:
DoAdvect<vtkm::worklet::flow::Streamline,
vtkm::worklet::flow::StreamlineResult,
vtkm::worklet::flow::RK4Integrator>(
velField, ds, seedArray, stepSize, maxSteps, result);
vecField, ds, seedArray, stepSize, maxSteps, result);
}
else if (solverType == IntegrationSolverType::EULER_TYPE)
{
DoAdvect<vtkm::worklet::flow::Streamline,
vtkm::worklet::flow::StreamlineResult,
vtkm::worklet::flow::EulerIntegrator>(
velField, ds, seedArray, stepSize, maxSteps, result);
vecField, ds, seedArray, stepSize, maxSteps, result);
}
else
throw vtkm::cont::ErrorFilterExecution("Unsupported Integrator type");
@ -143,7 +174,7 @@ public:
class ResultType,
template <typename>
class SolverType>
static void DoAdvect(const VelocityFieldType& velField,
static void DoAdvect(const FieldType& vecField,
const vtkm::cont::DataSet& ds,
vtkm::cont::ArrayHandle<ParticleType>& seedArray,
vtkm::FloatDefault stepSize,
@ -154,7 +185,7 @@ public:
vtkm::worklet::flow::Stepper<SolverType<SteadyStateGridEvalType>, SteadyStateGridEvalType>;
WorkletType worklet;
SteadyStateGridEvalType eval(ds, velField);
SteadyStateGridEvalType eval(ds, vecField);
StepperType stepper(eval, stepSize);
result = worklet.Run(stepper, seedArray, maxSteps);
}
@ -173,23 +204,23 @@ VTKM_CONT inline void DataSetIntegratorSteadyState::DoAdvect(DSIHelperInfo<vtkm:
if (this->VecFieldType == VectorFieldType::VELOCITY_FIELD_TYPE)
{
using FieldType = vtkm::worklet::flow::VelocityField<ArrayType>;
FieldType velField;
this->GetVelocityField(velField);
FieldType vecField;
this->GetVelocityField(vecField);
using AHType = internal::AdvectHelper<internal::SteadyStateGridEvalType, vtkm::Particle>;
using AHType = internal::AdvectHelper<internal::VelocityFieldType, vtkm::Particle>;
if (this->IsParticleAdvectionResult())
{
vtkm::worklet::flow::ParticleAdvectionResult<vtkm::Particle> result;
AHType::Advect(
velField, this->DataSet, seedArray, stepSize, maxSteps, this->SolverType, result);
vecField, this->DataSet, seedArray, stepSize, maxSteps, this->SolverType, result);
this->UpdateResult(result, b);
}
else if (this->IsStreamlineResult())
{
vtkm::worklet::flow::StreamlineResult<vtkm::Particle> result;
AHType::Advect(
velField, this->DataSet, seedArray, stepSize, maxSteps, this->SolverType, result);
vecField, this->DataSet, seedArray, stepSize, maxSteps, this->SolverType, result);
this->UpdateResult(result, b);
}
else
@ -200,10 +231,42 @@ VTKM_CONT inline void DataSetIntegratorSteadyState::DoAdvect(DSIHelperInfo<vtkm:
}
VTKM_CONT inline void DataSetIntegratorSteadyState::DoAdvect(
DSIHelperInfo<vtkm::ChargedParticle>& vtkmNotUsed(b),
vtkm::FloatDefault vtkmNotUsed(stepSize),
vtkm::Id vtkmNotUsed(maxSteps))
DSIHelperInfo<vtkm::ChargedParticle>& b,
vtkm::FloatDefault stepSize,
vtkm::Id maxSteps)
{
using ArrayType = vtkm::cont::ArrayHandle<vtkm::Vec3f>;
auto copyFlag = (this->CopySeedArray ? vtkm::CopyFlag::On : vtkm::CopyFlag::Off);
auto seedArray = vtkm::cont::make_ArrayHandle(b.V, copyFlag);
if (this->VecFieldType == VectorFieldType::ELECTRO_MAGNETIC_FIELD_TYPE)
{
using FieldType = vtkm::worklet::flow::ElectroMagneticField<ArrayType>;
FieldType ebField;
this->GetElectroMagneticField(ebField);
using AHType = internal::AdvectHelper<internal::EBFieldType, vtkm::ChargedParticle>;
if (this->IsParticleAdvectionResult())
{
vtkm::worklet::flow::ParticleAdvectionResult<vtkm::ChargedParticle> result;
AHType::Advect(
ebField, this->DataSet, seedArray, stepSize, maxSteps, this->SolverType, result);
this->UpdateResult(result, b);
}
else if (this->IsStreamlineResult())
{
vtkm::worklet::flow::StreamlineResult<vtkm::ChargedParticle> result;
AHType::Advect(
ebField, this->DataSet, seedArray, stepSize, maxSteps, this->SolverType, result);
this->UpdateResult(result, b);
}
else
throw vtkm::cont::ErrorFilterExecution("Unsupported result type");
}
else
throw vtkm::cont::ErrorFilterExecution("Unsupported vector field type");
}
}

4
vtkm/filter/flow/internal/DataSetIntegratorUnsteadyState.h
View File

@ -254,9 +254,11 @@ VTKM_CONT inline void DataSetIntegratorUnsteadyState::DoAdvect(
vtkm::FloatDefault vtkmNotUsed(stepSize),
vtkm::Id vtkmNotUsed(maxSteps))
{
throw vtkm::cont::ErrorFilterExecution(
"Unsupported operation : charged particles and electromagnetic fielfs currently only supported "
"for steady state");
}
}
}
}

1
vtkm/filter/flow/testing/CMakeLists.txt
View File

@ -10,6 +10,7 @@
set(filter_unit_tests
UnitTestStreamlineFilter.cxx
UnitTestStreamlineFilterWarpX.cxx
UnitTestStreamSurfaceFilter.cxx
)
set(worklet_unit_tests

113
vtkm/filter/flow/testing/UnitTestStreamlineFilterWarpX.cxx Normal file
View File

@ -0,0 +1,113 @@
//============================================================================
// 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 <vtkm/Particle.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/filter/flow/Streamline.h>
#include <vtkm/io/VTKDataSetReader.h>
namespace
{
void GenerateChargedParticles(const vtkm::cont::ArrayHandle<vtkm::Vec3f>& pos,
const vtkm::cont::ArrayHandle<vtkm::Vec3f>& mom,
const vtkm::cont::ArrayHandle<vtkm::Float64>& mass,
const vtkm::cont::ArrayHandle<vtkm::Float64>& charge,
const vtkm::cont::ArrayHandle<vtkm::Float64>& weight,
vtkm::cont::ArrayHandle<vtkm::ChargedParticle>& seeds)
{
auto pPortal = pos.ReadPortal();
auto uPortal = mom.ReadPortal();
auto mPortal = mass.ReadPortal();
auto qPortal = charge.ReadPortal();
auto wPortal = weight.ReadPortal();
auto numValues = pos.GetNumberOfValues();
seeds.Allocate(numValues);
auto sPortal = seeds.WritePortal();
for (vtkm::Id i = 0; i < numValues; i++)
{
vtkm::ChargedParticle electron(
pPortal.Get(i), i, mPortal.Get(i), qPortal.Get(i), wPortal.Get(i), uPortal.Get(i));
sPortal.Set(i, electron);
}
}
void TestStreamlineFilters()
{
std::string particleFile = vtkm::cont::testing::Testing::DataPath("misc/warpXparticles.vtk");
std::string fieldFile = vtkm::cont::testing::Testing::DataPath("misc/warpXfields.vtk");
using SeedsType = vtkm::cont::ArrayHandle<vtkm::ChargedParticle>;
SeedsType seeds;
vtkm::io::VTKDataSetReader seedsReader(particleFile);
vtkm::cont::DataSet seedsData = seedsReader.ReadDataSet();
vtkm::cont::ArrayHandle<vtkm::Vec3f> pos, mom;
vtkm::cont::ArrayHandle<vtkm::Float64> mass, charge, w;
seedsData.GetCoordinateSystem().GetDataAsDefaultFloat().AsArrayHandle(pos);
seedsData.GetField("Momentum").GetDataAsDefaultFloat().AsArrayHandle(mom);
seedsData.GetField("Mass").GetData().AsArrayHandle(mass);
seedsData.GetField("Charge").GetData().AsArrayHandle(charge);
seedsData.GetField("Weighting").GetData().AsArrayHandle(w);
GenerateChargedParticles(pos, mom, mass, charge, w, seeds);
vtkm::io::VTKDataSetReader dataReader(fieldFile);
vtkm::cont::DataSet dataset = dataReader.ReadDataSet();
vtkm::cont::UnknownCellSet cells = dataset.GetCellSet();
vtkm::cont::CoordinateSystem coords = dataset.GetCoordinateSystem();
auto bounds = coords.GetBounds();
std::cout << "Bounds : " << bounds << std::endl;
using Structured3DType = vtkm::cont::CellSetStructured<3>;
Structured3DType castedCells;
cells.AsCellSet(castedCells);
auto dims = castedCells.GetSchedulingRange(vtkm::TopologyElementTagPoint());
vtkm::Vec3f spacing = { static_cast<vtkm::FloatDefault>(bounds.X.Length()) / (dims[0] - 1),
static_cast<vtkm::FloatDefault>(bounds.Y.Length()) / (dims[1] - 1),
static_cast<vtkm::FloatDefault>(bounds.Z.Length()) / (dims[2] - 1) };
std::cout << spacing << std::endl;
constexpr static vtkm::FloatDefault SPEED_OF_LIGHT =
static_cast<vtkm::FloatDefault>(2.99792458e8);
spacing = spacing * spacing;
vtkm::Id steps = 50;
vtkm::FloatDefault length = static_cast<vtkm::FloatDefault>(
1.0 / (SPEED_OF_LIGHT * vtkm::Sqrt(1. / spacing[0] + 1. / spacing[1] + 1. / spacing[2])));
std::cout << "CFL length : " << length << std::endl;
vtkm::filter::flow::Streamline streamline;
streamline.SetStepSize(length);
streamline.SetNumberOfSteps(steps);
streamline.SetSeeds(seeds);
streamline.SetVectorFieldType(vtkm::filter::flow::VectorFieldType::ELECTRO_MAGNETIC_FIELD_TYPE);
streamline.SetEField("E");
streamline.SetBField("B");
auto output = streamline.Execute(dataset);
VTKM_TEST_ASSERT(output.GetNumberOfCoordinateSystems() == 1,
"Wrong number of coordinate systems in the output dataset");
VTKM_TEST_ASSERT(output.GetCoordinateSystem().GetNumberOfPoints() == 2550,
"Wrong number of coordinates");
VTKM_TEST_ASSERT(output.GetCellSet().GetNumberOfCells() == 50, "Wrong number of cells");
}
}
int UnitTestStreamlineFilterWarpX(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(TestStreamlineFilters, argc, argv);
}

16
vtkm/filter/geometry_refinement/Tetrahedralize.cxx
View File

@ -14,17 +14,6 @@
namespace
{
struct DeduceCellSet
{
template <typename CellSetType>
void operator()(const CellSetType& cellset,
vtkm::worklet::Tetrahedralize& worklet,
vtkm::cont::CellSetSingleType<>& outCellSet) const
{
outCellSet = worklet.Run(cellset);
}
};
VTKM_CONT bool DoMapField(vtkm::cont::DataSet& result,
const vtkm::cont::Field& field,
const vtkm::worklet::Tetrahedralize& worklet)
@ -62,11 +51,12 @@ namespace geometry_refinement
{
VTKM_CONT vtkm::cont::DataSet Tetrahedralize::DoExecute(const vtkm::cont::DataSet& input)
{
const vtkm::cont::UnknownCellSet& cells = input.GetCellSet();
const vtkm::cont::UnknownCellSet& inCellSet = input.GetCellSet();
vtkm::cont::CellSetSingleType<> outCellSet;
vtkm::worklet::Tetrahedralize worklet;
vtkm::cont::CastAndCall(cells, DeduceCellSet{}, worklet, outCellSet);
vtkm::cont::CastAndCall(inCellSet,
[&](const auto& concrete) { outCellSet = worklet.Run(concrete); });
auto mapper = [&](auto& result, const auto& f) { DoMapField(result, f, worklet); };
// create the output dataset (without a CoordinateSystem).

41
vtkm/filter/geometry_refinement/Triangulate.cxx
View File

@ -14,41 +14,6 @@
namespace
{
class DeduceCellSetTriangulate
{
vtkm::worklet::Triangulate& Worklet;
vtkm::cont::CellSetSingleType<>& OutCellSet;
public:
DeduceCellSetTriangulate(vtkm::worklet::Triangulate& worklet,
vtkm::cont::CellSetSingleType<>& outCellSet)
: Worklet(worklet)
, OutCellSet(outCellSet)
{
}
template <typename CellSetType>
void operator()(const CellSetType& vtkmNotUsed(cellset)) const
{
}
};
template <>
void DeduceCellSetTriangulate::operator()(const vtkm::cont::CellSetExplicit<>& cellset) const
{
this->OutCellSet = Worklet.Run(cellset);
}
template <>
void DeduceCellSetTriangulate::operator()(const vtkm::cont::CellSetStructured<2>& cellset) const
{
this->OutCellSet = Worklet.Run(cellset);
}
template <>
void DeduceCellSetTriangulate::operator()(const vtkm::cont::CellSetStructured<3>& cellset) const
{
this->OutCellSet = Worklet.Run(cellset);
}
//-----------------------------------------------------------------------------
VTKM_CONT bool DoMapField(vtkm::cont::DataSet& result,
const vtkm::cont::Field& field,
@ -87,13 +52,13 @@ namespace geometry_refinement
{
VTKM_CONT vtkm::cont::DataSet Triangulate::DoExecute(const vtkm::cont::DataSet& input)
{
const vtkm::cont::UnknownCellSet& cells = input.GetCellSet();
const vtkm::cont::UnknownCellSet& inCellSet = input.GetCellSet();
vtkm::cont::CellSetSingleType<> outCellSet;
vtkm::worklet::Triangulate worklet;
DeduceCellSetTriangulate triangulate(worklet, outCellSet);
vtkm::cont::CastAndCall(cells, triangulate);
vtkm::cont::CastAndCall(inCellSet,
[&](const auto& concrete) { outCellSet = worklet.Run(concrete); });
auto mapper = [&](auto& result, const auto& f) { DoMapField(result, f, worklet); };
// create the output dataset (without a CoordinateSystem).

30
vtkm/filter/resampling/worklet/Probe.h
View File

@ -238,36 +238,6 @@ public:
return result;
}
//============================================================================
template <typename T>
class MapCellField : public vtkm::worklet::WorkletMapField
{
public:
T InvalidValue;
MapCellField(const T& invalidValue)
: InvalidValue(invalidValue)
{
}
using ControlSignature = void(FieldIn cellIds, WholeArrayIn inputField, FieldOut result);
using ExecutionSignature = void(_1, _2, _3);
template <typename InputFieldPortalType>
VTKM_EXEC void operator()(vtkm::Id cellId,
const InputFieldPortalType& in,
typename InputFieldPortalType::ValueType& out) const
{
if (cellId != -1)
{
out = in.Get(cellId);
}
else
{
out = this->InvalidValue;
}
}
};
vtkm::cont::ArrayHandle<vtkm::Id> GetCellIds() const { return this->CellIds; }
//============================================================================

31
vtkm/filter/vector_analysis/worklet/SurfaceNormals.h
View File

@ -112,25 +112,6 @@ public:
void SetNormalize(bool value) { this->Normalize = value; }
bool GetNormalize() const { return this->Normalize; }
template <typename CellSetType,
typename CoordsCompType,
typename CoordsStorageType,
typename NormalCompType>
void Run(const CellSetType& cellset,
const vtkm::cont::ArrayHandle<vtkm::Vec<CoordsCompType, 3>, CoordsStorageType>& points,
vtkm::cont::ArrayHandle<vtkm::Vec<NormalCompType, 3>>& normals)
{
if (this->Normalize)
{
vtkm::worklet::DispatcherMapTopology<Worklet<>>().Invoke(cellset, points, normals);
}
else
{
vtkm::worklet::DispatcherMapTopology<Worklet<detail::PassThrough>>().Invoke(
cellset, points, normals);
}
}
template <typename CellSetType, typename PointsType, typename NormalCompType>
void Run(const CellSetType& cellset,
const PointsType& points,
@ -193,18 +174,6 @@ public:
{
vtkm::worklet::DispatcherMapTopology<Worklet>().Invoke(cellset, faceNormals, pointNormals);
}
template <typename CellSetType,
typename FaceNormalTypeList,
typename FaceNormalStorageList,
typename NormalCompType>
void Run(
const CellSetType& cellset,
const vtkm::cont::UncertainArrayHandle<FaceNormalTypeList, FaceNormalStorageList>& faceNormals,
vtkm::cont::ArrayHandle<vtkm::Vec<NormalCompType, 3>>& pointNormals)
{
vtkm::worklet::DispatcherMapTopology<Worklet>().Invoke(cellset, faceNormals, pointNormals);
}
};
}
} // vtkm::worklet

14
vtkm/internal/VariantImpl.h
View File

@ -352,6 +352,14 @@ public:
return (this->Index >= 0) && (this->Index < NumberOfTypes);
}
/// Returns true if this `Variant` stores the given type
///
template <typename T>
VTK_M_DEVICE bool IsType() const
{
return (this->GetIndex() == this->GetIndexOf<T>());
}
Variant() = default;
~Variant() = default;
Variant(const Variant&) = default;
@ -373,7 +381,7 @@ public:
template <typename T>
VTK_M_DEVICE Variant& operator=(const T& src)
{
if (this->GetIndex() == this->GetIndexOf<T>())
if (this->IsType<T>())
{
this->Get<T>() = src;
}
@ -474,14 +482,14 @@ private:
template <typename T>
VTK_M_DEVICE T& GetImpl(std::true_type)
{
VTKM_ASSERT(this->GetIndexOf<T>() == this->GetIndex());
VTKM_ASSERT(this->IsType<T>());
return detail::VariantUnionGet<IndexOf<T>::value>(this->Storage);
}
template <typename T>
VTK_M_DEVICE const T& GetImpl(std::true_type) const
{
VTKM_ASSERT(this->GetIndexOf<T>() == this->GetIndex());
VTKM_ASSERT(this->IsType<T>());
return detail::VariantUnionGet<IndexOf<T>::value>(this->Storage);
}

2
vtkm/io/CMakeLists.txt
View File

@ -48,6 +48,7 @@ set(sources
ImageWriterBase.cxx
ImageWriterPNG.cxx
ImageWriterPNM.cxx
PixelTypes.cxx
VTKDataSetReader.cxx
VTKDataSetReaderBase.cxx
VTKDataSetWriter.cxx
@ -82,6 +83,7 @@ vtkm_library(
TEMPLATE_SOURCES ${template_sources}
)
# CMAKE 3.12 does not know about PRIVATE TARGET_OBJECTS
target_link_libraries(vtkm_io PUBLIC vtkm_cont PRIVATE vtkm_lodepng)
if (VTKm_ENABLE_HDF5_IO)
target_include_directories(vtkm_io PRIVATE $<BUILD_INTERFACE:${HDF5_INCLUDE_DIR}>)

17
vtkm/io/ImageReaderPNG.cxx
View File

@ -12,6 +12,10 @@
#include <vtkm/io/PixelTypes.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <vtkm/thirdparty/lodepng/vtkmlodepng/lodepng.h>
VTKM_THIRDPARTY_POST_INCLUDE
namespace
{
@ -23,12 +27,13 @@ vtkm::io::ImageReaderBase::ColorArrayType ReadFromPNG(const std::string& fileNam
{
unsigned char* imageData;
unsigned uwidth, uheight;
vtkm::png::lodepng_decode_file(&imageData,
&uwidth,
&uheight,
fileName.c_str(),
PixelType::PNG_COLOR_TYPE,
PixelType::BIT_DEPTH);
vtkm::png::lodepng_decode_file(
&imageData,
&uwidth,
&uheight,
fileName.c_str(),
static_cast<vtkm::png::LodePNGColorType>(PixelType::GetColorType()),
PixelType::GetBitDepth());
width = static_cast<vtkm::Id>(uwidth);
height = static_cast<vtkm::Id>(uheight);

13
vtkm/io/ImageWriterPNG.cxx
View File

@ -56,12 +56,13 @@ void ImageWriterPNG::WriteToFile(vtkm::Id width, vtkm::Id height, const ColorArr
}
}
vtkm::png::lodepng_encode_file(this->FileName.c_str(),
imageData.data(),
static_cast<unsigned>(width),
static_cast<unsigned>(height),
PixelType::PNG_COLOR_TYPE,
PixelType::BIT_DEPTH);
vtkm::png::lodepng_encode_file(
this->FileName.c_str(),
imageData.data(),
static_cast<unsigned>(width),
static_cast<unsigned>(height),
static_cast<vtkm::png::LodePNGColorType>(PixelType::GetColorType()),
PixelType::GetBitDepth());
}
}
} // namespace vtkm::io

17
vtkm/source/Source.cxx → vtkm/io/PixelTypes.cxx
View File

@ -8,17 +8,12 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/source/Source.h>
#include <vtkm/io/PixelTypes.h>
namespace vtkm
{
namespace source
{
VTKM_THIRDPARTY_PRE_INCLUDE
#include <vtkm/thirdparty/lodepng/vtkmlodepng/lodepng.h>
VTKM_THIRDPARTY_POST_INCLUDE
//Fix the vtable for Source to be in the vtkm_source library
Source::Source() = default;
Source::~Source() = default;
int vtkm::io::internal::GreyColorType = static_cast<int>(vtkm::png::LodePNGColorType::LCT_GREY);
} // namespace source
} // namespace vtkm
int vtkm::io::internal::RGBColorType = static_cast<int>(vtkm::png::LodePNGColorType::LCT_RGB);

27
vtkm/io/PixelTypes.h
View File

@ -13,15 +13,20 @@
#include <vtkm/Types.h>
#include <vtkm/VecTraits.h>
VTKM_THIRDPARTY_PRE_INCLUDE
#include <vtkm/thirdparty/lodepng/vtkmlodepng/lodepng.h>
VTKM_THIRDPARTY_POST_INCLUDE
namespace vtkm
{
namespace io
{
namespace internal
{
extern int GreyColorType;
extern int RGBColorType;
}
// ----------------------------------------------------------------------
// Define custom SFINAE structures to calculate the VTKM types associated
// with provided BitDepths
@ -75,6 +80,10 @@ public:
static constexpr vtkm::IdComponent MAX_COLOR_VALUE = (1 << BitDepth) - 1;
static constexpr vtkm::IdComponent NUM_CHANNELS = Superclass::NUM_COMPONENTS;
static constexpr vtkm::IdComponent BYTES_PER_PIXEL = NUM_CHANNELS * NUM_BYTES;
static constexpr vtkm::IdComponent GetBitDepth()
{
return BasePixel<BitDepth, Channels>::BIT_DEPTH;
}
using Superclass::Superclass;
BasePixel() = default;
@ -131,8 +140,6 @@ public:
// RGB values are stored in a vtkm::Vec<ComponentType, 3>
using Superclass = BasePixel<BitDepth, 3>;
using ComponentType = typename Superclass::ComponentType;
static constexpr vtkm::png::LodePNGColorType PNG_COLOR_TYPE =
vtkm::png::LodePNGColorType::LCT_RGB;
using Superclass::Superclass;
RGBPixel() = default;
@ -145,6 +152,9 @@ public:
virtual ~RGBPixel() = default;
// Get the implementation specific color type (Gray|RGB) value
static int GetColorType();
ComponentType Diff(const Superclass& pixel) const override;
vtkm::Vec4f_32 ToVec4f() const override;
@ -169,8 +179,6 @@ public:
// in order to simplify the pixel helper functions
using Superclass = BasePixel<BitDepth, 1>;
using ComponentType = typename Superclass::ComponentType;
static constexpr vtkm::png::LodePNGColorType PNG_COLOR_TYPE =
vtkm::png::LodePNGColorType::LCT_GREY;
using Superclass::Superclass;
GreyPixel() = default;
@ -182,6 +190,9 @@ public:
virtual ~GreyPixel() = default;
// Get the implementation specific color type (Gray|RGB) value
static int GetColorType();
ComponentType Diff(const Superclass& pixel) const override;
vtkm::Vec4f_32 ToVec4f() const override;

11
vtkm/io/PixelTypes.hxx
View File

@ -78,6 +78,17 @@ vtkm::Vec4f_32 GreyPixel<B>::ToVec4f() const
1);
}
template <const vtkm::Id B>
int RGBPixel<B>::GetColorType()
{
return internal::RGBColorType;
}
template <const vtkm::Id B>
int GreyPixel<B>::GetColorType()
{
return internal::GreyColorType;
}
} // namespace io
} // namespace vtkm

2
vtkm/io/testing/CMakeLists.txt
View File

@ -16,7 +16,7 @@ set(unit_tests
UnitTestVTKDataSetWriter.cxx
)
set(unit_test_libraries vtkm_lodepng vtkm_io)
set(unit_test_libraries vtkm_io)
if(VTKm_ENABLE_RENDERING)
list(APPEND unit_tests

2
vtkm/io/testing/UnitTestPixelTypes.cxx
View File

@ -106,6 +106,7 @@ void TestGreyPixelConstructors()
"Incorrect Conversion");
VTKM_TEST_ASSERT(vtkm::Vec<vtkm::UInt16, 1>(10) == pixel_7, "Bad Vec4f_32 construction");
VTKM_TEST_ASSERT(GreyPixel<16>::GetBitDepth() == 16, "Bad BitDepth");
VTKM_TEST_ASSERT(GreyPixel<16>::BIT_DEPTH == 16, "Bad BitDepth");
VTKM_TEST_ASSERT(GreyPixel<16>::NUM_BYTES == 2, "Bad NumBytes");
VTKM_TEST_ASSERT(GreyPixel<16>::MAX_COLOR_VALUE == 65535, "Bad NumBytes");
@ -146,6 +147,7 @@ void TestRGBPixelConstructors()
"Incorrect Conversion");
VTKM_TEST_ASSERT(vtkm::Vec<vtkm::UInt16, 3>(10, 10, 10) == pixel_8, "Bad Vec4f_32 construction");
VTKM_TEST_ASSERT(RGBPixel<16>::GetBitDepth() == 16, "Bad BitDepth");
VTKM_TEST_ASSERT(RGBPixel<16>::BIT_DEPTH == 16, "Bad BitDepth");
VTKM_TEST_ASSERT(RGBPixel<16>::NUM_BYTES == 2, "Bad NumBytes");
VTKM_TEST_ASSERT(RGBPixel<16>::MAX_COLOR_VALUE == 65535, "Bad NumBytes");

1
vtkm/source/CMakeLists.txt
View File

@ -20,7 +20,6 @@ set(headers
set(device_sources
Amr.cxx
Oscillator.cxx
Source.cxx
Tangle.cxx
Wavelet.cxx
PerlinNoise.cxx

4
vtkm/source/Oscillator.cxx
View File

@ -139,7 +139,7 @@ private:
vtkm::VecVariable<internal::Oscillator, MAX_OSCILLATORS> PeriodicOscillators;
vtkm::VecVariable<internal::Oscillator, MAX_OSCILLATORS> DampedOscillators;
vtkm::VecVariable<internal::Oscillator, MAX_OSCILLATORS> DecayingOscillators;
vtkm::FloatDefault Time;
vtkm::FloatDefault Time{};
}; // OscillatorSource
} // internal
@ -194,7 +194,7 @@ void Oscillator::AddDecaying(vtkm::FloatDefault x,
//-----------------------------------------------------------------------------
vtkm::cont::DataSet Oscillator::Execute() const
vtkm::cont::DataSet Oscillator::DoExecute() const
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);

10
vtkm/source/Oscillator.h
View File

@ -34,13 +34,13 @@ class VTKM_SOURCE_EXPORT Oscillator final : public vtkm::source::Source
public:
///Construct a Oscillator with Cell Dimensions
VTKM_CONT
Oscillator(vtkm::Id3 dims);
explicit Oscillator(vtkm::Id3 dims);
// We can not declare default destructor here since compiler does not know how
// to create one for the Worklet at this point yet. However, the implementation
// in Oscillator.cxx does have ~Oscillator() = default;
VTKM_CONT
~Oscillator();
~Oscillator() override;
VTKM_CONT
void SetTime(vtkm::FloatDefault time);
@ -69,13 +69,13 @@ public:
vtkm::FloatDefault omega,
vtkm::FloatDefault zeta);
VTKM_CONT vtkm::cont::DataSet Execute() const;
private:
VTKM_CONT vtkm::cont::DataSet DoExecute() const override;
vtkm::Id3 Dims;
std::unique_ptr<internal::OscillatorSource> Worklet;
};
}
}
#endif // vtk_m_source_Oscillator_h
#endif // vtk_m_source_OscillatorSource_h

18
vtkm/source/PerlinNoise.cxx
View File

@ -11,7 +11,7 @@
#include <random>
#include <vtkm/VectorAnalysis.h>
#include <vtkm/filter/FilterField.h>
#include <vtkm/filter/NewFilterField.h>
#include <vtkm/source/PerlinNoise.h>
#include <vtkm/worklet/WorkletMapTopology.h>
@ -124,7 +124,7 @@ struct PerlinNoiseWorklet : public vtkm::worklet::WorkletVisitPointsWithCells
vtkm::Id Repeat;
};
class PerlinNoiseField : public vtkm::filter::FilterField<PerlinNoiseField>
class PerlinNoiseField : public vtkm::filter::NewFilterField
{
public:
VTKM_CONT PerlinNoiseField(vtkm::IdComponent tableSize, vtkm::Id seed)
@ -135,22 +135,16 @@ public:
this->SetUseCoordinateSystemAsField(true);
}
template <typename FieldType, typename DerivedPolicy>
VTKM_CONT vtkm::cont::DataSet DoExecute(
const vtkm::cont::DataSet& input,
const FieldType& vtkmNotUsed(field),
const vtkm::filter::FieldMetadata& fieldMetadata,
vtkm::filter::PolicyBase<DerivedPolicy> vtkmNotUsed(policy))
private:
VTKM_CONT vtkm::cont::DataSet DoExecute(const vtkm::cont::DataSet& input) override
{
vtkm::cont::ArrayHandle<vtkm::FloatDefault> noise;
PerlinNoiseWorklet worklet{ this->TableSize };
this->Invoke(
worklet, input.GetCellSet(), input.GetCoordinateSystem(), this->Permutations, noise);
return vtkm::filter::CreateResult(input, noise, this->GetOutputFieldName(), fieldMetadata);
return this->CreateResultFieldPoint(input, this->GetOutputFieldName(), noise);
}
private:
VTKM_CONT void GeneratePermutations()
{
std::mt19937_64 rng;
@ -206,7 +200,7 @@ PerlinNoise::PerlinNoise(vtkm::Id3 dims, vtkm::Vec3f origin, vtkm::IdComponent s
{
}
vtkm::cont::DataSet PerlinNoise::Execute() const
vtkm::cont::DataSet PerlinNoise::DoExecute() const
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);

4
vtkm/source/PerlinNoise.h
View File

@ -39,9 +39,9 @@ public:
void SetSeed(vtkm::IdComponent seed) { this->Seed = seed; }
vtkm::cont::DataSet Execute() const override;
private:
vtkm::cont::DataSet DoExecute() const override;
vtkm::Id3 Dims;
vtkm::Vec3f Origin;
vtkm::IdComponent Seed;

10
vtkm/source/Source.h
View File

@ -14,6 +14,7 @@
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/source/vtkm_source_export.h>
namespace vtkm
{
namespace source
@ -23,14 +24,13 @@ class VTKM_SOURCE_EXPORT Source
{
public:
VTKM_CONT
Source();
virtual ~Source() = default;
VTKM_CONT
virtual ~Source();
virtual vtkm::cont::DataSet Execute() const = 0;
vtkm::cont::DataSet Execute() const { return this->DoExecute(); }
protected:
virtual vtkm::cont::DataSet DoExecute() const = 0;
vtkm::cont::Invoker Invoke;
};

2
vtkm/source/Tangle.cxx
View File

@ -60,7 +60,7 @@ public:
};
} // namespace tangle
vtkm::cont::DataSet Tangle::Execute() const
vtkm::cont::DataSet Tangle::DoExecute() const
{
VTKM_LOG_SCOPE_FUNCTION(vtkm::cont::LogLevel::Perf);

4
vtkm/source/Tangle.h
View File

@ -42,9 +42,9 @@ public:
{
}
vtkm::cont::DataSet Execute() const;
private:
vtkm::cont::DataSet DoExecute() const override;
vtkm::Id3 Dims;
};
} //namespace source

Some files were not shown because too many files have changed in this diff Show More