Brick Library 0.1
Performance-portable stencil datalayout & codegen
brick-dpc.h
Go to the documentation of this file.
1
8#ifndef BRICK_BRICK_DPC_H
9#define BRICK_BRICK_DPC_H
10
11#include <CL/sycl.hpp>
12#include "brick.h"
13
14#ifndef NDEBUG
15static sycl::property_list properties{sycl::property::queue::enable_profiling()};
16static sycl::queue gpu_queue = sycl::queue(sycl::gpu_selector(), properties);
17static sycl::queue cpu_queue = sycl::queue(sycl::host_selector(), properties);
18#else
19static sycl::queue gpu_queue = sycl::queue(sycl::gpu_selector());
20static sycl::queue cpu_queue = sycl::queue(sycl::host_selector());
21#endif
22
27};
28
32};
33
41template <typename T>
42inline dpcError_t dpc_malloc(T **buffer, size_t size) {
43 T *ptr = (T *) sycl::malloc_device(size, gpu_queue);
44 if (ptr == nullptr) {
45 return malloc_failed;
46 }
47 gpu_queue.memset(ptr, 0, size).wait_and_throw();
48 (*buffer) = ptr;
49 return dpc_success;
50}
51
61template <typename T>
62inline dpcError_t dpc_memcpy(T *dst, T *ptr, size_t size, dpcMemcpyKind type) {
63 assert(type == dpcMemcpyHostToDevice);
64 gpu_queue.memcpy((void *) dst, (void *) ptr, size).wait_and_throw();
65 return dpc_success;
66}
67
74template <typename T>
75inline dpcError_t dpc_free(T *ptr) {
76 sycl::free((void *) ptr, gpu_queue);
77 return dpc_success;
78}
79
80static const char *dpc_get_error_string(dpcError_t e) {
81 switch(e) {
82 case dpc_success: return "DPC Success";
83 case malloc_failed: return "Failed to allocate memory";
84 case memcpy_failed: return "Failed to perform memcpy";
85 default: return "Unknown error";
86 }
87}
88
89#define gpuMalloc(p, s) dpc_malloc(p, s)
90#define gpuMemcpy(d, p, s, k) dpc_memcpy(d, p, s, k)
91#define gpuFree(p) dpc_free(p)
92#define gpuMemcpyKind dpcMemcpyKind
93#define gpuMemcpyHostToDevice dpcMemcpyHostToDevice
94#define gpuMemcpyDeviceToHost dpcMemcpyDeviceToHost
95#define gpuSuccess dpc_success
96#define gpuGetErrorString(e) dpc_get_error_string(e)
97
99template<typename...>
100struct DPCBrick;
101
113template<
114 unsigned ... BDims,
115 unsigned ... Folds>
116struct DPCBrick<Dim<BDims...>, Dim<Folds...> > {
117 typedef DPCBrick<Dim<BDims...>, Dim<Folds...> > mytype;
118 typedef BrickInfo<sizeof...(BDims)> myBrickInfo;
119
120 static constexpr unsigned VECLEN = cal_size<Folds...>::value;
121 static constexpr unsigned BRICKSIZE = cal_size<BDims...>::value;
122
124 size_t step;
126
128 FORCUDA
129 inline _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
130 typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type> operator[](unsigned b) {
131 return _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
132 typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>(this, b, 0, 0, 0);
133 }
134
136 template<unsigned ... Offsets>
137 FORCUDA
138 inline bElem *neighbor(unsigned b) {
139 unsigned off = cal_offs<sizeof...(BDims), Offsets...>::value;
140 return &dat[bInfo->adj[b][off] * step];
141 }
142
143 DPCBrick(myBrickInfo *bInfo, bElem *bData, size_t bStep, unsigned offset) : bInfo(bInfo) {
144 dat = bData + offset;
145 step = bStep;
146 }
147};
148
149#include "brick-gpu.h"
150
151#endif //BRICK_BRICK_GPU_H
dpcError_t
Definition: brick-dpc.h:23
@ dpc_success
Definition: brick-dpc.h:24
@ memcpy_failed
Definition: brick-dpc.h:25
@ malloc_failed
Definition: brick-dpc.h:26
dpcMemcpyKind
Definition: brick-dpc.h:29
@ dpcMemcpyHostToDevice
Definition: brick-dpc.h:30
@ dpcMemcpyDeviceToHost
Definition: brick-dpc.h:31
dpcError_t dpc_malloc(T **buffer, size_t size)
Interface to allocate memory on a DPCPP GPU with a similar footprint to CUDA.
Definition: brick-dpc.h:42
dpcError_t dpc_memcpy(T *dst, T *ptr, size_t size, dpcMemcpyKind type)
Copy data from host to DPCPP GPU. If data must be returned to the host after kernel execution,...
Definition: brick-dpc.h:62
dpcError_t dpc_free(T *ptr)
Free allocated memory on a DPCPP GPU.
Definition: brick-dpc.h:75
This file should not be directly included. It defines instructions for using bricklib with a GPU,...
Main header for bricks.
#define FORCUDA
Overloaded attributes for potentially GPU-usable functions (in place of host device etc....
Definition: brick.h:25
Metadata related to bricks.
Definition: brick.h:97
adjlist adj
Adjacency list.
Definition: brick.h:101
Brick data structure.
Definition: brick-dpc.h:116
bElem * dat
Offsetted memory (BrickStorage)
Definition: brick-dpc.h:125
size_t step
Spacing between bricks in unit of bElem (BrickStorage)
Definition: brick-dpc.h:124
myBrickInfo * bInfo
Pointer to (possibly shared) metadata.
Definition: brick-dpc.h:123
DPCBrick(myBrickInfo *bInfo, bElem *bData, size_t bStep, unsigned offset)
Definition: brick-dpc.h:143
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-dpc.h:130
DPCBrick< Dim< BDims... >, Dim< Folds... > > mytype
Shorthand for this struct's type.
Definition: brick-dpc.h:117
FORCUDA bElem * neighbor(unsigned b)
Return the adjacency list of brick b
Definition: brick-dpc.h:138
Generic base template, see Brick< Dim< BDims... >, Dim< Folds... > >
Definition: brick-dpc.h:100
Empty template to specify an n-D list.
Definition: brick.h:131
Generic base template for Accessing brick elements using [].
Definition: brick.h:215
Generic base template for Calculating the offset within the adjacency list.
Definition: brick.h:174
Generic base template for Calculate the product of n numbers in a template.
Definition: brick.h:143
#define bElem
Basic datatype for all brick elements.
Definition: vecscatter.h:13