Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-08 13:23:51 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Fixing unit test

Browse Source
This commit is contained in:
ayenpure 2018-08-30 10:19:00 -07:00
commit 22ca8bce15
141 changed files with 3858 additions and 2061 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
Utilities/Scripts/devAlgoBenchSummary.py → Utilities/Scripts/benchCompare.py
View File

@ -3,14 +3,11 @@
# Compares the output from BenchmarkDeviceAdapter from the serial
# device to a parallel device and prints a table containing the results.
#
# While this was written for the device adapter algorithms, it could be used
# to compare any VTKM benchmark output.
#
# Example usage:
#
# $ BenchmarkDeviceAdapter_SERIAL > serial.out
# $ BenchmarkDeviceAdapter_TBB > tbb.out
# $ devAlgoBenchSummary.py serial.out tbb.out
# $ benchCompare.py serial.out tbb.out
#
#
# The number of threads (optional -- only used to generate the "Warn" column)

111
Utilities/Scripts/benchSummary.py Executable file
View File

@ -0,0 +1,111 @@
#!/usr/bin/env python
#
# Prints a concise summary of a benchmark output as a TSV blob.
#
# Example usage:
#
# $ BenchmarkXXX_DEVICE > bench.out
# $ benchSummary.py bench.out
#
# Options SortByType, SortByName, or SortByMean may be passed after the
# filename to sort the output by the indicated quantity. If no sort option
# is provided, the output order matches the input. If multiple options are
# specified, the list will be sorted repeatedly in the order requested.
import re
import sys
assert(len(sys.argv) >= 2)
# Parses "*** vtkm::Float64 ***************" --> vtkm::Float64
typeParser = re.compile("\\*{3} ([^*]+) \\*{15}")
# Parses "Benchmark 'Benchmark name' results:" --> Benchmark name
nameParser = re.compile("Benchmark '([^-]+)' results:")
# Parses "mean = 0.0125s" --> 0.0125
meanParser = re.compile("\\s+mean = ([0-9.Ee+-]+)s")
# Parses "std dev = 0.0125s" --> 0.0125
stdDevParser = re.compile("\\s+std dev = ([naN0-9.Ee+-]+)s")
filename = sys.argv[1]
benchFile = open(filename, 'r')
sortOpt = None
if len(sys.argv) > 2:
sortOpt = sys.argv[2:]
class BenchKey:
def __init__(self, name_, type_):
self.name = name_
self.type = type_
def __eq__(self, other):
return self.name == other.name and self.type == other.type
def __lt__(self, other):
if self.name < other.name: return True
elif self.name > other.name: return False
else: return self.type < other.type
def __hash__(self):
return (self.name + self.type).__hash__()
class BenchData:
def __init__(self, mean_, stdDev_):
self.mean = mean_
self.stdDev = stdDev_
def parseFile(f, benchmarks):
type = ""
bench = ""
mean = -1.
stdDev = -1.
for line in f:
typeRes = typeParser.match(line)
if typeRes:
type = typeRes.group(1)
continue
nameRes = nameParser.match(line)
if nameRes:
name = nameRes.group(1)
continue
meanRes = meanParser.match(line)
if meanRes:
mean = float(meanRes.group(1))
continue
stdDevRes = stdDevParser.match(line)
if stdDevRes:
stdDev = float(stdDevRes.group(1))
# stdDev is always the last parse for a given benchmark, add entry now
benchmarks[BenchKey(name, type)] = BenchData(mean, stdDev)
mean = -1.
stdDev = -1.
continue
benchmarks = {}
parseFile(benchFile, benchmarks)
# Sort keys by type:
keys = benchmarks.keys()
if sortOpt:
for opt in sortOpt:
if opt.lower() == "sortbytype":
keys = sorted(keys, key=lambda k: k.type)
elif opt.lower() == "sortbyname":
keys = sorted(keys, key=lambda k: k.name)
elif opt.lower() == "sortbymean":
keys = sorted(keys, key=lambda k: benchmarks[k].mean)
print("# Summary: (%s)"%filename)
print("%-9s\t%-9s\t%-s"%("Mean", "Stdev", "Benchmark (type)"))
for key in keys:
data = benchmarks[key]
print("%9.6f\t%9.6f\t%s (%s)"%(data.mean, data.stdDev, key.name, key.type))

156
Utilities/Scripts/benchSummaryWithBaselines.py Executable file
View File

@ -0,0 +1,156 @@
#!/usr/bin/env python
#
# Prints a concise summary of a benchmark output as a TSV blob. Benchmarks are
# expected to have "Baseline" in the name, and a matching benchmark with the
# same name but Baseline replaced with something else. For example,
#
# Baseline benchmark name: "Some benchmark: Baseline, Size=4"
# Test benchmark name: "Some benchmark: Blahblah, Size=4"
#
# The output will print the baseline, test, and overhead times for the
# benchmarks.
#
# Example usage:
#
# $ BenchmarkXXX_DEVICE > bench.out
# $ benchSummaryWithBaselines.py bench.out
#
# Options SortByType, SortByName, SortByOverhead, or SortByRatio
# (testtime/baseline) may be passed after the filename to sort the output by
# the indicated quantity. If no sort option is provided, the output order
# matches the input. If multiple options are specified, the list will be sorted
# repeatedly in the order requested.
import re
import sys
assert(len(sys.argv) >= 2)
# Parses "*** vtkm::Float64 ***************" --> vtkm::Float64
typeParser = re.compile("\\*{3} ([^*]+) \\*{15}")
# Parses "Benchmark 'Benchmark name' results:" --> Benchmark name
nameParser = re.compile("Benchmark '([^-]+)' results:")
# Parses "mean = 0.0125s" --> 0.0125
meanParser = re.compile("\\s+mean = ([0-9.Ee+-]+)s")
# Parses "std dev = 0.0125s" --> 0.0125
stdDevParser = re.compile("\\s+std dev = ([naN0-9.Ee+-]+)s")
# Parses "SomeText Baseline Other Text" --> ("SomeText ", " Other Text")
baselineParser = re.compile("(.*)Baseline(.*)")
filename = sys.argv[1]
benchFile = open(filename, 'r')
sortOpt = None
if len(sys.argv) > 2:
sortOpt = sys.argv[2:]
class BenchKey:
def __init__(self, name_, type_):
self.name = name_
self.type = type_
def __eq__(self, other):
return self.name == other.name and self.type == other.type
def __lt__(self, other):
if self.name < other.name: return True
elif self.name > other.name: return False
else: return self.type < other.type
def __hash__(self):
return (self.name + self.type).__hash__()
class BenchData:
def __init__(self, mean_, stdDev_):
self.mean = mean_
self.stdDev = stdDev_
def parseFile(f, benchmarks):
type = ""
bench = ""
mean = -1.
stdDev = -1.
for line in f:
typeRes = typeParser.match(line)
if typeRes:
type = typeRes.group(1)
continue
nameRes = nameParser.match(line)
if nameRes:
name = nameRes.group(1)
continue
meanRes = meanParser.match(line)
if meanRes:
mean = float(meanRes.group(1))
continue
stdDevRes = stdDevParser.match(line)
if stdDevRes:
stdDev = float(stdDevRes.group(1))
# stdDev is always the last parse for a given benchmark, add entry now
benchmarks[BenchKey(name, type)] = BenchData(mean, stdDev)
mean = -1.
stdDev = -1.
continue
class BaselinedBenchData:
def __init__(self, baseline, test):
self.baseline = baseline.mean
self.test = test.mean
self.overhead = test.mean - baseline.mean
def findBaselines(benchmarks):
result = {}
for baseKey in benchmarks.keys():
# Look for baseline entries
baselineRes = baselineParser.match(baseKey.name)
if baselineRes:
prefix = baselineRes.group(1)
suffix = baselineRes.group(2)
# Find the test entry matching the baseline:
for testKey in benchmarks.keys():
if baseKey.type != testKey.type: # Need same type
continue
if baseKey.name == testKey.name: # Skip the base key
continue
if testKey.name.startswith(prefix) and testKey.name.endswith(suffix):
newName = (prefix + suffix).replace(", ,", ",")
newKey = BenchKey(newName, testKey.type)
newVal = BaselinedBenchData(benchmarks[baseKey], benchmarks[testKey])
result[newKey] = newVal
return result
benchmarks = {}
parseFile(benchFile, benchmarks)
benchmarks = findBaselines(benchmarks)
# Sort keys by type:
keys = benchmarks.keys()
if sortOpt:
for opt in sortOpt:
if opt.lower() == "sortbytype":
keys = sorted(keys, key=lambda k: k.type)
elif opt.lower() == "sortbyname":
keys = sorted(keys, key=lambda k: k.name)
elif opt.lower() == "sortbyoverhead":
keys = sorted(keys, key=lambda k: benchmarks[k].overhead)
elif opt.lower() == "sortbyratio":
keys = sorted(keys, key=lambda k: benchmarks[k].overhead / benchmarks[k].baseline)
print("# Summary: (%s)"%filename)
print("%-9s\t%-9s\t%-9s\t%-9s\t%-s"%("Baseline", "TestTime", "Overhead", "Test/Base", "Benchmark (type)"))
for key in keys:
data = benchmarks[key]
print("%9.6f\t%9.6f\t%9.6f\t%9.6f\t%s (%s)"%(data.baseline, data.test,
data.overhead, data.test / data.baseline, key.name, key.type))

581
benchmarking/BenchmarkAtomicArray.cxx Normal file
View File

@ -0,0 +1,581 @@
//============================================================================
// 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.
//
// Copyright 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2018 UT-Battelle, LLC.
// Copyright 2018 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include "Benchmarker.h"
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/AtomicArray.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
#include <vtkm/cont/Timer.h>
#include <vtkm/exec/FunctorBase.h>
#include <iomanip>
#include <sstream>
#include <string>
namespace vtkm
{
namespace benchmarking
{
// This is 32x larger than the largest array size.
static constexpr vtkm::Id NumWrites = 33554432; // 2^25
#define MAKE_ATOMIC_BENCHMARKS(Name, Class) \
VTKM_MAKE_BENCHMARK(Name##1, Class, 1); \
VTKM_MAKE_BENCHMARK(Name##8, Class, 8); \
VTKM_MAKE_BENCHMARK(Name##32, Class, 32); \
VTKM_MAKE_BENCHMARK(Name##512, Class, 512); \
VTKM_MAKE_BENCHMARK(Name##2048, Class, 2048); \
VTKM_MAKE_BENCHMARK(Name##32768, Class, 32768); \
VTKM_MAKE_BENCHMARK(Name##1048576, Class, 1048576)
#define RUN_ATOMIC_BENCHMARKS(Name) \
VTKM_RUN_BENCHMARK(Name##1, vtkm::cont::AtomicArrayTypeListTag{}); \
VTKM_RUN_BENCHMARK(Name##8, vtkm::cont::AtomicArrayTypeListTag{}); \
VTKM_RUN_BENCHMARK(Name##32, vtkm::cont::AtomicArrayTypeListTag{}); \
VTKM_RUN_BENCHMARK(Name##512, vtkm::cont::AtomicArrayTypeListTag{}); \
VTKM_RUN_BENCHMARK(Name##2048, vtkm::cont::AtomicArrayTypeListTag{}); \
VTKM_RUN_BENCHMARK(Name##32768, vtkm::cont::AtomicArrayTypeListTag{}); \
VTKM_RUN_BENCHMARK(Name##1048576, vtkm::cont::AtomicArrayTypeListTag{})
template <class Device>
class BenchmarkAtomicArray
{
public:
using Algo = vtkm::cont::DeviceAdapterAlgorithm<Device>;
using Timer = vtkm::cont::Timer<Device>;
// Benchmarks AtomicArray::Add such that each work index writes to adjacent
// indices.
template <typename ValueType>
struct BenchAddSeq
{
vtkm::Id ArraySize;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, PortalType portal)
: ArraySize(arraySize)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const { this->Portal.Add(i % this->ArraySize, 1); }
};
BenchAddSeq(vtkm::Id arraySize)
: ArraySize(arraySize)
{
this->Data.PrepareForOutput(this->ArraySize, Device{});
}
VTKM_CONT
vtkm::Float64 operator()()
{
vtkm::cont::AtomicArray<ValueType> array(this->Data);
auto portal = array.PrepareForExecution(Device{});
Worker<decltype(portal)> worker{ this->ArraySize, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "Add (Seq, Atomic, " << std::setw(7) << std::setfill('0') << this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(AddSeq, BenchAddSeq);
// Provides a non-atomic baseline for BenchAddSeq
template <typename ValueType>
struct BenchAddSeqBaseline
{
vtkm::Id ArraySize;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, PortalType portal)
: ArraySize(arraySize)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const
{
vtkm::Id idx = i % this->ArraySize;
this->Portal.Set(idx, this->Portal.Get(idx) + 1);
}
};
BenchAddSeqBaseline(vtkm::Id arraySize)
: ArraySize(arraySize)
{
}
VTKM_CONT
vtkm::Float64 operator()()
{
auto portal = this->Data.PrepareForOutput(this->ArraySize, Device{});
Worker<decltype(portal)> worker{ this->ArraySize, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "Add (Seq, Baseline, " << std::setw(7) << std::setfill('0') << this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(AddSeqBase, BenchAddSeqBaseline);
// Benchmarks AtomicArray::Add such that each work index writes to a strided
// index ( floor(i / stride) + stride * (i % stride)
template <typename ValueType>
struct BenchAddStride
{
vtkm::Id ArraySize;
vtkm::Id Stride;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
vtkm::Id Stride;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, vtkm::Id stride, PortalType portal)
: ArraySize(arraySize)
, Stride(stride)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const
{
vtkm::Id idx = (i / this->Stride + this->Stride * (i % this->Stride)) % this->ArraySize;
this->Portal.Add(idx % this->ArraySize, 1);
}
};
BenchAddStride(vtkm::Id arraySize, vtkm::Id stride = 32)
: ArraySize(arraySize)
, Stride(stride)
{
this->Data.PrepareForOutput(this->ArraySize, Device{});
}
VTKM_CONT
vtkm::Float64 operator()()
{
vtkm::cont::AtomicArray<ValueType> array(this->Data);
auto portal = array.PrepareForExecution(Device{});
Worker<decltype(portal)> worker{ this->ArraySize, this->Stride, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "Add (Stride=" << this->Stride << ", Atomic, " << std::setw(7) << std::setfill('0')
<< this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(AddStride, BenchAddStride);
// Non-atomic baseline for AddStride
template <typename ValueType>
struct BenchAddStrideBaseline
{
vtkm::Id ArraySize;
vtkm::Id Stride;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
vtkm::Id Stride;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, vtkm::Id stride, PortalType portal)
: ArraySize(arraySize)
, Stride(stride)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const
{
vtkm::Id idx = (i / this->Stride + this->Stride * (i % this->Stride)) % this->ArraySize;
this->Portal.Set(idx, this->Portal.Get(idx) + 1);
}
};
BenchAddStrideBaseline(vtkm::Id arraySize, vtkm::Id stride = 32)
: ArraySize(arraySize)
, Stride(stride)
{
}
VTKM_CONT
vtkm::Float64 operator()()
{
auto portal = this->Data.PrepareForOutput(this->ArraySize, Device{});
Worker<decltype(portal)> worker{ this->ArraySize, this->Stride, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "Add (Stride=" << this->Stride << ", Baseline, " << std::setw(7) << std::setfill('0')
<< this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(AddStrideBase, BenchAddStrideBaseline);
// Benchmarks AtomicArray::CompareAndSwap such that each work index writes to adjacent
// indices.
template <typename ValueType>
struct BenchCASSeq
{
vtkm::Id ArraySize;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, PortalType portal)
: ArraySize(arraySize)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const
{
vtkm::Id idx = i % this->ArraySize;
ValueType val = static_cast<ValueType>(i);
// Get the old val with a no-op
ValueType oldVal = this->Portal.Add(idx, static_cast<ValueType>(0));
ValueType assumed = static_cast<ValueType>(0);
do
{
assumed = oldVal;
oldVal = this->Portal.CompareAndSwap(idx, assumed + val, assumed);
} while (assumed != oldVal);
}
};
BenchCASSeq(vtkm::Id arraySize)
: ArraySize(arraySize)
{
this->Data.PrepareForOutput(this->ArraySize, Device{});
}
VTKM_CONT
vtkm::Float64 operator()()
{
vtkm::cont::AtomicArray<ValueType> array(this->Data);
auto portal = array.PrepareForExecution(Device{});
Worker<decltype(portal)> worker{ this->ArraySize, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "CAS (Seq, Atomic, " << std::setw(7) << std::setfill('0') << this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(CASSeq, BenchCASSeq);
// Provides a non-atomic baseline for BenchCASSeq
template <typename ValueType>
struct BenchCASSeqBaseline
{
vtkm::Id ArraySize;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, PortalType portal)
: ArraySize(arraySize)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const
{
vtkm::Id idx = i % this->ArraySize;
ValueType val = static_cast<ValueType>(i);
ValueType oldVal = this->Portal.Get(idx);
this->Portal.Set(idx, oldVal + val);
}
};
BenchCASSeqBaseline(vtkm::Id arraySize)
: ArraySize(arraySize)
{
}
VTKM_CONT
vtkm::Float64 operator()()
{
auto portal = this->Data.PrepareForOutput(this->ArraySize, Device{});
Worker<decltype(portal)> worker{ this->ArraySize, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "CAS (Seq, Baseline, " << std::setw(7) << std::setfill('0') << this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(CASSeqBase, BenchCASSeqBaseline);
// Benchmarks AtomicArray::CompareAndSwap such that each work index writes to
// a strided index:
// ( floor(i / stride) + stride * (i % stride)
template <typename ValueType>
struct BenchCASStride
{
vtkm::Id ArraySize;
vtkm::Id Stride;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
vtkm::Id Stride;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, vtkm::Id stride, PortalType portal)
: ArraySize(arraySize)
, Stride(stride)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const
{
vtkm::Id idx = (i / this->Stride + this->Stride * (i % this->Stride)) % this->ArraySize;
ValueType val = static_cast<ValueType>(i);
// Get the old val with a no-op
ValueType oldVal = this->Portal.Add(idx, static_cast<ValueType>(0));
ValueType assumed = static_cast<ValueType>(0);
do
{
assumed = oldVal;
oldVal = this->Portal.CompareAndSwap(idx, assumed + val, assumed);
} while (assumed != oldVal);
}
};
BenchCASStride(vtkm::Id arraySize, vtkm::Id stride = 32)
: ArraySize(arraySize)
, Stride(stride)
{
this->Data.PrepareForOutput(this->ArraySize, Device{});
}
VTKM_CONT
vtkm::Float64 operator()()
{
vtkm::cont::AtomicArray<ValueType> array(this->Data);
auto portal = array.PrepareForExecution(Device{});
Worker<decltype(portal)> worker{ this->ArraySize, this->Stride, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "CAS (Stride=" << this->Stride << ", Atomic, " << std::setw(7) << std::setfill('0')
<< this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(CASStride, BenchCASStride);
// Non-atomic baseline for CASStride
template <typename ValueType>
struct BenchCASStrideBaseline
{
vtkm::Id ArraySize;
vtkm::Id Stride;
vtkm::cont::ArrayHandle<ValueType> Data;
template <typename PortalType>
struct Worker : public vtkm::exec::FunctorBase
{
vtkm::Id ArraySize;
vtkm::Id Stride;
PortalType Portal;
VTKM_CONT
Worker(vtkm::Id arraySize, vtkm::Id stride, PortalType portal)
: ArraySize(arraySize)
, Stride(stride)
, Portal(portal)
{
}
VTKM_EXEC
void operator()(vtkm::Id i) const
{
vtkm::Id idx = (i / this->Stride + this->Stride * (i % this->Stride)) % this->ArraySize;
ValueType val = static_cast<ValueType>(i);
ValueType oldVal = this->Portal.Get(idx);
this->Portal.Set(idx, oldVal + val);
}
};
BenchCASStrideBaseline(vtkm::Id arraySize, vtkm::Id stride = 32)
: ArraySize(arraySize)
, Stride(stride)
{
}
VTKM_CONT
vtkm::Float64 operator()()
{
auto portal = this->Data.PrepareForOutput(this->ArraySize, Device{});
Worker<decltype(portal)> worker{ this->ArraySize, this->Stride, portal };
Timer timer;
Algo::Schedule(worker, NumWrites);
return timer.GetElapsedTime();
}
VTKM_CONT
std::string Description() const
{
std::ostringstream desc;
desc << "CAS (Stride=" << this->Stride << ", Baseline, " << std::setw(7) << std::setfill('0')
<< this->ArraySize << ")";
return desc.str();
}
};
MAKE_ATOMIC_BENCHMARKS(CASStrideBase, BenchCASStrideBaseline);
static void Run()
{
RUN_ATOMIC_BENCHMARKS(AddSeq);
RUN_ATOMIC_BENCHMARKS(AddSeqBase);
RUN_ATOMIC_BENCHMARKS(AddStride);
RUN_ATOMIC_BENCHMARKS(AddStrideBase);
RUN_ATOMIC_BENCHMARKS(CASSeq);
RUN_ATOMIC_BENCHMARKS(CASSeqBase);
RUN_ATOMIC_BENCHMARKS(CASStride);
RUN_ATOMIC_BENCHMARKS(CASStrideBase);
}
};
}
} // end namespace vtkm::benchmarking
int main(int, char* [])
{
using Device = VTKM_DEFAULT_DEVICE_ADAPTER_TAG;
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
tracker.ForceDevice(Device{});
try
{
vtkm::benchmarking::BenchmarkAtomicArray<Device>::Run();
}
catch (std::exception& e)
{
std::cerr << "Benchmark encountered an exception: " << e.what() << "\n";
return 1;
}
return 0;
}

34
benchmarking/BenchmarkFieldAlgorithms.cxx
View File

@ -631,8 +631,9 @@ private:
vtkm::cont::ArrayHandle<Value> result;
Timer timer;
vtkm::worklet::DispatcherMapField<InterpolateField, DeviceAdapterTag>().Invoke(
this->EdgePairHandle, this->WeightHandle, this->FieldHandle, result);
vtkm::worklet::DispatcherMapField<InterpolateField> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(this->EdgePairHandle, this->WeightHandle, this->FieldHandle, result);
return timer.GetElapsedTime();
}
@ -663,8 +664,9 @@ private:
vtkm::cont::ArrayHandle<Value> result;
Timer timer;
vtkm::worklet::DispatcherMapField<InterpolateField, DeviceAdapterTag>().Invoke(
dedges, dweight, dfield, result);
vtkm::worklet::DispatcherMapField<InterpolateField> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(dedges, dweight, dfield, result);
return timer.GetElapsedTime();
}
@ -719,7 +721,7 @@ private:
vtkm::Float64 operator()()
{
using EvalWorklet = EvaluateImplicitFunction<vtkm::Sphere>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet>;
auto handle = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
auto function =
@ -727,7 +729,9 @@ private:
EvalWorklet eval(function);
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
EvalDispatcher dispatcher(eval);
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(this->Internal.Points, this->Internal.Result);
return timer.GetElapsedTime();
}
@ -755,13 +759,15 @@ private:
vtkm::Float64 operator()()
{
using EvalWorklet = EvaluateImplicitFunction<vtkm::ImplicitFunction>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet>;
auto sphere = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
EvalWorklet eval(sphere.PrepareForExecution(DeviceAdapterTag()));
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
EvalDispatcher dispatcher(eval);
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(this->Internal.Points, this->Internal.Result);
return timer.GetElapsedTime();
}
@ -789,7 +795,7 @@ private:
vtkm::Float64 operator()()
{
using EvalWorklet = Evaluate2ImplicitFunctions<vtkm::Sphere, vtkm::Sphere>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet>;
auto h1 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
auto h2 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere2);
@ -798,7 +804,9 @@ private:
EvalWorklet eval(f1, f2);
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
EvalDispatcher dispatcher(eval);
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(this->Internal.Points, this->Internal.Result);
return timer.GetElapsedTime();
}
@ -827,7 +835,7 @@ private:
{
using EvalWorklet =
Evaluate2ImplicitFunctions<vtkm::ImplicitFunction, vtkm::ImplicitFunction>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet, DeviceAdapterTag>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvalWorklet>;
auto s1 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere1);
auto s2 = vtkm::cont::make_ImplicitFunctionHandle(Internal.Sphere2);
@ -835,7 +843,9 @@ private:
s2.PrepareForExecution(DeviceAdapterTag()));
vtkm::cont::Timer<DeviceAdapterTag> timer;
EvalDispatcher(eval).Invoke(this->Internal.Points, this->Internal.Result);
EvalDispatcher dispatcher(eval);
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(this->Internal.Points, this->Internal.Result);
return timer.GetElapsedTime();
}

3
benchmarking/BenchmarkFilters.cxx
View File

@ -934,7 +934,8 @@ void CreateFields(bool needPointScalars, bool needCellScalars, bool needPointVec
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> pvecs;
PointVectorGenerator worklet(bounds);
vtkm::worklet::DispatcherMapField<PointVectorGenerator, Device> dispatch(worklet);
vtkm::worklet::DispatcherMapField<PointVectorGenerator> dispatch(worklet);
dispatch.SetDevice(Device());
dispatch.Invoke(points, pvecs);
InputDataSet.AddField(
vtkm::cont::Field("GeneratedPointVectors", vtkm::cont::Field::Association::POINTS, pvecs));

1
benchmarking/CMakeLists.txt
View File

@ -60,6 +60,7 @@ endfunction()
set(benchmarks
BenchmarkArrayTransfer
BenchmarkAtomicArray
BenchmarkCopySpeeds
BenchmarkDeviceAdapter
BenchmarkFieldAlgorithms

26
docs/changelog/agorithms_support_runtime_device.md Normal file
View File

@ -0,0 +1,26 @@
# vtkm::cont::Algorithm now can be told which device to use at runtime
The `vtkm::cont::Algorithm` has been extended to support the user specifying
which device to use at runtime previously Algorithm would only use the first
enabled device, requiring users to modify the `vtkm::cont::GlobalRuntimeDeviceTracker`
if they wanted a specific device used.
To select a specific device with vtkm::cont::Algorithm pass the `vtkm::cont::DeviceAdapterId`
as the first parameter.
```cpp
vtkm::cont::ArrayHandle<double> values;
//call with no tag, will run on first enabled device
auto result = vtkm::cont::Algorithm::Reduce(values, 0.0);
//call with an explicit device tag, will only run on serial
vtkm::cont::DeviceAdapterTagSerial serial;
result = vtkm::cont::Algorithm::Reduce(serial, values, 0.0);
//call with an runtime device tag, will only run on serial
vtkm::cont::DeviceAdapterId device = serial;
result = vtkm::cont::Algorithm::Reduce(device, values, 0.0);
```

23
docs/changelog/array-handle-view.md Normal file
View File

@ -0,0 +1,23 @@
# Add `ArrayHandleView` fancy array
Added a new class named `ArrayHandleView` that allows you to get a subset
of an array. You use the `ArrayHandleView` by giving it a target array, a
starting index, and a length. Here is a simple example of usage:
``` cpp
vtkm::cont::ArrayHandle<vtkm::Id> sourceArray;
vtkm::cont::ArrayCopy(vtkm::cont::ArrayHandleIndex(10), sourceArray);
// sourceArray has [0, 1, 2, 3, 4, 5, 6, 7, 8, 9]
vtkm::cont::ArrayHandleView<vtkm::cont::ArrayHandle<vtkm::Id>>
viewArray(sourceArray, 3, 5);
// viewArray has [3, 4, 5, 6, 7]
```
There is also a convenience `make_ArraHandleView` function to create view
arrays. The following makes the same view array as before.
``` cpp
auto viewArray = vtkm::cont::make_ArrayHandleView(sourceArray, 3, 5);
```

0
docs/cellsetexplicit-cache-celltopoint.md → docs/changelog/cellsetexplicit-cache-celltopoint.md
View File

15
docs/changelog/dispatcher-auto-device.md Normal file
View File

@ -0,0 +1,15 @@
# Make DispatcherBase invoke using a TryExecute
Rather than force all dispatchers to be templated on a device adapter,
instead use a TryExecute internally within the invoke to select a device
adapter.
Because this removes the need to declare a device when invoking a worklet,
this commit also removes the need to declare a device in several other
areas of the code.
This changes touches quite a bit a code. The first pass of the change
usually does the minimum amount of work, which is to change the
compile-time specification of the device to a run-time call to `SetDevice`
on the dispatcher. Although functionally equivalent, it might mean calling
`TryExecute` within itself.

6
examples/game_of_life/GameOfLife.cxx
View File

@ -139,7 +139,7 @@ public:
this->PrintedDeviceMsg = true;
}
using DispatcherType = vtkm::worklet::DispatcherPointNeighborhood<UpdateLifeState, Device>;
using DispatcherType = vtkm::worklet::DispatcherPointNeighborhood<UpdateLifeState>;
vtkm::cont::ArrayHandle<vtkm::UInt8> state;
@ -153,7 +153,9 @@ public:
input.GetField("state", vtkm::cont::Field::Association::POINTS).GetData().CopyTo(prevstate);
//Update the game state
DispatcherType().Invoke(vtkm::filter::ApplyPolicy(cells, policy), prevstate, state, colors);
DispatcherType dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(vtkm::filter::ApplyPolicy(cells, policy), prevstate, state, colors);
//save the results
vtkm::cont::DataSet output;

2
examples/multi_backend/IOGenerator.cxx
View File

@ -53,7 +53,7 @@ vtkm::cont::DataSet make_test3DImageData(vtkm::Id3 dims)
vtkm::cont::DataSet ds = Builder::Create(dims);
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::FloatDefault, 3>> field;
vtkm::worklet::DispatcherMapField<WaveField, vtkm::cont::DeviceAdapterTagSerial> dispatcher;
vtkm::worklet::DispatcherMapField<WaveField> dispatcher;
dispatcher.Invoke(ds.GetCoordinateSystem(), field);
FieldAdd::AddPointField(ds, "vec_field", field);

425
vtkm/cont/Algorithm.h
View File

@ -317,30 +317,71 @@ struct UpperBoundsFunctor
struct Algorithm
{
template <typename T, typename U, class CIn, class COut>
VTKM_CONT static void Copy(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<U, COut>& output)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::CopyFunctor(), input, output);
}
template <typename T, typename U, class CIn, class COut>
VTKM_CONT static void Copy(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<U, COut>& output)
{
vtkm::cont::TryExecute(detail::CopyFunctor(), input, output);
Copy(vtkm::cont::DeviceAdapterTagAny(), input, output);
}
template <typename T, typename U, class CIn, class CStencil, class COut>
VTKM_CONT static void CopyIf(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<U, CStencil>& stencil,
vtkm::cont::ArrayHandle<T, COut>& output)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::CopyIfFunctor(), input, stencil, output);
}
template <typename T, typename U, class CIn, class CStencil, class COut>
VTKM_CONT static void CopyIf(const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<U, CStencil>& stencil,
vtkm::cont::ArrayHandle<T, COut>& output)
{
vtkm::cont::TryExecute(detail::CopyIfFunctor(), input, stencil, output);
CopyIf(vtkm::cont::DeviceAdapterTagAny(), input, stencil, output);
}
template <typename T, typename U, class CIn, class CStencil, class COut, class UnaryPredicate>
VTKM_CONT static void CopyIf(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<U, CStencil>& stencil,
vtkm::cont::ArrayHandle<T, COut>& output,
UnaryPredicate unary_predicate)
{
vtkm::cont::TryExecuteOnDevice(
devId, detail::CopyIfFunctor(), input, stencil, output, unary_predicate);
}
template <typename T, typename U, class CIn, class CStencil, class COut, class UnaryPredicate>
VTKM_CONT static void CopyIf(const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<U, CStencil>& stencil,
vtkm::cont::ArrayHandle<T, COut>& output,
UnaryPredicate unary_predicate)
{
vtkm::cont::TryExecute(detail::CopyIfFunctor(), input, stencil, output, unary_predicate);
CopyIf(vtkm::cont::DeviceAdapterTagAny(), input, stencil, output, unary_predicate);
}
template <typename T, typename U, class CIn, class COut>
VTKM_CONT static bool CopySubRange(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::Id inputStartIndex,
vtkm::Id numberOfElementsToCopy,
vtkm::cont::ArrayHandle<U, COut>& output,
vtkm::Id outputIndex = 0)
{
detail::CopySubRangeFunctor functor;
vtkm::cont::TryExecuteOnDevice(
devId, functor, input, inputStartIndex, numberOfElementsToCopy, output, outputIndex);
return functor.valid;
}
template <typename T, typename U, class CIn, class COut>
VTKM_CONT static bool CopySubRange(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::Id inputStartIndex,
@ -348,54 +389,124 @@ struct Algorithm
vtkm::cont::ArrayHandle<U, COut>& output,
vtkm::Id outputIndex = 0)
{
detail::CopySubRangeFunctor functor;
vtkm::cont::TryExecute(
functor, input, inputStartIndex, numberOfElementsToCopy, output, outputIndex);
return functor.valid;
return CopySubRange(vtkm::cont::DeviceAdapterTagAny(),
input,
inputStartIndex,
numberOfElementsToCopy,
output,
outputIndex);
}
template <typename T, class CIn, class CVal, class COut>
VTKM_CONT static void LowerBounds(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::LowerBoundsFunctor(), input, values, output);
}
template <typename T, class CIn, class CVal, class COut>
VTKM_CONT static void LowerBounds(const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output)
{
vtkm::cont::TryExecute(detail::LowerBoundsFunctor(), input, values, output);
LowerBounds(vtkm::cont::DeviceAdapterTagAny(), input, values, output);
}
template <typename T, class CIn, class CVal, class COut, class BinaryCompare>
VTKM_CONT static void LowerBounds(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecuteOnDevice(
devId, detail::LowerBoundsFunctor(), input, values, output, binary_compare);
}
template <typename T, class CIn, class CVal, class COut, class BinaryCompare>
VTKM_CONT static void LowerBounds(const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecute(detail::LowerBoundsFunctor(), input, values, output, binary_compare);
LowerBounds(vtkm::cont::DeviceAdapterTagAny(), input, values, output, binary_compare);
}
template <class CIn, class COut>
VTKM_CONT static void LowerBounds(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<vtkm::Id, CIn>& input,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& values_output)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::LowerBoundsFunctor(), input, values_output);
}
template <class CIn, class COut>
VTKM_CONT static void LowerBounds(const vtkm::cont::ArrayHandle<vtkm::Id, CIn>& input,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& values_output)
{
vtkm::cont::TryExecute(detail::LowerBoundsFunctor(), input, values_output);
LowerBounds(vtkm::cont::DeviceAdapterTagAny(), input, values_output);
}
template <typename T, typename U, class CIn>
VTKM_CONT static U Reduce(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue)
{
detail::ReduceFunctor<U> functor;
vtkm::cont::TryExecuteOnDevice(devId, functor, input, initialValue);
return functor.result;
}
template <typename T, typename U, class CIn>
VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input, U initialValue)
{
detail::ReduceFunctor<U> functor;
vtkm::cont::TryExecute(functor, input, initialValue);
return functor.result;
return Reduce(vtkm::cont::DeviceAdapterTagAny(), input, initialValue);
}
template <typename T, typename U, class CIn, class BinaryFunctor>
VTKM_CONT static U Reduce(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue,
BinaryFunctor binary_functor)
{
detail::ReduceFunctor<U> functor;
vtkm::cont::TryExecuteOnDevice(devId, functor, input, initialValue, binary_functor);
return functor.result;
}
template <typename T, typename U, class CIn, class BinaryFunctor>
VTKM_CONT static U Reduce(const vtkm::cont::ArrayHandle<T, CIn>& input,
U initialValue,
BinaryFunctor binary_functor)
{
detail::ReduceFunctor<U> functor;
vtkm::cont::TryExecute(functor, input, initialValue, binary_functor);
return functor.result;
return Reduce(vtkm::cont::DeviceAdapterTagAny(), input, initialValue, binary_functor);
}
template <typename T,
typename U,
class CKeyIn,
class CValIn,
class CKeyOut,
class CValOut,
class BinaryFunctor>
VTKM_CONT static void ReduceByKey(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CKeyIn>& keys,
const vtkm::cont::ArrayHandle<U, CValIn>& values,
vtkm::cont::ArrayHandle<T, CKeyOut>& keys_output,
vtkm::cont::ArrayHandle<U, CValOut>& values_output,
BinaryFunctor binary_functor)
{
vtkm::cont::TryExecuteOnDevice(devId,
detail::ReduceByKeyFunctor(),
keys,
values,
keys_output,
values_output,
binary_functor);
}
template <typename T,
typename U,
class CKeyIn,
@ -409,39 +520,81 @@ struct Algorithm
vtkm::cont::ArrayHandle<U, CValOut>& values_output,
BinaryFunctor binary_functor)
{
vtkm::cont::TryExecute(
detail::ReduceByKeyFunctor(), keys, values, keys_output, values_output, binary_functor);
ReduceByKey(
vtkm::cont::DeviceAdapterTagAny(), keys, values, keys_output, values_output, binary_functor);
}
template <typename T, class CIn, class COut>
VTKM_CONT static T ScanInclusive(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
detail::ScanInclusiveResultFunctor<T> functor;
vtkm::cont::TryExecuteOnDevice(devId, functor, input, output);
return functor.result;
}
template <typename T, class CIn, class COut>
VTKM_CONT static T ScanInclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
detail::ScanInclusiveResultFunctor<T> functor;
vtkm::cont::TryExecute(functor, input, output);
return functor.result;
return ScanInclusive(vtkm::cont::DeviceAdapterTagAny(), input, output);
}
template <typename T, class CIn, class COut>
VTKM_CONT static T StreamingScanExclusive(vtkm::cont::DeviceAdapterId devId,
const vtkm::Id numBlocks,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
detail::StreamingScanExclusiveFunctor<T> functor;
vtkm::cont::TryExecuteOnDevice(devId, functor, numBlocks, input, output);
return functor.result;
}
template <typename T, class CIn, class COut>
VTKM_CONT static T StreamingScanExclusive(const vtkm::Id numBlocks,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
detail::StreamingScanExclusiveFunctor<T> functor;
vtkm::cont::TryExecute(functor, numBlocks, input, output);
return functor.result;
return StreamingScanExclusive(vtkm::cont::DeviceAdapterTagAny(), numBlocks, input, output);
}
template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanInclusive(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binary_functor)
{
detail::ScanInclusiveResultFunctor<T> functor;
vtkm::cont::TryExecuteOnDevice(devId, functor, input, output, binary_functor);
return functor.result;
}
template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanInclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binary_functor)
{
detail::ScanInclusiveResultFunctor<T> functor;
vtkm::cont::TryExecute(functor, input, output, binary_functor);
return functor.result;
return ScanInclusive(vtkm::cont::DeviceAdapterTagAny(), input, output, binary_functor);
}
template <typename T,
typename U,
typename KIn,
typename VIn,
typename VOut,
typename BinaryFunctor>
VTKM_CONT static void ScanInclusiveByKey(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& values_output,
BinaryFunctor binary_functor)
{
vtkm::cont::TryExecuteOnDevice(
devId, detail::ScanInclusiveByKeyFunctor(), keys, values, values_output, binary_functor);
}
template <typename T,
typename U,
typename KIn,
@ -453,38 +606,84 @@ struct Algorithm
vtkm::cont::ArrayHandle<U, VOut>& values_output,
BinaryFunctor binary_functor)
{
vtkm::cont::TryExecute(
detail::ScanInclusiveByKeyFunctor(), keys, values, values_output, binary_functor);
ScanInclusiveByKey(
vtkm::cont::DeviceAdapterTagAny(), keys, values, values_output, binary_functor);
}
template <typename T, typename U, typename KIn, typename VIn, typename VOut>
VTKM_CONT static void ScanInclusiveByKey(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& values_output)
{
vtkm::cont::TryExecuteOnDevice(
devId, detail::ScanInclusiveByKeyFunctor(), keys, values, values_output);
}
template <typename T, typename U, typename KIn, typename VIn, typename VOut>
VTKM_CONT static void ScanInclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& values_output)
{
vtkm::cont::TryExecute(detail::ScanInclusiveByKeyFunctor(), keys, values, values_output);
ScanInclusiveByKey(vtkm::cont::DeviceAdapterTagAny(), keys, values, values_output);
}
template <typename T, class CIn, class COut>
VTKM_CONT static T ScanExclusive(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
detail::ScanExclusiveFunctor<T> functor;
vtkm::cont::TryExecuteOnDevice(devId, functor, input, output);
return functor.result;
}
template <typename T, class CIn, class COut>
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output)
{
detail::ScanExclusiveFunctor<T> functor;
vtkm::cont::TryExecute(functor, input, output);
return functor.result;
return ScanExclusive(vtkm::cont::DeviceAdapterTagAny(), input, output);
}
template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanExclusive(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binaryFunctor,
const T& initialValue)
{
detail::ScanExclusiveFunctor<T> functor;
vtkm::cont::TryExecuteOnDevice(devId, functor, input, output, binaryFunctor, initialValue);
return functor.result;
}
template <typename T, class CIn, class COut, class BinaryFunctor>
VTKM_CONT static T ScanExclusive(const vtkm::cont::ArrayHandle<T, CIn>& input,
vtkm::cont::ArrayHandle<T, COut>& output,
BinaryFunctor binaryFunctor,
const T& initialValue)
{
detail::ScanExclusiveFunctor<T> functor;
vtkm::cont::TryExecute(functor, input, output, binaryFunctor, initialValue);
return functor.result;
return ScanExclusive(
vtkm::cont::DeviceAdapterTagAny(), input, output, binaryFunctor, initialValue);
}
template <typename T, typename U, typename KIn, typename VIn, typename VOut, class BinaryFunctor>
VTKM_CONT static void ScanExclusiveByKey(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& output,
const U& initialValue,
BinaryFunctor binaryFunctor)
{
vtkm::cont::TryExecuteOnDevice(devId,
detail::ScanExclusiveByKeyFunctor(),
keys,
values,
output,
initialValue,
binaryFunctor);
}
template <typename T, typename U, typename KIn, typename VIn, typename VOut, class BinaryFunctor>
VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values,
@ -492,60 +691,139 @@ struct Algorithm
const U& initialValue,
BinaryFunctor binaryFunctor)
{
vtkm::cont::TryExecute(
detail::ScanExclusiveByKeyFunctor(), keys, values, output, initialValue, binaryFunctor);
ScanExclusiveByKey(
vtkm::cont::DeviceAdapterTagAny(), keys, values, output, initialValue, binaryFunctor);
}
template <typename T, typename U, class KIn, typename VIn, typename VOut>
VTKM_CONT static void ScanExclusiveByKey(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& output)
{
vtkm::cont::TryExecuteOnDevice(
devId, detail::ScanExclusiveByKeyFunctor(), keys, values, output);
}
template <typename T, typename U, class KIn, typename VIn, typename VOut>
VTKM_CONT static void ScanExclusiveByKey(const vtkm::cont::ArrayHandle<T, KIn>& keys,
const vtkm::cont::ArrayHandle<U, VIn>& values,
vtkm::cont::ArrayHandle<U, VOut>& output)
{
vtkm::cont::TryExecute(detail::ScanExclusiveByKeyFunctor(), keys, values, output);
ScanExclusiveByKey(vtkm::cont::DeviceAdapterTagAny(), keys, values, output);
}
template <class Functor>
VTKM_CONT static void Schedule(vtkm::cont::DeviceAdapterId devId,
Functor functor,
vtkm::Id numInstances)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::ScheduleFunctor(), functor, numInstances);
}
template <class Functor>
VTKM_CONT static void Schedule(Functor functor, vtkm::Id numInstances)
{
vtkm::cont::TryExecute(detail::ScheduleFunctor(), functor, numInstances);
Schedule(vtkm::cont::DeviceAdapterTagAny(), functor, numInstances);
}
template <class Functor>
VTKM_CONT static void Schedule(vtkm::cont::DeviceAdapterId devId,
Functor functor,
vtkm::Id3 rangeMax)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::ScheduleFunctor(), functor, rangeMax);
}
template <class Functor>
VTKM_CONT static void Schedule(Functor functor, vtkm::Id3 rangeMax)
{
vtkm::cont::TryExecute(detail::ScheduleFunctor(), functor, rangeMax);
Schedule(vtkm::cont::DeviceAdapterTagAny(), functor, rangeMax);
}
template <typename T, class Storage>
VTKM_CONT static void Sort(vtkm::cont::DeviceAdapterId devId,
vtkm::cont::ArrayHandle<T, Storage>& values)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::SortFunctor(), values);
}
template <typename T, class Storage>
VTKM_CONT static void Sort(vtkm::cont::ArrayHandle<T, Storage>& values)
{
vtkm::cont::TryExecute(detail::SortFunctor(), values);
Sort(vtkm::cont::DeviceAdapterTagAny(), values);
}
template <typename T, class Storage, class BinaryCompare>
VTKM_CONT static void Sort(vtkm::cont::DeviceAdapterId devId,
vtkm::cont::ArrayHandle<T, Storage>& values,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::SortFunctor(), values, binary_compare);
}
template <typename T, class Storage, class BinaryCompare>
VTKM_CONT static void Sort(vtkm::cont::ArrayHandle<T, Storage>& values,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecute(detail::SortFunctor(), values, binary_compare);
Sort(vtkm::cont::DeviceAdapterTagAny(), values, binary_compare);
}
template <typename T, typename U, class StorageT, class StorageU>
VTKM_CONT static void SortByKey(vtkm::cont::DeviceAdapterId devId,
vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::SortByKeyFunctor(), keys, values);
}
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)
{
vtkm::cont::TryExecute(detail::SortByKeyFunctor(), keys, values);
SortByKey(vtkm::cont::DeviceAdapterTagAny(), keys, values);
}
template <typename T, typename U, class StorageT, class StorageU, class BinaryCompare>
VTKM_CONT static void SortByKey(vtkm::cont::DeviceAdapterId devId,
vtkm::cont::ArrayHandle<T, StorageT>& keys,
vtkm::cont::ArrayHandle<U, StorageU>& values,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::SortByKeyFunctor(), keys, values, binary_compare);
}
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)
{
vtkm::cont::TryExecute(detail::SortByKeyFunctor(), keys, values, binary_compare);
SortByKey(vtkm::cont::DeviceAdapterTagAny(), keys, values, binary_compare);
}
VTKM_CONT static void Synchronize() { vtkm::cont::TryExecute(detail::SynchronizeFunctor()); }
VTKM_CONT static void Synchronize(vtkm::cont::DeviceAdapterId devId)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::SynchronizeFunctor());
}
VTKM_CONT static void Synchronize() { Synchronize(vtkm::cont::DeviceAdapterTagAny()); }
template <typename T,
typename U,
typename V,
typename StorageT,
typename StorageU,
typename StorageV,
typename BinaryFunctor>
VTKM_CONT static void Transform(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, StorageT>& input1,
const vtkm::cont::ArrayHandle<U, StorageU>& input2,
vtkm::cont::ArrayHandle<V, StorageV>& output,
BinaryFunctor binaryFunctor)
{
vtkm::cont::TryExecuteOnDevice(
devId, detail::TransformFunctor(), input1, input2, output, binaryFunctor);
}
template <typename T,
typename U,
typename V,
@ -558,44 +836,87 @@ struct Algorithm
vtkm::cont::ArrayHandle<V, StorageV>& output,
BinaryFunctor binaryFunctor)
{
vtkm::cont::TryExecute(detail::TransformFunctor(), input1, input2, output, binaryFunctor);
Transform(vtkm::cont::DeviceAdapterTagAny(), input1, input2, output, binaryFunctor);
}
template <typename T, class Storage>
VTKM_CONT static void Unique(vtkm::cont::DeviceAdapterId devId,
vtkm::cont::ArrayHandle<T, Storage>& values)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::UniqueFunctor(), values);
}
template <typename T, class Storage>
VTKM_CONT static void Unique(vtkm::cont::ArrayHandle<T, Storage>& values)
{
vtkm::cont::TryExecute(detail::UniqueFunctor(), values);
Unique(vtkm::cont::DeviceAdapterTagAny(), values);
}
template <typename T, class Storage, class BinaryCompare>
VTKM_CONT static void Unique(vtkm::cont::DeviceAdapterId devId,
vtkm::cont::ArrayHandle<T, Storage>& values,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::UniqueFunctor(), values, binary_compare);
}
template <typename T, class Storage, class BinaryCompare>
VTKM_CONT static void Unique(vtkm::cont::ArrayHandle<T, Storage>& values,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecute(detail::UniqueFunctor(), values, binary_compare);
Unique(vtkm::cont::DeviceAdapterTagAny(), values, binary_compare);
}
template <typename T, class CIn, class CVal, class COut>
VTKM_CONT static void UpperBounds(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::UpperBoundsFunctor(), input, values, output);
}
template <typename T, class CIn, class CVal, class COut>
VTKM_CONT static void UpperBounds(const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output)
{
vtkm::cont::TryExecute(detail::UpperBoundsFunctor(), input, values, output);
UpperBounds(vtkm::cont::DeviceAdapterTagAny(), input, values, output);
}
template <typename T, class CIn, class CVal, class COut, class BinaryCompare>
VTKM_CONT static void UpperBounds(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecuteOnDevice(
devId, detail::UpperBoundsFunctor(), input, values, output, binary_compare);
}
template <typename T, class CIn, class CVal, class COut, class BinaryCompare>
VTKM_CONT static void UpperBounds(const vtkm::cont::ArrayHandle<T, CIn>& input,
const vtkm::cont::ArrayHandle<T, CVal>& values,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& output,
BinaryCompare binary_compare)
{
vtkm::cont::TryExecute(detail::UpperBoundsFunctor(), input, values, output, binary_compare);
UpperBounds(vtkm::cont::DeviceAdapterTagAny(), input, values, output, binary_compare);
}
template <class CIn, class COut>
VTKM_CONT static void UpperBounds(vtkm::cont::DeviceAdapterId devId,
const vtkm::cont::ArrayHandle<vtkm::Id, CIn>& input,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& values_output)
{
vtkm::cont::TryExecuteOnDevice(devId, detail::UpperBoundsFunctor(), input, values_output);
}
template <class CIn, class COut>
VTKM_CONT static void UpperBounds(const vtkm::cont::ArrayHandle<vtkm::Id, CIn>& input,
vtkm::cont::ArrayHandle<vtkm::Id, COut>& values_output)
{
vtkm::cont::TryExecute(detail::UpperBoundsFunctor(), input, values_output);
UpperBounds(vtkm::cont::DeviceAdapterTagAny(), input, values_output);
}
};
}

