Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in
Toggle navigation
Menu
Open sidebar
Sbalzarini Lab
S
Software
P
Parallel Computing
OpenFPM
openfpm_data
Commits
c2cab188
Commit
c2cab188
authored
Dec 19, 2021
by
incardon
Browse files
Fixing to be mutable
parent
6c458d31
Pipeline
#4043
failed with stages
in 9 minutes and 13 seconds
Changes
5
Pipelines
1
Hide whitespace changes
Inline
Side-by-side
src/Grid/Geometry/grid_smb.hpp
View file @
c2cab188
...
...
@@ -150,6 +150,33 @@ public:
return
blockLinId
*
blockSize
+
localLinId
;
}
/*! \brief Linearize the coordinate index
*
* The linearization is given by getting the block indexes and the local coordinate indexes
*
* Linearize the block index (blockLinId), linearize the local index (localLinId) and return
* blockLinId and offset
*
* \param coord coordinates
*
* \return linearized index
*
*/
template
<
typename
indexT_
>
inline
__host__
__device__
void
LinId
(
const
grid_key_dx
<
dim
,
indexT_
>
coord
,
indexT
&
blockLinId
,
int
&
localLinId
)
const
{
//todo: Check (in debug mode only) that the coordinates passed here are valid and not overflowing dimensions (???)
blockLinId
=
coord
.
get
(
dim
-
1
)
/
blockEdgeSize
;
localLinId
=
coord
.
get
(
dim
-
1
)
%
blockEdgeSize
;
for
(
int
d
=
dim
-
2
;
d
>=
0
;
--
d
)
{
blockLinId
*=
blockSz
[
d
];
localLinId
*=
blockEdgeSize
;
blockLinId
+=
coord
.
get
(
d
)
/
blockEdgeSize
;
localLinId
+=
coord
.
get
(
d
)
%
blockEdgeSize
;
}
}
inline
__host__
__device__
grid_key_dx
<
dim
,
int
>
InvLinId
(
const
indexT
linId
)
const
{
indexT
blockLinId
=
linId
/
blockSize
;
...
...
src/SparseGridGpu/BlockMapGpu.hpp
View file @
c2cab188
...
...
@@ -78,6 +78,15 @@ public:
}
}
auto
get
(
unsigned
int
linId
)
const
->
const
decltype
(
blockMap
.
get
(
0
))
&
{
typedef
BlockTypeOf
<
AggregateBlockT
,
0
>
BlockT
;
unsigned
int
blockId
=
linId
/
BlockT
::
size
;
unsigned
int
offset
=
linId
%
BlockT
::
size
;
auto
&
aggregate
=
blockMap
.
get
(
blockId
);
return
aggregate
;
}
/*! \brief insert data, host version
*
* \tparam property id
...
...
@@ -100,6 +109,24 @@ public:
return
block
[
offset
];
}
/*! \brief insert data, host version
*
* \tparam property id
*
* \param linId linearized id block + local linearization
*
* \return a reference to the data
*
*/
auto
insert_o
(
unsigned
int
linId
)
->
decltype
(
blockMap
.
insert
(
0
))
{
typedef
BlockTypeOf
<
AggregateBlockT
,
0
>
BlockT
;
unsigned
int
blockId
=
linId
/
BlockT
::
size
;
unsigned
int
offset
=
linId
%
BlockT
::
size
;
auto
&
aggregate
=
blockMap
.
insert
(
blockId
);
return
aggregate
;
}
/*! \brief insert a block + flush, host version
*
* \tparam property id
...
...
src/SparseGridGpu/SparseGridGpu.hpp
View file @
c2cab188
...
...
@@ -111,6 +111,30 @@ struct aggregate_add<aggregate<types ...>>
/////////////
template
<
typename
enc_type
>
class
encap_data_block
{
int
offset
;
enc_type
enc
;
public:
encap_data_block
(
int
offset
,
const
enc_type
&
enc
)
:
offset
(
offset
),
enc
(
enc
)
{}
encap_data_block
operator
=
(
const
encap_data_block
<
enc_type
>
&
enc
)
{
copy_cpu_encap_single
<
encap_data_block
<
enc_type
>>
cp
(
enc
,
*
this
);
boost
::
mpl
::
for_each_ref
<
boost
::
mpl
::
range_c
<
int
,
0
,
enc_type
::
T_type
::
max_prop
>
>
(
cp
);
return
*
this
;
}
};
/////////////
enum
StencilMode
{
STENCIL_MODE_INPLACE
=
1
,
...
...
@@ -1790,6 +1814,55 @@ public:
return
private_get_data_array
().
template
get
<
p
>(
coord
.
get_cnk_pos_id
())[
coord
.
get_data_id
()];
}
/*! \brief Return the index array of the blocks
*
* \return the index arrays of the blocks
*
*/
auto
private_get_data_array
()
->
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
())
&
{
return
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
();
}
/*! \brief Return the data array of the blocks
*
* \return the index arrays of the blocks
*
*/
auto
private_get_data_array
()
const
->
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
())
{
return
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
();
}
/*! \brief Get an element using the point coordinates
*
* \param coord point coordinates
*
* \return the element
*
*/
template
<
typename
CoordT
>
auto
get_o
(
const
grid_key_dx
<
dim
,
CoordT
>
&
coord
)
const
->
encap_data_block
<
typename
std
::
remove_const
<
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
get
(
0
))
>::
type
>
{
int
offset
;
indexT
lin
;
gridGeometry
.
LinId
(
coord
,
lin
,
offset
);
return
encap_data_block
<
typename
std
::
remove_const
<
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
get
(
0
))
>::
type
>
(
offset
,
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
get
(
lin
));
}
/*! \brief Get an element using sparse_grid_gpu_index (using this index it guarantee that the point exist)
*
* \param element
*
* \return the element
*
*/
auto
get_o
(
const
sparse_grid_gpu_index
<
self
>
&
coord
)
const
->
encap_data_block
<
typename
std
::
remove_const
<
decltype
(
private_get_data_array
().
get
(
0
))
>::
type
>
{
return
encap_data_block
<
typename
std
::
remove_const
<
decltype
(
private_get_data_array
().
get
(
0
))
>::
type
>
(
coord
.
get_data_id
(),
private_get_data_array
().
get
(
coord
.
get_cnk_pos_id
()));
}
/*! \brief This function check if keep geometry is possible for this grid
*
* \return true if skip labelling is possible
...
...
@@ -1833,6 +1906,16 @@ public:
return
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
template
insert
<
p
>(
gridGeometry
.
LinId
(
coord
));
}
template
<
typename
CoordT
>
auto
insert_o
(
const
CoordT
&
coord
)
->
encap_data_block
<
typename
std
::
remove_const
<
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
insert_o
(
0
))
>::
type
>
{
indexT
ind
;
int
offset
;
gridGeometry
.
LinId
(
coord
,
ind
,
offset
);
return
encap_data_block
<
typename
std
::
remove_const
<
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
insert_o
(
0
))
>::
type
>
(
offset
,
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
insert_o
(
ind
));
}
/*! \brief construct link between levels
*
* \praram grid_up grid level up
...
...
@@ -3484,31 +3567,11 @@ public:
* \return the index arrays of the blocks
*
*/
auto
private_get_data_array
()
->
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
())
&
{
return
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
();
}
/*! \brief Return the index array of the blocks
*
* \return the index arrays of the blocks
*
*/
auto
private_get_index_array
()
->
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getIndexBuffer
())
&
auto
private_get_index_array
()
->
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getIndexBuffer
())
{
return
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getIndexBuffer
();
}
/*! \brief Return the index array of the blocks
*
* \return the index arrays of the blocks
*
*/
auto
private_get_data_array
()
const
->
decltype
(
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
())
&
{
return
BlockMapGpu
<
AggregateInternalT
,
threadBlockSize
,
indexT
,
layout_base
>::
blockMap
.
getDataBuffer
();
}
/*! \brief Return the index array of the blocks
*
* \return the index arrays of the blocks
...
...
src/Vector/cuda/map_vector_cuda_ker.cuh
View file @
c2cab188
...
...
@@ -130,7 +130,7 @@ namespace openfpm
unsigned
int
v_size
;
//! 1-D static grid
grid_gpu_ker
<
1
,
T_
,
layout_base
,
grid_sm
<
1
,
void
>>
base
;
mutable
grid_gpu_ker
<
1
,
T_
,
layout_base
,
grid_sm
<
1
,
void
>>
base
;
/*! \brief Check that the key is inside the grid
*
...
...
src/util/multi_array_openfpm/multi_array_ref_openfpm.hpp
View file @
c2cab188
/*
* multi_array_ref_openfpm.hpp
*
* This is an heavily modified boost::multi_array version
*
* Created on: Jun 29, 2018
* Author: i-bird
*/
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment