Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-10-05 01:49:02 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

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

Browse Source
This commit is contained in:
Dave Pugmire 2023-03-15 09:23:03 -04:00
commit 6db8c316ad
159 changed files with 4681 additions and 2023 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
.gitlab-ci.yml
View File

@ -38,7 +38,7 @@
- .docker_image
.ubuntu1804_cuda_kokkos: &ubuntu1804_cuda_kokkos
image: "kitware/vtkm:ci-ubuntu1804_cuda11_kokkos-20220407"
image: "kitware/vtkm:ci-ubuntu1804_cuda11_kokkos-20230125"
extends:
- .docker_image
@ -53,12 +53,12 @@
- .docker_image
.ubuntu2004_kokkos: &ubuntu2004_kokkos
image: "kitware/vtkm:ci-ubuntu2004_kokkos-20210916"
image: "kitware/vtkm:ci-ubuntu2004_kokkos-20230125"
extends:
- .docker_image
.ubuntu2004_hip_kokkos: &ubuntu2004_hip_kokkos
image: "kitware/vtkm:ci-ubuntu2004_hip_kokkos-20220620"
image: "kitware/vtkm:ci-ubuntu2004_hip_kokkos-20230220"
extends:
- .docker_image
@ -96,7 +96,7 @@
when: on_success
- when: never
.run_spock_ci: &run_spock_ci
.run_crusher_ci: &run_crusher_ci
rules:
- if: '$CI_PROJECT_PATH == "ci/csc331_crusher/dev/vtk-m"'
when: on_success
@ -214,7 +214,7 @@ stages:
include:
- local: '/.gitlab/ci/ascent.yml'
- local: '/.gitlab/ci/spock.yml'
- local: '/.gitlab/ci/crusher.yml'
- local: '/.gitlab/ci/centos7.yml'
- local: '/.gitlab/ci/centos8.yml'
- local: '/.gitlab/ci/doxygen.yml'

41
.gitlab/ci/ascent.yml
View File

@ -2,7 +2,7 @@
.ascent_gcc_cuda:
variables:
CCACHE_BASEDIR: "/gpfs/wolf/"
CCACHE_DIR: "/gpfs/wolf/proj-shared/csc331/vtk-m/ci/ccache/"
CCACHE_DIR: "/gpfs/wolf/csc331/scratch/vbolea/ci/ccache"
# -isystem= is not affected by CCACHE_BASEDIR, thus we must ignore it
CCACHE_IGNOREOPTIONS: "-isystem=*"
CCACHE_NOHASHDIR: "true"
@ -12,7 +12,7 @@
CMAKE_BUILD_TYPE: "RelWithDebInfo"
CMAKE_GENERATOR: "Ninja"
CUDAHOSTCXX: "g++"
CUSTOM_CI_BUILDS_DIR: "/gpfs/wolf/proj-shared/csc331/vtk-m/ci/runtime"
CUSTOM_CI_BUILDS_DIR: "/gpfs/wolf/csc331/scratch/vbolea/ci/vtk-m"
VTKM_SETTINGS: cuda+ascent+ccache
JOB_MODULES: >-
@ -30,14 +30,18 @@
module purge
module load ${JOB_MODULES}
module list
export PATH="/gpfs/wolf/proj-shared/csc331/vtk-m/ci/utils:$PATH"
export PATH="/gpfs/wolf/csc331/scratch/vbolea/ci/utils:$PATH"
build:ascent_gcc_cuda:
stage: build
tags: [olcf, ascent, nobatch]
tags: [olcf, ascent, batch]
extends:
- .ascent_gcc_cuda
- .run_ascent_ci
- .cmake_build_artifacts
variables:
SCHEDULER_PARAMETERS: -P CSC331 -W 2:00 -nnodes 1 -alloc_flags smt1
timeout: 125 minutes
before_script:
- *setup_env_ecpci
- ccache -z
@ -46,14 +50,16 @@ build:ascent_gcc_cuda:
- git-lfs install
- git-lfs pull lfs
script:
- CTEST_MAX_PARALLELISM=32 cmake -V -P .gitlab/ci/config/gitlab_ci_setup.cmake
# Each Ascent (Summit) node has 172 threads (43 cores). SMT1 is needed to
# avoid L1 cache pollution among different processes. Thus, using 40 cores
# seems a reasonable choice which leaves a couple of cores for system processes.
- CTEST_MAX_PARALLELISM=40 cmake -V -P .gitlab/ci/config/gitlab_ci_setup.cmake
- ctest -VV -S .gitlab/ci/ctest_configure.cmake
artifacts:
expire_in: 24 hours
when: always
paths:
- build/
timeout: 10 minutes
- GITLAB_CI_EMULATION=1 jsrun -n1 -a1 -g1 -c40 -bpacked:40 ctest -VV -S .gitlab/ci/ctest_build.cmake
after_script:
- *setup_env_ecpci
- ccache -s
- ctest -VV -S .gitlab/ci/ctest_submit_build.cmake
test:ascent_gcc_cuda:
stage: test
@ -67,17 +73,18 @@ test:ascent_gcc_cuda:
dependencies:
- build:ascent_gcc_cuda
variables:
SCHEDULER_PARAMETERS: -P CSC331 -W 2:00 -nnodes 1 -alloc_flags gpudefault
# For tests we want to use a small number of proccesses, for some reason
# a higher parallelism number tend to results in test malfunctions.
CTEST_MAX_PARALLELISM: 4
# We need this to skip ctest_submit from being run inside a jsrun job
GITLAB_CI_EMULATION: 1
SCHEDULER_PARAMETERS: -P CSC331 -W 1:00 -nnodes 1 -alloc_flags gpudefault
timeout: 65 minutes
before_script:
- *setup_env_ecpci
script:
- jsrun -n1 -a1 -g1 -c42 ctest -VV -S .gitlab/ci/ctest_build.cmake
- CTEST_MAX_PARALLELISM=4 jsrun -n1 -a1 -g1 -c8 ctest -VV -S .gitlab/ci/ctest_test.cmake
- cmake -V -P .gitlab/ci/config/gitlab_ci_setup.cmake
- jsrun -n1 -a1 -g1 -c1 ctest -VV -S .gitlab/ci/ctest_test.cmake
after_script:
- *setup_env_ecpci
- ccache -s
- ctest -VV -S .gitlab/ci/ctest_submit_build.cmake
- ctest -VV -S .gitlab/ci/ctest_submit_test.cmake
timeout: 120 minutes

2
.gitlab/ci/centos8.yml
View File

@ -45,7 +45,7 @@ build:centos8:
- build
- vtkm
- docker
- linux
- linux-x86_64
extends:
- .centos8
- .cmake_build_linux

30
.gitlab/ci/config/kokkos.sh Executable file
View File

@ -0,0 +1,30 @@
#!/usr/bin/env bash
set -x
WORKDIR="$1"
VERSION="$2"
shift 2
if [ ! -d "$WORKDIR" ] || [ -z "$VERSION" ]
then
echo "[E] missing args: Invoke as .gitlab/ci/config/kokkos.sh <WORKDIR> <VERSION> [extra_args]"
exit 1
fi
# Build and install Kokkos
curl -L "https://github.com/kokkos/kokkos/archive/refs/tags/$VERSION.tar.gz" \
| tar -C "$WORKDIR" -xzf -
cmake -S "$WORKDIR/kokkos-$VERSION" -B "$WORKDIR/kokkos_build" \
"-DCMAKE_BUILD_TYPE:STRING=release" \
"-DCMAKE_CXX_COMPILER_LAUNCHER=ccache" \
"-DCMAKE_CXX_STANDARD:STRING=14" \
"-DCMAKE_POSITION_INDEPENDENT_CODE:BOOL=ON" \
"-DKokkos_ENABLE_HIP:BOOL=ON" \
"-DKokkos_ENABLE_HIP_RELOCATABLE_DEVICE_CODE:BOOL=OFF" \
"-DKokkos_ENABLE_SERIAL:BOOL=ON" \
$*
cmake --build "$WORKDIR/kokkos_build"
cmake --install "$WORKDIR/kokkos_build"

55
.gitlab/ci/spock.yml → .gitlab/ci/crusher.yml
View File

@ -1,9 +1,9 @@
# Ad-hoc build that runs in the ECP Hardware, concretely in OLCF Spock.
.spock_gcc_hip:
.crusher_gcc_hip:
variables:
CCACHE_BASEDIR: "/gpfs/alpine/world-shared/csc331/"
CCACHE_DIR: "/gpfs/alpine/world-shared/csc331/vtk-m/ci/ccache"
CUSTOM_CI_BUILDS_DIR: "/gpfs/alpine/world-shared/csc331/vtk-m/ci/runtime"
CCACHE_BASEDIR: "/gpfs/alpine/csc331/scratch/"
CCACHE_DIR: "/gpfs/alpine/csc331/scratch/vbolea/ci/vtk-m/ccache"
CUSTOM_CI_BUILDS_DIR: "/gpfs/alpine/csc331/scratch/vbolea/ci/vtk-m/runtime"
# -isystem= is not affected by CCACHE_BASEDIR, thus we must ignore it
CCACHE_IGNOREOPTIONS: "-isystem=*"
@ -12,35 +12,25 @@
CMAKE_BUILD_TYPE: "RelWithDebInfo"
CMAKE_GENERATOR: "Ninja"
CMAKE_PREFIX_PATH: "$CI_BUILDS_DIR/kokkos_install"
# This is needed for the smoke_test, while we use rocm 5 to build VTK-m the
# smoke_test needs 4.5 since Kokkos is built agains rocm 4.5
LD_LIBRARY_PATH: "/opt/rocm-4.5.0/lib:${LD_LIBRARY_PATH}"
LIBRARY_PATH: "/opt/rocm-4.5.0/lib:${LIBRARY_PATH}"
KOKKOS_OPTS: >-
-DCMAKE_INSTALL_PREFIX:PATH=$CI_BUILDS_DIR/kokkos_install
-DCMAKE_CXX_COMPILER:FILEPATH=/opt/rocm-4.5.0/hip/bin/hipcc
-DKokkos_ARCH_VEGA90A:BOOL=ON
# While Kokkos and VTK-m uses ROCm 4.5.0 runtime/sdk, we need to build
# VTK-m with HIPCC from ROCM 5
CMAKE_HIP_COMPILER: "/opt/rocm-default/llvm/bin/clang++"
Kokkos_CXX_COMPILER: "/opt/rocm-default/llvm/bin/clang++"
CMAKE_HIP_ARCHITECTURES: "gfx908"
CC: gcc
CXX: g++
# DefApps/default;craype;rocm;gcc should be loaded first
JOB_MODULES: >-
DefApps/default
craype-accel-amd-gfx90a
rocm/5
rocm/4.5.0
gcc/10
cmake/3.22
git
git-lfs
kokkos/3.5.00
lsf-tools
ninja
spectrum-mpi
zstd
VTKM_SETTINGS: kokkos+hip+spock+ccache+no_rendering
VTKM_SETTINGS: kokkos+hip+gfx90a+crusher+ccache+no_rendering
interruptible: true
.setup_env_ecpci: &setup_env_ecpci |
@ -49,12 +39,12 @@
module list
export PATH="${CCACHE_INSTALL_DIR}/ccache:$PATH"
build:spock_gcc_hip:
build:crusher_gcc_hip:
stage: build
tags: [spock, shell]
tags: [crusher, shell]
extends:
- .spock_gcc_hip
- .run_spock_ci
- .crusher_gcc_hip
- .run_crusher_ci
before_script:
- *setup_env_ecpci
- mkdir -p "$CCACHE_INSTALL_DIR"
@ -62,6 +52,7 @@ build:spock_gcc_hip:
- cmake -VV -P .gitlab/ci/config/ccache.cmake
- ccache -z
- ccache -s
- .gitlab/ci/config/kokkos.sh "$CI_BUILDS_DIR" "3.7.01" $KOKKOS_OPTS
- git remote add lfs https://gitlab.kitware.com/vtk/vtk-m.git
- git fetch lfs
@ -77,18 +68,18 @@ build:spock_gcc_hip:
- build/
- .gitlab/ccache/ccache
test:spock_gcc_hip:
test:crusher_gcc_hip:
stage: test
tags: [spock, slurm]
tags: [crusher, slurm]
extends:
- .spock_gcc_hip
- .run_spock_ci
- .crusher_gcc_hip
- .run_crusher_ci
needs:
- build:spock_gcc_hip
- build:crusher_gcc_hip
dependencies:
- build:spock_gcc_hip
- build:crusher_gcc_hip
variables:
SCHEDULER_PARAMETERS: "-ACSC331 -pecp -t120 --nice=0 -c32 --gpus=4 -N 1"
SCHEDULER_PARAMETERS: "-ACSC331 -pbatch -t120 --nice=0 -c32 --gpus=4 -N 1"
# We need this to skip ctest_submit from being run inside a jsrun job
GITLAB_CI_EMULATION: 1
# Tests errors to address due to different env in Spock

9
.gitlab/ci/docker/ubuntu1804/kokkos-cuda/Dockerfile
View File

@ -25,13 +25,14 @@ RUN mkdir /opt/cmake/ && \
ENV PATH "/opt/cmake/bin:${PATH}"
# Build and install Kokkos
ARG KOKKOS_VERSION=3.7.01
RUN mkdir -p /opt/kokkos/build && \
cd /opt/kokkos/build && \
curl -L https://github.com/kokkos/kokkos/archive/refs/tags/3.4.01.tar.gz > kokkos-3.4.01.tar.gz && \
tar -xf kokkos-3.4.01.tar.gz && \
curl -L https://github.com/kokkos/kokkos/archive/refs/tags/$KOKKOS_VERSION.tar.gz > kokkos-$KOKKOS_VERSION.tar.gz && \
tar -xf kokkos-$KOKKOS_VERSION.tar.gz && \
mkdir bld && cd bld && \
CXX=/opt/kokkos/build/kokkos-3.4.01/bin/nvcc_wrapper \
cmake -B . -S ../kokkos-3.4.01 \
CXX=/opt/kokkos/build/kokkos-$KOKKOS_VERSION/bin/nvcc_wrapper \
cmake -B . -S ../kokkos-$KOKKOS_VERSION \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=/opt/kokkos \
-DCMAKE_CXX_FLAGS=-fPIC \

3
.gitlab/ci/docker/ubuntu2004/kokkos-hip/Dockerfile
View File

@ -14,6 +14,7 @@ RUN apt update && \
ninja-build \
rsync \
ssh \
rocthrust-dev \
&& \
apt clean
@ -44,7 +45,7 @@ ENV CMAKE_PREFIX_PATH "/opt/rocm/lib/cmake:/opt/rocm/lib:${CMAKE_PREFIX_PATH}"
ENV CMAKE_GENERATOR "Ninja"
# Build and install Kokkos
ARG KOKKOS_VERSION=3.6.00
ARG KOKKOS_VERSION=3.7.01
COPY kokkos_cmake_config.cmake kokkos_cmake_config.cmake
RUN curl -L https://github.com/kokkos/kokkos/archive/refs/tags/$KOKKOS_VERSION.tar.gz | tar -xzf - && \
cmake -S kokkos-$KOKKOS_VERSION -B build -C kokkos_cmake_config.cmake && \

7
.gitlab/ci/docker/ubuntu2004/kokkos/Dockerfile
View File

@ -31,11 +31,12 @@ RUN mkdir /opt/cmake/ && \
ENV PATH "${PATH}:/opt/cmake/bin"
# Build and install Kokkos
ARG KOKKOS_VERSION=3.7.01
RUN mkdir -p /opt/kokkos/build && \
cd /opt/kokkos/build && \
curl -L https://github.com/kokkos/kokkos/archive/refs/tags/3.4.01.tar.gz > kokkos-3.4.01.tar.gz && \
tar -xf kokkos-3.4.01.tar.gz && \
curl -L https://github.com/kokkos/kokkos/archive/refs/tags/$KOKKOS_VERSION.tar.gz > kokkos-$KOKKOS_VERSION.tar.gz && \
tar -xf kokkos-$KOKKOS_VERSION.tar.gz && \
mkdir bld && cd bld && \
cmake -GNinja -DCMAKE_INSTALL_PREFIX=/opt/kokkos -DCMAKE_CXX_FLAGS=-fPIC -DKokkos_ENABLE_SERIAL=ON ../kokkos-3.4.01 &&\
cmake -GNinja -DCMAKE_INSTALL_PREFIX=/opt/kokkos -DCMAKE_CXX_FLAGS=-fPIC -DKokkos_ENABLE_SERIAL=ON ../kokkos-$KOKKOS_VERSION &&\
ninja all && \
ninja install

10
CMake/VTKmDeviceAdapters.cmake
View File

@ -322,7 +322,7 @@ endfunction()
if(VTKm_ENABLE_KOKKOS AND NOT TARGET vtkm_kokkos)
cmake_minimum_required(VERSION 3.13 FATAL_ERROR)
find_package(Kokkos REQUIRED)
find_package(Kokkos 3.7 REQUIRED)
# We must empty this property for every kokkos backend device since it
# contains a generator expresion which breaks some of our users builds.
@ -357,6 +357,14 @@ if(VTKm_ENABLE_KOKKOS AND NOT TARGET vtkm_kokkos)
add_library(vtkm_kokkos_hip INTERFACE)
set_property(TARGET vtkm_kokkos_hip PROPERTY EXPORT_NAME kokkos_hip)
install(TARGETS vtkm_kokkos_hip EXPORT ${VTKm_EXPORT_NAME})
# Make sure rocthrust is available if requested
if(VTKm_ENABLE_KOKKOS_THRUST)
find_package(rocthrust)
if(NOT rocthrust_FOUND)
message(FATAL_ERROR "rocthrust not found. Please set VTKm_ENABLE_KOKKOS_THRUST to OFF.")
endif()
endif()
endif()
add_library(vtkm_kokkos INTERFACE IMPORTED GLOBAL)

41
CMake/testing/VTKmPerformanceTest.cmake
View File

@ -79,6 +79,24 @@ function(add_benchmark_test benchmark)
set(test_name "PerformanceTest${VTKm_PERF_NAME}")
###TEST INVOKATIONS##########################################################
if (NOT TEST PerformanceTestFetch)
add_test(NAME "PerformanceTestFetch"
COMMAND ${CMAKE_COMMAND}
"-DVTKm_PERF_REPO=${VTKm_PERF_REPO}"
"-DVTKm_SOURCE_DIR=${VTKm_SOURCE_DIR}"
"-DVTKm_PERF_REMOTE_URL=${VTKm_PERF_REMOTE_URL}"
-P "${VTKm_SOURCE_DIR}/CMake/testing/VTKmPerformanceTestFetch.cmake"
)
set_property(TEST PerformanceTestFetch PROPERTY FIXTURES_SETUP "FixturePerformanceTestSetup")
endif()
if (NOT TEST PerformanceTestCleanUp)
add_test(NAME "PerformanceTestCleanUp"
COMMAND ${CMAKE_COMMAND} -E rm -rf "${VTKm_PERF_REPO}"
)
set_property(TEST PerformanceTestCleanUp PROPERTY FIXTURES_CLEANUP "FixturePerformanceTestCleanUp")
endif()
add_test(NAME "${test_name}Run"
COMMAND ${CMAKE_COMMAND}
"-DVTKm_PERF_BENCH_DEVICE=Any"
@ -93,14 +111,6 @@ function(add_benchmark_test benchmark)
-P "${VTKm_SOURCE_DIR}/CMake/testing/VTKmPerformanceTestRun.cmake"
)
add_test(NAME "${test_name}Fetch"
COMMAND ${CMAKE_COMMAND}
"-DVTKm_PERF_REPO=${VTKm_PERF_REPO}"
"-DVTKm_SOURCE_DIR=${VTKm_SOURCE_DIR}"
"-DVTKm_PERF_REMOTE_URL=${VTKm_PERF_REMOTE_URL}"
-P "${VTKm_SOURCE_DIR}/CMake/testing/VTKmPerformanceTestFetch.cmake"
)
add_test(NAME "${test_name}Upload"
COMMAND ${CMAKE_COMMAND}
"-DVTKm_PERF_REPO=${VTKm_PERF_REPO}"
@ -122,22 +132,21 @@ function(add_benchmark_test benchmark)
-P "${VTKm_SOURCE_DIR}/CMake/testing/VTKmPerformanceTestReport.cmake"
)
add_test(NAME "${test_name}CleanUp"
COMMAND ${CMAKE_COMMAND} -E rm -rf "${VTKm_PERF_REPO}"
)
###TEST PROPERTIES###########################################################
set_property(TEST ${test_name}Upload PROPERTY DEPENDS ${test_name}Report)
set_property(TEST ${test_name}Report PROPERTY DEPENDS ${test_name}Run)
set_property(TEST ${test_name}Report PROPERTY FIXTURES_REQUIRED "FixturePerformanceTestSetup")
set_property(TEST ${test_name}Upload PROPERTY FIXTURES_REQUIRED "FixturePerformanceTestCleanUp")
set_tests_properties("${test_name}Report" "${test_name}Upload"
PROPERTIES
FIXTURE_REQUIRED "${test_name}Run;${test_name}Fetch"
FIXTURE_CLEANUP "${test_name}CleanUp"
REQUIRED_FILES "${VTKm_PERF_COMPARE_JSON}")
set_tests_properties("${test_name}Run"
"${test_name}Report"
"${test_name}Upload"
"${test_name}Fetch"
"${test_name}CleanUp"
"PerformanceTestFetch"
"PerformanceTestCleanUp"
PROPERTIES RUN_SERIAL ON)
set_tests_properties(${test_name}Run PROPERTIES TIMEOUT 1800)

9
CMake/testing/VTKmPerformanceTestFetch.cmake
View File

@ -10,8 +10,17 @@
include(${VTKm_SOURCE_DIR}/CMake/testing/VTKmPerformanceTestLib.cmake)
REQUIRE_FLAG("VTKm_SOURCE_DIR")
REQUIRE_FLAG_MUTABLE("VTKm_PERF_REPO")
REQUIRE_FLAG_MUTABLE("VTKm_PERF_REMOTE_URL")
set(upstream_url "https://gitlab.kitware.com/vtk/vtk-m.git")
file(REMOVE_RECURSE vtk-m-benchmark-records)
execute(COMMAND /usr/bin/git clone -b records ${VTKm_PERF_REMOTE_URL} ${VTKm_PERF_REPO})
# Fetch VTK-m main git repo objects, this is needed to ensure that when running the CI
# from a fork project of VTK-m it will have access to the latest git commits in
# the upstream vtk-m git repo.
execute(COMMAND /usr/bin/git -C ${VTKm_SOURCE_DIR} remote add upstream ${upstream_url})
execute(COMMAND /usr/bin/git -C ${VTKm_SOURCE_DIR} fetch upstream)

2
CMake/testing/VTKmPerformanceTestReport.cmake
View File

