6 #include "util/cuda_launch.hpp" 7 #include <boost/shared_array.hpp> 9 #include <initializer_list> 11 #include "memory/memory.hpp" 12 #include "Space/Shape/Box.hpp" 13 #include "grid_key.hpp" 15 #include "util/mathutil.hpp" 16 #include "iterators/stencil_type.hpp" 25 template<
typename SparseGr
id_type>
26 __device__ __host__
bool check(SparseGrid_type & sggt,
unsigned int dataBlockPos,
unsigned int offset)
32 template<
unsigned int dim,
typename T>
42 template<
typename SparseGr
idGpu_type>
43 __device__ __host__
bool check(SparseGridGpu_type & sggt,
unsigned int dataBlockId,
unsigned int offset)
45 auto key = sggt.getCoord(dataBlockId,offset);
69 static bool valid(
size_t v_id,
size_t sz)
94 static bool valid(
size_t v_id,
size_t sz)
100 template<
unsigned int dim>
103 #if defined(CUDA_GPU) 113 return wthr.x * wthr.y * wthr.z;
118 return thr.x * thr.y * thr.z;
124 template<
unsigned int dim>
130 {tot_work *= ite.wthr.x * ite.thr.x;}
133 tot_work *= ite.wthr.x * ite.thr.x;
134 tot_work *= ite.wthr.y * ite.thr.y;
138 tot_work *= ite.wthr.x * ite.thr.x;
139 tot_work *= ite.wthr.y * ite.thr.y;
140 tot_work *= ite.wthr.z * ite.thr.z;
143 return tot_work != 0;
147 template<
unsigned int N,
typename T>
class grid_sm;
149 template<
unsigned int dim,
typename gr
id_sm_type,
typename T>
156 template<
unsigned int dim,
typename stencil=no_stencil,
typename linearizer = gr
id_sm<dim,
void>,
typename warn=pr
int_warning_on_adjustment<dim,linearizer>>
class grid_key_dx_iterator_sub;
165 template<
unsigned int N,
typename T>
199 for (
size_t i = 1 ; i < N ; i++)
230 for (
size_t i = 1 ; i < N ; i++)
258 for (
size_t i = 1 ; i < N ; i++)
274 template<
typename a,
typename ...lT>
275 __device__ __host__
inline mem_id
Lin_impl(a v,lT...t)
const 281 template<
typename a> __device__ __host__
inline mem_id
Lin_impl(a v)
const 312 for (
size_t i = 0 ; i < N ; i++)
342 for (
size_t i = 0 ; i < N ; i++)
360 for (
size_t i = 0 ; i < N ; i++)
369 inline size_t totalSize(
const size_t sz)
373 for (
size_t i = 0 ; i < N ; i++)
383 inline size_t totalSize(
const size_t (&
sz)[N])
387 for (
size_t i = 0 ; i < N ; i++)
433 template<
typename check=NoCheck,
typename ids_type>
440 if (check::valid(gk.
get(0) + sum_id[0],
sz[0]) ==
false)
443 lid = gk.
get(0) + sum_id[0];
446 for (mem_id i = 1 ; i < N ; i++)
450 if (check::valid(gk.
get(i) + sum_id[i],
sz[i]) ==
false)
469 template<
typename check=NoCheck,
typename ids_type>
476 if (bc[0] == NON_PERIODIC)
478 if (check::valid(gk.
get(0) + sum_id[0],
sz[0]) ==
false)
481 lid = gk.
get(0) + sum_id[0];
485 lid = openfpm::math::positive_modulo(gk.
get(0) + sum_id[0],
sz[0]);
488 for (mem_id i = 1 ; i < N ; i++)
493 if (bc[i] == NON_PERIODIC)
495 if (check::valid(gk.
get(i) + sum_id[i],
sz[i]) ==
false)
502 lid += (openfpm::math::positive_modulo(gk.
get(i) + sum_id[i],
sz[i])) *
sz_s[i-1];
520 for (mem_id i = 1 ; i < N ; i++)
537 __device__ __host__
inline mem_id
LinId(
const size_t (& k)[N])
const 540 for (mem_id i = 1 ; i < N ; i++)
561 for (mem_id i = 1 ; i < N ; i++)
575 template<
typename a,
typename ...lT,
typename enabler =
typename std::enable_if<
sizeof...(lT) == N-1>::type >
576 __device__ __host__
inline mem_id
Lin(a v,lT...t)
const 595 for (mem_id i = 0 ; i < N ; i++)
615 inline mem_id
LinId(mem_id *
id)
const 619 for (mem_id i = 1 ; i < N ; i++)
637 __device__ __host__
inline size_t size()
const 652 for (
size_t i = 0 ; i < N ; i++)
673 for (
size_t i = 0 ; i < N ; i++)
675 if (
sz[i] != g.
sz[i])
684 for (
size_t i = 0 ; i < N ; i++)
716 inline size_t size_s(
unsigned int i)
const 730 __device__ __host__
inline size_t size(
unsigned int i)
const 758 #if defined(CUDA_GPU) 766 template<
typename T2>
769 return getGPUIterator_impl<N>(*
this,key1,key2,n_thr);
778 struct ite_gpu<N> getGPUIterator(
size_t n_thr = default_kernel_wg_threads_)
const 783 for (
size_t i = 0 ; i < N ; i++)
789 return getGPUIterator_impl<N>(*
this,k1,k2,n_thr);
805 for (
size_t i = 0 ; i < N ; i++)
824 std::stringstream str;
826 for (
size_t i = 0 ; i < N ; i++)
827 str <<
"sz[" << i <<
"]=" <<
size(i) <<
" ";
833 template <
unsigned int,
typename>
friend class grid_sm;
837 template<
unsigned int dim,
typename gr
id_sm_type,
typename T>
841 for (
size_t i = 0 ; i < dim ; i++)
842 {tot_work *= key2.
get(i) - key1.
get(i) + 1;}
844 size_t n = (tot_work <= n_thr)?openfpm::math::round_big_2(tot_work):n_thr;
871 {ig.thr.x = ig.thr.x << 1;}
872 else if (dir % 3 == 1)
873 {ig.thr.y = ig.thr.y << 1;}
874 else if (dir % 3 == 2)
875 {ig.thr.z = ig.thr.z << 1;}
884 {ig.wthr.x = (key2.
get(0) - key1.
get(0) + 1) / ig.thr.x + (((key2.
get(0) - key1.
get(0) + 1)%ig.thr.x != 0)?1:0);}
887 {ig.wthr.y = (key2.
get(1) - key1.
get(1) + 1) / ig.thr.y + (((key2.
get(1) - key1.
get(1) + 1)%ig.thr.y != 0)?1:0);}
895 for (
size_t i = 2 ; i < dim ; i++)
896 {ig.wthr.z *= (key2.
get(i) - key1.
get(i) + 1) / ig.thr.z + (((key2.
get(i) - key1.
get(i) + 1)%ig.thr.z != 0)?1:0);}
903 if (dim >= 1 && ig.wthr.x == 1)
904 {ig.thr.x = (key2.
get(0) - key1.
get(0) + 1);}
906 if (dim >= 2 && ig.wthr.y == 1)
907 {ig.thr.y = key2.
get(1) - key1.
get(1) + 1;}
909 if (dim == 3 && ig.wthr.z == 1)
910 {ig.thr.z = key2.
get(2) - key1.
get(2) + 1;}
912 for (
size_t i = 0 ; i < dim ; i++)
914 ig.start.set_d(i,key1.
get(i));
915 ig.stop.set_d(i,key2.
get(i));
957 for(
unsigned int i = 0 ; i < dim ; i++)
983 template<
typename a,
typename ...T>
void set(a v, T...t)
1028 k[
sizeof...(T)] = v;
1029 invert_assign(t...);
1032 template<
typename a,
typename ...T>
void invert_assign(a v)
1037 void invert_assign()
1082 :dim(n),
sz(
sz),gk(n)
1086 for (
size_t i = 0 ; i < dim ; i++)
1088 gk.
set_d(i,dim-i-1);
1109 for ( ; i < dim-1 ; i++)
1111 size_t id = gk.
get(i);
1130 for (
unsigned int s = 0 ; s <= i+1 ; s++)
1132 gk.
set_d(i+1-s,
id+1+s);
1159 if (gk.
get(dim-1) < static_cast<mem_id>(
sz-dim+1))
Declaration print_warning_on_adjustment.
grid_sm()
Default constructor.
size_t sz_s[N]
size of the grid on each stride (used for linearization)
__device__ __host__ mem_id Lin_impl(a v) const
Linearize a set of index.
__device__ __host__ grid_key_dx< N > InvLinId(mem_id id) const
Construct.
Iterator_g_const(size_t n, size_t sz)
Constructor.
std::string toString() const
Produce a string from the object.
__device__ __host__ mem_id Lin_impl(a v, lT...t) const
linearize an arbitrary set of index
__device__ __host__ size_t size() const
Return the size of the grid.
void set_d(size_t i, mem_id id)
Set the i index.
grid_key_dx_r(size_t dim)
constructor
void set(a v, T...t)
set the grid key from a list of numbers
__device__ __host__ T getLow(int i) const
get the i-coordinate of the low bound interval of the box
__device__ __host__ mem_id LinId(const grid_key_dx< N, ids_type > &gk) const
Linearization of the grid_key_dx.
grid_key_dx_iterator_sub< N > getSubIterator(const grid_key_dx< N > &start, const grid_key_dx< N > &stop) const
Return a sub-grid iterator.
Emulate grid_key_dx with runtime dimensionality.
size_t size_tot
total number of the elements in the grid
Box< N, size_t > box
Box enclosing the grid.
static bool valid(size_t v_id, size_t sz)
Check if vertex exist.
__device__ __host__ index_type get(index_type i) const
Get the i index.
__device__ __host__ mem_id LinId(const size_t(&k)[N]) const
Linearization of the grid_key_dx.
mem_id get(size_t i)
get the i index
grid_key_dx_r(grid_key_dx_r &key)
constructor from another key
size_t sz[N]
size of the grid
bool operator==(const grid_sm< N, T > &g)
Check if the two grid_sm are the same.
bool operator!=(const grid_sm< N, T > &g)
Check if the two grid_sm are the same.
void invert_assign(a v, T...t)
Recursively invert the assignment.
__device__ __host__ void setHigh(int i, T val)
set the high interval of the box
mem_id LinId(const grid_key_dx< N, ids_type > &gk, const char sum_id[N]) const
Linearization of the grid_key_dx with a specified shift.
void Initialize()
Initialize the basic structure.
mem_id LinIdPtr(size_t *k) const
Linearization of the set of indexes.
const Box< N, size_t > getBoxKey() const
Return the box enclosing the grid.
grid_sm(const size_t &sz)
Construct a grid of a specified size.
__device__ __host__ mem_id Lin(a v, lT...t) const
linearize an arbitrary set of index
static bool valid(size_t v_id, size_t sz)
No check is performed.
void swap(grid_sm< N, T > &g)
swap the grid_sm informations
grid_sm(const grid_sm< N, S > &g)
construct a grid from another grid
__device__ __host__ void setLow(int i, T val)
set the low interval of the box
size_t sz
size of the grid (the grid is assumed a square so equal on each dimension)
size_t getDim()
Get the dimensionality of the iterator.
const size_t(& getSize() const)[N]
Return the size of the grid as an array.
__device__ __host__ size_t size(unsigned int i) const
void Initialize(const size_t(&sz)[N])
Initialize the basic structure.
size_t size_s(unsigned int i) const
This class represent an N-dimensional box.
friend class grid_sm
It simply mean that all the classes grid are friend of all its specialization.
Iterator_g_const & operator++()
Get the next element.
mem_id LinId(const grid_key_dx< N, ids_type > &gk, const char sum_id[N], const size_t(&bc)[N]) const
Linearization of the grid_key_dx with a specified shift.
__device__ __host__ bool isInsideKey(const KeyType &k) const
Check if the point is inside the region (Border included)
grid_key_dx_r & get()
Return the actual key.
Declaration grid_key_dx_iterator_sub.
Class to check if the edge can be created or not.
bool isNext()
Check if there is the next element.
__device__ __host__ grid_sm< N, T > & operator=(const grid_sm< N, T > &g)
Copy the grid from another grid.
Box< N, size_t > getBox() const
Return the box enclosing the grid.
__device__ __host__ void set_d(index_type i, index_type id)
Set the i index.
mem_id * k
structure that store all the index
mem_id LinId(mem_id *id) const
Linearization of an array of mem_id (long int)
Class to check if the edge can be created or not.
__device__ __host__ T getHigh(int i) const
get the high interval of the box
void Initialize(const size_t sz)
Initialize the basic structure.
void setDimensions(const size_t(&dims)[N])
Reset the dimension of the grid.
size_t getDim()
Get the dimensionality of the key.
grid_sm(const size_t(&sz)[N])
Construct a grid of a specified size.