2
vtkm/cont/ArrayHandle.h
View File

@ -498,7 +498,7 @@ public:
{
return this->Internals->ExecutionArrayValid
? this->Internals->ExecutionArray->GetDeviceAdapterId()
: DeviceAdapterIdUndefined{};
: DeviceAdapterTagUndefined{};
}
struct VTKM_ALWAYS_EXPORT InternalStruct

310
vtkm/cont/ArrayHandleView.h Normal file
View File

@ -0,0 +1,310 @@
//============================================================================
// 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.
//
// Copyright 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2018 UT-Battelle, LLC.
// Copyright 2018 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#ifndef vtk_m_cont_ArrayHandleView_h
#define vtk_m_cont_ArrayHandleView_h
#include <vtkm/Assert.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/ArrayPortal.h>
namespace vtkm
{
namespace cont
{
namespace internal
{
template <typename TargetPortalType>
class ArrayPortalView
{
public:
using ValueType = typename TargetPortalType::ValueType;
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
ArrayPortalView() {}
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
ArrayPortalView(const TargetPortalType& targetPortal, vtkm::Id startIndex, vtkm::Id numValues)
: TargetPortal(targetPortal)
, StartIndex(startIndex)
, NumValues(numValues)
{
}
VTKM_SUPPRESS_EXEC_WARNINGS
template <typename OtherPortalType>
VTKM_EXEC_CONT ArrayPortalView(const ArrayPortalView<OtherPortalType>& otherPortal)
: TargetPortal(otherPortal.GetTargetPortal())
, StartIndex(otherPortal.GetStartIndex())
, NumValues(otherPortal.GetNumberOfValues())
{
}
VTKM_EXEC_CONT
vtkm::Id GetNumberOfValues() const { return this->NumValues; }
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
ValueType Get(vtkm::Id index) const { return this->TargetPortal.Get(index + this->StartIndex); }
VTKM_SUPPRESS_EXEC_WARNINGS
VTKM_EXEC_CONT
void Set(vtkm::Id index, const ValueType& value) const
{
this->TargetPortal.Set(index + this->StartIndex, value);
}
VTKM_EXEC_CONT
const TargetPortalType& GetTargetPortal() const { return this->TargetPortal; }
VTKM_EXEC_CONT
vtkm::Id GetStartIndex() const { return this->StartIndex; }
private:
TargetPortalType TargetPortal;
vtkm::Id StartIndex;
vtkm::Id NumValues;
};
} // namespace internal
template <typename ArrayHandleType>
struct StorageTagView
{
};
namespace internal
{
template <typename ArrayHandleType>
class Storage<typename ArrayHandleType::ValueType, StorageTagView<ArrayHandleType>>
{
public:
using ValueType = typename ArrayHandleType::ValueType;
using PortalType = ArrayPortalView<typename ArrayHandleType::PortalControl>;
using PortalConstType = ArrayPortalView<typename ArrayHandleType::PortalConstControl>;
VTKM_CONT
Storage()
: Valid(false)
{
}
VTKM_CONT
Storage(const ArrayHandleType& array, vtkm::Id startIndex, vtkm::Id numValues)
: Array(array)
, StartIndex(startIndex)
, NumValues(numValues)
, Valid(true)
{
VTKM_ASSERT(this->StartIndex >= 0);
VTKM_ASSERT((this->StartIndex + this->NumValues) <= this->Array.GetNumberOfValues());
}
VTKM_CONT
PortalType GetPortal()
{
VTKM_ASSERT(this->Valid);
return PortalType(this->Array.GetPortalControl(), this->StartIndex, this->NumValues);
}
VTKM_CONT
PortalConstType GetPortalConst() const
{
VTKM_ASSERT(this->Valid);
return PortalConstType(this->Array.GetPortalConstControl(), this->StartIndex, this->NumValues);
}
VTKM_CONT
vtkm::Id GetNumberOfValues() const { return this->NumValues; }
VTKM_CONT
void Allocate(vtkm::Id vtkmNotUsed(numberOfValues))
{
throw vtkm::cont::ErrorInternal("ArrayHandleView should not be allocated explicitly. ");
}
VTKM_CONT
void Shrink(vtkm::Id numberOfValues)
{
VTKM_ASSERT(this->Valid);
if (numberOfValues > this->NumValues)
{
throw vtkm::cont::ErrorBadValue("Shrink method cannot be used to grow array.");
}
this->NumValues = numberOfValues;
}
VTKM_CONT
void ReleaseResources()
{
VTKM_ASSERT(this->Valid);
this->Array.ReleaseResources();
}
// Requried for later use in ArrayTransfer class.
VTKM_CONT
const ArrayHandleType& GetArray() const
{
VTKM_ASSERT(this->Valid);
return this->Array;
}
VTKM_CONT
vtkm::Id GetStartIndex() const { return this->StartIndex; }
private:
ArrayHandleType Array;
vtkm::Id StartIndex;
vtkm::Id NumValues;
bool Valid;
};
template <typename ArrayHandleType, typename Device>
class ArrayTransfer<typename ArrayHandleType::ValueType, StorageTagView<ArrayHandleType>, Device>
{
public:
using ValueType = typename ArrayHandleType::ValueType;
private:
using StorageTag = StorageTagView<ArrayHandleType>;
using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;
public:
using PortalControl = typename StorageType::PortalType;
using PortalConstControl = typename StorageType::PortalConstType;
using PortalExecution =
ArrayPortalView<typename ArrayHandleType::template ExecutionTypes<Device>::Portal>;
using PortalConstExecution =
ArrayPortalView<typename ArrayHandleType::template ExecutionTypes<Device>::PortalConst>;
VTKM_CONT
ArrayTransfer(StorageType* storage)
: Array(storage->GetArray())
, StartIndex(storage->GetStartIndex())
, NumValues(storage->GetNumberOfValues())
{
}
VTKM_CONT
vtkm::Id GetNumberOfValues() const { return this->NumValues; }
VTKM_CONT
PortalConstExecution PrepareForInput(bool vtkmNotUsed(updateData))
{
return PortalConstExecution(
this->Array.PrepareForInput(Device()), this->StartIndex, this->NumValues);
}
VTKM_CONT
PortalExecution PrepareForInPlace(bool vtkmNotUsed(updateData))
{
return PortalExecution(
this->Array.PrepareForInPlace(Device()), this->StartIndex, this->NumValues);
}
VTKM_CONT
PortalExecution PrepareForOutput(vtkm::Id numberOfValues)
{
if (numberOfValues != this->GetNumberOfValues())
{
throw vtkm::cont::ErrorBadValue(
"An ArrayHandleView can be used as an output array, "
"but it cannot be resized. Make sure the index array is sized "
"to the appropriate length before trying to prepare for output.");
}
// We cannot practically allocate ValueArray because we do not know the
// range of indices. We try to check by seeing if ValueArray has no
// entries, which clearly indicates that it is not allocated. Otherwise,
// we have to assume the allocation is correct.
if ((numberOfValues > 0) && (this->Array.GetNumberOfValues() < 1))
{
throw vtkm::cont::ErrorBadValue(
"The value array must be pre-allocated before it is used for the "
"output of ArrayHandlePermutation.");
}
return PortalExecution(this->Array.PrepareForOutput(this->Array.GetNumberOfValues(), Device()),
this->StartIndex,
this->NumValues);
}
VTKM_CONT
void RetrieveOutputData(StorageType* vtkmNotUsed(storage)) const
{
// No implementation necessary
}
VTKM_CONT
void Shrink(vtkm::Id numberOfValues) { this->NumValues = numberOfValues; }
VTKM_CONT
void ReleaseResources() { this->Array.ReleaseResourcesExecution(); }
private:
ArrayHandleType Array;
vtkm::Id StartIndex;
vtkm::Id NumValues;
};
} // namespace internal
template <typename ArrayHandleType>
class ArrayHandleView : public vtkm::cont::ArrayHandle<typename ArrayHandleType::ValueType,
StorageTagView<ArrayHandleType>>
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
public:
VTKM_ARRAY_HANDLE_SUBCLASS(ArrayHandleView,
(ArrayHandleView<ArrayHandleType>),
(vtkm::cont::ArrayHandle<typename ArrayHandleType::ValueType,
StorageTagView<ArrayHandleType>>));
private:
using StorageType = vtkm::cont::internal::Storage<ValueType, StorageTag>;
public:
VTKM_CONT
ArrayHandleView(const ArrayHandleType& array, vtkm::Id startIndex, vtkm::Id numValues)
: Superclass(StorageType(array, startIndex, numValues))
{
}
};
template <typename ArrayHandleType>
ArrayHandleView<ArrayHandleType> make_ArrayHandleView(const ArrayHandleType& array,
vtkm::Id startIndex,
vtkm::Id numValues)
{
VTKM_IS_ARRAY_HANDLE(ArrayHandleType);
return ArrayHandleView<ArrayHandleType>(array, startIndex, numValues);
}
}
} // namespace vtkm::cont
#endif //vtk_m_cont_ArrayHandleView_h

17
vtkm/cont/BoundingIntervalHierarchy.hxx
View File