@ -19,7 +19,7 @@ REQUIRE_FLAG_MUTABLE("VTKm_PERF_ALPHA")
REQUIRE_FLAG_MUTABLE("VTKm_PERF_DIST")
###FIND MOST RECENT BASELINE####################################################
execute(COMMAND /usr/bin/git -C "${VTKm_SOURCE_DIR}" merge-base origin/master @
execute(COMMAND /usr/bin/git -C "${VTKm_SOURCE_DIR}" merge-base upstream/master @
OUTPUT_VARIABLE GIT_BASE_COMMIT)
string(STRIP "${GIT_BASE_COMMIT}" GIT_BASE_COMMIT)

5
CMakeLists.txt
View File

@ -227,6 +227,11 @@ include(VTKmBuildType)
# Include the vtk-m wrappers
include(VTKmWrappers)
# By default: Set VTKm_ENABLE_KOKKOS_THRUST to ON if VTKm_ENABLE_KOKKOS is ON, otherwise
# disable it (or if the user explicitly turns this option OFF)
cmake_dependent_option(VTKm_ENABLE_KOKKOS_THRUST "Enable Kokkos thrust support (only valid with CUDA and HIP)"
ON "VTKm_ENABLE_KOKKOS;Kokkos_ENABLE_CUDA OR Kokkos_ENABLE_HIP" OFF)
# Create vtkm_compiler_flags library. This is an interface library that
# holds all the C++ compiler flags that are needed for consumers and
# when building VTK-m.

2
LICENSE.txt
View File

@ -1,7 +1,7 @@
VTKm License Version 2.0
========================================================================
Copyright (c) 2014-2022
Copyright (c) 2014-2023
Kitware Inc.,
National Technology & Engineering Solutions of Sandia, LLC (NTESS),
UT-Battelle, LLC.,

7
README.md
View File

@ -1,3 +1,6 @@
![GitLab tag](https://img.shields.io/gitlab/v/tag/vtk/vtk-m?color=red&gitlab_url=https%3A%2F%2Fgitlab.kitware.com&include_prereleases&sort=semver)
![Spack version](https://img.shields.io/spack/v/vtk-m.svg)
# VTK-m #
VTK-m is a toolkit of scientific visualization algorithms for emerging
@ -64,7 +67,7 @@ effort.
VTK-m Requires:
+ C++11 Compiler. VTK-m has been confirmed to work with the following
+ C++14 Compiler. VTK-m has been confirmed to work with the following
+ GCC 5.4+
+ Clang 5.0+
+ XCode 5.0+
@ -76,6 +79,8 @@ VTK-m Requires:
Optional dependencies are:
+ Kokkos Device Adapter
+ [Kokkos](https://kokkos.github.io/) 3.7+
+ CUDA Device Adapter
+ [Cuda Toolkit 9.2, >= 10.2](https://developer.nvidia.com/cuda-toolkit)
+ Note CUDA >= 10.2 is required on Windows

49
benchmarking/BenchmarkInSitu.cxx
View File

@ -982,6 +982,55 @@ void ParseBenchmarkOptions(int& argc, char** argv)
std::cerr << "Using data set dimensions = " << DataSetDim << std::endl;
std::cerr << "Using image size = " << ImageSize << "x" << ImageSize << std::endl;
// Now go back through the arg list and remove anything that is not in the list of
// unknown options or non-option arguments.
int destArg = 1;
// This is copy/pasted from vtkm::cont::Initialize(), should probably be abstracted eventually:
for (int srcArg = 1; srcArg < argc; ++srcArg)
{
std::string thisArg{ argv[srcArg] };
bool copyArg = false;
// Special case: "--" gets removed by optionparser but should be passed.
if (thisArg == "--")
{
copyArg = true;
}
for (const option::Option* opt = options[UNKNOWN]; !copyArg && opt != nullptr;
opt = opt->next())
{
if (thisArg == opt->name)
{
copyArg = true;
}
if ((opt->arg != nullptr) && (thisArg == opt->arg))
{
copyArg = true;
}
// Special case: optionparser sometimes removes a single "-" from an option
if (thisArg.substr(1) == opt->name)
{
copyArg = true;
}
}
for (int nonOpt = 0; !copyArg && nonOpt < commandLineParse.nonOptionsCount(); ++nonOpt)
{
if (thisArg == commandLineParse.nonOption(nonOpt))
{
copyArg = true;
}
}
if (copyArg)
{
if (destArg != srcArg)
{
argv[destArg] = argv[srcArg];
}
++destArg;
}
}
argc = destArg;
}
// Adding a const char* or std::string to a vector of char* is harder than it sounds.

3
data/data/curvilinear/kitchen.vtk Normal file
View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:29d6de0eb33dfb792405d3903b1ede382cac417c29bb5d7e17e60be384db42cc
size 1940269

3
data/data/rectilinear/fishtank.vtk
View File

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

3
data/data/rectilinear/fishtank_double_ascii.vtk
View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:2bb3d36ea5ecef5e7ef1057d0dddebbc590424915083091ead3dac2928000524
size 2904465

3
data/data/rectilinear/fishtank_double_big_endian.vtk
View File

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

3
data/data/rectilinear/fusion.vtk
View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:2cbdf56fd5445ddc5b6bc05507b8825fb8d74fe1ccce894bde03e5ff2ecf5fb6
size 525141

3
data/data/rectilinear/magField.vtk
View File

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

3
data/data/uniform/rotate-vectors.vtk Normal file
View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:13241631e0d65d17d7bad2be44b52c0352797135fb585647f5c821e439ff9817
size 862894

2
docs/CodingConventions.md
View File

@ -181,7 +181,7 @@ for (auto v : vector)
avoided in class, method, and function scopes (fully qualified
namespace references are preferred).
+ All code must be valid by the C++11 specifications. It must also
+ All code must be valid by the C++14 specifications. It must also
compile with Microsoft Visual Studio 2015.
+ New code must include regression tests that will run on the dashboards.

4
docs/NewRelease.md.tmpl
View File

@ -40,6 +40,8 @@ $endif\
## Create update branch
- [ ] Freeze the release branch (In Gitlab VTK-m page)
- Settings/Repository/Protected Branches: Release; "allowed to push:No one"
- [ ] Create update branch `git checkout -b update-to-$(VERSION)`
$if(PATCH == 0 and RC == "-rc1")\
- [ ] Bring as a second parent the history of master (Solve conflicts always
@ -104,6 +106,8 @@ $endif\
- [ ] `Do: merge`
- Push tags
- [ ] `git push origin v$(VERSION)$(RC)`
- [ ] Unfreeze the release branch (In Gitlab VTK-m page)
- Settings/Repository/Protected Branches: Release; "allowed to push: Maintainers"
## Update Spack
- [ ] Update Spack package: https://github.com/spack/spack/blob/develop/var/spack/repos/builtin/packages/vtk-m/package.py

10
docs/ReleaseRoadmap.md
View File

@ -1,13 +1,13 @@
# Minor Release Roadmap
# Release Roadmap
| Version | Date | Delay (days) | Life-cycle (*planned) | End of Support |
| --------- | ------------ | ------- | ----------- | ---------------- |
| 1.7.0 | 2021-12-01 | +8 | Long Term | 2022-12-01 |
| 1.8.0 | 2022-06-01 | +14 | Long Term | 2023-06-01 |
| 1.9.0 | 2022-09-01 | +41 | Short Term | next release |
| 2.0.0 | 2022-12-01 | | Long Term* | TBD |
| 2.1.0 | 2023-03-01 | | Short Term* | TBD |
| 2.2.0 | 2023-06-01 | | Long Term* | TBD |
| 1.9.0 | 2022-09-01 | +41 | Short Term | 2023-06-01 |
| 2.0.0 | 2022-12-01 | +52 | Long Term* | 2023-12-01 |
| 2.1.0 | 2023-06-01 | | Short Term* | TBD |
| 2.2.0 | 2023-09-01 | | Long Term* | TBD |
## Legend

75
docs/changelog/2.0/release-notes.md
View File

@ -21,10 +21,13 @@ VTK-m 2.0 Release Notes
- Attach compressed ZFP data as WholeDataSet field
4. [Execution Environment](#Execution-Environment)
- Removed ExecutionWholeArray class
- Add initial support for aborting execution
5. [Worklets and Filters](#Worklets-and-Filters)
- Correct particle density filter output field
- Rename NewFilter base classes to Filter
- Fix handling of cell fields in Tube filter
- Fix setting fields to pass in Filter when setting mode
- Respect Filter::PassCoordinateSystem flag in filters creating coordinate systems.
6. [Build](#Build)
- More performance test options
- Output complete list of libraries for external Makefiles
@ -32,6 +35,7 @@ VTK-m 2.0 Release Notes
7. [Other](#Other)
- Expose the Variant helper class
- Fix VTKM_LOG_SCOPE
- Clarify field index ordering in Doxygen
# Core
@ -240,6 +244,29 @@ that `ExecutionWholeArray` had was that it provided an subscript operator
(somewhat incorrectly). Thus, any use of '[..]' to index the array portal
have to be changed to use the `Get` method.
## Add initial support for aborting execution
VTK-m now has preliminary support for aborting execution. The per-thread instances of
`RuntimeDeviceTracker` have a functor called `AbortChecker`. This functor can be set using
`RuntimeDeviceTracker::SetAbortChecker()` and cleared by `RuntimeDeviceTracker::ClearAbortChecker()`
The abort checker functor should return `true` if an abort is requested for the thread,
otherwise, it should return `false`.
Before launching a new task, `TaskExecute` calls the functor to see if an abort is requested,
and If so, throws an exception of type `vtkm::cont::ErrorUserAbort`.
Any code that wants to use the abort feature, should set an appropriate `AbortChecker`
functor for the target thread. Then any piece of code that has parts that can execute on
the device should be put under a `try-catch` block. Any clean-up that is required for an
aborted execution should be handled in a `catch` block that handles exceptions of type
`vtkm::cont::ErrorUserAbort`.
The limitation of this implementation is that it is control-side only. The check for abort
is done before launching a new device task. Once execution has begun on the device, there is
currently no way to abort that. Therefore, this feature is only useful for aborting code
that is made up of several smaller device task launches (Which is the case for most
worklets and filters in VTK-m)
# Worklets and Filters
## Correct particle density filter output field
@ -282,6 +309,43 @@ would also get around some problems with the implementation that was
removed here when mixing polylines with other cell types and degenerate
lines.)
## Fix setting fields to pass in `Filter` when setting mode
The `Filter` class has several version of the `SetFieldsToPass` method that
works in conjunction with the `FieldSelection` object to specify which
fields are mapped. For example, the user might have code like this to pass
all fields except those named `pointvar` and `cellvar`:
``` cpp
filter.SetFieldsToPass({ "pointvar", "cellvar" },
vtkm::filter::FieldSelection::Mode::Exclude);
```
This previously worked by implicitly creating a `FieldSelection` object
using the `std::initializer_list` filled with the 2 strings. This would
then be passed to the `SetFieldsToPass` method, which would capture the
`FieldSelection` object and change the mode.
This stopped working in a recent change to `FieldSelection` where each
entry is given its own mode. With this new class, the `FieldSelection`
constructor would capture each field in the default mode (`Select`) and
then change the default mode to `Exclude`. However, the already set modes
kept their `Select` status, which is not what is intended.
This behavior is fixed by adding `SetFieldToPass` overloads that capture
both the `initializer_list` and the `Mode` and then constructs the
`FieldSelection` correctly.
## Respect `Filter::PassCoordinateSystem` flag in filters creating coordinate systems
The `Filter` class has a `PassCoordinateSystem` flag that specifies whether
coordinate systems should be passed regardless of whether the associated
field is passed. However, if a filter created its output with the
`CreateResultCoordinateSystem` method this flag was ignored, and the
provided coordinate system was always passed. This might not be what the
user intended, so this method has been fixed to first check the
`PassCoordinateSystem` flag before setting the coordinates on the output.
# Build
## More performance test options
@ -375,3 +439,14 @@ This was not what was happening. The second log message was being printed
immediately after the first. This is because the scope was taken inside of
the `LogScope` method. The macro has been rewritten to put the tracking in
the right scope.
## Clarify field index ordering in Doxygen
The fields in a `DataSet` are indexed from `0` to `GetNumberOfFields() - 1`.
It is natural to assume that the fields will be indexed in the order that
they are added, but they are not. Rather, the indexing is arbitrary and can
change any time a field is added to the dataset.
To make this more clear, Doxygen documentation is added to the `DataSet`
methods to inform users to not make any assumptions about the order of
field indexing.

3
docs/changelog/2.0/shrink-filter Normal file
View File

@ -0,0 +1,3 @@
# New Shrink filter
The Shrink filter shrinks the cells of a DataSet towards their centroid, computed as the average position of the cell points. This filter disconnects the cells, duplicating the points connected to multiple cells. The resulting CellSet is always an `ExplicitCellSet`.

18
docs/changelog/array-range-fast-paths.md Normal file
View File

@ -0,0 +1,18 @@
# Fast paths for `ArrayRangeCompute` fixed
The precompiled `ArrayRangeCompute` function was not following proper fast
paths for special arrays. For example, when computing the range of an
`ArrayHandleUniformPointCoordinates`, the ranges should be taken from the
origin and spacing of the special array. However, the precompiled version
was calling the generic range computation, which was doing an unnecessary
reduction over the entire array. These fast paths have been fixed.
These mistakes in the code were caused by quirks in how templated method
overloading works. To prevent this mistake from happening again in the
precompiled `ArrayRangeCompute` function and elsewhere, all templated forms
of `ArrayRangeCompute` have been deprecated. Most will call
`ArrayRangeCompute` with no issues. For those that need the templated
version, `ArrayRangeComputeTemplate` replaces the old templated
`ArrayRangeCompute`. There is exactly one templated declaration of
`ArrayRangeComputeTemplate` that uses a class, `ArrayRangeComputeImpl`,
with partial specialization to ensure the correct form is used.

3
docs/changelog/composite-vector-filter.md Normal file
View File

@ -0,0 +1,3 @@
# New Composite Vector filter
The composite vector filter combines multiple scalar fields into a single vector field. Scalar fields are selected as the active input fields, and the combined vector field is set at the output. The current composite vector filter only supports 2d and 3d scalar field composition. Users may use `vtkm::cont::make_ArrayHandleCompositeVector` to execute a more flexible scalar field composition.

10
docs/changelog/document-field-index-order.md
View File

@ -1,10 +0,0 @@
# Clarify field index ordering in Doxygen
The fields in a `DataSet` are indexed from `0` to `GetNumberOfFields() - 1`.
It is natural to assume that the fields will be indexed in the order that
they are added, but they are not. Rather, the indexing is arbitrary and can
change any time a field is added to the dataset.
To make this more clear, Doxygen documentation is added to the `DataSet`
methods to inform users to not make any assumptions about the order of
field indexing.

27
docs/changelog/interp-any-field.md Normal file
View File

@ -0,0 +1,27 @@
# Update filters' field map and execute to work on any field type
Several filters implemented their map field by checking for common field
types and interpolated those. Although there was a float fallback to catch
odd component types, there were still a couple of issues. First, it meant
that several types got converted to `vtkm::FloatDefault`, which is often at
odds with how VTK handles it. Second, it does not handle all `Vec` lengths,
so it is still possible to drop fields.
The map field functions for these filters have been changed to support all
possible types. This is done by using the extract component functionality
to get data from any type of array. The following filters have been
updated. In some circumstances where it makes sense, a simple float
fallback is used.
* `CleanGrid`
* `CellAverage`
* `ClipWithField`
* `ClipWithImplicitFunction`
* `Contour`
* `MIRFilter`
* `NDHistogram`
* `ParticleDensityCloudInCell`
* `ParticleDensityNearestGridPoint`
* `PointAverage`
* `Probe`
* `VectorMagnitude`

36
docs/changelog/output-vec-like-fix.md Normal file
View File

@ -0,0 +1,36 @@
# Support using arrays with dynamic Vec-likes as output arrays
When you use an `ArrayHandle` as an output array in a worklet (for example,
as a `FieldOut`), the fetch operation does not read values from the array
during the `Load`. Instead, it just constructs a new object. This makes
sense as an output array is expected to have garbage in it anyway.
This is a problem for some special arrays that contain `Vec`-like objects
that are sized dynamically. For example, if you use an
`ArrayHandleGroupVecVariable`, each entry is a dynamically sized `Vec`. The
array is referenced by creating a special version of `Vec` that holds a
reference to the array portal and an index. Components are retrieved and
set by accessing the memory in the array portal. This allows us to have a
dynamically sized `Vec` in the execution environment without having to
allocate within the worklet.
The problem comes when we want to use one of these arrays with `Vec`-like
objects for an output. The typical fetch fails because you cannot construct
one of these `Vec`-like objects without an array portal to bind it to. In
these cases, we need the fetch to create the `Vec`-like object by reading
it from the array. Even though the data will be garbage, you get the
necessary buffer into the array (and nothing more).
Previously, the problem was fixed by creating partial specializations of
the `Fetch` for these `ArrayHandle`s. This worked OK as long as you were
using the array directly. However, the approach failed if the `ArrayHandle`
was wrapped in another `ArrayHandle` (for example, if an `ArrayHandleView`
was applied to an `ArrayHandleGroupVecVariable`).
To get around this problem and simplify things, the basic `Fetch` for
direct output arrays is changed to handle all cases where the values in the
`ArrayHandle` cannot be directly constructed. A compile-time check of the
array's value type is checked with `std::is_default_constructible`. If it
can be constructed, then the array is not accessed. If it cannot be
constructed, then it grabs a value out of the array.

4
docs/changelog/require-kokkos-3-7.md Normal file
View File

@ -0,0 +1,4 @@
# Require Kokkos 3.7
The minimum version of Kokkos supported is now set to Kokkos 3.7. This is
to synchronize with the development of the Kokkos team.

18
docs/changelog/resize-extracted-component.md Normal file
View File

@ -0,0 +1,18 @@
# Added ability to resize strided arrays from ArrayExtractComponent
Previously, it was not possible to resize an `ArrayHandleStride` because
the operation is a bit ambiguous. The actual array is likely to be padded
by some amount, and there could be an unknown amount of space skipped at
the beginning.
However, there is a good reason to want to resize `ArrayHandleStride`. This
is the array used to implement the `ArrayExtractComponent` feature, and
this in turn is used when extracting arrays from an `UnknownArrayHandle`
whether independent or as an `ArrayHandleRecombineVec`.
The problem really happens when you create an array of an unknown type in
an `UnknownArrayHandle` (such as with `NewInstance`) and then use that as
an output to a worklet. Sure, you could use `ArrayHandle::Allocate` to
resize before getting the array, but that is awkward for programers.
Instead, allow the extracted arrays to be resized as normal output arrays
would be.

23
docs/changelog/runtime-vec-array.md Normal file
View File

@ -0,0 +1,23 @@
# Added `ArrayHandleRuntimeVec` to specify vector sizes at runtime.
The new `ArrayHandleRuntimeVec` is a fancy `ArrayHandle` allows you to
specify a basic array of `Vec`s where the number of components of the `Vec`
are not known until runtime. (It can also optionally specify scalars.) The
behavior is much like that of `ArrayHandleGroupVecVariable` except that its
representation is much more constrained. This constrained representation
allows it to be automatically converted to an `ArrayHandleBasic` with the
proper `Vec` value type. This allows one part of code (such as a file
reader) to create an array with any `Vec` size, and then that array can be
fed to an algorithm that expects an `ArrayHandleBasic` of a certain value
type.
The `UnknownArrayHandle` has also been updated to allow
`ArrayHandleRuntimeVec` to work interchangeably with basic `ArrayHandle`.
If an `ArrayHandleRuntimeVec` is put into an `UnknownArrayHandle`, it can
be later retrieved as an `ArrayHandleBasic` as long as the base component
type matches and it has the correct amount of components. This means that
an array can be created as an `ArrayHandleRuntimeVec` and be used with any
filters or most other features designed to operate on basic `ArrayHandle`s.
Likewise, an array added as a basic `ArrayHandle` can be retrieved in an
`ArrayHandleRuntimeVec`. This makes it easier to pull arrays from VTK-m and
place them in external structures (such as `vtkDataArray`).

22
docs/changelog/serialize-dataset.md Normal file
View File

@ -0,0 +1,22 @@
# Simplified serialization of DataSet objects
`vtkm::cont::DataSet` is a dynamic object that can hold cell sets and
fields of many different types, none of which are known until runtime. This
causes a problem with serialization, which has to know what type to compile
the serialization for, particularly when unserializing the type at the
receiving end. The original implementation "solved" the problem by creating
a secondary wrapper object that was templated on types of field arrays and
cell sets that might be serialized. This is not a great solution as it
punts the problem to algorithm developers.
This problem has been completely solved for fields, as it is possible to
serialize most types of arrays without knowing their type now. You still
need to iterate over every possible `CellSet` type, but there are not that
many `CellSet`s that are practically encountered. Thus, there is now a
direct implementation of `Serialization` for `DataSet` that covers all the
data types you are likely to encounter.
The old `SerializableDataSet` has been deprecated. In the unlikely event an
algorithm needs to transfer a non-standard type of `CellSet` (such as a
permuted cell set), it can use the replacement `DataSetWithCellSetTypes`,
which just specifies the cell set types.

3
docs/changelog/statistics-filter.md Normal file
View File

@ -0,0 +1,3 @@
# New Statistics filter
The statistics filter computes the descriptive statistics of the fields specified by users based on `DescriptiveStatistics`. Users can set `RequiredStatsList` to specify which statistics will be stored in the output data set.

18
docs/changelog/unknownarray-runtime-components.md Normal file
View File

@ -0,0 +1,18 @@
# Added support for getting vec sizes of unknown arrays when runtime selected
The `GetNumberOfComponents` and `GetNumberOfComponentsFlat` methods in
`UnknownArrayHandle` have been updated to correctly report the number of
components in special `ArrayHandle`s where the `Vec` sizes of the values
are not selected until runtime.
Previously, these methods always reported 0 because the value type could
not report the size of the `Vec`. The lookup has been modified to query the
`ArrayHandle`'s `Storage` for the number of components where supported.
Note that this only works on `Storage` that provides a method to get the
runtime `Vec` size. If that is not provided, as will be the case if the
number of components can vary from one value to the next, it will still
report 0.
This feature is implemented by looking for a method named
`GetNumberOfComponents` is the `Storage` class for the `ArrayHandle`. If
this method is found, it is used to query the size at runtime.

10
examples/particle_advection/ParticleAdvection.cxx
View File

@ -16,11 +16,11 @@
#include <vtkm/io/VTKDataSetWriter.h>
// Example computing streamlines.
// An example vector field is available in the vtk-m data directory: magField.vtk
// An example vector field is available in the vtk-m data directory: rotate-vectors.vtk
// Example usage:
// this will advect 200 particles 50 steps using a step size of 0.01
// this will advect 200 particles 50 steps using a step size of 0.05
//
// Particle_Advection <path-to-data-dir>/magField.vtk vec 200 50 0.01 output.vtk
// Particle_Advection <path-to-data-dir>/rotate-vectors.vtk rotate 200 50 0.05 output.vtk
//
int main(int argc, char** argv)
@ -28,10 +28,10 @@ int main(int argc, char** argv)
auto opts = vtkm::cont::InitializeOptions::DefaultAnyDevice;
auto config = vtkm::cont::Initialize(argc, argv, opts);
if (argc < 8)
if (argc < 7)
{
std::cerr << "Usage: " << argv[0]
<< "dataFile varName numSeeds numSteps stepSize outputFile [options]" << std::endl;
<< " dataFile varName numSeeds numSteps stepSize outputFile [options]" << std::endl;
std::cerr << "where options are: " << std::endl << config.Usage << std::endl;
return -1;
}

10
examples/redistribute_points/RedistributePoints.cxx
View File

@ -133,8 +133,7 @@ public:
this->Decomposer.fill_bounds(bds, target.gid);
auto extractedDS = this->Extract(*block, bds);
// TODO: Need a better way to serialize DataSet. See issue #725.
rp.enqueue(target, vtkm::cont::SerializableDataSet<>(extractedDS));
rp.enqueue(target, extractedDS);
}
// clear our dataset.
*block = vtkm::cont::DataSet();
@ -149,10 +148,9 @@ public:
auto target = rp.in_link().target(cc);
if (rp.incoming(target.gid).size() > 0)
{
// TODO: Need a better way to serialize DataSet. See issue #725.
vtkm::cont::SerializableDataSet<> sds;
rp.dequeue(target.gid, sds);
receives.push_back(sds.DataSet);
vtkm::cont::DataSet incomingDS;
rp.dequeue(target.gid, incomingDS);
receives.push_back(incomingDS);
numValues += receives.back().GetCoordinateSystem(0).GetNumberOfPoints();
}
}

6
examples/streamline_mpi/StreamlineMPI.cxx
View File

@ -75,11 +75,11 @@ void LoadData(std::string& fname, std::vector<vtkm::cont::DataSet>& dataSets, in
}
// Example computing streamlines.
// An example vector field is available in the vtk-m data directory: magField.vtk
// An example vector field is available in the vtk-m data directory: rotate-vectors.vtk
// Example usage:
// this will advect 200 particles 50 steps using a step size of 0.01
// this will advect 200 particles 50 steps using a step size of 0.05
//
// Particle_Advection <path-to-data-dir>/magField.vtk vec 200 50 0.01 output.vtk
// Particle_Advection <path-to-data-dir>/rotate-vectors.vtk vec 200 50 0.05 output.vtk
//
int main(int argc, char** argv)

44
tutorial/CMakeLists.txt
View File

@ -9,11 +9,10 @@
##============================================================================
#add the directory that contains the VTK-m config file to the cmake
#path so that our examples can find VTK-m
#path so that our tutorial can find VTK-m
#Normally when running CMake, you need to set the VTKm_DIR to
#find the VTKmConfig.cmake file. Because we already know where this
#file is, we can add the location to look to CMAKE_PREFIX_PATH.
set(CMAKE_PREFIX_PATH ${VTKm_BINARY_DIR}/${VTKm_INSTALL_CONFIG_DIR})
cmake_minimum_required(VERSION 3.12...3.15 FATAL_ERROR)
@ -22,20 +21,41 @@ project(VTKm_tut)
#Find the VTK-m package
find_package(VTKm REQUIRED QUIET)
if (VTKm_ENABLE_TUTORIALS)
# VTK-m tutorial targets expect vtkm libraries to be namespaced with the prefix vtkm::.
# However, as the tutorial can also be built as part of the VTK-m code,
# those prefix are not added to the targets (This happens during the
# installation). To workaround this issue here, we create IMPORTED libs linking
# to the vtkm libraries used by the tutorial targets with the expected vtkm:: prefix.
vtkm_module_get_list(module_list)
foreach(tgt IN LISTS module_list)
if(TARGET ${tgt})
# The reason of creating this phony IMPORTED libraries instead of making
# ALIAS libraries is that ALIAS libraries are GLOBAL whereas IMPORTED are
# local at the directory level where they are created. We do not want these
# phony targets to be visible outside of the tutorial directory.
vtkm_target_mangle(tgt_name_mangled ${tgt})
add_library("vtkm::${tgt_name_mangled}" INTERFACE IMPORTED)
target_link_libraries("vtkm::${tgt_name_mangled}" INTERFACE ${tgt})
endif()
endforeach()
endif()
add_executable(io io.cxx)
target_link_libraries(io vtkm_filter vtkm_io)
target_link_libraries(io vtkm::io)
add_executable(contour contour.cxx)
target_link_libraries(contour vtkm_filter vtkm_io)
target_link_libraries(contour vtkm::filter_core vtkm::filter_contour vtkm::io)
add_executable(contour_two_fields contour_two_fields.cxx)
target_link_libraries(contour_two_fields vtkm_filter vtkm_io)
target_link_libraries(contour_two_fields vtkm::filter_core vtkm::filter_contour vtkm::io)
add_executable(two_filters two_filters.cxx)
target_link_libraries(two_filters vtkm_filter vtkm_io)
target_link_libraries(two_filters vtkm::filter_core vtkm::filter_contour vtkm::io)
add_executable(mag_grad mag_grad.cxx)
target_link_libraries(mag_grad vtkm_filter vtkm_io)
target_link_libraries(mag_grad vtkm::filter_core vtkm::filter_vector_analysis vtkm::io)
# Because mag_grad.cxx creates a worklet with code that
# runs on a GPU, it needs additional information.
vtkm_add_target_information(mag_grad
@ -44,23 +64,23 @@ vtkm_add_target_information(mag_grad
if (VTKm_ENABLE_RENDERING)
add_executable(rendering rendering.cxx)
target_link_libraries(rendering vtkm_filter vtkm_io vtkm_rendering)
target_link_libraries(rendering vtkm::io vtkm::rendering)
endif ()
add_executable(error_handling error_handling.cxx)
target_link_libraries(error_handling vtkm_filter vtkm_io)
target_link_libraries(error_handling vtkm::filter_core vtkm::filter_contour vtkm::io)
add_executable(logging logging.cxx)
target_link_libraries(logging vtkm_filter vtkm_io)
target_link_libraries(logging vtkm::io)
add_executable(point_to_cell point_to_cell.cxx)
target_link_libraries(point_to_cell vtkm_cont vtkm_filter vtkm_io)
target_link_libraries(point_to_cell vtkm::worklet vtkm::filter_core vtkm::io)
vtkm_add_target_information(point_to_cell
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES point_to_cell.cxx)
add_executable(extract_edges extract_edges.cxx)
target_link_libraries(extract_edges vtkm_cont vtkm_filter vtkm_io)
target_link_libraries(extract_edges vtkm::cont vtkm::filter_core vtkm::filter_contour vtkm::worklet vtkm::io)
vtkm_add_target_information(extract_edges
DROP_UNUSED_SYMBOLS MODIFY_CUDA_FLAGS
DEVICE_SOURCES extract_edges.cxx)

24
vtkm/Swap.h
View File

@ -24,21 +24,31 @@ namespace vtkm
/// Performs a swap operation. Safe to call from cuda code.
#if defined(VTKM_CUDA)
// CUDA 12 adds a `cub::Swap` function that creates ambiguity with `vtkm::Swap`.
// This happens when a function from the `cub` namespace is called with an object of a class
// defined in the `vtkm` namespace as an argument. If that function has an unqualified call to
// `Swap`, it results in ADL being used, causing the templated functions `cub::Swap` and
// `vtkm::Swap` to conflict.
#if defined(VTKM_CUDA_VERSION_MAJOR) && (VTKM_CUDA_VERSION_MAJOR >= 12) && \
defined(VTKM_CUDA_DEVICE_PASS)
using cub::Swap;
#else
template <typename T>
VTKM_EXEC_CONT void Swap(T& a, T& b)
VTKM_EXEC_CONT inline void Swap(T& a, T& b)
{
using namespace thrust;
using thrust::swap;
swap(a, b);
}
#endif
#elif defined(VTKM_HIP)
template <typename T>
__host__ void Swap(T& a, T& b)
__host__ inline void Swap(T& a, T& b)
{
using namespace std;
using std::swap;
swap(a, b);
}
template <typename T>
__device__ void Swap(T& a, T& b)
__device__ inline void Swap(T& a, T& b)
{
T temp = a;
a = b;
@ -46,9 +56,9 @@ __device__ void Swap(T& a, T& b)
}
#else
template <typename T>
VTKM_EXEC_CONT void Swap(T& a, T& b)
VTKM_EXEC_CONT inline void Swap(T& a, T& b)
{
using namespace std;
using std::swap;
swap(a, b);
}
#endif

58
vtkm/Types.h
View File

@ -1426,26 +1426,11 @@ static inline VTKM_EXEC_CONT vtkm::VecCConst<T> make_VecC(const T* array, vtkm::
namespace detail
{
template <typename T>
struct DotType
{
//results when < 32bit can be float if somehow we are using float16/float8, otherwise is
// int32 or uint32 depending on if it signed or not.
using float_type = vtkm::Float32;
using integer_type =
typename std::conditional<std::is_signed<T>::value, vtkm::Int32, vtkm::UInt32>::type;
using promote_type =
typename std::conditional<std::is_integral<T>::value, integer_type, float_type>::type;
using type =
typename std::conditional<(sizeof(T) < sizeof(vtkm::Float32)), promote_type, T>::type;
};
template <typename T>
static inline VTKM_EXEC_CONT typename DotType<typename T::ComponentType>::type vec_dot(const T& a,
const T& b)
static inline VTKM_EXEC_CONT auto vec_dot(const T& a, const T& b)
{
using U = typename DotType<typename T::ComponentType>::type;
U result = a[0] * b[0];
auto result = a[0] * b[0];
for (vtkm::IdComponent i = 1; i < a.GetNumberOfComponents(); ++i)
{
result = result + a[i] * b[i];
@ -1453,50 +1438,44 @@ static inline VTKM_EXEC_CONT typename DotType<typename T::ComponentType>::type v
return result;
}
template <typename T, vtkm::IdComponent Size>
static inline VTKM_EXEC_CONT typename DotType<T>::type vec_dot(const vtkm::Vec<T, Size>& a,
const vtkm::Vec<T, Size>& b)
static inline VTKM_EXEC_CONT auto vec_dot(const vtkm::Vec<T, Size>& a, const vtkm::Vec<T, Size>& b)
{
using U = typename DotType<T>::type;
U result = a[0] * b[0];
auto result = a[0] * b[0];
for (vtkm::IdComponent i = 1; i < Size; ++i)
{
result = result + a[i] * b[i];
}
return result;
}
}
} // namespace detail
template <typename T>
static inline VTKM_EXEC_CONT auto Dot(const T& a, const T& b) -> decltype(detail::vec_dot(a, b))
static inline VTKM_EXEC_CONT auto Dot(const T& a, const T& b)
{
return detail::vec_dot(a, b);
}
template <typename T>
static inline VTKM_EXEC_CONT typename detail::DotType<T>::type Dot(const vtkm::Vec<T, 2>& a,
const vtkm::Vec<T, 2>& b)
static inline VTKM_EXEC_CONT auto Dot(const vtkm::Vec<T, 2>& a, const vtkm::Vec<T, 2>& b)
{
return (a[0] * b[0]) + (a[1] * b[1]);
}
template <typename T>
static inline VTKM_EXEC_CONT typename detail::DotType<T>::type Dot(const vtkm::Vec<T, 3>& a,
const vtkm::Vec<T, 3>& b)
static inline VTKM_EXEC_CONT auto Dot(const vtkm::Vec<T, 3>& a, const vtkm::Vec<T, 3>& b)
{
return (a[0] * b[0]) + (a[1] * b[1]) + (a[2] * b[2]);
}
template <typename T>
static inline VTKM_EXEC_CONT typename detail::DotType<T>::type Dot(const vtkm::Vec<T, 4>& a,
const vtkm::Vec<T, 4>& b)
static inline VTKM_EXEC_CONT auto Dot(const vtkm::Vec<T, 4>& a, const vtkm::Vec<T, 4>& b)
{
return (a[0] * b[0]) + (a[1] * b[1]) + (a[2] * b[2]) + (a[3] * b[3]);
}
// Integer types of a width less than an integer get implicitly casted to
// an integer when doing a multiplication.
#define VTK_M_SCALAR_DOT(stype) \
static inline VTKM_EXEC_CONT detail::DotType<stype>::type dot(stype a, stype b) \
{ \
return a * b; \
} /* LEGACY */ \
static inline VTKM_EXEC_CONT detail::DotType<stype>::type Dot(stype a, stype b) { return a * b; }
#define VTK_M_SCALAR_DOT(stype) \
static inline VTKM_EXEC_CONT auto dot(stype a, stype b) { return a * b; } /* LEGACY */ \
static inline VTKM_EXEC_CONT auto Dot(stype a, stype b) { return a * b; }
VTK_M_SCALAR_DOT(vtkm::Int8)
VTK_M_SCALAR_DOT(vtkm::UInt8)
VTK_M_SCALAR_DOT(vtkm::Int16)
@ -1515,20 +1494,17 @@ static inline VTKM_EXEC_CONT auto dot(const T& a, const T& b) -> decltype(detail
return vtkm::Dot(a, b);
}
template <typename T>
static inline VTKM_EXEC_CONT typename detail::DotType<T>::type dot(const vtkm::Vec<T, 2>& a,
const vtkm::Vec<T, 2>& b)
static inline VTKM_EXEC_CONT auto dot(const vtkm::Vec<T, 2>& a, const vtkm::Vec<T, 2>& b)
{
return vtkm::Dot(a, b);
}
template <typename T>
static inline VTKM_EXEC_CONT typename detail::DotType<T>::type dot(const vtkm::Vec<T, 3>& a,
const vtkm::Vec<T, 3>& b)
static inline VTKM_EXEC_CONT auto dot(const vtkm::Vec<T, 3>& a, const vtkm::Vec<T, 3>& b)
{
return vtkm::Dot(a, b);
}
template <typename T>
static inline VTKM_EXEC_CONT typename detail::DotType<T>::type dot(const vtkm::Vec<T, 4>& a,
const vtkm::Vec<T, 4>& b)
static inline VTKM_EXEC_CONT auto dot(const vtkm::Vec<T, 4>& a, const vtkm::Vec<T, 4>& b)
{
return vtkm::Dot(a, b);
}

22
vtkm/VecFlat.h
View File

@ -22,18 +22,20 @@ namespace internal
{
template <typename T,
typename MultipleComponents = typename vtkm::VecTraits<T>::HasMultipleComponents>
typename MultipleComponents =
typename vtkm::internal::SafeVecTraits<T>::HasMultipleComponents>
struct TotalNumComponents;
template <typename T>
struct TotalNumComponents<T, vtkm::VecTraitsTagMultipleComponents>
{
VTKM_STATIC_ASSERT_MSG(
(std::is_same<typename vtkm::VecTraits<T>::IsSizeStatic, vtkm::VecTraitsTagSizeStatic>::value),
(std::is_same<typename vtkm::internal::SafeVecTraits<T>::IsSizeStatic,
vtkm::VecTraitsTagSizeStatic>::value),
"vtkm::VecFlat can only be used with Vec types with a static number of components.");
using ComponentType = typename vtkm::VecTraits<T>::ComponentType;
using ComponentType = typename vtkm::internal::SafeVecTraits<T>::ComponentType;
static constexpr vtkm::IdComponent value =
vtkm::VecTraits<T>::NUM_COMPONENTS * TotalNumComponents<ComponentType>::value;
vtkm::internal::SafeVecTraits<T>::NUM_COMPONENTS * TotalNumComponents<ComponentType>::value;
};
template <typename T>
@ -43,7 +45,7 @@ struct TotalNumComponents<T, vtkm::VecTraitsTagSingleComponent>
};
template <typename T>
using FlattenVec = vtkm::Vec<typename vtkm::VecTraits<T>::BaseComponentType,
using FlattenVec = vtkm::Vec<typename vtkm::internal::SafeVecTraits<T>::BaseComponentType,
vtkm::internal::TotalNumComponents<T>::value>;
template <typename T>
@ -62,10 +64,10 @@ VTKM_EXEC_CONT T GetFlatVecComponentImpl(const T& component,
}
template <typename T>
VTKM_EXEC_CONT typename vtkm::VecTraits<T>::BaseComponentType
VTKM_EXEC_CONT typename vtkm::internal::SafeVecTraits<T>::BaseComponentType
GetFlatVecComponentImpl(const T& vec, vtkm::IdComponent index, std::false_type vtkmNotUsed(isBase))
{
using Traits = vtkm::VecTraits<T>;
using Traits = vtkm::internal::SafeVecTraits<T>;
using ComponentType = typename Traits::ComponentType;
using BaseComponentType = typename Traits::BaseComponentType;
@ -78,7 +80,7 @@ GetFlatVecComponentImpl(const T& vec, vtkm::IdComponent index, std::false_type v
} // namespace detail
template <typename T>
VTKM_EXEC_CONT typename vtkm::VecTraits<T>::BaseComponentType GetFlatVecComponent(
VTKM_EXEC_CONT typename vtkm::internal::SafeVecTraits<T>::BaseComponentType GetFlatVecComponent(
const T& vec,
vtkm::IdComponent index)
{
@ -112,7 +114,7 @@ VTKM_EXEC_CONT void CopyVecNestedToFlatImpl(const NestedVecType& nestedVec,
vtkm::Vec<T, N>& flatVec,
vtkm::IdComponent flatOffset)
{
using Traits = vtkm::VecTraits<NestedVecType>;
using Traits = vtkm::internal::SafeVecTraits<NestedVecType>;
using ComponentType = typename Traits::ComponentType;
constexpr vtkm::IdComponent subSize = TotalNumComponents<ComponentType>::value;
@ -174,7 +176,7 @@ VTKM_EXEC_CONT void CopyVecFlatToNestedImpl(const vtkm::Vec<T, N>& flatVec,
vtkm::IdComponent flatOffset,
NestedVecType& nestedVec)
{
using Traits = vtkm::VecTraits<NestedVecType>;
using Traits = vtkm::internal::SafeVecTraits<NestedVecType>;
using ComponentType = typename Traits::ComponentType;
constexpr vtkm::IdComponent subSize = TotalNumComponents<ComponentType>::value;

46
vtkm/VecFromPortal.h
View File

@ -31,14 +31,6 @@ class VecFromPortal
public:
using ComponentType = typename std::remove_const<typename PortalType::ValueType>::type;
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
VecFromPortal()
: NumComponents(0)
, Offset(0)
{
}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
VecFromPortal(const PortalType& portal, vtkm::IdComponent numComponents = 0, vtkm::Id offset = 0)
@ -61,6 +53,18 @@ public:
}
}
template <vtkm::IdComponent N>
VTKM_EXEC_CONT operator vtkm::Vec<ComponentType, N>() const
{
vtkm::Vec<ComponentType, N> result;
this->CopyInto(result);
for (vtkm::IdComponent index = this->NumComponents; index < N; ++index)
{
result[index] = vtkm::TypeTraits<ComponentType>::ZeroInitialization();
}
return result;
}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
vtkm::internal::ArrayPortalValueReference<PortalType> operator[](vtkm::IdComponent index) const
@ -69,6 +73,20 @@ public:
index + this->Offset);
}
template <typename T, vtkm::IdComponent N>
VTKM_EXEC_CONT VecFromPortal& operator=(const vtkm::Vec<T, N>& src)
{
vtkm::IdComponent numComponents = vtkm::Min(N, this->NumComponents);
for (vtkm::IdComponent index = 0; index < numComponents; ++index)
{
this->Portal.Set(index + this->Offset, src[index]);
}
return *this;
}
VTKM_EXEC_CONT const PortalType& GetPortal() const { return this->Portal; }
VTKM_EXEC_CONT vtkm::Id GetOffset() const { return this->Offset; }
private:
PortalType Portal;
vtkm::IdComponent NumComponents;
@ -99,7 +117,8 @@ struct VecTraits<vtkm::VecFromPortal<PortalType>>
using VecType = vtkm::VecFromPortal<PortalType>;
using ComponentType = typename VecType::ComponentType;
using BaseComponentType = typename vtkm::VecTraits<ComponentType>::BaseComponentType;
using BaseComponentType =
typename vtkm::internal::SafeVecTraits<ComponentType>::BaseComponentType;
using HasMultipleComponents = vtkm::VecTraitsTagMultipleComponents;
using IsSizeStatic = vtkm::VecTraitsTagSizeVariable;
@ -117,6 +136,15 @@ struct VecTraits<vtkm::VecFromPortal<PortalType>>
return vector[componentIndex];
}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
static void SetComponent(const VecType& vector,
vtkm::IdComponent componentIndex,
const ComponentType& value)
{
vector[componentIndex] = value;
}
VTKM_SUPPRESS_EXEC_WARNINGS
template <vtkm::IdComponent destSize>
VTKM_EXEC_CONT static void CopyInto(const VecType& src, vtkm::Vec<ComponentType, destSize>& dest)

29
vtkm/VecTraits.h
View File

@ -45,16 +45,20 @@ struct VecTraitsTagSizeVariable
namespace internal
{
template <vtkm::IdComponent numComponents>
// Forward declaration
template <typename T>
struct SafeVecTraits;
template <vtkm::IdComponent numComponents, typename ComponentType>
struct VecTraitsMultipleComponentChooser
{
using Type = vtkm::VecTraitsTagMultipleComponents;
};
template <>
struct VecTraitsMultipleComponentChooser<1>
template <typename ComponentType>
struct VecTraitsMultipleComponentChooser<1, ComponentType>
{
using Type = vtkm::VecTraitsTagSingleComponent;
using Type = typename vtkm::internal::SafeVecTraits<ComponentType>::HasMultipleComponents;
};
} // namespace internal
@ -95,7 +99,7 @@ struct VTKM_NEVER_EXPORT VecTraits
/// is really just a scalar.
///
using HasMultipleComponents =
typename internal::VecTraitsMultipleComponentChooser<NUM_COMPONENTS>::Type;
typename internal::VecTraitsMultipleComponentChooser<NUM_COMPONENTS, ComponentType>::Type;
/// \brief A tag specifying whether the size of this vector is known at compile time.
///
@ -197,7 +201,8 @@ template <typename T, vtkm::IdComponent Size, typename NewT>
struct VecReplaceBaseComponentTypeGCC4or5
{
using type =
vtkm::Vec<typename vtkm::VecTraits<T>::template ReplaceBaseComponentType<NewT>, Size>;
vtkm::Vec<typename vtkm::internal::SafeVecTraits<T>::template ReplaceBaseComponentType<NewT>,
Size>;
};
} // namespace detail
@ -219,7 +224,8 @@ struct VTKM_NEVER_EXPORT VecTraits<vtkm::Vec<T, Size>>
/// Similar to ComponentType except that for nested vectors (e.g. Vec<Vec<T, M>, N>), it
/// returns the base scalar type at the end of the composition (T in this example).
///
using BaseComponentType = typename vtkm::VecTraits<ComponentType>::BaseComponentType;
using BaseComponentType =
typename vtkm::internal::SafeVecTraits<ComponentType>::BaseComponentType;
/// Number of components in the vector.
///
@ -235,7 +241,7 @@ struct VTKM_NEVER_EXPORT VecTraits<vtkm::Vec<T, Size>>
/// when a vector is really just a scalar.
///
using HasMultipleComponents =
typename internal::VecTraitsMultipleComponentChooser<NUM_COMPONENTS>::Type;
typename internal::VecTraitsMultipleComponentChooser<NUM_COMPONENTS, ComponentType>::Type;
/// A tag specifying whether the size of this vector is known at compile
/// time. If set to \c VecTraitsTagSizeStatic, then \c NUM_COMPONENTS is set.
@ -298,9 +304,10 @@ struct VTKM_NEVER_EXPORT VecTraits<vtkm::Vec<T, Size>>
typename detail::VecReplaceBaseComponentTypeGCC4or5<T, Size, NewComponentType>::type;
#else // !GCC <= 5
template <typename NewComponentType>
using ReplaceBaseComponentType = vtkm::Vec<
typename vtkm::VecTraits<ComponentType>::template ReplaceBaseComponentType<NewComponentType>,
Size>;
using ReplaceBaseComponentType =
vtkm::Vec<typename vtkm::internal::SafeVecTraits<
ComponentType>::template ReplaceBaseComponentType<NewComponentType>,
Size>;
#endif
///@}

73
vtkm/cont/ArrayHandleCartesianProduct.h
View File

@ -11,6 +11,8 @@
#define vtk_m_cont_ArrayHandleCartesianProduct_h
#include <vtkm/Assert.h>
#include <vtkm/Range.h>
#include <vtkm/VecTraits.h>
#include <vtkm/cont/ArrayExtractComponent.h>
#include <vtkm/cont/ArrayHandle.h>
@ -486,6 +488,77 @@ struct ArrayExtractComponentImpl<vtkm::cont::StorageTagCartesianProduct<STs...>>
}
};
template <typename S>
struct ArrayRangeComputeImpl;
template <typename ST1, typename ST2, typename ST3>
struct VTKM_CONT_EXPORT ArrayRangeComputeImpl<vtkm::cont::StorageTagCartesianProduct<ST1, ST2, ST3>>
{
template <typename T>
VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Range> operator()(
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>,
vtkm::cont::StorageTagCartesianProduct<ST1, ST2, ST3>>& input_,
vtkm::cont::DeviceAdapterId device) const
{
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(3);
auto resultPortal = result.WritePortal();
const vtkm::cont::ArrayHandleCartesianProduct<vtkm::cont::ArrayHandle<T, ST1>,
vtkm::cont::ArrayHandle<T, ST2>,
vtkm::cont::ArrayHandle<T, ST3>>& input = input_;
vtkm::cont::ArrayHandle<vtkm::Range> componentRangeArray;
vtkm::IdComponent index = 0;
vtkm::cont::ArrayHandle<T, ST1> firstArray = input.GetFirstArray();
componentRangeArray = vtkm::cont::internal::ArrayRangeComputeImpl<ST1>{}(firstArray, device);
vtkm::Id numSubComponents = componentRangeArray.GetNumberOfValues();
if (numSubComponents > 1)
{
result.Allocate(result.GetNumberOfValues() + numSubComponents - 1, vtkm::CopyFlag::On);
resultPortal = result.WritePortal();
}
auto componentRangePortal = componentRangeArray.ReadPortal();
for (vtkm::IdComponent subComponent = 0; subComponent < numSubComponents; ++subComponent)
{
resultPortal.Set(index, componentRangePortal.Get(subComponent));
++index;
}
vtkm::cont::ArrayHandle<T, ST2> secondArray = input.GetSecondArray();
componentRangeArray = vtkm::cont::internal::ArrayRangeComputeImpl<ST2>{}(secondArray, device);
numSubComponents = componentRangeArray.GetNumberOfValues();
if (numSubComponents > 1)
{
result.Allocate(result.GetNumberOfValues() + numSubComponents - 1, vtkm::CopyFlag::On);
resultPortal = result.WritePortal();
}
componentRangePortal = componentRangeArray.ReadPortal();
for (vtkm::IdComponent subComponent = 0; subComponent < numSubComponents; ++subComponent)
{
resultPortal.Set(index, componentRangePortal.Get(subComponent));
++index;
}
vtkm::cont::ArrayHandle<T, ST3> thirdArray = input.GetThirdArray();
componentRangeArray = vtkm::cont::internal::ArrayRangeComputeImpl<ST3>{}(thirdArray, device);
numSubComponents = componentRangeArray.GetNumberOfValues();
if (numSubComponents > 1)
{
result.Allocate(result.GetNumberOfValues() + numSubComponents - 1, vtkm::CopyFlag::On);
resultPortal = result.WritePortal();
}
componentRangePortal = componentRangeArray.ReadPortal();
for (vtkm::IdComponent subComponent = 0; subComponent < numSubComponents; ++subComponent)
{
resultPortal.Set(index, componentRangePortal.Get(subComponent));
++index;
}
return result;
}
};
} // namespace internal
}

33
vtkm/cont/ArrayHandleConstant.h
View File

@ -12,6 +12,9 @@
#include <vtkm/cont/ArrayHandleImplicit.h>
#include <vtkm/Range.h>
#include <vtkm/VecFlat.h>
namespace vtkm
{
namespace cont
@ -90,6 +93,36 @@ vtkm::cont::ArrayHandleConstant<T> make_ArrayHandleConstant(T value, vtkm::Id nu
{
return vtkm::cont::ArrayHandleConstant<T>(value, numberOfValues);
}
namespace internal
{
template <typename S>
struct ArrayRangeComputeImpl;
template <>
struct VTKM_CONT_EXPORT ArrayRangeComputeImpl<vtkm::cont::StorageTagConstant>
{
template <typename T>
VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Range> operator()(
const vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagConstant>& input,
vtkm::cont::DeviceAdapterId) const
{
auto value = vtkm::make_VecFlat(input.ReadPortal().Get(0));
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(value.GetNumberOfComponents());
auto resultPortal = result.WritePortal();
for (vtkm::IdComponent index = 0; index < value.GetNumberOfComponents(); ++index)
{
resultPortal.Set(index, vtkm::Range{ value[index], value[index] });
}
return result;
}
};
} // namespace internal
}
} // vtkm::cont

47
vtkm/cont/ArrayHandleCounting.h
View File

@ -12,6 +12,7 @@
#include <vtkm/cont/ArrayHandleImplicit.h>
#include <vtkm/Range.h>
#include <vtkm/TypeTraits.h>
#include <vtkm/VecTraits.h>
@ -152,6 +153,52 @@ make_ArrayHandleCounting(CountingValueType start, CountingValueType step, vtkm::
{
return vtkm::cont::ArrayHandleCounting<CountingValueType>(start, step, length);
}
namespace internal
{
template <typename S>
struct ArrayRangeComputeImpl;
template <>
struct VTKM_CONT_EXPORT ArrayRangeComputeImpl<vtkm::cont::StorageTagCounting>
{
template <typename T>
VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Range> operator()(
const vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagCounting>& input,
vtkm::cont::DeviceAdapterId) const
{
using Traits = vtkm::VecTraits<T>;
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(Traits::NUM_COMPONENTS);
auto portal = result.WritePortal();
if (portal.GetNumberOfValues() > 0)
{
T first = input.ReadPortal().Get(0);
T last = input.ReadPortal().Get(input.GetNumberOfValues() - 1);
for (vtkm::IdComponent cIndex = 0; cIndex < Traits::NUM_COMPONENTS; ++cIndex)
{
auto firstComponent = Traits::GetComponent(first, cIndex);
auto lastComponent = Traits::GetComponent(last, cIndex);
portal.Set(cIndex,
vtkm::Range(vtkm::Min(firstComponent, lastComponent),
vtkm::Max(firstComponent, lastComponent)));
}
}
else
{
// Array is empty
for (vtkm::IdComponent cIndex = 0; cIndex < Traits::NUM_COMPONENTS; ++cIndex)
{
portal.Set(cIndex, vtkm::Range{});
}
}
return result;
}
};
} // namespace internal
}
} // namespace vtkm::cont

39
vtkm/cont/ArrayHandleGroupVecVariable.h
View File

@ -47,7 +47,7 @@ public:
{
}
/// Copy constructor for any other ArrayPortalConcatenate with a portal type
/// Copy constructor for any other ArrayPortalGroupVecVariable with a portal type
/// that can be copied to this portal type. This allows us to do any type
/// casting that the portals do (like the non-const to const cast).
VTKM_SUPPRESS_EXEC_WARNINGS
@ -77,12 +77,19 @@ public:
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
void Set(vtkm::Id vtkmNotUsed(index), const ValueType& vtkmNotUsed(value)) const
void Set(vtkm::Id index, const ValueType& value) const
{
// The ValueType (VecFromPortal) operates on demand. Thus, if you set
// something in the value, it has already been passed to the array. Perhaps
// we should check to make sure that the value used matches the location
// you are trying to set in the array, but we don't do that.
if ((&value.GetPortal() == &this->ComponentsPortal) &&
(value.GetOffset() == this->OffsetsPortal.Get(index)))
{
// The ValueType (VecFromPortal) operates on demand. Thus, if you set
// something in the value, it has already been passed to the array.
}
else
{
// The value comes from somewhere else. Copy data in.
this->Get(index) = value;
}
}
VTKM_SUPPRESS_EXEC_WARNINGS
@ -237,13 +244,15 @@ public:
/// it contains three values of Vec-like objects with the data [0,1,2,3],
/// [4,5], and [6,7,8].
///
/// Note that this version of \c ArrayHandle breaks some of the assumptions
/// about \c ArrayHandle a little bit. Typically, there is exactly one type for
/// every value in the array, and this value is also the same between the
/// control and execution environment. However, this class uses \c
/// VecFromPortal it implement a Vec-like class that has a variable number of
/// values, and this type can change between control and execution
/// environments.
/// Note that caution should be used with `ArrayHandleRuntimeVec` because the
/// size of the `Vec` values is not known at compile time. Thus, the value
/// type of this array is forced to a special `VecFromPortal` class that can cause
/// surprises if treated as a `Vec`. In particular, the static `NUM_COMPONENTS`
/// expression does not exist. Furthermore, new variables of type `VecFromPortal`
/// cannot be created. This means that simple operators like `+` will not work
/// because they require an intermediate object to be created. (Equal operators
/// like `+=` do work because they are given an existing variable to place the
/// output.)
///
/// The offsets array is often derived from a list of sizes for each of the
/// entries. You can use the convenience function \c
@ -314,10 +323,6 @@ make_ArrayHandleGroupVecVariable(const ComponentsArrayHandleType& componentsArra
}
} // namespace vtkm::cont
//=============================================================================
// Specializations of worklet arguments using ArrayHandleGropuVecVariable
#include <vtkm/exec/arg/FetchTagArrayDirectOutArrayHandleGroupVecVariable.h>
//=============================================================================
// Specializations of serialization related classes
/// @cond SERIALIZATION

24
vtkm/cont/ArrayHandleIndex.h
View File

@ -10,6 +10,7 @@
#ifndef vtk_m_cont_ArrayHandleIndex_h
#define vtk_m_cont_ArrayHandleIndex_h
#include <vtkm/Range.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
namespace vtkm
@ -71,6 +72,29 @@ VTKM_CONT inline vtkm::cont::ArrayHandleIndex make_ArrayHandleIndex(vtkm::Id len
{
return vtkm::cont::ArrayHandleIndex(length);
}
namespace internal
{
template <typename S>
struct ArrayRangeComputeImpl;
template <>
struct VTKM_CONT_EXPORT ArrayRangeComputeImpl<vtkm::cont::StorageTagIndex>
{
VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Range> operator()(
const vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagIndex>& input,
vtkm::cont::DeviceAdapterId) const
{
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(1);
result.WritePortal().Set(0, vtkm::Range(0, input.GetNumberOfValues() - 1));
return result;
}
};
} // namespace internal
}
} // namespace vtkm::cont

73
vtkm/cont/ArrayHandleRecombineVec.h
View File

@ -15,8 +15,6 @@
#include <vtkm/cont/ArrayHandleStride.h>
#include <vtkm/cont/DeviceAdapterTag.h>
#include <vtkm/VecVariable.h>
#include <vtkm/internal/ArrayPortalValueReference.h>
namespace vtkm
@ -24,12 +22,18 @@ namespace vtkm
namespace internal
{
// Forward declaration
template <typename SourcePortalType>
class ArrayPortalRecombineVec;
template <typename PortalType>
class RecombineVec
{
vtkm::VecCConst<PortalType> Portals;
vtkm::Id Index;
friend vtkm::internal::ArrayPortalRecombineVec<PortalType>;
public:
using ComponentType = typename std::remove_const<typename PortalType::ValueType>::type;
@ -72,7 +76,14 @@ public:
VTKM_EXEC_CONT RecombineVec& operator=(const RecombineVec& src)
{
this->DoCopy(src);
if ((&this->Portals[0] != &src.Portals[0]) || (this->Index != src.Index))
{
this->DoCopy(src);
}
else
{
// Copying to myself. Do not need to do anything.
}
return *this;
}
@ -320,13 +331,26 @@ public:
VTKM_EXEC_CONT void Set(vtkm::Id index, const ValueType& value) const
{
// The ValueType is actually a reference back to the portals, and sets to it should
// already be set in the portal. Thus, we don't really need to do anything.
VTKM_ASSERT(value.GetIndex() == index);
if ((value.GetIndex() == index) && (value.Portals.GetPointer() == this->Portals))
{
// The ValueType is actually a reference back to the portals. If this reference is
// actually pointing back to the same index, we don't need to do anything.
}
else
{
this->DoCopy(index, value);
}
}
template <typename T>
VTKM_EXEC_CONT void Set(vtkm::Id index, const T& value) const
{
this->DoCopy(index, value);
}
private:
template <typename T>
VTKM_EXEC_CONT void DoCopy(vtkm::Id index, const T& value) const
{
using Traits = vtkm::VecTraits<T>;
VTKM_ASSERT(Traits::GetNumberOfComponents(value) == this->NumberOfComponents);
@ -408,12 +432,10 @@ class Storage<vtkm::internal::RecombineVec<ReadWritePortal>,
}
public:
VTKM_STORAGE_NO_RESIZE;
using ReadPortalType = vtkm::internal::ArrayPortalRecombineVec<ReadWritePortal>;
using WritePortalType = vtkm::internal::ArrayPortalRecombineVec<ReadWritePortal>;
VTKM_CONT static vtkm::IdComponent NumberOfComponents(
VTKM_CONT static vtkm::IdComponent GetNumberOfComponents(
const std::vector<vtkm::cont::internal::Buffer>& buffers)
{
return static_cast<vtkm::IdComponent>(
@ -426,6 +448,19 @@ public:
return SourceStorage::GetNumberOfValues(BuffersForComponent(buffers, 0));
}
VTKM_CONT static void ResizeBuffers(vtkm::Id numValues,
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::CopyFlag preserve,
vtkm::cont::Token& token)
{
vtkm::IdComponent numComponents = GetNumberOfComponents(buffers);
for (vtkm::IdComponent component = 0; component < numComponents; ++component)
{
SourceStorage::ResizeBuffers(
numValues, BuffersForComponent(buffers, component), preserve, token);
}
}
VTKM_CONT static void Fill(const std::vector<vtkm::cont::internal::Buffer>&,
const vtkm::internal::RecombineVec<ReadWritePortal>&,
vtkm::Id,
@ -440,7 +475,7 @@ public:
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token)
{
vtkm::IdComponent numComponents = NumberOfComponents(buffers);
vtkm::IdComponent numComponents = GetNumberOfComponents(buffers);
// The array portal needs a runtime-allocated array of portals for each component.
// We use the vtkm::cont::internal::Buffer object to allow us to allocate memory on the
@ -478,7 +513,7 @@ public:
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token)
{
vtkm::IdComponent numComponents = NumberOfComponents(buffers);
vtkm::IdComponent numComponents = GetNumberOfComponents(buffers);
// The array portal needs a runtime-allocated array of portals for each component.
// We use the vtkm::cont::internal::Buffer object to allow us to allocate memory on the
@ -549,9 +584,13 @@ public:
///
/// Note that caution should be used with `ArrayHandleRecombineVec` because the
/// size of the `Vec` values is not known at compile time. Thus, the value
/// type of this array is forced to a `VecVariable`, which can cause surprises
/// if treated as a `Vec`. In particular, the static `NUM_COMPONENTS` expression
/// does not exist.
/// type of this array is forced to a special `RecombineVec` class that can cause
/// surprises if treated as a `Vec`. In particular, the static `NUM_COMPONENTS`
/// expression does not exist. Furthermore, new variables of type `RecombineVec`
/// cannot be created. This means that simple operators like `+` will not work
/// because they require an intermediate object to be created. (Equal operators
/// like `+=` do work because they are given an existing variable to place the
/// output.)
///
template <typename ComponentType>
class ArrayHandleRecombineVec
@ -571,7 +610,7 @@ private:
public:
vtkm::IdComponent GetNumberOfComponents() const
{
return StorageType::NumberOfComponents(this->GetBuffers());
return StorageType::GetNumberOfComponents(this->GetBuffers());
}
vtkm::cont::ArrayHandleStride<ComponentType> GetComponentArray(
@ -618,8 +657,4 @@ struct ArrayExtractComponentImpl<vtkm::cont::internal::StorageTagRecombineVec>
}
} // namespace vtkm::cont
//=============================================================================
// Specializations of worklet arguments using ArrayHandleGropuVecVariable
#include <vtkm/exec/arg/FetchTagArrayDirectOutArrayHandleRecombineVec.h>
#endif //vtk_m_cont_ArrayHandleRecombineVec_h

504
vtkm/cont/ArrayHandleRuntimeVec.h Normal file
View File

@ -0,0 +1,504 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_cont_ArrayHandleRuntimeVec_h
#define vtk_m_cont_ArrayHandleRuntimeVec_h
#include <vtkm/cont/ArrayExtractComponent.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayHandleBasic.h>
#include <vtkm/cont/ArrayPortal.h>
#include <vtkm/cont/ErrorBadType.h>
#include <vtkm/Assert.h>
#include <vtkm/StaticAssert.h>
#include <vtkm/VecFromPortal.h>
#include <vtkm/VecTraits.h>
namespace vtkm
{
namespace internal
{
namespace detail
{
template <typename T>
struct UnrollVecImpl
{
using type = vtkm::Vec<T, 1>;
};
template <typename T, vtkm::IdComponent N>
struct UnrollVecImpl<vtkm::Vec<T, N>>
{
using subtype = typename UnrollVecImpl<T>::type;
using type = vtkm::Vec<typename subtype::ComponentType, subtype::NUM_COMPONENTS * N>;
};
} // namespace detail
// A helper class that unrolls a nested `Vec` to a single layer `Vec`. This is similar
// to `vtkm::VecFlat`, except that this only flattens `vtkm::Vec<T,N>` objects, and not
// any other `Vec`-like objects. The reason is that a `vtkm::Vec<T,N>` is the same as N
// consecutive `T` objects whereas the same may not be said about other `Vec`-like objects.
template <typename T>
using UnrollVec = typename detail::UnrollVecImpl<T>::type;
template <typename ComponentsPortalType>
class VTKM_ALWAYS_EXPORT ArrayPortalRuntimeVec
{
public:
using ComponentType = typename std::remove_const<typename ComponentsPortalType::ValueType>::type;
using ValueType = vtkm::VecFromPortal<ComponentsPortalType>;
ArrayPortalRuntimeVec() = default;
VTKM_EXEC_CONT ArrayPortalRuntimeVec(const ComponentsPortalType& componentsPortal,
vtkm::IdComponent numComponents)
: ComponentsPortal(componentsPortal)
, NumberOfComponents(numComponents)
{
}
/// Copy constructor for any other ArrayPortalRuntimeVec with a portal type
/// that can be copied to this portal type. This allows us to do any type
/// casting that the portals do (like the non-const to const cast).
template <typename OtherComponentsPortalType>
VTKM_EXEC_CONT ArrayPortalRuntimeVec(const ArrayPortalRuntimeVec<OtherComponentsPortalType>& src)
: ComponentsPortal(src.GetComponentsPortal())
, NumberOfComponents(src.GetNumberOfComponents())
{
}
VTKM_EXEC_CONT vtkm::Id GetNumberOfValues() const
{
return this->ComponentsPortal.GetNumberOfValues() / this->NumberOfComponents;
}
VTKM_EXEC_CONT ValueType Get(vtkm::Id index) const
{
return ValueType(
this->ComponentsPortal, this->NumberOfComponents, index * this->NumberOfComponents);
}
VTKM_EXEC_CONT void Set(vtkm::Id index, const ValueType& value) const
{
if ((&value.GetPortal() == &this->ComponentsPortal) &&
(value.GetOffset() == (index * this->NumberOfComponents)))
{
// The ValueType (VecFromPortal) operates on demand. Thus, if you set
// something in the value, it has already been passed to the array.
}
else
{
// The value comes from somewhere else. Copy data in.
this->Get(index) = value;
}
}
VTKM_EXEC_CONT const ComponentsPortalType& GetComponentsPortal() const
{
return this->ComponentsPortal;
}
VTKM_EXEC_CONT vtkm::IdComponent GetNumberOfComponents() const
{
return this->NumberOfComponents;
}
private:
ComponentsPortalType ComponentsPortal;
vtkm::IdComponent NumberOfComponents = 0;
};
}
} // namespace vtkm::internal
namespace vtkm
{
namespace cont
{
struct VTKM_ALWAYS_EXPORT StorageTagRuntimeVec
{
};
namespace internal
{
template <typename ComponentsPortal>
class Storage<vtkm::VecFromPortal<ComponentsPortal>, vtkm::cont::StorageTagRuntimeVec>
{
using ComponentType = typename ComponentsPortal::ValueType;
using ComponentsStorage =
vtkm::cont::internal::Storage<ComponentType, vtkm::cont::StorageTagBasic>;
VTKM_STATIC_ASSERT_MSG(
vtkm::internal::SafeVecTraits<ComponentType>::NUM_COMPONENTS == 1,
"ArrayHandleRuntimeVec only supports scalars grouped into a single Vec. Nested Vecs can "
"still be used with ArrayHandleRuntimeVec. The values are treated as flattened (like "
"with VecFlat).");
using ComponentsArray = vtkm::cont::ArrayHandle<ComponentType, StorageTagBasic>;
VTKM_STATIC_ASSERT_MSG(
(std::is_same<ComponentsPortal, typename ComponentsStorage::WritePortalType>::value),
"Used invalid ComponentsPortal type with expected ComponentsStorageTag.");
struct Info
{
vtkm::IdComponent NumberOfComponents;
};
VTKM_CONT static std::vector<vtkm::cont::internal::Buffer> ComponentsBuffers(
const std::vector<vtkm::cont::internal::Buffer>& buffers)
{
return std::vector<vtkm::cont::internal::Buffer>(buffers.begin() + 1, buffers.end());
}
public:
using ReadPortalType =
vtkm::internal::ArrayPortalRuntimeVec<typename ComponentsStorage::ReadPortalType>;
using WritePortalType =
vtkm::internal::ArrayPortalRuntimeVec<typename ComponentsStorage::WritePortalType>;
VTKM_CONT static vtkm::IdComponent GetNumberOfComponents(
const std::vector<vtkm::cont::internal::Buffer>& buffers)
{
return buffers[0].GetMetaData<Info>().NumberOfComponents;
}
VTKM_CONT static vtkm::Id GetNumberOfValues(
const std::vector<vtkm::cont::internal::Buffer>& buffers)
{
return ComponentsStorage::GetNumberOfValues(ComponentsBuffers(buffers)) /
GetNumberOfComponents(buffers);
}
VTKM_CONT static void ResizeBuffers(vtkm::Id numValues,
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::CopyFlag preserve,
vtkm::cont::Token& token)
{
ComponentsStorage::ResizeBuffers(
numValues * GetNumberOfComponents(buffers), ComponentsBuffers(buffers), preserve, token);
}
VTKM_CONT static void Fill(const std::vector<vtkm::cont::internal::Buffer>&,
const vtkm::VecFromPortal<ComponentsPortal>&,
vtkm::Id,
vtkm::Id,
vtkm::cont::Token&)
{
throw vtkm::cont::ErrorBadType("Fill not supported for ArrayHandleRuntimeVec.");
}
VTKM_CONT static ReadPortalType CreateReadPortal(
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token)
{
return ReadPortalType(
ComponentsStorage::CreateReadPortal(ComponentsBuffers(buffers), device, token),
GetNumberOfComponents(buffers));
}
VTKM_CONT static WritePortalType CreateWritePortal(
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::cont::DeviceAdapterId device,
vtkm::cont::Token& token)
{
return WritePortalType(
ComponentsStorage::CreateWritePortal(ComponentsBuffers(buffers), device, token),
GetNumberOfComponents(buffers));
}
VTKM_CONT static std::vector<vtkm::cont::internal::Buffer> CreateBuffers(
vtkm::IdComponent numComponents = 1,
const ComponentsArray& componentsArray = ComponentsArray{})
{
VTKM_LOG_IF_S(vtkm::cont::LogLevel::Warn,
(componentsArray.GetNumberOfValues() % numComponents) != 0,
"Array given to ArrayHandleRuntimeVec has size ("
<< componentsArray.GetNumberOfValues()
<< ") that is not divisible by the number of components selected ("
<< numComponents << ").");
Info info;
info.NumberOfComponents = numComponents;
return vtkm::cont::internal::CreateBuffers(info, componentsArray);
}
VTKM_CONT static ComponentsArray GetComponentsArray(
const std::vector<vtkm::cont::internal::Buffer>& buffers)
{
return ComponentsArray(ComponentsBuffers(buffers));
}
VTKM_CONT static void AsArrayHandleBasic(
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::cont::ArrayHandle<ComponentType, vtkm::cont::StorageTagBasic>& dest)
{
if (GetNumberOfComponents(buffers) != 1)
{
throw vtkm::cont::ErrorBadType(
"Attempted to pull a scalar array from an ArrayHandleRuntime that does not hold scalars.");
}
dest = GetComponentsArray(buffers);
}
template <vtkm::IdComponent N>
VTKM_CONT static void AsArrayHandleBasic(
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::cont::ArrayHandle<vtkm::Vec<ComponentType, N>, vtkm::cont::StorageTagBasic>& dest)
{
if (GetNumberOfComponents(buffers) != N)
{
throw vtkm::cont::ErrorBadType(
"Attempted to pull an array of Vecs of the wrong size from an ArrayHandleRuntime.");
}
dest = vtkm::cont::ArrayHandle<vtkm::Vec<ComponentType, N>, vtkm::cont::StorageTagBasic>(
ComponentsBuffers(buffers));
}
template <typename T, vtkm::IdComponent NInner, vtkm::IdComponent NOuter>
VTKM_CONT static void AsArrayHandleBasic(
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Vec<T, NInner>, NOuter>, vtkm::cont::StorageTagBasic>&
dest)
{
// Flatten the Vec by one level and attempt to get the array handle for that.
vtkm::cont::ArrayHandleBasic<vtkm::Vec<T, NInner * NOuter>> squashedArray;
AsArrayHandleBasic(buffers, squashedArray);
// Now unsquash the array by stealling the buffers and creating an array of the right type
dest =
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Vec<T, NInner>, NOuter>, vtkm::cont::StorageTagBasic>(
squashedArray.GetBuffers());
}
};
} // namespace internal
/// \brief Fancy array handle for a basic array with runtime selected vec size.
///
/// It is sometimes the case that you need to create an array of `Vec`s where
/// the number of components is not known until runtime. This is problematic
/// for normal `ArrayHandle`s because you have to specify the size of the `Vec`s
/// as a template parameter at compile time. `ArrayHandleRuntimeVec` can be used
/// in this case.
///
/// Note that caution should be used with `ArrayHandleRuntimeVec` because the
/// size of the `Vec` values is not known at compile time. Thus, the value
/// type of this array is forced to a special `VecFromPortal` class that can cause
/// surprises if treated as a `Vec`. In particular, the static `NUM_COMPONENTS`
/// expression does not exist. Furthermore, new variables of type `VecFromPortal`
/// cannot be created. This means that simple operators like `+` will not work
/// because they require an intermediate object to be created. (Equal operators
/// like `+=` do work because they are given an existing variable to place the
/// output.)
///
/// It is possible to provide an `ArrayHandleBasic` of the same component
/// type as the underlying storage for this array. In this case, the array
/// will be accessed much in the same manner as `ArrayHandleGroupVec`.
///
/// `ArrayHandleRuntimeVec` also allows you to convert the array to an
/// `ArrayHandleBasic` of the appropriate `Vec` type (or `component` type).
/// A runtime check will be performed to make sure the number of components
/// matches.
///
template <typename ComponentType>
class ArrayHandleRuntimeVec
: public vtkm::cont::ArrayHandle<
vtkm::VecFromPortal<typename ArrayHandleBasic<ComponentType>::WritePortalType>,
vtkm::cont::StorageTagRuntimeVec>
{
public:
VTKM_ARRAY_HANDLE_SUBCLASS(
ArrayHandleRuntimeVec,
(ArrayHandleRuntimeVec<ComponentType>),
(vtkm::cont::ArrayHandle<
vtkm::VecFromPortal<typename ArrayHandleBasic<ComponentType>::WritePortalType>,
vtkm::cont::StorageTagRuntimeVec>));
private:
using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;
using ComponentsArrayType = vtkm::cont::ArrayHandle<ComponentType, StorageTagBasic>;
public:
VTKM_CONT
ArrayHandleRuntimeVec(vtkm::IdComponent numComponents,
const ComponentsArrayType& componentsArray = ComponentsArrayType{})
: Superclass(StorageType::CreateBuffers(numComponents, componentsArray))
{
}
VTKM_CONT vtkm::IdComponent GetNumberOfComponents() const
{
return StorageType::GetNumberOfComponents(this->GetBuffers());
}
VTKM_CONT vtkm::cont::ArrayHandleBasic<ComponentType> GetComponentsArray() const
{
return StorageType::GetComponentsArray(this->GetBuffers());
}
///@{
/// \brief Converts the array to that of a basic array handle.
///
/// This method converts the `ArrayHandleRuntimeVec` to a simple `ArrayHandleBasic`.
/// This is useful if the `ArrayHandleRuntimeVec` is passed to a routine that works
/// on an array of a specific `Vec` size (or scalars). After a runtime check, the
/// array can be converted to a typical array and used as such.
template <typename ValueType>
void AsArrayHandleBasic(vtkm::cont::ArrayHandle<ValueType>& array) const
{
StorageType::AsArrayHandleBasic(this->GetBuffers(), array);
}
template <typename ArrayType>
ArrayType AsArrayHandleBasic() const
{
ArrayType array;
this->AsArrayHandleBasic(array);
return array;
}
///@}
};
/// \c make_ArrayHandleRuntimeVec is convenience function to generate an
/// ArrayHandleRuntimeVec. It takes in an ArrayHandle of values and an
/// array handle of offsets and returns an array handle with consecutive
/// entries grouped in a Vec.
///
template <typename T>
VTKM_CONT auto make_ArrayHandleRuntimeVec(
vtkm::IdComponent numComponents,
const vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>& componentsArray =
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>{})
{
using UnrolledVec = vtkm::internal::UnrollVec<T>;
using ComponentType = typename UnrolledVec::ComponentType;
// Use some dangerous magic to convert the basic array to its base component and create
// an ArrayHandleRuntimeVec from that.
vtkm::cont::ArrayHandle<ComponentType, vtkm::cont::StorageTagBasic> flatComponents(
componentsArray.GetBuffers());
return vtkm::cont::ArrayHandleRuntimeVec<ComponentType>(
numComponents * UnrolledVec::NUM_COMPONENTS, flatComponents);
}
template <typename T>
VTKM_CONT auto make_ArrayHandleRuntimeVec(
const vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagBasic>& componentsArray)
{
return make_ArrayHandleRuntimeVec(1, componentsArray);
}
namespace internal
{
template <>
struct ArrayExtractComponentImpl<vtkm::cont::StorageTagRuntimeVec>
{
template <typename T>
auto operator()(const vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagRuntimeVec>& src,
vtkm::IdComponent componentIndex,
vtkm::CopyFlag allowCopy) const
{
using ComponentType = typename T::ComponentType;
vtkm::cont::ArrayHandleRuntimeVec<ComponentType> array{ src };
constexpr vtkm::IdComponent NUM_SUB_COMPONENTS = vtkm::VecFlat<ComponentType>::NUM_COMPONENTS;
vtkm::cont::ArrayHandleStride<typename vtkm::VecTraits<T>::BaseComponentType> dest =
ArrayExtractComponentImpl<vtkm::cont::StorageTagBasic>{}(
array.GetComponentsArray(), componentIndex % NUM_SUB_COMPONENTS, allowCopy);
// Adjust stride and offset to expectations of grouped values
const vtkm::IdComponent numComponents = array.GetNumberOfComponents();
return vtkm::cont::ArrayHandleStride<typename vtkm::VecTraits<T>::BaseComponentType>(
dest.GetBasicArray(),
dest.GetNumberOfValues() / numComponents,
dest.GetStride() * numComponents,
dest.GetOffset() + (dest.GetStride() * (componentIndex / NUM_SUB_COMPONENTS)),
dest.GetModulo(),
dest.GetDivisor());
}
};
} // namespace internal
}
} // namespace vtkm::cont
//=============================================================================
// Specializations of serialization related classes
/// @cond SERIALIZATION
namespace vtkm
{
namespace cont
{
template <typename T>
struct SerializableTypeString<vtkm::cont::ArrayHandleRuntimeVec<T>>
{
static VTKM_CONT const std::string& Get()
{
static std::string name = "AH_RuntimeVec<" + SerializableTypeString<T>::Get() + ">";
return name;
}
};
template <typename VecType>
struct SerializableTypeString<vtkm::cont::ArrayHandle<VecType, vtkm::cont::StorageTagRuntimeVec>>
: SerializableTypeString<vtkm::cont::ArrayHandleRuntimeVec<typename VecType::ComponentType>>
{
};
}
} // vtkm::cont
namespace mangled_diy_namespace
{
template <typename T>
struct Serialization<vtkm::cont::ArrayHandleRuntimeVec<T>>
{
private:
using Type = vtkm::cont::ArrayHandleRuntimeVec<T>;
using BaseType = vtkm::cont::ArrayHandle<typename Type::ValueType, typename Type::StorageTag>;
public:
static VTKM_CONT void save(BinaryBuffer& bb, const BaseType& obj)
{
vtkmdiy::save(bb, Type(obj).GetNumberOfComponents());
vtkmdiy::save(bb, Type(obj).GetComponentsArray());
}
static VTKM_CONT void load(BinaryBuffer& bb, BaseType& obj)
{
vtkm::IdComponent numComponents;
vtkm::cont::ArrayHandleBasic<T> componentArray;
vtkmdiy::load(bb, numComponents);
vtkmdiy::load(bb, componentArray);
obj = vtkm::cont::make_ArrayHandleRuntimeVec(numComponents, componentArray);
}
};
template <typename VecType>
struct Serialization<vtkm::cont::ArrayHandle<VecType, vtkm::cont::StorageTagRuntimeVec>>
: Serialization<vtkm::cont::ArrayHandleRuntimeVec<typename VecType::ComponentType>>
{
};
} // diy
/// @endcond SERIALIZATION
#endif //vtk_m_cont_ArrayHandleRuntimeVec_h

81
vtkm/cont/ArrayHandleStride.h
View File

@ -158,8 +158,6 @@ class VTKM_ALWAYS_EXPORT Storage<T, vtkm::cont::StorageTagStride>
using StrideInfo = vtkm::internal::ArrayStrideInfo;
public:
VTKM_STORAGE_NO_RESIZE;
using ReadPortalType = vtkm::internal::ArrayPortalStrideRead<T>;
using WritePortalType = vtkm::internal::ArrayPortalStrideWrite<T>;
@ -174,6 +172,85 @@ public:
return GetInfo(buffers).NumberOfValues;
}
VTKM_CONT static void ResizeBuffers(vtkm::Id numValues,
const std::vector<vtkm::cont::internal::Buffer>& buffers,
vtkm::CopyFlag preserve,
vtkm::cont::Token& token)
{
StrideInfo& info = GetInfo(buffers);
if (info.NumberOfValues == numValues)
{
// Array resized to current size. Don't need to do anything.
return;
}
// Find the end index after dealing with the divsor and modulo.
auto lengthDivMod = [info](vtkm::Id length) -> vtkm::Id {
vtkm::Id resultLength = ((length - 1) / info.Divisor) + 1;
if ((info.Modulo > 0) && (info.Modulo < resultLength))
{
resultLength = info.Modulo;
}
return resultLength;
};
vtkm::Id lastStridedIndex = lengthDivMod(numValues);
vtkm::Id originalStride;
vtkm::Id originalOffset;
if (info.Stride > 0)
{
originalStride = info.Stride;
originalOffset = info.Offset;
}
else
{
// The stride is negative, which means we are counting backward. Here we have to be careful
// about the offset, which should move to push to the end of the array. We also need to
// be careful about multiplying by the stride.
originalStride = -info.Stride;
vtkm::Id originalSize = lengthDivMod(info.NumberOfValues);
// Because the stride is negative, we expect the offset to be at the end of the array.
// We will call the "real" offset the distance from that end.
originalOffset = originalSize - info.Offset - 1;
}
// If the offset is more than the stride, that means there are values skipped at the
// beginning of the array, and it is impossible to know exactly how many. In this case,
// we cannot know how to resize. (If this is an issue, we will have to change
// `ArrayHandleStride` to take resizing parameters.)
if (originalOffset >= originalStride)
{
if (numValues == 0)
{
// Array resized to zero. This can happen when releasing resources.
// Should we try to clear out the buffers, or avoid that for messing up shared buffers?
return;
}
throw vtkm::cont::ErrorBadAllocation(
"Cannot resize stride array with offset greater than stride (start of stride unknown).");
}
// lastIndex should be the index in the source array after each stride block. Assuming the
// offset is inside the first stride, this should be the end of the array regardless of
// offset.
vtkm::Id lastIndex = lastStridedIndex * originalStride;
buffers[1].SetNumberOfBytes(
vtkm::internal::NumberOfValuesToNumberOfBytes<T>(lastIndex), preserve, token);
info.NumberOfValues = numValues;
if (info.Stride < 0)
{
// As described above, when the stride is negative, we are counting backward. This means
// that the offset is actually relative to the end, so we need to adjust it to the new
// end of the array.
info.Offset = lastIndex - originalOffset - 1;
}
}
VTKM_CONT static void Fill(const std::vector<vtkm::cont::internal::Buffer>&,
const T&,
vtkm::Id,

22
vtkm/cont/ArrayHandleUniformPointCoordinates.cxx
View File

@ -86,6 +86,28 @@ ArrayExtractComponentImpl<vtkm::cont::StorageTagUniformPoints>::operator()(
}
}
VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Range>
ArrayRangeComputeImpl<vtkm::cont::StorageTagUniformPoints>::operator()(
const vtkm::cont::ArrayHandleUniformPointCoordinates& input,
vtkm::cont::DeviceAdapterId vtkmNotUsed(device)) const
{
vtkm::internal::ArrayPortalUniformPointCoordinates portal = input.ReadPortal();
// In this portal we know that the min value is the first entry and the
// max value is the last entry.
vtkm::Vec3f minimum = portal.Get(0);
vtkm::Vec3f maximum = portal.Get(portal.GetNumberOfValues() - 1);
vtkm::cont::ArrayHandle<vtkm::Range> rangeArray;
rangeArray.Allocate(3);
vtkm::cont::ArrayHandle<vtkm::Range>::WritePortalType outPortal = rangeArray.WritePortal();
outPortal.Set(0, vtkm::Range(minimum[0], maximum[0]));
outPortal.Set(1, vtkm::Range(minimum[1], maximum[1]));
outPortal.Set(2, vtkm::Range(minimum[2], maximum[2]));
return rangeArray;
}
} // namespace internal
}

12
vtkm/cont/ArrayHandleUniformPointCoordinates.h
View File

@ -10,6 +10,7 @@
#ifndef vtk_m_cont_ArrayHandleUniformPointCoordinates_h
#define vtk_m_cont_ArrayHandleUniformPointCoordinates_h
#include <vtkm/Range.h>
#include <vtkm/cont/ArrayExtractComponent.h>
#include <vtkm/cont/ArrayHandleImplicit.h>
#include <vtkm/internal/ArrayPortalUniformPointCoordinates.h>
@ -83,6 +84,17 @@ struct VTKM_CONT_EXPORT ArrayExtractComponentImpl<vtkm::cont::StorageTagUniformP
vtkm::CopyFlag allowCopy) const;
};
template <typename S>
struct ArrayRangeComputeImpl;
template <>
struct VTKM_CONT_EXPORT ArrayRangeComputeImpl<vtkm::cont::StorageTagUniformPoints>
{
VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Range> operator()(
const vtkm::cont::ArrayHandleUniformPointCoordinates& input,
vtkm::cont::DeviceAdapterId device) const;
};
} // namespace internal
}

