Brick Library  0.1
Performance-portable stencil datalayout & codegen
brick.h
Go to the documentation of this file.
1 
6 #ifndef BRICK_H
7 #define BRICK_H
8 
9 #include <stdlib.h>
10 #include <type_traits>
11 #include "vecscatter.h"
12 
14 #define ALIGN 2048
15 
17 #if defined(__CUDACC__) || defined(__HIP__)
18 #define FORCUDA __host__ __device__
19 #else
20 #define FORCUDA
21 #endif
22 
28 template<unsigned base, unsigned exp>
30 struct static_power {
31  static constexpr unsigned value = base * static_power<base, exp - 1>::value;
32 };
33 
35 template<unsigned base>
36 struct static_power<base, 0> {
37  static constexpr unsigned value = 1;
38 };
48 struct BrickStorage {
56  long chunks;
58  long step;
60  void *mmap_info = nullptr;
61 
63  static BrickStorage allocate(long chunks, long step) {
64  BrickStorage b;
65  b.chunks = chunks;
66  b.step = step;
67  b.dat = (bElem *) aligned_alloc(ALIGN, chunks * step * sizeof(bElem));
68  return b;
69  }
70 
72  static BrickStorage mmap_alloc(long chunks, long step);
73 
75  static BrickStorage mmap_alloc(long chunks, long step, void *mmap_fd, size_t offset);
76 };
77 
90 template<unsigned dims>
91 struct BrickInfo {
93  typedef unsigned (*adjlist)[static_power<3, dims>::value];
97  unsigned nbricks;
98 
103  explicit BrickInfo(unsigned nbricks) : nbricks(nbricks) {
104  adj = (adjlist) malloc(nbricks * static_power<3, dims>::value * sizeof(unsigned));
105  }
106 
108  BrickStorage allocate(long step) {
109  return BrickStorage::allocate(nbricks, step);
110  }
111 
114  return BrickStorage::mmap_alloc(nbricks, step);
115  }
116 
118  BrickStorage mmap_alloc(long step, void *mmap_fd, size_t offset) {
119  return BrickStorage::mmap_alloc(nbricks, step, mmap_fd, offset);
120  }
121 };
122 
124 template<unsigned ... Ds>
125 struct Dim {
126 };
127 
136 template<unsigned ... xs>
137 struct cal_size;
138 
143 template<unsigned x>
144 struct cal_size<x> {
145  static constexpr unsigned value = x;
146 };
147 
153 template<unsigned x, unsigned ... xs>
154 struct cal_size<x, xs...> {
155  static constexpr unsigned value = x * cal_size<xs ...>::value;
156 };
167 template<unsigned ... offs>
168 struct cal_offs;
169 
174 template<unsigned off>
175 struct cal_offs<1, off> {
176  static constexpr unsigned value = off;
177 };
178 
185 template<unsigned dim, unsigned off, unsigned ...offs>
186 struct cal_offs<dim, off, offs...> {
187  static constexpr unsigned value = off * static_power<3, dim - 1>::value + cal_offs<dim - 1, offs...>::value;
188 };
207 template<typename...>
210 
212 template<typename T,
213  unsigned D,
214  unsigned F>
215 struct _BrickAccessor<T, Dim<D>, Dim<F>, bool> {
216  T *par;
217 
218  unsigned b;
219  unsigned pos;
220  unsigned nvec;
221  unsigned wvec;
222 
223  FORCUDA
224  _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec) :
225  par(par), b(b), pos(pos), nvec(nvec), wvec(wvec) {
226  }
227 
228  FORCUDA
229  inline bElem &operator[](unsigned i) {
230  // change pos
231  unsigned dir = i + D;
232  unsigned d = pos * 3 + dir / D;
233  // new vec position
234  unsigned l = dir % D;
235  unsigned w = wvec * F + l % F;
236  unsigned n = nvec * (D / F) + l / F;
237  unsigned offset = n * par->VECLEN + w;
238 
239  return par->dat[par->bInfo->adj[b][d] * par->step + offset];
240  }
241 };
242 
251 template<typename T,
252  unsigned D,
253  unsigned F,
254  unsigned ... BDims,
255  unsigned ... Folds>
256 struct _BrickAccessor<T, Dim<D, BDims...>, Dim<F, Folds...>, bool> {
257  T *par;
258 
259  unsigned b;
260  unsigned pos;
261  unsigned nvec;
262  unsigned wvec;
263 
264  FORCUDA
265  _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec) :
266  par(par), b(b), pos(pos), nvec(nvec), wvec(wvec) {
267  }
268 
269  FORCUDA
270  inline _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>, bool> operator[](unsigned i) {
271  // change pos
272  unsigned dir = i + D;
273  unsigned d = pos * 3 + dir / D;
274  // new vec position
275  unsigned l = dir % D;
276  unsigned w = wvec * F + l % F;
277  unsigned n = nvec * (D / F) + l / F;
278  return _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>, bool>(par, b, d, n, w);
279  }
280 };
281 
290 template<typename T,
291  unsigned D,
292  unsigned ... BDims,
293  unsigned ... Folds>
294 struct _BrickAccessor<T, Dim<D, BDims...>, Dim<Folds...>, void> {
295  T *par;
296 
297  unsigned b;
298  unsigned pos;
299  unsigned nvec;
300  unsigned wvec;
301 
302  FORCUDA
303  _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec) :
304  par(par), b(b), pos(pos), nvec(nvec), wvec(wvec) {
305  }
306 
307  FORCUDA
308  inline _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>,
309  typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>
310  operator[](unsigned i) {
311  // change pos
312  unsigned dir = i + D;
313  unsigned d = pos * 3 + dir / D;
314  // new vec position
315  unsigned l = dir % D;
316  unsigned w = wvec;
317  unsigned n = nvec * D + l;
318  return _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>,
319  typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>(par, b, d, n, w);
320  }
321 };
332 template<typename...>
334 struct Brick;
335 
347 template<
348  unsigned ... BDims,
349  unsigned ... Folds>
350 struct Brick<Dim<BDims...>, Dim<Folds...> > {
351  typedef Brick<Dim<BDims...>, Dim<Folds...> > mytype;
352  typedef BrickInfo<sizeof...(BDims)> myBrickInfo;
353 
354  static constexpr unsigned VECLEN = cal_size<Folds...>::value;
355  static constexpr unsigned BRICKSIZE = cal_size<BDims...>::value;
356 
358  unsigned step;
360 
362  FORCUDA
363  inline _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
364  typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type> operator[](unsigned b) {
365  return _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
366  typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>(this, b, 0, 0, 0);
367  }
368 
370  template<unsigned ... Offsets>
371  FORCUDA
372  inline bElem *neighbor(unsigned b) {
373  unsigned off = cal_offs<sizeof...(BDims), Offsets...>::value;
374  return &dat[bInfo->adj[b][off] * step];
375  }
376 
383  FORCUDA
384  Brick(myBrickInfo *bInfo, const BrickStorage &bStorage, unsigned offset) : bInfo(bInfo) {
385  dat = bStorage.dat + offset;
386  step = (unsigned) bStorage.step;
387  }
388 };
391 #endif //BRICK_H
BrickStorage::dat
bElem * dat
Pointer holding brick data.
Definition: brick.h:50
BrickInfo::BrickInfo
BrickInfo(unsigned nbricks)
Creating an empty metadata consisting of the specified number of bricks.
Definition: brick.h:103
_BrickAccessor< T, Dim< D, BDims... >, Dim< F, Folds... >, bool >::pos
unsigned pos
Accumulative position within adjacency list.
Definition: brick.h:260
BrickInfo::mmap_alloc
BrickStorage mmap_alloc(long step)
Allocate a new brick storage BrickStorage::mmap_alloc(long, long)
Definition: brick.h:113
Brick< Dim< BDims... >, Dim< Folds... > >::dat
bElem * dat
Offsetted memory (BrickStorage)
Definition: brick.h:359
BrickInfo::mmap_alloc
BrickStorage mmap_alloc(long step, void *mmap_fd, size_t offset)
Allocate a new brick storage BrickStorage::mmap_alloc(long, long, void*, size_t)
Definition: brick.h:118
Brick< Dim< BDims... >, Dim< Folds... > >::Brick
FORCUDA Brick(myBrickInfo *bInfo, const BrickStorage &bStorage, unsigned offset)
Initialize a brick data structure.
Definition: brick.h:384
_BrickAccessor< T, Dim< D, BDims... >, Dim< Folds... >, void >::operator[]
FORCUDA _BrickAccessor< T, Dim< BDims... >, Dim< Folds... >, typename std::conditional< sizeof...(BDims)==sizeof...(Folds), bool, void >::type > operator[](unsigned i)
Definition: brick.h:310
static_power::value
static constexpr unsigned value
Definition: brick.h:31
Brick< Dim< BDims... >, Dim< Folds... > >::operator[]
FORCUDA _BrickAccessor< mytype, Dim< BDims... >, Dim< Folds... >, typename std::conditional< sizeof...(BDims)==sizeof...(Folds), bool, void >::type > operator[](unsigned b)
Indexing operator returns: Accessing brick elements using [].
Definition: brick.h:364
BrickStorage::step
long step
Size of a chunk in number of elements.
Definition: brick.h:58
base
Definition: base.py:1
bElem
#define bElem
Basic datatype for all brick elements.
Definition: vecscatter.h:12
_BrickAccessor< T, Dim< D >, Dim< F >, bool >::wvec
unsigned wvec
Position within a vector.
Definition: brick.h:221
_BrickAccessor< T, Dim< D, BDims... >, Dim< Folds... >, void >::par
T * par
parent Brick data structure reference
Definition: brick.h:295
_BrickAccessor< T, Dim< D >, Dim< F >, bool >::nvec
unsigned nvec
Which vector.
Definition: brick.h:220
_BrickAccessor< T, Dim< D >, Dim< F >, bool >::par
T * par
parent Brick data structure reference
Definition: brick.h:216
_BrickAccessor
Generic base template for Accessing brick elements using [].
Definition: brick.h:209
Brick< Dim< BDims... >, Dim< Folds... > >::neighbor
FORCUDA bElem * neighbor(unsigned b)
Return the adjacency list of brick b
Definition: brick.h:372
BrickInfo::allocate
BrickStorage allocate(long step)
Allocate a new brick storage BrickStorage::allocate()
Definition: brick.h:108
_BrickAccessor< T, Dim< D, BDims... >, Dim< F, Folds... >, bool >::par
T * par
parent Brick data structure reference
Definition: brick.h:257
_BrickAccessor< T, Dim< D >, Dim< F >, bool >::operator[]
FORCUDA bElem & operator[](unsigned i)
Definition: brick.h:229
Dim
Empty template to specify an n-D list.
Definition: brick.h:125
vecscatter.h
Interface to code generator.
_BrickAccessor< T, Dim< D, BDims... >, Dim< F, Folds... >, bool >::nvec
unsigned nvec
Which vector.
Definition: brick.h:261
BrickInfo::nbricks
unsigned nbricks
Number of bricks in this list.
Definition: brick.h:97
_BrickAccessor< T, Dim< D >, Dim< F >, bool >::pos
unsigned pos
Accumulative position within adjacency list.
Definition: brick.h:219
BrickInfo::adj
adjlist adj
Adjacency list.
Definition: brick.h:95
_BrickAccessor< T, Dim< D, BDims... >, Dim< Folds... >, void >::b
unsigned b
Reference (center) brick.
Definition: brick.h:297
FORCUDA
#define FORCUDA
Overloaded attributes for potentially GPU-usable functions (in place of host device etc....
Definition: brick.h:20
_BrickAccessor< T, Dim< D, BDims... >, Dim< F, Folds... >, bool >::wvec
unsigned wvec
Position within a vector.
Definition: brick.h:262
cal_offs
Generic base template for Calculating the offset within the adjacency list.
Definition: brick.h:168
_BrickAccessor< T, Dim< D, BDims... >, Dim< F, Folds... >, bool >::_BrickAccessor
FORCUDA _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec)
Definition: brick.h:265
Brick< Dim< BDims... >, Dim< Folds... > >
Brick data structure.
Definition: brick.h:350
Brick< Dim< BDims... >, Dim< Folds... > >::bInfo
myBrickInfo * bInfo
Pointer to (possibly shared) metadata.
Definition: brick.h:357
_BrickAccessor< T, Dim< D, BDims... >, Dim< Folds... >, void >::nvec
unsigned nvec
Which vector.
Definition: brick.h:299
_BrickAccessor< T, Dim< D, BDims... >, Dim< F, Folds... >, bool >::b
unsigned b
Reference (center) brick.
Definition: brick.h:259
ALIGN
#define ALIGN
BrickStorage allocation alignment.
Definition: brick.h:14
BrickStorage::chunks
long chunks
Number of chunks.
Definition: brick.h:56
static_power
Compute Statically compute exponentials.
Definition: brick.h:30
_BrickAccessor< T, Dim< D >, Dim< F >, bool >::b
unsigned b
Reference (center) brick.
Definition: brick.h:218
_BrickAccessor< T, Dim< D, BDims... >, Dim< Folds... >, void >::wvec
unsigned wvec
Position within a vector.
Definition: brick.h:300
BrickStorage::allocate
static BrickStorage allocate(long chunks, long step)
Allocation using *alloc.
Definition: brick.h:63
Brick
Generic base template, see Brick< Dim< BDims... >, Dim< Folds... > >
Definition: brick.h:334
BrickStorage
Initializing and holding the storage of bricks.
Definition: brick.h:48
_BrickAccessor< T, Dim< D, BDims... >, Dim< Folds... >, void >::_BrickAccessor
FORCUDA _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec)
Definition: brick.h:303
cal_size
Generic base template for Calculate the product of n numbers in a template.
Definition: brick.h:137
_BrickAccessor< T, Dim< D, BDims... >, Dim< Folds... >, void >::pos
unsigned pos
Accumulative position within adjacency list.
Definition: brick.h:298
Brick< Dim< BDims... >, Dim< Folds... > >::mytype
Brick< Dim< BDims... >, Dim< Folds... > > mytype
Shorthand for this struct's type.
Definition: brick.h:351
BrickStorage::mmap_info
void * mmap_info
MMAP data structure when using mmap as allocator.
Definition: brick.h:60
_BrickAccessor< T, Dim< D >, Dim< F >, bool >::_BrickAccessor
FORCUDA _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec)
Definition: brick.h:224
Brick< Dim< BDims... >, Dim< Folds... > >::step
unsigned step
Spacing between bricks in unit of bElem (BrickStorage)
Definition: brick.h:358
BrickInfo::adjlist
unsigned(* adjlist)[static_power< 3, dims >::value]
Adjacency list type.
Definition: brick.h:93
_BrickAccessor< T, Dim< D, BDims... >, Dim< F, Folds... >, bool >::operator[]
FORCUDA _BrickAccessor< T, Dim< BDims... >, Dim< Folds... >, bool > operator[](unsigned i)
Definition: brick.h:270
BrickInfo
Metadata related to bricks.
Definition: brick.h:91
BrickStorage::mmap_alloc
static BrickStorage mmap_alloc(long chunks, long step)
mmap allocator using default (new) file
Definition: memfd.cpp:94