@ -135,15 +135,16 @@ VTKM_CONT void BoundingIntervalHierarchy::CalculatePlaneSplitCost(
vtkm::cont::ArrayHandle<vtkm::FloatDefault> splitPlanes;
vtkm::worklet::spatialstructure::SplitPlaneCalculatorWorklet splitPlaneCalcWorklet(planeIndex,
numPlanes);
vtkm::worklet::DispatcherMapField<vtkm::worklet::spatialstructure::SplitPlaneCalculatorWorklet,
DeviceAdapter>
vtkm::worklet::DispatcherMapField<vtkm::worklet::spatialstructure::SplitPlaneCalculatorWorklet>
splitDispatcher(splitPlaneCalcWorklet);
splitDispatcher.SetDevice(DeviceAdapter());
splitDispatcher.Invoke(segmentRanges, splitPlanes);
// Check if a point is to the left of the split plane or right
vtkm::cont::ArrayHandle<vtkm::Id> isLEQOfSplitPlane, isROfSplitPlane;
vtkm::worklet::DispatcherMapField<vtkm::worklet::spatialstructure::LEQWorklet, DeviceAdapter>()
.Invoke(coords, splitPlanes, isLEQOfSplitPlane, isROfSplitPlane);
vtkm::worklet::DispatcherMapField<vtkm::worklet::spatialstructure::LEQWorklet> LEQDispatcher;
LEQDispatcher.SetDevice(DeviceAdapter());
LEQDispatcher.Invoke(coords, splitPlanes, isLEQOfSplitPlane, isROfSplitPlane);
// Count of points to the left
vtkm::cont::ArrayHandle<vtkm::Id> pointsToLeft;
@ -270,9 +271,11 @@ public:
//START_TIMER(s12);
CoordsArrayHandle centerXs, centerYs, centerZs;
RangeArrayHandle xRanges, yRanges, zRanges;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::spatialstructure::CellRangesExtracter,
DeviceAdapter>()
.Invoke(cellSet, points, xRanges, yRanges, zRanges, centerXs, centerYs, centerZs);
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::spatialstructure::CellRangesExtracter>
cellRangesExtracterDispatcher;
cellRangesExtracterDispatcher.SetDevice(DeviceAdapter());
cellRangesExtracterDispatcher.Invoke(
cellSet, points, xRanges, yRanges, zRanges, centerXs, centerYs, centerZs);
//PRINT_TIMER("1.2", s12);
bool done = false;

2
vtkm/cont/CMakeLists.txt
View File

@ -39,6 +39,7 @@ set(headers
ArrayHandleSwizzle.h
ArrayHandleTransform.h
ArrayHandleUniformPointCoordinates.h
ArrayHandleView.h
ArrayHandleVirtualCoordinates.h
ArrayHandleZip.h
ArrayPortal.h
@ -133,6 +134,7 @@ set(sources
FieldRangeGlobalCompute.cxx
internal/ArrayHandleBasicImpl.cxx
internal/ArrayManagerExecutionShareWithControl.cxx
internal/DeviceAdapterTag.cxx
internal/SimplePolymorphicContainer.cxx
MultiBlock.cxx
PresetColorTables.cxx

5
vtkm/cont/CellLocatorHelper.h
View File

@ -138,8 +138,9 @@ public:
auto coordinates =
this->Coordinates.GetData().Cast<vtkm::cont::ArrayHandleUniformPointCoordinates>();
auto cellset = this->CellSet.ResetCellSetList(StructuredCellSetList());
vtkm::worklet::DispatcherMapField<FindCellWorklet, DeviceAdapter>().Invoke(
points, cellset, coordinates, cellIds, parametricCoords);
vtkm::worklet::DispatcherMapField<FindCellWorklet> dispatcher;
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(points, cellset, coordinates, cellIds, parametricCoords);
}
else
{

44
vtkm/cont/CellLocatorTwoLevelUniformGrid.h
View File

@ -513,8 +513,9 @@ public:
// 2: For each cell, find the number of top level bins they intersect
vtkm::cont::ArrayHandle<vtkm::Id> binCounts;
CountBinsL1 countL1(ls.TopLevel);
vtkm::worklet::DispatcherMapTopology<CountBinsL1, DeviceAdapter>(countL1).Invoke(
cellset, coords, binCounts);
vtkm::worklet::DispatcherMapTopology<CountBinsL1> countL1Dispatcher(countL1);
countL1Dispatcher.SetDevice(DeviceAdapter());
countL1Dispatcher.Invoke(cellset, coords, binCounts);
// 3: Total number of unique (cell, bin) pairs (for pre-allocating arrays)
vtkm::Id numPairsL1 = Algorithm::ScanExclusive(binCounts, binCounts);
@ -523,8 +524,9 @@ public:
vtkm::cont::ArrayHandle<vtkm::Id> binIds;
binIds.Allocate(numPairsL1);
FindBinsL1 findL1(ls.TopLevel);
vtkm::worklet::DispatcherMapTopology<FindBinsL1, DeviceAdapter>(findL1).Invoke(
cellset, coords, binCounts, binIds);
vtkm::worklet::DispatcherMapTopology<FindBinsL1> findL1Dispatcher(findL1);
findL1Dispatcher.SetDevice(DeviceAdapter());
findL1Dispatcher.Invoke(cellset, coords, binCounts, binIds);
binCounts.ReleaseResources();
// 5: From above, find the number of cells that intersect each top level bin
@ -544,8 +546,9 @@ public:
Algorithm::Copy(vtkm::cont::make_ArrayHandleConstant(DimVec3(0), numberOfBins),
ls.LeafDimensions);
GenerateBinsL1 generateL1(ls.TopLevel.BinSize, this->DensityL2);
vtkm::worklet::DispatcherMapField<GenerateBinsL1, DeviceAdapter>(generateL1)
.Invoke(bins, cellsPerBin, ls.LeafDimensions);
vtkm::worklet::DispatcherMapField<GenerateBinsL1> generateL1Dispatcher(generateL1);
generateL1Dispatcher.SetDevice(DeviceAdapter());
generateL1Dispatcher.Invoke(bins, cellsPerBin, ls.LeafDimensions);
bins.ReleaseResources();
cellsPerBin.ReleaseResources();
@ -557,8 +560,9 @@ public:
// 8: For each cell, find the number of l2 bins they intersect
CountBinsL2 countL2(ls.TopLevel);
vtkm::worklet::DispatcherMapTopology<CountBinsL2, DeviceAdapter>(countL2).Invoke(
cellset, coords, ls.LeafDimensions, binCounts);
vtkm::worklet::DispatcherMapTopology<CountBinsL2> countL2Dispatcher(countL2);
countL2Dispatcher.SetDevice(DeviceAdapter());
countL2Dispatcher.Invoke(cellset, coords, ls.LeafDimensions, binCounts);
// 9: Total number of unique (cell, bin) pairs (for pre-allocating arrays)
vtkm::Id numPairsL2 = Algorithm::ScanExclusive(binCounts, binCounts);
@ -567,7 +571,9 @@ public:
binIds.Allocate(numPairsL2);
ls.CellIds.Allocate(numPairsL2);
FindBinsL2 findL2(ls.TopLevel);
vtkm::worklet::DispatcherMapTopology<FindBinsL2, DeviceAdapter>(findL2).Invoke(
vtkm::worklet::DispatcherMapTopology<FindBinsL2> findL2Dispatcher(findL2);
findL2Dispatcher.SetDevice(DeviceAdapter());
findL2Dispatcher.Invoke(
cellset, coords, ls.LeafDimensions, ls.LeafStartIndex, binCounts, binIds, ls.CellIds);
binCounts.ReleaseResources();
@ -587,8 +593,9 @@ public:
Algorithm::Copy(vtkm::cont::ArrayHandleConstant<vtkm::Id>(0, numberOfLeaves),
ls.CellStartIndex);
Algorithm::Copy(vtkm::cont::ArrayHandleConstant<vtkm::Id>(0, numberOfLeaves), ls.CellCount);
vtkm::worklet::DispatcherMapField<GenerateBinsL2, DeviceAdapter>().Invoke(
bins, cellsStart, cellsPerBin, ls.CellStartIndex, ls.CellCount);
vtkm::worklet::DispatcherMapField<GenerateBinsL2> dispatchGenerateBinsL2;
dispatchGenerateBinsL2.SetDevice(DeviceAdapter());
dispatchGenerateBinsL2.Invoke(bins, cellsStart, cellsPerBin, ls.CellStartIndex, ls.CellCount);
std::swap(this->LookupStructure, ls);
}
@ -708,13 +715,14 @@ public:
DeviceAdapter,
CellSetList cellSetTypes = CellSetList()) const
{
vtkm::worklet::DispatcherMapField<FindCellWorklet, DeviceAdapter>().Invoke(
points,
this->CellSet.ResetCellSetList(cellSetTypes),
this->Coordinates,
this->LookupStructure,
cellIds,
parametricCoords);
vtkm::worklet::DispatcherMapField<FindCellWorklet> dispatcher;
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(points,
this->CellSet.ResetCellSetList(cellSetTypes),
this->Coordinates,
this->LookupStructure,
cellIds,
parametricCoords);
}
private:

17
vtkm/cont/CellLocatorUniformGrid.h
View File

@ -48,17 +48,18 @@ public:
throw vtkm::cont::ErrorInternal("Cell set is not 3D structured type");
Bounds = coords.GetBounds();
Dims = cellSet.Cast<StructuredType>().GetSchedulingRange(vtkm::TopologyElementTagPoint());
vtkm::Vec<vtkm::Id, 3> dims =
cellSet.Cast<StructuredType>().GetSchedulingRange(vtkm::TopologyElementTagPoint());
RangeTransform[0] = static_cast<vtkm::FloatDefault>((Dims[0] - 1.0l) / (Bounds.X.Lenght()));
RangeTransform[1] = static_cast<vtkm::FloatDefault>((Dims[1] - 1.0l) / (Bounds.Y.Length()));
RangeTransform[2] = static_cast<vtkm::FloatDefault>((Dims[2] - 1.0l) / (Bounds.Z.Length()));
RangeTransform[0] = static_cast<vtkm::FloatDefault>((dims[0] - 1.0l) / (Bounds.X.Length()));
RangeTransform[1] = static_cast<vtkm::FloatDefault>((dims[1] - 1.0l) / (Bounds.Y.Length()));
RangeTransform[2] = static_cast<vtkm::FloatDefault>((dims[2] - 1.0l) / (Bounds.Z.Length()));
// Since we are calculating the cell Id, and the number of cells is
// 1 less than the number of points in each direction, the -1 from Dims
// 1 less than the number of points in each direction, the -1 from dims
// is necessary.
PlaneSize = (Dims[0] - 1) * (Dims[1] - 1);
RowSize = (Dims[0] - 1);
PlaneSize = (dims[0] - 1) * (dims[1] - 1);
RowSize = (dims[0] - 1);
}
struct PrepareForExecutionFunctor
@ -71,7 +72,6 @@ public:
using ExecutionType = vtkm::exec::CellLocatorUniformGrid<DeviceAdapter>;
ExecutionType* execObject =
new ExecutionType(contLocator.Bounds,
contLocator.Dims,
contLocator.RangeTransform,
contLocator.PlaneSize,
contLocator.RowSize,
@ -101,7 +101,6 @@ private:
using StructuredType = vtkm::cont::CellSetStructured<3>;
vtkm::Bounds Bounds;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::Vec<vtkm::FloatDefault, 3> RangeTransform;
vtkm::Id PlaneSize;
vtkm::Id RowSize;

8
vtkm/cont/CellSetExplicit.hxx
View File

@ -405,7 +405,7 @@ VTKM_CONT auto CellSetExplicit<ShapeStorageTag,
OffsetsStorageTag>::GetShapesArray(FromTopology, ToTopology) const
-> const typename ConnectivityChooser<FromTopology, ToTopology>::ShapeArrayType&
{
this->BuildConnectivity(vtkm::cont::DeviceAdapterIdAny{}, FromTopology(), ToTopology());
this->BuildConnectivity(vtkm::cont::DeviceAdapterTagAny{}, FromTopology(), ToTopology());
return this->GetConnectivity(FromTopology(), ToTopology()).Shapes;
}
@ -421,7 +421,7 @@ VTKM_CONT auto CellSetExplicit<ShapeStorageTag,
ToTopology) const -> const
typename ConnectivityChooser<FromTopology, ToTopology>::NumIndicesArrayType&
{
this->BuildConnectivity(vtkm::cont::DeviceAdapterIdAny{}, FromTopology(), ToTopology());
this->BuildConnectivity(vtkm::cont::DeviceAdapterTagAny{}, FromTopology(), ToTopology());
return this->GetConnectivity(FromTopology(), ToTopology()).NumIndices;
}
@ -437,7 +437,7 @@ VTKM_CONT auto CellSetExplicit<ShapeStorageTag,
ToTopology) const -> const
typename ConnectivityChooser<FromTopology, ToTopology>::ConnectivityArrayType&
{
this->BuildConnectivity(vtkm::cont::DeviceAdapterIdAny{}, FromTopology(), ToTopology());
this->BuildConnectivity(vtkm::cont::DeviceAdapterTagAny{}, FromTopology(), ToTopology());
return this->GetConnectivity(FromTopology(), ToTopology()).Connectivity;
}
@ -453,7 +453,7 @@ VTKM_CONT auto CellSetExplicit<ShapeStorageTag,
ToTopology) const -> const
typename ConnectivityChooser<FromTopology, ToTopology>::IndexOffsetArrayType&
{
this->BuildConnectivity(vtkm::cont::DeviceAdapterIdAny{}, FromTopology(), ToTopology());
this->BuildConnectivity(vtkm::cont::DeviceAdapterTagAny{}, FromTopology(), ToTopology());
return this->GetConnectivity(FromTopology(), ToTopology()).IndexOffsets;
}

10
vtkm/cont/CellSetPermutation.h
View File

@ -92,8 +92,9 @@ public:
vtkm::Id connectivityLength = vtkm::cont::DeviceAdapterAlgorithm<Device>::ScanExclusive(
vtkm::cont::make_ArrayHandleCast(conn.NumIndices, vtkm::Id()), conn.IndexOffsets);
conn.Connectivity.Allocate(connectivityLength);
vtkm::worklet::DispatcherMapTopology<WriteConnectivity, Device>().Invoke(
cellset, conn.IndexOffsets, conn.Connectivity);
vtkm::worklet::DispatcherMapTopology<WriteConnectivity> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(cellset, conn.IndexOffsets, conn.Connectivity);
return conn;
}
@ -126,8 +127,9 @@ public:
conn.NumIndices = make_ArrayHandleConstant(numPointsInCell, numberOfCells);
conn.IndexOffsets = ArrayHandleCounting<vtkm::Id>(0, numPointsInCell, numberOfCells);
conn.Connectivity.Allocate(connectivityLength);
vtkm::worklet::DispatcherMapTopology<WriteConnectivity, Device>().Invoke(
cellset, conn.IndexOffsets, conn.Connectivity);
vtkm::worklet::DispatcherMapTopology<WriteConnectivity> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(cellset, conn.IndexOffsets, conn.Connectivity);
return conn;
}

6
vtkm/cont/ColorTable.hxx
View File

@ -57,7 +57,8 @@ struct map_color_table
{
using namespace vtkm::worklet::colorconversion;
TransferFunction transfer(colors->PrepareForExecution(device));
vtkm::worklet::DispatcherMapField<TransferFunction, DeviceAdapter> dispatcher(transfer);
vtkm::worklet::DispatcherMapField<TransferFunction> dispatcher(transfer);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(std::forward<Args>(args)...);
return true;
}
@ -71,7 +72,8 @@ struct map_color_table_with_samples
Args&&... args) const
{
using namespace vtkm::worklet::colorconversion;
vtkm::worklet::DispatcherMapField<LookupTable, DeviceAdapter> dispatcher(lookupTable);
vtkm::worklet::DispatcherMapField<LookupTable> dispatcher(lookupTable);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(std::forward<Args>(args)...);
return true;
}

4
vtkm/cont/DataSetBuilderRectilinear.h
View File

@ -20,6 +20,7 @@
#ifndef vtk_m_cont_DataSetBuilderRectilinear_h
#define vtk_m_cont_DataSetBuilderRectilinear_h
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/ArrayHandleCartesianProduct.h>
#include <vtkm/cont/ArrayPortalToIterators.h>
#include <vtkm/cont/CoordinateSystem.h>
@ -43,8 +44,7 @@ class VTKM_CONT_EXPORT DataSetBuilderRectilinear
VTKM_CONT static void CopyInto(const vtkm::cont::ArrayHandle<T>& input,
vtkm::cont::ArrayHandle<U>& output)
{
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<vtkm::cont::DeviceAdapterTagSerial>;
Algorithm::Copy(input, output);
vtkm::cont::ArrayCopy(input, output);
}
template <typename T, typename U>

2
vtkm/cont/ErrorBadDevice.cxx
View File

@ -32,7 +32,7 @@ void throwFailedRuntimeDeviceTransfer(const std::string& className,
vtkm::cont::DeviceAdapterId deviceId)
{ //Should we support typeid() instead of className?
const std::string msg = "VTK-m was unable to transfer " + className + " to DeviceAdapter[id=" +
std::to_string(deviceId.GetValue()) +
std::to_string(deviceId.GetValue()) + ", name=" + deviceId.GetName() +
"]. This is generally caused by asking for execution on a DeviceAdapter that "
"isn't compiled into VTK-m. In the case of CUDA it can also be caused by accidentally "
"compiling source files as C++ files instead of CUDA.";

2
vtkm/cont/ErrorBadValue.h
View File

@ -36,7 +36,7 @@ class VTKM_ALWAYS_EXPORT ErrorBadValue : public Error
{
public:
ErrorBadValue(const std::string& message)
: Error(message)
: Error(message, true)
{
}
};

2
vtkm/cont/ErrorExecution.h
View File

@ -36,7 +36,7 @@ class VTKM_ALWAYS_EXPORT ErrorExecution : public vtkm::cont::Error
{
public:
ErrorExecution(const std::string& message)
: Error(message)
: Error(message, true)
{
}
};

78
vtkm/cont/PointLocatorUniformGrid.h
View File

@ -20,10 +20,9 @@
#ifndef vtk_m_cont_PointLocatorUniformGrid_h
#define vtk_m_cont_PointLocatorUniformGrid_h
#include <vtkm/cont/Algorithm.h>
#include <vtkm/cont/ArrayCopy.h>
#include <vtkm/cont/ArrayHandleCounting.h>
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/cont/DeviceAdapterAlgorithm.h>
#include <vtkm/cont/TryExecute.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
@ -77,62 +76,29 @@ public:
vtkm::Vec<vtkm::FloatDefault, 3> Dxdydz;
};
/// \brief Construct a 3D uniform grid for nearest neighbor search.
///
/// \param coords An ArrayHandle of x, y, z coordinates of input points.
/// \param device Tag for selecting device adapter
struct BuildFunctor
{
BuildFunctor(vtkm::cont::PointLocatorUniformGrid* self)
: Self(self)
{
}
template <typename Device>
bool operator()(Device)
{
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<Device>;
// Save training data points.
Algorithm::Copy(this->Self->GetCoordinates().GetData(), this->Self->coords);
// generate unique id for each input point
vtkm::cont::ArrayHandleCounting<vtkm::Id> pointCounting(
0, 1, this->Self->coords.GetNumberOfValues());
Algorithm::Copy(pointCounting, this->Self->pointIds);
// bin points into cells and give each of them the cell id.
BinPointsWorklet cellIdWorklet(this->Self->Min, this->Self->Max, this->Self->Dims);
vtkm::worklet::DispatcherMapField<BinPointsWorklet, Device> dispatchCellId(cellIdWorklet);
dispatchCellId.Invoke(this->Self->coords, this->Self->cellIds);
// Group points of the same cell together by sorting them according to the cell ids
Algorithm::SortByKey(this->Self->cellIds, this->Self->pointIds);
// for each cell, find the lower and upper bound of indices to the sorted point ids.
vtkm::cont::ArrayHandleCounting<vtkm::Id> cell_ids_counting(
0, 1, this->Self->Dims[0] * this->Self->Dims[1] * this->Self->Dims[2]);
Algorithm::UpperBounds(this->Self->cellIds, cell_ids_counting, this->Self->cellUpper);
Algorithm::LowerBounds(this->Self->cellIds, cell_ids_counting, this->Self->cellLower);
return true;
}
private:
vtkm::cont::PointLocatorUniformGrid* Self;
};
void Build() override
{
BuildFunctor functor(this);
// Save training data points.
vtkm::cont::ArrayCopy(this->GetCoordinates().GetData(), this->coords);
bool success = vtkm::cont::TryExecute(functor);
if (!success)
{
throw vtkm::cont::ErrorExecution("Could not build point locator structure");
}
};
// generate unique id for each input point
vtkm::cont::ArrayHandleCounting<vtkm::Id> pointCounting(0, 1, this->coords.GetNumberOfValues());
vtkm::cont::ArrayCopy(pointCounting, this->pointIds);
// bin points into cells and give each of them the cell id.
BinPointsWorklet cellIdWorklet(this->Min, this->Max, this->Dims);
vtkm::worklet::DispatcherMapField<BinPointsWorklet> dispatchCellId(cellIdWorklet);
dispatchCellId.Invoke(this->coords, this->cellIds);
// Group points of the same cell together by sorting them according to the cell ids
vtkm::cont::Algorithm::SortByKey(this->cellIds, this->pointIds);
// for each cell, find the lower and upper bound of indices to the sorted point ids.
vtkm::cont::ArrayHandleCounting<vtkm::Id> cell_ids_counting(
0, 1, this->Dims[0] * this->Dims[1] * this->Dims[2]);
vtkm::cont::Algorithm::UpperBounds(this->cellIds, cell_ids_counting, this->cellUpper);
vtkm::cont::Algorithm::LowerBounds(this->cellIds, cell_ids_counting, this->cellLower);
}
using HandleType = vtkm::cont::VirtualObjectHandle<vtkm::exec::PointLocator>;

1
vtkm/cont/RuntimeDeviceInformation.h
View File

@ -25,6 +25,7 @@
//Bring in each device adapters runtime class
#include <vtkm/cont/cuda/internal/DeviceAdapterRuntimeDetectorCuda.h>
#include <vtkm/cont/internal/DeviceAdapterError.h>
#include <vtkm/cont/openmp/internal/DeviceAdapterRuntimeDetectorOpenMP.h>
#include <vtkm/cont/serial/internal/DeviceAdapterRuntimeDetectorSerial.h>
#include <vtkm/cont/tbb/internal/DeviceAdapterRuntimeDetectorTBB.h>

93
vtkm/cont/RuntimeDeviceTracker.cxx
View File

@ -34,6 +34,34 @@
#include <sstream>
#include <thread>
namespace
{
struct VTKM_NEVER_EXPORT GetDeviceNameFunctor
{
vtkm::cont::DeviceAdapterNameType* Names;
VTKM_CONT
GetDeviceNameFunctor(vtkm::cont::DeviceAdapterNameType* names)
: Names(names)
{
std::fill_n(this->Names, VTKM_MAX_DEVICE_ADAPTER_ID, "InvalidDeviceId");
}
template <typename Device>
VTKM_CONT void operator()(Device device)
{
auto id = device.GetValue();
if (id > 0 && id < VTKM_MAX_DEVICE_ADAPTER_ID)
{
this->Names[id] = vtkm::cont::DeviceAdapterTraits<Device>::GetName();
}
}
};
} // end anon namespace
namespace vtkm
{
namespace cont
@ -45,6 +73,7 @@ namespace detail
struct RuntimeDeviceTrackerInternals
{
bool RuntimeValid[VTKM_MAX_DEVICE_ADAPTER_ID];
DeviceAdapterNameType DeviceNames[VTKM_MAX_DEVICE_ADAPTER_ID];
};
}
@ -52,6 +81,9 @@ VTKM_CONT
RuntimeDeviceTracker::RuntimeDeviceTracker()
: Internals(new detail::RuntimeDeviceTrackerInternals)
{
GetDeviceNameFunctor functor(this->Internals->DeviceNames);
vtkm::ListForEach(functor, VTKM_DEFAULT_DEVICE_ADAPTER_LIST_TAG());
this->Reset();
}
@ -61,31 +93,28 @@ RuntimeDeviceTracker::~RuntimeDeviceTracker()
}
VTKM_CONT
void RuntimeDeviceTracker::CheckDevice(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName) const
void RuntimeDeviceTracker::CheckDevice(vtkm::cont::DeviceAdapterId deviceId) const
{
if (!deviceId.IsValueValid())
{
std::stringstream message;
message << "Device '" << deviceName << "' has invalid ID of " << (int)deviceId.GetValue();
message << "Device '" << deviceId.GetName() << "' has invalid ID of "
<< (int)deviceId.GetValue();
throw vtkm::cont::ErrorBadValue(message.str());
}
}
VTKM_CONT
bool RuntimeDeviceTracker::CanRunOnImpl(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName) const
bool RuntimeDeviceTracker::CanRunOnImpl(vtkm::cont::DeviceAdapterId deviceId) const
{
this->CheckDevice(deviceId, deviceName);
this->CheckDevice(deviceId);
return this->Internals->RuntimeValid[deviceId.GetValue()];
}
VTKM_CONT
void RuntimeDeviceTracker::SetDeviceState(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName,
bool state)
void RuntimeDeviceTracker::SetDeviceState(vtkm::cont::DeviceAdapterId deviceId, bool state)
{
this->CheckDevice(deviceId, deviceName);
this->CheckDevice(deviceId);
this->Internals->RuntimeValid[deviceId.GetValue()] = state;
}
@ -135,18 +164,16 @@ void RuntimeDeviceTracker::DeepCopy(const vtkm::cont::RuntimeDeviceTracker& src)
}
VTKM_CONT
void RuntimeDeviceTracker::ForceDeviceImpl(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName,
bool runtimeExists)
void RuntimeDeviceTracker::ForceDeviceImpl(vtkm::cont::DeviceAdapterId deviceId, bool runtimeExists)
{
if (!runtimeExists)
{
std::stringstream message;
message << "Cannot force to device '" << deviceName
message << "Cannot force to device '" << deviceId.GetName()
<< "' because that device is not available on this system";
throw vtkm::cont::ErrorBadValue(message.str());
}
this->CheckDevice(deviceId, deviceName);
this->CheckDevice(deviceId);
std::fill_n(this->Internals->RuntimeValid, VTKM_MAX_DEVICE_ADAPTER_ID, false);
@ -177,5 +204,41 @@ vtkm::cont::RuntimeDeviceTracker GetGlobalRuntimeDeviceTracker()
return runtimeDeviceTracker;
#endif
}
VTKM_CONT
DeviceAdapterNameType RuntimeDeviceTracker::GetDeviceName(DeviceAdapterId device) const
{
auto id = device.GetValue();
if (id < 0)
{
switch (id)
{
case VTKM_DEVICE_ADAPTER_ERROR:
return vtkm::cont::DeviceAdapterTraits<vtkm::cont::DeviceAdapterTagError>::GetName();
case VTKM_DEVICE_ADAPTER_UNDEFINED:
return vtkm::cont::DeviceAdapterTraits<vtkm::cont::DeviceAdapterTagUndefined>::GetName();
default:
break;
}
}
else if (id >= VTKM_MAX_DEVICE_ADAPTER_ID)
{
switch (id)
{
case VTKM_DEVICE_ADAPTER_ANY:
return vtkm::cont::DeviceAdapterTraits<vtkm::cont::DeviceAdapterTagAny>::GetName();
default:
break;
}
}
else // id is valid:
{
return this->Internals->DeviceNames[id];
}
// Device 0 is invalid:
return this->Internals->DeviceNames[0];
}
}
} // namespace vtkm::cont

42
vtkm/cont/RuntimeDeviceTracker.h
View File

@ -64,8 +64,7 @@ public:
template <typename DeviceAdapterTag>
VTKM_CONT bool CanRunOn(DeviceAdapterTag device) const
{
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag>;
return this->CanRunOnImpl(device, Traits::GetName());
return this->CanRunOnImpl(device);
}
/// Report a failure to allocate memory on a device, this will flag the
@ -76,8 +75,7 @@ public:
VTKM_CONT void ReportAllocationFailure(DeviceAdapterTag device,
const vtkm::cont::ErrorBadAllocation&)
{
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag>;
this->SetDeviceState(device, Traits::GetName(), false);
this->SetDeviceState(device, false);
}
/// Report a failure to allocate memory on a device, this will flag the
@ -85,10 +83,9 @@ public:
/// the filter.
///
VTKM_CONT void ReportAllocationFailure(vtkm::cont::DeviceAdapterId deviceId,
const std::string& name,
const vtkm::cont::ErrorBadAllocation&)
{
this->SetDeviceState(deviceId, name, false);
this->SetDeviceState(deviceId, false);
}
//@{
@ -96,15 +93,13 @@ public:
template <typename DeviceAdapterTag>
VTKM_CONT void ReportBadDeviceFailure(DeviceAdapterTag device, const vtkm::cont::ErrorBadDevice&)
{
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag>;
this->SetDeviceState(device, Traits::GetName(), false);
this->SetDeviceState(device, false);
}
VTKM_CONT void ReportBadDeviceFailure(vtkm::cont::DeviceAdapterId deviceId,
const std::string& name,
const vtkm::cont::ErrorBadDevice&)
{
this->SetDeviceState(deviceId, name, false);
this->SetDeviceState(deviceId, false);
}
//@}
@ -114,9 +109,8 @@ public:
template <typename DeviceAdapterTag>
VTKM_CONT void ResetDevice(DeviceAdapterTag device)
{
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag>;
vtkm::cont::RuntimeDeviceInformation<DeviceAdapterTag> runtimeDevice;
this->SetDeviceState(device, Traits::GetName(), runtimeDevice.Exists());
this->SetDeviceState(device, runtimeDevice.Exists());
}
/// Reset the tracker to its default state for default devices.
@ -180,8 +174,7 @@ public:
template <typename DeviceAdapterTag>
VTKM_CONT void DisableDevice(DeviceAdapterTag device)
{
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag>;
this->SetDeviceState(device, Traits::GetName(), false);
this->SetDeviceState(device, false);
}
/// \brief Disable all devices except the specified one.
@ -199,35 +192,32 @@ public:
template <typename DeviceAdapterTag>
VTKM_CONT void ForceDevice(DeviceAdapterTag device)
{
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag>;
vtkm::cont::RuntimeDeviceInformation<DeviceAdapterTag> runtimeDevice;
this->ForceDeviceImpl(device, Traits::GetName(), runtimeDevice.Exists());
this->ForceDeviceImpl(device, runtimeDevice.Exists());
}
VTKM_CONT_EXPORT
VTKM_CONT
DeviceAdapterNameType GetDeviceName(DeviceAdapterId id) const;
private:
std::shared_ptr<detail::RuntimeDeviceTrackerInternals> Internals;
VTKM_CONT_EXPORT
VTKM_CONT
void CheckDevice(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName) const;
void CheckDevice(vtkm::cont::DeviceAdapterId deviceId) const;
VTKM_CONT_EXPORT
VTKM_CONT
bool CanRunOnImpl(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName) const;
bool CanRunOnImpl(vtkm::cont::DeviceAdapterId deviceId) const;
VTKM_CONT_EXPORT
VTKM_CONT
void SetDeviceState(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName,
bool state);
void SetDeviceState(vtkm::cont::DeviceAdapterId deviceId, bool state);
VTKM_CONT_EXPORT
VTKM_CONT
void ForceDeviceImpl(vtkm::cont::DeviceAdapterId deviceId,
const vtkm::cont::DeviceAdapterNameType& deviceName,
bool runtimeExists);
void ForceDeviceImpl(vtkm::cont::DeviceAdapterId deviceId, bool runtimeExists);
};
/// \brief Get the \c RuntimeDeviceTracker for the current thread.

13
vtkm/cont/TryExecute.cxx
View File