145
vtkm/cont/ArrayRangeCompute.cxx
View File

@ -12,118 +12,16 @@
#include <vtkm/TypeList.h>
namespace vtkm
{
namespace cont
{
void ThrowArrayRangeComputeFailed()
{
throw vtkm::cont::ErrorExecution("Failed to run ArrayRangeComputation on any device.");
}
#define VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(T, Storage) \
VTKM_CONT \
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( \
const vtkm::cont::ArrayHandle<T, Storage>& input, vtkm::cont::DeviceAdapterId device) \
{ \
return detail::ArrayRangeComputeImpl(input, device); \
} \
struct SwallowSemicolon
#define VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(T, N, Storage) \
VTKM_CONT \
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( \
const vtkm::cont::ArrayHandle<vtkm::Vec<T, N>, Storage>& input, \
vtkm::cont::DeviceAdapterId device) \
{ \
return detail::ArrayRangeComputeImpl(input, device); \
} \
struct SwallowSemicolon
#define VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_SCALAR_T(Storage) \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::Int8, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::UInt8, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::Int16, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::UInt16, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::Int32, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::UInt32, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::Int64, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::UInt64, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::Float32, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(vtkm::Float64, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(char, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(signed VTKM_UNUSED_INT_TYPE, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_T(unsigned VTKM_UNUSED_INT_TYPE, Storage)
#define VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC(N, Storage) \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Int8, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::UInt8, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Int16, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::UInt16, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Int32, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::UInt32, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Int64, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::UInt64, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Float32, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Float64, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(char, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(signed VTKM_UNUSED_INT_TYPE, N, Storage); \
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(unsigned VTKM_UNUSED_INT_TYPE, N, Storage)
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_SCALAR_T(vtkm::cont::StorageTagBasic);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC(2, vtkm::cont::StorageTagBasic);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC(3, vtkm::cont::StorageTagBasic);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC(4, vtkm::cont::StorageTagBasic);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC(2, vtkm::cont::StorageTagSOA);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC(3, vtkm::cont::StorageTagSOA);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC(4, vtkm::cont::StorageTagSOA);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_SCALAR_T(vtkm::cont::StorageTagStride);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Float32, 3, vtkm::cont::StorageTagXGCCoordinates);
VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC(vtkm::Float64, 3, vtkm::cont::StorageTagXGCCoordinates);
#undef VTKM_ARRAY_RANGE_COMPUTE_IMPL_T
#undef VTKM_ARRAY_RANGE_COMPUTE_IMPL_VEC
#undef VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_SCALAR_T
#undef VTKM_ARRAY_RANGE_COMPUTE_IMPL_ALL_VEC
// Special implementation for regular point coordinates, which are easy
// to determine.
VTKM_CONT
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle<vtkm::Vec3f,
vtkm::cont::ArrayHandleUniformPointCoordinates::StorageTag>& array,
vtkm::cont::DeviceAdapterId)
{
vtkm::internal::ArrayPortalUniformPointCoordinates portal = array.ReadPortal();
// In this portal we know that the min value is the first entry and the
// max value is the last entry.
vtkm::Vec3f minimum = portal.Get(0);
vtkm::Vec3f maximum = portal.Get(portal.GetNumberOfValues() - 1);
vtkm::cont::ArrayHandle<vtkm::Range> rangeArray;
rangeArray.Allocate(3);
vtkm::cont::ArrayHandle<vtkm::Range>::WritePortalType outPortal = rangeArray.WritePortal();
outPortal.Set(0, vtkm::Range(minimum[0], maximum[0]));
outPortal.Set(1, vtkm::Range(minimum[1], maximum[1]));
outPortal.Set(2, vtkm::Range(minimum[2], maximum[2]));
return rangeArray;
}
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagIndex>& input,
vtkm::cont::DeviceAdapterId)
{
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(1);
result.WritePortal().Set(0, vtkm::Range(0, input.GetNumberOfValues() - 1));
return result;
}
#include <vtkm/cont/ArrayHandleBasic.h>
#include <vtkm/cont/ArrayHandleCartesianProduct.h>
#include <vtkm/cont/ArrayHandleCompositeVector.h>
#include <vtkm/cont/ArrayHandleConstant.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/ArrayHandleSOA.h>
#include <vtkm/cont/ArrayHandleStride.h>
#include <vtkm/cont/ArrayHandleUniformPointCoordinates.h>
#include <vtkm/cont/ArrayHandleXGCCoordinates.h>
namespace
{
@ -152,7 +50,7 @@ struct ComputeRangeFunctor
vtkm::cont::DeviceAdapterId device,
vtkm::cont::ArrayHandle<vtkm::Range>& ranges) const
{
ranges = vtkm::cont::ArrayRangeCompute(array, device);
ranges = vtkm::cont::ArrayRangeComputeTemplate(array, device);
}
// Used with vtkm::ListForEach to get components
@ -172,7 +70,7 @@ struct ComputeRangeFunctor
{
vtkm::cont::ArrayHandleStride<T> componentArray = array.ExtractComponent<T>(componentI);
vtkm::cont::ArrayHandle<vtkm::Range> componentRange =
vtkm::cont::ArrayRangeCompute(componentArray, device);
vtkm::cont::ArrayRangeComputeTemplate(componentArray, device);
rangePortal.Set(componentI, componentRange.ReadPortal().Get(0));
}
success = true;
@ -191,6 +89,21 @@ vtkm::cont::ArrayHandle<vtkm::Range> ComputeForStorage(const vtkm::cont::Unknown
} // anonymous namespace
namespace vtkm
{
namespace cont
{
namespace internal
{
void ThrowArrayRangeComputeFailed()
{
throw vtkm::cont::ErrorExecution("Failed to run ArrayRangeComputation on any device.");
}
} // namespace internal
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(const vtkm::cont::UnknownArrayHandle& array,
vtkm::cont::DeviceAdapterId device)
{
@ -214,7 +127,7 @@ vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(const vtkm::cont::Unknown
{
vtkm::cont::ArrayHandleUniformPointCoordinates uniformPoints;
array.AsArrayHandle(uniformPoints);
return vtkm::cont::ArrayRangeCompute(uniformPoints, device);
return vtkm::cont::ArrayRangeComputeTemplate(uniformPoints, device);
}
using CartesianProductStorage =
vtkm::cont::StorageTagCartesianProduct<vtkm::cont::StorageTagBasic,
@ -234,7 +147,7 @@ vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(const vtkm::cont::Unknown
}
if (array.IsStorageType<vtkm::cont::StorageTagIndex>())
{
return ArrayRangeCompute(array.AsArrayHandle<vtkm::cont::ArrayHandleIndex>(), device);
return ArrayRangeComputeTemplate(array.AsArrayHandle<vtkm::cont::ArrayHandleIndex>(), device);
}
}
catch (vtkm::cont::ErrorBadType&)

179
vtkm/cont/ArrayRangeCompute.h
View File

@ -31,7 +31,6 @@ namespace vtkm
namespace cont
{
///@{
/// \brief Compute the range of the data in an array handle.
///
/// Given an `ArrayHandle`, this function computes the range (min and max) of
@ -48,179 +47,27 @@ namespace cont
/// Note that the ArrayRangeCompute.h header file contains only precompiled overloads
/// of ArrayRangeCompute. This is so that ArrayRangeCompute.h can be included in
/// code that does not use a device compiler. If you need to compute array ranges
/// for arbitrary `ArrayHandle`s not in this precompiled list, you need to include
/// ArrayRangeComputeTemplate.h. This contains a templated version of ArrayRangeCompute
/// that will compile for any `ArrayHandle` type not already handled.
/// for arbitrary `ArrayHandle`s not in this precompiled list, you need to use
/// `ArrayRangeComputeTemplate` (declared in `ArrayRangeComputeTemplate`), which
/// will compile for any `ArrayHandle` type not already handled.
///
VTKM_CONT_EXPORT vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::UnknownArrayHandle& array,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny{});
#define VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(T, Storage) \
VTKM_CONT_EXPORT \
VTKM_CONT \
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( \
const vtkm::cont::ArrayHandle<T, Storage>& input, \
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny())
#define VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(T, N, Storage) \
VTKM_CONT_EXPORT \
VTKM_CONT \
vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute( \
const vtkm::cont::ArrayHandle<vtkm::Vec<T, N>, Storage>& input, \
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny())
#define VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_SCALAR_T(Storage) \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::Int8, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::UInt8, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::Int16, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::UInt16, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::Int32, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::UInt32, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::Int64, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::UInt64, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::Float32, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(vtkm::Float64, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(char, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(signed VTKM_UNUSED_INT_TYPE, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T(unsigned VTKM_UNUSED_INT_TYPE, Storage)
#define VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC(N, Storage) \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Int8, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::UInt8, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Int16, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::UInt16, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Int32, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::UInt32, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Int64, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::UInt64, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float32, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float64, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(char, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(signed VTKM_UNUSED_INT_TYPE, N, Storage); \
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(unsigned VTKM_UNUSED_INT_TYPE, N, Storage)
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_SCALAR_T(vtkm::cont::StorageTagBasic);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC(2, vtkm::cont::StorageTagBasic);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC(3, vtkm::cont::StorageTagBasic);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC(4, vtkm::cont::StorageTagBasic);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC(2, vtkm::cont::StorageTagSOA);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC(3, vtkm::cont::StorageTagSOA);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC(4, vtkm::cont::StorageTagSOA);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_SCALAR_T(vtkm::cont::StorageTagStride);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float32, 3, vtkm::cont::StorageTagXGCCoordinates);
VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC(vtkm::Float64, 3, vtkm::cont::StorageTagXGCCoordinates);
#undef VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_T
#undef VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_VEC
#undef VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_SCALAR_T
#undef VTK_M_ARRAY_RANGE_COMPUTE_EXPORT_ALL_VEC
VTKM_CONT_EXPORT VTKM_CONT vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle<vtkm::Vec3f,
vtkm::cont::ArrayHandleUniformPointCoordinates::StorageTag>& array,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny());
// Implementation of cartesian products
template <typename T, typename ST1, typename ST2, typename ST3>
VTKM_CONT inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle<vtkm::Vec<T, 3>,
vtkm::cont::StorageTagCartesianProduct<ST1, ST2, ST3>>& input_,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny())
namespace internal
{
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(3);
vtkm::cont::ArrayHandle<vtkm::Range> componentRangeArray;
vtkm::Range componentRange;
vtkm::cont::ArrayHandleCartesianProduct<vtkm::cont::ArrayHandle<T, ST1>,
vtkm::cont::ArrayHandle<T, ST2>,
vtkm::cont::ArrayHandle<T, ST3>>
input = input_;
vtkm::cont::ArrayHandle<T, ST1> firstArray = input.GetFirstArray();
componentRangeArray = vtkm::cont::ArrayRangeCompute(firstArray, device);
componentRange = componentRangeArray.ReadPortal().Get(0);
result.WritePortal().Set(0, componentRange);
vtkm::cont::ArrayHandle<T, ST2> secondArray = input.GetSecondArray();
componentRangeArray = vtkm::cont::ArrayRangeCompute(secondArray, device);
componentRange = componentRangeArray.ReadPortal().Get(0);
result.WritePortal().Set(1, componentRange);
vtkm::cont::ArrayHandle<T, ST3> thirdArray = input.GetThirdArray();
componentRangeArray = vtkm::cont::ArrayRangeCompute(thirdArray, device);
componentRange = componentRangeArray.ReadPortal().Get(0);
result.WritePortal().Set(2, componentRange);
return result;
}
// Implementation of constant arrays
template <typename T>
VTKM_CONT inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagConstant>& input,
vtkm::cont::DeviceAdapterId vtkmNotUsed(device) = vtkm::cont::DeviceAdapterTagAny{})
{
using Traits = vtkm::VecTraits<T>;
const T value = vtkm::cont::ArrayHandleConstant<T>(input).GetValue();
vtkm::IdComponent numComponents = Traits::GetNumberOfComponents(value);
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(numComponents);
auto portal = result.WritePortal();
for (vtkm::IdComponent cIndex = 0; cIndex < numComponents; ++cIndex)
{
auto component = Traits::GetComponent(value, cIndex);
portal.Set(cIndex, vtkm::Range(component, component));
}
return result;
}
// Implementation of counting arrays
template <typename T>
VTKM_CONT inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagCounting>& input,
vtkm::cont::DeviceAdapterId vtkmNotUsed(device) = vtkm::cont::DeviceAdapterTagAny{})
{
using Traits = vtkm::VecTraits<T>;
vtkm::cont::ArrayHandle<vtkm::Range> result;
result.Allocate(Traits::NUM_COMPONENTS);
auto portal = result.WritePortal();
if (portal.GetNumberOfValues() > 0)
{
T first = input.ReadPortal().Get(0);
T last = input.ReadPortal().Get(input.GetNumberOfValues() - 1);
for (vtkm::IdComponent cIndex = 0; cIndex < Traits::NUM_COMPONENTS; ++cIndex)
{
auto firstComponent = Traits::GetComponent(first, cIndex);
auto lastComponent = Traits::GetComponent(last, cIndex);
portal.Set(cIndex,
vtkm::Range(vtkm::Min(firstComponent, lastComponent),
vtkm::Max(firstComponent, lastComponent)));
}
}
else
{
// Array is empty
for (vtkm::IdComponent cIndex = 0; cIndex < Traits::NUM_COMPONENTS; ++cIndex)
{
portal.Set(cIndex, vtkm::Range{});
}
}
return result;
}
// Implementation of index arrays
VTKM_CONT_EXPORT vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const vtkm::cont::ArrayHandle<vtkm::Id, vtkm::cont::StorageTagIndex>& input,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny{});
///@}
VTKM_CONT_EXPORT void ThrowArrayRangeComputeFailed();
} // namespace internal
VTKM_DEPRECATED(2.1, "Moved to vtkm::cont::internal.")
inline void ThrowArrayRangeComputeFailed()
{
internal::ThrowArrayRangeComputeFailed();
}
}
} // namespace vtkm::cont

