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 <memory>
12#include "vecscatter.h"
13
15#define ALIGN 2048
16
18#if defined(__CUDACC__) || defined(__HIP__)
19#define FORCUDA __host__ __device__
20#else
21#define FORCUDA
22#endif
23
30template<unsigned base, unsigned exp>
32 static constexpr unsigned value = base * static_power<base, exp - 1>::value;
33};
34
36template<unsigned base>
37struct static_power<base, 0> {
38 static constexpr unsigned value = 1;
39};
51 std::shared_ptr<bElem> dat;
57 long chunks;
59 size_t step;
61 void *mmap_info = nullptr;
62
64 static BrickStorage allocate(long chunks, size_t step) {
66 b.chunks = chunks;
67 b.step = step;
68 b.dat = std::shared_ptr<bElem>((bElem*)aligned_alloc(ALIGN, chunks * step * sizeof(bElem)), free);
69 return b;
70 }
71
73 static BrickStorage mmap_alloc(long chunks, long step);
74
76 static BrickStorage mmap_alloc(long chunks, long step, void *mmap_fd, size_t offset);
77};
78
91template<unsigned dims>
92struct BrickInfo {
98 unsigned nbricks;
99
104 explicit BrickInfo(unsigned nbricks) : nbricks(nbricks) {
105 adj = (adjlist) malloc(nbricks * static_power<3, dims>::value * sizeof(unsigned));
106 }
107
110 return BrickStorage::allocate(nbricks, step);
111 }
112
115 return BrickStorage::mmap_alloc(nbricks, step);
116 }
117
119 BrickStorage mmap_alloc(long step, void *mmap_fd, size_t offset) {
120 return BrickStorage::mmap_alloc(nbricks, step, mmap_fd, offset);
121 }
122};
123
125template<unsigned ... Ds>
126struct Dim {
127};
128
137template<unsigned ... xs>
138struct cal_size;
139
144template<unsigned x>
145struct cal_size<x> {
146 static constexpr unsigned value = x;
147};
148
154template<unsigned x, unsigned ... xs>
155struct cal_size<x, xs...> {
156 static constexpr unsigned value = x * cal_size<xs ...>::value;
157};
168template<unsigned ... offs>
169struct cal_offs;
170
175template<unsigned off>
176struct cal_offs<1, off> {
177 static constexpr unsigned value = off;
178};
179
186template<unsigned dim, unsigned off, unsigned ...offs>
187struct cal_offs<dim, off, offs...> {
188 static constexpr unsigned value = off * static_power<3, dim - 1>::value + cal_offs<dim - 1, offs...>::value;
189};
209template<typename...>
211
213template<typename T,
214 unsigned D,
215 unsigned F>
216struct _BrickAccessor<T, Dim<D>, Dim<F>, bool> {
217 T *par;
218
219 unsigned b;
220 unsigned pos;
221 unsigned nvec;
222 unsigned wvec;
223
224 FORCUDA
225 _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec) :
226 par(par), b(b), pos(pos), nvec(nvec), wvec(wvec) {
227 }
228
229 FORCUDA
230 inline bElem &operator[](unsigned i) {
231 // change pos
232 unsigned dir = i + D;
233 unsigned d = pos * 3 + dir / D;
234 // new vec position
235 unsigned l = dir % D;
236 unsigned w = wvec * F + l % F;
237 unsigned n = nvec * (D / F) + l / F;
238 unsigned offset = n * par->VECLEN + w;
239
240 return par->dat[par->bInfo->adj[b][d] * par->step + offset];
241 }
242};
243
252template<typename T,
253 unsigned D,
254 unsigned F,
255 unsigned ... BDims,
256 unsigned ... Folds>
257struct _BrickAccessor<T, Dim<D, BDims...>, Dim<F, Folds...>, bool> {
258 T *par;
259
260 unsigned b;
261 unsigned pos;
262 unsigned nvec;
263 unsigned wvec;
264
265 FORCUDA
266 _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec) :
267 par(par), b(b), pos(pos), nvec(nvec), wvec(wvec) {
268 }
269
270 FORCUDA
271 inline _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>, bool> operator[](unsigned i) {
272 // change pos
273 unsigned dir = i + D;
274 unsigned d = pos * 3 + dir / D;
275 // new vec position
276 unsigned l = dir % D;
277 unsigned w = wvec * F + l % F;
278 unsigned n = nvec * (D / F) + l / F;
279 return _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>, bool>(par, b, d, n, w);
280 }
281};
282
291template<typename T,
292 unsigned D,
293 unsigned ... BDims,
294 unsigned ... Folds>
295struct _BrickAccessor<T, Dim<D, BDims...>, Dim<Folds...>, void> {
296 T *par;
297
298 unsigned b;
299 unsigned pos;
300 unsigned nvec;
301 unsigned wvec;
302
303 FORCUDA
304 _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec) :
305 par(par), b(b), pos(pos), nvec(nvec), wvec(wvec) {
306 }
307
308 FORCUDA
309 inline _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>,
310 typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>
311 operator[](unsigned i) {
312 // change pos
313 unsigned dir = i + D;
314 unsigned d = pos * 3 + dir / D;
315 // new vec position
316 unsigned l = dir % D;
317 unsigned w = wvec;
318 unsigned n = nvec * D + l;
319 return _BrickAccessor<T, Dim<BDims...>, Dim<Folds...>,
320 typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>(par, b, d, n, w);
321 }
322};
334template<typename...>
335struct Brick;
336
348template<
349 unsigned ... BDims,
350 unsigned ... Folds>
351struct Brick<Dim<BDims...>, Dim<Folds...> > {
352 typedef Brick<Dim<BDims...>, Dim<Folds...> > mytype;
353 typedef BrickInfo<sizeof...(BDims)> myBrickInfo;
354
355 static constexpr unsigned VECLEN = cal_size<Folds...>::value;
356 static constexpr unsigned BRICKSIZE = cal_size<BDims...>::value;
357
359 size_t step;
362
364 FORCUDA
365 inline _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
366 typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type> operator[](unsigned b) {
367 return _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
368 typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>(this, b, 0, 0, 0);
369 }
370
372 template<unsigned ... Offsets>
373 FORCUDA
374 inline bElem *neighbor(unsigned b) {
375 unsigned off = cal_offs<sizeof...(BDims), Offsets...>::value;
376 return &dat[bInfo->adj[b][off] * step];
377 }
378
385 Brick(myBrickInfo *bInfo, const BrickStorage &brickStorage, unsigned offset) : bInfo(bInfo) {
386 bStorage = brickStorage;
387 dat = bStorage.dat.get() + offset;
388 step = (unsigned) bStorage.step;
389 }
390};
393#endif //BRICK_H
#define FORCUDA
Overloaded attributes for potentially GPU-usable functions (in place of host device etc....
Definition: brick.h:21
#define ALIGN
BrickStorage allocation alignment.
Definition: brick.h:15
Definition: base.py:1
Metadata related to bricks.
Definition: brick.h:92
unsigned(* adjlist)[static_power< 3, dims >::value]
Adjacency list type.
Definition: brick.h:94
unsigned nbricks
Number of bricks in this list.
Definition: brick.h:98
adjlist adj
Adjacency list.
Definition: brick.h:96
BrickInfo(unsigned nbricks)
Creating an empty metadata consisting of the specified number of bricks.
Definition: brick.h:104
BrickStorage mmap_alloc(long step)
Allocate a new brick storage BrickStorage::mmap_alloc(long, long)
Definition: brick.h:114
BrickStorage allocate(long step)
Allocate a new brick storage BrickStorage::allocate()
Definition: brick.h:109
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:119
Initializing and holding the storage of bricks.
Definition: brick.h:49
std::shared_ptr< bElem > dat
Pointer holding brick data.
Definition: brick.h:51
long chunks
Number of chunks.
Definition: brick.h:57
static BrickStorage allocate(long chunks, size_t step)
Allocation using *alloc.
Definition: brick.h:64
void * mmap_info
MMAP data structure when using mmap as allocator.
Definition: brick.h:61
static BrickStorage mmap_alloc(long chunks, long step)
mmap allocator using default (new) file
Definition: memfd.cpp:103
size_t step
Size of a chunk in number of elements.
Definition: brick.h:59
Brick data structure.
Definition: brick.h:351
BrickStorage bStorage
Definition: brick.h:361
Brick< Dim< BDims... >, Dim< Folds... > > mytype
Shorthand for this struct's type.
Definition: brick.h:352
FORCUDA bElem * neighbor(unsigned b)
Return the adjacency list of brick b
Definition: brick.h:374
size_t step
Spacing between bricks in unit of bElem (BrickStorage)
Definition: brick.h:359
myBrickInfo * bInfo
Pointer to (possibly shared) metadata.
Definition: brick.h:358
bElem * dat
Offsetted memory (BrickStorage)
Definition: brick.h:360
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:366
Brick(myBrickInfo *bInfo, const BrickStorage &brickStorage, unsigned offset)
Initialize a brick data structure.
Definition: brick.h:385
Generic base template, see Brick< Dim< BDims... >, Dim< Folds... > >
Definition: brick.h:335
Empty template to specify an n-D list.
Definition: brick.h:126
unsigned nvec
Which vector.
Definition: brick.h:262
unsigned wvec
Position within a vector.
Definition: brick.h:263
T * par
parent Brick data structure reference
Definition: brick.h:258
unsigned pos
Accumulative position within adjacency list.
Definition: brick.h:261
FORCUDA _BrickAccessor< T, Dim< BDims... >, Dim< Folds... >, bool > operator[](unsigned i)
Definition: brick.h:271
FORCUDA _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec)
Definition: brick.h:266
unsigned b
Reference (center) brick.
Definition: brick.h:260
FORCUDA _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec)
Definition: brick.h:304
unsigned nvec
Which vector.
Definition: brick.h:300
unsigned b
Reference (center) brick.
Definition: brick.h:298
unsigned pos
Accumulative position within adjacency list.
Definition: brick.h:299
T * par
parent Brick data structure reference
Definition: brick.h:296
FORCUDA _BrickAccessor< T, Dim< BDims... >, Dim< Folds... >, typename std::conditional< sizeof...(BDims)==sizeof...(Folds), bool, void >::type > operator[](unsigned i)
Definition: brick.h:311
unsigned wvec
Position within a vector.
Definition: brick.h:301
T * par
parent Brick data structure reference
Definition: brick.h:217
FORCUDA bElem & operator[](unsigned i)
Definition: brick.h:230
unsigned wvec
Position within a vector.
Definition: brick.h:222
FORCUDA _BrickAccessor(T *par, unsigned b, unsigned pos, unsigned nvec, unsigned wvec)
Definition: brick.h:225
unsigned b
Reference (center) brick.
Definition: brick.h:219
unsigned nvec
Which vector.
Definition: brick.h:221
unsigned pos
Accumulative position within adjacency list.
Definition: brick.h:220
Generic base template for Accessing brick elements using [].
Definition: brick.h:210
Generic base template for Calculating the offset within the adjacency list.
Definition: brick.h:169
Generic base template for Calculate the product of n numbers in a template.
Definition: brick.h:138
Compute Statically compute exponentials.
Definition: brick.h:31
static constexpr unsigned value
Definition: brick.h:32
Interface to code generator.
#define bElem
Basic datatype for all brick elements.
Definition: vecscatter.h:13