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

Remove VTKM_EXEC_CONSTANT

Browse Source 
If a global static array is declared with VTKM_EXEC_CONSTANT and the code
is compiled by nvcc (for multibackend code) then the array is only accesible
on the GPU. If for some reason a worklet fails on the cuda backend and it is
re-executed on any of the CPU backends, it will continue to fail.

We couldn't find a simple way to declare the array once and have it available
on both CPU and GPU. The approach we are using here is to declare the arrays
as static inside some "Get" function which is marked as VTKM_EXEC_CONT.
This commit is contained in:
Sujin Philip 2017-12-05 13:42:03 -05:00
parent 1d58eeaa9e
commit e28309f09b
15 changed files with 795 additions and 711 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
vtkm/VecAxisAlignedPointCoordinates.h
View File

@ -56,12 +56,6 @@ struct VecAxisAlignedPointCoordinatesNumComponents<3>
static const vtkm::IdComponent NUM_COMPONENTS = 8; static const vtkm::IdComponent NUM_COMPONENTS = 8;
}; };
VTKM_EXEC_CONSTANT
const vtkm::FloatDefault VecAxisAlignedPointCoordinatesOffsetTable[8][3] = {
{ 0.0f, 0.0f, 0.0f }, { 1.0f, 0.0f, 0.0f }, { 1.0f, 1.0f, 0.0f }, { 0.0f, 1.0f, 0.0f },
{ 0.0f, 0.0f, 1.0f }, { 1.0f, 0.0f, 1.0f }, { 1.0f, 1.0f, 1.0f }, { 0.0f, 1.0f, 1.0f }
};
} // namespace detail } // namespace detail
/// \brief An implicit vector for point coordinates in axis aligned cells. For /// \brief An implicit vector for point coordinates in axis aligned cells. For
@ -110,7 +104,12 @@ public:
VTKM_EXEC_CONT VTKM_EXEC_CONT
ComponentType operator[](vtkm::IdComponent index) const ComponentType operator[](vtkm::IdComponent index) const
{ {
const vtkm::FloatDefault* offset = detail::VecAxisAlignedPointCoordinatesOffsetTable[index]; static const vtkm::FloatDefault VecAxisAlignedPointCoordinatesOffsetTable[8][3] = {
{ 0.0f, 0.0f, 0.0f }, { 1.0f, 0.0f, 0.0f }, { 1.0f, 1.0f, 0.0f }, { 0.0f, 1.0f, 0.0f },
{ 0.0f, 0.0f, 1.0f }, { 1.0f, 0.0f, 1.0f }, { 1.0f, 1.0f, 1.0f }, { 0.0f, 1.0f, 1.0f }
};
const auto& offset = VecAxisAlignedPointCoordinatesOffsetTable[index];
return ComponentType(this->Origin[0] + offset[0] * this->Spacing[0], return ComponentType(this->Origin[0] + offset[0] * this->Spacing[0],
this->Origin[1] + offset[1] * this->Spacing[1], this->Origin[1] + offset[1] * this->Spacing[1],
this->Origin[2] + offset[2] * this->Spacing[2]); this->Origin[2] + offset[2] * this->Spacing[2]);

471
vtkm/exec/CellEdge.h
View File

@ -17,8 +17,8 @@
// Laboratory (LANL), the U.S. Government retains certain rights in // Laboratory (LANL), the U.S. Government retains certain rights in
// this software. // this software.
//============================================================================ //============================================================================
#ifndef vtk_m_exec_CellFaces_h #ifndef vtk_m_exec_CellEdge_h
#define vtk_m_exec_CellFaces_h #define vtk_m_exec_CellEdge_h
#include <vtkm/Assert.h> #include <vtkm/Assert.h>
#include <vtkm/CellShape.h> #include <vtkm/CellShape.h>
@ -35,224 +35,241 @@ namespace exec
namespace detail namespace detail
{ {
static const vtkm::IdComponent MAX_NUM_EDGES = 12; class CellEdgeTables
{
public:
static const vtkm::IdComponent MAX_NUM_EDGES = 12;
VTKM_EXEC_CONSTANT private:
static const vtkm::IdComponent NumEdges[vtkm::NUMBER_OF_CELL_SHAPES] = { struct Tables
0, // 0: CELL_SHAPE_EMPTY {
0, // 1: CELL_SHAPE_VERTEX vtkm::IdComponent NumEdges[vtkm::NUMBER_OF_CELL_SHAPES];
0, // 2: Unused vtkm::IdComponent PointsInEdge[vtkm::NUMBER_OF_CELL_SHAPES][MAX_NUM_EDGES][2];
0, // 3: CELL_SHAPE_LINE };
0, // 4: Unused
3, // 5: CELL_SHAPE_TRIANGLE
0, // 6: Unused
-1, // 7: CELL_SHAPE_POLYGON ---special case---
0, // 8: Unused
4, // 9: CELL_SHAPE_QUAD
6, // 10: CELL_SHAPE_TETRA
0, // 11: Unused
12, // 12: CELL_SHAPE_HEXAHEDRON
9, // 13: CELL_SHAPE_WEDGE
8 // 14: CELL_SHAPE_PYRAMID
};
VTKM_EXEC_CONSTANT public:
static const vtkm::IdComponent PointsInEdge[vtkm::NUMBER_OF_CELL_SHAPES][MAX_NUM_EDGES][2] = { VTKM_EXEC_CONT static const Tables& Get()
// 0: CELL_SHAPE_EMPTY {
{ { -1, -1 }, static const Tables table = { {
{ -1, -1 }, // NumEdges
{ -1, -1 }, 0, // 0: CELL_SHAPE_EMPTY
{ -1, -1 }, 0, // 1: CELL_SHAPE_VERTEX
{ -1, -1 }, 0, // 2: Unused
{ -1, -1 }, 0, // 3: CELL_SHAPE_LINE
{ -1, -1 }, 0, // 4: Unused
{ -1, -1 }, 3, // 5: CELL_SHAPE_TRIANGLE
{ -1, -1 }, 0, // 6: Unused
{ -1, -1 }, -1, // 7: CELL_SHAPE_POLYGON ---special case---
{ -1, -1 }, 0, // 8: Unused
{ -1, -1 } }, 4, // 9: CELL_SHAPE_QUAD
// 1: CELL_SHAPE_VERTEX 6, // 10: CELL_SHAPE_TETRA
{ { -1, -1 }, 0, // 11: Unused
{ -1, -1 }, 12, // 12: CELL_SHAPE_HEXAHEDRON
{ -1, -1 }, 9, // 13: CELL_SHAPE_WEDGE
{ -1, -1 }, 8 // 14: CELL_SHAPE_PYRAMID
{ -1, -1 }, },
{ -1, -1 },
{ -1, -1 }, {
{ -1, -1 }, // PointsInEdge
{ -1, -1 }, // 0: CELL_SHAPE_EMPTY
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 2: Unused { -1, -1 },
{ { -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 1: CELL_SHAPE_VERTEX
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 3: CELL_SHAPE_LINE { -1, -1 },
{ { -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 2: Unused
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 4: Unused { -1, -1 },
{ { -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 3: CELL_SHAPE_LINE
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 5: CELL_SHAPE_TRIANGLE { -1, -1 },
{ { 0, 1 }, { -1, -1 },
{ 1, 2 }, { -1, -1 },
{ 2, 0 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 4: Unused
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 6: Unused { -1, -1 },
{ { -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 5: CELL_SHAPE_TRIANGLE
{ -1, -1 }, { { 0, 1 },
{ -1, -1 }, { 1, 2 },
{ -1, -1 } }, { 2, 0 },
// 7: CELL_SHAPE_POLYGON --- special case --- { -1, -1 },
{ { -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 6: Unused
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 8: Unused { -1, -1 },
{ { -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 7: CELL_SHAPE_POLYGON --- special case ---
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 9: CELL_SHAPE_QUAD { -1, -1 },
{ { 0, 1 }, { -1, -1 },
{ 1, 2 }, { -1, -1 },
{ 2, 3 }, { -1, -1 },
{ 3, 0 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 8: Unused
{ -1, -1 }, { { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 } }, { -1, -1 },
// 10: CELL_SHAPE_TETRA { -1, -1 },
{ { 0, 1 }, { -1, -1 },
{ 1, 2 }, { -1, -1 },
{ 2, 0 }, { -1, -1 },
{ 0, 3 }, { -1, -1 },
{ 1, 3 }, { -1, -1 },
{ 2, 3 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 9: CELL_SHAPE_QUAD
{ -1, -1 }, { { 0, 1 },
{ -1, -1 }, { 1, 2 },
{ -1, -1 } }, { 2, 3 },
// 11: Unused { 3, 0 },
{ { -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 },
{ -1, -1 }, { -1, -1 } },
{ -1, -1 }, // 10: CELL_SHAPE_TETRA
{ -1, -1 }, { { 0, 1 },
{ -1, -1 }, { 1, 2 },
{ -1, -1 } }, { 2, 0 },
// 12: CELL_SHAPE_HEXAHEDRON { 0, 3 },
{ { 0, 1 }, { 1, 3 },
{ 1, 2 }, { 2, 3 },
{ 3, 2 }, { -1, -1 },
{ 0, 3 }, { -1, -1 },
{ 4, 5 }, { -1, -1 },
{ 5, 6 }, { -1, -1 },
{ 7, 6 }, { -1, -1 },
{ 4, 7 }, { -1, -1 } },
{ 0, 4 }, // 11: Unused
{ 1, 5 }, { { -1, -1 },
{ 3, 7 }, { -1, -1 },
{ 2, 6 } }, { -1, -1 },
// 13: CELL_SHAPE_WEDGE { -1, -1 },
{ { 0, 1 }, { -1, -1 },
{ 1, 2 }, { -1, -1 },
{ 2, 0 }, { -1, -1 },
{ 3, 4 }, { -1, -1 },
{ 4, 5 }, { -1, -1 },
{ 5, 3 }, { -1, -1 },
{ 0, 3 }, { -1, -1 },
{ 1, 4 }, { -1, -1 } },
{ 2, 5 }, // 12: CELL_SHAPE_HEXAHEDRON
{ -1, -1 }, { { 0, 1 },
{ -1, -1 }, { 1, 2 },
{ -1, -1 } }, { 3, 2 },
// 14: CELL_SHAPE_PYRAMID { 0, 3 },
{ { 0, 1 }, { 4, 5 },
{ 1, 2 }, { 5, 6 },
{ 2, 3 }, { 7, 6 },
{ 3, 0 }, { 4, 7 },
{ 0, 4 }, { 0, 4 },
{ 1, 4 }, { 1, 5 },
{ 2, 4 }, { 3, 7 },
{ 3, 4 }, { 2, 6 } },
{ -1, -1 }, // 13: CELL_SHAPE_WEDGE
{ -1, -1 }, { { 0, 1 },
{ -1, -1 }, { 1, 2 },
{ -1, -1 } }, { 2, 0 },
{ 3, 4 },
{ 4, 5 },
{ 5, 3 },
{ 0, 3 },
{ 1, 4 },
{ 2, 5 },
{ -1, -1 },
{ -1, -1 },
{ -1, -1 } },
// 14: CELL_SHAPE_PYRAMID
{ { 0, 1 },
{ 1, 2 },
{ 2, 3 },
{ 3, 0 },
{ 0, 4 },
{ 1, 4 },
{ 2, 4 },
{ 3, 4 },
{ -1, -1 },
{ -1, -1 },
{ -1, -1 },
{ -1, -1 } },
} };
return table;
}
}; };
} // namespace detail } // namespace detail
@ -264,7 +281,7 @@ static inline VTKM_EXEC vtkm::IdComponent CellEdgeNumberOfEdges(vtkm::IdComponen
{ {
(void)numPoints; // Silence compiler warnings. (void)numPoints; // Silence compiler warnings.
VTKM_ASSERT(numPoints == vtkm::CellTraits<CellShapeTag>::NUM_POINTS); VTKM_ASSERT(numPoints == vtkm::CellTraits<CellShapeTag>::NUM_POINTS);
return detail::NumEdges[CellShapeTag::Id]; return detail::CellEdgeTables::Get().NumEdges[CellShapeTag::Id];
} }
static inline VTKM_EXEC vtkm::IdComponent CellEdgeNumberOfEdges(vtkm::IdComponent numPoints, static inline VTKM_EXEC vtkm::IdComponent CellEdgeNumberOfEdges(vtkm::IdComponent numPoints,
@ -286,7 +303,7 @@ static inline VTKM_EXEC vtkm::IdComponent CellEdgeNumberOfEdges(
} }
else else
{ {
return detail::NumEdges[shape.Id]; return detail::CellEdgeTables::Get().NumEdges[shape.Id];
} }
} }
@ -298,15 +315,15 @@ static inline VTKM_EXEC vtkm::Vec<vtkm::IdComponent, 2> CellEdgeLocalIndices(
const vtkm::exec::FunctorBase& worklet) const vtkm::exec::FunctorBase& worklet)
{ {
VTKM_ASSUME(edgeIndex >= 0); VTKM_ASSUME(edgeIndex >= 0);
VTKM_ASSUME(edgeIndex < detail::MAX_NUM_EDGES); VTKM_ASSUME(edgeIndex < detail::CellEdgeTables::MAX_NUM_EDGES);
if (edgeIndex >= vtkm::exec::CellEdgeNumberOfEdges(numPoints, shape, worklet)) if (edgeIndex >= vtkm::exec::CellEdgeNumberOfEdges(numPoints, shape, worklet))
{ {
worklet.RaiseError("Invalid edge number."); worklet.RaiseError("Invalid edge number.");
return vtkm::Vec<vtkm::IdComponent, 2>(0); return vtkm::Vec<vtkm::IdComponent, 2>(0);
} }
return vtkm::make_Vec(detail::PointsInEdge[CellShapeTag::Id][edgeIndex][0], return vtkm::make_Vec(detail::CellEdgeTables::Get().PointsInEdge[CellShapeTag::Id][edgeIndex][0],
detail::PointsInEdge[CellShapeTag::Id][edgeIndex][1]); detail::CellEdgeTables::Get().PointsInEdge[CellShapeTag::Id][edgeIndex][1]);
} }
static inline VTKM_EXEC vtkm::Vec<vtkm::IdComponent, 2> CellEdgeLocalIndices( static inline VTKM_EXEC vtkm::Vec<vtkm::IdComponent, 2> CellEdgeLocalIndices(
@ -336,7 +353,7 @@ static inline VTKM_EXEC vtkm::Vec<vtkm::IdComponent, 2> CellEdgeLocalIndices(
const vtkm::exec::FunctorBase& worklet) const vtkm::exec::FunctorBase& worklet)
{ {
VTKM_ASSUME(edgeIndex >= 0); VTKM_ASSUME(edgeIndex >= 0);
VTKM_ASSUME(edgeIndex < detail::MAX_NUM_EDGES); VTKM_ASSUME(edgeIndex < detail::CellEdgeTables::MAX_NUM_EDGES);
if (shape.Id == vtkm::CELL_SHAPE_POLYGON) if (shape.Id == vtkm::CELL_SHAPE_POLYGON)
{ {
@ -344,14 +361,14 @@ static inline VTKM_EXEC vtkm::Vec<vtkm::IdComponent, 2> CellEdgeLocalIndices(
} }
else else
{ {
if (edgeIndex >= detail::NumEdges[shape.Id]) if (edgeIndex >= detail::CellEdgeTables::Get().NumEdges[shape.Id])
{ {
worklet.RaiseError("Invalid edge number."); worklet.RaiseError("Invalid edge number.");
return vtkm::Vec<vtkm::IdComponent, 2>(0); return vtkm::Vec<vtkm::IdComponent, 2>(0);
} }
return vtkm::make_Vec(detail::PointsInEdge[shape.Id][edgeIndex][0], return vtkm::make_Vec(detail::CellEdgeTables::Get().PointsInEdge[shape.Id][edgeIndex][0],
detail::PointsInEdge[shape.Id][edgeIndex][1]); detail::CellEdgeTables::Get().PointsInEdge[shape.Id][edgeIndex][1]);
} }
} }

324
vtkm/exec/CellFace.h
View File

@ -34,157 +34,174 @@ namespace exec
namespace detail namespace detail
{ {
static const vtkm::IdComponent MAX_FACE_SIZE = 4; class CellFaceTables
static const vtkm::IdComponent MAX_NUM_FACES = 6; {
public:
static const vtkm::IdComponent MAX_FACE_SIZE = 4;
static const vtkm::IdComponent MAX_NUM_FACES = 6;
VTKM_EXEC_CONSTANT private:
static const vtkm::IdComponent NumFaces[vtkm::NUMBER_OF_CELL_SHAPES] = { struct Tables
0, // 0: CELL_SHAPE_EMPTY {
0, // 1: CELL_SHAPE_VERTEX vtkm::IdComponent NumFaces[vtkm::NUMBER_OF_CELL_SHAPES];
0, // 2: Unused vtkm::IdComponent NumPointsInFace[vtkm::NUMBER_OF_CELL_SHAPES][MAX_NUM_FACES];
0, // 3: CELL_SHAPE_LINE vtkm::IdComponent PointsInFace[vtkm::NUMBER_OF_CELL_SHAPES][MAX_NUM_FACES][MAX_FACE_SIZE];
0, // 4: Unused };
0, // 5: CELL_SHAPE_TRIANGLE
0, // 6: Unused public:
0, // 7: CELL_SHAPE_POLYGON VTKM_EXEC_CONT
0, // 8: Unused static const Tables& Get()
0, // 9: CELL_SHAPE_QUAD {
4, // 10: CELL_SHAPE_TETRA static const Tables table = { // NumFaces
0, // 11: Unused {
6, // 12: CELL_SHAPE_HEXAHEDRON 0, // 0: CELL_SHAPE_EMPTY
5, // 13: CELL_SHAPE_WEDGE 0, // 1: CELL_SHAPE_VERTEX
5 // 14: CELL_SHAPE_PYRAMID 0, // 2: Unused
0, // 3: CELL_SHAPE_LINE
0, // 4: Unused
0, // 5: CELL_SHAPE_TRIANGLE
0, // 6: Unused
0, // 7: CELL_SHAPE_POLYGON
0, // 8: Unused
0, // 9: CELL_SHAPE_QUAD
4, // 10: CELL_SHAPE_TETRA
0, // 11: Unused
6, // 12: CELL_SHAPE_HEXAHEDRON
5, // 13: CELL_SHAPE_WEDGE
5 // 14: CELL_SHAPE_PYRAMID
},
// NumPointsInFace
{
{ -1, -1, -1, -1, -1, -1 }, // 0: CELL_SHAPE_EMPTY
{ -1, -1, -1, -1, -1, -1 }, // 1: CELL_SHAPE_VERTEX
{ -1, -1, -1, -1, -1, -1 }, // 2: Unused
{ -1, -1, -1, -1, -1, -1 }, // 3: CELL_SHAPE_LINE
{ -1, -1, -1, -1, -1, -1 }, // 4: Unused
{ -1, -1, -1, -1, -1, -1 }, // 5: CELL_SHAPE_TRIANGLE
{ -1, -1, -1, -1, -1, -1 }, // 6: Unused
{ -1, -1, -1, -1, -1, -1 }, // 7: CELL_SHAPE_POLYGON
{ -1, -1, -1, -1, -1, -1 }, // 8: Unused
{ -1, -1, -1, -1, -1, -1 }, // 9: CELL_SHAPE_QUAD
{ 3, 3, 3, 3, -1, -1 }, // 10: CELL_SHAPE_TETRA
{ -1, -1, -1, -1, -1, -1 }, // 11: Unused
{ 4, 4, 4, 4, 4, 4 }, // 12: CELL_SHAPE_HEXAHEDRON
{ 3, 3, 4, 4, 4, -1 }, // 13: CELL_SHAPE_WEDGE
{ 4, 3, 3, 3, 3, -1 } // 14: CELL_SHAPE_PYRAMID
},
// PointsInFace
{ // 0: CELL_SHAPE_EMPTY
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 1: CELL_SHAPE_VERTEX
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 2: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 3: CELL_SHAPE_LINE
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 4: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 5: CELL_SHAPE_TRIANGLE
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 6: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 7: CELL_SHAPE_POLYGON
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 8: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 9: CELL_SHAPE_QUAD
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 10: CELL_SHAPE_TETRA
{ { 0, 1, 3, -1 },
{ 1, 2, 3, -1 },
{ 2, 0, 3, -1 },
{ 0, 2, 1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 11: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 12: CELL_SHAPE_HEXAHEDRON
{ { 0, 4, 7, 3 },
{ 1, 2, 6, 5 },
{ 0, 1, 5, 4 },
{ 3, 7, 6, 2 },
{ 0, 3, 2, 1 },
{ 4, 5, 6, 7 } },
// 13: CELL_SHAPE_WEDGE
{ { 0, 1, 2, -1 },
{ 3, 5, 4, -1 },
{ 0, 3, 4, 1 },
{ 1, 4, 5, 2 },
{ 2, 5, 3, 0 },
{ -1, -1, -1, -1 } },
// 14: CELL_SHAPE_PYRAMID
{ { 0, 3, 2, 1 },
{ 0, 1, 4, -1 },
{ 1, 2, 4, -1 },
{ 2, 3, 4, -1 },
{ 3, 0, 4, -1 },
{ -1, -1, -1, -1 } } }
};
return table;
}
}; };
VTKM_EXEC_CONSTANT
static const vtkm::IdComponent NumPointsInFace[vtkm::NUMBER_OF_CELL_SHAPES][MAX_NUM_FACES] = {
{ -1, -1, -1, -1, -1, -1 }, // 0: CELL_SHAPE_EMPTY
{ -1, -1, -1, -1, -1, -1 }, // 1: CELL_SHAPE_VERTEX
{ -1, -1, -1, -1, -1, -1 }, // 2: Unused
{ -1, -1, -1, -1, -1, -1 }, // 3: CELL_SHAPE_LINE
{ -1, -1, -1, -1, -1, -1 }, // 4: Unused
{ -1, -1, -1, -1, -1, -1 }, // 5: CELL_SHAPE_TRIANGLE
{ -1, -1, -1, -1, -1, -1 }, // 6: Unused
{ -1, -1, -1, -1, -1, -1 }, // 7: CELL_SHAPE_POLYGON
{ -1, -1, -1, -1, -1, -1 }, // 8: Unused
{ -1, -1, -1, -1, -1, -1 }, // 9: CELL_SHAPE_QUAD
{ 3, 3, 3, 3, -1, -1 }, // 10: CELL_SHAPE_TETRA
{ -1, -1, -1, -1, -1, -1 }, // 11: Unused
{ 4, 4, 4, 4, 4, 4 }, // 12: CELL_SHAPE_HEXAHEDRON
{ 3, 3, 4, 4, 4, -1 }, // 13: CELL_SHAPE_WEDGE
{ 4, 3, 3, 3, 3, -1 } // 14: CELL_SHAPE_PYRAMID
};
VTKM_EXEC_CONSTANT
static const vtkm::IdComponent PointsInFace[vtkm::NUMBER_OF_CELL_SHAPES][MAX_NUM_FACES]
[MAX_FACE_SIZE] = {
// 0: CELL_SHAPE_EMPTY
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 1: CELL_SHAPE_VERTEX
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 2: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 3: CELL_SHAPE_LINE
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 4: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 5: CELL_SHAPE_TRIANGLE
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 6: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 7: CELL_SHAPE_POLYGON
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 8: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 9: CELL_SHAPE_QUAD
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 10: CELL_SHAPE_TETRA
{ { 0, 1, 3, -1 },
{ 1, 2, 3, -1 },
{ 2, 0, 3, -1 },
{ 0, 2, 1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 11: Unused
{ { -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 },
{ -1, -1, -1, -1 } },
// 12: CELL_SHAPE_HEXAHEDRON
{ { 0, 4, 7, 3 },
{ 1, 2, 6, 5 },
{ 0, 1, 5, 4 },
{ 3, 7, 6, 2 },
{ 0, 3, 2, 1 },
{ 4, 5, 6, 7 } },
// 13: CELL_SHAPE_WEDGE
{ { 0, 1, 2, -1 },
{ 3, 5, 4, -1 },
{ 0, 3, 4, 1 },
{ 1, 4, 5, 2 },
{ 2, 5, 3, 0 },
{ -1, -1, -1, -1 } },
// 14: CELL_SHAPE_PYRAMID
{ { 0, 3, 2, 1 },
{ 0, 1, 4, -1 },
{ 1, 2, 4, -1 },
{ 2, 3, 4, -1 },
{ 3, 0, 4, -1 },
{ -1, -1, -1, -1 } }
};
} // namespace detail } // namespace detail
template <typename CellShapeTag> template <typename CellShapeTag>
@ -192,7 +209,7 @@ static inline VTKM_EXEC vtkm::IdComponent CellFaceNumberOfFaces(CellShapeTag sha
const vtkm::exec::FunctorBase&) const vtkm::exec::FunctorBase&)
{ {
(void)shape; //C4100 false positive workaround (void)shape; //C4100 false positive workaround
return detail::NumFaces[shape.Id]; return detail::CellFaceTables::Get().NumFaces[shape.Id];
} }
template <typename CellShapeTag> template <typename CellShapeTag>
@ -202,14 +219,14 @@ static inline VTKM_EXEC vtkm::IdComponent CellFaceNumberOfPoints(
const vtkm::exec::FunctorBase& worklet) const vtkm::exec::FunctorBase& worklet)
{ {
VTKM_ASSUME(faceIndex >= 0); VTKM_ASSUME(faceIndex >= 0);
VTKM_ASSUME(faceIndex < detail::MAX_NUM_FACES); VTKM_ASSUME(faceIndex < detail::CellFaceTables::MAX_NUM_FACES);
if (faceIndex >= vtkm::exec::CellFaceNumberOfFaces(shape, worklet)) if (faceIndex >= vtkm::exec::CellFaceNumberOfFaces(shape, worklet))
{ {
worklet.RaiseError("Invalid face number."); worklet.RaiseError("Invalid face number.");
return 0; return 0;
} }
return detail::NumPointsInFace[shape.Id][faceIndex]; return detail::CellFaceTables::Get().NumPointsInFace[shape.Id][faceIndex];
} }
template <typename CellShapeTag> template <typename CellShapeTag>
@ -218,7 +235,7 @@ static inline VTKM_EXEC vtkm::UInt8 CellFaceShape(vtkm::IdComponent faceIndex,
const vtkm::exec::FunctorBase& worklet) const vtkm::exec::FunctorBase& worklet)
{ {
VTKM_ASSUME(faceIndex >= 0); VTKM_ASSUME(faceIndex >= 0);
VTKM_ASSUME(faceIndex < detail::MAX_NUM_FACES); VTKM_ASSUME(faceIndex < detail::CellFaceTables::MAX_NUM_FACES);
switch (CellFaceNumberOfPoints(faceIndex, shape, worklet)) switch (CellFaceNumberOfPoints(faceIndex, shape, worklet))
{ {
case 3: case 3:
@ -244,7 +261,8 @@ static inline VTKM_EXEC vtkm::VecCConst<vtkm::IdComponent> CellFaceLocalIndices(
return vtkm::VecCConst<vtkm::IdComponent>(); return vtkm::VecCConst<vtkm::IdComponent>();
} }
return vtkm::make_VecC(detail::PointsInFace[shape.Id][faceIndex], numPointsInFace); return vtkm::make_VecC(detail::CellFaceTables::Get().PointsInFace[shape.Id][faceIndex],
numPointsInFace);
} }
/// \brief Returns a canonical identifer for a cell face /// \brief Returns a canonical identifer for a cell face

2
vtkm/internal/ExportMacros.h
View File

@ -34,12 +34,10 @@
#else #else
#define VTKM_SUPPRESS_EXEC_WARNINGS #pragma hd_warning_disable #define VTKM_SUPPRESS_EXEC_WARNINGS #pragma hd_warning_disable
#endif #endif
#define VTKM_EXEC_CONSTANT __device__ __constant__
#else #else
#define VTKM_EXEC #define VTKM_EXEC
#define VTKM_EXEC_CONT #define VTKM_EXEC_CONT
#define VTKM_SUPPRESS_EXEC_WARNINGS #define VTKM_SUPPRESS_EXEC_WARNINGS
#define VTKM_EXEC_CONSTANT
#endif #endif
#define VTKM_CONT #define VTKM_CONT

98
vtkm/rendering/raytracing/CellIntersector.h
View File

@ -47,8 +47,10 @@ VTKM_EXEC_CONT inline void IntersectZoo(T xpoints[8],
vtkm::Int32 kx, ky, kz; vtkm::Int32 kx, ky, kz;
WaterTight<T> intersector; WaterTight<T> intersector;
intersector.FindDir(dir, sx, sy, sz, kx, ky, kz); intersector.FindDir(dir, sx, sy, sz, kx, ky, kz);
const vtkm::Int32 tableOffset = ZooLookUp[CellTypeLookUp[shapeType]][0]; const vtkm::Int32 tableOffset =
const vtkm::Int32 numTriangles = ZooLookUp[CellTypeLookUp[shapeType]][1]; CellTables::Get().ZooLookUp[CellTables::Get().CellTypeLookUp[shapeType]][0];
const vtkm::Int32 numTriangles =
CellTables::Get().ZooLookUp[CellTables::Get().CellTypeLookUp[shapeType]][1];
// Decompose each face into two triangles // Decompose each face into two triangles
for (int i = 0; i < 6; ++i) for (int i = 0; i < 6; ++i)
distances[i] = -1.; distances[i] = -1.;
@ -56,16 +58,16 @@ VTKM_EXEC_CONT inline void IntersectZoo(T xpoints[8],
{ {
const vtkm::Int32 offset = tableOffset + i; const vtkm::Int32 offset = tableOffset + i;
vtkm::Vec<T, 3> a, c, b; vtkm::Vec<T, 3> a, c, b;
a[0] = xpoints[ZooTable[offset][1]]; a[0] = xpoints[CellTables::Get().ZooTable[offset][1]];
a[1] = ypoints[ZooTable[offset][1]]; a[1] = ypoints[CellTables::Get().ZooTable[offset][1]];
a[2] = zpoints[ZooTable[offset][1]]; a[2] = zpoints[CellTables::Get().ZooTable[offset][1]];
b[0] = xpoints[ZooTable[offset][2]]; b[0] = xpoints[CellTables::Get().ZooTable[offset][2]];
b[1] = ypoints[ZooTable[offset][2]]; b[1] = ypoints[CellTables::Get().ZooTable[offset][2]];
b[2] = zpoints[ZooTable[offset][2]]; b[2] = zpoints[CellTables::Get().ZooTable[offset][2]];
c[0] = xpoints[ZooTable[offset][3]]; c[0] = xpoints[CellTables::Get().ZooTable[offset][3]];
c[1] = ypoints[ZooTable[offset][3]]; c[1] = ypoints[CellTables::Get().ZooTable[offset][3]];
c[2] = zpoints[ZooTable[offset][3]]; c[2] = zpoints[CellTables::Get().ZooTable[offset][3]];
const vtkm::Int32 faceId = ZooTable[offset][0]; const vtkm::Int32 faceId = CellTables::Get().ZooTable[offset][0];
T distance = -1.f; T distance = -1.f;
T uNotUsed, vNotUsed; T uNotUsed, vNotUsed;
@ -111,18 +113,18 @@ VTKM_EXEC_CONT inline void IntersectHex(T xpoints[8],
for (int i = 0; i < 6; ++i) for (int i = 0; i < 6; ++i)
{ {
vtkm::Vec<T, 3> a, c, b, d; vtkm::Vec<T, 3> a, c, b, d;
a[0] = xpoints[ShapesFaceList[i][1]]; a[0] = xpoints[CellTables::Get().ShapesFaceList[i][1]];
a[1] = ypoints[ShapesFaceList[i][1]]; a[1] = ypoints[CellTables::Get().ShapesFaceList[i][1]];
a[2] = zpoints[ShapesFaceList[i][1]]; a[2] = zpoints[CellTables::Get().ShapesFaceList[i][1]];
b[0] = xpoints[ShapesFaceList[i][2]]; b[0] = xpoints[CellTables::Get().ShapesFaceList[i][2]];
b[1] = ypoints[ShapesFaceList[i][2]]; b[1] = ypoints[CellTables::Get().ShapesFaceList[i][2]];
b[2] = zpoints[ShapesFaceList[i][2]]; b[2] = zpoints[CellTables::Get().ShapesFaceList[i][2]];
c[0] = xpoints[ShapesFaceList[i][3]]; c[0] = xpoints[CellTables::Get().ShapesFaceList[i][3]];
c[1] = ypoints[ShapesFaceList[i][3]]; c[1] = ypoints[CellTables::Get().ShapesFaceList[i][3]];
c[2] = zpoints[ShapesFaceList[i][3]]; c[2] = zpoints[CellTables::Get().ShapesFaceList[i][3]];
d[0] = xpoints[ShapesFaceList[i][4]]; d[0] = xpoints[CellTables::Get().ShapesFaceList[i][4]];
d[1] = ypoints[ShapesFaceList[i][4]]; d[1] = ypoints[CellTables::Get().ShapesFaceList[i][4]];
d[2] = zpoints[ShapesFaceList[i][4]]; d[2] = zpoints[CellTables::Get().ShapesFaceList[i][4]];
T distance = -1.f; T distance = -1.f;
distances[i] = distance; //init to -1 distances[i] = distance; //init to -1
@ -189,19 +191,20 @@ VTKM_EXEC_CONT inline void IntersectTet(T xpoints[8],
WaterTight<T> intersector; WaterTight<T> intersector;
intersector.FindDir(dir, sx, sy, sz, kx, ky, kz); intersector.FindDir(dir, sx, sy, sz, kx, ky, kz);
const vtkm::Int32 tableOffset = FaceLookUp[CellTypeLookUp[CELL_SHAPE_TETRA]][0]; const vtkm::Int32 tableOffset =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[CELL_SHAPE_TETRA]][0];
for (vtkm::Int32 i = 0; i < 4; ++i) for (vtkm::Int32 i = 0; i < 4; ++i)
{ {
vtkm::Vec<T, 3> a, c, b; vtkm::Vec<T, 3> a, c, b;
a[0] = xpoints[ShapesFaceList[i + tableOffset][1]]; a[0] = xpoints[CellTables::Get().ShapesFaceList[i + tableOffset][1]];
a[1] = ypoints[ShapesFaceList[i + tableOffset][1]]; a[1] = ypoints[CellTables::Get().ShapesFaceList[i + tableOffset][1]];
a[2] = zpoints[ShapesFaceList[i + tableOffset][1]]; a[2] = zpoints[CellTables::Get().ShapesFaceList[i + tableOffset][1]];
b[0] = xpoints[ShapesFaceList[i + tableOffset][2]]; b[0] = xpoints[CellTables::Get().ShapesFaceList[i + tableOffset][2]];
b[1] = ypoints[ShapesFaceList[i + tableOffset][2]]; b[1] = ypoints[CellTables::Get().ShapesFaceList[i + tableOffset][2]];
b[2] = zpoints[ShapesFaceList[i + tableOffset][2]]; b[2] = zpoints[CellTables::Get().ShapesFaceList[i + tableOffset][2]];
c[0] = xpoints[ShapesFaceList[i + tableOffset][3]]; c[0] = xpoints[CellTables::Get().ShapesFaceList[i + tableOffset][3]];
c[1] = ypoints[ShapesFaceList[i + tableOffset][3]]; c[1] = ypoints[CellTables::Get().ShapesFaceList[i + tableOffset][3]];
c[2] = zpoints[ShapesFaceList[i + tableOffset][3]]; c[2] = zpoints[CellTables::Get().ShapesFaceList[i + tableOffset][3]];
T distance = -1.f; T distance = -1.f;
distances[i] = distance; //init to -1 distances[i] = distance; //init to -1
@ -245,23 +248,24 @@ VTKM_EXEC_CONT inline void IntersectWedge(T xpoints[8],
WaterTight<T> intersector; WaterTight<T> intersector;
intersector.FindDir(dir, sx, sy, sz, kx, ky, kz); intersector.FindDir(dir, sx, sy, sz, kx, ky, kz);
// TODO: try two sepate loops to see performance impact // TODO: try two sepate loops to see performance impact
const vtkm::Int32 tableOffset = FaceLookUp[CellTypeLookUp[CELL_SHAPE_WEDGE]][0]; const vtkm::Int32 tableOffset =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[CELL_SHAPE_WEDGE]][0];
// Decompose each face into two triangles // Decompose each face into two triangles
for (int i = 0; i < 5; ++i) for (int i = 0; i < 5; ++i)
{ {
vtkm::Vec<T, 3> a, c, b, d; vtkm::Vec<T, 3> a, c, b, d;
a[0] = xpoints[ShapesFaceList[i + tableOffset][1]]; a[0] = xpoints[CellTables::Get().ShapesFaceList[i + tableOffset][1]];
a[1] = ypoints[ShapesFaceList[i + tableOffset][1]]; a[1] = ypoints[CellTables::Get().ShapesFaceList[i + tableOffset][1]];
a[2] = zpoints[ShapesFaceList[i + tableOffset][1]]; a[2] = zpoints[CellTables::Get().ShapesFaceList[i + tableOffset][1]];
b[0] = xpoints[ShapesFaceList[i + tableOffset][2]]; b[0] = xpoints[CellTables::Get().ShapesFaceList[i + tableOffset][2]];
b[1] = ypoints[ShapesFaceList[i + tableOffset][2]]; b[1] = ypoints[CellTables::Get().ShapesFaceList[i + tableOffset][2]];
b[2] = zpoints[ShapesFaceList[i + tableOffset][2]]; b[2] = zpoints[CellTables::Get().ShapesFaceList[i + tableOffset][2]];
c[0] = xpoints[ShapesFaceList[i + tableOffset][3]]; c[0] = xpoints[CellTables::Get().ShapesFaceList[i + tableOffset][3]];
c[1] = ypoints[ShapesFaceList[i + tableOffset][3]]; c[1] = ypoints[CellTables::Get().ShapesFaceList[i + tableOffset][3]];
c[2] = zpoints[ShapesFaceList[i + tableOffset][3]]; c[2] = zpoints[CellTables::Get().ShapesFaceList[i + tableOffset][3]];
d[0] = xpoints[ShapesFaceList[i + tableOffset][4]]; d[0] = xpoints[CellTables::Get().ShapesFaceList[i + tableOffset][4]];
d[1] = ypoints[ShapesFaceList[i + tableOffset][4]]; d[1] = ypoints[CellTables::Get().ShapesFaceList[i + tableOffset][4]];
d[2] = zpoints[ShapesFaceList[i + tableOffset][4]]; d[2] = zpoints[CellTables::Get().ShapesFaceList[i + tableOffset][4]];
T distance = -1.f; T distance = -1.f;
distances[i] = distance; //init to -1 distances[i] = distance; //init to -1

186
vtkm/rendering/raytracing/CellTables.h
View File

@ -30,100 +30,114 @@ namespace rendering
namespace raytracing namespace raytracing
{ {
//LookUp of Shapes to FaceLookUp class CellTables
VTKM_EXEC_CONSTANT {
static vtkm::Int32 CellTypeLookUp[15] = { private:
4, // 0 Nothing struct Tables
4, // 1 Vertex {
4, // 2 (Not Used) Poly Vertex vtkm::Int32 CellTypeLookUp[15];
4, // 3 Line vtkm::Int32 FaceLookUp[5][3];
4, // 4 (Not Used) Poly Line vtkm::Int32 ShapesFaceList[20][5];
4, // 5 Triangle vtkm::Int32 ZooTable[30][4];
4, // 6 (not used) triangle strip vtkm::Int32 ZooLookUp[5][2];
4, // 7 Polygon };
4, // 8 (Not used)Pixel
4, // 9 Quad
1, // 10 Tetra
4, // 11 (Not used) Voxel
0, // 12 Hex
2, // 13 Wedge
3 // 14 Pyramid
};
VTKM_EXEC_CONSTANT public:
static vtkm::Int32 FaceLookUp[5][3] = { VTKM_EXEC_CONT static const Tables& Get()
{ 0, 6, 8 }, //hex offset into shapes face list, num faces and number of Indices {
{ 6, 4, 4 }, //tet static const Tables tables = {
{ 10, 5, 6 }, //wedge // CellTypeLookUp: LookUp of Shapes to FaceLookUp
{ 15, 5, 5 }, //pyramid {
{ -1, 0, 0 } //unsupported shape 4, // 0 Nothing
}; 4, // 1 Vertex
4, // 2 (Not Used) Poly Vertex
4, // 3 Line
4, // 4 (Not Used) Poly Line
4, // 5 Triangle
4, // 6 (not used) triangle strip
4, // 7 Polygon
4, // 8 (Not used)Pixel
4, // 9 Quad
1, // 10 Tetra
4, // 11 (Not used) Voxel
0, // 12 Hex
2, // 13 Wedge
3 // 14 Pyramid
},
// The convention for the faces is that looking from the outside of // FaceLookUp
// the shape at a face, triangles should wind CCW. {
// Quads are broken up by {4=quad,a,b,c,d}: { 0, 6, 8 }, //hex offset into shapes face list, num faces and number of Indices
// t1 = abc and t2 = acd. Indices of the face are ordered CW, and the mapping { 6, 4, 4 }, //tet
// of t1 and t2 become CCW. { 10, 5, 6 }, //wedge
// Since we know the triangle winding, we could tell { 15, 5, 5 }, //pyramid
// if we hit an inside face or outside face. { -1, 0, 0 } //unsupported shape
VTKM_EXEC_CONSTANT },
static vtkm::Int32 ShapesFaceList[20][5] = {
//hex
{ 4, 0, 1, 5, 4 }, //face 0
{ 4, 1, 2, 6, 5 },
{ 4, 3, 7, 6, 2 },
{ 4, 0, 4, 7, 3 },
{ 4, 0, 3, 2, 1 },
{ 4, 4, 5, 6, 7 }, //face 5
//tet // ShapesFaceList:
{ 3, 0, 3, 1, -1 }, // The convention for the faces is that looking from the outside of
{ 3, 1, 2, 3, -1 }, // the shape at a face, triangles should wind CCW.
{ 3, 0, 2, 3, -1 }, // Quads are broken up by {4=quad,a,b,c,d}:
{ 3, 0, 2, 1, -1 }, // t1 = abc and t2 = acd. Indices of the face are ordered CW, and the mapping
// of t1 and t2 become CCW.
// Since we know the triangle winding, we could tell
// if we hit an inside face or outside face.
{ //hex
{ 4, 0, 1, 5, 4 }, //face 0
{ 4, 1, 2, 6, 5 },
{ 4, 3, 7, 6, 2 },
{ 4, 0, 4, 7, 3 },
{ 4, 0, 3, 2, 1 },
{ 4, 4, 5, 6, 7 }, //face 5
//wedge //tet
{ 3, 0, 1, 2, -1 }, { 3, 0, 3, 1, -1 },
{ 3, 3, 5, 4, -1 }, { 3, 1, 2, 3, -1 },
{ 4, 3, 0, 2, 5 }, { 3, 0, 2, 3, -1 },
{ 4, 1, 4, 5, 2 }, { 3, 0, 2, 1, -1 },
{ 4, 0, 3, 4, 1 },
//pyramid //wedge
{ 3, 0, 4, 1, -1 }, { 3, 0, 1, 2, -1 },
{ 3, 1, 2, 4, -1 }, { 3, 3, 5, 4, -1 },
{ 3, 2, 3, 4, -1 }, { 4, 3, 0, 2, 5 },
{ 3, 0, 4, 3, -1 }, { 4, 1, 4, 5, 2 },
{ 4, 3, 2, 1, 0 } { 4, 0, 3, 4, 1 },
}; //pyramid
{ 3, 0, 4, 1, -1 },
{ 3, 1, 2, 4, -1 },
{ 3, 2, 3, 4, -1 },
{ 3, 0, 4, 3, -1 },
{ 4, 3, 2, 1, 0 } },
// Test of zoo table. // ZooTable:
// Format (faceNumber, triangle) // Test of zoo table.
// // Format (faceNumber, triangle)
VTKM_EXEC_CONSTANT //
static vtkm::Int32 ZooTable[30][4] = { { { 0, 0, 1, 5 }, // hex
{ 0, 0, 1, 5 }, // hex { 0, 0, 5, 4 }, { 1, 1, 2, 6 }, { 1, 1, 6, 5 }, { 2, 3, 7, 6 }, { 2, 3, 6, 2 },
{ 0, 0, 5, 4 }, { 1, 1, 2, 6 }, { 1, 1, 6, 5 }, { 2, 3, 7, 6 }, { 2, 3, 6, 2 }, { 3, 0, 4, 7 }, { 3, 0, 7, 3 }, { 4, 0, 3, 2 }, { 4, 0, 2, 1 }, { 5, 4, 5, 6 },
{ 3, 0, 4, 7 }, { 3, 0, 7, 3 }, { 4, 0, 3, 2 }, { 4, 0, 2, 1 }, { 5, 4, 5, 6 }, { 5, 4, 6, 7 }, { 0, 0, 3, 1 }, // Tet
{ 5, 4, 6, 7 }, { 0, 0, 3, 1 }, // Tet { 1, 1, 2, 3 }, { 2, 0, 2, 3 }, { 3, 0, 2, 1 }, { 0, 0, 1, 2 }, // Wedge
{ 1, 1, 2, 3 }, { 2, 0, 2, 3 }, { 3, 0, 2, 1 }, { 0, 0, 1, 2 }, // Wedge { 1, 3, 5, 4 }, { 2, 3, 0, 2 }, { 2, 3, 2, 5 }, { 3, 1, 4, 5 }, { 3, 1, 5, 2 },
{ 1, 3, 5, 4 }, { 2, 3, 0, 2 }, { 2, 3, 2, 5 }, { 3, 1, 4, 5 }, { 3, 1, 5, 2 }, { 4, 0, 3, 4 }, { 4, 0, 4, 1 }, { 0, 0, 4, 1 }, // Pyramid
{ 4, 0, 3, 4 }, { 4, 0, 4, 1 }, { 0, 0, 4, 1 }, // Pyramid { 1, 1, 2, 4 }, { 2, 2, 3, 4 }, { 3, 0, 4, 3 }, { 4, 3, 2, 1 }, { 4, 3, 1, 0 } },
{ 1, 1, 2, 4 }, { 2, 2, 3, 4 }, { 3, 0, 4, 3 }, { 4, 3, 2, 1 }, { 4, 3, 1, 0 }
};
// // ZooLookUp:
// Offset into zoo table and the // Offset into zoo table and the
// number of triangles for the shape // number of triangles for the shape
// //
VTKM_EXEC_CONSTANT {
static vtkm::Int32 ZooLookUp[5][2] = { { 0, 12 }, //hex offset into shapes face list, num faces and number of Indices
{ 0, 12 }, //hex offset into shapes face list, num faces and number of Indices { 12, 4 }, //tet
{ 12, 4 }, //tet { 16, 8 }, //wedge
{ 16, 8 }, //wedge { 24, 6 }, //pyramid
{ 24, 6 }, //pyramid { -1, 0 } //unsupported shape
{ -1, 0 } //unsupported shape }
};
return tables;
}
}; };
} // namespace raytracing } // namespace raytracing

6
vtkm/rendering/raytracing/ConnectivityTracer.h
View File

@ -281,7 +281,8 @@ public:
} }
const vtkm::UInt8 cellShape = MeshConn.GetCellShape(currentCell); const vtkm::UInt8 cellShape = MeshConn.GetCellShape(currentCell);
Intersector.IntersectCell(xpoints, ypoints, zpoints, dir, origin, distances, cellShape); Intersector.IntersectCell(xpoints, ypoints, zpoints, dir, origin, distances, cellShape);
const vtkm::Int32 numFaces = FaceLookUp[CellTypeLookUp[cellShape]][1]; const vtkm::Int32 numFaces =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[cellShape]][1];
//vtkm::Int32 minFace = 6; //vtkm::Int32 minFace = 6;
vtkm::Int32 maxFace = -1; vtkm::Int32 maxFace = -1;
@ -415,7 +416,8 @@ public:
const vtkm::UInt8 cellShape = MeshConn.GetCellShape(currentCell); const vtkm::UInt8 cellShape = MeshConn.GetCellShape(currentCell);
Intersector.IntersectCell(xpoints, ypoints, zpoints, dir, origin, distances, cellShape); Intersector.IntersectCell(xpoints, ypoints, zpoints, dir, origin, distances, cellShape);
const vtkm::Int32 numFaces = FaceLookUp[CellTypeLookUp[cellShape]][1]; const vtkm::Int32 numFaces =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[cellShape]][1];
//vtkm::Int32 minFace = 6; //vtkm::Int32 minFace = 6;
vtkm::Int32 maxFace = -1; vtkm::Int32 maxFace = -1;

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

@ -110,19 +110,19 @@ public:
vtkm::Int32 tableOffset = 0; vtkm::Int32 tableOffset = 0;
if (shapeType == vtkm::CELL_SHAPE_TETRA) if (shapeType == vtkm::CELL_SHAPE_TETRA)
{ {
tableOffset = FaceLookUp[1][0]; tableOffset = CellTables::Get().FaceLookUp[1][0];
} }
else if (shapeType == vtkm::CELL_SHAPE_HEXAHEDRON) else if (shapeType == vtkm::CELL_SHAPE_HEXAHEDRON)
{ {
tableOffset = FaceLookUp[0][0]; tableOffset = CellTables::Get().FaceLookUp[0][0];
} }
else if (shapeType == vtkm::CELL_SHAPE_WEDGE) else if (shapeType == vtkm::CELL_SHAPE_WEDGE)
{ {
tableOffset = FaceLookUp[2][0]; tableOffset = CellTables::Get().FaceLookUp[2][0];
} }
else if (shapeType == vtkm::CELL_SHAPE_PYRAMID) else if (shapeType == vtkm::CELL_SHAPE_PYRAMID)
{ {
tableOffset = FaceLookUp[3][0]; tableOffset = CellTables::Get().FaceLookUp[3][0];
} }
else else
printf("Error shape not recognized %d\n", (int)shapeType); printf("Error shape not recognized %d\n", (int)shapeType);
@ -190,8 +190,8 @@ public:
vtkm::Id shape2Offset = vtkm::Id shape2Offset =
GetShapeOffset(shapes.Get(cellId2)) + faceIdPairs.Get(currentIndex)[1]; GetShapeOffset(shapes.Get(cellId2)) + faceIdPairs.Get(currentIndex)[1];
vtkm::Int32 icount1 = ShapesFaceList[shape1Offset][0]; vtkm::Int32 icount1 = CellTables::Get().ShapesFaceList[shape1Offset][0];
vtkm::Int32 icount2 = ShapesFaceList[shape2Offset][0]; vtkm::Int32 icount2 = CellTables::Get().ShapesFaceList[shape2Offset][0];
//Check to see if we have the same number of idices //Check to see if we have the same number of idices
if (icount1 != icount2) if (icount1 != icount2)
continue; continue;
@ -201,22 +201,24 @@ public:
vtkm::Vec<vtkm::Id, 4> indices1; vtkm::Vec<vtkm::Id, 4> indices1;
vtkm::Vec<vtkm::Id, 4> indices2; vtkm::Vec<vtkm::Id, 4> indices2;
for (vtkm::Int32 i = 1; i <= ShapesFaceList[shape1Offset][0]; ++i) for (vtkm::Int32 i = 1; i <= CellTables::Get().ShapesFaceList[shape1Offset][0]; ++i)
{ {
BOUNDS_CHECK(offsets, cellId1); BOUNDS_CHECK(offsets, cellId1);
BOUNDS_CHECK(offsets, cellId2); BOUNDS_CHECK(offsets, cellId2);
BOUNDS_CHECK(connectivity, (offsets.Get(cellId1) + ShapesFaceList[shape1Offset][i])); BOUNDS_CHECK(connectivity,
BOUNDS_CHECK(connectivity, (offsets.Get(cellId2) + ShapesFaceList[shape2Offset][i])); (offsets.Get(cellId1) + CellTables::Get().ShapesFaceList[shape1Offset][i]));
indices1[i - 1] = BOUNDS_CHECK(connectivity,
connectivity.Get(offsets.Get(cellId1) + ShapesFaceList[shape1Offset][i]); (offsets.Get(cellId2) + CellTables::Get().ShapesFaceList[shape2Offset][i]));
indices2[i - 1] = indices1[i - 1] = connectivity.Get(offsets.Get(cellId1) +
connectivity.Get(offsets.Get(cellId2) + ShapesFaceList[shape2Offset][i]); CellTables::Get().ShapesFaceList[shape1Offset][i]);
indices2[i - 1] = connectivity.Get(offsets.Get(cellId2) +
CellTables::Get().ShapesFaceList[shape2Offset][i]);
} }
bool isEqual = true; bool isEqual = true;
for (vtkm::Int32 i = 0; i < ShapesFaceList[shape1Offset][0]; ++i) for (vtkm::Int32 i = 0; i < CellTables::Get().ShapesFaceList[shape1Offset][0]; ++i)
{ {
if (!IsIn(indices1[i], indices2, ShapesFaceList[shape1Offset][0])) if (!IsIn(indices1[i], indices2, CellTables::Get().ShapesFaceList[shape1Offset][0]))
isEqual = false; isEqual = false;
} }
@ -285,7 +287,8 @@ public:
vtkm::Id offset = shapeOffsets.Get(cellId); vtkm::Id offset = shapeOffsets.Get(cellId);
BOUNDS_CHECK(shapes, cellId); BOUNDS_CHECK(shapes, cellId);
vtkm::Int32 shapeId = static_cast<vtkm::Int32>(shapes.Get(cellId)); vtkm::Int32 shapeId = static_cast<vtkm::Int32>(shapes.Get(cellId));
vtkm::Int32 shapesFaceOffset = FaceLookUp[CellTypeLookUp[shapeId]][0]; vtkm::Int32 shapesFaceOffset =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[shapeId]][0];
if (shapesFaceOffset == -1) if (shapesFaceOffset == -1)
{ {
printf("Unsupported Shape Type %d\n", shapeId); printf("Unsupported Shape Type %d\n", shapeId);
@ -294,12 +297,12 @@ public:
vtkm::Vec<vtkm::Id, 4> faceIndices(-1, -1, -1, -1); vtkm::Vec<vtkm::Id, 4> faceIndices(-1, -1, -1, -1);
vtkm::Int32 tableIndex = static_cast<vtkm::Int32>(shapesFaceOffset + faceIdPair[1]); vtkm::Int32 tableIndex = static_cast<vtkm::Int32>(shapesFaceOffset + faceIdPair[1]);
const vtkm::Int32 numIndices = ShapesFaceList[tableIndex][0]; const vtkm::Int32 numIndices = CellTables::Get().ShapesFaceList[tableIndex][0];
for (vtkm::Int32 i = 1; i <= numIndices; ++i) for (vtkm::Int32 i = 1; i <= numIndices; ++i)
{ {
BOUNDS_CHECK(indices, offset + ShapesFaceList[tableIndex][i]); BOUNDS_CHECK(indices, offset + CellTables::Get().ShapesFaceList[tableIndex][i]);
faceIndices[i - 1] = indices.Get(offset + ShapesFaceList[tableIndex][i]); faceIndices[i - 1] = indices.Get(offset + CellTables::Get().ShapesFaceList[tableIndex][i]);
} }
vtkm::Vec<vtkm::Id, 4> triangle; vtkm::Vec<vtkm::Id, 4> triangle;
triangle[0] = cellId; triangle[0] = cellId;
@ -344,18 +347,6 @@ public:
} }
}; //class WriteFaceConn }; //class WriteFaceConn
// Each one of size segments will process
// one face of the hex and domain
VTKM_EXEC_CONSTANT
static vtkm::Int32 SegmentToFace[6] = { 0, 2, 1, 3, 4, 5 };
// Each face/segment has 2 varying dimensions
VTKM_EXEC_CONSTANT
static vtkm::Int32 SegmentDirections[6][2] = { { 0, 2 }, //face 0 and 2 have the same directions
{ 0, 2 }, { 1, 2 }, //1 and 3
{ 1, 2 }, { 0, 1 }, // 4 and 5
{ 0, 1 } };
class StructuredExternalTriangles : public vtkm::worklet::WorkletMapField class StructuredExternalTriangles : public vtkm::worklet::WorkletMapField
{ {
protected: protected:
@ -395,6 +386,18 @@ public:
template <typename TrianglePortalType> template <typename TrianglePortalType>
VTKM_EXEC inline void operator()(const vtkm::Id& index, TrianglePortalType& triangles) const VTKM_EXEC inline void operator()(const vtkm::Id& index, TrianglePortalType& triangles) const
{ {
// Each one of size segments will process
// one face of the hex and domain
static const vtkm::Int32 SegmentToFace[6] = { 0, 2, 1, 3, 4, 5 };
// Each face/segment has 2 varying dimensions
static const vtkm::Int32 SegmentDirections[6][2] = {
{ 0, 2 }, //face 0 and 2 have the same directions
{ 0, 2 }, { 1, 2 }, //1 and 3
{ 1, 2 }, { 0, 1 }, // 4 and 5
{ 0, 1 }
};
// //
// We get one index per extenal face // We get one index per extenal face
// //
@ -449,7 +452,8 @@ public:
// Look up the offset into the face list for each cell type // Look up the offset into the face list for each cell type
// This should always be zero, but in case this table changes I don't // This should always be zero, but in case this table changes I don't
// want to break anything. // want to break anything.
vtkm::Int32 shapesFaceOffset = FaceLookUp[CellTypeLookUp[CELL_SHAPE_HEXAHEDRON]][0]; vtkm::Int32 shapesFaceOffset =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[CELL_SHAPE_HEXAHEDRON]][0];
vtkm::Vec<vtkm::Id, 4> faceIndices; vtkm::Vec<vtkm::Id, 4> faceIndices;
vtkm::Int32 tableIndex = shapesFaceOffset + cellFace; vtkm::Int32 tableIndex = shapesFaceOffset + cellFace;
@ -457,7 +461,7 @@ public:
// Load the face // Load the face
for (vtkm::Int32 i = 1; i <= 4; ++i) for (vtkm::Int32 i = 1; i <= 4; ++i)
{ {
faceIndices[i - 1] = cellIndices[ShapesFaceList[tableIndex][i]]; faceIndices[i - 1] = cellIndices[CellTables::Get().ShapesFaceList[tableIndex][i]];
} }
const vtkm::Id outputOffset = index * 2; const vtkm::Id outputOffset = index * 2;
vtkm::Vec<vtkm::Id, 4> triangle; vtkm::Vec<vtkm::Id, 4> triangle;
@ -492,14 +496,15 @@ public:
vtkm::Id cellId = faceIdPair[0]; vtkm::Id cellId = faceIdPair[0];
vtkm::Id cellFace = faceIdPair[1]; vtkm::Id cellFace = faceIdPair[1];
vtkm::Int32 shapeType = static_cast<vtkm::Int32>(shapes.Get(cellId)); vtkm::Int32 shapeType = static_cast<vtkm::Int32>(shapes.Get(cellId));
vtkm::Int32 faceStartIndex = FaceLookUp[CellTypeLookUp[shapeType]][0]; vtkm::Int32 faceStartIndex =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[shapeType]][0];
if (faceStartIndex == -1) if (faceStartIndex == -1)
{ {
//Unsupported Shape Type this should never happen //Unsupported Shape Type this should never happen
triangleCount = 0; triangleCount = 0;
return; return;
} }
vtkm::Int32 faceType = ShapesFaceList[faceStartIndex + cellFace][0]; vtkm::Int32 faceType = CellTables::Get().ShapesFaceList[faceStartIndex + cellFace][0];
// The face will either have 4 or 3 indices, so quad or tri // The face will either have 4 or 3 indices, so quad or tri
triangleCount = (faceType == 4) ? 2 : 1; triangleCount = (faceType == 4) ? 2 : 1;

7
vtkm/rendering/raytracing/MeshConnectivityStructures.h
View File

@ -278,7 +278,8 @@ public:
inline vtkm::Int32 GetCellIndices(vtkm::Id cellIndices[8], const vtkm::Id& cellId) const inline vtkm::Int32 GetCellIndices(vtkm::Id cellIndices[8], const vtkm::Id& cellId) const
{ {
const vtkm::Int32 shapeId = static_cast<vtkm::Int32>(ShapesPortal.Get(cellId)); const vtkm::Int32 shapeId = static_cast<vtkm::Int32>(ShapesPortal.Get(cellId));
const vtkm::Int32 numIndices = FaceLookUp[CellTypeLookUp[shapeId]][2]; const vtkm::Int32 numIndices =
CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[shapeId]][2];
BOUNDS_CHECK(CellOffsetsPortal, cellId); BOUNDS_CHECK(CellOffsetsPortal, cellId);
const vtkm::Id cellOffset = CellOffsetsPortal.Get(cellId); const vtkm::Id cellOffset = CellOffsetsPortal.Get(cellId);
@ -367,7 +368,7 @@ public:
Cellset.GetShapesArray(vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell()); Cellset.GetShapesArray(vtkm::TopologyElementTagPoint(), vtkm::TopologyElementTagCell());
ShapeId = shapes.GetPortalConstControl().Get(0); ShapeId = shapes.GetPortalConstControl().Get(0);
NumIndices = FaceLookUp[CellTypeLookUp[ShapeId]][2]; NumIndices = CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[ShapeId]][2];
if (NumIndices == 0) if (NumIndices == 0)
{ {
@ -377,7 +378,7 @@ public:
throw vtkm::cont::ErrorBadValue(message.str()); throw vtkm::cont::ErrorBadValue(message.str());
} }
vtkm::Id start = 0; vtkm::Id start = 0;
NumFaces = FaceLookUp[CellTypeLookUp[ShapeId]][1]; NumFaces = CellTables::Get().FaceLookUp[CellTables::Get().CellTypeLookUp[ShapeId]][1];
vtkm::Id numCells = CellConnectivity.GetPortalConstControl().GetNumberOfValues(); vtkm::Id numCells = CellConnectivity.GetPortalConstControl().GetNumberOfValues();
CellOffsets = vtkm::cont::make_ArrayHandleCounting<vtkm::Id>(start, NumIndices, numCells); CellOffsets = vtkm::cont::make_ArrayHandleCounting<vtkm::Id>(start, NumIndices, numCells);

26
vtkm/rendering/raytracing/MortonCodes.h
View File

@ -172,30 +172,26 @@ public:
{ {
vtkm::Int32 faceCount; vtkm::Int32 faceCount;
vtkm::Int32 tableOffset; vtkm::Int32 tableOffset;
// suppress some unused variables warnings
(void)(CellTypeLookUp);
(void)(ZooTable);
(void)(ZooLookUp);
if (cellShape.Id == vtkm::CELL_SHAPE_TETRA) if (cellShape.Id == vtkm::CELL_SHAPE_TETRA)
{ {
faceCount = FaceLookUp[1][1]; faceCount = CellTables::Get().FaceLookUp[1][1];
tableOffset = FaceLookUp[1][0]; tableOffset = CellTables::Get().FaceLookUp[1][0];
} }
else if (cellShape.Id == vtkm::CELL_SHAPE_HEXAHEDRON) else if (cellShape.Id == vtkm::CELL_SHAPE_HEXAHEDRON)
{ {
faceCount = FaceLookUp[0][1]; faceCount = CellTables::Get().FaceLookUp[0][1];
tableOffset = FaceLookUp[0][0]; tableOffset = CellTables::Get().FaceLookUp[0][0];
} }
else if (cellShape.Id == vtkm::CELL_SHAPE_WEDGE) else if (cellShape.Id == vtkm::CELL_SHAPE_WEDGE)
{ {
faceCount = FaceLookUp[2][1]; faceCount = CellTables::Get().FaceLookUp[2][1];
tableOffset = FaceLookUp[2][0]; tableOffset = CellTables::Get().FaceLookUp[2][0];
} }
else if (cellShape.Id == vtkm::CELL_SHAPE_PYRAMID) else if (cellShape.Id == vtkm::CELL_SHAPE_PYRAMID)
{ {
faceCount = FaceLookUp[3][1]; faceCount = CellTables::Get().FaceLookUp[3][1];
tableOffset = FaceLookUp[3][0]; tableOffset = CellTables::Get().FaceLookUp[3][0];
} }
else else
{ {
@ -218,11 +214,11 @@ public:
// we currently check for. // we currently check for.
vtkm::Vec<vtkm::Id, 4> faceIndices; vtkm::Vec<vtkm::Id, 4> faceIndices;
faceIndices[3] = -1; faceIndices[3] = -1;
const vtkm::Id indiceCount = //Number of indices this face has
ShapesFaceList[tableOffset + i][0]; //Number of indices this face has const vtkm::Id indiceCount = CellTables::Get().ShapesFaceList[tableOffset + i][0];
for (vtkm::Int32 j = 1; j <= indiceCount; j++) for (vtkm::Int32 j = 1; j <= indiceCount; j++)
{ {
faceIndices[j - 1] = cellIndices[ShapesFaceList[tableOffset + i][j]]; faceIndices[j - 1] = cellIndices[CellTables::Get().ShapesFaceList[tableOffset + i][j]];
} }
//sort the indices in descending order //sort the indices in descending order
Sort4(faceIndices); Sort4(faceIndices);

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

@ -38,8 +38,11 @@ namespace raytracing
namespace namespace
{ {
VTKM_EXEC_CONSTANT static vtkm::Int32 END_FLAG2 = -1000000000;
VTKM_EXEC_CONSTANT static vtkm::Float32 EPSILON2 = 0.0001f; enum : vtkm::Int32
{
END_FLAG2 = -1000000000
};
} }
template <typename TriIntersector> template <typename TriIntersector>
@ -99,6 +102,8 @@ public:
const Precision& originY, const Precision& originY,
const Precision& originZ) const const Precision& originZ) const
{ {
const vtkm::Float32 EPSILON2 = 0.0001f;
vtkm::Vec<Precision, 3> e1 = b - a; vtkm::Vec<Precision, 3> e1 = b - a;
vtkm::Vec<Precision, 3> e2 = c - a; vtkm::Vec<Precision, 3> e2 = c - a;
@ -774,7 +779,7 @@ public:
vtkm::Int32 todo[64]; vtkm::Int32 todo[64];
vtkm::Int32 stackptr = 0; vtkm::Int32 stackptr = 0;
vtkm::Int32 barrier = (vtkm::Int32)END_FLAG2; vtkm::Int32 barrier = END_FLAG2;
currentNode = 0; currentNode = 0;
todo[stackptr] = barrier; todo[stackptr] = barrier;
@ -923,7 +928,7 @@ public:
vtkm::Int32 todo[64]; vtkm::Int32 todo[64];
vtkm::Int32 stackptr = 0; vtkm::Int32 stackptr = 0;
vtkm::Int32 barrier = (vtkm::Int32)END_FLAG2; vtkm::Int32 barrier = END_FLAG2;
currentNode = 0; currentNode = 0;
todo[stackptr] = barrier; todo[stackptr] = barrier;

202
vtkm/testing/TestingMath.h
View File

@ -43,56 +43,88 @@
namespace UnitTestMathNamespace namespace UnitTestMathNamespace
{ {
const vtkm::IdComponent NUM_NUMBERS = 5; class Lists
VTKM_EXEC_CONSTANT {
const vtkm::Float64 NumberList[NUM_NUMBERS] = { 0.25, 0.5, 1.0, 2.0, 3.75 }; public:
static const vtkm::IdComponent NUM_NUMBERS = 5;
VTKM_EXEC_CONSTANT private:
const vtkm::Float64 AngleList[NUM_NUMBERS] = { 0.643501108793284, // angle for 3, 4, 5 triangle. struct Data
0.78539816339745, // pi/4 {
0.5235987755983, // pi/6 vtkm::Float64 NumberList[NUM_NUMBERS];
1.0471975511966, // pi/3 vtkm::Float64 AngleList[NUM_NUMBERS];
0.0 }; vtkm::Float64 OppositeList[NUM_NUMBERS];
VTKM_EXEC_CONSTANT vtkm::Float64 AdjacentList[NUM_NUMBERS];
const vtkm::Float64 OppositeList[NUM_NUMBERS] = { 3.0, vtkm::Float64 HypotenuseList[NUM_NUMBERS];
1.0, vtkm::Float64 NumeratorList[NUM_NUMBERS];
1.0, vtkm::Float64 DenominatorList[NUM_NUMBERS];
1.732050807568877 /*sqrt(3)*/, vtkm::Float64 FModRemainderList[NUM_NUMBERS];
0.0 }; vtkm::Float64 RemainderList[NUM_NUMBERS];
VTKM_EXEC_CONSTANT vtkm::Int64 QuotientList[NUM_NUMBERS];
const vtkm::Float64 AdjacentList[NUM_NUMBERS] = { 4.0, vtkm::Float64 XList[NUM_NUMBERS];
1.0, vtkm::Float64 FractionalList[NUM_NUMBERS];
1.732050807568877 /*sqrt(3)*/, vtkm::Float64 FloorList[NUM_NUMBERS];
1.0, vtkm::Float64 CeilList[NUM_NUMBERS];
1.0 }; vtkm::Float64 RoundList[NUM_NUMBERS];
VTKM_EXEC_CONSTANT };
const vtkm::Float64 HypotenuseList[NUM_NUMBERS] = { 5.0,
1.414213562373095 /*sqrt(2)*/,
2.0,
2.0,
1.0 };
VTKM_EXEC_CONSTANT public:
const vtkm::Float64 NumeratorList[NUM_NUMBERS] = { 6.5, 5.8, 9.3, 77.0, 0.1 }; VTKM_EXEC_CONT static const Data& Get()
VTKM_EXEC_CONSTANT {
const vtkm::Float64 DenominatorList[NUM_NUMBERS] = { 2.3, 1.6, 3.1, 19.0, 0.4 }; static const Data data = {
VTKM_EXEC_CONSTANT // NumberList
const vtkm::Float64 FModRemainderList[NUM_NUMBERS] = { 1.9, 1.0, 0.0, 1.0, 0.1 }; { 0.25, 0.5, 1.0, 2.0, 3.75 },
VTKM_EXEC_CONSTANT
const vtkm::Float64 RemainderList[NUM_NUMBERS] = { -0.4, -0.6, 0.0, 1.0, 0.1 };
VTKM_EXEC_CONSTANT
const vtkm::Int64 QuotientList[NUM_NUMBERS] = { 3, 4, 3, 4, 0 };
VTKM_EXEC_CONSTANT // AngleList
const vtkm::Float64 XList[NUM_NUMBERS] = { 4.6, 0.1, 73.4, 55.0, 3.75 }; { 0.643501108793284, // angle for 3, 4, 5 triangle.
VTKM_EXEC_CONSTANT 0.78539816339745, // pi/4
const vtkm::Float64 FractionalList[NUM_NUMBERS] = { 0.6, 0.1, 0.4, 0.0, 0.75 }; 0.5235987755983, // pi/6
VTKM_EXEC_CONSTANT 1.0471975511966, // pi/3
const vtkm::Float64 FloorList[NUM_NUMBERS] = { 4.0, 0.0, 73.0, 55.0, 3.0 }; 0.0 },
VTKM_EXEC_CONSTANT
const vtkm::Float64 CeilList[NUM_NUMBERS] = { 5.0, 1.0, 74.0, 55.0, 4.0 }; // OppositeList
VTKM_EXEC_CONSTANT { 3.0, 1.0, 1.0, 1.732050807568877 /*sqrt(3)*/, 0.0 },
const vtkm::Float64 RoundList[NUM_NUMBERS] = { 5.0, 0.0, 73.0, 55.0, 4.0 };
// AdjacentList
{ 4.0, 1.0, 1.732050807568877 /*sqrt(3)*/, 1.0, 1.0 },
// HypotenuseList
{ 5.0, 1.414213562373095 /*sqrt(2)*/, 2.0, 2.0, 1.0 },
// NumeratorList
{ 6.5, 5.8, 9.3, 77.0, 0.1 },
// DenominatorList
{ 2.3, 1.6, 3.1, 19.0, 0.4 },
// FModRemainderList
{ 1.9, 1.0, 0.0, 1.0, 0.1 },
// RemainderList
{ -0.4, -0.6, 0.0, 1.0, 0.1 },
// QuotientList
{ 3, 4, 3, 4, 0 },
// XList
{ 4.6, 0.1, 73.4, 55.0, 3.75 },
// FractionalList
{ 0.6, 0.1, 0.4, 0.0, 0.75 },
// FloorList
{ 4.0, 0.0, 73.0, 55.0, 3.0 },
// CeilList
{ 5.0, 1.0, 74.0, 55.0, 4.0 },
// RoundList
{ 5.0, 0.0, 73.0, 55.0, 4.0 },
};
return data;
}
};
//----------------------------------------------------------------------------- //-----------------------------------------------------------------------------
template <typename T> template <typename T>
@ -130,9 +162,9 @@ struct ScalarFieldTests : public vtkm::exec::FunctorBase
void TestPow() const void TestPow() const
{ {
// std::cout << "Running power tests." << std::endl; // std::cout << "Running power tests." << std::endl;
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS; index++)
{ {
T x = static_cast<T>(NumberList[index]); T x = static_cast<T>(Lists::Get().NumberList[index]);
T powx = vtkm::Pow(x, static_cast<T>(2.0)); T powx = vtkm::Pow(x, static_cast<T>(2.0));
T sqrx = x * x; T sqrx = x * x;
VTKM_MATH_ASSERT(test_equal(powx, sqrx), "Power gave wrong result."); VTKM_MATH_ASSERT(test_equal(powx, sqrx), "Power gave wrong result.");
@ -206,13 +238,13 @@ struct ScalarFieldTests : public vtkm::exec::FunctorBase
void TestRemainders() const void TestRemainders() const
{ {
// std::cout << "Testing remainders." << std::endl; // std::cout << "Testing remainders." << std::endl;
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS; index++)
{ {
T numerator = static_cast<T>(NumeratorList[index]); T numerator = static_cast<T>(Lists::Get().NumeratorList[index]);
T denominator = static_cast<T>(DenominatorList[index]); T denominator = static_cast<T>(Lists::Get().DenominatorList[index]);
T fmodremainder = static_cast<T>(FModRemainderList[index]); T fmodremainder = static_cast<T>(Lists::Get().FModRemainderList[index]);
T remainder = static_cast<T>(RemainderList[index]); T remainder = static_cast<T>(Lists::Get().RemainderList[index]);
vtkm::Int64 quotient = QuotientList[index]; vtkm::Int64 quotient = Lists::Get().QuotientList[index];
VTKM_MATH_ASSERT(test_equal(vtkm::FMod(numerator, denominator), fmodremainder), VTKM_MATH_ASSERT(test_equal(vtkm::FMod(numerator, denominator), fmodremainder),
"Bad FMod remainder."); "Bad FMod remainder.");
@ -229,13 +261,13 @@ struct ScalarFieldTests : public vtkm::exec::FunctorBase
void TestRound() const void TestRound() const
{ {
// std::cout << "Testing round." << std::endl; // std::cout << "Testing round." << std::endl;
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS; index++)
{ {
T x = static_cast<T>(XList[index]); T x = static_cast<T>(Lists::Get().XList[index]);
T fractional = static_cast<T>(FractionalList[index]); T fractional = static_cast<T>(Lists::Get().FractionalList[index]);
T floor = static_cast<T>(FloorList[index]); T floor = static_cast<T>(Lists::Get().FloorList[index]);
T ceil = static_cast<T>(CeilList[index]); T ceil = static_cast<T>(Lists::Get().CeilList[index]);
T round = static_cast<T>(RoundList[index]); T round = static_cast<T>(Lists::Get().RoundList[index]);
T intPart; T intPart;
VTKM_MATH_ASSERT(test_equal(vtkm::ModF(x, intPart), fractional), VTKM_MATH_ASSERT(test_equal(vtkm::ModF(x, intPart), fractional),
@ -312,7 +344,7 @@ struct ScalarVectorFieldTests : public vtkm::exec::FunctorBase
{ {
// std::cout << "Testing normal trig functions." << std::endl; // std::cout << "Testing normal trig functions." << std::endl;
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS - NUM_COMPONENTS + 1; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS - NUM_COMPONENTS + 1; index++)
{ {
VectorType angle; VectorType angle;
VectorType opposite; VectorType opposite;
@ -321,16 +353,21 @@ struct ScalarVectorFieldTests : public vtkm::exec::FunctorBase
for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++) for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++)
{ {
Traits::SetComponent( Traits::SetComponent(
angle, componentIndex, static_cast<ComponentType>(AngleList[componentIndex + index])); angle,
Traits::SetComponent(opposite, componentIndex,
componentIndex, static_cast<ComponentType>(Lists::Get().AngleList[componentIndex + index]));
static_cast<ComponentType>(OppositeList[componentIndex + index])); Traits::SetComponent(
Traits::SetComponent(adjacent, opposite,
componentIndex, componentIndex,
static_cast<ComponentType>(AdjacentList[componentIndex + index])); static_cast<ComponentType>(Lists::Get().OppositeList[componentIndex + index]));
Traits::SetComponent(hypotenuse, Traits::SetComponent(
componentIndex, adjacent,
static_cast<ComponentType>(HypotenuseList[componentIndex + index])); componentIndex,
static_cast<ComponentType>(Lists::Get().AdjacentList[componentIndex + index]));
Traits::SetComponent(
hypotenuse,
componentIndex,
static_cast<ComponentType>(Lists::Get().HypotenuseList[componentIndex + index]));
} }
VTKM_MATH_ASSERT(test_equal(vtkm::Sin(angle), opposite / hypotenuse), "Sin failed test."); VTKM_MATH_ASSERT(test_equal(vtkm::Sin(angle), opposite / hypotenuse), "Sin failed test.");
@ -364,13 +401,15 @@ struct ScalarVectorFieldTests : public vtkm::exec::FunctorBase
const VectorType zero(0); const VectorType zero(0);
const VectorType half(0.5); const VectorType half(0.5);
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS - NUM_COMPONENTS + 1; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS - NUM_COMPONENTS + 1; index++)
{ {
VectorType x; VectorType x;
for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++) for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++)
{ {
Traits::SetComponent( Traits::SetComponent(
x, componentIndex, static_cast<ComponentType>(AngleList[componentIndex + index])); x,
componentIndex,
static_cast<ComponentType>(Lists::Get().AngleList[componentIndex + index]));
} }
const VectorType minusX = zero - x; const VectorType minusX = zero - x;
@ -391,13 +430,14 @@ struct ScalarVectorFieldTests : public vtkm::exec::FunctorBase
template <typename FunctionType> template <typename FunctionType>
VTKM_EXEC void RaiseToTest(FunctionType function, ComponentType exponent) const VTKM_EXEC void RaiseToTest(FunctionType function, ComponentType exponent) const
{ {
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS - NUM_COMPONENTS + 1; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS - NUM_COMPONENTS + 1; index++)
{ {
VectorType original; VectorType original;
VectorType raiseresult; VectorType raiseresult;
for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++) for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++)
{ {
ComponentType x = static_cast<ComponentType>(NumberList[componentIndex + index]); ComponentType x =
static_cast<ComponentType>(Lists::Get().NumberList[componentIndex + index]);
Traits::SetComponent(original, componentIndex, x); Traits::SetComponent(original, componentIndex, x);
Traits::SetComponent(raiseresult, componentIndex, vtkm::Pow(x, exponent)); Traits::SetComponent(raiseresult, componentIndex, vtkm::Pow(x, exponent));
} }
@ -462,13 +502,14 @@ struct ScalarVectorFieldTests : public vtkm::exec::FunctorBase
ComponentType exponentbias = 0.0, ComponentType exponentbias = 0.0,
ComponentType resultbias = 0.0) const ComponentType resultbias = 0.0) const
{ {
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS - NUM_COMPONENTS + 1; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS - NUM_COMPONENTS + 1; index++)
{ {
VectorType original; VectorType original;
VectorType raiseresult; VectorType raiseresult;
for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++) for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++)
{ {
ComponentType x = static_cast<ComponentType>(NumberList[componentIndex + index]); ComponentType x =
static_cast<ComponentType>(Lists::Get().NumberList[componentIndex + index]);
Traits::SetComponent(original, componentIndex, x); Traits::SetComponent(original, componentIndex, x);
Traits::SetComponent( Traits::SetComponent(
raiseresult, componentIndex, vtkm::Pow(base, x + exponentbias) + resultbias); raiseresult, componentIndex, vtkm::Pow(base, x + exponentbias) + resultbias);
@ -533,14 +574,15 @@ struct ScalarVectorFieldTests : public vtkm::exec::FunctorBase
ComponentType base, ComponentType base,
ComponentType bias = 0.0) const ComponentType bias = 0.0) const
{ {
for (vtkm::IdComponent index = 0; index < NUM_NUMBERS - NUM_COMPONENTS + 1; index++) for (vtkm::IdComponent index = 0; index < Lists::NUM_NUMBERS - NUM_COMPONENTS + 1; index++)
{ {
VectorType basevector(base); VectorType basevector(base);
VectorType original; VectorType original;
VectorType biased; VectorType biased;
for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++) for (vtkm::IdComponent componentIndex = 0; componentIndex < NUM_COMPONENTS; componentIndex++)
{ {
ComponentType x = static_cast<ComponentType>(NumberList[componentIndex + index]); ComponentType x =
static_cast<ComponentType>(Lists::Get().NumberList[componentIndex + index]);
Traits::SetComponent(original, componentIndex, x); Traits::SetComponent(original, componentIndex, x);
Traits::SetComponent(biased, componentIndex, x + bias); Traits::SetComponent(biased, componentIndex, x + bias);
} }

39
vtkm/worklet/internal/ClipTables.h
View File

@ -19456,26 +19456,6 @@ static vtkm::UInt8 CellEdges[CELL_EDGES_SIZE] = {
}; };
#undef X #undef X
// index into ClipTablesIndices for each shape
VTKM_EXEC_CONSTANT
static vtkm::Int16 CellIndexLookup[vtkm::NUMBER_OF_CELL_SHAPES] = {
-1, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
0, // 1 = vtkm::CELL_SHAPE_VERTEX
-1, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
2, // 3 = vtkm::CELL_SHAPE_LINE
-1, // 4 = vtkm::CELL_SHAPE_POLY_LINE
6, // 5 = vtkm::CELL_SHAPE_TRIANGLE
-1, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
-1, // 8 = vtkm::CELL_SHAPE_PIXEL
14, // 9 = vtkm::CELL_SHAPE_QUAD
30, // 10 = vtkm::CELL_SHAPE_TETRA
-1, // 11 = vtkm::CELL_SHAPE_VOXEL
46, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
302, // 13 = vtkm::CELL_SHAPE_WEDGE
366 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
class ClipTables class ClipTables
{ {
public: public:
@ -19488,6 +19468,25 @@ public:
VTKM_EXEC VTKM_EXEC
vtkm::Id GetCaseIndex(vtkm::Id shape, vtkm::Id caseId) const vtkm::Id GetCaseIndex(vtkm::Id shape, vtkm::Id caseId) const
{ {
// index into ClipTablesIndices for each shape
static vtkm::Int16 CellIndexLookup[vtkm::NUMBER_OF_CELL_SHAPES] = {
-1, // 0 = vtkm::CELL_SHAPE_EMPTY_CELL
0, // 1 = vtkm::CELL_SHAPE_VERTEX
-1, // 2 = vtkm::CELL_SHAPE_POLY_VERTEX
2, // 3 = vtkm::CELL_SHAPE_LINE
-1, // 4 = vtkm::CELL_SHAPE_POLY_LINE
6, // 5 = vtkm::CELL_SHAPE_TRIANGLE
-1, // 6 = vtkm::CELL_SHAPE_TRIANGLE_STRIP
-1, // 7 = vtkm::CELL_SHAPE_POLYGON
-1, // 8 = vtkm::CELL_SHAPE_PIXEL
14, // 9 = vtkm::CELL_SHAPE_QUAD
30, // 10 = vtkm::CELL_SHAPE_TETRA
-1, // 11 = vtkm::CELL_SHAPE_VOXEL
46, // 12 = vtkm::CELL_SHAPE_HEXAHEDRON
302, // 13 = vtkm::CELL_SHAPE_WEDGE
366 // 14 = vtkm::CELL_SHAPE_PYRAMID
};
vtkm::Id index = CellIndexLookup[shape]; vtkm::Id index = CellIndexLookup[shape];
return this->ClipTablesIndicesPortal.Get(index + caseId); return this->ClipTablesIndicesPortal.Get(index + caseId);
} }

28
vtkm/worklet/tetrahedralize/TetrahedralizeStructured.h
View File

@ -40,17 +40,6 @@ namespace vtkm
namespace worklet namespace worklet
{ {
namespace detail
{
VTKM_EXEC_CONSTANT
const static vtkm::IdComponent StructuredTetrahedronIndices[2][5][4] = {
{ { 0, 1, 3, 4 }, { 1, 4, 5, 6 }, { 1, 4, 6, 3 }, { 1, 3, 6, 2 }, { 3, 6, 7, 4 } },
{ { 2, 1, 5, 0 }, { 0, 2, 3, 7 }, { 2, 5, 6, 7 }, { 0, 7, 4, 5 }, { 0, 2, 7, 5 } }
};
} // namespace detail
/// \brief Compute the tetrahedralize cells for a uniform grid data set /// \brief Compute the tetrahedralize cells for a uniform grid data set
template <typename DeviceAdapter> template <typename DeviceAdapter>
class TetrahedralizeStructured class TetrahedralizeStructured
@ -82,6 +71,11 @@ public:
ConnectivityOutVec& connectivityOut, ConnectivityOutVec& connectivityOut,
const ThreadIndicesType threadIndices) const const ThreadIndicesType threadIndices) const
{ {
const static vtkm::IdComponent StructuredTetrahedronIndices[2][5][4] = {
{ { 0, 1, 3, 4 }, { 1, 4, 5, 6 }, { 1, 4, 6, 3 }, { 1, 3, 6, 2 }, { 3, 6, 7, 4 } },
{ { 2, 1, 5, 0 }, { 0, 2, 3, 7 }, { 2, 5, 6, 7 }, { 0, 7, 4, 5 }, { 0, 2, 7, 5 } }
};
vtkm::Id3 inputIndex = threadIndices.GetInputIndex3D(); vtkm::Id3 inputIndex = threadIndices.GetInputIndex3D();
// Calculate the type of tetrahedron generated because it alternates // Calculate the type of tetrahedron generated because it alternates
@ -89,14 +83,10 @@ public:
vtkm::IdComponent visitIndex = threadIndices.GetVisitIndex(); vtkm::IdComponent visitIndex = threadIndices.GetVisitIndex();
connectivityOut[0] = connectivityOut[0] = connectivityIn[StructuredTetrahedronIndices[indexType][visitIndex][0]];
connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][0]]; connectivityOut[1] = connectivityIn[StructuredTetrahedronIndices[indexType][visitIndex][1]];
connectivityOut[1] = connectivityOut[2] = connectivityIn[StructuredTetrahedronIndices[indexType][visitIndex][2]];
connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][1]]; connectivityOut[3] = connectivityIn[StructuredTetrahedronIndices[indexType][visitIndex][3]];
connectivityOut[2] =
connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][2]];
connectivityOut[3] =
connectivityIn[detail::StructuredTetrahedronIndices[indexType][visitIndex][3]];
} }
}; };

16
vtkm/worklet/triangulate/TriangulateStructured.h
View File

@ -40,14 +40,6 @@ namespace vtkm
namespace worklet namespace worklet
{ {
namespace detail
{
VTKM_EXEC_CONSTANT
const static vtkm::IdComponent StructuredTriangleIndices[2][3] = { { 0, 1, 2 }, { 0, 2, 3 } };
} // namespace detail
/// \brief Compute the triangulate cells for a uniform grid data set /// \brief Compute the triangulate cells for a uniform grid data set
template <typename DeviceAdapter> template <typename DeviceAdapter>
class TriangulateStructured class TriangulateStructured
@ -79,9 +71,11 @@ public:
ConnectivityOutVec& connectivityOut, ConnectivityOutVec& connectivityOut,
vtkm::IdComponent visitIndex) const vtkm::IdComponent visitIndex) const
{ {
connectivityOut[0] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][0]]; const static vtkm::IdComponent StructuredTriangleIndices[2][3] = { { 0, 1, 2 }, { 0, 2, 3 } };
connectivityOut[1] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][1]];
connectivityOut[2] = connectivityIn[detail::StructuredTriangleIndices[visitIndex][2]]; connectivityOut[0] = connectivityIn[StructuredTriangleIndices[visitIndex][0]];
connectivityOut[1] = connectivityIn[StructuredTriangleIndices[visitIndex][1]];
connectivityOut[2] = connectivityIn[StructuredTriangleIndices[visitIndex][2]];
} }
}; };