37
vtkm/cont/ArrayRangeComputeTemplate.h
View File

@ -13,6 +13,7 @@
#include <vtkm/cont/ArrayRangeCompute.h>
#include <vtkm/BinaryOperators.h>
#include <vtkm/Deprecated.h>
#include <vtkm/VecTraits.h>
#include <vtkm/cont/Algorithm.h>
@ -42,8 +43,13 @@ struct ArrayRangeComputeFunctor
}
};
} // namespace detail
namespace internal
{
template <typename T, typename S>
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeImpl(
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeGeneric(
const vtkm::cont::ArrayHandle<T, S>& input,
vtkm::cont::DeviceAdapterId device)
{
@ -74,7 +80,7 @@ inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeImpl(
initial[1] = T(std::numeric_limits<CT>::lowest());
const bool rangeComputed = vtkm::cont::TryExecuteOnDevice(
device, detail::ArrayRangeComputeFunctor{}, input, initial, result);
device, vtkm::cont::detail::ArrayRangeComputeFunctor{}, input, initial, result);
if (!rangeComputed)
{
ThrowArrayRangeComputeFailed();
@ -93,17 +99,38 @@ inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeImpl(
return range;
}
} // namespace detail
template <typename S>
struct ArrayRangeComputeImpl
{
template <typename T>
vtkm::cont::ArrayHandle<vtkm::Range> operator()(const vtkm::cont::ArrayHandle<T, S>& input,
vtkm::cont::DeviceAdapterId device) const
{
return vtkm::cont::internal::ArrayRangeComputeGeneric(input, device);
}
};
} // namespace internal
template <typename ArrayHandleType>
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeComputeTemplate(
const ArrayHandleType& input,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny{})
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
return detail::ArrayRangeComputeImpl(input, device);
return internal::ArrayRangeComputeImpl<typename ArrayHandleType::StorageTag>{}(input, device);
}
template <typename ArrayHandleType>
VTKM_DEPRECATED(2.1, "Use precompiled ArrayRangeCompute or ArrayRangeComputeTemplate.")
inline vtkm::cont::ArrayHandle<vtkm::Range> ArrayRangeCompute(
const ArrayHandleType& input,
vtkm::cont::DeviceAdapterId device = vtkm::cont::DeviceAdapterTagAny{})
{
return ArrayRangeComputeTemplate(input, device);
}
}
} // namespace vtkm::cont

2
vtkm/cont/CMakeLists.txt
View File

@ -38,6 +38,7 @@ set(headers
ArrayHandleRandomStandardNormal.h
ArrayHandleRandomUniformBits.h
ArrayHandleRandomUniformReal.h
ArrayHandleRuntimeVec.h
ArrayHandleSOA.h
ArrayHandleStride.h
ArrayHandleSwizzle.h
@ -94,6 +95,7 @@ set(headers
ErrorExecution.h
ErrorFilterExecution.h
ErrorInternal.h
ErrorUserAbort.h
ExecutionAndControlObjectBase.h
ExecutionObjectBase.h
Field.h

31
vtkm/cont/DataSet.cxx
View File

@ -8,6 +8,10 @@
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#include <vtkm/cont/CellSetExplicit.h>
#include <vtkm/cont/CellSetExtrude.h>
#include <vtkm/cont/CellSetSingleType.h>
#include <vtkm/cont/CellSetStructured.h>
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/Logging.h>
@ -341,3 +345,30 @@ void DataSet::ConvertToExpected()
} // namespace cont
} // namespace vtkm
namespace mangled_diy_namespace
{
using SerializedCellSetTypes = vtkm::ListAppend<VTKM_DEFAULT_CELL_SET_LIST,
vtkm::List<vtkm::cont::CellSetStructured<1>,
vtkm::cont::CellSetStructured<2>,
vtkm::cont::CellSetStructured<3>,
vtkm::cont::CellSetExplicit<>,
vtkm::cont::CellSetSingleType<>,
vtkm::cont::CellSetExtrude>>;
using DefaultDataSetWithCellTypes = vtkm::cont::DataSetWithCellSetTypes<SerializedCellSetTypes>;
void Serialization<vtkm::cont::DataSet>::save(BinaryBuffer& bb, const vtkm::cont::DataSet& obj)
{
vtkmdiy::save(bb, DefaultDataSetWithCellTypes{ obj });
}
void Serialization<vtkm::cont::DataSet>::load(BinaryBuffer& bb, vtkm::cont::DataSet& obj)
{
DefaultDataSetWithCellTypes data;
vtkmdiy::load(bb, data);
obj = data.DataSet;
}
} // namespace mangled_diy_namespace

71
vtkm/cont/DataSet.h
View File

@ -420,37 +420,70 @@ namespace vtkm
namespace cont
{
template <typename FieldTypeList = VTKM_DEFAULT_TYPE_LIST,
typename CellSetTypesList = VTKM_DEFAULT_CELL_SET_LIST>
struct SerializableDataSet
/// \brief Specify cell sets to use when serializing a `DataSet`.
///
/// Usually when serializing a `DataSet`, it uses a fixed set of standard
/// `CellSet` types to serialize. If you are writing an algorithm with a
/// custom `CellSet`, you can specify the `CellSet`(s) as the template
/// parameter for this class (either as a list of `CellSet`s or in a
/// single `vtkm::List` parameter).
///
template <typename... CellSetTypes>
struct DataSetWithCellSetTypes
{
SerializableDataSet() = default;
vtkm::cont::DataSet DataSet;
explicit SerializableDataSet(const vtkm::cont::DataSet& dataset)
DataSetWithCellSetTypes() = default;
explicit DataSetWithCellSetTypes(const vtkm::cont::DataSet& dataset)
: DataSet(dataset)
{
}
vtkm::cont::DataSet DataSet;
};
template <typename... CellSetTypes>
struct DataSetWithCellSetTypes<vtkm::List<CellSetTypes...>>
: DataSetWithCellSetTypes<CellSetTypes...>
{
using DataSetWithCellSetTypes<CellSetTypes...>::DataSetWithCellSetTypes;
};
template <typename FieldTypeList = VTKM_DEFAULT_TYPE_LIST,
typename CellSetTypesList = VTKM_DEFAULT_CELL_SET_LIST>
struct VTKM_DEPRECATED(
2.1,
"Serialize DataSet directly or use DataSetWithCellSetTypes for weird CellSets.")
SerializableDataSet : DataSetWithCellSetTypes<CellSetTypesList>
{
using DataSetWithCellSetTypes<CellSetTypesList>::DataSetWithCellSetTypes;
};
}
} // vtkm::cont
namespace mangled_diy_namespace
{
template <typename FieldTypeList, typename CellSetTypesList>
struct Serialization<vtkm::cont::SerializableDataSet<FieldTypeList, CellSetTypesList>>
template <>
struct VTKM_CONT_EXPORT Serialization<vtkm::cont::DataSet>
{
static VTKM_CONT void foo();
static VTKM_CONT void save(BinaryBuffer& bb, const vtkm::cont::DataSet& obj);
static VTKM_CONT void load(BinaryBuffer& bb, vtkm::cont::DataSet& obj);
};
template <typename... CellSetTypes>
struct Serialization<vtkm::cont::DataSetWithCellSetTypes<CellSetTypes...>>
{
private:
using Type = vtkm::cont::SerializableDataSet<FieldTypeList, CellSetTypesList>;
using Type = vtkm::cont::DataSetWithCellSetTypes<CellSetTypes...>;
public:
static VTKM_CONT void save(BinaryBuffer& bb, const Type& serializable)
{
const auto& dataset = serializable.DataSet;
vtkmdiy::save(bb, dataset.GetCellSet().ResetCellSetList(CellSetTypesList{}));
vtkmdiy::save(bb, dataset.GetCellSet().ResetCellSetList(vtkm::List<CellSetTypes...>{}));
vtkm::IdComponent numberOfFields = dataset.GetNumberOfFields();
vtkmdiy::save(bb, numberOfFields);
@ -472,7 +505,7 @@ public:
auto& dataset = serializable.DataSet;
dataset = {}; // clear
vtkm::cont::UncertainCellSet<CellSetTypesList> cells;
vtkm::cont::UncertainCellSet<vtkm::List<CellSetTypes...>> cells;
vtkmdiy::load(bb, cells);
dataset.SetCellSet(cells);
@ -496,6 +529,20 @@ public:
}
};
template <typename... CellSetTypes>
struct Serialization<vtkm::cont::DataSetWithCellSetTypes<vtkm::List<CellSetTypes...>>>
: Serialization<vtkm::cont::DataSetWithCellSetTypes<CellSetTypes...>>
{
};
VTKM_DEPRECATED_SUPPRESS_BEGIN
template <typename FieldTypeList, typename CellSetTypesList>
struct Serialization<vtkm::cont::SerializableDataSet<FieldTypeList, CellSetTypesList>>
: Serialization<vtkm::cont::DataSetWithCellSetTypes<CellSetTypesList>>
{
};
VTKM_DEPRECATED_SUPPRESS_END
} // diy
/// @endcond SERIALIZATION

42
vtkm/cont/ErrorUserAbort.h Normal file
View File

@ -0,0 +1,42 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_cont_ErrorUserAbort_h
#define vtk_m_cont_ErrorUserAbort_h
#include <vtkm/cont/Error.h>
namespace vtkm
{
namespace cont
{
VTKM_SILENCE_WEAK_VTABLE_WARNING_START
/// This class is thrown when vtk-m detects a request for aborting execution
/// in the current thread
///
class VTKM_ALWAYS_EXPORT ErrorUserAbort : public Error
{
public:
ErrorUserAbort()
: Error(Message, true)
{
}
private:
static constexpr const char* Message = "User abort detected.";
};
VTKM_SILENCE_WEAK_VTABLE_WARNING_END
}
} // namespace vtkm::cont
#endif // vtk_m_cont_ErrorUserAbort_h

14
vtkm/cont/RuntimeDeviceInformation.cxx
View File

@ -97,13 +97,6 @@ public:
throw vtkm::cont::ErrorBadDevice("Tried to set the number of threads on an invalid device");
}
VTKM_CONT virtual vtkm::cont::internal::RuntimeDeviceConfigReturnCode SetNumaRegions(
const vtkm::Id&) override final
{
throw vtkm::cont::ErrorBadDevice(
"Tried to set the number of numa regions on an invalid device");
}
VTKM_CONT virtual vtkm::cont::internal::RuntimeDeviceConfigReturnCode SetDeviceInstance(
const vtkm::Id&) override final
{
@ -116,13 +109,6 @@ public:
throw vtkm::cont::ErrorBadDevice("Tried to get the number of threads on an invalid device");
}
VTKM_CONT virtual vtkm::cont::internal::RuntimeDeviceConfigReturnCode GetNumaRegions(
vtkm::Id&) const override final
{
throw vtkm::cont::ErrorBadDevice(
"Tried to get the number of numa regions on an invalid device");
}
VTKM_CONT virtual vtkm::cont::internal::RuntimeDeviceConfigReturnCode GetDeviceInstance(
vtkm::Id&) const override final
{

109
vtkm/cont/RuntimeDeviceTracker.cxx
View File

@ -13,6 +13,7 @@
#include <vtkm/cont/ErrorBadValue.h>
#include <algorithm>
#include <array>
#include <map>
#include <mutex>
#include <sstream>
@ -28,40 +29,13 @@ 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 ResetRuntimeAllowed() { this->RuntimeAllowed.fill(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];
std::array<bool, VTKM_MAX_DEVICE_ADAPTER_ID> RuntimeAllowed;
bool ThreadFriendlyMemAlloc = false;
std::function<bool()> AbortChecker;
};
}
@ -99,7 +73,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->GetRuntimeAllowed(static_cast<std::size_t>(i)))
if (this->Internals->RuntimeAllowed[static_cast<std::size_t>(i)])
{
return true;
}
@ -109,14 +83,14 @@ bool RuntimeDeviceTracker::CanRunOn(vtkm::cont::DeviceAdapterId deviceId) const
else
{
this->CheckDevice(deviceId);
return this->Internals->GetRuntimeAllowed(deviceId.GetValue());
return this->Internals->RuntimeAllowed[static_cast<std::size_t>(deviceId.GetValue())];
}
}
VTKM_CONT
bool RuntimeDeviceTracker::GetThreadFriendlyMemAlloc() const
{
return this->Internals->GetThreadFriendlyMemAlloc();
return this->Internals->ThreadFriendlyMemAlloc;
}
VTKM_CONT
@ -124,13 +98,13 @@ void RuntimeDeviceTracker::SetDeviceState(vtkm::cont::DeviceAdapterId deviceId,
{
this->CheckDevice(deviceId);
this->Internals->SetRuntimeAllowed(deviceId.GetValue(), state);
this->Internals->RuntimeAllowed[static_cast<std::size_t>(deviceId.GetValue())] = state;
}
VTKM_CONT
void RuntimeDeviceTracker::SetThreadFriendlyMemAlloc(bool state)
{
this->Internals->SetThreadFriendlyMemAlloc(state);
this->Internals->ThreadFriendlyMemAlloc = state;
}
VTKM_CONT void RuntimeDeviceTracker::ResetDevice(vtkm::cont::DeviceAdapterId deviceId)
@ -163,7 +137,7 @@ void RuntimeDeviceTracker::Reset()
if (device.IsValueValid())
{
const bool state = runtimeDevice.Exists(device);
this->Internals->SetRuntimeAllowed(device.GetValue(), state);
this->Internals->RuntimeAllowed[static_cast<std::size_t>(device.GetValue())] = state;
}
}
this->LogEnabledDevices();
@ -203,14 +177,36 @@ void RuntimeDeviceTracker::ForceDevice(DeviceAdapterId deviceId)
}
this->Internals->ResetRuntimeAllowed();
this->Internals->SetRuntimeAllowed(deviceId.GetValue(), runtimeExists);
this->Internals->RuntimeAllowed[static_cast<std::size_t>(deviceId.GetValue())] = runtimeExists;
this->LogEnabledDevices();
}
}
VTKM_CONT void RuntimeDeviceTracker::CopyStateFrom(const vtkm::cont::RuntimeDeviceTracker& tracker)
{
*(this->Internals) = tracker.Internals;
*(this->Internals) = *tracker.Internals;
}
VTKM_CONT
void RuntimeDeviceTracker::SetAbortChecker(const std::function<bool()>& func)
{
this->Internals->AbortChecker = func;
}
VTKM_CONT
bool RuntimeDeviceTracker::CheckForAbortRequest() const
{
if (this->Internals->AbortChecker)
{
return this->Internals->AbortChecker();
}
return false;
}
VTKM_CONT
void RuntimeDeviceTracker::ClearAbortChecker()
{
this->Internals->AbortChecker = nullptr;
}
VTKM_CONT
@ -247,25 +243,12 @@ void RuntimeDeviceTracker::LogEnabledDevices() const
}
VTKM_CONT
ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(vtkm::cont::DeviceAdapterId device,
RuntimeDeviceTrackerMode mode)
: RuntimeDeviceTracker(GetRuntimeDeviceTracker().Internals, false)
, SavedState(new detail::RuntimeDeviceTrackerInternals(this->Internals))
ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(
const vtkm::cont::RuntimeDeviceTracker& tracker)
: RuntimeDeviceTracker(tracker.Internals, false)
, SavedState(new detail::RuntimeDeviceTrackerInternals(*this->Internals))
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Entering scoped runtime region");
if (mode == RuntimeDeviceTrackerMode::Force)
{
this->ForceDevice(device);
}
else if (mode == RuntimeDeviceTrackerMode::Enable)
{
this->ResetDevice(device);
}
else if (mode == RuntimeDeviceTrackerMode::Disable)
{
this->DisableDevice(device);
}
}
VTKM_CONT
@ -273,11 +256,8 @@ ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(
vtkm::cont::DeviceAdapterId device,
RuntimeDeviceTrackerMode mode,
const vtkm::cont::RuntimeDeviceTracker& tracker)
: RuntimeDeviceTracker(tracker.Internals, false)
, SavedState(new detail::RuntimeDeviceTrackerInternals(this->Internals))
: ScopedRuntimeDeviceTracker(tracker)
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Entering scoped runtime region");
if (mode == RuntimeDeviceTrackerMode::Force)
{
this->ForceDevice(device);
@ -292,20 +272,19 @@ ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(
}
}
VTKM_CONT
ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(
VTKM_CONT ScopedRuntimeDeviceTracker::ScopedRuntimeDeviceTracker(
const std::function<bool()>& abortChecker,
const vtkm::cont::RuntimeDeviceTracker& tracker)
: RuntimeDeviceTracker(tracker.Internals, false)
, SavedState(new detail::RuntimeDeviceTrackerInternals(this->Internals))
: ScopedRuntimeDeviceTracker(tracker)
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Entering scoped runtime region");
this->SetAbortChecker(abortChecker);
}
VTKM_CONT
ScopedRuntimeDeviceTracker::~ScopedRuntimeDeviceTracker()
{
VTKM_LOG_S(vtkm::cont::LogLevel::DevicesEnabled, "Leaving scoped runtime region");
*(this->Internals) = this->SavedState.get();
*(this->Internals) = *this->SavedState;
this->LogEnabledDevices();
}

150
vtkm/cont/RuntimeDeviceTracker.h
View File

@ -17,6 +17,7 @@
#include <vtkm/cont/ErrorBadDevice.h>
#include <vtkm/cont/RuntimeDeviceInformation.h>
#include <functional>
#include <memory>
namespace vtkm
@ -123,6 +124,18 @@ public:
///
VTKM_CONT void CopyStateFrom(const vtkm::cont::RuntimeDeviceTracker& tracker);
///@{
/// \brief Set/Clear the abort checker functor.
///
/// If set the abort checker functor is called by \c TryExecute before scheduling
/// a task on a device from the associated the thread. If the functor returns
/// \e true, an exception is thrown.
VTKM_CONT void SetAbortChecker(const std::function<bool()>& func);
VTKM_CONT void ClearAbortChecker();
///@}
VTKM_CONT bool CheckForAbortRequest() const;
VTKM_CONT void PrintSummary(std::ostream& out) const;
private:
@ -149,82 +162,7 @@ private:
void LogEnabledDevices() const;
};
enum struct RuntimeDeviceTrackerMode
{
Force,
Enable,
Disable
};
/// A class that can be used to determine or modify which device adapter
/// VTK-m algorithms should be run on. This class captures the state
/// of the per-thread device adapter and will revert any changes applied
/// during its lifetime on destruction.
///
///
struct VTKM_CONT_EXPORT ScopedRuntimeDeviceTracker : public vtkm::cont::RuntimeDeviceTracker
{
/// Construct a ScopedRuntimeDeviceTracker where the state of the active devices
/// for the current thread are determined by the parameters to the constructor.
///
/// 'Force'
/// - Force-Enable the provided single device adapter
/// - Force-Enable all device adapters when using vtkm::cont::DeviceAdaterTagAny
/// 'Enable'
/// - Enable the provided single device adapter if it was previously disabled
/// - Enable all device adapters that are currently disabled when using
/// vtkm::cont::DeviceAdaterTagAny
/// 'Disable'
/// - Disable the provided single device adapter
/// - Disable all device adapters when using vtkm::cont::DeviceAdaterTagAny
///
/// Constructor is not thread safe
VTKM_CONT ScopedRuntimeDeviceTracker(
vtkm::cont::DeviceAdapterId device,
RuntimeDeviceTrackerMode mode = RuntimeDeviceTrackerMode::Force);
/// Construct a ScopedRuntimeDeviceTracker associated with the thread
/// associated with the provided tracker. The active devices
/// for the current thread are determined by the parameters to the constructor.
///
/// 'Force'
/// - Force-Enable the provided single device adapter
/// - Force-Enable all device adapters when using vtkm::cont::DeviceAdaterTagAny
/// 'Enable'
/// - Enable the provided single device adapter if it was previously disabled
/// - Enable all device adapters that are currently disabled when using
/// vtkm::cont::DeviceAdaterTagAny
/// 'Disable'
/// - Disable the provided single device adapter
/// - Disable all device adapters when using vtkm::cont::DeviceAdaterTagAny
///
/// Any modifications to the ScopedRuntimeDeviceTracker will effect what
/// ever thread the \c tracker is associated with, which might not be
/// the thread which ScopedRuntimeDeviceTracker was constructed on.
///
/// Constructor is not thread safe
VTKM_CONT ScopedRuntimeDeviceTracker(vtkm::cont::DeviceAdapterId device,
RuntimeDeviceTrackerMode mode,
const vtkm::cont::RuntimeDeviceTracker& tracker);
/// Construct a ScopedRuntimeDeviceTracker associated with the thread
/// associated with the provided tracker.
///
/// Any modifications to the ScopedRuntimeDeviceTracker will effect what
/// ever thread the \c tracker is associated with, which might not be
/// the thread which ScopedRuntimeDeviceTracker was constructed on.
///
/// Constructor is not thread safe
VTKM_CONT ScopedRuntimeDeviceTracker(const vtkm::cont::RuntimeDeviceTracker& tracker);
/// Destructor is not thread safe
VTKM_CONT ~ScopedRuntimeDeviceTracker();
private:
std::unique_ptr<detail::RuntimeDeviceTrackerInternals> SavedState;
};
///----------------------------------------------------------------------------
/// \brief Get the \c RuntimeDeviceTracker for the current thread.
///
/// Many features in VTK-m will attempt to run algorithms on the "best
@ -236,6 +174,66 @@ private:
VTKM_CONT_EXPORT
VTKM_CONT
vtkm::cont::RuntimeDeviceTracker& GetRuntimeDeviceTracker();
enum struct RuntimeDeviceTrackerMode
{
Force,
Enable,
Disable
};
///----------------------------------------------------------------------------
/// A class to create a scoped runtime device tracker object. This object captures the state
/// of the per-thread device tracker and will revert any changes applied
/// during its lifetime on destruction.
///
struct VTKM_CONT_EXPORT ScopedRuntimeDeviceTracker : public vtkm::cont::RuntimeDeviceTracker
{
/// Construct a ScopedRuntimeDeviceTracker associated with the thread,
/// associated with the provided tracker (defaults to current thread's tracker).
///
/// Any modifications to the ScopedRuntimeDeviceTracker will effect what
/// ever thread the \c tracker is associated with, which might not be
/// the thread on which the ScopedRuntimeDeviceTracker was constructed.
///
/// Constructors are not thread safe
/// @{
///
VTKM_CONT ScopedRuntimeDeviceTracker(
const vtkm::cont::RuntimeDeviceTracker& tracker = GetRuntimeDeviceTracker());
/// Use this constructor to modify the state of the device adapters associated with
/// the provided tracker. Use \p mode with \p device as follows:
///
/// 'Force' (default)
/// - Force-Enable the provided single device adapter
/// - Force-Enable all device adapters when using vtkm::cont::DeviceAdaterTagAny
/// 'Enable'
/// - Enable the provided single device adapter if it was previously disabled
/// - Enable all device adapters that are currently disabled when using
/// vtkm::cont::DeviceAdaterTagAny
/// 'Disable'
/// - Disable the provided single device adapter
/// - Disable all device adapters when using vtkm::cont::DeviceAdaterTagAny
///
VTKM_CONT ScopedRuntimeDeviceTracker(
vtkm::cont::DeviceAdapterId device,
RuntimeDeviceTrackerMode mode = RuntimeDeviceTrackerMode::Force,
const vtkm::cont::RuntimeDeviceTracker& tracker = GetRuntimeDeviceTracker());
/// Use this constructor to set the abort checker functor for the provided tracker.
///
VTKM_CONT ScopedRuntimeDeviceTracker(
const std::function<bool()>& abortChecker,
const vtkm::cont::RuntimeDeviceTracker& tracker = GetRuntimeDeviceTracker());
/// Destructor is not thread safe
VTKM_CONT ~ScopedRuntimeDeviceTracker();
private:
std::unique_ptr<detail::RuntimeDeviceTrackerInternals> SavedState;
};
}
} // namespace vtkm::cont

8
vtkm/cont/TryExecute.cxx
View File

@ -12,6 +12,7 @@
#include <vtkm/cont/ErrorBadType.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/ErrorFilterExecution.h>
#include <vtkm/cont/ErrorUserAbort.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
namespace vtkm
@ -55,6 +56,13 @@ VTKM_CONT_EXPORT void HandleTryExecuteException(vtkm::cont::DeviceAdapterId devi
VTKM_LOG_TRYEXECUTE_FAIL("ErrorBadValue (" << e.GetMessage() << ")", functorName, deviceId);
throw;
}
catch (vtkm::cont::ErrorUserAbort& e)
{
VTKM_LOG_S(vtkm::cont::LogLevel::Info,
e.GetMessage() << " Aborting: " << functorName << ", on device "
<< deviceId.GetName());
throw;
}
catch (vtkm::cont::Error& e)
{
VTKM_LOG_TRYEXECUTE_FAIL(e.GetMessage(), functorName, deviceId);

6
vtkm/cont/TryExecute.h
View File

@ -12,6 +12,7 @@
#include <vtkm/cont/DeviceAdapterList.h>
#include <vtkm/cont/DeviceAdapterTag.h>
#include <vtkm/cont/ErrorUserAbort.h>
#include <vtkm/cont/Logging.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
@ -40,6 +41,11 @@ inline bool TryExecuteIfValid(std::true_type,
{
try
{
if (tracker.CheckForAbortRequest())
{
throw vtkm::cont::ErrorUserAbort{};
}
return f(tag, std::forward<Args>(args)...);
}
catch (...)

145
vtkm/cont/UnknownArrayHandle.cxx
View File

@ -18,6 +18,7 @@
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/ArrayHandlePermutation.h>
#include <vtkm/cont/ArrayHandleReverse.h>
#include <vtkm/cont/ArrayHandleRuntimeVec.h>
#include <vtkm/cont/ArrayHandleSOA.h>
#include <vtkm/cont/ArrayHandleUniformPointCoordinates.h>
#include <vtkm/cont/ErrorBadType.h>
@ -45,12 +46,16 @@ struct AllVecImpl<N, vtkm::List<Scalars...>>
template <vtkm::IdComponent N>
using AllVec = typename AllVecImpl<N, vtkm::TypeListBaseC>::type;
template <typename T>
using IsBasicStorage = std::is_same<vtkm::cont::StorageTagBasic, T>;
template <typename List>
using RemoveBasicStorage = vtkm::ListRemoveIf<List, IsBasicStorage>;
using UnknownSerializationTypes =
vtkm::ListAppend<vtkm::TypeListBaseC, AllVec<2>, AllVec<3>, AllVec<4>>;
using UnknownSerializationStorage =
vtkm::ListAppend<VTKM_DEFAULT_STORAGE_LIST,
vtkm::List<vtkm::cont::StorageTagBasic,
vtkm::cont::StorageTagCartesianProduct<vtkm::cont::StorageTagBasic,
using UnknownSerializationSpecializedStorage =
vtkm::ListAppend<RemoveBasicStorage<VTKM_DEFAULT_STORAGE_LIST>,
vtkm::List<vtkm::cont::StorageTagCartesianProduct<vtkm::cont::StorageTagBasic,
vtkm::cont::StorageTagBasic,
vtkm::cont::StorageTagBasic>,
vtkm::cont::StorageTagConstant,
@ -265,7 +270,7 @@ vtkm::IdComponent UnknownArrayHandle::GetNumberOfComponents() const
{
if (this->Container)
{
return this->Container->NumberOfComponents();
return this->Container->NumberOfComponents(this->Container->ArrayHandlePointer);
}
else
{
@ -277,7 +282,7 @@ VTKM_CONT vtkm::IdComponent UnknownArrayHandle::GetNumberOfComponentsFlat() cons
{
if (this->Container)
{
return this->Container->NumberOfComponentsFlat();
return this->Container->NumberOfComponentsFlat(this->Container->ArrayHandlePointer);
}
else
{
@ -432,30 +437,82 @@ std::string SerializableTypeString<vtkm::cont::UnknownArrayHandle>::Get()
}
} // namespace vtkm::cont
namespace mangled_diy_namespace
namespace
{
void Serialization<vtkm::cont::UnknownArrayHandle>::save(BinaryBuffer& bb,
const vtkm::cont::UnknownArrayHandle& obj)
enum struct SerializedArrayType : vtkm::UInt8
{
BasicArray = 0,
SpecializedStorage
};
struct SaveBasicArray
{
template <typename ComponentType>
VTKM_CONT void operator()(ComponentType,
mangled_diy_namespace::BinaryBuffer& bb,
const vtkm::cont::UnknownArrayHandle& obj,
bool& saved)
{
// Basic arrays and arrays with compatible layouts can be loaed/saved as an
// ArrayHandleRuntimeVec. Thus, we can load/save them all with one routine.
using ArrayType = vtkm::cont::ArrayHandleRuntimeVec<ComponentType>;
if (!saved && obj.CanConvert<ArrayType>())
{
ArrayType array = obj.AsArrayHandle<ArrayType>();
vtkmdiy::save(bb, SerializedArrayType::BasicArray);
vtkmdiy::save(bb, vtkm::cont::TypeToString<ComponentType>());
vtkmdiy::save(bb, array);
saved = true;
}
}
};
struct LoadBasicArray
{
template <typename ComponentType>
VTKM_CONT void operator()(ComponentType,
mangled_diy_namespace::BinaryBuffer& bb,
vtkm::cont::UnknownArrayHandle& obj,
const std::string& componentTypeString,
bool& loaded)
{
if (!loaded && (componentTypeString == vtkm::cont::TypeToString<ComponentType>()))
{
vtkm::cont::ArrayHandleRuntimeVec<ComponentType> array;
vtkmdiy::load(bb, array);
obj = array;
loaded = true;
}
}
};
VTKM_CONT void SaveSpecializedArray(mangled_diy_namespace::BinaryBuffer& bb,
const vtkm::cont::UnknownArrayHandle& obj)
{
vtkm::IdComponent numComponents = obj.GetNumberOfComponents();
switch (numComponents)
{
case 1:
vtkmdiy::save(bb, SerializedArrayType::SpecializedStorage);
vtkmdiy::save(bb, numComponents);
vtkmdiy::save(bb, obj.ResetTypes<vtkm::TypeListBaseC, UnknownSerializationStorage>());
vtkmdiy::save(bb,
obj.ResetTypes<vtkm::TypeListBaseC, UnknownSerializationSpecializedStorage>());
break;
case 2:
vtkmdiy::save(bb, SerializedArrayType::SpecializedStorage);
vtkmdiy::save(bb, numComponents);
vtkmdiy::save(bb, obj.ResetTypes<AllVec<2>, UnknownSerializationStorage>());
vtkmdiy::save(bb, obj.ResetTypes<AllVec<2>, UnknownSerializationSpecializedStorage>());
break;
case 3:
vtkmdiy::save(bb, SerializedArrayType::SpecializedStorage);
vtkmdiy::save(bb, numComponents);
vtkmdiy::save(bb, obj.ResetTypes<AllVec<3>, UnknownSerializationStorage>());
vtkmdiy::save(bb, obj.ResetTypes<AllVec<3>, UnknownSerializationSpecializedStorage>());
break;
case 4:
vtkmdiy::save(bb, SerializedArrayType::SpecializedStorage);
vtkmdiy::save(bb, numComponents);
vtkmdiy::save(bb, obj.ResetTypes<AllVec<4>, UnknownSerializationStorage>());
vtkmdiy::save(bb, obj.ResetTypes<AllVec<4>, UnknownSerializationSpecializedStorage>());
break;
default:
throw vtkm::cont::ErrorBadType(
@ -465,16 +522,17 @@ void Serialization<vtkm::cont::UnknownArrayHandle>::save(BinaryBuffer& bb,
}
}
void Serialization<vtkm::cont::UnknownArrayHandle>::load(BinaryBuffer& bb,
vtkm::cont::UnknownArrayHandle& obj)
VTKM_CONT void LoadSpecializedArray(mangled_diy_namespace::BinaryBuffer& bb,
vtkm::cont::UnknownArrayHandle& obj)
{
vtkm::IdComponent numComponents;
vtkmdiy::load(bb, numComponents);
vtkm::cont::UncertainArrayHandle<vtkm::TypeListBaseC, UnknownSerializationStorage> array1;
vtkm::cont::UncertainArrayHandle<AllVec<2>, UnknownSerializationStorage> array2;
vtkm::cont::UncertainArrayHandle<AllVec<3>, UnknownSerializationStorage> array3;
vtkm::cont::UncertainArrayHandle<AllVec<4>, UnknownSerializationStorage> array4;
vtkm::cont::UncertainArrayHandle<vtkm::TypeListBaseC, UnknownSerializationSpecializedStorage>
array1;
vtkm::cont::UncertainArrayHandle<AllVec<2>, UnknownSerializationSpecializedStorage> array2;
vtkm::cont::UncertainArrayHandle<AllVec<3>, UnknownSerializationSpecializedStorage> array3;
vtkm::cont::UncertainArrayHandle<AllVec<4>, UnknownSerializationSpecializedStorage> array4;
switch (numComponents)
{
@ -499,4 +557,53 @@ void Serialization<vtkm::cont::UnknownArrayHandle>::load(BinaryBuffer& bb,
}
}
} // anonymous namespace
namespace mangled_diy_namespace
{
void Serialization<vtkm::cont::UnknownArrayHandle>::save(BinaryBuffer& bb,
const vtkm::cont::UnknownArrayHandle& obj)
{
bool saved = false;
// First, try serializing basic arrays (which we can do for any Vec size).
vtkm::ListForEach(SaveBasicArray{}, vtkm::TypeListBaseC{}, bb, obj, saved);
// If that did not work, try one of the specialized arrays.
if (!saved)
{
SaveSpecializedArray(bb, obj);
}
}
void Serialization<vtkm::cont::UnknownArrayHandle>::load(BinaryBuffer& bb,
vtkm::cont::UnknownArrayHandle& obj)
{
SerializedArrayType arrayType;
vtkmdiy::load(bb, arrayType);
switch (arrayType)
{
case SerializedArrayType::BasicArray:
{
std::string componentTypeString;
vtkmdiy::load(bb, componentTypeString);
bool loaded = false;
vtkm::ListForEach(
LoadBasicArray{}, vtkm::TypeListBaseC{}, bb, obj, componentTypeString, loaded);
if (!loaded)
{
throw vtkm::cont::ErrorInternal("Failed to load basic array. Unexpected buffer values.");
}
break;
}
case SerializedArrayType::SpecializedStorage:
LoadSpecializedArray(bb, obj);
break;
default:
throw vtkm::cont::ErrorInternal("Got inappropriate enumeration value for loading array.");
}
}
} // namespace mangled_diy_namespace

263
vtkm/cont/UnknownArrayHandle.h
View File

@ -17,6 +17,7 @@
#include <vtkm/cont/ArrayHandleCast.h>
#include <vtkm/cont/ArrayHandleMultiplexer.h>
#include <vtkm/cont/ArrayHandleRecombineVec.h>
#include <vtkm/cont/ArrayHandleRuntimeVec.h>
#include <vtkm/cont/ArrayHandleStride.h>
#include <vtkm/cont/StorageList.h>
@ -42,6 +43,14 @@ void UnknownAHDelete(void* mem)
delete arrayHandle;
}
template <typename T, typename S>
const std::vector<vtkm::cont::internal::Buffer>& UnknownAHBuffers(void* mem)
{
using AH = vtkm::cont::ArrayHandle<T, S>;
AH* arrayHandle = reinterpret_cast<AH*>(mem);
return arrayHandle->GetBuffers();
}
template <typename T, typename S>
void* UnknownAHNewInstance()
{
@ -56,49 +65,88 @@ vtkm::Id UnknownAHNumberOfValues(void* mem)
return arrayHandle->GetNumberOfValues();
}
template <typename T, typename StaticSize = typename vtkm::internal::SafeVecTraits<T>::IsSizeStatic>
struct UnknownAHNumberOfComponentsImpl;
template <typename T>
struct UnknownAHNumberOfComponentsImpl<T, vtkm::VecTraitsTagSizeStatic>
// Uses SFINAE to use Storage<>::GetNumberOfComponents if it exists, or the VecTraits otherwise
template <typename T, typename S>
inline auto UnknownAHNumberOfComponentsImpl(void* mem)
-> decltype(vtkm::cont::internal::Storage<T, S>::GetNumberOfComponents(
std::vector<vtkm::cont::internal::Buffer>()))
{
static constexpr vtkm::IdComponent Value = vtkm::internal::SafeVecTraits<T>::NUM_COMPONENTS;
};
template <typename T>
struct UnknownAHNumberOfComponentsImpl<T, vtkm::VecTraitsTagSizeVariable>
{
static constexpr vtkm::IdComponent Value = 0;
};
template <typename T>
vtkm::IdComponent UnknownAHNumberOfComponents()
{
return UnknownAHNumberOfComponentsImpl<T>::Value;
using AH = vtkm::cont::ArrayHandle<T, S>;
AH* arrayHandle = reinterpret_cast<AH*>(mem);
return vtkm::cont::internal::Storage<T, S>::GetNumberOfComponents(arrayHandle->GetBuffers());
}
template <typename T,
typename = typename vtkm::internal::SafeVecTraits<T>::IsSizeStatic,
typename = vtkm::HasVecTraits<T>>
struct UnknownAHNumberOfComponentsFlatImpl;
template <typename T>
struct UnknownAHNumberOfComponentsFlatImpl<T, vtkm::VecTraitsTagSizeStatic, std::true_type>
// Uses SFINAE to use the number of compnents in VecTraits.
// Note that this will conflict with the above overloaded function if the storage has a
// GetNumberOfComponents method and VecTraits has a static size. However, I cannot think
// of a use case for the storage to report the number of components for a static data type.
// If that happens, this implementation will need to be modified.
template <typename T, typename S>
inline auto UnknownAHNumberOfComponentsImpl(void*)
-> decltype(vtkm::internal::SafeVecTraits<T>::NUM_COMPONENTS)
{
static constexpr vtkm::IdComponent Value = vtkm::VecFlat<T>::NUM_COMPONENTS;
};
template <typename T>
struct UnknownAHNumberOfComponentsFlatImpl<T, vtkm::VecTraitsTagSizeVariable, std::true_type>
{
static constexpr vtkm::IdComponent Value = 0;
};
template <typename T>
struct UnknownAHNumberOfComponentsFlatImpl<T, vtkm::VecTraitsTagSizeStatic, std::false_type>
{
static constexpr vtkm::IdComponent Value = 1;
};
static constexpr vtkm::IdComponent numComponents =
vtkm::internal::SafeVecTraits<T>::NUM_COMPONENTS;
return numComponents;
}
template <typename T>
vtkm::IdComponent UnknownAHNumberOfComponentsFlat()
// Fallback for when there is no way to determine the number of components. (This could be
// because each value could have a different number of components.
template <typename T, typename S>
inline vtkm::IdComponent UnknownAHNumberOfComponentsImpl(...)
{
return UnknownAHNumberOfComponentsFlatImpl<T>::Value;
return 0;
}
template <typename T, typename S>
vtkm::IdComponent UnknownAHNumberOfComponents(void* mem)
{
return UnknownAHNumberOfComponentsImpl<T, S>(mem);
}
// Uses SFINAE to use Storage<>::GetNumberOfComponents if it exists, or the VecTraits otherwise
template <typename T, typename S>
inline auto UnknownAHNumberOfComponentsFlatImpl(void* mem)
-> decltype(vtkm::cont::internal::Storage<T, S>::GetNumberOfComponents(
std::vector<vtkm::cont::internal::Buffer>()))
{
using AH = vtkm::cont::ArrayHandle<T, S>;
AH* arrayHandle = reinterpret_cast<AH*>(mem);
// Making an assumption here that `T` is a `Vec`-like object that `GetNumberOfComponents`
// will report on how many each has. Further assuming that the components of `T` are
// static. If a future `ArrayHandle` type violates this, this code will have to become
// more complex.
return (vtkm::cont::internal::Storage<T, S>::GetNumberOfComponents(arrayHandle->GetBuffers()) *
vtkm::VecFlat<typename vtkm::internal::SafeVecTraits<T>::ComponentType>::NUM_COMPONENTS);
}
// Uses SFINAE to use the number of compnents in VecTraits.
// Note that this will conflict with the above overloaded function if the storage has a
// GetNumberOfComponents method and VecTraits has a static size. However, I cannot think
// of a use case for the storage to report the number of components for a static data type.
// If that happens, this implementation will need to be modified.
template <typename T, typename S>
inline auto UnknownAHNumberOfComponentsFlatImpl(void*)
-> decltype(vtkm::VecTraits<T>::NUM_COMPONENTS)
{
// static constexpr vtkm::IdComponent numComponents = vtkm::VecFlat<T>::NUM_COMPONENTS;
// return numComponents;
return vtkm::VecFlat<T>::NUM_COMPONENTS;
}
// Fallback for when there is no way to determine the number of components. (This could be
// because each value could have a different number of components or just that VecTraits
// are not defined.) Since it cannot be flattened, just return the same as num components.
template <typename T, typename S>
inline vtkm::IdComponent UnknownAHNumberOfComponentsFlatImpl(...)
{
return UnknownAHNumberOfComponentsImpl<T, S>(static_cast<void*>(nullptr));
}
template <typename T, typename S>
vtkm::IdComponent UnknownAHNumberOfComponentsFlat(void* mem)
{
return UnknownAHNumberOfComponentsFlatImpl<T, S>(mem);
}
template <typename T, typename S>
@ -220,6 +268,9 @@ struct VTKM_CONT_EXPORT UnknownAHContainer
using DeleteType = void(void*);
DeleteType* DeleteFunction;
using BuffersType = const std::vector<vtkm::cont::internal::Buffer>&(void*);
BuffersType* Buffers;
using NewInstanceType = void*();
NewInstanceType* NewInstance;
@ -230,7 +281,7 @@ struct VTKM_CONT_EXPORT UnknownAHContainer
using NumberOfValuesType = vtkm::Id(void*);
NumberOfValuesType* NumberOfValues;
using NumberOfComponentsType = vtkm::IdComponent();
using NumberOfComponentsType = vtkm::IdComponent(void*);
NumberOfComponentsType* NumberOfComponents;
NumberOfComponentsType* NumberOfComponentsFlat;
@ -344,12 +395,13 @@ inline UnknownAHContainer::UnknownAHContainer(const vtkm::cont::ArrayHandle<T, S
, BaseComponentType(
UnknownAHComponentInfo::Make<typename vtkm::internal::SafeVecTraits<T>::BaseComponentType>())
, DeleteFunction(detail::UnknownAHDelete<T, S>)
, Buffers(detail::UnknownAHBuffers<T, S>)
, NewInstance(detail::UnknownAHNewInstance<T, S>)
, NewInstanceBasic(detail::UnknownAHNewInstanceBasic<T>)
, NewInstanceFloatBasic(detail::UnknownAHNewInstanceFloatBasic<T>)
, NumberOfValues(detail::UnknownAHNumberOfValues<T, S>)
, NumberOfComponents(detail::UnknownAHNumberOfComponents<T>)
, NumberOfComponentsFlat(detail::UnknownAHNumberOfComponentsFlat<T>)
, NumberOfComponents(detail::UnknownAHNumberOfComponents<T, S>)
, NumberOfComponentsFlat(detail::UnknownAHNumberOfComponentsFlat<T, S>)
, Allocate(detail::UnknownAHAllocate<T, S>)
, ShallowCopy(detail::UnknownAHShallowCopy<T, S>)
, DeepCopy(detail::UnknownAHDeepCopy<T, S>)
@ -607,13 +659,10 @@ public:
#ifdef VTKM_MSVC
VTKM_DEPRECATED_SUPPRESS_BEGIN
#endif
///@{
/// Returns this array cast appropriately and stored in the given `ArrayHandle` type.
/// Throws an `ErrorBadType` if the stored array cannot be stored in the given array type.
/// Use the `CanConvert` method to determine if the array can be returned with the given type.
///
private:
template <typename T, typename S>
VTKM_CONT void AsArrayHandle(vtkm::cont::ArrayHandle<T, S>& array) const
VTKM_CONT void BaseAsArrayHandle(vtkm::cont::ArrayHandle<T, S>& array) const
{
using ArrayType = vtkm::cont::ArrayHandle<T, S>;
if (!this->IsType<ArrayType>())
@ -625,6 +674,21 @@ public:
array = *reinterpret_cast<ArrayType*>(this->Container->ArrayHandlePointer);
}
public:
///@{
/// Returns this array cast appropriately and stored in the given `ArrayHandle` type.
/// Throws an `ErrorBadType` if the stored array cannot be stored in the given array type.
/// Use the `CanConvert` method to determine if the array can be returned with the given type.
///
template <typename T, typename S>
VTKM_CONT void AsArrayHandle(vtkm::cont::ArrayHandle<T, S>& array) const
{
this->BaseAsArrayHandle(array);
}
template <typename T>
VTKM_CONT void AsArrayHandle(vtkm::cont::ArrayHandle<T>& array) const;
template <typename T, typename... Ss>
VTKM_CONT void AsArrayHandle(
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagMultiplexer<Ss...>>& array) const;
@ -638,6 +702,26 @@ public:
this->AsArrayHandle<ContainedArrayType>());
}
template <typename T>
VTKM_CONT void AsArrayHandle(
vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagRuntimeVec>& array) const
{
using BaseT = typename T::ComponentType;
if (this->IsStorageType<vtkm::cont::StorageTagBasic>() && this->IsBaseComponentType<BaseT>())
{
// Reinterpret the basic array as components, and then wrap that in a runtime vec
// with the correct amount of components.
vtkm::cont::ArrayHandle<BaseT, vtkm::cont::StorageTagBasic> basicArray(
this->Container->Buffers(this->Container->ArrayHandlePointer));
array =
vtkm::cont::ArrayHandleRuntimeVec<BaseT>(this->GetNumberOfComponentsFlat(), basicArray);
}
else
{
this->BaseAsArrayHandle(array);
}
}
template <typename ArrayType>
VTKM_CONT ArrayType AsArrayHandle() const
{
@ -877,6 +961,19 @@ struct UnknownArrayHandleCanConvert
}
};
template <typename T>
struct UnknownArrayHandleCanConvert<T, vtkm::cont::StorageTagBasic>
{
VTKM_CONT bool operator()(const vtkm::cont::UnknownArrayHandle& array) const
{
using UnrolledVec = vtkm::internal::UnrollVec<T>;
return (array.IsType<vtkm::cont::ArrayHandleBasic<T>>() ||
(array.IsStorageType<vtkm::cont::StorageTagRuntimeVec>() &&
array.IsBaseComponentType<typename UnrolledVec::ComponentType>() &&
UnrolledVec::NUM_COMPONENTS == array.GetNumberOfComponentsFlat()));
}
};
template <typename TargetT, typename SourceT, typename SourceS>
struct UnknownArrayHandleCanConvert<TargetT, vtkm::cont::StorageTagCast<SourceT, SourceS>>
{
@ -907,10 +1004,22 @@ struct UnknownArrayHandleCanConvert<T, vtkm::cont::StorageTagMultiplexer<Ss...>>
}
};
template <typename T>
struct UnknownArrayHandleCanConvert<T, vtkm::cont::StorageTagRuntimeVec>
{
VTKM_CONT bool operator()(const vtkm::cont::UnknownArrayHandle& array) const
{
using BaseComponentType = typename T::ComponentType;
return (array.IsType<vtkm::cont::ArrayHandle<T, vtkm::cont::StorageTagRuntimeVec>>() ||
(array.IsStorageType<vtkm::cont::StorageTagBasic>() &&
array.IsBaseComponentType<BaseComponentType>()));
}
};
} // namespace detail
template <typename ArrayHandleType>
VTKM_CONT bool UnknownArrayHandle::CanConvert() const
VTKM_CONT inline bool UnknownArrayHandle::CanConvert() const
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
@ -921,6 +1030,66 @@ VTKM_CONT bool UnknownArrayHandle::CanConvert() const
namespace detail
{
template <typename T,
vtkm::IdComponent = vtkm::internal::SafeVecTraits<
typename vtkm::internal::UnrollVec<T>::ComponentType>::NUM_COMPONENTS>
struct UnknownArrayHandleRuntimeVecAsBasic
{
VTKM_CONT bool operator()(const vtkm::cont::UnknownArrayHandle*,
const detail::UnknownAHContainer*,
vtkm::cont::ArrayHandle<T>&) const
{
// This version only gets called if T contains a `Vec`-like object that is not a strict `Vec`.
// This is rare but could happen. In this case, the type cannot be stored in an
// `ArrayHandleRuntimeVec` and therefore the load can never happen, so just ignore.
return false;
}
};
template <typename T>
struct UnknownArrayHandleRuntimeVecAsBasic<T, 1>
{
VTKM_CONT bool operator()(const vtkm::cont::UnknownArrayHandle* self,
const detail::UnknownAHContainer* container,
vtkm::cont::ArrayHandle<T>& array) const
{
using UnrolledVec = vtkm::internal::UnrollVec<T>;
using ComponentType = typename UnrolledVec::ComponentType;
if (self->IsStorageType<vtkm::cont::StorageTagRuntimeVec>() &&
self->IsBaseComponentType<ComponentType>() &&
UnrolledVec::NUM_COMPONENTS == self->GetNumberOfComponentsFlat())
{
// Pull out the components array out of the buffers. The array might not match exactly
// the array put in, but the buffer should still be consistent with the array (which works
// because the size of a basic array is based on the number of bytes in the buffer).
using RuntimeVecType = typename vtkm::cont::ArrayHandleRuntimeVec<ComponentType>::ValueType;
using StorageRuntimeVec =
vtkm::cont::internal::Storage<RuntimeVecType, vtkm::cont::StorageTagRuntimeVec>;
StorageRuntimeVec::AsArrayHandleBasic(container->Buffers(container->ArrayHandlePointer),
array);
return true;
}
else
{
return false;
}
}
};
} // namespace detail
template <typename T>
VTKM_CONT inline void UnknownArrayHandle::AsArrayHandle(vtkm::cont::ArrayHandle<T>& array) const
{
if (!detail::UnknownArrayHandleRuntimeVecAsBasic<T>{}(this, this->Container.get(), array))
{
this->BaseAsArrayHandle(array);
}
}
namespace detail
{
struct UnknownArrayHandleMultiplexerCastTry
{
template <typename T, typename S, typename... Ss>

15
vtkm/cont/internal/RuntimeDeviceConfiguration.cxx
View File

@ -90,11 +90,6 @@ void RuntimeDeviceConfigurationBase::Initialize(
[&](const vtkm::Id& value) { return this->SetThreads(value); },
"SetThreads",
this->GetDevice().GetName());
InitializeOption(
configOptions.VTKmNumaRegions,
[&](const vtkm::Id& value) { return this->SetNumaRegions(value); },
"SetNumaRegions",
this->GetDevice().GetName());
InitializeOption(
configOptions.VTKmDeviceInstance,
[&](const vtkm::Id& value) { return this->SetDeviceInstance(value); },
@ -117,11 +112,6 @@ RuntimeDeviceConfigReturnCode RuntimeDeviceConfigurationBase::SetThreads(const v
return RuntimeDeviceConfigReturnCode::INVALID_FOR_DEVICE;
}
RuntimeDeviceConfigReturnCode RuntimeDeviceConfigurationBase::SetNumaRegions(const vtkm::Id&)
{
return RuntimeDeviceConfigReturnCode::INVALID_FOR_DEVICE;
}
RuntimeDeviceConfigReturnCode RuntimeDeviceConfigurationBase::SetDeviceInstance(const vtkm::Id&)
{
return RuntimeDeviceConfigReturnCode::INVALID_FOR_DEVICE;
@ -132,11 +122,6 @@ RuntimeDeviceConfigReturnCode RuntimeDeviceConfigurationBase::GetThreads(vtkm::I
return RuntimeDeviceConfigReturnCode::INVALID_FOR_DEVICE;
}
RuntimeDeviceConfigReturnCode RuntimeDeviceConfigurationBase::GetNumaRegions(vtkm::Id&) const
{
return RuntimeDeviceConfigReturnCode::INVALID_FOR_DEVICE;
}
RuntimeDeviceConfigReturnCode RuntimeDeviceConfigurationBase::GetDeviceInstance(vtkm::Id&) const
{
return RuntimeDeviceConfigReturnCode::INVALID_FOR_DEVICE;

2
vtkm/cont/internal/RuntimeDeviceConfiguration.h
View File

@ -54,13 +54,11 @@ public:
/// A method should return INVALID_FOR_DEVICE if the overriden device does not
/// support the particular set method.
VTKM_CONT virtual RuntimeDeviceConfigReturnCode SetThreads(const vtkm::Id& value);
VTKM_CONT virtual RuntimeDeviceConfigReturnCode SetNumaRegions(const vtkm::Id& value);
VTKM_CONT virtual RuntimeDeviceConfigReturnCode SetDeviceInstance(const vtkm::Id& value);
/// The following public methods are overriden in each individual device and store the
/// values that were set via the above Set* methods for the given device.
VTKM_CONT virtual RuntimeDeviceConfigReturnCode GetThreads(vtkm::Id& value) const;
VTKM_CONT virtual RuntimeDeviceConfigReturnCode GetNumaRegions(vtkm::Id& value) const;
VTKM_CONT virtual RuntimeDeviceConfigReturnCode GetDeviceInstance(vtkm::Id& value) const;
/// The following public methods should be overriden as needed for each individual device

16
vtkm/cont/internal/RuntimeDeviceConfigurationOptions.cxx
View File

@ -31,13 +31,13 @@ void AppendOptionDescriptors(std::vector<option::Descriptor>& usage,
"vtkm-num-threads",
option::VtkmArg::Required,
" --vtkm-num-threads <dev> \tSets the number of threads to use for the selected device" });
usage.push_back(
{ useOptionIndex ? static_cast<uint32_t>(option::OptionIndex::NUMA_REGIONS) : 1,
0,
"",
"vtkm-numa-regions",
option::VtkmArg::Required,
" --vtkm-numa-regions <dev> \tSets the number of numa regions when using kokkos/OpenMP" });
usage.push_back({ useOptionIndex ? static_cast<uint32_t>(option::OptionIndex::NUMA_REGIONS) : 1,
0,
"",
"vtkm-numa-regions",
option::VtkmArg::Required,
" --vtkm-numa-regions <dev> \tSets the number of numa regions when using "
"kokkos/OpenMP (deprecated, has no effect)" });
usage.push_back(
{ useOptionIndex ? static_cast<uint32_t>(option::OptionIndex::DEVICE_INSTANCE) : 2,
0,
@ -51,7 +51,6 @@ void AppendOptionDescriptors(std::vector<option::Descriptor>& usage,
RuntimeDeviceConfigurationOptions::RuntimeDeviceConfigurationOptions(const bool& useOptionIndex)
: VTKmNumThreads(useOptionIndex ? option::OptionIndex::NUM_THREADS : 0, "VTKM_NUM_THREADS")
, VTKmNumaRegions(useOptionIndex ? option::OptionIndex::NUMA_REGIONS : 1, "VTKM_NUMA_REGIONS")
, VTKmDeviceInstance(useOptionIndex ? option::OptionIndex::DEVICE_INSTANCE : 2,
"VTKM_DEVICE_INSTANCE")
, Initialized(false)
@ -100,7 +99,6 @@ RuntimeDeviceConfigurationOptions::~RuntimeDeviceConfigurationOptions() noexcept
void RuntimeDeviceConfigurationOptions::Initialize(const option::Option* options)
{
this->VTKmNumThreads.Initialize(options);
this->VTKmNumaRegions.Initialize(options);
this->VTKmDeviceInstance.Initialize(options);
this->Initialized = true;
}

1
vtkm/cont/internal/RuntimeDeviceConfigurationOptions.h
View File

@ -48,7 +48,6 @@ public:
VTKM_CONT bool IsInitialized() const;
RuntimeDeviceOption VTKmNumThreads;
RuntimeDeviceOption VTKmNumaRegions;
RuntimeDeviceOption VTKmDeviceInstance;
protected:

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

@ -32,6 +32,21 @@ VTKM_THIRDPARTY_POST_INCLUDE
#include <type_traits>
#if KOKKOS_VERSION_MAJOR > 3 || (KOKKOS_VERSION_MAJOR == 3 && KOKKOS_VERSION_MINOR >= 7)
#define VTKM_VOLATILE
#else
#define VTKM_VOLATILE volatile
#endif
#if defined(VTKM_ENABLE_KOKKOS_THRUST) && (defined(__HIP__) || defined(__CUDA__))
#define VTKM_USE_KOKKOS_THRUST
#endif
#if defined(VTKM_USE_KOKKOS_THRUST)
#include <thrust/device_ptr.h>
#include <thrust/sort.h>
#endif
namespace vtkm
{
namespace internal
@ -256,7 +271,7 @@ private:
}
KOKKOS_INLINE_FUNCTION
void join(volatile value_type& dst, const volatile value_type& src) const
void join(VTKM_VOLATILE value_type& dst, const VTKM_VOLATILE value_type& src) const
{
dst = this->Operator(dst, src);
}
@ -726,6 +741,13 @@ private:
template <typename T>
VTKM_CONT static void SortImpl(vtkm::cont::ArrayHandle<T>& values, vtkm::SortLess, std::true_type)
{
// In Kokkos 3.7, we have noticed some errors when sorting with zero-length arrays (which
// should do nothing). There is no check, and the bin size computation gets messed up.
if (values.GetNumberOfValues() <= 1)
{
return;
}
vtkm::cont::Token token;
auto portal = values.PrepareForInPlace(vtkm::cont::DeviceAdapterTagKokkos{}, token);
kokkos::internal::KokkosViewExec<T> view(portal.GetArray(), portal.GetNumberOfValues());
@ -757,6 +779,88 @@ public:
SortImpl(values, comp, typename std::is_scalar<T>::type{});
}
protected:
// Kokkos currently (11/10/2022) does not support a sort_by_key operator
// so instead we are using thrust if and only if HIP or CUDA are the backends for Kokkos
#if defined(VTKM_USE_KOKKOS_THRUST)
template <typename T, typename U, typename BinaryCompare>
VTKM_CONT static std::enable_if_t<(std::is_same<BinaryCompare, vtkm::SortLess>::value ||
std::is_same<BinaryCompare, vtkm::SortGreater>::value)>
SortByKeyImpl(vtkm::cont::ArrayHandle<T>& keys,
vtkm::cont::ArrayHandle<U>& values,
BinaryCompare,
std::true_type,
std::true_type)
{
vtkm::cont::Token token;
auto keys_portal = keys.PrepareForInPlace(vtkm::cont::DeviceAdapterTagKokkos{}, token);
auto values_portal = values.PrepareForInPlace(vtkm::cont::DeviceAdapterTagKokkos{}, token);
kokkos::internal::KokkosViewExec<T> keys_view(keys_portal.GetArray(),
keys_portal.GetNumberOfValues());
kokkos::internal::KokkosViewExec<U> values_view(values_portal.GetArray(),
values_portal.GetNumberOfValues());
thrust::device_ptr<T> keys_begin(keys_view.data());
thrust::device_ptr<T> keys_end(keys_view.data() + keys_view.size());
thrust::device_ptr<U> values_begin(values_view.data());
if (std::is_same<BinaryCompare, vtkm::SortLess>::value)
{
thrust::sort_by_key(keys_begin, keys_end, values_begin, thrust::less<T>());
}
else
{
thrust::sort_by_key(keys_begin, keys_end, values_begin, thrust::greater<T>());
}
}
#endif
template <typename T,
typename U,
class StorageT,
class StorageU,
class BinaryCompare,
typename ValidKeys,
typename ValidValues>
VTKM_CONT static void SortByKeyImpl(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values,
BinaryCompare binary_compare,
ValidKeys,
ValidValues)
{
// Default to general algorithm
Superclass::SortByKey(keys, values, binary_compare);
}
public:
template <typename T, typename U, class StorageT, class StorageU>
VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values)
{
// Make sure not to use the general algorithm here since
// it will use Sort algorithm instead of SortByKey
SortByKey(keys, values, internal::DefaultCompareFunctor());
}
template <typename T, typename U, class StorageT, class StorageU, class BinaryCompare>
VTKM_CONT static void SortByKey(vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values,
BinaryCompare binary_compare)
{
// If T or U are not scalar types, or the BinaryCompare is not supported
// then the general algorithm is called, otherwise we will run thrust
SortByKeyImpl(keys,
values,
binary_compare,
typename std::is_scalar<T>::type{},
typename std::is_scalar<U>::type{});
}
//----------------------------------------------------------------------------
VTKM_CONT static void Synchronize()
{
vtkm::cont::kokkos::internal::GetExecutionSpaceInstance().fence();
@ -787,4 +891,6 @@ public:
}
} // namespace vtkm::cont
#undef VTKM_VOLATILE
#endif //vtk_m_cont_kokkos_internal_DeviceAdapterAlgorithmKokkos_h

27
vtkm/cont/kokkos/internal/RuntimeDeviceConfigurationKokkos.h
View File

@ -95,22 +95,7 @@ public:
return RuntimeDeviceConfigReturnCode::NOT_APPLIED;
}
this->KokkosArguments.insert(this->KokkosArguments.begin(),
"--kokkos-threads=" + std::to_string(value));
return RuntimeDeviceConfigReturnCode::SUCCESS;
}
VTKM_CONT virtual RuntimeDeviceConfigReturnCode SetNumaRegions(
const vtkm::Id& value) override final
{
if (Kokkos::is_initialized())
{
VTKM_LOG_S(vtkm::cont::LogLevel::Warn,
"SetNumaRegions was called but Kokkos was already initailized! Updates will not "
"be applied.");
return RuntimeDeviceConfigReturnCode::NOT_APPLIED;
}
this->KokkosArguments.insert(this->KokkosArguments.begin(),
"--kokkos-numa=" + std::to_string(value));
"--kokkos-num-threads=" + std::to_string(value));
return RuntimeDeviceConfigReturnCode::SUCCESS;
}
@ -131,13 +116,7 @@ public:
VTKM_CONT virtual RuntimeDeviceConfigReturnCode GetThreads(vtkm::Id& value) const override final
{
return GetArgFromList(this->KokkosArguments, "--kokkos-threads", value);
}
VTKM_CONT virtual RuntimeDeviceConfigReturnCode GetNumaRegions(
vtkm::Id& value) const override final
{
return GetArgFromList(this->KokkosArguments, "--kokkos-numa", value);
return GetArgFromList(this->KokkosArguments, "--kokkos-num-threads", value);
}
VTKM_CONT virtual RuntimeDeviceConfigReturnCode GetDeviceInstance(
@ -149,7 +128,7 @@ public:
protected:
/// Store a copy of the current arguments when initializing the Kokkos subsystem later
/// Appends a copy of the argv values in the KokkosArguments vector: this assumes the
/// argv values contain kokkos command line arguments (like --kokkos-threads, etc)
/// argv values contain kokkos command line arguments (like --kokkos-num-threads, etc)
VTKM_CONT virtual void ParseExtraArguments(int& argc, char* argv[]) override final
{
if (argc > 0 && argv)

39
vtkm/cont/kokkos/testing/UnitTestKokkosRuntimeDeviceConfiguration.cxx
View File

@ -26,22 +26,8 @@ TestingRuntimeDeviceConfiguration<vtkm::cont::DeviceAdapterTagKokkos>::TestRunti
{
int argc;
char** argv;
vtkm::cont::testing::Testing::MakeArgs(argc, argv, "--kokkos-numa=4");
vtkm::cont::testing::Testing::SetEnv("KOKKOS_DEVICE_ID", "0");
vtkm::cont::testing::Testing::MakeArgs(argc, argv, "--kokkos-print-configuration");
auto deviceOptions = TestingRuntimeDeviceConfiguration::DefaultInitializeConfigOptions();
bool threw = false;
try
{
RuntimeDeviceInformation{}.GetRuntimeConfiguration(
DeviceAdapterTagKokkos(), deviceOptions, argc, argv);
}
catch (const std::runtime_error& e)
{
threw = true;
}
VTKM_TEST_ASSERT(threw,
"GetRuntimeConfiguration should have thrown, env KOKKOS_DEVICE_ID didn't match");
VTKM_TEST_ASSERT(!Kokkos::is_initialized(), "Kokkos should not be initialized at this point");
deviceOptions.VTKmDeviceInstance.SetOption(0);
internal::RuntimeDeviceConfigurationBase& config =
RuntimeDeviceInformation{}.GetRuntimeConfiguration(
@ -54,28 +40,21 @@ TestingRuntimeDeviceConfiguration<vtkm::cont::DeviceAdapterTagKokkos>::TestRunti
"Failed to get set threads");
VTKM_TEST_ASSERT(testValue == 8,
"Set threads does not match expected value: 8 != " + std::to_string(testValue));
VTKM_TEST_ASSERT(config.GetNumaRegions(testValue) ==
internal::RuntimeDeviceConfigReturnCode::SUCCESS,
"Failed to get set numa regions");
VTKM_TEST_ASSERT(testValue == 4,
"Set numa regions does not match expected value: 4 != " +
std::to_string(testValue));
VTKM_TEST_ASSERT(config.GetDeviceInstance(testValue) ==
internal::RuntimeDeviceConfigReturnCode::SUCCESS,
"Failed to get set device instance");
VTKM_TEST_ASSERT(testValue == 0,
"Set device instance does not match expected value: 0 != " +
std::to_string(testValue));
// Ensure that with kokkos we can't re-initialize or set values after the first initialize
// Should pop up a few warnings in the test logs
std::cout
<< "Ensure that with kokkos we can't re-initialize or set values after the first initialize"
<< std::endl;
std::cout << "This should pop up a few warnings in the test logs" << std::endl;
deviceOptions.VTKmNumThreads.SetOption(16);
deviceOptions.VTKmNumaRegions.SetOption(2);
deviceOptions.VTKmDeviceInstance.SetOption(5);
config.Initialize(deviceOptions);
VTKM_TEST_ASSERT(config.SetThreads(1) == internal::RuntimeDeviceConfigReturnCode::NOT_APPLIED,
"Shouldn't be able to set threads after kokkos is initalized");
VTKM_TEST_ASSERT(config.SetNumaRegions(1) == internal::RuntimeDeviceConfigReturnCode::NOT_APPLIED,
"Shouldn't be able to set numa regions after kokkos is initalized");
VTKM_TEST_ASSERT(config.SetDeviceInstance(1) ==
internal::RuntimeDeviceConfigReturnCode::NOT_APPLIED,
"Shouldn't be able to set device instnace after kokkos is initalized");
@ -85,20 +64,12 @@ TestingRuntimeDeviceConfiguration<vtkm::cont::DeviceAdapterTagKokkos>::TestRunti
"Failed to get set threads");
VTKM_TEST_ASSERT(testValue == 8,
"Set threads does not match expected value: 8 != " + std::to_string(testValue));
VTKM_TEST_ASSERT(config.GetNumaRegions(testValue) ==
internal::RuntimeDeviceConfigReturnCode::SUCCESS,
"Failed to get set numa regions");
VTKM_TEST_ASSERT(testValue == 4,
"Set numa regions does not match expected value: 4 != " +
std::to_string(testValue));
VTKM_TEST_ASSERT(config.GetDeviceInstance(testValue) ==
internal::RuntimeDeviceConfigReturnCode::SUCCESS,
"Failed to get set device instance");
VTKM_TEST_ASSERT(testValue == 0,
"Set device instance does not match expected value: 0 != " +
std::to_string(testValue));
vtkm::cont::testing::Testing::UnsetEnv("KOKKOS_DEVICE_ID");
}
} // namespace vtkm::cont::testing

7
vtkm/cont/testing/CMakeLists.txt
View File

@ -85,6 +85,7 @@ set(unit_tests_device
UnitTestArrayHandleRandomStandardNormal.cxx
UnitTestArrayHandleRandomUniformReal.cxx
UnitTestArrayHandleRecombineVec.cxx
UnitTestArrayHandleRuntimeVec.cxx
UnitTestArrayHandleSOA.cxx
UnitTestArrayHandleSwizzle.cxx
UnitTestArrayHandleTransform.cxx
@ -122,6 +123,12 @@ if(TARGET vtkm_filter_field_conversion)
)
endif()
if(TARGET vtkm_filter_contour)
list(APPEND unit_tests
UnitTestAbort.cxx
)
endif()
vtkm_unit_tests(SOURCES ${unit_tests} DEVICE_SOURCES ${unit_tests_device})
#add distributed tests i.e.test to run with MPI

9
vtkm/cont/testing/MakeTestDataSet.h
View File

@ -11,8 +11,15 @@
#ifndef vtk_m_cont_testing_MakeTestDataSet_h
#define vtk_m_cont_testing_MakeTestDataSet_h
// The relative path of Testing.h is unknown, the only thing that we can assume
// is that it is located in the same directory as this header file. This is
// because the testing directory is reserved for test executables and not
// libraries, the vtkm_cont_testing module has to put this file in
// vtkm/cont/testlib instead of vtkm/cont/testing where you normally would
// expect it.
#include "Testing.h"
#include <vtkm/cont/DataSet.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/cont/testlib/vtkm_cont_testing_export.h>

1
vtkm/cont/testing/TestingRuntimeDeviceConfiguration.h
View File

@ -35,7 +35,6 @@ struct TestingRuntimeDeviceConfiguration
{
internal::RuntimeDeviceConfigurationOptions runtimeDeviceOptions{};
runtimeDeviceOptions.VTKmNumThreads.SetOption(8);
runtimeDeviceOptions.VTKmNumaRegions.SetOption(0);
runtimeDeviceOptions.VTKmDeviceInstance.SetOption(2);
runtimeDeviceOptions.Initialize(nullptr);
VTKM_TEST_ASSERT(runtimeDeviceOptions.IsInitialized(),

105
vtkm/cont/testing/UnitTestAbort.cxx Normal file
View File

@ -0,0 +1,105 @@
//============================================================================
// 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/cont/ErrorUserAbort.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
#include <vtkm/cont/testing/Testing.h>
#include <vtkm/filter/contour/Contour.h>
#include <vtkm/source/Wavelet.h>
#include <sstream>
namespace
{
// A function that checks for abort request.
// This function will be called by `TryExecute` befaure lauching a device task
// to check if abort has been requested.
// For this test case, we are using a simple logic of returning true for the
// `abortAt`th check.
// If this test is failing, one of the things to check would be to see if the
// `Contour` filter has changed such that it no longer has atleast `abortAt`
// task invocations.
bool ShouldAbort()
{
static int abortCheckCounter = 0;
static constexpr int abortAt = 5;
if (++abortCheckCounter >= abortAt)
{
std::cout << "Abort check " << abortCheckCounter << ": true\n";
return true;
}
std::cout << "Abort check " << abortCheckCounter << ": false\n";
return false;
}
int TestAbort()
{
vtkm::source::Wavelet wavelet;
wavelet.SetExtent(vtkm::Id3(-15), vtkm::Id3(16));
auto input = wavelet.Execute();
auto range = input.GetField("RTData").GetRange().ReadPortal().Get(0);
std::vector<vtkm::Float64> isovals;
static constexpr int numDivs = 5;
for (int i = 1; i < numDivs - 1; ++i)
{
auto v = range.Min +
(static_cast<vtkm::Float64>(i) *
((range.Max - range.Min) / static_cast<vtkm::Float64>(numDivs)));
isovals.push_back(v);
}
vtkm::filter::contour::Contour contour;
contour.SetActiveField("RTData");
contour.SetIsoValues(isovals);
// First we will run the filter with the abort function set
std::cout << "Run #1 with the abort function set\n";
try
{
vtkm::cont::ScopedRuntimeDeviceTracker tracker(ShouldAbort);
auto result = contour.Execute(input);
// execution shouldn't reach here
VTKM_TEST_FAIL("Error: filter execution was not aborted. Result: ",
result.GetNumberOfPoints(),
" points and ",
result.GetNumberOfCells(),
" triangles");
}
catch (const vtkm::cont::ErrorUserAbort&)
{
std::cout << "Execution was successfully aborted\n";
}
// Now run the filter without the abort function
std::cout << "Run #2 without the abort function set\n";
try
{
auto result = contour.Execute(input);
std::cout << "Success: filter execution was not aborted. Result: " << result.GetNumberOfPoints()
<< " points and " << result.GetNumberOfCells() << " triangles\n";
}
catch (const vtkm::cont::ErrorUserAbort&)
{
VTKM_TEST_FAIL("Execution was unexpectedly aborted");
}
return 0;
}
} // anon namespace
int UnitTestAbort(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(TestAbort, argc, argv);
}

147
vtkm/cont/testing/UnitTestArrayExtractComponent.cxx
View File

@ -17,6 +17,7 @@
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/ArrayHandleMultiplexer.h>
#include <vtkm/cont/ArrayHandleReverse.h>
#include <vtkm/cont/ArrayHandleRuntimeVec.h>
#include <vtkm/cont/ArrayHandleSOA.h>
#include <vtkm/cont/ArrayHandleUniformPointCoordinates.h>
#include <vtkm/cont/ArrayHandleView.h>
@ -25,88 +26,132 @@
#include <vtkm/cont/testing/Testing.h>
#include <algorithm>
#include <random>
#include <vector>
namespace
{
constexpr vtkm::Id ARRAY_SIZE = 10;
template <typename T>
vtkm::IdComponent GetTotalNumComponents(const T& vec)
{
using VTraits = vtkm::VecTraits<T>;
if (std::is_same<typename VTraits::ComponentType, typename VTraits::BaseComponentType>::value)
{
return VTraits::GetNumberOfComponents(vec);
}
else
{
return VTraits::GetNumberOfComponents(vec) *
GetTotalNumComponents(VTraits::GetComponent(vec, 0));
}
}
// VecFlat.h has something similar, but it only works with static Vec sizes. It might make sense
// to move this somewhere else later
template <typename BaseComponentType>
struct GetVecFlatIndexImpl
{
template <typename VecType>
VTKM_CONT BaseComponentType operator()(const VecType& vec, vtkm::IdComponent index) const
{
const vtkm::IdComponent subSize = GetTotalNumComponents(vec[0]);
return (*this)(vec[index / subSize], index % subSize);
}
VTKM_CONT BaseComponentType operator()(const BaseComponentType& component,
vtkm::IdComponent index) const
{
VTKM_ASSERT(index == 0);
return component;
}
};
template <typename T>
auto GetVecFlatIndex(const T& vec, vtkm::IdComponent index)
{
return GetVecFlatIndexImpl<typename vtkm::VecTraits<T>::BaseComponentType>{}(vec, index);
}
template <typename T, typename S>
void CheckInputArray(const vtkm::cont::ArrayHandle<T, S>& originalArray,
vtkm::CopyFlag allowCopy = vtkm::CopyFlag::Off)
{
//std::cout << " Checking input array type "
// << vtkm::cont::TypeToString<vtkm::cont::ArrayHandle<T, S>>() << std::endl;
//std::cout << " Original array: ";
//vtkm::cont::printSummary_ArrayHandle(originalArray, std::cout);
using FlatVec = vtkm::VecFlat<T>;
using ComponentType = typename FlatVec::ComponentType;
for (vtkm::IdComponent componentId = 0; componentId < FlatVec::NUM_COMPONENTS; ++componentId)
auto originalPortal = originalArray.ReadPortal();
using ComponentType = typename vtkm::VecTraits<T>::BaseComponentType;
const vtkm::IdComponent numComponents = GetTotalNumComponents(originalPortal.Get(0));
for (vtkm::IdComponent componentId = 0; componentId < numComponents; ++componentId)
{
vtkm::cont::ArrayHandleStride<ComponentType> componentArray =
vtkm::cont::ArrayExtractComponent(originalArray, componentId, allowCopy);
//std::cout << " Component " << componentId << ": ";
//vtkm::cont::printSummary_ArrayHandle(componentArray, std::cout);
auto originalPortal = originalArray.ReadPortal();
auto componentPortal = componentArray.ReadPortal();
VTKM_TEST_ASSERT(originalPortal.GetNumberOfValues() == componentPortal.GetNumberOfValues());
for (vtkm::Id arrayIndex = 0; arrayIndex < originalArray.GetNumberOfValues(); ++arrayIndex)
{
auto originalValue = vtkm::make_VecFlat(originalPortal.Get(arrayIndex));
auto originalValue = GetVecFlatIndex(originalPortal.Get(arrayIndex), componentId);
ComponentType componentValue = componentPortal.Get(arrayIndex);
VTKM_TEST_ASSERT(test_equal(originalValue[componentId], componentValue));
VTKM_TEST_ASSERT(test_equal(originalValue, componentValue));
}
}
}
template <typename T, typename S>
void CheckOutputArray(const vtkm::cont::ArrayHandle<T, S>& originalArray)
void CheckOutputArray(
const vtkm::cont::ArrayHandle<T, S>& originalArray,
const vtkm::cont::ArrayHandle<T, S>& outputArray = vtkm::cont::ArrayHandle<T, S>{})
{
CheckInputArray(originalArray);
//std::cout << " Checking output array type "
// << vtkm::cont::TypeToString<vtkm::cont::ArrayHandle<T, S>>() << std::endl;
using ComponentType = typename vtkm::VecTraits<T>::BaseComponentType;
const vtkm::IdComponent numComponents = GetTotalNumComponents(originalArray.ReadPortal().Get(0));
//std::cout << " Original array: ";
//vtkm::cont::printSummary_ArrayHandle(originalArray, std::cout);
vtkm::cont::ArrayHandle<T, S> outputArray;
outputArray.Allocate(originalArray.GetNumberOfValues());
using FlatVec = vtkm::VecFlat<T>;
using ComponentType = typename FlatVec::ComponentType;
constexpr vtkm::IdComponent numComponents = FlatVec::NUM_COMPONENTS;
// Extract all the stride arrays first, and then allocate them later. This tests to
// to make sure that the independent allocation of all the extracted arrays are consistent
// and correct.
std::vector<std::pair<vtkm::cont::ArrayHandleStride<ComponentType>,
vtkm::cont::ArrayHandleStride<ComponentType>>>
componentArrays;
componentArrays.reserve(static_cast<std::size_t>(numComponents));
for (vtkm::IdComponent componentId = 0; componentId < numComponents; ++componentId)
{
vtkm::cont::ArrayHandleStride<ComponentType> inComponentArray =
vtkm::cont::ArrayExtractComponent(originalArray, numComponents - componentId - 1);
vtkm::cont::ArrayHandleStride<ComponentType> outComponentArray =
vtkm::cont::ArrayExtractComponent(outputArray, componentId, vtkm::CopyFlag::Off);
componentArrays.emplace_back(
vtkm::cont::ArrayExtractComponent(originalArray, numComponents - componentId - 1),
vtkm::cont::ArrayExtractComponent(outputArray, componentId, vtkm::CopyFlag::Off));
}
auto inPortal = inComponentArray.ReadPortal();
auto outPortal = outComponentArray.WritePortal();
VTKM_TEST_ASSERT(inComponentArray.GetNumberOfValues() == originalArray.GetNumberOfValues());
VTKM_TEST_ASSERT(outComponentArray.GetNumberOfValues() == originalArray.GetNumberOfValues());
// Shuffle the component arrays to ensure the allocation/copy can occur in any order.
std::random_device rd;
std::default_random_engine rng(rd());
std::shuffle(componentArrays.begin(), componentArrays.end(), rng);
for (auto& inOutArrays : componentArrays)
{
inOutArrays.second.Allocate(originalArray.GetNumberOfValues());
auto inPortal = inOutArrays.first.ReadPortal();
auto outPortal = inOutArrays.second.WritePortal();
VTKM_TEST_ASSERT(inPortal.GetNumberOfValues() == originalArray.GetNumberOfValues());
VTKM_TEST_ASSERT(outPortal.GetNumberOfValues() == originalArray.GetNumberOfValues());
for (vtkm::Id arrayIndex = 0; arrayIndex < originalArray.GetNumberOfValues(); ++arrayIndex)
{
outPortal.Set(arrayIndex, inPortal.Get(arrayIndex));
}
}
//std::cout << " Output array: ";
//vtkm::cont::printSummary_ArrayHandle(outputArray, std::cout);
auto inPortal = originalArray.ReadPortal();
auto outPortal = outputArray.ReadPortal();
for (vtkm::Id arrayIndex = 0; arrayIndex < originalArray.GetNumberOfValues(); ++arrayIndex)
{
FlatVec inValue = vtkm::make_VecFlat(inPortal.Get(arrayIndex));
FlatVec outValue = vtkm::make_VecFlat(outPortal.Get(arrayIndex));
auto inValue = inPortal.Get(arrayIndex);
auto outValue = outPortal.Get(arrayIndex);
for (vtkm::IdComponent componentId = 0; componentId < numComponents; ++componentId)
{
VTKM_TEST_ASSERT(test_equal(inValue[componentId], outValue[numComponents - componentId - 1]));
VTKM_TEST_ASSERT(test_equal(GetVecFlatIndex(inValue, componentId),
GetVecFlatIndex(outValue, numComponents - componentId - 1)));
}
}
}
@ -146,6 +191,7 @@ void DoTest()
constexpr vtkm::Id STRIDE = 7;
vtkm::cont::ArrayHandleBasic<vtkm::Vec3f> originalArray;
originalArray.Allocate(ARRAY_SIZE * STRIDE);
SetPortal(originalArray.WritePortal());
for (vtkm::Id offset = 0; offset < STRIDE; ++offset)
{
vtkm::cont::ArrayHandleStride<vtkm::Vec3f> strideArray(
@ -157,9 +203,10 @@ void DoTest()
{
std::cout << "ArrayHandleGroupVec" << std::endl;
vtkm::cont::ArrayHandle<vtkm::Vec3f> array;
array.Allocate(ARRAY_SIZE * 2);
array.Allocate(ARRAY_SIZE * 4);
SetPortal(array.WritePortal());
CheckOutputArray(vtkm::cont::make_ArrayHandleGroupVec<2>(array));
CheckOutputArray(vtkm::cont::make_ArrayHandleGroupVec<4>(array));
}
{
@ -173,7 +220,23 @@ void DoTest()
auto compositeArray = vtkm::cont::make_ArrayHandleCompositeVector(array0, array1);
CheckOutputArray(compositeArray);
CheckOutputArray(vtkm::cont::make_ArrayHandleExtractComponent(compositeArray, 1));
// Note that when the extracted component array gets allocated, it only allocates the
// array it was given. This is a weird case when using `ArrayHandleExtractComponent`
// on something that has multiple arrays as input. It works fine if all components get
// extracted and updated, but can cause issues if only one is resized. In this case
// just test the input.
CheckInputArray(vtkm::cont::make_ArrayHandleExtractComponent(compositeArray, 1));
}
{
std::cout << "ArrayHandleRuntimeVec" << std::endl;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> array;
array.Allocate(ARRAY_SIZE * 4);
SetPortal(array.WritePortal());
CheckOutputArray(vtkm::cont::make_ArrayHandleRuntimeVec(2, array),
vtkm::cont::ArrayHandleRuntimeVec<vtkm::FloatDefault>(2));
CheckOutputArray(vtkm::cont::make_ArrayHandleRuntimeVec(4, array),
vtkm::cont::ArrayHandleRuntimeVec<vtkm::FloatDefault>(4));
}
{

9
vtkm/cont/testing/UnitTestArrayHandleRecombineVec.cxx
View File

@ -10,6 +10,7 @@
#include <vtkm/cont/ArrayHandleRecombineVec.h>
#include <vtkm/cont/ArrayHandleReverse.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/worklet/WorkletMapField.h>
@ -69,7 +70,6 @@ struct TestRecombineVecAsOutput
SetPortal(baseArray.WritePortal());
vtkm::cont::ArrayHandle<T> outputArray;
outputArray.Allocate(ARRAY_SIZE); // Cannot resize after recombine
using VTraits = vtkm::VecTraits<T>;
vtkm::cont::ArrayHandleRecombineVec<typename VTraits::ComponentType> recombinedArray;
@ -78,12 +78,15 @@ struct TestRecombineVecAsOutput
recombinedArray.AppendComponentArray(vtkm::cont::ArrayExtractComponent(outputArray, cIndex));
}
VTKM_TEST_ASSERT(recombinedArray.GetNumberOfComponents() == VTraits::NUM_COMPONENTS);
VTKM_TEST_ASSERT(recombinedArray.GetNumberOfValues() == ARRAY_SIZE);
vtkm::cont::Invoker invoke;
invoke(PassThrough{}, baseArray, recombinedArray);
VTKM_TEST_ASSERT(test_equal_ArrayHandles(baseArray, outputArray));
// Try outputing to a recombine vec inside of another fancy ArrayHandle.
auto reverseOutput = vtkm::cont::make_ArrayHandleReverse(recombinedArray);
invoke(PassThrough{}, baseArray, reverseOutput);
VTKM_TEST_ASSERT(test_equal_ArrayHandles(baseArray, reverseOutput));
}
};

248
vtkm/cont/testing/UnitTestArrayHandleRuntimeVec.cxx Normal file
View File

@ -0,0 +1,248 @@
//============================================================================
// 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/cont/ArrayHandleRuntimeVec.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/Invoker.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/cont/testing/Testing.h>
namespace
{
constexpr vtkm::Id ARRAY_SIZE = 10;
struct UnusualType
{
vtkm::Id X;
};
} // anonymous namespace
namespace detail
{
template <>
struct TestValueImpl<UnusualType>
{
VTKM_EXEC_CONT UnusualType operator()(vtkm::Id index) const
{
return { TestValue(index, decltype(UnusualType::X){}) };
}
};
template <>
struct TestEqualImpl<UnusualType, UnusualType>
{
VTKM_EXEC_CONT bool operator()(UnusualType value1,
UnusualType value2,
vtkm::Float64 tolerance) const
{
return test_equal(value1.X, value2.X, tolerance);
}
};
} // namespace detail
namespace
{
struct PassThrough : vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn, FieldOut);
using ExecutionSignature = void(_1, _2);
template <typename InValue, typename OutValue>
VTKM_EXEC void operator()(const InValue& inValue, OutValue& outValue) const
{
vtkm::IdComponent inIndex = 0;
vtkm::IdComponent outIndex = 0;
this->FlatCopy(inValue, inIndex, outValue, outIndex);
}
template <typename InValue, typename OutValue>
VTKM_EXEC void FlatCopy(const InValue& inValue,
vtkm::IdComponent& inIndex,
OutValue& outValue,
vtkm::IdComponent& outIndex) const
{
using VTraitsIn = vtkm::internal::SafeVecTraits<InValue>;
using VTraitsOut = vtkm::internal::SafeVecTraits<OutValue>;
VTraitsOut::SetComponent(outValue, outIndex, VTraitsIn::GetComponent(inValue, inIndex));
inIndex++;
outIndex++;
}
template <typename InComponent, vtkm::IdComponent InN, typename OutValue>
VTKM_EXEC void FlatCopy(const vtkm::Vec<InComponent, InN>& inValue,
vtkm::IdComponent& inIndex,
OutValue& outValue,
vtkm::IdComponent& outIndex) const
{
VTKM_ASSERT(inIndex == 0);
for (vtkm::IdComponent i = 0; i < InN; ++i)
{
FlatCopy(inValue[i], inIndex, outValue, outIndex);
inIndex = 0;
}
}
template <typename InValue, typename OutComponent, vtkm::IdComponent OutN>
VTKM_EXEC void FlatCopy(const InValue& inValue,
vtkm::IdComponent& inIndex,
vtkm::Vec<OutComponent, OutN>& outValue,
vtkm::IdComponent& outIndex) const
{
VTKM_ASSERT(outIndex == 0);
for (vtkm::IdComponent i = 0; i < OutN; ++i)
{
OutComponent outComponent;
FlatCopy(inValue, inIndex, outComponent, outIndex);
outValue[i] = outComponent;
outIndex = 0;
}
}
};
template <vtkm::IdComponent NUM_COMPONENTS>
struct TestRuntimeVecAsInput
{
template <typename ComponentType>
VTKM_CONT void operator()(ComponentType) const
{
using ValueType = vtkm::Vec<ComponentType, NUM_COMPONENTS>;
vtkm::cont::ArrayHandle<ComponentType> baseArray;
baseArray.Allocate(ARRAY_SIZE * NUM_COMPONENTS);
SetPortal(baseArray.WritePortal());
auto runtimeVecArray = vtkm::cont::make_ArrayHandleRuntimeVec(NUM_COMPONENTS, baseArray);
VTKM_TEST_ASSERT(runtimeVecArray.GetNumberOfValues() == ARRAY_SIZE,
"Group array reporting wrong array size.");
vtkm::cont::ArrayHandle<ValueType> resultArray;
vtkm::cont::Invoker{}(PassThrough{}, runtimeVecArray, resultArray);
VTKM_TEST_ASSERT(resultArray.GetNumberOfValues() == ARRAY_SIZE, "Got bad result array size.");
//verify that the control portal works
vtkm::Id totalIndex = 0;
auto resultPortal = resultArray.ReadPortal();
for (vtkm::Id index = 0; index < ARRAY_SIZE; ++index)
{
const ValueType result = resultPortal.Get(index);
for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++)
{
const ComponentType expectedValue = TestValue(totalIndex, ComponentType());
VTKM_TEST_ASSERT(test_equal(result[componentIndex], expectedValue),
"Result array got wrong value.");
totalIndex++;
}
}
//verify that you can get the data as a basic array
vtkm::cont::ArrayHandle<vtkm::Vec<ComponentType, NUM_COMPONENTS>> flatComponents;
runtimeVecArray.AsArrayHandleBasic(flatComponents);
VTKM_TEST_ASSERT(test_equal_ArrayHandles(
flatComponents, vtkm::cont::make_ArrayHandleGroupVec<NUM_COMPONENTS>(baseArray)));
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Vec<ComponentType, 1>, NUM_COMPONENTS>>
nestedComponents;
runtimeVecArray.AsArrayHandleBasic(nestedComponents);
auto flatPortal = flatComponents.ReadPortal();
auto nestedPortal = nestedComponents.ReadPortal();
for (vtkm::Id index = 0; index < flatPortal.GetNumberOfValues(); ++index)
{
VTKM_TEST_ASSERT(test_equal(vtkm::make_VecFlat(flatPortal.Get(index)),
vtkm::make_VecFlat(nestedPortal.Get(index))));
}
runtimeVecArray.ReleaseResources();
}
};
template <vtkm::IdComponent NUM_COMPONENTS>
struct TestRuntimeVecAsOutput
{
template <typename ComponentType>
VTKM_CONT void operator()(ComponentType) const
{
using ValueType = vtkm::Vec<ComponentType, NUM_COMPONENTS>;
vtkm::cont::ArrayHandle<ValueType> baseArray;
baseArray.Allocate(ARRAY_SIZE);
SetPortal(baseArray.WritePortal());
vtkm::cont::ArrayHandle<ComponentType> resultArray;
vtkm::cont::ArrayHandleRuntimeVec<ComponentType> runtimeVecArray(NUM_COMPONENTS, resultArray);
vtkm::cont::Invoker{}(PassThrough{}, baseArray, runtimeVecArray);
VTKM_TEST_ASSERT(runtimeVecArray.GetNumberOfValues() == ARRAY_SIZE,
"Group array reporting wrong array size.");
VTKM_TEST_ASSERT(resultArray.GetNumberOfValues() == ARRAY_SIZE * NUM_COMPONENTS,
"Got bad result array size.");
//verify that the control portal works
vtkm::Id totalIndex = 0;
auto resultPortal = resultArray.ReadPortal();
for (vtkm::Id index = 0; index < ARRAY_SIZE; ++index)
{
const ValueType expectedValue = TestValue(index, ValueType());
for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++)
{
const ComponentType result = resultPortal.Get(totalIndex);
VTKM_TEST_ASSERT(test_equal(result, expectedValue[componentIndex]),
"Result array got wrong value.");
totalIndex++;
}
}
}
};
void Run()
{
using HandleTypesToTest =
vtkm::List<vtkm::Id, vtkm::Vec2i_32, vtkm::FloatDefault, vtkm::Vec3f_64>;
using ScalarTypesToTest = vtkm::List<vtkm::UInt8, vtkm::FloatDefault>;
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleRuntimeVec(3) as Input" << std::endl;
vtkm::testing::Testing::TryTypes(TestRuntimeVecAsInput<3>(), HandleTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleRuntimeVec(4) as Input" << std::endl;
vtkm::testing::Testing::TryTypes(TestRuntimeVecAsInput<4>(), HandleTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleRuntimeVec(2) as Output" << std::endl;
vtkm::testing::Testing::TryTypes(TestRuntimeVecAsOutput<2>(), ScalarTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleRuntimeVec(3) as Output" << std::endl;
vtkm::testing::Testing::TryTypes(TestRuntimeVecAsOutput<3>(), ScalarTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleRuntimeVec(3) as Input with unusual type" << std::endl;
TestRuntimeVecAsInput<3>{}(UnusualType{});
}
} // anonymous namespace
int UnitTestArrayHandleRuntimeVec(int argc, char* argv[])
{
return vtkm::cont::testing::Testing::Run(Run, argc, argv);
}

96
vtkm/cont/testing/UnitTestArrayRangeCompute.cxx
View File

@ -11,17 +11,22 @@
#include <vtkm/cont/ArrayCopyDevice.h>
#include <vtkm/cont/ArrayHandleBasic.h>
#include <vtkm/cont/ArrayHandleCartesianProduct.h>
#include <vtkm/cont/ArrayHandleCast.h>
#include <vtkm/cont/ArrayHandleCompositeVector.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/ArrayHandleExtractComponent.h>
#include <vtkm/cont/ArrayHandleGroupVec.h>
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/ArrayHandleRandomUniformReal.h>
#include <vtkm/cont/ArrayHandleSOA.h>
#include <vtkm/cont/ArrayHandleStride.h>
#include <vtkm/cont/ArrayHandleUniformPointCoordinates.h>
#include <vtkm/cont/ArrayHandleView.h>
#include <vtkm/cont/ArrayHandleXGCCoordinates.h>
#include <vtkm/cont/ArrayRangeCompute.h>
#include <vtkm/Math.h>
#include <vtkm/VecTraits.h>
#include <vtkm/cont/testing/Testing.h>
@ -31,12 +36,12 @@ namespace
constexpr vtkm::Id ARRAY_SIZE = 20;
template <typename T, typename S>
void CheckRange(const vtkm::cont::ArrayHandle<T, S>& array, bool checkUnknown = true)
void VerifyRange(const vtkm::cont::ArrayHandle<T, S>& array,
const vtkm::cont::ArrayHandle<vtkm::Range>& computedRangeArray)
{
using Traits = vtkm::VecTraits<T>;
vtkm::IdComponent numComponents = Traits::NUM_COMPONENTS;
vtkm::cont::ArrayHandle<vtkm::Range> computedRangeArray = vtkm::cont::ArrayRangeCompute(array);
VTKM_TEST_ASSERT(computedRangeArray.GetNumberOfValues() == numComponents);
auto computedRangePortal = computedRangeArray.ReadPortal();
@ -54,28 +59,12 @@ void CheckRange(const vtkm::cont::ArrayHandle<T, S>& array, bool checkUnknown =
VTKM_TEST_ASSERT(!vtkm::IsNan(computedRange.Max));
VTKM_TEST_ASSERT(test_equal(expectedRange, computedRange));
}
}
if (checkUnknown)
{
computedRangeArray = vtkm::cont::ArrayRangeCompute(vtkm::cont::UnknownArrayHandle{ array });
VTKM_TEST_ASSERT(computedRangeArray.GetNumberOfValues() == numComponents);
computedRangePortal = computedRangeArray.ReadPortal();
portal = array.ReadPortal();
for (vtkm::IdComponent component = 0; component < numComponents; ++component)
{
vtkm::Range computedRange = computedRangePortal.Get(component);
vtkm::Range expectedRange;
for (vtkm::Id index = 0; index < portal.GetNumberOfValues(); ++index)
{
T value = portal.Get(index);
expectedRange.Include(Traits::GetComponent(value, component));
}
VTKM_TEST_ASSERT(!vtkm::IsNan(computedRange.Min));
VTKM_TEST_ASSERT(!vtkm::IsNan(computedRange.Max));
VTKM_TEST_ASSERT(test_equal(expectedRange, computedRange));
}
}
template <typename T, typename S>
void CheckRange(const vtkm::cont::ArrayHandle<T, S>& array)
{
VerifyRange(array, vtkm::cont::ArrayRangeCompute(array));
}
template <typename T, typename S>
@ -119,7 +108,7 @@ void TestSOAArray(vtkm::TypeTraitsScalarTag)
}
template <typename T>
void TestStrideArray(vtkm::TypeTraitsScalarTag)
void TestStrideArray()
{
std::cout << "Checking stride array" << std::endl;
vtkm::cont::ArrayHandleBasic<T> array;
@ -127,6 +116,16 @@ void TestStrideArray(vtkm::TypeTraitsScalarTag)
CheckRange(vtkm::cont::ArrayHandleStride<T>(array, ARRAY_SIZE / 2, 2, 1));
}
template <typename T>
void TestCastArray()
{
std::cout << "Checking cast array" << std::endl;
using CastType = typename vtkm::VecTraits<T>::template ReplaceBaseComponentType<vtkm::Float64>;
vtkm::cont::ArrayHandle<T> array;
FillArray(array);
CheckRange(vtkm::cont::make_ArrayHandleCast<CastType>(array));
}
template <typename T>
void TestCartesianProduct(vtkm::TypeTraitsScalarTag)
{
@ -151,7 +150,7 @@ void TestCartesianProduct(vtkm::TypeTraitsVectorTag)
template <typename T>
void TestComposite(vtkm::TypeTraitsScalarTag)
{
std::cout << "Checking composite vector" << std::endl;
std::cout << "Checking composite vector array" << std::endl;
vtkm::cont::ArrayHandleBasic<T> array0;
FillArray(array0);
@ -169,6 +168,32 @@ void TestComposite(vtkm::TypeTraitsVectorTag)
// Skip test.
}
template <typename T>
void TestGroup(vtkm::TypeTraitsScalarTag)
{
std::cout << "Checking group vec array" << std::endl;
vtkm::cont::ArrayHandleBasic<T> array;
FillArray(array);
CheckRange(vtkm::cont::make_ArrayHandleGroupVec<2>(array));
}
template <typename T>
void TestGroup(vtkm::TypeTraitsVectorTag)
{
// Skip test.
}
template <typename T>
void TestView()
{
std::cout << "Checking view array" << std::endl;
vtkm::cont::ArrayHandleBasic<T> array;
FillArray(array);
CheckRange(vtkm::cont::make_ArrayHandleView(array, 2, ARRAY_SIZE - 5));
}
template <typename T>
void TestConstant()
{
@ -205,15 +230,29 @@ void TestUniformPointCoords()
vtkm::cont::ArrayHandleUniformPointCoordinates(vtkm::Id3(ARRAY_SIZE, ARRAY_SIZE, ARRAY_SIZE)));
}
void TestXGCCoordinates()
{
std::cout << "Checking XGC coordinates array" << std::endl;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> array;
FillArray(array);
CheckRange(vtkm::cont::make_ArrayHandleXGCCoordinates(array, 4, true));
}
struct DoTestFunctor
{
template <typename T>
void operator()(T) const
{
typename vtkm::TypeTraits<T>::DimensionalityTag dimensionality{};
TestBasicArray<T>();
TestSOAArray<T>(typename vtkm::TypeTraits<T>::DimensionalityTag{});
TestCartesianProduct<T>(typename vtkm::TypeTraits<T>::DimensionalityTag{});
TestComposite<T>(typename vtkm::TypeTraits<T>::DimensionalityTag{});
TestSOAArray<T>(dimensionality);
TestStrideArray<T>();
TestCastArray<T>();
TestCartesianProduct<T>(dimensionality);
TestComposite<T>(dimensionality);
TestGroup<T>(dimensionality);
TestView<T>();
TestConstant<T>();
TestCounting<T>(typename std::is_signed<typename vtkm::VecTraits<T>::ComponentType>::type{});
}
@ -226,6 +265,7 @@ void DoTest()
std::cout << "*** Specific arrays *****************" << std::endl;
TestIndex();
TestUniformPointCoords();
TestXGCCoordinates();
}
} // anonymous namespace

20
vtkm/cont/testing/UnitTestBitField.cxx
View File

@ -24,6 +24,25 @@
#include <cstdio>
#if defined(KOKKOS_ENABLE_SYCL)
#define DEVICE_ASSERT_MSG(cond, message) \
do \
{ \
if (!(cond)) \
{ \
return false; \
} \
} while (false)
#define DEVICE_ASSERT(cond) \
do \
{ \
if (!(cond)) \
{ \
return false; \
} \
} while (false)
#else
#define DEVICE_ASSERT_MSG(cond, message) \
do \
{ \
@ -47,6 +66,7 @@
return false; \
} \
} while (false)
#endif
// Test with some trailing bits in partial last word:
#define NUM_BITS \

15
vtkm/cont/testing/UnitTestDataSetPermutation.cxx
View File

@ -105,7 +105,10 @@ void TestDataSet_Explicit()
//run a basic for-each topology algorithm on this
vtkm::cont::ArrayHandle<vtkm::Float32> result;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellAverage> dispatcher;
dispatcher.Invoke(subset, dataSet.GetField("pointvar"), result);
dispatcher.Invoke(
subset,
dataSet.GetField("pointvar").GetData().AsArrayHandle<vtkm::cont::ArrayHandle<vtkm::Float32>>(),
result);
//iterate same cell 4 times
vtkm::Float32 expected[4] = { 30.1667f, 30.1667f, 30.1667f, 30.1667f };
@ -139,7 +142,10 @@ void TestDataSet_Structured2D()
//run a basic for-each topology algorithm on this
vtkm::cont::ArrayHandle<vtkm::Float32> result;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellAverage> dispatcher;
dispatcher.Invoke(subset, dataSet.GetField("pointvar"), result);
dispatcher.Invoke(
subset,
dataSet.GetField("pointvar").GetData().AsArrayHandle<vtkm::cont::ArrayHandle<vtkm::Float32>>(),
result);
vtkm::Float32 expected[4] = { 40.1f, 40.1f, 40.1f, 40.1f };
auto resultPortal = result.ReadPortal();
@ -172,7 +178,10 @@ void TestDataSet_Structured3D()
//run a basic for-each topology algorithm on this
vtkm::cont::ArrayHandle<vtkm::Float32> result;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellAverage> dispatcher;
dispatcher.Invoke(subset, dataSet.GetField("pointvar"), result);
dispatcher.Invoke(
subset,
dataSet.GetField("pointvar").GetData().AsArrayHandle<vtkm::cont::ArrayHandle<vtkm::Float32>>(),
result);
vtkm::Float32 expected[4] = { 70.2125f, 70.2125f, 70.2125f, 70.2125f };
auto resultPortal = result.ReadPortal();

12
vtkm/cont/testing/UnitTestInitialize.cxx
View File

@ -146,16 +146,8 @@ void InitializeRuntimeDeviceConfigurationWithArgs()
{
int argc;
char** argv;
vtkm::cont::testing::Testing::MakeArgsAddProgramName(argc,
argv,
"--vtkm-device",
"Any",
"--vtkm-num-threads",
"100",
"--vtkm-numa-regions",
"4",
"--vtkm-device-instance",
"2");
vtkm::cont::testing::Testing::MakeArgsAddProgramName(
argc, argv, "--vtkm-device", "Any", "--vtkm-num-threads", "100", "--vtkm-device-instance", "2");
vtkm::cont::Initialize(argc, argv);
CheckArgs(argc, argv);
}

22
vtkm/cont/testing/UnitTestRuntimeConfigurationOptions.cxx
View File

@ -187,11 +187,9 @@ void TestConfigOptionValues(const internal::RuntimeDeviceConfigurationOptions& c
VTKM_TEST_ASSERT(configOptions.IsInitialized(), "runtime config options should be initialized");
VTKM_TEST_ASSERT(configOptions.VTKmNumThreads.IsSet(), "num threads should be set");
VTKM_TEST_ASSERT(configOptions.VTKmNumaRegions.IsSet(), "numa regions should be set");
VTKM_TEST_ASSERT(configOptions.VTKmDeviceInstance.IsSet(), "device instance should be set");
VTKM_TEST_ASSERT(configOptions.VTKmNumThreads.GetValue() == 100, "num threads should == 100");
VTKM_TEST_ASSERT(configOptions.VTKmNumaRegions.GetValue() == 2, "numa regions should == 2");
VTKM_TEST_ASSERT(configOptions.VTKmDeviceInstance.GetValue() == 1, "device instance should == 1");
}
@ -211,14 +209,8 @@ void TestRuntimeDeviceConfigurationOptions()
int argc;
char** argv;
vtkm::cont::testing::Testing::MakeArgs(argc,
argv,
"--vtkm-num-threads",
"100",
"--vtkm-numa-regions",
"2",
"--vtkm-device-instance",
"1");
vtkm::cont::testing::Testing::MakeArgs(
argc, argv, "--vtkm-num-threads", "100", "--vtkm-device-instance", "1");
auto options = GetOptions(argc, argv, usage);
VTKM_TEST_ASSERT(!configOptions.IsInitialized(),
@ -230,14 +222,8 @@ void TestRuntimeDeviceConfigurationOptions()
{
int argc;
char** argv;
vtkm::cont::testing::Testing::MakeArgs(argc,
argv,
"--vtkm-num-threads",
"100",
"--vtkm-numa-regions",
"2",
"--vtkm-device-instance",
"1");
vtkm::cont::testing::Testing::MakeArgs(
argc, argv, "--vtkm-num-threads", "100", "--vtkm-device-instance", "1");
internal::RuntimeDeviceConfigurationOptions configOptions(argc, argv);
TestConfigOptionValues(configOptions);
}

17
vtkm/cont/testing/UnitTestSerializationArrayHandle.cxx
View File

@ -22,6 +22,7 @@
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/ArrayHandlePermutation.h>
#include <vtkm/cont/ArrayHandleReverse.h>
#include <vtkm/cont/ArrayHandleRuntimeVec.h>
#include <vtkm/cont/ArrayHandleSOA.h>
#include <vtkm/cont/ArrayHandleSwizzle.h>
#include <vtkm/cont/ArrayHandleTransform.h>
@ -303,6 +304,19 @@ struct TestArrayHandleGroupVecVariable
}
};
struct TestArrayHandleRuntimeVec
{
template <typename T>
void operator()(T) const
{
auto numComps = RandomValue<vtkm::IdComponent>::Make(1, 5);
auto flat = RandomArrayHandle<T>::Make(ArraySize * numComps);
auto array = vtkm::cont::make_ArrayHandleRuntimeVec(numComps, flat);
RunTest(array);
RunTest(MakeTestUnknownArrayHandle(array));
}
};
void TestArrayHandleIndex()
{
auto size = RandomValue<vtkm::Id>::Make(2, 10);
@ -413,6 +427,9 @@ void TestArrayHandleSerialization()
std::cout << "Testing ArrayHandleGroupVecVariable\n";
vtkm::testing::Testing::TryTypes(TestArrayHandleGroupVecVariable(), TestTypesList());
std::cout << "Testing ArrayHandleRuntimeVec\n";
vtkm::testing::Testing::TryTypes(TestArrayHandleRuntimeVec(), TestTypesList());
std::cout << "Testing ArrayHandleIndex\n";
TestArrayHandleIndex();

16
vtkm/cont/testing/UnitTestSerializationDataSet.cxx
View File

@ -22,16 +22,26 @@ using CellSetTypes = vtkm::List<vtkm::cont::CellSetExplicit<>,
vtkm::cont::CellSetStructured<2>,
vtkm::cont::CellSetStructured<3>>;
using DataSetWrapper = vtkm::cont::SerializableDataSet<FieldTypeList, CellSetTypes>;
using DataSetWrapper = vtkm::cont::DataSetWithCellSetTypes<CellSetTypes>;
VTKM_CONT void TestEqualDataSet(const DataSetWrapper& ds1, const DataSetWrapper& ds2)
VTKM_CONT void TestEqualDataSetWrapper(const DataSetWrapper& ds1, const DataSetWrapper& ds2)
{
VTKM_TEST_ASSERT(test_equal_DataSets(ds1.DataSet, ds2.DataSet, CellSetTypes{}));
}
VTKM_CONT void TestEqualDataSet(const vtkm::cont::DataSet& ds1, const vtkm::cont::DataSet& ds2)
{
VTKM_TEST_ASSERT(test_equal_DataSets(ds1, ds2, CellSetTypes{}));
}
void RunTest(const vtkm::cont::DataSet& ds)
{
TestSerialization(DataSetWrapper(ds), TestEqualDataSet);
VTKM_DEPRECATED_SUPPRESS_BEGIN
TestSerialization(vtkm::cont::SerializableDataSet<FieldTypeList, CellSetTypes>(ds),
TestEqualDataSetWrapper);
VTKM_DEPRECATED_SUPPRESS_END
TestSerialization(DataSetWrapper(ds), TestEqualDataSetWrapper);
TestSerialization(ds, TestEqualDataSet);
}
void TestDataSetSerialization()

94
vtkm/cont/testing/UnitTestUnknownArrayHandle.cxx
View File

@ -16,6 +16,7 @@
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/ArrayHandleGroupVecVariable.h>
#include <vtkm/cont/ArrayHandleMultiplexer.h>
#include <vtkm/cont/ArrayHandleRuntimeVec.h>
#include <vtkm/TypeTraits.h>
@ -350,6 +351,43 @@ void TryAsMultiplexer(vtkm::cont::UnknownArrayHandle sourceArray)
#endif
}
struct SimpleRecombineCopy
{
template <typename T>
void operator()(const vtkm::cont::ArrayHandleRecombineVec<T>& inputArray,
const vtkm::cont::UnknownArrayHandle& output) const
{
vtkm::cont::ArrayHandleRecombineVec<T> outputArray =
output.ExtractArrayFromComponents<T>(vtkm::CopyFlag::Off);
vtkm::Id size = inputArray.GetNumberOfValues();
outputArray.Allocate(size);
auto inputPortal = inputArray.ReadPortal();
auto outputPortal = outputArray.WritePortal();
for (vtkm::Id index = 0; index < size; ++index)
{
outputPortal.Set(index, inputPortal.Get(index));
}
}
};
template <typename T>
void TryExtractArray(const vtkm::cont::UnknownArrayHandle& originalArray)
{
// This check should already have been performed by caller, but just in case.
CheckUnknownArray<vtkm::List<T>, VTKM_DEFAULT_STORAGE_LIST>(originalArray,
vtkm::VecTraits<T>::NUM_COMPONENTS);
std::cout << "Create new instance of array." << std::endl;
vtkm::cont::UnknownArrayHandle newArray = originalArray.NewInstanceBasic();
std::cout << "Do CastAndCallWithExtractedArray." << std::endl;
originalArray.CastAndCallWithExtractedArray(SimpleRecombineCopy{}, newArray);
CheckUnknownArray<vtkm::List<T>, VTKM_DEFAULT_STORAGE_LIST>(newArray,
vtkm::VecTraits<T>::NUM_COMPONENTS);
}
template <typename T>
void TryDefaultType()
{
@ -360,6 +398,8 @@ void TryDefaultType()
TryNewInstance<T>(array);
TryAsMultiplexer<T>(array);
TryExtractArray<T>(array);
}
struct TryBasicVTKmType
@ -503,6 +543,57 @@ void TrySetMultiplexerArray()
CheckUnknownArray<vtkm::List<vtkm::Id>, vtkm::List<VTKM_DEFAULT_STORAGE_TAG>>(unknownArray, 1);
}
template <typename T>
void TryConvertRuntimeVec()
{
using BasicArrayType = vtkm::cont::ArrayHandle<T>;
using BasicComponentType = typename vtkm::VecFlat<T>::ComponentType;
constexpr vtkm::IdComponent numFlatComponents = vtkm::VecFlat<T>::NUM_COMPONENTS;
using RuntimeArrayType = vtkm::cont::ArrayHandleRuntimeVec<BasicComponentType>;
std::cout << " Get basic array as ArrayHandleRuntimeVec" << std::endl;
BasicArrayType inputArray;
inputArray.Allocate(ARRAY_SIZE);
SetPortal(inputArray.WritePortal());
vtkm::cont::UnknownArrayHandle unknownWithBasic{ inputArray };
VTKM_TEST_ASSERT(unknownWithBasic.GetNumberOfComponentsFlat() == numFlatComponents);
VTKM_TEST_ASSERT(unknownWithBasic.CanConvert<RuntimeArrayType>());
RuntimeArrayType runtimeArray = unknownWithBasic.AsArrayHandle<RuntimeArrayType>();
// Hack to convert the array handle to a flat array to make it easy to check the runtime array
vtkm::cont::ArrayHandle<vtkm::VecFlat<T>> flatInput{ inputArray.GetBuffers() };
VTKM_TEST_ASSERT(test_equal_ArrayHandles(flatInput, runtimeArray));
std::cout << " Get ArrayHandleRuntimeVec as basic array" << std::endl;
vtkm::cont::UnknownArrayHandle unknownWithRuntimeVec{ runtimeArray };
VTKM_TEST_ASSERT(unknownWithRuntimeVec.GetNumberOfComponentsFlat() == numFlatComponents);
VTKM_TEST_ASSERT(unknownWithRuntimeVec.CanConvert<RuntimeArrayType>());
VTKM_TEST_ASSERT(unknownWithRuntimeVec.CanConvert<BasicArrayType>());
BasicArrayType outputArray = unknownWithRuntimeVec.AsArrayHandle<BasicArrayType>();
VTKM_TEST_ASSERT(test_equal_ArrayHandles(inputArray, outputArray));
}
void TryConvertRuntimeVec()
{
std::cout << " Scalar array." << std::endl;
TryConvertRuntimeVec<vtkm::FloatDefault>();
std::cout << " Equivalent scalar." << std::endl;
TryConvertRuntimeVec<VTKM_UNUSED_INT_TYPE>();
std::cout << " Basic Vec." << std::endl;
TryConvertRuntimeVec<vtkm::Id3>();
std::cout << " Vec of Vecs." << std::endl;
TryConvertRuntimeVec<vtkm::Vec<vtkm::Vec2f, 3>>();
std::cout << " Vec of Vecs of Vecs." << std::endl;
TryConvertRuntimeVec<vtkm::Vec<vtkm::Vec<vtkm::Id4, 3>, 2>>();
}
struct DefaultTypeFunctor
{
template <typename T>
@ -534,6 +625,9 @@ void TestUnknownArrayHandle()
std::cout << "Try setting ArrayHandleMultiplexer" << std::endl;
TrySetMultiplexerArray();
std::cout << "Try converting between ArrayHandleRuntimeVec and basic array" << std::endl;
TryConvertRuntimeVec();
}
} // anonymous namespace

1
vtkm/cont/vtkm.module
View File

@ -9,6 +9,7 @@ DEPENDS
OPTIONAL_DEPENDS
vtkm_loguru
TEST_OPTIONAL_DEPENDS
vtkm_filter_contour
vtkm_filter_field_conversion
TEST_DEPENDS
vtkm_source

2
vtkm/exec/arg/CMakeLists.txt
View File

@ -19,8 +19,6 @@ set(headers
FetchTagArrayDirectIn.h
FetchTagArrayDirectInOut.h
FetchTagArrayDirectOut.h
FetchTagArrayDirectOutArrayHandleGroupVecVariable.h
FetchTagArrayDirectOutArrayHandleRecombineVec.h
FetchTagArrayNeighborhoodIn.h
FetchTagArrayTopologyMapIn.h
FetchTagExecObject.h

38
vtkm/exec/arg/FetchTagArrayDirectOut.h
View File

@ -13,6 +13,8 @@
#include <vtkm/exec/arg/AspectTagDefault.h>
#include <vtkm/exec/arg/Fetch.h>
#include <type_traits>
namespace vtkm
{
namespace exec
@ -35,14 +37,15 @@ struct Fetch<vtkm::exec::arg::FetchTagArrayDirectOut,
vtkm::exec::arg::AspectTagDefault,
ExecObjectType>
{
using ValueType = typename ExecObjectType::ValueType;
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename ThreadIndicesType>
VTKM_EXEC auto Load(const ThreadIndicesType&, const ExecObjectType&) const ->
typename ExecObjectType::ValueType
VTKM_EXEC ValueType Load(const ThreadIndicesType& indices,
const ExecObjectType& arrayPortal) const
{
// Load is a no-op for this fetch.
using ValueType = typename ExecObjectType::ValueType;
return ValueType();
return this->DoLoad(
indices, arrayPortal, typename std::is_default_constructible<ValueType>::type{});
}
VTKM_SUPPRESS_EXEC_WARNINGS
@ -51,9 +54,32 @@ struct Fetch<vtkm::exec::arg::FetchTagArrayDirectOut,
const ExecObjectType& arrayPortal,
const T& value) const
{
using ValueType = typename ExecObjectType::ValueType;
arrayPortal.Set(indices.GetOutputIndex(), static_cast<ValueType>(value));
}
private:
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename ThreadIndicesType>
VTKM_EXEC ValueType DoLoad(const ThreadIndicesType&, const ExecObjectType&, std::true_type) const
{
// Load is a no-op for this fetch.
return ValueType();
}
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename ThreadIndicesType>
VTKM_EXEC ValueType DoLoad(const ThreadIndicesType& indices,
const ExecObjectType& arrayPortal,
std::false_type) const
{
// Cannot create a ValueType object, so pull one out of the array portal. This may seem
// weird because an output array often has garbage in it. However, this case can happen
// with special arrays with Vec-like values that reference back to the array memory.
// For example, with ArrayHandleRecombineVec, the values are actual objects that point
// back to the array for on demand reading and writing. You need the buffer established
// by the array even if there is garbage in that array.
return arrayPortal.Get(indices.GetOutputIndex());
}
};
}
}

