Skip to content
GitLab
Projects
Groups
Snippets
Help
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
tnl-dev
Project overview
Project overview
Details
Activity
Releases
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Issues
32
Issues
32
List
Boards
Labels
Service Desk
Milestones
Merge Requests
1
Merge Requests
1
Operations
Operations
Incidents
Analytics
Analytics
Repository
Value Stream
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Create a new issue
Commits
Issue Boards
Open sidebar
TNL
tnl-dev
Commits
1187b1be
Commit
1187b1be
authored
Dec 31, 2015
by
Tomáš Oberhuber
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
Adding operator function.
Removing parameter mesh from method getValue in operators.
parent
99f2203b
Changes
26
Hide whitespace changes
Inline
Side-by-side
Showing
26 changed files
with
421 additions
and
339 deletions
+421
-339
ChangeLog
ChangeLog
+5
-0
src/functions/tnlOperatorFunction.h
src/functions/tnlOperatorFunction.h
+67
-0
src/functions/tnlOperatorFunction_impl.h
src/functions/tnlOperatorFunction_impl.h
+52
-0
src/mesh/grids/tnlGridEntity.h
src/mesh/grids/tnlGridEntity.h
+4
-1
src/mesh/grids/tnlTraverser_Grid1D_impl.h
src/mesh/grids/tnlTraverser_Grid1D_impl.h
+18
-18
src/mesh/grids/tnlTraverser_Grid2D_impl.h
src/mesh/grids/tnlTraverser_Grid2D_impl.h
+27
-33
src/mesh/grids/tnlTraverser_Grid3D_impl.h
src/mesh/grids/tnlTraverser_Grid3D_impl.h
+44
-45
src/operators/diffusion/nonlinear-diffusion-operators/tnlFiniteVolumeNonlinearOperator.h
...ar-diffusion-operators/tnlFiniteVolumeNonlinearOperator.h
+3
-6
src/operators/diffusion/nonlinear-diffusion-operators/tnlFiniteVolumeNonlinearOperator_impl.h
...ffusion-operators/tnlFiniteVolumeNonlinearOperator_impl.h
+49
-50
src/operators/diffusion/nonlinear-diffusion-operators/tnlOneSideDiffNonlinearOperator.h
...ear-diffusion-operators/tnlOneSideDiffNonlinearOperator.h
+3
-6
src/operators/diffusion/nonlinear-diffusion-operators/tnlOneSideDiffNonlinearOperator_impl.h
...iffusion-operators/tnlOneSideDiffNonlinearOperator_impl.h
+60
-69
src/operators/diffusion/tnlLinearDiffusion.h
src/operators/diffusion/tnlLinearDiffusion.h
+3
-6
src/operators/diffusion/tnlLinearDiffusion_impl.h
src/operators/diffusion/tnlLinearDiffusion_impl.h
+9
-12
src/operators/diffusion/tnlNonlinearDiffusion.h
src/operators/diffusion/tnlNonlinearDiffusion.h
+5
-8
src/operators/diffusion/tnlNonlinearDiffusion_impl.h
src/operators/diffusion/tnlNonlinearDiffusion_impl.h
+6
-9
src/operators/operator-Q/tnlFiniteVolumeOperatorQ.h
src/operators/operator-Q/tnlFiniteVolumeOperatorQ.h
+1
-3
src/operators/operator-Q/tnlFiniteVolumeOperatorQ_impl.h
src/operators/operator-Q/tnlFiniteVolumeOperatorQ_impl.h
+7
-8
src/operators/operator-Q/tnlOneSideDiffOperatorQ.h
src/operators/operator-Q/tnlOneSideDiffOperatorQ.h
+11
-13
src/operators/operator-Q/tnlOneSideDiffOperatorQ_impl.h
src/operators/operator-Q/tnlOneSideDiffOperatorQ_impl.h
+19
-23
src/operators/tnlDirichletBoundaryConditions.h
src/operators/tnlDirichletBoundaryConditions.h
+3
-3
src/operators/tnlDirichletBoundaryConditions_impl.h
src/operators/tnlDirichletBoundaryConditions_impl.h
+2
-2
src/operators/tnlExactOperatorEvaluator.h
src/operators/tnlExactOperatorEvaluator.h
+2
-2
src/operators/tnlNeumannBoundaryConditions.h
src/operators/tnlNeumannBoundaryConditions.h
+8
-8
src/operators/tnlNeumannBoundaryConditions_impl.h
src/operators/tnlNeumannBoundaryConditions_impl.h
+9
-9
src/solvers/pde/tnlBoundaryConditionsSetter.h
src/solvers/pde/tnlBoundaryConditionsSetter.h
+2
-2
src/solvers/pde/tnlExplicitUpdater.h
src/solvers/pde/tnlExplicitUpdater.h
+2
-3
No files found.
ChangeLog
View file @
1187b1be
* tnlBoundaryConditionsSetter was implemented
- it allows to set boundary conditions to given function
- in explicit solvers, it must be called explicitly now after calling tnlExplicitUpdater
2015-12-25 Tomas Oberhuber <tomas.oberhuber@fjfi.cvut.cz>
* significant changes in grids
- the central structure is now grid entity rather than grid itself
...
...
src/functions/tnlOperatorFunction.h
0 → 100644
View file @
1187b1be
/***************************************************************************
tnlOperatorFunction.h - description
-------------------
begin : Dec 31, 2015
copyright : (C) 2015 by oberhuber
email : tomas.oberhuber@fjfi.cvut.cz
***************************************************************************/
/***************************************************************************
* *
* This program is free software; you can redistribute it and/or modify *
* it under the terms of the GNU General Public License as published by *
* the Free Software Foundation; either version 2 of the License, or *
* (at your option) any later version. *
* *
***************************************************************************/
#ifndef TNLOPERATORFUNCTION_H
#define TNLOPERATORFUNCTION_H
#include <type_traits>
#include <core/tnlCuda.h>
template
<
typename
Operator
,
typename
MeshFunction
,
typename
BoundaryConditions
>
class
tnlOperatorFunction
{
static_assert
(
std
::
is_same
<
typename
Operator
::
MeshType
,
typename
MeshFunction
::
MeshType
>::
value
,
"Operator and MeshFunction have different mesh type."
);
static_assert
(
std
::
is_same
<
typename
Operator
::
MeshType
,
typename
BoundaryConditions
::
MeshType
>::
value
,
"Operator and BoundaryConditions have different mesh type."
);
public:
typedef
Operator
OperatorType
;
typedef
MeshFunction
MeshFunctionType
;
typedef
BoundaryConditions
BoundaryConditionsType
;
typedef
typename
OperatorType
::
MeshType
MeshType
;
typedef
typename
OperatorType
::
RealType
RealType
;
typedef
typename
OperatorType
::
DeviceType
DeviceType
;
typedef
typename
OperatorType
::
IndexType
IndexType
;
tnlOperatorFunction
(
const
OperatorType
&
operator_
,
const
MeshFunctionType
&
function
,
const
BoundaryConditionsType
&
boundaryConditions
);
template
<
typename
MeshEntity
>
__cuda_callable__
RealType
operator
()(
const
MeshEntity
&
meshEntity
,
const
RealType
&
time
);
protected:
const
Operator
&
operator_
;
const
MeshFunction
&
function
;
const
BoundaryConditions
&
boundaryConditions
;
};
#endif
/* TNLOPERATORFUNCTION_H */
src/functions/tnlOperatorFunction_impl.h
0 → 100644
View file @
1187b1be
/***************************************************************************
tnlOperatorFunction_impl.h - description
-------------------
begin : Dec 31, 2015
copyright : (C) 2015 by oberhuber
email : tomas.oberhuber@fjfi.cvut.cz
***************************************************************************/
/***************************************************************************
* *
* This program is free software; you can redistribute it and/or modify *
* it under the terms of the GNU General Public License as published by *
* the Free Software Foundation; either version 2 of the License, or *
* (at your option) any later version. *
* *
***************************************************************************/
#ifndef TNLOPERATORFUNCTION_IMPL_H
#define TNLOPERATORFUNCTION_IMPL_H
template
<
typename
Operator
,
typename
Function
>
tnlOperatorFunction
<
Operator
,
Function
>::
tnlOperatorFunction
(
const
OperatorType
&
operator_
,
const
MeshFunctionType
&
function
,
const
BoundaryConditionsType
&
boundaryConditions
)
:
operator_
(
operator_
),
function
(
function
),
boundaryConditions
(
boundaryConditions
)
{
}
template
<
typename
Operator
,
typename
Function
>
template
<
typename
MeshEntity
>
__cuda_callable__
typename
tnlOperatorFunction
<
Operator
,
Function
>::
RealType
tnlOperatorFunction
<
Operator
,
Function
>::
operator
()(
const
MeshEntity
&
meshEntity
,
const
RealType
&
time
)
{
static_assert
(
MeshEntity
::
entityDimensions
==
Function
::
getMeshEntityDimensions
(),
"Wrong mesh entity dimensions."
);
if
(
meshEntity
.
isBoundaryEntity
()
)
return
boundaryConditions
.
getValue
(
meshEntity
,
function
.
getData
(),
time
);
return
operator_
.
getValue
(
meshEntity
,
function
.
getData
(),
time
);
}
#endif
/* TNLOPERATORFUNCTION_IMPL_H */
src/mesh/grids/tnlGridEntity.h
View file @
1187b1be
...
...
@@ -51,6 +51,7 @@ class tnlGridEntity< tnlGrid< Dimensions, Real, Device, Index >, EntityDimension
public:
typedef
tnlGrid
<
Dimensions
,
Real
,
Device
,
Index
>
GridType
;
typedef
GridType
MeshType
;
typedef
typename
GridType
::
IndexType
IndexType
;
typedef
typename
GridType
::
CoordinatesType
CoordinatesType
;
typedef
Config
ConfigType
;
...
...
@@ -168,6 +169,7 @@ class tnlGridEntity< tnlGrid< Dimensions, Real, Device, Index >, Dimensions, Con
public:
typedef
tnlGrid
<
Dimensions
,
Real
,
Device
,
Index
>
GridType
;
typedef
GridType
MeshType
;
typedef
typename
GridType
::
IndexType
IndexType
;
typedef
typename
GridType
::
CoordinatesType
CoordinatesType
;
typedef
typename
GridType
::
VertexType
VertexType
;
...
...
@@ -278,7 +280,8 @@ class tnlGridEntity< tnlGrid< Dimensions, Real, Device, Index >, 0, Config >
{
public:
typedef
tnlGrid
<
Dimensions
,
Real
,
Device
,
Index
>
GridType
;
typedef
tnlGrid
<
Dimensions
,
Real
,
Device
,
Index
>
GridType
;
typedef
GridType
MeshType
;
typedef
typename
GridType
::
IndexType
IndexType
;
typedef
typename
GridType
::
CoordinatesType
CoordinatesType
;
typedef
typename
GridType
::
VertexType
VertexType
;
...
...
src/mesh/grids/tnlTraverser_Grid1D_impl.h
View file @
1187b1be
...
...
@@ -31,8 +31,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Boundary cells
*/
const
int
CellDimensions
=
GridType
::
meshDimensions
;
//typename GridType::template GridEntity< CellDimensions > entity( grid );
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -58,8 +57,7 @@ processInteriorEntities( const GridType& grid,
/****
* Interior cells
*/
const
int
CellDimensions
=
GridType
::
meshDimensions
;
//typename GridType::template GridEntity< CellDimensions > cell( grid );
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
GridEntity
cell
(
grid
);
const
IndexType
&
xSize
=
grid
.
getDimensions
().
x
();
...
...
@@ -85,8 +83,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Boundary vertices
*/
const
int
VerticesDimensions
=
0
;
//typename GridType::template GridEntity< VerticesDimensions > entity( grid );
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -112,10 +109,9 @@ processInteriorEntities( const GridType& grid,
/****
* Interior vertices
*/
const
int
VerticesDimensions
=
0
;
//typename GridType::template GridEntity< VerticesDimensions > entity( grid );
GridEntity
entity
(
grid
);
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
const
IndexType
&
xSize
=
grid
.
getDimensions
().
x
();
for
(
coordinates
.
x
()
=
1
;
coordinates
.
x
()
<
xSize
;
coordinates
.
x
()
++
)
...
...
@@ -134,6 +130,7 @@ processInteriorEntities( const GridType& grid,
#ifdef HAVE_CUDA
template
<
typename
Real
,
typename
Index
,
typename
GridEntity
,
typename
UserData
,
typename
EntitiesProcessor
,
bool
processAllEntities
,
...
...
@@ -142,11 +139,10 @@ __global__ void tnlTraverserGrid1DCells( const tnlGrid< 1, Real, tnlCuda, Index
UserData
*
userData
,
Index
gridXIdx
)
{
typedef
tnlGrid
<
1
,
Real
,
tnlCuda
,
Index
>
GridType
;
const
int
CellDimensions
=
GridType
::
meshDimensions
;
//typename GridType::template GridEntity< CellDimensions > entity( *grid );
GridEntity
entity
(
*
grid
);
typedef
tnlGrid
<
1
,
Real
,
tnlCuda
,
Index
>
GridType
;
typedef
typename
GridType
::
CoordinatesType
CoordinatesType
;
GridEntity
entity
(
*
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
const
Index
index
=
(
gridXIdx
*
tnlCuda
::
getMaxGridSize
()
+
blockIdx
.
x
)
*
blockDim
.
x
+
threadIdx
.
x
;
...
...
@@ -178,6 +174,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Boundary conditions
*/
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -187,7 +184,7 @@ processBoundaryEntities( const GridType& grid,
const
IndexType
cudaXGrids
=
tnlCuda
::
getNumberOfGrids
(
cudaBlocks
.
x
);
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
tnlTraverserGrid1DCells
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
true
>
tnlTraverserGrid1DCells
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
true
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -213,6 +210,7 @@ processInteriorEntities( const GridType& grid,
/****
* Interior cells
*/
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
checkCudaDevice
;
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -227,7 +225,7 @@ processInteriorEntities( const GridType& grid,
{
if
(
gridXIdx
==
cudaXGrids
-
1
)
cudaGridSize
.
x
=
cudaBlocks
.
x
%
tnlCuda
::
getMaxGridSize
();
tnlTraverserGrid1DCells
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
false
>
tnlTraverserGrid1DCells
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
false
>
<<<
cudaGridSize
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -243,6 +241,7 @@ processInteriorEntities( const GridType& grid,
#ifdef HAVE_CUDA
template
<
typename
Real
,
typename
Index
,
typename
GridEntity
,
typename
UserData
,
typename
EntitiesProcessor
,
bool
processAllEntities
,
...
...
@@ -252,7 +251,6 @@ __global__ void tnlTraverserGrid1DVertices( const tnlGrid< 1, Real, tnlCuda, Ind
Index
gridXIdx
)
{
typedef
tnlGrid
<
1
,
Real
,
tnlCuda
,
Index
>
GridType
;
//typename GridType::template GridEntity< GridType::Vertices > vertex;
GridEntity
vertex
(
*
grid
);
const
Index
&
xSize
=
grid
->
getDimensions
().
x
();
...
...
@@ -290,6 +288,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Boundary vertices
*/
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -299,7 +298,7 @@ processBoundaryEntities( const GridType& grid,
const
IndexType
cudaXGrids
=
tnlCuda
::
getNumberOfGrids
(
cudaBlocks
.
x
);
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
tnlTraverserGrid1DVertices
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
true
>
tnlTraverserGrid1DVertices
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
true
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -327,6 +326,7 @@ processInteriorEntities( const GridType& grid,
/****
* Interior vertices
*/
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
checkCudaDevice
;
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -341,7 +341,7 @@ processInteriorEntities( const GridType& grid,
{
if
(
gridXIdx
==
cudaXGrids
-
1
)
cudaGridSize
.
x
=
cudaBlocks
.
x
%
tnlCuda
::
getMaxGridSize
();
tnlTraverserGrid1DVertices
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
false
>
tnlTraverserGrid1DVertices
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
false
>
<<<
cudaGridSize
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
src/mesh/grids/tnlTraverser_Grid2D_impl.h
View file @
1187b1be
...
...
@@ -32,8 +32,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Traversing boundary cells
*/
const
int
CellDimensions
=
GridType
::
meshDimensions
;
//typename GridType::template GridEntity< CellDimensions > entity( grid );
static_assert
(
GridEntity
::
entityDimensions
==
2
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -73,9 +72,8 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior cells
*/
const
int
CellDimensions
=
GridType
::
meshDimensions
;
//typename GridType::template GridEntity< CellDimensions > entity( grid );
static_assert
(
GridEntity
::
entityDimensions
==
2
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -105,9 +103,8 @@ processBoundaryEntities( const GridType& grid,
{
/****
* Traversing boundary faces
*/
const
int
FaceDimensions
=
GridType
::
meshDimensions
-
1
;
//typedef typename GridType::template GridEntity< FaceDimensions > EntityType;
*/
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
typedef
typename
GridEntity
::
EntityOrientationType
EntityOrientationType
;
GridEntity
entity
(
grid
);
...
...
@@ -151,8 +148,7 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior faces
*/
const
int
FaceDimensions
=
GridType
::
meshDimensions
-
1
;
//typedef typename GridType::template GridEntity< FaceDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
typedef
typename
GridEntity
::
EntityOrientationType
EntityOrientationType
;
GridEntity
entity
(
grid
);
...
...
@@ -195,8 +191,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Traversing boundary vertices
*/
const
int
VertexDimensions
=
0
;
//typedef typename GridType::template GridEntity< VertexDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
typedef
typename
GridEntity
::
EntityOrientationType
EntityOrientation
;
GridEntity
entity
(
grid
);
...
...
@@ -237,8 +232,7 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior vertices
*/
const
int
VertexDimensions
=
0
;
//typedef typename GridType::template GridEntity< VertexDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
typedef
typename
GridEntity
::
EntityOrientationType
EntityOrientation
;
GridEntity
entity
(
grid
);
...
...
@@ -262,10 +256,10 @@ processInteriorEntities( const GridType& grid,
* CUDA Specializations
*
*/
#ifdef HAVE_CUDA
template
<
typename
Real
,
typename
Index
,
typename
GridEntity
,
typename
UserData
,
typename
EntitiesProcessor
,
bool
processAllEntities
,
...
...
@@ -276,13 +270,8 @@ __global__ void tnlTraverserGrid2DCells( const tnlGrid< 2, Real, tnlCuda, Index
const
Index
gridYIdx
)
{
typedef
tnlGrid
<
2
,
Real
,
tnlCuda
,
Index
>
GridType
;
const
int
CellDimensions
=
GridType
::
meshDimensions
;
GridEntity
entity
(
*
grid
);
typedef
typename
GridType
::
CoordinatesType
CoordinatesType
;
//CoordinatesType& coordinates = entity.getCoordinates();
/*const Index& xSize = grid->getDimensions().x();
const Index& ySize = grid->getDimensions().y();*/
entity
.
getCoordinates
().
x
()
=
(
gridXIdx
*
tnlCuda
::
getMaxGridSize
()
+
blockIdx
.
x
)
*
blockDim
.
x
+
threadIdx
.
x
;
entity
.
getCoordinates
().
y
()
=
(
gridYIdx
*
tnlCuda
::
getMaxGridSize
()
+
blockIdx
.
y
)
*
blockDim
.
y
+
threadIdx
.
y
;
...
...
@@ -303,6 +292,7 @@ __global__ void tnlTraverserGrid2DCells( const tnlGrid< 2, Real, tnlCuda, Index
template
<
typename
Real
,
typename
Index
,
typename
GridEntity
,
typename
UserData
,
typename
EntitiesProcessor
,
bool
processAllEntities
,
...
...
@@ -316,11 +306,11 @@ __global__ void tnlTraverserGrid2DFaces( const tnlGrid< 2, Real, tnlCuda, Index
{
typedef
tnlGrid
<
2
,
Real
,
tnlCuda
,
Index
>
GridType
;
const
int
FaceDimensions
=
GridType
::
meshDimensions
-
1
;
//typedef typename GridType::template GridEntity< GridType::Cells > EntityType;
GridEntity
entity
(
*
grid
);
typedef
typename
GridType
::
CoordinatesType
CoordinatesType
;
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
entity
.
setOrientation
(
typename
EntityType
::
EntityOrientationType
(
nx
,
ny
)
);
entity
.
setOrientation
(
typename
GridEntity
::
EntityOrientationType
(
nx
,
ny
)
);
const
Index
&
xSize
=
grid
->
getDimensions
().
x
();
const
Index
&
ySize
=
grid
->
getDimensions
().
y
();
...
...
@@ -344,6 +334,7 @@ __global__ void tnlTraverserGrid2DFaces( const tnlGrid< 2, Real, tnlCuda, Index
template
<
typename
Real
,
typename
Index
,
typename
GridEntity
,
typename
UserData
,
typename
EntitiesProcessor
,
bool
processAllEntities
,
...
...
@@ -354,8 +345,6 @@ __global__ void tnlTraverserGrid2DVertices( const tnlGrid< 2, Real, tnlCuda, Ind
const
Index
gridYIdx
)
{
typedef
tnlGrid
<
2
,
Real
,
tnlCuda
,
Index
>
GridType
;
const
int
VertexDimensions
=
0
;
//typedef typename GridType::template GridEntity< VertexDimensions > EntityType;
typedef
typename
GridType
::
CoordinatesType
CoordinatesType
;
GridEntity
entity
(
*
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -396,6 +385,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Boundary conditions
*/
static_assert
(
GridEntity
::
entityDimensions
==
2
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -409,7 +399,7 @@ processBoundaryEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DCells
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
true
>
tnlTraverserGrid2DCells
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
true
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -435,6 +425,7 @@ processInteriorEntities( const GridType& grid,
/****
* Interior cells
*/
static_assert
(
GridEntity
::
entityDimensions
==
2
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -448,7 +439,7 @@ processInteriorEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DCells
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
false
>
tnlTraverserGrid2DCells
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
false
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -475,6 +466,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Boundary faces
*/
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -492,7 +484,7 @@ processBoundaryEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DFaces
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
true
>
tnlTraverserGrid2DFaces
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
true
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -514,7 +506,7 @@ processBoundaryEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DFaces
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
true
>
tnlTraverserGrid2DFaces
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
true
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -542,7 +534,7 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior faces
*/
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -560,7 +552,7 @@ processInteriorEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DFaces
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
false
>
tnlTraverserGrid2DFaces
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
false
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -582,7 +574,7 @@ processInteriorEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DFaces
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
false
>
tnlTraverserGrid2DFaces
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
false
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -609,6 +601,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Traversing boundary vertices
*/
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -622,7 +615,7 @@ processBoundaryEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DVertices
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
true
>
tnlTraverserGrid2DVertices
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
true
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
@@ -649,6 +642,7 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior vertices
*/
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
GridType
*
kernelGrid
=
tnlCuda
::
passToDevice
(
grid
);
UserData
*
kernelUserData
=
tnlCuda
::
passToDevice
(
userData
);
...
...
@@ -662,7 +656,7 @@ processInteriorEntities( const GridType& grid,
for
(
IndexType
gridXIdx
=
0
;
gridXIdx
<
cudaXGrids
;
gridXIdx
++
)
for
(
IndexType
gridYIdx
=
0
;
gridYIdx
<
cudaYGrids
;
gridYIdx
++
)
{
tnlTraverserGrid2DVertices
<
Real
,
Index
,
UserData
,
EntitiesProcessor
,
false
,
false
>
tnlTraverserGrid2DVertices
<
Real
,
Index
,
GridEntity
,
UserData
,
EntitiesProcessor
,
false
,
false
>
<<<
cudaBlocks
,
cudaBlockSize
>>>
(
kernelGrid
,
kernelUserData
,
...
...
src/mesh/grids/tnlTraverser_Grid3D_impl.h
View file @
1187b1be
...
...
@@ -31,8 +31,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Boundary conditions
*/
const
int
CellDimensions
=
GridType
::
meshDimensions
;
//typename GridType::template GridEntity< CellDimensions > entity( grid );
static_assert
(
GridEntity
::
entityDimensions
==
3
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -86,8 +85,7 @@ processInteriorEntities( const GridType& grid,
/****
* Interior cells
*/
const
int
CellDimensions
=
GridType
::
meshDimensions
;
//typename GridType::template GridEntity< CellDimensions > entity( grid );
static_assert
(
GridEntity
::
entityDimensions
==
3
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -121,8 +119,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Traversing boundary faces
*/
const
int
FaceDimensions
=
GridType
::
meshDimensions
-
1
;
//typedef typename GridType::template GridEntity< FaceDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
2
,
"The entity has wrong dimensions."
);
GridEntity
entity
;
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -181,8 +178,7 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior faces
*/
const
int
FaceDimensions
=
GridType
::
meshDimensions
-
1
;
//typedef typename GridType::template GridEntity< FaceDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
2
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -238,8 +234,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Traversing boundary edges
*/
const
int
EdgeDimensions
=
1
;
//typedef typename GridType::template GridEntity< EdgeDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -334,8 +329,7 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior edges
*/
const
int
EdgeDimensions
=
1
;
//typedef typename GridType::template GridEntity< EdgeDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
1
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -390,8 +384,7 @@ processBoundaryEntities( const GridType& grid,
/****
* Traversing boundary vertices
*/
const
int
VertexDimensions
=
0
;
//typedef typename GridType::template GridEntity< VertexDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
GridEntity
entity
;
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -448,8 +441,7 @@ processInteriorEntities( const GridType& grid,
/****
* Traversing interior vertices
*/
const
int
VertexDimensions
=
0
;
//typedef typename GridType::template GridEntity< VertexDimensions > EntityType;
static_assert
(
GridEntity
::
entityDimensions
==
0
,
"The entity has wrong dimensions."
);
GridEntity
entity
(
grid
);
CoordinatesType
&
coordinates
=
entity
.
getCoordinates
();
...
...
@@ -470,7 +462,6 @@ processInteriorEntities( const GridType& grid,
}
}
/***
*
* CUDA Specializations
...
...
@@ -479,6 +470,7 @@ processInteriorEntities( const GridType& grid,
#ifdef HAVE_CUDA
template
<
typename
Real
,
typename
Index
,
typename
GridEntity
,