Brick Library 0.1
Performance-portable stencil datalayout & codegen
brick-cuda.h
Go to the documentation of this file.
1
6#ifndef BRICK_BRICK_CUDA_H
7#define BRICK_BRICK_CUDA_H
8
9#include <cassert>
10#include <brick.h>
11#include <cuda_runtime.h>
12
16#ifndef NDEBUG
17#define cudaCheck(x) x
18#else
19
20#include <cstdio>
21
22#define cudaCheck(x) _cudaCheck(x, #x ,__FILE__, __LINE__)
23#endif
24
25
27template<typename T>
28void _cudaCheck(T e, const char *func, const char *call, const int line) {
29 if (e != cudaSuccess) {
30 printf("\"%s\" at %d in %s\n\treturned %d\n-> %s\n", func, line, call, (int) e, cudaGetErrorString(e));
31 exit(EXIT_FAILURE);
32 }
33}
34
42template<unsigned dims>
43BrickInfo<dims> movBrickInfo(BrickInfo<dims> &bInfo, cudaMemcpyKind kind) {
44 assert(kind == cudaMemcpyHostToDevice || kind == cudaMemcpyDeviceToHost);
45
46 // Make a copy
47 BrickInfo<dims> ret = bInfo;
48 size_t size = bInfo.nbricks * static_power<3, dims>::value * sizeof(unsigned);
49
50 if (kind == cudaMemcpyHostToDevice) {
51 cudaCheck(cudaMalloc(&ret.adj, size));
52 } else {
53 ret.adj = (typename BrickInfo<dims>::adjlist) malloc(size);
54 }
55 cudaCheck(cudaMemcpy(ret.adj, bInfo.adj, size, kind));
56 return ret;
57}
58
65inline BrickStorage movBrickStorage(BrickStorage &bStorage, cudaMemcpyKind kind) {
66 assert(kind == cudaMemcpyHostToDevice || kind == cudaMemcpyDeviceToHost);
67
68 bool isToDevice = (kind == cudaMemcpyHostToDevice);
69 // Make a copy
70 BrickStorage ret = bStorage;
71 size_t size = bStorage.step * bStorage.chunks * sizeof(bElem);
72 bElem *datptr;
73 if (isToDevice) {
74 cudaCheck(cudaMalloc(&datptr, size));
75 } else {
76 datptr = (bElem *) malloc(size);
77 }
78 cudaCheck(cudaMemcpy(datptr, bStorage.dat.get(), size, kind));
79 if (isToDevice) {
80 ret.dat = std::shared_ptr<bElem>(datptr, [](bElem *p) { cudaFree(p); });
81 } else {
82 ret.dat = std::shared_ptr<bElem>(datptr, [](bElem *p) { free(p); });
83 }
84 return ret;
85}
86
87#include "dev_shl.h"
88
89#endif //BRICK_BRICK_CUDA_H
void _cudaCheck(T e, const char *func, const char *call, const int line)
Internal for cudaCheck(x)
Definition: brick-cuda.h:28
BrickInfo< dims > movBrickInfo(BrickInfo< dims > &bInfo, cudaMemcpyKind kind)
Moving BrickInfo to or from GPU (allocate new)
Definition: brick-cuda.h:43
BrickStorage movBrickStorage(BrickStorage &bStorage, cudaMemcpyKind kind)
Moving BrickStorage to or from GPU (allocate new)
Definition: brick-cuda.h:65
#define cudaCheck(x)
Check the return of CUDA calls, do nothing during release build.
Definition: brick-cuda.h:17
Main header for bricks.
Implementation for various shuffle implementations.
Definition: func.py:1
p
Definition: printer.py:127
Metadata related to bricks.
Definition: brick.h:92
unsigned nbricks
Number of bricks in this list.
Definition: brick.h:98
adjlist adj
Adjacency list.
Definition: brick.h:96
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
size_t step
Size of a chunk in number of elements.
Definition: brick.h:59
Compute Statically compute exponentials.
Definition: brick.h:31
#define bElem
Basic datatype for all brick elements.
Definition: vecscatter.h:13