@ -32,7 +32,6 @@ namespace detail
{
void HandleTryExecuteException(vtkm::cont::DeviceAdapterId deviceId,
const std::string& name,
vtkm::cont::RuntimeDeviceTracker& tracker)
{
try
@ -45,12 +44,12 @@ void HandleTryExecuteException(vtkm::cont::DeviceAdapterId deviceId,
std::cerr << "caught ErrorBadAllocation " << e.GetMessage() << std::endl;
//currently we only consider OOM errors worth disabling a device for
//than we fallback to another device
tracker.ReportAllocationFailure(deviceId, name, e);
tracker.ReportAllocationFailure(deviceId, e);
}
catch (vtkm::cont::ErrorBadDevice& e)
{
std::cerr << "caught ErrorBadDevice: " << e.GetMessage() << std::endl;
tracker.ReportBadDeviceFailure(deviceId, name, e);
tracker.ReportBadDeviceFailure(deviceId, e);
}
catch (vtkm::cont::ErrorBadType& e)
{
@ -58,11 +57,11 @@ void HandleTryExecuteException(vtkm::cont::DeviceAdapterId deviceId,
//deferring to another device adapter?
std::cerr << "caught ErrorBadType : " << e.GetMessage() << std::endl;
}
catch (vtkm::cont::ErrorBadValue& e)
catch (vtkm::cont::ErrorBadValue&)
{
//should bad value errors should stop the filter, instead of deferring
//to another device adapter?
std::cerr << "caught ErrorBadValue : " << e.GetMessage() << std::endl;
// Should bad values be deferred to another device? Seems unlikely they will succeed.
// Re-throw instead.
throw;
}
catch (vtkm::cont::Error& e)
{

8
vtkm/cont/TryExecute.h
View File

@ -33,7 +33,6 @@ namespace detail
{
VTKM_CONT_EXPORT void HandleTryExecuteException(vtkm::cont::DeviceAdapterId,
const std::string&,
vtkm::cont::RuntimeDeviceTracker&);
template <typename DeviceTag, typename Functor, typename... Args>
@ -44,7 +43,7 @@ inline bool TryExecuteIfValid(std::true_type,
vtkm::cont::RuntimeDeviceTracker& tracker,
Args&&... args)
{
if ((tag == devId || devId == DeviceAdapterIdAny()) && tracker.CanRunOn(tag))
if ((tag == devId || devId == DeviceAdapterTagAny()) && tracker.CanRunOn(tag))
{
try
{
@ -52,8 +51,7 @@ inline bool TryExecuteIfValid(std::true_type,
}
catch (...)
{
using Traits = vtkm::cont::DeviceAdapterTraits<DeviceTag>;
detail::HandleTryExecuteException(tag, Traits::GetName(), tracker);
detail::HandleTryExecuteException(tag, tracker);
}
}
@ -344,7 +342,7 @@ template <typename Functor, typename... Args>
VTKM_CONT bool TryExecute(Functor&& functor, Args&&... args)
{
return TryExecuteOnDevice(
vtkm::cont::DeviceAdapterIdAny(), std::forward<Functor>(functor), std::forward<Args>(args)...);
vtkm::cont::DeviceAdapterTagAny(), std::forward<Functor>(functor), std::forward<Args>(args)...);
}

6
vtkm/cont/VirtualObjectHandle.h
View File

@ -131,9 +131,9 @@ public:
if (!this->Internals->Transfers[static_cast<std::size_t>(deviceId.GetValue())])
{
std::string msg =
"VTK-m was asked to transfer a VirtualObjectHandle for execution on DeviceAdapter " +
std::to_string(deviceId.GetValue()) +
". It can't as this VirtualObjectHandle was not constructed/bound with this "
"VTK-m was asked to transfer a VirtualObjectHandle for execution on DeviceAdapter '" +
deviceId.GetName() + "' (" + std::to_string(deviceId.GetValue()) +
"). It can't as this VirtualObjectHandle was not constructed/bound with this "
"DeviceAdapter in the list of valid DeviceAdapters.";
throw vtkm::cont::ErrorBadDevice(msg);
}

9
vtkm/cont/cuda/internal/DeviceAdapterAtomicArrayImplementationCuda.h
View File

@ -65,6 +65,10 @@ public:
T* lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
return this->vtkmAtomicAdd(lockedValue, value);
#else
// Shut up, compiler
(void)index;
(void)value;
throw vtkm::cont::ErrorExecution("AtomicArray used in control environment, "
"or incorrect array implementation used "
"for device.");
@ -85,6 +89,11 @@ public:
T* lockedValue = ::thrust::raw_pointer_cast(this->Portal.GetIteratorBegin() + index);
return this->vtkmCompareAndSwap(lockedValue, newValue, oldValue);
#else
// Shut up, compiler.
(void)index;
(void)newValue;
(void)oldValue;
throw vtkm::cont::ErrorExecution("AtomicArray used in control environment, "
"or incorrect array implementation used "
"for device.");

5
vtkm/cont/cuda/testing/UnitTestCudaMath.cu
View File

@ -92,7 +92,7 @@ struct TriggerICE : public vtkm::worklet::WorkletMapField
template <class ValueType>
ValueType operator()(const ValueType& bad, const ValueType& sane, const vtkm::Id sequenceId) const
{
return bad + sane * sequenceId;
return bad + sane * static_cast<ValueType>(sequenceId);
}
#endif
};
@ -156,7 +156,8 @@ void RunEdgeCases()
auto bad = vtkm::cont::make_ArrayHandle(badvalues);
auto sane = vtkm::cont::make_ArrayHandle(sanevalues);
decltype(sane) result;
vtkm::worklet::DispatcherMapField<TriggerICE, Device> dispatcher;
vtkm::worklet::DispatcherMapField<TriggerICE> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(bad, sane, result);
auto portal = result.GetPortalConstControl();

2
vtkm/cont/internal/ArrayHandleBasicImpl.cxx
View File

@ -250,7 +250,7 @@ bool ArrayHandleImpl::PrepareForDevice(DeviceAdapterId devId, vtkm::UInt64 sizeO
DeviceAdapterId ArrayHandleImpl::GetDeviceAdapterId() const
{
return this->ExecutionArrayValid ? this->ExecutionInterface->GetDeviceId()
: DeviceAdapterIdUndefined{};
: DeviceAdapterTagUndefined{};
}

35
vtkm/cont/internal/DeviceAdapterTag.cxx Normal file
View File

@ -0,0 +1,35 @@
//============================================================================
// 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.
//
// Copyright 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2018 UT-Battelle, LLC.
// Copyright 2018 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/internal/DeviceAdapterTag.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
namespace vtkm
{
namespace cont
{
DeviceAdapterNameType DeviceAdapterId::GetName() const
{
return vtkm::cont::GetGlobalRuntimeDeviceTracker().GetDeviceName(*this);
}
}
} // end namespace vtkm::cont

35
vtkm/cont/internal/DeviceAdapterTag.h
View File

@ -25,6 +25,8 @@
#include <vtkm/internal/Configure.h>
#include <vtkm/internal/ExportMacros.h>
#include <vtkm/cont/vtkm_cont_export.h>
#include <string>
#define VTKM_DEVICE_ADAPTER_ERROR -2
@ -42,6 +44,8 @@ namespace vtkm
namespace cont
{
using DeviceAdapterNameType = std::string;
struct DeviceAdapterId
{
constexpr bool operator==(DeviceAdapterId other) const { return this->Value == other.Value; }
@ -55,6 +59,9 @@ struct DeviceAdapterId
constexpr vtkm::Int8 GetValue() const { return this->Value; }
VTKM_CONT_EXPORT
DeviceAdapterNameType GetName() const;
protected:
constexpr explicit DeviceAdapterId(vtkm::Int8 id)
: Value(id)
@ -65,27 +72,6 @@ private:
vtkm::Int8 Value;
};
// Represents when using TryExecute that the functor
// can be executed on any device instead of a specific
// one
struct DeviceAdapterIdAny : DeviceAdapterId
{
constexpr DeviceAdapterIdAny()
: DeviceAdapterId(127)
{
}
};
struct DeviceAdapterIdUndefined : DeviceAdapterId
{
constexpr DeviceAdapterIdUndefined()
: DeviceAdapterId(VTKM_DEVICE_ADAPTER_UNDEFINED)
{
}
};
using DeviceAdapterNameType = std::string;
template <typename DeviceAdapter>
struct DeviceAdapterTraits;
}
@ -139,6 +125,13 @@ struct DeviceAdapterTraits;
} \
}
// Represents when using TryExecute that the functor
// can be executed on any device instead of a specific
// one
VTKM_VALID_DEVICE_ADAPTER(Any, VTKM_DEVICE_ADAPTER_ANY)
VTKM_INVALID_DEVICE_ADAPTER(Undefined, VTKM_DEVICE_ADAPTER_UNDEFINED)
/// Checks that the argument is a proper device adapter tag. This is a handy
/// concept check for functions and classes to make sure that a template
/// argument is actually a device adapter tag. (You can get weird errors

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

@ -73,6 +73,7 @@ set(unit_tests
UnitTestFieldRangeCompute.cxx
UnitTestMultiBlock.cxx
UnitTestRuntimeDeviceInformation.cxx
UnitTestRuntimeDeviceNames.cxx
UnitTestStorageBasic.cxx
UnitTestStorageImplicit.cxx
UnitTestStorageListTag.cxx

22
vtkm/cont/testing/TestingArrayHandleVirtualCoordinates.h
View File

@ -63,12 +63,11 @@ struct DoubleWorklet : public vtkm::worklet::WorkletMapField
}
};
template <typename T, typename S, typename DeviceAdapter>
template <typename T, typename S>
inline void TestVirtualAccess(const vtkm::cont::ArrayHandle<T, S>& in,
vtkm::cont::ArrayHandle<T>& out,
DeviceAdapter)
vtkm::cont::ArrayHandle<T>& out)
{
vtkm::worklet::DispatcherMapField<CopyWorklet, DeviceAdapter>().Invoke(
vtkm::worklet::DispatcherMapField<CopyWorklet>().Invoke(
vtkm::cont::ArrayHandleVirtualCoordinates(in), vtkm::cont::ArrayHandleVirtualCoordinates(out));
VTKM_TEST_ASSERT(test_equal_portals(in.GetPortalConstControl(), out.GetPortalConstControl()),
@ -103,11 +102,11 @@ private:
{
a1.GetPortalControl().Set(i, TestValue(i, PointType()));
}
TestVirtualAccess(a1, out, DeviceAdapter{});
TestVirtualAccess(a1, out);
std::cout << "Testing ArrayHandleUniformPointCoordinates as input\n";
auto t = vtkm::cont::ArrayHandleUniformPointCoordinates(vtkm::Id3(4, 4, 4));
TestVirtualAccess(t, out, DeviceAdapter{});
TestVirtualAccess(t, out);
std::cout << "Testing ArrayHandleCartesianProduct as input\n";
vtkm::cont::ArrayHandle<vtkm::FloatDefault> c1, c2, c3;
@ -121,13 +120,12 @@ private:
c2.GetPortalControl().Set(i, p[1]);
c3.GetPortalControl().Set(i, p[2]);
}
TestVirtualAccess(
vtkm::cont::make_ArrayHandleCartesianProduct(c1, c2, c3), out, DeviceAdapter{});
TestVirtualAccess(vtkm::cont::make_ArrayHandleCartesianProduct(c1, c2, c3), out);
std::cout << "Testing resources releasing on ArrayHandleVirtualCoordinates\n";
vtkm::cont::ArrayHandleVirtualCoordinates virtualC =
vtkm::cont::ArrayHandleVirtualCoordinates(a1);
vtkm::worklet::DispatcherMapField<DoubleWorklet, DeviceAdapter>().Invoke(a1);
vtkm::worklet::DispatcherMapField<DoubleWorklet>().Invoke(a1);
virtualC.ReleaseResourcesExecution();
VTKM_TEST_ASSERT(a1.GetNumberOfValues() == length,
"ReleaseResourcesExecution"
@ -143,7 +141,11 @@ private:
}
public:
static int Run() { return vtkm::cont::testing::Testing::Run(TestAll); }
static int Run()
{
vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(DeviceAdapter());
return vtkm::cont::testing::Testing::Run(TestAll);
}
};
}
}

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

@ -22,6 +22,7 @@
#include <vtkm/TypeTraits.h>
#include <vtkm/cont/ArrayHandle.h>
#include <vtkm/cont/RuntimeDeviceTracker.h>
#include <vtkm/worklet/DispatcherMapField.h>
#include <vtkm/worklet/WorkletMapField.h>
@ -135,7 +136,7 @@ private:
using Algorithm = vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapterTag>;
using DispatcherPassThrough = vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag>;
using DispatcherPassThrough = vtkm::worklet::DispatcherMapField<PassThrough>;
struct VerifyEmptyArrays
{
template <typename T>
@ -455,7 +456,11 @@ private:
};
public:
static VTKM_CONT int Run() { return vtkm::cont::testing::Testing::Run(TryArrayHandleType()); }
static VTKM_CONT int Run()
{
vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(DeviceAdapterTag());
return vtkm::cont::testing::Testing::Run(TryArrayHandleType());
}
};
}
}

5
vtkm/cont/testing/TestingCellLocatorTwoLevelUniformGrid.h
View File