69
vtkm/exec/arg/FetchTagArrayDirectOutArrayHandleGroupVecVariable.h
View File

@ -1,69 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_exec_arg_FetchTagArrayDirectOutArrayHandleGroupVecVariable_h
#define vtk_m_exec_arg_FetchTagArrayDirectOutArrayHandleGroupVecVariable_h
#include <vtkm/exec/arg/FetchTagArrayDirectOut.h>
// We need to override the fetch for output fields using
// ArrayPortalGroupVecVariable because this portal does not behave like most
// ArrayPortals. Usually you ignore the Load and implement the Store. But if
// you ignore the Load, the VecFromPortal gets no portal to set values into.
// Instead, you need to implement the Load to point to the array portal. You
// can also ignore the Store because the data is already set in the array at
// that point.
// This file is included from ArrayHandleGroupVecVariable.h
namespace vtkm
{
namespace exec
{
namespace arg
{
// We need to override the fetch for output fields using
// ArrayPortalGroupVecVariable because this portal does not behave like most
// ArrayPortals. Usually you ignore the Load and implement the Store. But if
// you ignore the Load, the VecFromPortal gets no portal to set values into.
// Instead, you need to implement the Load to point to the array portal. You
// can also ignore the Store because the data is already set in the array at
// that point.
template <typename ComponentsPortalType, typename OffsetsPortalType>
struct Fetch<vtkm::exec::arg::FetchTagArrayDirectOut,
vtkm::exec::arg::AspectTagDefault,
vtkm::internal::ArrayPortalGroupVecVariable<ComponentsPortalType, OffsetsPortalType>>
{
using ExecObjectType =
vtkm::internal::ArrayPortalGroupVecVariable<ComponentsPortalType, OffsetsPortalType>;
using ValueType = typename ExecObjectType::ValueType;
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename ThreadIndicesType>
VTKM_EXEC ValueType Load(const ThreadIndicesType& indices,
const ExecObjectType& arrayPortal) const
{
return arrayPortal.Get(indices.GetOutputIndex());
}
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename ThreadIndicesType>
VTKM_EXEC void Store(const ThreadIndicesType&, const ExecObjectType&, const ValueType&) const
{
// We can actually ignore this because the VecFromPortal will already have
// set new values in the array.
}
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_FetchTagArrayDirectOutArrayHandleGroupVecVariable_h

46
vtkm/exec/arg/FetchTagArrayDirectOutArrayHandleRecombineVec.h
View File

@ -1,46 +0,0 @@
//============================================================================
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
//============================================================================
#ifndef vtk_m_exec_arg_FetchTagArrayDirectOutArrayHandleRecombineVec_h
#define vtk_m_exec_arg_FetchTagArrayDirectOutArrayHandleRecombineVec_h
#include <vtkm/exec/arg/FetchTagArrayDirectInOut.h>
#include <vtkm/exec/arg/FetchTagArrayDirectOut.h>
// The `Fetch` for direct array out breaks for `ArrayHandleRecombineVec` because the `Load`
// method attempts to create a `vtkm::internal::RecombineVec` with a default constructor,
// which does not exist. Instead, have the direct out `Fetch` behave like the direct in/out
// `Fetch`, which loads the initial value from the array. The actual load will not load the
// data but rather set up the portals in the returned object, which is necessary for the
// later `Store` to work anyway.
// This file is included from ArrayHandleRecombineVec.h
namespace vtkm
{
namespace exec
{
namespace arg
{
template <typename SourcePortalType>
struct Fetch<vtkm::exec::arg::FetchTagArrayDirectOut,
vtkm::exec::arg::AspectTagDefault,
vtkm::internal::ArrayPortalRecombineVec<SourcePortalType>>
: Fetch<vtkm::exec::arg::FetchTagArrayDirectInOut,
vtkm::exec::arg::AspectTagDefault,
vtkm::internal::ArrayPortalRecombineVec<SourcePortalType>>
{
};
}
}
} // namespace vtkm::exec::arg
#endif //vtk_m_exec_arg_FetchTagArrayDirectOutArrayHandleRecombineVec_h

14
vtkm/exec/arg/FetchTagArrayTopologyMapIn.h
View File

@ -227,15 +227,13 @@ struct Fetch<vtkm::exec::arg::FetchTagArrayTopologyMapIn,
const vtkm::Id offset2 = (xgcidx.Planes[1] * xgcidx.NumberOfPointsPerPlane);
using ValueType = vtkm::Vec<typename ExecObjectType::ValueType, 6>;
ValueType result;
result[0] = portal.Get(offset1 + xgcidx.PointIds[0][0]);
result[1] = portal.Get(offset1 + xgcidx.PointIds[0][1]);
result[2] = portal.Get(offset1 + xgcidx.PointIds[0][2]);
result[3] = portal.Get(offset2 + xgcidx.PointIds[1][0]);
result[4] = portal.Get(offset2 + xgcidx.PointIds[1][1]);
result[5] = portal.Get(offset2 + xgcidx.PointIds[1][2]);
return result;
return ValueType(portal.Get(offset1 + xgcidx.PointIds[0][0]),
portal.Get(offset1 + xgcidx.PointIds[0][1]),
portal.Get(offset1 + xgcidx.PointIds[0][2]),
portal.Get(offset2 + xgcidx.PointIds[1][0]),
portal.Get(offset2 + xgcidx.PointIds[1][1]),
portal.Get(offset2 + xgcidx.PointIds[1][2]));
}

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