@ -118,7 +118,7 @@ vtkm::cont::DataSet MakeTestDataSet(const vtkm::Vec<vtkm::Id, DIMENSIONS>& dims,
cellset = vtkm::worklet::Triangulate().Run(uniformCs, device);
break;
case 3:
cellset = vtkm::worklet::Tetrahedralize().Run(uniformCs, device);
cellset = vtkm::worklet::Tetrahedralize().Run(uniformCs);
break;
default:
VTKM_ASSERT(false);
@ -181,8 +181,9 @@ void GenerateRandomInput(const vtkm::cont::DataSet& ds,
pcoords.GetPortalControl().Set(i, pc);
}
vtkm::worklet::DispatcherMapTopology<ParametricToWorldCoordinates, DeviceAdapter> dispatcher(
vtkm::worklet::DispatcherMapTopology<ParametricToWorldCoordinates> dispatcher(
ParametricToWorldCoordinates::MakeScatter(cellIds));
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(ds.GetCellSet(), ds.GetCoordinateSystem().GetData(), pcoords, wcoords);
}

5
vtkm/cont/testing/TestingCellLocatorUniformGrid.h
View File

@ -134,8 +134,9 @@ public:
vtkm::cont::ArrayHandle<PointType> parametric;
vtkm::cont::ArrayHandle<bool> match;
LocatorWorklet worklet(bounds, dims);
vtkm::worklet::DispatcherMapField<LocatorWorklet, DeviceAdapter> dispather(worklet);
dispather.Invoke(points, locator, cellIds, parametric, match);
vtkm::worklet::DispatcherMapField<LocatorWorklet> dispatcher(worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(points, locator, cellIds, parametric, match);
auto matchPortal = match.GetPortalConstControl();
for (vtkm::Id index = 0; index < match.GetNumberOfValues(); index++)

3
vtkm/cont/testing/TestingColorTable.h
View File

@ -461,7 +461,8 @@ public:
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::UInt8, 4>> colors;
TransferFunction transfer(table.PrepareForExecution(DeviceAdapterTag{}));
vtkm::worklet::DispatcherMapField<TransferFunction, DeviceAdapterTag> dispatcher(transfer);
vtkm::worklet::DispatcherMapField<TransferFunction> dispatcher(transfer);
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(samples, colors);
const vtkm::Vec<vtkm::UInt8, 4> correct_sampling_points[nvals] = { { 14, 28, 31, 255 },

3
vtkm/cont/testing/TestingDataSetSingleType.h
View File

@ -139,7 +139,8 @@ private:
//run a basic for-each topology algorithm on this
vtkm::cont::ArrayHandle<vtkm::Float32> result;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellAverage, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellAverage> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(cellset, dataSet.GetField("pointvar"), result);
vtkm::Float32 expected[3] = { 20.1333f, 30.1667f, 40.2333f };

140
vtkm/cont/testing/TestingFancyArrayHandles.h
View File

@ -36,6 +36,7 @@
#include <vtkm/cont/ArrayHandleIndex.h>
#include <vtkm/cont/ArrayHandlePermutation.h>
#include <vtkm/cont/ArrayHandleTransform.h>
#include <vtkm/cont/ArrayHandleView.h>
#include <vtkm/cont/ArrayHandleZip.h>
#include <vtkm/worklet/DispatcherMapField.h>
@ -167,7 +168,7 @@ private:
vtkm::cont::ArrayHandle<vtkm::Vec<ValueType, 3>> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(composite, result);
//verify that the control portal works
@ -195,7 +196,7 @@ private:
vtkm::cont::make_ArrayHandleConstant(value, ARRAY_SIZE);
vtkm::cont::ArrayHandle<ValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(constant, result);
vtkm::cont::printSummary_ArrayHandle(constant, std::cout);
@ -230,7 +231,7 @@ private:
vtkm::cont::make_ArrayHandleCounting(start, ValueType(1), length);
vtkm::cont::ArrayHandle<ValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(counting, result);
vtkm::cont::printSummary_ArrayHandle(counting, std::cout);
@ -266,7 +267,7 @@ private:
vtkm::cont::ArrayHandle<ValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(implicit, result);
//verify that the control portal works
@ -319,7 +320,7 @@ private:
vtkm::cont::ArrayHandle<ValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(concatenate, result);
//verify that the control portal works
@ -373,7 +374,7 @@ private:
vtkm::cont::ArrayHandle<ValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(permutation, result);
//verify that the control portal works
@ -392,6 +393,52 @@ private:
}
};
struct TestViewAsInput
{
template <typename ValueType>
VTKM_CONT void operator()(const ValueType vtkmNotUsed(v)) const
{
const vtkm::Id length = ARRAY_SIZE;
using FunctorType = ::fancy_array_detail::IndexSquared<ValueType>;
using ValueHandleType = vtkm::cont::ArrayHandleImplicit<FunctorType>;
using ViewHandleType = vtkm::cont::ArrayHandleView<ValueHandleType>;
FunctorType functor;
for (vtkm::Id start_pos = 0; start_pos < length; start_pos += length / 4)
{
const vtkm::Id counting_length = length - start_pos;
ValueHandleType implicit = vtkm::cont::make_ArrayHandleImplicit(functor, length);
ViewHandleType view =
vtkm::cont::make_ArrayHandleView(implicit, start_pos, counting_length);
vtkm::cont::printSummary_ArrayHandle(view, std::cout);
std::cout << std::endl;
vtkm::cont::ArrayHandle<ValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(view, result);
//verify that the control portal works
for (vtkm::Id i = 0; i < counting_length; ++i)
{
const vtkm::Id value_index = i;
const vtkm::Id key_index = start_pos + i;
const ValueType result_v = result.GetPortalConstControl().Get(value_index);
const ValueType correct_value = implicit.GetPortalConstControl().Get(key_index);
const ValueType control_value = view.GetPortalConstControl().Get(value_index);
VTKM_TEST_ASSERT(test_equal(result_v, correct_value), "Implicit Handle Failed");
VTKM_TEST_ASSERT(test_equal(result_v, control_value), "Implicit Handle Failed");
}
}
}
};
struct TestTransformAsInput
{
template <typename ValueType>
@ -414,7 +461,7 @@ private:
vtkm::cont::ArrayHandle<ValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(transformed, result);
//verify that the control portal works
@ -456,7 +503,7 @@ private:
vtkm::cont::ArrayHandle<OutputValueType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(countingTransformed, result);
//verify that the control portal works
@ -485,7 +532,7 @@ private:
vtkm::cont::make_ArrayHandleCast(input, CastToType());
vtkm::cont::ArrayHandle<CastToType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(castArray, result);
vtkm::cont::printSummary_ArrayHandle(castArray, std::cout);
@ -529,7 +576,7 @@ private:
vtkm::cont::ArrayHandle<ValueType> resultArray;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(groupArray, resultArray);
VTKM_TEST_ASSERT(resultArray.GetNumberOfValues() == ARRAY_SIZE, "Got bad result array size.");
@ -568,7 +615,7 @@ private:
vtkm::cont::ArrayHandleGroupVec<vtkm::cont::ArrayHandle<ComponentType>, NUM_COMPONENTS>
groupArray(resultArray);
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(baseArray, groupArray);
vtkm::cont::printSummary_ArrayHandle(groupArray, std::cout);
@ -649,7 +696,7 @@ private:
vtkm::cont::make_ArrayHandleGroupVecVariable(sourceArray, offsetsArray), std::cout);
std::cout << std::endl;
vtkm::worklet::DispatcherMapField<GroupVariableInputWorklet, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<GroupVariableInputWorklet> dispatcher;
dispatcher.Invoke(vtkm::cont::make_ArrayHandleGroupVecVariable(sourceArray, offsetsArray));
}
};
@ -695,7 +742,7 @@ private:
vtkm::cont::ArrayHandle<ComponentType> sourceArray;
sourceArray.Allocate(sourceArraySize);
vtkm::worklet::DispatcherMapField<GroupVariableOutputWorklet, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<GroupVariableOutputWorklet> dispatcher;
dispatcher.Invoke(vtkm::cont::ArrayHandleIndex(ARRAY_SIZE),
vtkm::cont::make_ArrayHandleGroupVecVariable(sourceArray, offsetsArray));
@ -739,7 +786,7 @@ private:
vtkm::cont::ArrayHandle<PairType> result;
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(zip, result);
//verify that the control portal works
@ -776,7 +823,7 @@ private:
DiscardHandleType discard;
discard.Allocate(length);
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(input, discard);
// No output to verify since none is stored in memory. Just checking that
@ -812,7 +859,7 @@ private:
KeyHandleType counting = vtkm::cont::make_ArrayHandleCounting<vtkm::Id>(length, 1, length);
PermutationHandleType permutation = vtkm::cont::make_ArrayHandlePermutation(counting, values);
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(input, permutation);
vtkm::cont::printSummary_ArrayHandle(permutation, std::cout);
@ -829,6 +876,47 @@ private:
}
};
struct TestViewAsOutput
{
template <typename ValueType>
VTKM_CONT void operator()(const ValueType vtkmNotUsed(v)) const
{
const vtkm::Id length = ARRAY_SIZE;
using ValueHandleType = vtkm::cont::ArrayHandle<ValueType>;
using ViewHandleType = vtkm::cont::ArrayHandleView<ValueHandleType>;
using ComponentType = typename vtkm::VecTraits<ValueType>::ComponentType;
vtkm::cont::ArrayHandle<ValueType> input;
using Portal = typename vtkm::cont::ArrayHandle<ValueType>::PortalControl;
input.Allocate(length);
Portal inputPortal = input.GetPortalControl();
for (vtkm::Id i = 0; i < length; ++i)
{
inputPortal.Set(i, ValueType(ComponentType(i)));
}
ValueHandleType values;
values.Allocate(length * 2);
ViewHandleType view = vtkm::cont::make_ArrayHandleView(values, length, length);
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(input, view);
vtkm::cont::printSummary_ArrayHandle(view, std::cout);
std::cout << std::endl;
//verify that the control portal works
for (vtkm::Id i = 0; i < length; ++i)
{
const ValueType result_v = view.GetPortalConstControl().Get(i);
const ValueType correct_value = ValueType(ComponentType(i));
VTKM_TEST_ASSERT(test_equal(result_v, correct_value),
"Permutation Handle Failed As Output");
}
}
};
struct TestZipAsOutput
{
template <typename KeyType, typename ValueType>
@ -853,7 +941,7 @@ private:
vtkm::cont::ArrayHandle<ValueType>>
result_zip = vtkm::cont::make_ArrayHandleZip(result_keys, result_values);
vtkm::worklet::DispatcherMapField<PassThrough, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<PassThrough> dispatcher;
dispatcher.Invoke(input, result_zip);
vtkm::cont::printSummary_ArrayHandle(result_zip, std::cout);
@ -886,7 +974,7 @@ private:
vtkm::cont::ArrayHandle<ValueType> outputValues;
outputValues.Allocate(ARRAY_SIZE);
vtkm::worklet::DispatcherMapField<InplaceFunctorPair, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<InplaceFunctorPair> dispatcher;
dispatcher.Invoke(vtkm::cont::make_ArrayHandleZip(inputValues, outputValues));
vtkm::cont::printSummary_ArrayHandle(outputValues, std::cout);
@ -949,6 +1037,11 @@ private:
vtkm::testing::Testing::TryTypes(
TestingFancyArrayHandles<DeviceAdapterTag>::TestPermutationAsInput(), HandleTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleView as Input" << std::endl;
vtkm::testing::Testing::TryTypes(
TestingFancyArrayHandles<DeviceAdapterTag>::TestViewAsInput(), HandleTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleTransform as Input" << std::endl;
vtkm::testing::Testing::TryTypes(
@ -1007,6 +1100,11 @@ private:
vtkm::testing::Testing::TryTypes(
TestingFancyArrayHandles<DeviceAdapterTag>::TestPermutationAsOutput(), HandleTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleView as Output" << std::endl;
vtkm::testing::Testing::TryTypes(
TestingFancyArrayHandles<DeviceAdapterTag>::TestViewAsOutput(), HandleTypesToTest());
std::cout << "-------------------------------------------" << std::endl;
std::cout << "Testing ArrayHandleDiscard as Output" << std::endl;
vtkm::testing::Testing::TryTypes(
@ -1034,7 +1132,11 @@ public:
/// all the fancy array handles that vtkm supports. Returns an
/// error code that can be returned from the main function of a test.
///
static VTKM_CONT int Run() { return vtkm::cont::testing::Testing::Run(TestAll()); }
static VTKM_CONT int Run()
{
vtkm::cont::GetGlobalRuntimeDeviceTracker().ForceDevice(DeviceAdapterTag());
return vtkm::cont::testing::Testing::Run(TestAll());
}
};
}
}

6
vtkm/cont/testing/TestingImplicitFunction.h
View File

@ -71,10 +71,12 @@ void EvaluateOnCoordinates(vtkm::cont::CoordinateSystem points,
vtkm::cont::ArrayHandle<vtkm::FloatDefault>& values,
DeviceAdapter device)
{
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvaluateImplicitFunction, DeviceAdapter>;
using EvalDispatcher = vtkm::worklet::DispatcherMapField<EvaluateImplicitFunction>;
EvaluateImplicitFunction eval(function.PrepareForExecution(device));
EvalDispatcher(eval).Invoke(points, values);
EvalDispatcher dispatcher(eval);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(points, values);
}
template <std::size_t N>

10
vtkm/cont/testing/TestingPointLocatorUniformGrid.h
View File

@ -144,16 +144,18 @@ public:
vtkm::cont::ArrayHandle<vtkm::FloatDefault> nnDis_Handle;
PointLocatorUniformGridWorklet pointLocatorUniformGridWorklet;
vtkm::worklet::DispatcherMapField<PointLocatorUniformGridWorklet, DeviceAdapter>
locatorDispatcher(pointLocatorUniformGridWorklet);
vtkm::worklet::DispatcherMapField<PointLocatorUniformGridWorklet> locatorDispatcher(
pointLocatorUniformGridWorklet);
locatorDispatcher.SetDevice(DeviceAdapter());
locatorDispatcher.Invoke(qc_Handle, locator, nnId_Handle, nnDis_Handle);
// brute force
vtkm::cont::ArrayHandle<vtkm::Id> bfnnId_Handle;
vtkm::cont::ArrayHandle<vtkm::Float32> bfnnDis_Handle;
NearestNeighborSearchBruteForce3DWorklet nnsbf3dWorklet;
vtkm::worklet::DispatcherMapField<NearestNeighborSearchBruteForce3DWorklet, DeviceAdapter>
nnsbf3DDispatcher(nnsbf3dWorklet);
vtkm::worklet::DispatcherMapField<NearestNeighborSearchBruteForce3DWorklet> nnsbf3DDispatcher(
nnsbf3dWorklet);
nnsbf3DDispatcher.SetDevice(DeviceAdapter());
nnsbf3DDispatcher.Invoke(
qc_Handle, vtkm::cont::make_ArrayHandle(coordi), bfnnId_Handle, bfnnDis_Handle);

88
vtkm/cont/testing/UnitTestRuntimeDeviceNames.cxx Normal file
View File

@ -0,0 +1,88 @@
//============================================================================
// 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.
//
// Copyright 2018 National Technology & Engineering Solutions of Sandia, LLC (NTESS).
// Copyright 2018 UT-Battelle, LLC.
// Copyright 2018 Los Alamos National Security.
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
// Laboratory (LANL), the U.S. Government retains certain rights in
// this software.
//============================================================================
#include <vtkm/cont/RuntimeDeviceTracker.h>
#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
#include <vtkm/cont/internal/DeviceAdapterError.h>
#include <vtkm/cont/openmp/DeviceAdapterOpenMP.h>
#include <vtkm/cont/serial/DeviceAdapterSerial.h>
#include <vtkm/cont/tbb/DeviceAdapterTBB.h>
#include <vtkm/cont/testing/Testing.h>
namespace
{
// Invalid tag for testing. Returns the default "InvalidDeviceId" from
// vtkm::cont::RuntimeDeviceTracker::GetName.
struct VTKM_ALWAYS_EXPORT DeviceAdapterTagInvalidDeviceId : vtkm::cont::DeviceAdapterId
{
constexpr DeviceAdapterTagInvalidDeviceId()
: DeviceAdapterId(VTKM_MAX_DEVICE_ADAPTER_ID)
{
}
};
template <typename Tag>
void TestName(const std::string& name, Tag tag, vtkm::cont::DeviceAdapterId id)
{
auto tracker = vtkm::cont::GetGlobalRuntimeDeviceTracker();
#if 0
std::cerr << "Expected: " << name << "\n"
<< "\t" << id.GetName() << "\n"
<< "\t" << tag.GetName() << "\n"
<< "\t" << tracker.GetDeviceName(id) << "\n"
<< "\t" << tracker.GetDeviceName(tag) << "\n";
#endif
VTKM_TEST_ASSERT(id.GetName() == name, "Id::GetName() failed.");
VTKM_TEST_ASSERT(tag.GetName() == name, "Tag::GetName() failed.");
VTKM_TEST_ASSERT(tracker.GetDeviceName(id) == name, "RTDeviceTracker::GetDeviceName(Id) failed.");
VTKM_TEST_ASSERT(tracker.GetDeviceName(tag) == name,
"RTDeviceTracker::GetDeviceName(Tag) failed.");
}
void TestNames()
{
DeviceAdapterTagInvalidDeviceId invalidTag;
vtkm::cont::DeviceAdapterTagError errorTag;
vtkm::cont::DeviceAdapterTagUndefined undefinedTag;
vtkm::cont::DeviceAdapterTagSerial serialTag;
vtkm::cont::DeviceAdapterTagTBB tbbTag;
vtkm::cont::DeviceAdapterTagOpenMP openmpTag;
vtkm::cont::DeviceAdapterTagCuda cudaTag;
TestName("InvalidDeviceId", invalidTag, invalidTag);
TestName("Error", errorTag, errorTag);
TestName("Undefined", undefinedTag, undefinedTag);
TestName("Serial", serialTag, serialTag);
TestName("TBB", tbbTag, tbbTag);
TestName("OpenMP", openmpTag, openmpTag);
TestName("Cuda", cudaTag, cudaTag);
}
} // end anon namespace
int UnitTestRuntimeDeviceNames(int, char* [])
{
return vtkm::cont::testing::Testing::Run(TestNames);
}

3
vtkm/exec/CellLocatorUniformGrid.h
View File

@ -43,7 +43,6 @@ class CellLocatorUniformGrid : public vtkm::exec::CellLocator
public:
VTKM_CONT
CellLocatorUniformGrid(const vtkm::Bounds& bounds,
const vtkm::Vec<vtkm::Id, 3>& dims,
const vtkm::Vec<vtkm::FloatDefault, 3> rangeTransform,
const vtkm::Id planeSize,
const vtkm::Id rowSize,
@ -51,7 +50,6 @@ public:
const vtkm::cont::ArrayHandleVirtualCoordinates& coords,
DeviceAdapter)
: Bounds(bounds)
, Dims(dims)
, RangeTransform(rangeTransform)
, PlaneSize(planeSize)
, RowSize(rowSize)
@ -101,7 +99,6 @@ private:
DeviceAdapter>::PortalConst;
vtkm::Bounds Bounds;
vtkm::Vec<vtkm::Id, 3> Dims;
vtkm::Vec<vtkm::FloatDefault, 3> RangeTransform;
vtkm::Id PlaneSize;
vtkm::Id RowSize;

4
vtkm/filter/CellAverage.hxx
View File

@ -55,8 +55,8 @@ inline VTKM_CONT vtkm::cont::DataSet CellAverage::DoExecute(
//If the input is implicit, we should know what to fall back to
vtkm::cont::ArrayHandle<T> outArray;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellAverage, DeviceAdapter> dispatcher(
this->Worklet);
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellAverage> dispatcher(this->Worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(vtkm::filter::ApplyPolicy(cellSet, policy), inField, outArray);

4
vtkm/filter/CellMeasures.hxx
View File

@ -50,8 +50,8 @@ inline VTKM_CONT vtkm::cont::DataSet CellMeasures<IntegrationType>::DoExecute(
const auto& cellset = input.GetCellSet(this->GetActiveCellSetIndex());
vtkm::cont::ArrayHandle<T> outArray;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellMeasure<IntegrationType>, DeviceAdapter>
dispatcher;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::CellMeasure<IntegrationType>> dispatcher;
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(vtkm::filter::ApplyPolicy(cellset, policy), points, outArray);
vtkm::cont::DataSet result;

3
vtkm/filter/CrossProduct.hxx
View File

@ -38,7 +38,8 @@ struct CrossProductFunctor
template <typename PrimaryFieldType, typename SecondaryFieldType>
void operator()(const SecondaryFieldType& secondaryField, const PrimaryFieldType& primaryField)
{
vtkm::worklet::DispatcherMapField<vtkm::worklet::CrossProduct, DeviceAdapter> dispatcher;
vtkm::worklet::DispatcherMapField<vtkm::worklet::CrossProduct> dispatcher;
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(primaryField,
vtkm::cont::make_ArrayHandleCast<vtkm::Vec<T, 3>>(secondaryField),
this->OutArray);

3
vtkm/filter/DotProduct.hxx
View File

@ -39,7 +39,8 @@ struct DotProductFunctor
template <typename PrimaryFieldType, typename SecondaryFieldType>
void operator()(const SecondaryFieldType& secondaryField, const PrimaryFieldType& primaryField)
{
vtkm::worklet::DispatcherMapField<vtkm::worklet::DotProduct, DeviceAdapter> dispatcher;
vtkm::worklet::DispatcherMapField<vtkm::worklet::DotProduct> dispatcher;
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(primaryField,
vtkm::cont::make_ArrayHandleCast<vtkm::Vec<T, 3>>(secondaryField),
this->OutArray);

4
vtkm/filter/OscillatorSource.hxx
View File

@ -84,8 +84,8 @@ inline VTKM_CONT vtkm::cont::DataSet OscillatorSource::DoExecute(
const DeviceAdapter&)
{
vtkm::cont::ArrayHandle<vtkm::Float64> outArray;
vtkm::worklet::DispatcherMapField<vtkm::worklet::OscillatorSource, DeviceAdapter> dispatcher(
this->Worklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::OscillatorSource> dispatcher(this->Worklet);
dispatcher.SetDevice(DeviceAdapter());
//todo, we need to use the policy to determine the valid conversions
//that the dispatcher should do

4
vtkm/filter/PointAverage.hxx
View File

@ -55,8 +55,8 @@ inline VTKM_CONT vtkm::cont::DataSet PointAverage::DoExecute(
//If the input is implicit, we should know what to fall back to
vtkm::cont::ArrayHandle<T> outArray;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::PointAverage, DeviceAdapter> dispatcher(
this->Worklet);
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::PointAverage> dispatcher(this->Worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(vtkm::filter::ApplyPolicy(cellSet, policy), inField, outArray);

4
vtkm/filter/PointElevation.hxx
View File

@ -63,8 +63,8 @@ inline VTKM_CONT vtkm::cont::DataSet PointElevation::DoExecute(
const DeviceAdapter&)
{
vtkm::cont::ArrayHandle<vtkm::Float64> outArray;
vtkm::worklet::DispatcherMapField<vtkm::worklet::PointElevation, DeviceAdapter> dispatcher(
this->Worklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::PointElevation> dispatcher(this->Worklet);
dispatcher.SetDevice(DeviceAdapter());
//todo, we need to use the policy to determine the valid conversions
//that the dispatcher should do

4
vtkm/filter/PointTransform.hxx
View File

@ -127,8 +127,8 @@ inline VTKM_CONT vtkm::cont::DataSet PointTransform<S>::DoExecute(
const DeviceAdapter&)
{
vtkm::cont::ArrayHandle<T> outArray;
vtkm::worklet::DispatcherMapField<vtkm::worklet::PointTransform<S>, DeviceAdapter> dispatcher(
this->Worklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::PointTransform<S>> dispatcher(this->Worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(field, outArray);

2
vtkm/filter/Tetrahedralize.hxx
View File

@ -39,7 +39,7 @@ public:
template <typename CellSetType>
void operator()(const CellSetType& cellset) const
{
this->OutCellSet = Worklet.Run(cellset, DeviceAdapter());
this->OutCellSet = Worklet.Run(cellset);
}
};
}

4
vtkm/filter/VectorMagnitude.hxx
View File

@ -48,8 +48,8 @@ inline VTKM_CONT vtkm::cont::DataSet VectorMagnitude::DoExecute(
using ReturnType = typename ::vtkm::detail::FloatingPointReturnType<T>::Type;
vtkm::cont::ArrayHandle<ReturnType> outArray;
vtkm::worklet::DispatcherMapField<vtkm::worklet::Magnitude, DeviceAdapter> dispatcher(
this->Worklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::Magnitude> dispatcher(this->Worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(field, outArray);

154
vtkm/rendering/Canvas.cxx
View File

@ -61,33 +61,6 @@ struct ClearBuffers : public vtkm::worklet::WorkletMapField
}
}; // struct ClearBuffers
struct ClearBuffersExecutor
{
using ColorBufferType = vtkm::rendering::Canvas::ColorBufferType;
using DepthBufferType = vtkm::rendering::Canvas::DepthBufferType;
ColorBufferType ColorBuffer;
DepthBufferType DepthBuffer;
VTKM_CONT
ClearBuffersExecutor(const ColorBufferType& colorBuffer, const DepthBufferType& depthBuffer)
: ColorBuffer(colorBuffer)
, DepthBuffer(depthBuffer)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device) const
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
ClearBuffers worklet;
vtkm::worklet::DispatcherMapField<ClearBuffers, Device> dispatcher(worklet);
dispatcher.Invoke(this->ColorBuffer, this->DepthBuffer);
return true;
}
}; // struct ClearBuffersExecutor
struct BlendBackground : public vtkm::worklet::WorkletMapField
{
vtkm::Vec<vtkm::Float32, 4> BackgroundColor;
@ -114,32 +87,6 @@ struct BlendBackground : public vtkm::worklet::WorkletMapField
}
}; // struct BlendBackground
struct BlendBackgroundExecutor
{
using ColorBufferType = vtkm::rendering::Canvas::ColorBufferType;
ColorBufferType ColorBuffer;
BlendBackground Worklet;
VTKM_CONT
BlendBackgroundExecutor(const ColorBufferType& colorBuffer,
const vtkm::Vec<vtkm::Float32, 4>& backgroundColor)
: ColorBuffer(colorBuffer)
, Worklet(backgroundColor)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device) const
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::worklet::DispatcherMapField<BlendBackground, Device> dispatcher(this->Worklet);
dispatcher.Invoke(this->ColorBuffer);
return true;
}
}; // struct BlendBackgroundExecutor
struct DrawColorSwatch : public vtkm::worklet::WorkletMapField
{
using ControlSignature = void(FieldIn<>, WholeArrayInOut<>);
@ -257,81 +204,6 @@ struct DrawColorBar : public vtkm::worklet::WorkletMapField
bool Horizontal;
}; // struct DrawColorBar
struct ColorSwatchExecutor
{
VTKM_CONT
ColorSwatchExecutor(vtkm::Id2 dims,
vtkm::Id2 xBounds,
vtkm::Id2 yBounds,
const vtkm::Vec<vtkm::Float32, 4>& color,
const vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Float32, 4>>& colorBuffer)
: Dims(dims)
, XBounds(xBounds)
, YBounds(yBounds)
, Color(color)
, ColorBuffer(colorBuffer)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device) const
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::Id totalPixels = (XBounds[1] - XBounds[0]) * (YBounds[1] - YBounds[0]);
vtkm::cont::ArrayHandleCounting<vtkm::Id> iterator(0, 1, totalPixels);
vtkm::worklet::DispatcherMapField<DrawColorSwatch, Device> dispatcher(
DrawColorSwatch(this->Dims, this->XBounds, this->YBounds, Color));
dispatcher.Invoke(iterator, this->ColorBuffer);
return true;
}
vtkm::Id2 Dims;
vtkm::Id2 XBounds;
vtkm::Id2 YBounds;
const vtkm::Vec<vtkm::Float32, 4>& Color;
const vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Float32, 4>>& ColorBuffer;
}; // struct ColorSwatchExecutor
struct ColorBarExecutor
{
VTKM_CONT
ColorBarExecutor(vtkm::Id2 dims,
vtkm::Id2 xBounds,
vtkm::Id2 yBounds,
bool horizontal,
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::UInt8, 4>>& colorMap,
const vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Float32, 4>>& colorBuffer)
: Dims(dims)
, XBounds(xBounds)
, YBounds(yBounds)
, Horizontal(horizontal)
, ColorMap(colorMap)
, ColorBuffer(colorBuffer)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device) const
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::Id totalPixels = (XBounds[1] - XBounds[0]) * (YBounds[1] - YBounds[0]);
vtkm::cont::ArrayHandleCounting<vtkm::Id> iterator(0, 1, totalPixels);
vtkm::worklet::DispatcherMapField<DrawColorBar, Device> dispatcher(
DrawColorBar(this->Dims, this->XBounds, this->YBounds, this->Horizontal));
dispatcher.Invoke(iterator, this->ColorBuffer, this->ColorMap);
return true;
}
vtkm::Id2 Dims;
vtkm::Id2 XBounds;
vtkm::Id2 YBounds;
bool Horizontal;
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::UInt8, 4>>& ColorMap;
const vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Float32, 4>>& ColorBuffer;
}; // struct ColorSwatchExecutor
} // namespace internal
struct Canvas::CanvasInternals
@ -441,9 +313,8 @@ void Canvas::Activate()
void Canvas::Clear()
{
// TODO: Should the rendering library support policies or some other way to
// configure with custom devices?
vtkm::cont::TryExecute(internal::ClearBuffersExecutor(GetColorBuffer(), GetDepthBuffer()));
vtkm::worklet::DispatcherMapField<internal::ClearBuffers>().Invoke(this->GetColorBuffer(),
this->GetDepthBuffer());
}
void Canvas::Finish()
@ -452,8 +323,9 @@ void Canvas::Finish()
void Canvas::BlendBackground()
{
vtkm::cont::TryExecute(
internal::BlendBackgroundExecutor(GetColorBuffer(), GetBackgroundColor().Components));
vtkm::worklet::DispatcherMapField<internal::BlendBackground>(
this->GetBackgroundColor().Components)
.Invoke(this->GetColorBuffer());
}
void Canvas::ResizeBuffers(vtkm::Id width, vtkm::Id height)
@ -491,8 +363,12 @@ void Canvas::AddColorSwatch(const vtkm::Vec<vtkm::Float64, 2>& point0,
y[1] = static_cast<vtkm::Id>(((point2[1] + 1.) / 2.) * height + .5);
vtkm::Id2 dims(this->GetWidth(), this->GetHeight());
vtkm::cont::TryExecute(
internal::ColorSwatchExecutor(dims, x, y, color.Components, this->GetColorBuffer()));
vtkm::Id totalPixels = (x[1] - x[0]) * (y[1] - y[0]);
vtkm::cont::ArrayHandleCounting<vtkm::Id> iterator(0, 1, totalPixels);
vtkm::worklet::DispatcherMapField<internal::DrawColorSwatch> dispatcher(
internal::DrawColorSwatch(dims, x, y, color.Components));
dispatcher.Invoke(iterator, this->GetColorBuffer());
}
void Canvas::AddColorSwatch(const vtkm::Float64 x0,
@ -552,8 +428,12 @@ void Canvas::AddColorBar(const vtkm::Bounds& bounds,
colorTable.Sample(static_cast<vtkm::Int32>(numSamples), colorMap);
vtkm::Id2 dims(this->GetWidth(), this->GetHeight());
vtkm::cont::TryExecute(
internal::ColorBarExecutor(dims, x, y, horizontal, colorMap, this->GetColorBuffer()));
vtkm::Id totalPixels = (x[1] - x[0]) * (y[1] - y[0]);
vtkm::cont::ArrayHandleCounting<vtkm::Id> iterator(0, 1, totalPixels);
vtkm::worklet::DispatcherMapField<internal::DrawColorBar> dispatcher(
internal::DrawColorBar(dims, x, y, horizontal));
dispatcher.Invoke(iterator, this->GetColorBuffer(), colorMap);
}
void Canvas::AddColorBar(vtkm::Float32 x,

55
vtkm/rendering/CanvasRayTracer.cxx
View File

@ -110,57 +110,24 @@ public:
}
}; //class SurfaceConverter
template <typename Precision>
struct WriteFunctor
{
protected:
vtkm::rendering::CanvasRayTracer* Canvas;
const vtkm::rendering::raytracing::Ray<Precision>& Rays;
const vtkm::cont::ArrayHandle<Precision>& Colors;
const vtkm::rendering::Camera& CameraView;
vtkm::Matrix<vtkm::Float32, 4, 4> ViewProjMat;
public:
VTKM_CONT
WriteFunctor(vtkm::rendering::CanvasRayTracer* canvas,
const vtkm::rendering::raytracing::Ray<Precision>& rays,
const vtkm::cont::ArrayHandle<Precision>& colors,
const vtkm::rendering::Camera& camera)
: Canvas(canvas)
, Rays(rays)
, Colors(colors)
, CameraView(camera)
{
ViewProjMat = vtkm::MatrixMultiply(
CameraView.CreateProjectionMatrix(Canvas->GetWidth(), Canvas->GetHeight()),
CameraView.CreateViewMatrix());
}
template <typename Device>
VTKM_CONT bool operator()(Device)
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::worklet::DispatcherMapField<SurfaceConverter, Device>(SurfaceConverter(ViewProjMat))
.Invoke(Rays.PixelIdx,
Colors,
Rays.Distance,
Rays.Origin,
Rays.Dir,
Canvas->GetDepthBuffer(),
Canvas->GetColorBuffer());
return true;
}
};
template <typename Precision>
VTKM_CONT void WriteToCanvas(const vtkm::rendering::raytracing::Ray<Precision>& rays,
const vtkm::cont::ArrayHandle<Precision>& colors,
const vtkm::rendering::Camera& camera,
vtkm::rendering::CanvasRayTracer* canvas)
{
WriteFunctor<Precision> functor(canvas, rays, colors, camera);
vtkm::Matrix<vtkm::Float32, 4, 4> viewProjMat =
vtkm::MatrixMultiply(camera.CreateProjectionMatrix(canvas->GetWidth(), canvas->GetHeight()),
camera.CreateViewMatrix());
vtkm::cont::TryExecute(functor);
vtkm::worklet::DispatcherMapField<SurfaceConverter>(SurfaceConverter(viewProjMat))
.Invoke(rays.PixelIdx,
colors,
rays.Distance,
rays.Origin,
rays.Dir,
canvas->GetDepthBuffer(),
canvas->GetColorBuffer());
//Force the transfer so the vectors contain data from device
canvas->GetColorBuffer().GetPortalControl().Get(0);

5
vtkm/rendering/MapperGL.cxx
View File

@ -179,8 +179,9 @@ struct MapColorAndVerticesInvokeFunctor
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
auto colorHandle = this->ColorTable.GetHandleForExecution();
MapColorAndVertices worklet(colorHandle->PrepareForExecution(device), this->SMin, this->SDiff);
MapColorAndVertices worklet(
this->ColorTable.PrepareForExecution(device), this->SMin, this->SDiff);
vtkm::worklet::DispatcherMapField<MapColorAndVertices, Device> dispatcher(worklet);
vtkm::cont::ArrayHandleIndex indexArray(this->TriangleIndices.GetNumberOfValues());

91
vtkm/rendering/MapperWireframer.cxx
View File

@ -59,21 +59,6 @@ public:
}
}; // conn
struct ConnFunctor
{
template <typename Device>
VTKM_CONT bool operator()(Device,
vtkm::cont::ArrayHandleCounting<vtkm::Id>& iter,
vtkm::cont::ArrayHandle<vtkm::Id>& conn)
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::worklet::DispatcherMapField<CreateConnectivity, Device>(CreateConnectivity())
.Invoke(iter, conn);
return true;
}
};
class Convert1DCoordinates : public vtkm::worklet::WorkletMapField
{
private:
@ -123,26 +108,6 @@ public:
}
}; // convert coords
struct ConvertFunctor
{
template <typename Device, typename CoordType, typename ScalarType>
VTKM_CONT bool operator()(Device,
CoordType coords,
ScalarType scalars,
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Float32, 3>>& outCoords,
vtkm::cont::ArrayHandle<vtkm::Float32>& outScalars,
bool logY,
bool logX)
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::worklet::DispatcherMapField<Convert1DCoordinates, Device>(
Convert1DCoordinates(logY, logX))
.Invoke(coords, scalars, outCoords, outScalars);
return true;
}
};
#if defined(VTKM_MSVC)
#pragma warning(push)
#pragma warning(disable : 4127) //conditional expression is constant
@ -176,10 +141,10 @@ struct EdgesExtracter : public vtkm::worklet::WorkletMapPointToCell
using ScatterType = vtkm::worklet::ScatterCounting;
VTKM_CONT
template <typename CountArrayType, typename DeviceTag>
static ScatterType MakeScatter(const CountArrayType& counts, DeviceTag device)
template <typename CountArrayType>
static ScatterType MakeScatter(const CountArrayType& counts)
{
return ScatterType(counts, device);
return ScatterType(counts);
}
template <typename CellShapeTag, typename PointIndexVecType, typename EdgeIndexVecType>
@ -212,33 +177,6 @@ struct EdgesExtracter : public vtkm::worklet::WorkletMapPointToCell
#if defined(VTKM_MSVC)
#pragma warning(pop)
#endif
struct ExtractUniqueEdges
{
vtkm::cont::DynamicCellSet CellSet;
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Id, 2>> EdgeIndices;
VTKM_CONT
ExtractUniqueEdges(const vtkm::cont::DynamicCellSet& cellSet)
: CellSet(cellSet)
{
}
template <typename DeviceTag>
VTKM_CONT bool operator()(DeviceTag)
{
VTKM_IS_DEVICE_ADAPTER_TAG(DeviceTag);
vtkm::cont::ArrayHandle<vtkm::IdComponent> counts;
vtkm::worklet::DispatcherMapTopology<EdgesCounter, DeviceTag>().Invoke(CellSet, counts);
vtkm::worklet::DispatcherMapTopology<EdgesExtracter, DeviceTag> extractDispatcher(
EdgesExtracter::MakeScatter(counts, DeviceTag()));
extractDispatcher.Invoke(CellSet, EdgeIndices);
vtkm::cont::DeviceAdapterAlgorithm<DeviceTag>::template Sort<vtkm::Id2>(EdgeIndices);
vtkm::cont::DeviceAdapterAlgorithm<DeviceTag>::template Unique<vtkm::Id2>(EdgeIndices);
return true;
}
}; // struct ExtractUniqueEdges
} // namespace
struct MapperWireframer::InternalsType
@ -340,13 +278,9 @@ void MapperWireframer::RenderCells(const vtkm::cont::DynamicCellSet& inCellSet,
//
// Convert the cell set into something we can draw
//
vtkm::cont::TryExecute(ConvertFunctor(),
coords.GetData(),
inScalarField.GetData(),
newCoords,
newScalars,
this->LogarithmY,
this->LogarithmX);
vtkm::worklet::DispatcherMapField<Convert1DCoordinates>(
Convert1DCoordinates(this->LogarithmY, this->LogarithmX))
.Invoke(coords.GetData(), inScalarField.GetData(), newCoords, newScalars);
actualCoords = vtkm::cont::CoordinateSystem("coords", newCoords);
actualField = vtkm::cont::Field(
@ -357,7 +291,7 @@ void MapperWireframer::RenderCells(const vtkm::cont::DynamicCellSet& inCellSet,
vtkm::cont::ArrayHandleCounting<vtkm::Id> iter =
vtkm::cont::make_ArrayHandleCounting(vtkm::Id(0), vtkm::Id(1), numCells);
conn.Allocate(numCells * 2);
vtkm::cont::TryExecute(ConnFunctor(), iter, conn);
vtkm::worklet::DispatcherMapField<CreateConnectivity>(CreateConnectivity()).Invoke(iter, conn);
vtkm::cont::CellSetSingleType<> newCellSet("cells");
newCellSet.Fill(newCoords.GetNumberOfValues(), vtkm::CELL_SHAPE_LINE, 2, conn);
@ -391,9 +325,14 @@ void MapperWireframer::RenderCells(const vtkm::cont::DynamicCellSet& inCellSet,
}
// Extract unique edges from the cell set.
ExtractUniqueEdges extracter(cellSet);
vtkm::cont::TryExecute(extracter);
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Id, 2>> edgeIndices = extracter.EdgeIndices;
vtkm::cont::ArrayHandle<vtkm::IdComponent> counts;
vtkm::cont::ArrayHandle<vtkm::Vec<vtkm::Id, 2>> edgeIndices;
vtkm::worklet::DispatcherMapTopology<EdgesCounter>().Invoke(cellSet, counts);
vtkm::worklet::DispatcherMapTopology<EdgesExtracter> extractDispatcher(
EdgesExtracter::MakeScatter(counts));
extractDispatcher.Invoke(cellSet, edgeIndices);
vtkm::cont::Algorithm::template Sort<vtkm::Id2>(edgeIndices);
vtkm::cont::Algorithm::template Unique<vtkm::Id2>(edgeIndices);
Wireframer renderer(
this->Internals->Canvas, this->Internals->ShowInternalZones, this->Internals->IsOverlay);

3
vtkm/rendering/TextRenderer.cxx
View File

@ -165,7 +165,8 @@ struct RenderBitmapFontExecutor
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::worklet::DispatcherMapField<RenderBitmapFont, Device> dispatcher(Worklet);
vtkm::worklet::DispatcherMapField<RenderBitmapFont> dispatcher(Worklet);
dispatcher.SetDevice(Device());
dispatcher.Invoke(
ScreenCoords, TextureCoords, FontTexture.GetExecObjectFactory(), ColorBuffer, DepthBuffer);
return true;

10
vtkm/rendering/Wireframer.h
View File

@ -554,12 +554,14 @@ private:
ColorMap,
FrameBuffer,
Camera.GetClippingRange());
vtkm::worklet::DispatcherMapField<EdgePlotter<DeviceTag>, DeviceTag>(plotter).Invoke(
PointIndices, Coordinates, ScalarField.GetData());
vtkm::worklet::DispatcherMapField<EdgePlotter<DeviceTag>> plotterDispatcher(plotter);
plotterDispatcher.SetDevice(DeviceTag());
plotterDispatcher.Invoke(PointIndices, Coordinates, ScalarField.GetData());
BufferConverter converter;
vtkm::worklet::DispatcherMapField<BufferConverter, DeviceTag>(converter).Invoke(
FrameBuffer, Canvas->GetDepthBuffer(), Canvas->GetColorBuffer());
vtkm::worklet::DispatcherMapField<BufferConverter> converterDispatcher(converter);
converterDispatcher.SetDevice(DeviceTag());
converterDispatcher.Invoke(FrameBuffer, Canvas->GetDepthBuffer(), Canvas->GetColorBuffer());
}
VTKM_CONT

121
vtkm/rendering/raytracing/BoundingVolumeHierarchy.cxx
View File

@ -608,7 +608,8 @@ VTKM_CONT void LinearBVHBuilder::SortAABBS(
//create array of indexes to be sorted with morton codes
vtkm::cont::ArrayHandle<vtkm::Id> iterator;
iterator.PrepareForOutput(bvh.GetNumberOfPrimitives(), Device());
vtkm::worklet::DispatcherMapField<CountingIterator, Device> iteratorDispatcher;
vtkm::worklet::DispatcherMapField<CountingIterator> iteratorDispatcher;
iteratorDispatcher.SetDevice(Device());
iteratorDispatcher.Invoke(iterator);
//std::cout<<"\n\n\n";
@ -623,56 +624,77 @@ VTKM_CONT void LinearBVHBuilder::SortAABBS(
tempStorage = new vtkm::cont::ArrayHandle<vtkm::Float32>();
//xmins
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>, Device>(
GatherFloat32<Device>(*bvh.xmins, *tempStorage, arraySize))
.Invoke(iterator);
{
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>> dispatcher(
GatherFloat32<Device>(*bvh.xmins, *tempStorage, arraySize));
dispatcher.SetDevice(Device());
dispatcher.Invoke(iterator);
}
tempPtr = bvh.xmins;
bvh.xmins = tempStorage;
tempStorage = tempPtr;
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>, Device>(
GatherFloat32<Device>(*bvh.ymins, *tempStorage, arraySize))
.Invoke(iterator);
{
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>> dispatcher(
GatherFloat32<Device>(*bvh.ymins, *tempStorage, arraySize));
dispatcher.SetDevice(Device());
dispatcher.Invoke(iterator);
}
tempPtr = bvh.ymins;
bvh.ymins = tempStorage;
tempStorage = tempPtr;
//zmins
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>, Device>(
GatherFloat32<Device>(*bvh.zmins, *tempStorage, arraySize))
.Invoke(iterator);
{
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>> dispatcher(
GatherFloat32<Device>(*bvh.zmins, *tempStorage, arraySize));
dispatcher.SetDevice(Device());
dispatcher.Invoke(iterator);
}
tempPtr = bvh.zmins;
bvh.zmins = tempStorage;
tempStorage = tempPtr;
//xmaxs
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>, Device>(
GatherFloat32<Device>(*bvh.xmaxs, *tempStorage, arraySize))
.Invoke(iterator);
{
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>> dispatcher(
GatherFloat32<Device>(*bvh.xmaxs, *tempStorage, arraySize));
dispatcher.SetDevice(Device());
dispatcher.Invoke(iterator);
}
tempPtr = bvh.xmaxs;
bvh.xmaxs = tempStorage;
tempStorage = tempPtr;
//ymaxs
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>, Device>(
GatherFloat32<Device>(*bvh.ymaxs, *tempStorage, arraySize))
.Invoke(iterator);
{
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>> dispatcher(
GatherFloat32<Device>(*bvh.ymaxs, *tempStorage, arraySize));
dispatcher.SetDevice(Device());
dispatcher.Invoke(iterator);
}
tempPtr = bvh.ymaxs;
bvh.ymaxs = tempStorage;
tempStorage = tempPtr;
//zmaxs
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>, Device>(
GatherFloat32<Device>(*bvh.zmaxs, *tempStorage, arraySize))
.Invoke(iterator);
{
vtkm::worklet::DispatcherMapField<GatherFloat32<Device>> dispatcher(
GatherFloat32<Device>(*bvh.zmaxs, *tempStorage, arraySize));
dispatcher.SetDevice(Device());
dispatcher.Invoke(iterator);
}
tempPtr = bvh.zmaxs;
bvh.zmaxs = tempStorage;
tempStorage = tempPtr;
vtkm::worklet::DispatcherMapField<GatherVecCast<Device>, Device>(
GatherVecCast<Device>(triangleIndices, outputTriangleIndices, arraySize))
.Invoke(iterator);
{
vtkm::worklet::DispatcherMapField<GatherVecCast<Device>> dispatcher(
GatherVecCast<Device>(triangleIndices, outputTriangleIndices, arraySize));
dispatcher.SetDevice(Device());
dispatcher.Invoke(iterator);
}
delete tempStorage;
} // method SortAABBs
@ -699,15 +721,16 @@ VTKM_CONT void LinearBVHBuilder::RunOnDevice(LinearBVH& linearBVH, Device device
BVHData bvh(numBBoxes, device);
vtkm::cont::Timer<Device> timer;
vtkm::worklet::DispatcherMapField<FindAABBs, Device>(FindAABBs())
.Invoke(triangleIndices,
*bvh.xmins,
*bvh.ymins,
*bvh.zmins,
*bvh.xmaxs,
*bvh.ymaxs,
*bvh.zmaxs,
coordsHandle);
vtkm::worklet::DispatcherMapField<FindAABBs> findAABBDispatcher{ (FindAABBs{}) };
findAABBDispatcher.SetDevice(Device());
findAABBDispatcher.Invoke(triangleIndices,
*bvh.xmins,
*bvh.ymins,
*bvh.zmins,
*bvh.xmaxs,
*bvh.ymaxs,
*bvh.zmaxs,
coordsHandle);
vtkm::Float64 time = timer.GetElapsedTime();
logger->AddLogData("find_aabb", time);
@ -742,10 +765,11 @@ VTKM_CONT void LinearBVHBuilder::RunOnDevice(LinearBVH& linearBVH, Device device
}
//Generate the morton codes
vtkm::worklet::DispatcherMapField<MortonCodeAABB, Device>(
MortonCodeAABB(inverseExtent, minExtent))
.Invoke(
*bvh.xmins, *bvh.ymins, *bvh.zmins, *bvh.xmaxs, *bvh.ymaxs, *bvh.zmaxs, bvh.mortonCodes);
vtkm::worklet::DispatcherMapField<MortonCodeAABB> mortonCodeAABBDispatcher(
MortonCodeAABB(inverseExtent, minExtent));
mortonCodeAABBDispatcher.SetDevice(Device());
mortonCodeAABBDispatcher.Invoke(
*bvh.xmins, *bvh.ymins, *bvh.zmins, *bvh.xmaxs, *bvh.ymaxs, *bvh.zmaxs, bvh.mortonCodes);
time = timer.GetElapsedTime();
logger->AddLogData("morton_codes", time);
@ -759,9 +783,10 @@ VTKM_CONT void LinearBVHBuilder::RunOnDevice(LinearBVH& linearBVH, Device device
logger->AddLogData("sort_aabbs", time);
timer.Reset();
vtkm::worklet::DispatcherMapField<TreeBuilder<Device>, Device>(
TreeBuilder<Device>(bvh.mortonCodes, bvh.parent, bvh.GetNumberOfPrimitives()))
.Invoke(bvh.leftChild, bvh.rightChild);
vtkm::worklet::DispatcherMapField<TreeBuilder<Device>> treeBuilderDispatcher(
TreeBuilder<Device>(bvh.mortonCodes, bvh.parent, bvh.GetNumberOfPrimitives()));
treeBuilderDispatcher.SetDevice(Device());
treeBuilderDispatcher.Invoke(bvh.leftChild, bvh.rightChild);
time = timer.GetElapsedTime();
logger->AddLogData("build_tree", time);
@ -772,15 +797,23 @@ VTKM_CONT void LinearBVHBuilder::RunOnDevice(LinearBVH& linearBVH, Device device
vtkm::cont::ArrayHandle<vtkm::Int32> counters;
counters.PrepareForOutput(bvh.GetNumberOfPrimitives() - 1, Device());
vtkm::Int32 zero = 0;
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Int32>, Device>(MemSet<vtkm::Int32>(zero))
.Invoke(counters);
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Int32>> resetCountersDispatcher(
(MemSet<vtkm::Int32>(zero)));
resetCountersDispatcher.SetDevice(Device());
resetCountersDispatcher.Invoke(counters);
vtkm::cont::AtomicArray<vtkm::Int32> atomicCounters(counters);
vtkm::worklet::DispatcherMapField<PropagateAABBs<Device>, Device>(
PropagateAABBs<Device>(
bvh.parent, bvh.leftChild, bvh.rightChild, primitiveCount, linearBVH.FlatBVH, atomicCounters))
.Invoke(*bvh.xmins, *bvh.ymins, *bvh.zmins, *bvh.xmaxs, *bvh.ymaxs, *bvh.zmaxs);
vtkm::worklet::DispatcherMapField<PropagateAABBs<Device>> propagateAABBDispatcher(
PropagateAABBs<Device>(bvh.parent,
bvh.leftChild,
bvh.rightChild,
primitiveCount,
linearBVH.FlatBVH,
atomicCounters));
propagateAABBDispatcher.SetDevice(Device());
propagateAABBDispatcher.Invoke(
*bvh.xmins, *bvh.ymins, *bvh.zmins, *bvh.xmaxs, *bvh.ymaxs, *bvh.zmaxs);
time = timer.GetElapsedTime();
logger->AddLogData("propagate_aabbs", time);

166
vtkm/rendering/raytracing/Camera.cxx
View File

@ -20,6 +20,7 @@
#include <vtkm/VectorAnalysis.h>
#include <vtkm/cont/Algorithm.h>
#include <vtkm/cont/ErrorBadValue.h>
#include <vtkm/cont/Timer.h>
#include <vtkm/cont/TryExecute.h>
@ -659,69 +660,38 @@ struct Camera::CreateRaysFunctor
}
};
struct Camera::PixelDataFunctor
{
vtkm::rendering::raytracing::Camera* Self;
const vtkm::cont::CoordinateSystem& Coords;
vtkm::Int32& ActivePixels;
vtkm::Float32& AveDistPerRay;
VTKM_CONT
PixelDataFunctor(vtkm::rendering::raytracing::Camera* self,
const vtkm::cont::CoordinateSystem& coords,
vtkm::Int32& activePixels,
vtkm::Float32& aveDistPerRay)
: Self(self)
, Coords(coords)
, ActivePixels(activePixels)
, AveDistPerRay(aveDistPerRay)
{
}
template <typename Device>
VTKM_CONT bool operator()(Device)
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::Bounds boundingBox = Coords.GetBounds();
Self->FindSubset(boundingBox);
//Reset the camera look vector
Self->Look = Self->LookAt - Self->Position;
vtkm::Normalize(Self->Look);
const int size = Self->SubsetWidth * Self->SubsetHeight;
vtkm::cont::ArrayHandle<vtkm::Float32> dists;
vtkm::cont::ArrayHandle<vtkm::Int32> hits;
dists.PrepareForOutput(size, Device());
hits.PrepareForOutput(size, Device());
//Create the ray direction
vtkm::worklet::DispatcherMapField<PixelData, Device>(PixelData(Self->Width,
Self->Height,
Self->FovX,
Self->FovY,
Self->Look,
Self->Up,
Self->Zoom,
Self->SubsetWidth,
Self->SubsetMinX,
Self->SubsetMinY,
Self->Position,
boundingBox))
.Invoke(hits, dists); //X Y Z
ActivePixels = vtkm::cont::DeviceAdapterAlgorithm<Device>::Reduce(hits, vtkm::Int32(0));
AveDistPerRay = vtkm::cont::DeviceAdapterAlgorithm<Device>::Reduce(dists, vtkm::Float32(0)) /
vtkm::Float32(ActivePixels);
return true;
}
};
void Camera::GetPixelData(const vtkm::cont::CoordinateSystem& coords,
vtkm::Int32& activePixels,
vtkm::Float32& aveRayDistance)
{
vtkm::Bounds boundingBox = coords.GetBounds();
this->FindSubset(boundingBox);
//Reset the camera look vector
this->Look = this->LookAt - this->Position;
vtkm::Normalize(this->Look);
const int size = this->SubsetWidth * this->SubsetHeight;
vtkm::cont::ArrayHandle<vtkm::Float32> dists;
vtkm::cont::ArrayHandle<vtkm::Int32> hits;
dists.Allocate(size);
hits.Allocate(size);
PixelDataFunctor functor(this, coords, activePixels, aveRayDistance);
vtkm::cont::TryExecute(functor);
//Create the ray direction
vtkm::worklet::DispatcherMapField<PixelData>(PixelData(this->Width,
this->Height,
this->FovX,
this->FovY,
this->Look,
this->Up,
this->Zoom,
this->SubsetWidth,
this->SubsetMinX,
this->SubsetMinY,
this->Position,
boundingBox))
.Invoke(hits, dists); //X Y Z
activePixels = vtkm::cont::Algorithm::Reduce(hits, vtkm::Int32(0));
aveRayDistance =
vtkm::cont::Algorithm::Reduce(dists, vtkm::Float32(0)) / vtkm::Float32(activePixels);
}
VTKM_CONT
@ -758,18 +728,20 @@ VTKM_CONT void Camera::CreateRaysOnDevice(Ray<Precision>& rays,
Precision infinity;
GetInfinity(infinity);
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(MemSet<Precision>(infinity))
.Invoke(rays.MaxDistance);
vtkm::worklet::DispatcherMapField<MemSet<Precision>> memSetInfDispatcher(
(MemSet<Precision>(infinity)));
memSetInfDispatcher.SetDevice(Device());
memSetInfDispatcher.Invoke(rays.MaxDistance);
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(MemSet<Precision>(0.f))
.Invoke(rays.MinDistance);
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(MemSet<Precision>(0.f))
.Invoke(rays.Distance);
vtkm::worklet::DispatcherMapField<MemSet<Precision>> memSet0Dispatcher((MemSet<Precision>(0.f)));
memSet0Dispatcher.SetDevice(Device());
memSet0Dispatcher.Invoke(rays.MinDistance);
memSet0Dispatcher.Invoke(rays.Distance);
//Reset the Rays Hit Index to -2
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Id>, Device>(MemSet<vtkm::Id>(-2))
.Invoke(rays.HitIdx);
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Id>> memSetM2Dispatcher((MemSet<vtkm::Id>(-2)));
memSetM2Dispatcher.SetDevice(Device());
memSetM2Dispatcher.Invoke(rays.HitIdx);
vtkm::Float64 time = timer.GetElapsedTime();
logger->AddLogData("camera_memset", time);
@ -781,25 +753,26 @@ VTKM_CONT void Camera::CreateRaysOnDevice(Ray<Precision>& rays,
if (ortho)
{
vtkm::worklet::DispatcherMapField<Ortho2DRayGen, Device>(Ortho2DRayGen(this->Width,
this->Height,
this->Zoom,
this->SubsetWidth,
this->SubsetMinX,
this->SubsetMinY,
this->CameraView))
.Invoke(rays.DirX,
rays.DirY,
rays.DirZ,
rays.OriginX,
rays.OriginY,
rays.OriginZ,
rays.PixelIdx); //X Y Z
vtkm::worklet::DispatcherMapField<Ortho2DRayGen> dispatcher(Ortho2DRayGen(this->Width,
this->Height,
this->Zoom,
this->SubsetWidth,
this->SubsetMinX,
this->SubsetMinY,
this->CameraView));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.DirX,
rays.DirY,
rays.DirZ,
rays.OriginX,
rays.OriginY,
rays.OriginZ,
rays.PixelIdx); //X Y Z
}
else
{
//Create the ray direction
vtkm::worklet::DispatcherMapField<PerspectiveRayGen, Device>(
vtkm::worklet::DispatcherMapField<PerspectiveRayGen> dispatcher(
PerspectiveRayGen(this->Width,
this->Height,
this->FovX,
@ -809,23 +782,24 @@ VTKM_CONT void Camera::CreateRaysOnDevice(Ray<Precision>& rays,
this->Zoom,
this->SubsetWidth,
this->SubsetMinX,
this->SubsetMinY))
.Invoke(rays.DirX,
rays.DirY,
rays.DirZ,
rays.PixelIdx); //X Y Z
this->SubsetMinY));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.DirX, rays.DirY, rays.DirZ, rays.PixelIdx); //X Y Z
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(
MemSet<Precision>(this->Position[0]))
.Invoke(rays.OriginX);
vtkm::worklet::DispatcherMapField<MemSet<Precision>> memSetPos0Dispatcher(
MemSet<Precision>(this->Position[0]));
memSetPos0Dispatcher.SetDevice(Device());
memSetPos0Dispatcher.Invoke(rays.OriginX);
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(
MemSet<Precision>(this->Position[1]))
.Invoke(rays.OriginY);
vtkm::worklet::DispatcherMapField<MemSet<Precision>> memSetPos1Dispatcher(
MemSet<Precision>(this->Position[1]));
memSetPos1Dispatcher.SetDevice(Device());
memSetPos1Dispatcher.Invoke(rays.OriginY);
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(
MemSet<Precision>(this->Position[2]))
.Invoke(rays.OriginZ);
vtkm::worklet::DispatcherMapField<MemSet<Precision>> memSetPos2Dispatcher(
MemSet<Precision>(this->Position[2]));
memSetPos2Dispatcher.SetDevice(Device());
memSetPos2Dispatcher.Invoke(rays.OriginZ);
}
time = timer.GetElapsedTime();

50
vtkm/rendering/raytracing/ChannelBuffer.cxx
View File

@ -39,26 +39,6 @@ namespace rendering
namespace raytracing
{
template <typename Precision, typename OpWorklet>
struct BufferMathFunctor
{
ChannelBuffer<Precision>* Self;
const ChannelBuffer<Precision>* Other;
BufferMathFunctor(ChannelBuffer<Precision>* self, const ChannelBuffer<Precision>* other)
: Self(self)
, Other(other)
{
}
template <typename Device>
bool operator()(Device vtkmNotUsed(device))
{
vtkm::worklet::DispatcherMapField<OpWorklet, Device>(OpWorklet())
.Invoke(Other->Buffer, Self->Buffer);
return true;
}
};
class BufferAddition : public vtkm::worklet::WorkletMapField
{
public:
@ -150,8 +130,7 @@ void ChannelBuffer<Precision>::AddBuffer(const ChannelBuffer<Precision>& other)
if (this->Size != other.GetSize())
throw vtkm::cont::ErrorBadValue("ChannelBuffer add: size must be equal");
BufferMathFunctor<Precision, BufferAddition> functor(this, &other);
vtkm::cont::TryExecute(functor);
vtkm::worklet::DispatcherMapField<BufferAddition>().Invoke(other.Buffer, this->Buffer);
}
template <typename Precision>
@ -162,8 +141,7 @@ void ChannelBuffer<Precision>::MultiplyBuffer(const ChannelBuffer<Precision>& ot
if (this->Size != other.GetSize())
throw vtkm::cont::ErrorBadValue("ChannelBuffer add: size must be equal");
BufferMathFunctor<Precision, BufferMultiply> functor(this, &other);
vtkm::cont::TryExecute(functor);
vtkm::worklet::DispatcherMapField<BufferMultiply>().Invoke(other.Buffer, this->Buffer);
}
template <typename Precision>
@ -221,9 +199,10 @@ struct ExtractChannelFunctor
bool operator()(Device device)
{
Output.PrepareForOutput(Self->GetSize(), device);
vtkm::worklet::DispatcherMapField<ExtractChannel, Device>(
ExtractChannel(Self->GetNumChannels(), Channel))
.Invoke(Output, Self->Buffer);
vtkm::worklet::DispatcherMapField<ExtractChannel> dispatcher(
ExtractChannel(Self->GetNumChannels(), Channel));
dispatcher.SetDevice(Device());
dispatcher.Invoke(Output, Self->Buffer);
return true;
}
};
@ -307,8 +286,9 @@ struct ExpandFunctorSignature
Output->Buffer.PrepareForOutput(totalSize, device);
ChannelBufferOperations::InitChannels(*Output, Signature, device);
vtkm::worklet::DispatcherMapField<Expand, Device>(Expand(NumChannels))
.Invoke(Input, SparseIndexes, Output->Buffer);
vtkm::worklet::DispatcherMapField<Expand> dispatcher((Expand(NumChannels)));
dispatcher.SetDevice(Device());
dispatcher.Invoke(Input, SparseIndexes, Output->Buffer);
return true;
}
@ -347,8 +327,9 @@ struct ExpandFunctor
Output->Buffer.PrepareForOutput(totalSize, device);
ChannelBufferOperations::InitConst(*Output, InitVal, device);
vtkm::worklet::DispatcherMapField<Expand, Device>(Expand(NumChannels))
.Invoke(Input, SparseIndexes, Output->Buffer);
vtkm::worklet::DispatcherMapField<Expand> dispatcher((Expand(NumChannels)));
dispatcher.SetDevice(Device());
dispatcher.Invoke(Input, SparseIndexes, Output->Buffer);
return true;
}
@ -410,9 +391,10 @@ struct NormalizeFunctor
asField.GetRange(&range);
Precision minScalar = static_cast<Precision>(range.Min);
Precision maxScalar = static_cast<Precision>(range.Max);
vtkm::worklet::DispatcherMapField<NormalizeBuffer<Precision>, Device>(
NormalizeBuffer<Precision>(minScalar, maxScalar, Invert))
.Invoke(Input);
vtkm::worklet::DispatcherMapField<NormalizeBuffer<Precision>> dispatcher(
NormalizeBuffer<Precision>(minScalar, maxScalar, Invert));
dispatcher.SetDevice(Device());
dispatcher.Invoke(Input);
return true;
}

20
vtkm/rendering/raytracing/ChannelBufferOperations.h
View File

@ -116,9 +116,10 @@ public:
vtkm::cont::ArrayHandle<Precision> compactedBuffer;
compactedBuffer.PrepareForOutput(newSize * buffer.NumChannels, Device());
vtkm::worklet::DispatcherMapField<detail::CompactBuffer, Device>(
detail::CompactBuffer(buffer.NumChannels))
.Invoke(masks, buffer.Buffer, offsets, compactedBuffer);
vtkm::worklet::DispatcherMapField<detail::CompactBuffer> dispatcher(
detail::CompactBuffer(buffer.NumChannels));
dispatcher.SetDevice(Device());
dispatcher.Invoke(masks, buffer.Buffer, offsets, compactedBuffer);
buffer.Buffer = compactedBuffer;
buffer.Size = newSize;
}
@ -133,16 +134,19 @@ public:
std::string msg = "ChannelBuffer: number of bins in sourse signature must match NumChannels";
throw vtkm::cont::ErrorBadValue(msg);
}
vtkm::worklet::DispatcherMapField<detail::InitBuffer, Device>(
detail::InitBuffer(buffer.NumChannels))
.Invoke(buffer.Buffer, sourceSignature);
vtkm::worklet::DispatcherMapField<detail::InitBuffer> initBufferDispatcher(
detail::InitBuffer(buffer.NumChannels));
initBufferDispatcher.SetDevice(Device());
initBufferDispatcher.Invoke(buffer.Buffer, sourceSignature);
}
template <typename Device, typename Precision>
static void InitConst(ChannelBuffer<Precision>& buffer, const Precision value, Device)
{
vtkm::worklet::DispatcherMapField<MemSet<Precision>, Device>(MemSet<Precision>(value))
.Invoke(buffer.Buffer);
vtkm::worklet::DispatcherMapField<MemSet<Precision>> memSetDispatcher(
(MemSet<Precision>(value)));
memSetDispatcher.SetDevice(Device());
memSetDispatcher.Invoke(buffer.Buffer);
}
};
}

163
vtkm/rendering/raytracing/ConnectivityTracer.hxx
View File

@ -101,19 +101,24 @@ void RayTracking<FloatType>::Init(const vtkm::Id size,
//
// Set the initial Distances
//
vtkm::worklet::DispatcherMapField<CopyAndOffset<FloatType>, Device>(
CopyAndOffset<FloatType>(0.0f))
.Invoke(distances, *EnterDist);
vtkm::worklet::DispatcherMapField<CopyAndOffset<FloatType>> resetDistancesDispatcher(
CopyAndOffset<FloatType>(0.0f));
resetDistancesDispatcher.SetDevice(Device());
resetDistancesDispatcher.Invoke(distances, *EnterDist);
//
// Init the exit faces. This value is used to load the next cell
// base on the cell and face it left
//
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Int32>, Device>(MemSet<vtkm::Int32>(-1))
.Invoke(ExitFace);
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Int32>> initExitFaceDispatcher(
MemSet<vtkm::Int32>(-1));
initExitFaceDispatcher.SetDevice(Device());
initExitFaceDispatcher.Invoke(ExitFace);
vtkm::worklet::DispatcherMapField<MemSet<FloatType>, Device>(MemSet<FloatType>(-1))
.Invoke(*ExitDist);
vtkm::worklet::DispatcherMapField<MemSet<FloatType>> initExitDistanceDispatcher(
MemSet<FloatType>(-1));
initExitDistanceDispatcher.SetDevice(Device());
initExitDistanceDispatcher.Invoke(*ExitDist);
}
template <typename FloatType>
@ -956,15 +961,16 @@ void ConnectivityTracer<CellType, ConnectivityType>::IntersectCell(
{
using LocateC = LocateCell<CellType, FloatType, MeshConnExec<ConnectivityType, Device>>;
vtkm::cont::Timer<Device> timer;
vtkm::worklet::DispatcherMapField<LocateC, Device>(LocateC(MeshConn))
.Invoke(rays.HitIdx,
this->MeshConn.GetCoordinates(),
rays.Dir,
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.ExitFace,
rays.Status,
rays.Origin);
vtkm::worklet::DispatcherMapField<LocateC> dispatcher((LocateC(MeshConn)));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.HitIdx,
this->MeshConn.GetCoordinates(),
rays.Dir,
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.ExitFace,
rays.Status,
rays.Origin);
if (this->CountRayStatus)
RaysLost = RayOperations::GetStatusCount(rays, RAY_LOST, Device());
@ -978,11 +984,11 @@ void ConnectivityTracer<CellType, ConnectivityType>::AccumulatePathLengths(
detail::RayTracking<FloatType>& tracker,
Device)
{
vtkm::worklet::DispatcherMapField<AddPathLengths<FloatType>, Device>(AddPathLengths<FloatType>())
.Invoke(rays.Status,
*(tracker.EnterDist),
*(tracker.ExitDist),
rays.GetBuffer("path_lengths").Buffer);
vtkm::worklet::DispatcherMapField<AddPathLengths<FloatType>> dispatcher(
(AddPathLengths<FloatType>()));
dispatcher.SetDevice(Device());
dispatcher.Invoke(
rays.Status, *(tracker.EnterDist), *(tracker.ExitDist), rays.GetBuffer("path_lengths").Buffer);
}
template <vtkm::Int32 CellType, typename ConnectivityType>
@ -995,15 +1001,16 @@ void ConnectivityTracer<CellType, ConnectivityType>::FindLostRays(
using RayB = RayBumper<CellType, FloatType, Device, MeshConnExec<ConnectivityType, Device>>;
vtkm::cont::Timer<Device> timer;
vtkm::worklet::DispatcherMapField<RayB, Device>(
RayB(rays.DirX, rays.DirY, rays.DirZ, this->MeshConn))
.Invoke(rays.HitIdx,
this->MeshConn.GetCoordinates(),
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.ExitFace,
rays.Status,
rays.Origin);
vtkm::worklet::DispatcherMapField<RayB> dispatcher(
RayB(rays.DirX, rays.DirY, rays.DirZ, this->MeshConn));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.HitIdx,
this->MeshConn.GetCoordinates(),
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.ExitFace,
rays.Status,
rays.Origin);
this->LostRayTime += timer.GetElapsedTime();
}
@ -1025,38 +1032,40 @@ void ConnectivityTracer<CellType, ConnectivityType>::SampleCells(
if (FieldAssocPoints)
{
vtkm::worklet::DispatcherMapField<SampleP, Device>(
vtkm::worklet::DispatcherMapField<SampleP> dispatcher(
SampleP(this->SampleDistance,
vtkm::Float32(this->ScalarBounds.Min),
vtkm::Float32(this->ScalarBounds.Max),
this->ColorMap,
rays.Buffers.at(0).Buffer,
this->MeshConn))
.Invoke(rays.HitIdx,
this->MeshConn.GetCoordinates(),
this->ScalarField.GetData(),
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.CurrentDistance,
rays.Dir,
rays.Status,
rays.Origin);
this->MeshConn));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.HitIdx,
this->MeshConn.GetCoordinates(),
this->ScalarField.GetData(),
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.CurrentDistance,
rays.Dir,
rays.Status,
rays.Origin);
}
else
{
vtkm::worklet::DispatcherMapField<SampleC, Device>(
vtkm::worklet::DispatcherMapField<SampleC> dispatcher(
SampleC(this->SampleDistance,
vtkm::Float32(this->ScalarBounds.Min),
vtkm::Float32(this->ScalarBounds.Max),
this->ColorMap,
rays.Buffers.at(0).Buffer,
this->MeshConn))
.Invoke(rays.HitIdx,
this->ScalarField.GetData(),
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.CurrentDistance,
rays.Status);
this->MeshConn));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.HitIdx,
this->ScalarField.GetData(),
*(tracker.EnterDist),
*(tracker.ExitDist),
tracker.CurrentDistance,
rays.Status);
}
this->SampleTime += timer.GetElapsedTime();
@ -1075,29 +1084,31 @@ void ConnectivityTracer<CellType, ConnectivityType>::IntegrateCells(
bool divideEmisByAbsorp = false;
vtkm::cont::ArrayHandle<FloatType> absorp = rays.Buffers.at(0).Buffer;
vtkm::cont::ArrayHandle<FloatType> emission = rays.GetBuffer("emission").Buffer;
vtkm::worklet::DispatcherMapField<IntegrateEmission<FloatType>, Device>(
IntegrateEmission<FloatType>(rays.Buffers.at(0).GetNumChannels(), divideEmisByAbsorp))
.Invoke(rays.Status,
*(tracker.EnterDist),
*(tracker.ExitDist),
rays.Distance,
this->ScalarField.GetData(),
this->EmissionField.GetData(),
absorp,
emission,
rays.HitIdx);
vtkm::worklet::DispatcherMapField<IntegrateEmission<FloatType>> dispatcher(
IntegrateEmission<FloatType>(rays.Buffers.at(0).GetNumChannels(), divideEmisByAbsorp));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status,
*(tracker.EnterDist),
*(tracker.ExitDist),
rays.Distance,
this->ScalarField.GetData(),
this->EmissionField.GetData(),
absorp,
emission,
rays.HitIdx);
}
else
{
vtkm::worklet::DispatcherMapField<Integrate<FloatType>, Device>(
Integrate<FloatType>(rays.Buffers.at(0).GetNumChannels()))
.Invoke(rays.Status,
*(tracker.EnterDist),
*(tracker.ExitDist),
rays.Distance,
this->ScalarField.GetData(),
rays.Buffers.at(0).Buffer,
rays.HitIdx);
vtkm::worklet::DispatcherMapField<Integrate<FloatType>> dispatcher(
Integrate<FloatType>(rays.Buffers.at(0).GetNumChannels()));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status,
*(tracker.EnterDist),
*(tracker.ExitDist),
rays.Distance,
this->ScalarField.GetData(),
rays.Buffers.at(0).Buffer,
rays.HitIdx);
}
IntegrateTime += timer.GetElapsedTime();
@ -1136,9 +1147,10 @@ template <typename FloatType, typename Device>
void ConnectivityTracer<CellType, ConnectivityType>::OffsetMinDistances(Ray<FloatType>& rays,
Device)
{
vtkm::worklet::DispatcherMapField<AdvanceRay<FloatType>, Device>(
AdvanceRay<FloatType>(FloatType(0.001)))
.Invoke(rays.Status, rays.MinDistance);
vtkm::worklet::DispatcherMapField<AdvanceRay<FloatType>> dispatcher(
AdvanceRay<FloatType>(FloatType(0.001)));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status, rays.MinDistance);
}
template <vtkm::Int32 CellType, typename ConnectivityType>
@ -1253,9 +1265,10 @@ void ConnectivityTracer<CellType, ConnectivityType>::RenderOnDevice(Ray<FloatTyp
{
vtkm::cont::ArrayHandleCounting<vtkm::Id> pCounter(0, 1, rays.NumRays);
vtkm::worklet::DispatcherMapField<IdentifyMissedRay, Device>(
IdentifyMissedRay(rays.DebugWidth, rays.DebugHeight, this->BackgroundColor))
.Invoke(pCounter, rays.Buffers.at(0).Buffer);
vtkm::worklet::DispatcherMapField<IdentifyMissedRay> dispatcher(
IdentifyMissedRay(rays.DebugWidth, rays.DebugHeight, this->BackgroundColor));
dispatcher.SetDevice(Device());
dispatcher.Invoke(pCounter, rays.Buffers.at(0).Buffer);
}
vtkm::Float64 renderTime = renderTimer.GetElapsedTime();
this->LogTimers();

56
vtkm/rendering/raytracing/MeshConnectivityBuilder.h
View File

@ -571,8 +571,9 @@ public:
// scatter the coonectivity into the original order
vtkm::worklet::DispatcherMapField<WriteFaceConn, DeviceAdapter>(WriteFaceConn())
.Invoke(cellFaceId, this->FaceOffsets, faceConnectivity);
vtkm::worklet::DispatcherMapField<WriteFaceConn> dispatcherWriteFaceConn;
dispatcherWriteFaceConn.SetDevice(DeviceAdapter());
dispatcherWriteFaceConn.Invoke(cellFaceId, this->FaceOffsets, faceConnectivity);
FaceConnectivity = faceConnectivity;
@ -628,8 +629,9 @@ public:
this->ExtractExternalFaces(cellFaceId, faceConnectivity, shapes, conn, shapeOffsets);
// scatter the coonectivity into the original order
vtkm::worklet::DispatcherMapField<WriteFaceConn, DeviceAdapter>(WriteFaceConn())
.Invoke(cellFaceId, this->FaceOffsets, faceConnectivity);
vtkm::worklet::DispatcherMapField<WriteFaceConn> dispatcher{ (WriteFaceConn{}) };
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(cellFaceId, this->FaceOffsets, faceConnectivity);
FaceConnectivity = faceConnectivity;
OutsideTriangles = externalTriangles;
@ -653,10 +655,11 @@ public:
triangles.PrepareForOutput(numFaces * 2, DeviceAdapter());
vtkm::cont::ArrayHandleCounting<vtkm::Id> counting(0, 1, numFaces);
vtkm::worklet::DispatcherMapField<StructuredExternalTriangles, DeviceAdapter>(
vtkm::worklet::DispatcherMapField<StructuredExternalTriangles> dispatcher(
StructuredExternalTriangles(cellSetStructured.PrepareForInput(
DeviceAdapter(), vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell())))
.Invoke(counting, triangles);
DeviceAdapter(), vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell())));
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(counting, triangles);
vtkm::Float64 time = timer.GetElapsedTime();
Logger::GetInstance()->AddLogData("structured_external_faces", time);
@ -701,8 +704,9 @@ protected:
// Count the number of total faces in the cell set
vtkm::cont::ArrayHandle<vtkm::Id> facesPerCell;
vtkm::worklet::DispatcherMapField<CountFaces, DeviceAdapter>(CountFaces())
.Invoke(shapes, facesPerCell);
vtkm::worklet::DispatcherMapField<CountFaces> countFacesDispatcher{ (CountFaces{}) };
countFacesDispatcher.SetDevice(DeviceAdapter());
countFacesDispatcher.Invoke(shapes, facesPerCell);
vtkm::Id totalFaces = 0;
totalFaces =
@ -744,9 +748,10 @@ protected:
vtkm::cont::ArrayHandle<vtkm::UInt32> faceMortonCodes;
cellFaceId.PrepareForOutput(totalFaces, DeviceAdapter());
faceMortonCodes.PrepareForOutput(totalFaces, DeviceAdapter());
vtkm::worklet::DispatcherMapTopology<MortonCodeFace, DeviceAdapter>(
MortonCodeFace(inverseExtent, minPoint))
.Invoke(cellSet, coordinates, cellOffsets, faceMortonCodes, cellFaceId);
vtkm::worklet::DispatcherMapTopology<MortonCodeFace> mortonCodeFaceDispatcher(
MortonCodeFace(inverseExtent, minPoint));
mortonCodeFaceDispatcher.SetDevice(DeviceAdapter());
mortonCodeFaceDispatcher.Invoke(cellSet, coordinates, cellOffsets, faceMortonCodes, cellFaceId);
// Sort the "faces" (i.e., morton codes)
vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>::SortByKey(faceMortonCodes, cellFaceId);
@ -755,11 +760,15 @@ protected:
faceConnectivity.PrepareForOutput(totalFaces, DeviceAdapter());
// Initialize All connecting faces to -1 (connects to nothing)
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Id>, DeviceAdapter>(MemSet<vtkm::Id>(-1))
.Invoke(faceConnectivity);
vtkm::worklet::DispatcherMapField<MemSet<vtkm::Id>> memSetDispatcher(MemSet<vtkm::Id>(-1));
memSetDispatcher.SetDevice(DeviceAdapter());
memSetDispatcher.Invoke(faceConnectivity);
vtkm::worklet::DispatcherMapField<MortonNeighbor, DeviceAdapter>(MortonNeighbor())
.Invoke(faceMortonCodes, cellFaceId, conn, shapes, shapeOffsets, faceConnectivity);
vtkm::worklet::DispatcherMapField<MortonNeighbor> mortonNeighborDispatcher{ (
MortonNeighbor{}) };
mortonNeighborDispatcher.SetDevice(DeviceAdapter());
mortonNeighborDispatcher.Invoke(
faceMortonCodes, cellFaceId, conn, shapes, shapeOffsets, faceConnectivity);
vtkm::Float64 time = timer.GetElapsedTime();
Logger::GetInstance()->AddLogData("gen_face_conn", time);
@ -792,9 +801,10 @@ protected:
trianglesPerExternalFace.PrepareForOutput(numExternalFaces, DeviceAdapter());
vtkm::worklet::DispatcherMapField<CountExternalTriangles, DeviceAdapter>(
CountExternalTriangles())
.Invoke(externalFacePairs, shapes, trianglesPerExternalFace);
vtkm::worklet::DispatcherMapField<CountExternalTriangles> countExternalDispatcher{ (
CountExternalTriangles{}) };
countExternalDispatcher.SetDevice(DeviceAdapter());
countExternalDispatcher.Invoke(externalFacePairs, shapes, trianglesPerExternalFace);
vtkm::cont::ArrayHandle<vtkm::Id> externalTriangleOffsets;
vtkm::cont::DeviceAdapterAlgorithm<DeviceAdapter>::ScanExclusive(trianglesPerExternalFace,
@ -807,9 +817,11 @@ protected:
externalTriangles.PrepareForOutput(totalExternalTriangles, DeviceAdapter());
//count the number triangles in the external faces
vtkm::worklet::DispatcherMapField<ExternalTriangles, DeviceAdapter>(ExternalTriangles())
.Invoke(
externalFacePairs, shapes, shapeOffsets, conn, externalTriangles, externalTriangleOffsets);
vtkm::worklet::DispatcherMapField<ExternalTriangles> externalTriDispatcher{ (
ExternalTriangles{}) };
externalTriDispatcher.SetDevice(DeviceAdapter());
externalTriDispatcher.Invoke(
externalFacePairs, shapes, shapeOffsets, conn, externalTriangles, externalTriangleOffsets);
vtkm::Float64 time = timer.GetElapsedTime();

41
vtkm/rendering/raytracing/RayOperations.cxx
View File

@ -27,41 +27,6 @@ namespace rendering
namespace raytracing
{
template <typename T>
struct MapCanvasFunctor
{
const vtkm::Matrix<vtkm::Float32, 4, 4> Inverse;
const vtkm::Id Width;
const vtkm::Id Height;
const vtkm::cont::ArrayHandle<vtkm::Float32>& DepthBuffer;
Ray<T>& Rays;
vtkm::Vec<vtkm::Float32, 3> Origin;
MapCanvasFunctor(Ray<T>& rays,
const vtkm::Matrix<vtkm::Float32, 4, 4> inverse,
const vtkm::Id width,
const vtkm::Id height,
const vtkm::cont::ArrayHandle<vtkm::Float32>& depthBuffer,
const vtkm::Vec<vtkm::Float32, 3>& origin)
: Inverse(inverse)
, Width(width)
, Height(height)
, DepthBuffer(depthBuffer)
, Rays(rays)
, Origin(origin)
{
}
template <typename Device>
bool operator()(Device)
{
vtkm::worklet::DispatcherMapField<detail::RayMapCanvas, Device>(
detail::RayMapCanvas(Inverse, Width, Height, Origin))
.Invoke(Rays.PixelIdx, Rays.MaxDistance, DepthBuffer);
return true;
}
};
void RayOperations::MapCanvasToRays(Ray<vtkm::Float32>& rays,
const vtkm::rendering::Camera& camera,
const vtkm::rendering::CanvasRayTracer& canvas)
@ -75,9 +40,9 @@ void RayOperations::MapCanvasToRays(Ray<vtkm::Float32>& rays,
if (!valid)
throw vtkm::cont::ErrorBadValue("Inverse Invalid");
MapCanvasFunctor<vtkm::Float32> functor(
rays, inverse, width, height, canvas.GetDepthBuffer(), camera.GetPosition());
vtkm::cont::TryExecute(functor);
vtkm::worklet::DispatcherMapField<detail::RayMapCanvas>(
detail::RayMapCanvas(inverse, width, height, camera.GetPosition()))
.Invoke(rays.PixelIdx, rays.MaxDistance, canvas.GetDepthBuffer());
}
}
}

45
vtkm/rendering/raytracing/RayOperations.h
View File

@ -117,8 +117,10 @@ public:
template <typename Device, typename T>
static void ResetStatus(Ray<T>& rays, vtkm::UInt8 status, Device)
{
vtkm::worklet::DispatcherMapField<MemSet<vtkm::UInt8>, Device>(MemSet<vtkm::UInt8>(status))
.Invoke(rays.Status);
vtkm::worklet::DispatcherMapField<MemSet<vtkm::UInt8>> dispatcher(
(MemSet<vtkm::UInt8>(status)));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status);
}
//
@ -129,8 +131,10 @@ public:
template <typename Device, typename T>
static void UpdateRayStatus(Ray<T>& rays, Device)
{
vtkm::worklet::DispatcherMapField<detail::RayStatusFilter, Device>(detail::RayStatusFilter())
.Invoke(rays.HitIdx, rays.Status);
vtkm::worklet::DispatcherMapField<detail::RayStatusFilter> dispatcher{ (
detail::RayStatusFilter{}) };
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.HitIdx, rays.Status);
}
static void MapCanvasToRays(Ray<vtkm::Float32>& rays,
@ -146,9 +150,10 @@ public:
vtkm::cont::ArrayHandle<vtkm::UInt8> masks;
vtkm::worklet::DispatcherMapField<ManyMask<vtkm::UInt8, 2>, Device>(
ManyMask<vtkm::UInt8, 2>(maskValues))
.Invoke(rays.Status, masks);
vtkm::worklet::DispatcherMapField<ManyMask<vtkm::UInt8, 2>> dispatcher{ (
ManyMask<vtkm::UInt8, 2>{ maskValues }) };
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status, masks);
vtkm::cont::ArrayHandleCast<vtkm::Id, vtkm::cont::ArrayHandle<vtkm::UInt8>> castedMasks(masks);
const vtkm::Id initVal = 0;
vtkm::Id count = vtkm::cont::DeviceAdapterAlgorithm<Device>::Reduce(castedMasks, initVal);
@ -168,8 +173,10 @@ public:
statusUInt8 = static_cast<vtkm::UInt8>(status);
vtkm::cont::ArrayHandle<vtkm::UInt8> masks;
vtkm::worklet::DispatcherMapField<Mask<vtkm::UInt8>, Device>(Mask<vtkm::UInt8>(statusUInt8))
.Invoke(rays.Status, masks);
vtkm::worklet::DispatcherMapField<Mask<vtkm::UInt8>> dispatcher{ (
Mask<vtkm::UInt8>{ statusUInt8 }) };
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status, masks);
vtkm::cont::ArrayHandleCast<vtkm::Id, vtkm::cont::ArrayHandle<vtkm::UInt8>> castedMasks(masks);
const vtkm::Id initVal = 0;
vtkm::Id count = vtkm::cont::DeviceAdapterAlgorithm<Device>::Reduce(castedMasks, initVal);
@ -187,9 +194,10 @@ public:
vtkm::cont::ArrayHandle<vtkm::UInt8> masks;
vtkm::worklet::DispatcherMapField<ManyMask<vtkm::UInt8, 3>, Device>(
ManyMask<vtkm::UInt8, 3>(maskValues))
.Invoke(rays.Status, masks);
vtkm::worklet::DispatcherMapField<ManyMask<vtkm::UInt8, 3>> dispatcher{ (
ManyMask<vtkm::UInt8, 3>{ maskValues }) };
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status, masks);
vtkm::cont::ArrayHandleCast<vtkm::Id, vtkm::cont::ArrayHandle<vtkm::UInt8>> castedMasks(masks);
const vtkm::Id initVal = 0;
vtkm::Id count = vtkm::cont::DeviceAdapterAlgorithm<Device>::Reduce(castedMasks, initVal);
@ -205,8 +213,10 @@ public:
vtkm::UInt8 statusUInt8 = static_cast<vtkm::UInt8>(RAY_ACTIVE);
vtkm::cont::ArrayHandle<vtkm::UInt8> masks;
vtkm::worklet::DispatcherMapField<Mask<vtkm::UInt8>, Device>(Mask<vtkm::UInt8>(statusUInt8))
.Invoke(rays.Status, masks);
vtkm::worklet::DispatcherMapField<Mask<vtkm::UInt8>> dispatcher{ (
Mask<vtkm::UInt8>{ statusUInt8 }) };
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Status, masks);
vtkm::cont::ArrayHandle<T> emptyHandle;
@ -328,9 +338,10 @@ public:
template <typename Device, typename T>
static void CopyDistancesToMin(Ray<T> rays, Device, const T offset = 0.f)
{
vtkm::worklet::DispatcherMapField<CopyAndOffsetMask<T>, Device>(
CopyAndOffsetMask<T>(offset, RAY_EXITED_MESH))
.Invoke(rays.Distance, rays.MinDistance, rays.Status);
vtkm::worklet::DispatcherMapField<CopyAndOffsetMask<T>> dispatcher{ (
CopyAndOffsetMask<T>{ offset, RAY_EXITED_MESH }) };
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Distance, rays.MinDistance, rays.Status);
}
};
}

49
vtkm/rendering/raytracing/RayTracer.cxx
View File

@ -237,22 +237,27 @@ public:
throw vtkm::cont::ErrorBadValue("Field not accociated with cell set or points");
bool isAssocPoints = scalarField.GetAssociation() == vtkm::cont::Field::Association::POINTS;
vtkm::worklet::DispatcherMapField<CalculateNormals, Device>(CalculateNormals(bvh.LeafNodes))
.Invoke(rays.HitIdx, rays.Dir, rays.NormalX, rays.NormalY, rays.NormalZ, coordsHandle);
vtkm::worklet::DispatcherMapField<CalculateNormals> calcNormalsDispatcher(
CalculateNormals(bvh.LeafNodes));
calcNormalsDispatcher.SetDevice(Device());
calcNormalsDispatcher.Invoke(
rays.HitIdx, rays.Dir, rays.NormalX, rays.NormalY, rays.NormalZ, coordsHandle);
if (isAssocPoints)
{
vtkm::worklet::DispatcherMapField<LerpScalar<Precision>, Device>(
vtkm::worklet::DispatcherMapField<LerpScalar<Precision>> lerpScalarDispatcher(
LerpScalar<Precision>(
bvh.LeafNodes, vtkm::Float32(scalarRange.Min), vtkm::Float32(scalarRange.Max)))
.Invoke(rays.HitIdx, rays.U, rays.V, rays.Scalar, scalarField);
bvh.LeafNodes, vtkm::Float32(scalarRange.Min), vtkm::Float32(scalarRange.Max)));
lerpScalarDispatcher.SetDevice(Device());
lerpScalarDispatcher.Invoke(rays.HitIdx, rays.U, rays.V, rays.Scalar, scalarField);
}
else
{
vtkm::worklet::DispatcherMapField<NodalScalar<Precision>, Device>(
vtkm::worklet::DispatcherMapField<NodalScalar<Precision>> nodalScalarDispatcher(
NodalScalar<Precision>(
bvh.LeafNodes, vtkm::Float32(scalarRange.Min), vtkm::Float32(scalarRange.Max)))
.Invoke(rays.HitIdx, rays.Scalar, scalarField);
bvh.LeafNodes, vtkm::Float32(scalarRange.Min), vtkm::Float32(scalarRange.Max)));
nodalScalarDispatcher.SetDevice(Device());
nodalScalarDispatcher.Invoke(rays.HitIdx, rays.Scalar, scalarField);
}
} // Run
@ -373,10 +378,11 @@ public:
vtkm::Vec<vtkm::Float32, 3> scale(2, 2, 2);
vtkm::Vec<vtkm::Float32, 3> lightPosition = camera.GetPosition() + scale * camera.GetUp();
const vtkm::Int32 colorMapSize = vtkm::Int32(colorMap.GetNumberOfValues());
vtkm::worklet::DispatcherMapField<MapScalarToColor, Device>(
MapScalarToColor(
colorMap, colorMapSize, lightPosition, camera.GetPosition(), camera.GetLookAt()))
.Invoke(rays.HitIdx, rays.Scalar, rays.Normal, rays.Intersection, rays.Buffers.at(0).Buffer);
vtkm::worklet::DispatcherMapField<MapScalarToColor> dispatcher(MapScalarToColor(
colorMap, colorMapSize, lightPosition, camera.GetPosition(), camera.GetLookAt()));
dispatcher.SetDevice(Device());
dispatcher.Invoke(
rays.HitIdx, rays.Scalar, rays.Normal, rays.Intersection, rays.Buffers.at(0).Buffer);
}
}; // class SurfaceColor
@ -487,15 +493,16 @@ void RayTracer::RenderOnDevice(Ray<Precision>& rays, Device)
timer.Reset();
// Find the intersection point from hit distance
vtkm::worklet::DispatcherMapField<detail::IntersectionPoint, Device>(
detail::IntersectionPoint())
.Invoke(rays.HitIdx,
rays.Distance,
rays.Dir,
rays.Origin,
rays.IntersectionX,
rays.IntersectionY,
rays.IntersectionZ);
vtkm::worklet::DispatcherMapField<detail::IntersectionPoint> intersectionPointDispatcher{ (
detail::IntersectionPoint{}) };
intersectionPointDispatcher.SetDevice(Device());
intersectionPointDispatcher.Invoke(rays.HitIdx,
rays.Distance,
rays.Dir,
rays.Origin,
rays.IntersectionX,
rays.IntersectionY,
rays.IntersectionZ);
time = timer.GetElapsedTime();
logger->AddLogData("find_point", time);

46
vtkm/rendering/raytracing/TriangleIntersector.h
View File

@ -1049,16 +1049,17 @@ public:
template <typename DynamicCoordType, typename Precision>
VTKM_CONT void run(Ray<Precision>& rays, LinearBVH& bvh, DynamicCoordType coordsHandle)
{
vtkm::worklet::DispatcherMapField<Intersector, Device>(Intersector(false, bvh))
.Invoke(rays.Dir,
rays.Origin,
rays.Distance,
rays.MinDistance,
rays.MaxDistance,
rays.U,
rays.V,
rays.HitIdx,
coordsHandle);
vtkm::worklet::DispatcherMapField<Intersector> dispatcher(Intersector(false, bvh));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Dir,
rays.Origin,
rays.Distance,
rays.MinDistance,
rays.MaxDistance,
rays.U,
rays.V,
rays.HitIdx,
coordsHandle);
}
template <typename DynamicCoordType, typename Precision>
@ -1068,21 +1069,24 @@ public:
bool returnCellIndex)
{
vtkm::worklet::DispatcherMapField<IntersectorHitIndex<Precision>, Device>(
IntersectorHitIndex<Precision>(false, bvh))
.Invoke(rays.Dir,
rays.HitIdx,
coordsHandle,
rays.Distance,
rays.MinDistance,
rays.MaxDistance,
rays.Origin);
vtkm::worklet::DispatcherMapField<IntersectorHitIndex<Precision>> dispatcher(
IntersectorHitIndex<Precision>(false, bvh));
dispatcher.SetDevice(Device());
dispatcher.Invoke(rays.Dir,
rays.HitIdx,
coordsHandle,
rays.Distance,
rays.MinDistance,
rays.MaxDistance,
rays.Origin);
// Normally we return the index of the triangle hit,
// but in some cases we are only interested in the cell
if (returnCellIndex)
{
vtkm::worklet::DispatcherMapField<CellIndexFilter, Device>(CellIndexFilter(bvh))
.Invoke(rays.HitIdx);
vtkm::worklet::DispatcherMapField<CellIndexFilter> cellIndexFilterDispatcher(
(CellIndexFilter(bvh)));
cellIndexFilterDispatcher.SetDevice(Device());
cellIndexFilterDispatcher.Invoke(rays.HitIdx);
}
// Update ray status
RayOperations::UpdateRayStatus(rays, Device());

77
vtkm/rendering/raytracing/VolumeRendererStructured.cxx
View File

@ -818,8 +818,11 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
}
vtkm::cont::Timer<Device> timer;
vtkm::worklet::DispatcherMapField<CalcRayStart, Device>(CalcRayStart(this->SpatialExtent))
.Invoke(rays.Dir, rays.MinDistance, rays.Distance, rays.MaxDistance, rays.Origin);
vtkm::worklet::DispatcherMapField<CalcRayStart> calcRayStartDispatcher(
CalcRayStart(this->SpatialExtent));
calcRayStartDispatcher.SetDevice(Device());
calcRayStartDispatcher.Invoke(
rays.Dir, rays.MinDistance, rays.Distance, rays.MaxDistance, rays.Origin);
vtkm::Float64 time = timer.GetElapsedTime();
logger->AddLogData("calc_ray_start", time);
@ -840,18 +843,19 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
if (isAssocPoints)
{
vtkm::worklet::DispatcherMapField<Sampler<Device, UniformLocator<Device>>, Device>(
vtkm::worklet::DispatcherMapField<Sampler<Device, UniformLocator<Device>>> samplerDispatcher(
Sampler<Device, UniformLocator<Device>>(ColorMap,
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator))
.Invoke(rays.Dir,
rays.Origin,
rays.MinDistance,
rays.MaxDistance,
rays.Buffers.at(0).Buffer,
*ScalarField);
locator));
samplerDispatcher.SetDevice(Device());
samplerDispatcher.Invoke(rays.Dir,
rays.Origin,
rays.MinDistance,
rays.MaxDistance,
rays.Buffers.at(0).Buffer,
*ScalarField);
}
else
{
@ -876,34 +880,37 @@ void VolumeRendererStructured::RenderOnDevice(vtkm::rendering::raytracing::Ray<P
RectilinearLocator<Device> locator(vertices, Cellset);
if (isAssocPoints)
{
vtkm::worklet::DispatcherMapField<Sampler<Device, RectilinearLocator<Device>>, Device>(
Sampler<Device, RectilinearLocator<Device>>(ColorMap,
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator))
.Invoke(rays.Dir,
rays.Origin,
rays.MinDistance,
rays.MaxDistance,
rays.Buffers.at(0).Buffer,
*ScalarField);
vtkm::worklet::DispatcherMapField<Sampler<Device, RectilinearLocator<Device>>>
samplerDispatcher(
Sampler<Device, RectilinearLocator<Device>>(ColorMap,
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator));
samplerDispatcher.SetDevice(Device());
samplerDispatcher.Invoke(rays.Dir,
rays.Origin,
rays.MinDistance,
rays.MaxDistance,
rays.Buffers.at(0).Buffer,
*ScalarField);
}
else
{
vtkm::worklet::DispatcherMapField<SamplerCellAssoc<Device, RectilinearLocator<Device>>,
Device>(
SamplerCellAssoc<Device, RectilinearLocator<Device>>(ColorMap,
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator))
.Invoke(rays.Dir,
rays.Origin,
rays.MinDistance,
rays.MaxDistance,
rays.Buffers.at(0).Buffer,
*ScalarField);
vtkm::worklet::DispatcherMapField<SamplerCellAssoc<Device, RectilinearLocator<Device>>>
rectilinearLocatorDispatcher(
SamplerCellAssoc<Device, RectilinearLocator<Device>>(ColorMap,
vtkm::Float32(ScalarRange.Min),
vtkm::Float32(ScalarRange.Max),
SampleDistance,
locator));
rectilinearLocatorDispatcher.SetDevice(Device());
rectilinearLocatorDispatcher.Invoke(rays.Dir,
rays.Origin,
rays.MinDistance,
rays.MaxDistance,
rays.Buffers.at(0).Buffer,
*ScalarField);
}
}

8
vtkm/worklet/AverageByKey.h
View File

@ -78,7 +78,8 @@ struct AverageByKey
{
VTKM_IS_DEVICE_ADAPTER_TAG(Device);
vtkm::worklet::DispatcherReduceByKey<AverageWorklet, Device> dispatcher;
vtkm::worklet::DispatcherReduceByKey<AverageWorklet> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(keys, inValues, outAverages);
}
@ -175,8 +176,9 @@ struct AverageByKey
keyArraySorted, inputZipHandle, outputKeyArray, outputZipHandle, vtkm::Add());
// get average
DispatcherMapField<DivideWorklet, DeviceAdapter>().Invoke(
sumArray, countArray, outputValueArray);
DispatcherMapField<DivideWorklet> dispatcher;
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(sumArray, countArray, outputValueArray);
}
};
}

6
vtkm/worklet/CellDeepCopy.h
View File

@ -89,7 +89,8 @@ struct CellDeepCopy
vtkm::cont::ArrayHandle<vtkm::IdComponent, NumIndicesStorage> numIndices;
vtkm::worklet::DispatcherMapTopology<CountCellPoints, Device> countDispatcher;
vtkm::worklet::DispatcherMapTopology<CountCellPoints> countDispatcher;
countDispatcher.SetDevice(Device());
countDispatcher.Invoke(inCellSet, numIndices);
vtkm::cont::ArrayHandle<vtkm::UInt8, ShapeStorage> shapes;
@ -100,7 +101,8 @@ struct CellDeepCopy
vtkm::cont::ConvertNumComponentsToOffsets(numIndices, offsets, connectivitySize);
connectivity.Allocate(connectivitySize);
vtkm::worklet::DispatcherMapTopology<PassCellStructure, Device> passDispatcher;
vtkm::worklet::DispatcherMapTopology<PassCellStructure> passDispatcher;
passDispatcher.SetDevice(Device());
passDispatcher.Invoke(
inCellSet, shapes, vtkm::cont::make_ArrayHandleGroupVecVariable(connectivity, offsets));

25
vtkm/worklet/Clip.h
View File

@ -496,8 +496,9 @@ public:
vtkm::cont::ArrayHandle<ClipStats> stats;
ComputeStats<DeviceAdapter> computeStats(value, clipTablesDevicePortal, invert);
DispatcherMapTopology<ComputeStats<DeviceAdapter>, DeviceAdapter>(computeStats)
.Invoke(cellSet, scalars, clipTableIdxs, stats);
DispatcherMapTopology<ComputeStats<DeviceAdapter>> computeStatsDispatcher(computeStats);
computeStatsDispatcher.SetDevice(DeviceAdapter());
computeStatsDispatcher.Invoke(cellSet, scalars, clipTableIdxs, stats);
// compute offsets for each invocation
ClipStats zero;
@ -522,15 +523,17 @@ public:
this->CellIdMap.Allocate(total.NumberOfCells);
GenerateCellSet<DeviceAdapter> generateCellSet(value, clipTablesDevicePortal);
DispatcherMapTopology<GenerateCellSet<DeviceAdapter>, DeviceAdapter>(generateCellSet)
.Invoke(cellSet,
scalars,
clipTableIdxs,
cellSetIndices,
outConnectivity,
newPoints,
newPointsConnectivityReverseMap,
this->CellIdMap);
DispatcherMapTopology<GenerateCellSet<DeviceAdapter>> generateCellSetDispatcher(
generateCellSet);
generateCellSetDispatcher.SetDevice(DeviceAdapter());
generateCellSetDispatcher.Invoke(cellSet,
scalars,
clipTableIdxs,
cellSetIndices,
outConnectivity,
newPoints,
newPointsConnectivityReverseMap,
this->CellIdMap);
cellSetIndices.ReleaseResources();
// Step 3. remove duplicates from the list of new points

12
vtkm/worklet/CoordinateSystemTransform.h
View File

@ -192,12 +192,14 @@ public:
{
if (CartesianToSpherical)
{
vtkm::worklet::DispatcherMapField<detail::CarToSphere<T>, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<detail::CarToSphere<T>> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(inPoints, outPoints);
}
else
{
vtkm::worklet::DispatcherMapField<detail::SphereToCar<T>, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<detail::SphereToCar<T>> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(inPoints, outPoints);
}
}
@ -209,12 +211,14 @@ public:
{
if (CartesianToSpherical)
{
vtkm::worklet::DispatcherMapField<detail::CarToSphere<T>, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<detail::CarToSphere<T>> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(inPoints, outPoints);
}
else
{
vtkm::worklet::DispatcherMapField<detail::SphereToCar<T>, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherMapField<detail::SphereToCar<T>> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
dispatcher.Invoke(inPoints, outPoints);
}
}

14
vtkm/worklet/DispatcherMapField.h
View File

@ -20,7 +20,6 @@
#ifndef vtk_m_worklet_Dispatcher_MapField_h
#define vtk_m_worklet_Dispatcher_MapField_h
#include <vtkm/cont/DeviceAdapter.h>
#include <vtkm/worklet/WorkletMapField.h>
#include <vtkm/worklet/internal/DispatcherBase.h>
@ -31,16 +30,15 @@ namespace worklet
/// \brief Dispatcher for worklets that inherit from \c WorkletMapField.
///
template <typename WorkletType, typename Device = VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
template <typename WorkletType>
class DispatcherMapField
: public vtkm::worklet::internal::DispatcherBase<DispatcherMapField<WorkletType, Device>,
: public vtkm::worklet::internal::DispatcherBase<DispatcherMapField<WorkletType>,
WorkletType,
vtkm::worklet::WorkletMapField>
{
using Superclass =
vtkm::worklet::internal::DispatcherBase<DispatcherMapField<WorkletType, Device>,
WorkletType,
vtkm::worklet::WorkletMapField>;
using Superclass = vtkm::worklet::internal::DispatcherBase<DispatcherMapField<WorkletType>,
WorkletType,
vtkm::worklet::WorkletMapField>;
using ScatterType = typename Superclass::ScatterType;
public:
@ -80,7 +78,7 @@ public:
// A MapField is a pretty straightforward dispatch. Once we know the number
// of invocations, the superclass can take care of the rest.
this->BasicInvoke(invocation, numInstances, Device());
this->BasicInvoke(invocation, numInstances);
}
};
}

9
vtkm/worklet/DispatcherMapTopology.h
View File

@ -32,14 +32,14 @@ namespace worklet
/// \brief Dispatcher for worklets that inherit from \c WorkletMapTopology.
///
template <typename WorkletType, typename Device = VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
template <typename WorkletType>
class DispatcherMapTopology
: public vtkm::worklet::internal::DispatcherBase<DispatcherMapTopology<WorkletType, Device>,
: public vtkm::worklet::internal::DispatcherBase<DispatcherMapTopology<WorkletType>,
WorkletType,
vtkm::worklet::detail::WorkletMapTopologyBase>
{
using Superclass =
vtkm::worklet::internal::DispatcherBase<DispatcherMapTopology<WorkletType, Device>,
vtkm::worklet::internal::DispatcherBase<DispatcherMapTopology<WorkletType>,
WorkletType,
vtkm::worklet::detail::WorkletMapTopologyBase>;
using ScatterType = typename Superclass::ScatterType;
@ -81,8 +81,7 @@ public:
// Now that we have the input domain, we can extract the range of the
// scheduling and call BadicInvoke.
this->BasicInvoke(
invocation, internal::scheduling_range(inputDomain, SchedulingRangeType{}), Device());
this->BasicInvoke(invocation, internal::scheduling_range(inputDomain, SchedulingRangeType{}));
}
};
}

8
vtkm/worklet/DispatcherPointNeighborhood.h
View File

@ -33,14 +33,14 @@ namespace worklet
/// \brief Dispatcher for worklets that inherit from \c WorkletPointNeighborhood.
///
template <typename WorkletType, typename Device = VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
template <typename WorkletType>
class DispatcherPointNeighborhood
: public vtkm::worklet::internal::DispatcherBase<DispatcherPointNeighborhood<WorkletType, Device>,
: public vtkm::worklet::internal::DispatcherBase<DispatcherPointNeighborhood<WorkletType>,
WorkletType,
vtkm::worklet::WorkletPointNeighborhoodBase>
{
using Superclass =
vtkm::worklet::internal::DispatcherBase<DispatcherPointNeighborhood<WorkletType, Device>,
vtkm::worklet::internal::DispatcherBase<DispatcherPointNeighborhood<WorkletType>,
WorkletType,
vtkm::worklet::WorkletPointNeighborhoodBase>;
using ScatterType = typename Superclass::ScatterType;
@ -82,7 +82,7 @@ public:
// This is pretty straightforward dispatch. Once we know the number
// of invocations, the superclass can take care of the rest.
this->BasicInvoke(invocation, inputRange, Device());
this->BasicInvoke(invocation, inputRange);
}
};
}

13
vtkm/worklet/DispatcherReduceByKey.h
View File

@ -33,16 +33,15 @@ namespace worklet
/// \brief Dispatcher for worklets that inherit from \c WorkletReduceByKey.
///
template <typename WorkletType, typename Device = VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
template <typename WorkletType>
class DispatcherReduceByKey
: public vtkm::worklet::internal::DispatcherBase<DispatcherReduceByKey<WorkletType, Device>,
: public vtkm::worklet::internal::DispatcherBase<DispatcherReduceByKey<WorkletType>,
WorkletType,
vtkm::worklet::WorkletReduceByKey>
{
using Superclass =
vtkm::worklet::internal::DispatcherBase<DispatcherReduceByKey<WorkletType, Device>,
WorkletType,
vtkm::worklet::WorkletReduceByKey>;
using Superclass = vtkm::worklet::internal::DispatcherBase<DispatcherReduceByKey<WorkletType>,
WorkletType,
vtkm::worklet::WorkletReduceByKey>;
using ScatterType = typename Superclass::ScatterType;
public:
@ -83,7 +82,7 @@ public:
// Now that we have the input domain, we can extract the range of the
// scheduling and call BadicInvoke.
this->BasicInvoke(invocation, inputDomain.GetInputRange(), Device());
this->BasicInvoke(invocation, inputDomain.GetInputRange());
}
};
}

55
vtkm/worklet/DispatcherStreamingMapField.h
View File

@ -33,7 +33,22 @@ namespace worklet
namespace detail
{
template <typename ControlInterface, typename Device>
struct DispatcherStreamingTryExecuteFunctor
{
template <typename Device, typename DispatcherBaseType, typename Invocation, typename RangeType>
VTKM_CONT bool operator()(Device device,
const DispatcherBaseType* self,
Invocation& invocation,
const RangeType& dimensions,
const RangeType& globalIndexOffset)
{
self->InvokeTransportParameters(
invocation, dimensions, globalIndexOffset, self->Scatter.GetOutputRange(dimensions), device);
return true;
}
};
template <typename ControlInterface>
struct DispatcherStreamingMapFieldTransformFunctor
{
vtkm::Id BlockIndex;
@ -116,7 +131,7 @@ struct DispatcherStreamingMapFieldTransformFunctor
}
};
template <typename ControlInterface, typename Device>
template <typename ControlInterface>
struct DispatcherStreamingMapFieldTransferFunctor
{
VTKM_CONT
@ -163,14 +178,14 @@ struct DispatcherStreamingMapFieldTransferFunctor
/// \brief Dispatcher for worklets that inherit from \c WorkletMapField.
///
template <typename WorkletType, typename Device = VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
template <typename WorkletType>
class DispatcherStreamingMapField
: public vtkm::worklet::internal::DispatcherBase<DispatcherStreamingMapField<WorkletType, Device>,
: public vtkm::worklet::internal::DispatcherBase<DispatcherStreamingMapField<WorkletType>,
WorkletType,
vtkm::worklet::WorkletMapField>
{
using Superclass =
vtkm::worklet::internal::DispatcherBase<DispatcherStreamingMapField<WorkletType, Device>,
vtkm::worklet::internal::DispatcherBase<DispatcherStreamingMapField<WorkletType>,
WorkletType,
vtkm::worklet::WorkletMapField>;
using ScatterType = typename Superclass::ScatterType;
@ -199,17 +214,23 @@ public:
VTKM_CONT
void SetNumberOfBlocks(vtkm::Id numberOfBlocks) { NumberOfBlocks = numberOfBlocks; }
template <typename Invocation, typename DeviceAdapter>
friend struct detail::DispatcherStreamingTryExecuteFunctor;
template <typename Invocation>
VTKM_CONT void BasicInvoke(Invocation& invocation,
vtkm::Id numInstances,
vtkm::Id globalIndexOffset,
DeviceAdapter device) const
vtkm::Id globalIndexOffset) const
{
this->InvokeTransportParameters(invocation,
numInstances,
globalIndexOffset,
this->Scatter.GetOutputRange(numInstances),
device);
bool success = vtkm::cont::TryExecuteOnDevice(this->GetDevice(),
detail::DispatcherStreamingTryExecuteFunctor(),
this,
invocation,
numInstances,
globalIndexOffset);
if (!success)
{
throw vtkm::cont::ErrorExecution("Failed to execute worklet on any device.");
}
}
template <typename Invocation>
@ -232,11 +253,9 @@ public:
blockSize += 1;
using TransformFunctorType =
detail::DispatcherStreamingMapFieldTransformFunctor<typename Invocation::ControlInterface,
Device>;
detail::DispatcherStreamingMapFieldTransformFunctor<typename Invocation::ControlInterface>;
using TransferFunctorType =
detail::DispatcherStreamingMapFieldTransferFunctor<typename Invocation::ControlInterface,
Device>;
detail::DispatcherStreamingMapFieldTransferFunctor<typename Invocation::ControlInterface>;
for (vtkm::Id block = 0; block < NumberOfBlocks; block++)
{
@ -255,7 +274,7 @@ public:
using ChangedType = typename Invocation::template ChangeParametersType<ReportedType>::type;
ChangedType changedParams = invocation.ChangeParameters(newParams);
this->BasicInvoke(changedParams, numberOfInstances, globalIndexOffset, Device());
this->BasicInvoke(changedParams, numberOfInstances, globalIndexOffset);
// Loop over parameters again to sync results for this block into control array
using ParameterInterfaceType2 = typename ChangedType::ParameterInterface;

35
vtkm/worklet/ExternalFaces.h
View File

@ -737,8 +737,9 @@ public:
// Create a worklet to count the number of external faces on each cell
vtkm::cont::ArrayHandle<vtkm::IdComponent> numExternalFaces;
vtkm::worklet::DispatcherMapTopology<NumExternalFacesPerStructuredCell, DeviceAdapter>
vtkm::worklet::DispatcherMapTopology<NumExternalFacesPerStructuredCell>
numExternalFacesDispatcher((NumExternalFacesPerStructuredCell(MinPoint, MaxPoint)));
numExternalFacesDispatcher.SetDevice(DeviceAdapter());
numExternalFacesDispatcher.Invoke(inCellSet, numExternalFaces, coordData);
@ -761,9 +762,10 @@ public:
// information to.
faceConnectivity.Allocate(connectivitySize);
vtkm::worklet::DispatcherMapTopology<BuildConnectivityStructured, DeviceAdapter>
vtkm::worklet::DispatcherMapTopology<BuildConnectivityStructured>
buildConnectivityStructuredDispatcher(BuildConnectivityStructured(MinPoint, MaxPoint),
scatterCellToExternalFace);
buildConnectivityStructuredDispatcher.SetDevice(DeviceAdapter());
buildConnectivityStructuredDispatcher.Invoke(
inCellSet,
@ -797,7 +799,8 @@ public:
//Create a worklet to map the number of faces to each cell
vtkm::cont::ArrayHandle<vtkm::IdComponent> facesPerCell;
vtkm::worklet::DispatcherMapTopology<NumFacesPerCell, DeviceAdapter> numFacesDispatcher;
vtkm::worklet::DispatcherMapTopology<NumFacesPerCell> numFacesDispatcher;
numFacesDispatcher.SetDevice(DeviceAdapter());
numFacesDispatcher.Invoke(inCellSet, facesPerCell);
@ -813,7 +816,8 @@ public:
if (this->PassPolyData)
{
vtkm::cont::ArrayHandle<vtkm::IdComponent> isPolyDataCell;
vtkm::worklet::DispatcherMapTopology<IsPolyDataCell, DeviceAdapter> isPolyDataCellDispatcher;
vtkm::worklet::DispatcherMapTopology<IsPolyDataCell> isPolyDataCellDispatcher;
isPolyDataCellDispatcher.SetDevice(DeviceAdapter());
isPolyDataCellDispatcher.Invoke(inCellSet, isPolyDataCell);
@ -823,16 +827,18 @@ public:
if (scatterPolyDataCells.GetOutputRange(inCellSet.GetNumberOfCells()) != 0)
{
vtkm::worklet::DispatcherMapTopology<CountPolyDataCellPoints, DeviceAdapter>
vtkm::worklet::DispatcherMapTopology<CountPolyDataCellPoints>
countPolyDataCellPointsDispatcher(scatterPolyDataCells);
countPolyDataCellPointsDispatcher.SetDevice(DeviceAdapter());
countPolyDataCellPointsDispatcher.Invoke(inCellSet, polyDataPointCount);
vtkm::cont::ConvertNumComponentsToOffsets(
polyDataPointCount, polyDataOffsets, polyDataConnectivitySize);
vtkm::worklet::DispatcherMapTopology<PassPolyDataCells, DeviceAdapter>
passPolyDataCellsDispatcher(scatterPolyDataCells);
vtkm::worklet::DispatcherMapTopology<PassPolyDataCells> passPolyDataCellsDispatcher(
scatterPolyDataCells);
passPolyDataCellsDispatcher.SetDevice(DeviceAdapter());
polyDataConnectivity.Allocate(polyDataConnectivitySize);
@ -869,23 +875,25 @@ public:
vtkm::cont::ArrayHandle<vtkm::HashType> faceHashes;
vtkm::cont::ArrayHandle<vtkm::Id> originCells;
vtkm::cont::ArrayHandle<vtkm::IdComponent> originFaces;
vtkm::worklet::DispatcherMapTopology<FaceHash, DeviceAdapter> faceHashDispatcher(
scatterCellToFace);
vtkm::worklet::DispatcherMapTopology<FaceHash> faceHashDispatcher(scatterCellToFace);
faceHashDispatcher.SetDevice(DeviceAdapter());
faceHashDispatcher.Invoke(inCellSet, faceHashes, originCells, originFaces);
vtkm::worklet::Keys<vtkm::HashType> faceKeys(faceHashes, DeviceAdapter());
vtkm::cont::ArrayHandle<vtkm::IdComponent> faceOutputCount;
vtkm::worklet::DispatcherReduceByKey<FaceCounts, DeviceAdapter> faceCountDispatcher;
vtkm::worklet::DispatcherReduceByKey<FaceCounts> faceCountDispatcher;
faceCountDispatcher.SetDevice(DeviceAdapter());
faceCountDispatcher.Invoke(faceKeys, inCellSet, originCells, originFaces, faceOutputCount);
auto scatterCullInternalFaces = NumPointsPerFace::MakeScatter(faceOutputCount, DeviceAdapter());
PointCountArrayType facePointCount;
vtkm::worklet::DispatcherReduceByKey<NumPointsPerFace, DeviceAdapter> pointsPerFaceDispatcher(
vtkm::worklet::DispatcherReduceByKey<NumPointsPerFace> pointsPerFaceDispatcher(
scatterCullInternalFaces);
pointsPerFaceDispatcher.SetDevice(DeviceAdapter());
pointsPerFaceDispatcher.Invoke(faceKeys, inCellSet, originCells, originFaces, facePointCount);
@ -900,8 +908,9 @@ public:
// information to.
faceConnectivity.Allocate(connectivitySize);
vtkm::worklet::DispatcherReduceByKey<BuildConnectivity, DeviceAdapter>
buildConnectivityDispatcher(scatterCullInternalFaces);
vtkm::worklet::DispatcherReduceByKey<BuildConnectivity> buildConnectivityDispatcher(
scatterCullInternalFaces);
buildConnectivityDispatcher.SetDevice(DeviceAdapter());
vtkm::cont::ArrayHandle<vtkm::Id> faceToCellIdMap;

3
vtkm/worklet/ExtractGeometry.h
View File

@ -169,7 +169,8 @@ public:
extractInside,
extractBoundaryCells,
extractOnlyBoundaryCells);
DispatcherMapTopology<ExtractCellsByVOI, DeviceAdapter> dispatcher(worklet);
DispatcherMapTopology<ExtractCellsByVOI> dispatcher(worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(cellSet, coordinates, passFlags);
vtkm::cont::ArrayHandleCounting<vtkm::Id> indices =

3
vtkm/worklet/ExtractPoints.h
View File

@ -107,7 +107,8 @@ public:
vtkm::cont::ArrayHandle<bool> passFlags;
ExtractPointsByVOI worklet(implicitFunction.PrepareForExecution(device), extractInside);
DispatcherMapTopology<ExtractPointsByVOI, DeviceAdapter> dispatcher(worklet);
DispatcherMapTopology<ExtractPointsByVOI> dispatcher(worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(cellSet, coordinates, passFlags);
vtkm::cont::ArrayHandleCounting<vtkm::Id> indices =

5
vtkm/worklet/FieldEntropy.h
View File

@ -91,8 +91,9 @@ public:
///// calculate information content of each bin using self-define worklet /////
vtkm::cont::ArrayHandle<vtkm::Float64> informationContent;
SetBinInformationContent binWorklet(static_cast<vtkm::Float64>(freqSum));
vtkm::worklet::DispatcherMapField<SetBinInformationContent, DeviceAdapter>
setBinInformationContentDispatcher(binWorklet);
vtkm::worklet::DispatcherMapField<SetBinInformationContent> setBinInformationContentDispatcher(
binWorklet);
setBinInformationContentDispatcher.SetDevice(DeviceAdapter());
setBinInformationContentDispatcher.Invoke(binArray, informationContent);
///// calculate entropy by summing up information conetent of all bins /////

8
vtkm/worklet/FieldHistogram.h
View File

@ -166,8 +166,9 @@ public:
// Worklet to set the bin number for each data value
SetHistogramBin<FieldType> binWorklet(numberOfBins, fieldMinValue, fieldDelta);
vtkm::worklet::DispatcherMapField<SetHistogramBin<FieldType>, DeviceAdapter>
setHistogramBinDispatcher(binWorklet);
vtkm::worklet::DispatcherMapField<SetHistogramBin<FieldType>> setHistogramBinDispatcher(
binWorklet);
setHistogramBinDispatcher.SetDevice(DeviceAdapter());
setHistogramBinDispatcher.Invoke(fieldArray, binIndex);
// Sort the resulting bin array for counting
@ -179,7 +180,8 @@ public:
DeviceAlgorithms::UpperBounds(binIndex, binCounter, totalCount);
// Difference between adjacent items is the bin count
vtkm::worklet::DispatcherMapField<AdjacentDifference, DeviceAdapter> dispatcher;
vtkm::worklet::DispatcherMapField<AdjacentDifference> dispatcher;
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(binCounter, totalCount, binArray);
//update the users data

6
vtkm/worklet/FieldStatistics.h
View File

@ -152,8 +152,9 @@ public:
pow4Array.Allocate(dataSize);
// Raw moments via Worklet
vtkm::worklet::DispatcherMapField<CalculatePowers, DeviceAdapter> calculatePowersDispatcher(
vtkm::worklet::DispatcherMapField<CalculatePowers> calculatePowersDispatcher(
CalculatePowers(4));
calculatePowersDispatcher.SetDevice(DeviceAdapter());
calculatePowersDispatcher.Invoke(fieldArray, pow1Array, pow2Array, pow3Array, pow4Array);
// Accumulate the results using ScanInclusive
@ -163,8 +164,9 @@ public:
statinfo.rawMoment[FOURTH] = DeviceAlgorithms::ScanInclusive(pow4Array, pow4Array) / numValues;
// Subtract the mean from every value and leave in tempArray
vtkm::worklet::DispatcherMapField<SubtractConst, DeviceAdapter> subtractConstDispatcher(
vtkm::worklet::DispatcherMapField<SubtractConst> subtractConstDispatcher(
SubtractConst(statinfo.mean));
subtractConstDispatcher.SetDevice(DeviceAdapter());
subtractConstDispatcher.Invoke(fieldArray, tempArray);
// Calculate sums of powers on the (value - mean) array

19
vtkm/worklet/Gradient.h
View File

@ -60,7 +60,8 @@ struct DeducedPointGrad
template <typename CellSetType>
void operator()(const CellSetType& cellset) const
{
vtkm::worklet::DispatcherMapTopology<PointGradient<T>, Device> dispatcher;
vtkm::worklet::DispatcherMapTopology<PointGradient<T>> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(cellset, //topology to iterate on a per point basis
cellset, //whole cellset in
*this->Points,
@ -70,7 +71,8 @@ struct DeducedPointGrad
void operator()(const vtkm::cont::CellSetStructured<3>& cellset) const
{
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>, Device> dispatcher;
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(cellset, //topology to iterate on a per point basis
*this->Points,
*this->Field,
@ -81,7 +83,8 @@ struct DeducedPointGrad
void operator()(const vtkm::cont::CellSetPermutation<vtkm::cont::CellSetStructured<3>,
PermIterType>& cellset) const
{
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>, Device> dispatcher;
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(cellset, //topology to iterate on a per point basis
*this->Points,
*this->Field,
@ -90,7 +93,8 @@ struct DeducedPointGrad
void operator()(const vtkm::cont::CellSetStructured<2>& cellset) const
{
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>, Device> dispatcher;
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(cellset, //topology to iterate on a per point basis
*this->Points,
*this->Field,
@ -101,7 +105,8 @@ struct DeducedPointGrad
void operator()(const vtkm::cont::CellSetPermutation<vtkm::cont::CellSetStructured<2>,
PermIterType>& cellset) const
{
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>, Device> dispatcher;
vtkm::worklet::DispatcherPointNeighborhood<StructuredPointGradient<T>> dispatcher;
dispatcher.SetDevice(Device());
dispatcher.Invoke(cellset, //topology to iterate on a per point basis
*this->Points,
*this->Field,
@ -273,11 +278,11 @@ public:
DeviceAdapter)
{
using DispatcherType =
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::gradient::CellGradient<T>, DeviceAdapter>;
vtkm::worklet::DispatcherMapTopology<vtkm::worklet::gradient::CellGradient<T>>;
vtkm::worklet::gradient::CellGradient<T> worklet;
DispatcherType dispatcher(worklet);
dispatcher.SetDevice(DeviceAdapter());
dispatcher.Invoke(cells, coords, field, extraOutput);
return extraOutput.Gradient;

17
vtkm/worklet/KernelSplatter.h
View File

@ -431,8 +431,8 @@ struct KernelSplatterFilterUniformGrid
IdHandleType localNeighborIds;
GetFootprint footprint_worklet(origin_, spacing_, pointDimensions, kernel_);
vtkm::worklet::DispatcherMapField<GetFootprint, DeviceAdapter> footprintDispatcher(
footprint_worklet);
vtkm::worklet::DispatcherMapField<GetFootprint> footprintDispatcher(footprint_worklet);
footprintDispatcher.SetDevice(DeviceAdapter());
START_TIMER_BLOCK(GetFootprint)
footprintDispatcher.Invoke(
@ -495,7 +495,8 @@ struct KernelSplatterFilterUniformGrid
IdPermType offsets(neighbor2SplatId, numNeighborsExclusiveSum);
debug::OutputArrayDebug(offsets, "offsets");
vtkm::worklet::DispatcherMapField<ComputeLocalNeighborId, DeviceAdapter> idDispatcher;
vtkm::worklet::DispatcherMapField<ComputeLocalNeighborId> idDispatcher;
idDispatcher.SetDevice(DeviceAdapter());
START_TIMER_BLOCK(idDispatcher)
idDispatcher.Invoke(modulii, offsets, localNeighborIds);
END_TIMER_BLOCK(idDispatcher)
@ -528,8 +529,8 @@ struct KernelSplatterFilterUniformGrid
FloatHandleType splatValues;
GetSplatValue splatterDispatcher_worklet(origin_, spacing_, pointDimensions, kernel_);
vtkm::worklet::DispatcherMapField<GetSplatValue, DeviceAdapter> splatterDispatcher(
splatterDispatcher_worklet);
vtkm::worklet::DispatcherMapField<GetSplatValue> splatterDispatcher(splatterDispatcher_worklet);
splatterDispatcher.SetDevice(DeviceAdapter());
START_TIMER_BLOCK(GetSplatValue)
splatterDispatcher.Invoke(ptSplatPoints,
@ -582,7 +583,8 @@ struct KernelSplatterFilterUniformGrid
// initialize each field value to zero to begin with
//---------------------------------------------------------------
IdCountingType indexArray(vtkm::Id(0), 1, numVolumePoints);
vtkm::worklet::DispatcherMapField<zero_voxel, DeviceAdapter> zeroDispatcher;
vtkm::worklet::DispatcherMapField<zero_voxel> zeroDispatcher;
zeroDispatcher.SetDevice(DeviceAdapter());
zeroDispatcher.Invoke(indexArray, scalarSplatOutput);
//
indexArray.ReleaseResources();
@ -591,7 +593,8 @@ struct KernelSplatterFilterUniformGrid
// Scatter operation to write the previously-computed splat
// value sums into their corresponding entries in the output array
//---------------------------------------------------------------
vtkm::worklet::DispatcherMapField<UpdateVoxelSplats, DeviceAdapter> scatterDispatcher;
vtkm::worklet::DispatcherMapField<UpdateVoxelSplats> scatterDispatcher;
scatterDispatcher.SetDevice(DeviceAdapter());
START_TIMER_BLOCK(UpdateVoxelSplats)
scatterDispatcher.Invoke(uniqueVoxelIds, voxelSplatSums, scalarSplatOutput);

32
vtkm/worklet/MarchingCubes.h
View File

@ -469,7 +469,8 @@ void MergeDuplicates(const vtkm::cont::ArrayHandle<KeyType, KeyStorage>& origina
input_keys.ReleaseResources();
{
vtkm::worklet::DispatcherReduceByKey<MergeDuplicateValues, DeviceAdapterTag> dispatcher;
vtkm::worklet::DispatcherReduceByKey<MergeDuplicateValues> dispatcher;
dispatcher.SetDevice(DeviceAdapterTag());
vtkm::cont::ArrayHandle<vtkm::Id> writeCells;
vtkm::cont::ArrayHandle<vtkm::FloatDefault> writeWeights;
dispatcher.Invoke(keys, weights, cellids, writeWeights, writeCells);
@ -483,7 +484,8 @@ void MergeDuplicates(const vtkm::cont::ArrayHandle<KeyType, KeyStorage>& origina
uniqueKeys, original_keys, connectivity, marchingcubes::MultiContourLess());
//update the edge ids
vtkm::worklet::DispatcherMapField<CopyEdgeIds, DeviceAdapterTag> edgeDispatcher;
vtkm::worklet::DispatcherMapField<CopyEdgeIds> edgeDispatcher;
edgeDispatcher.SetDevice(DeviceAdapterTag());
edgeDispatcher.Invoke(uniqueKeys, edgeIds);
}
@ -677,13 +679,15 @@ struct GenerateNormalsDeduced
// The final normal is interpolated from the two gradient values and stored
// in the normals array.
//
vtkm::worklet::DispatcherMapTopology<NormalsWorkletPass1, DeviceAdapterTag>
dispatcherNormalsPass1(NormalsWorkletPass1::MakeScatter(*edges));
vtkm::worklet::DispatcherMapTopology<NormalsWorkletPass1> dispatcherNormalsPass1(
NormalsWorkletPass1::MakeScatter(*edges));
dispatcherNormalsPass1.SetDevice(DeviceAdapterTag());
dispatcherNormalsPass1.Invoke(
*cellset, *cellset, coordinates, marchingcubes::make_ScalarField(*field), *normals);
vtkm::worklet::DispatcherMapTopology<NormalsWorkletPass2, DeviceAdapterTag>
dispatcherNormalsPass2(NormalsWorkletPass2::MakeScatter(*edges));
vtkm::worklet::DispatcherMapTopology<NormalsWorkletPass2> dispatcherNormalsPass2(
NormalsWorkletPass2::MakeScatter(*edges));
dispatcherNormalsPass2.SetDevice(DeviceAdapterTag());
dispatcherNormalsPass2.Invoke(
*cellset, *cellset, coordinates, marchingcubes::make_ScalarField(*field), *weights, *normals);
}
@ -797,8 +801,8 @@ public:
{
using vtkm::worklet::marchingcubes::MapPointField;
MapPointField applyToField;
vtkm::worklet::DispatcherMapField<MapPointField, DeviceAdapter> applyFieldDispatcher(
applyToField);
vtkm::worklet::DispatcherMapField<MapPointField> applyFieldDispatcher(applyToField);
applyFieldDispatcher.SetDevice(DeviceAdapter());
vtkm::cont::ArrayHandle<ValueType> output;
applyFieldDispatcher.Invoke(
@ -940,12 +944,10 @@ private:
using vtkm::worklet::marchingcubes::MapPointField;
// Setup the Dispatcher Typedefs
using ClassifyDispatcher =
typename vtkm::worklet::DispatcherMapTopology<ClassifyCell<ValueType>, DeviceAdapter>;
using ClassifyDispatcher = vtkm::worklet::DispatcherMapTopology<ClassifyCell<ValueType>>;
using GenerateDispatcher =
typename vtkm::worklet::DispatcherMapTopology<EdgeWeightGenerate<ValueType, DeviceAdapter>,
DeviceAdapter>;
vtkm::worklet::DispatcherMapTopology<EdgeWeightGenerate<ValueType, DeviceAdapter>>;
vtkm::cont::ArrayHandle<ValueType> isoValuesHandle =
vtkm::cont::make_ArrayHandle(isovalues, numIsoValues);
@ -957,6 +959,7 @@ private:
{
ClassifyCell<ValueType> classifyCell;
ClassifyDispatcher classifyCellDispatcher(classifyCell);
classifyCellDispatcher.SetDevice(DeviceAdapter());
classifyCellDispatcher.Invoke(
isoValuesHandle, inputField, cells, numOutputTrisPerCell, this->NumTrianglesTable);
}
@ -983,6 +986,7 @@ private:
EdgeWeightGenerate<ValueType, DeviceAdapter> weightGenerate(metaData);
GenerateDispatcher edgeDispatcher(weightGenerate, scatter);
edgeDispatcher.SetDevice(DeviceAdapter());
edgeDispatcher.Invoke(
cells,
//cast to a scalar field if not one, as cellderivative only works on those
@ -1034,8 +1038,8 @@ private:
//generate the vertices's
MapPointField applyToField;
vtkm::worklet::DispatcherMapField<MapPointField, DeviceAdapter> applyFieldDispatcher(
applyToField);
vtkm::worklet::DispatcherMapField<MapPointField> applyFieldDispatcher(applyToField);
applyFieldDispatcher.SetDevice(DeviceAdapter());
applyFieldDispatcher.Invoke(
this->InterpolationEdgeIds, this->InterpolationWeights, coordinateSystem, vertices);

4
vtkm/worklet/NDimsEntropy.h
View File

@ -80,9 +80,9 @@ public:
vtkm::cont::ArrayHandle<vtkm::Float64> informationContent;
vtkm::worklet::histogram::SetBinInformationContent binWorklet(
static_cast<vtkm::Float64>(freqSum));
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::SetBinInformationContent,
DeviceAdapter>
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::SetBinInformationContent>
setBinInformationContentDispatcher(binWorklet);
setBinInformationContentDispatcher.SetDevice(DeviceAdapter());
setBinInformationContentDispatcher.Invoke(freqs, informationContent);
///// calculate entropy by summing up information conetent of all bins /////

30
vtkm/worklet/NDimsHistMarginalization.h
View File

@ -94,8 +94,9 @@ public:
numMarginalVariables++;
const vtkm::Id nFieldBins = numberOfBins.GetPortalControl().Get(i);
vtkm::worklet::histogram::To1DIndex binWorklet(nFieldBins);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::To1DIndex, DeviceAdapter>
to1DIndexDispatcher(binWorklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::To1DIndex> to1DIndexDispatcher(
binWorklet);
to1DIndexDispatcher.SetDevice(DeviceAdapter());
size_t vecIndex = static_cast<size_t>(i);
to1DIndexDispatcher.Invoke(binId[vecIndex], bin1DIndex, bin1DIndex);
}
@ -107,9 +108,9 @@ public:
conditionFunc
};
conditionalFreqWorklet.setVar(i);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConditionalFreq<BinaryCompare>,
DeviceAdapter>
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConditionalFreq<BinaryCompare>>
cfDispatcher(conditionalFreqWorklet);
cfDispatcher.SetDevice(DeviceAdapter());
size_t vecIndex = static_cast<size_t>(i);
cfDispatcher.Invoke(binId[vecIndex], freqs, freqs);
}
@ -138,11 +139,11 @@ public:
{
const vtkm::Id nFieldBins = numberOfBins.GetPortalControl().Get(i);
vtkm::worklet::histogram::ConvertHistBinToND binWorklet(nFieldBins);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConvertHistBinToND,
DeviceAdapter>
ConvertHistBinToNDDispatcher(binWorklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConvertHistBinToND>
convertHistBinToNDDispatcher(binWorklet);
convertHistBinToNDDispatcher.SetDevice(DeviceAdapter());
size_t vecIndex = static_cast<size_t>(marginalVarIdx);
ConvertHistBinToNDDispatcher.Invoke(
convertHistBinToNDDispatcher.Invoke(
sparseMarginal1DBinId, sparseMarginal1DBinId, marginalBinId[vecIndex]);
marginalVarIdx--;
}
@ -181,8 +182,9 @@ public:
numMarginalVariables++;
const vtkm::Id nFieldBins = numberOfBins.GetPortalControl().Get(i);
vtkm::worklet::histogram::To1DIndex binWorklet(nFieldBins);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::To1DIndex, DeviceAdapter>
to1DIndexDispatcher(binWorklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::To1DIndex> to1DIndexDispatcher(
binWorklet);
to1DIndexDispatcher.SetDevice(DeviceAdapter());
size_t vecIndex = static_cast<size_t>(i);
to1DIndexDispatcher.Invoke(binId[vecIndex], bin1DIndex, bin1DIndex);
}
@ -203,11 +205,11 @@ public:
{
const vtkm::Id nFieldBins = numberOfBins.GetPortalControl().Get(i);
vtkm::worklet::histogram::ConvertHistBinToND binWorklet(nFieldBins);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConvertHistBinToND,
DeviceAdapter>
ConvertHistBinToNDDispatcher(binWorklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConvertHistBinToND>
convertHistBinToNDDispatcher(binWorklet);
convertHistBinToNDDispatcher.SetDevice(DeviceAdapter());
size_t vecIndex = static_cast<size_t>(marginalVarIdx);
ConvertHistBinToNDDispatcher.Invoke(bin1DIndex, bin1DIndex, marginalBinId[vecIndex]);
convertHistBinToNDDispatcher.Invoke(bin1DIndex, bin1DIndex, marginalBinId[vecIndex]);
marginalVarIdx--;
}
}

7
vtkm/worklet/NDimsHistogram.h
View File

@ -109,10 +109,11 @@ public:
{
const vtkm::Id nFieldBins = NumberOfBins[static_cast<size_t>(i)];
vtkm::worklet::histogram::ConvertHistBinToND binWorklet(nFieldBins);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConvertHistBinToND, DeviceAdapter>
ConvertHistBinToNDDispatcher(binWorklet);
vtkm::worklet::DispatcherMapField<vtkm::worklet::histogram::ConvertHistBinToND>
convertHistBinToNDDispatcher(binWorklet);
convertHistBinToNDDispatcher.SetDevice(DeviceAdapter());
size_t vectorId = static_cast<size_t>(i);
ConvertHistBinToNDDispatcher.Invoke(Bin1DIndex, Bin1DIndex, binId[vectorId]);
convertHistBinToNDDispatcher.Invoke(Bin1DIndex, Bin1DIndex, binId[vectorId]);
}
}

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