brick-dpc.h

Go to the documentation of this file.

 
#ifndef BRICK_BRICK_DPC_H
#define BRICK_BRICK_DPC_H
 
#include <CL/sycl.hpp>
#include "brick.h"
 
#ifndef NDEBUG
static sycl::property_list properties{sycl::property::queue::enable_profiling()};
static sycl::queue gpu_queue = sycl::queue(sycl::gpu_selector(), properties);
static sycl::queue cpu_queue = sycl::queue(sycl::host_selector(), properties);
#else
static sycl::queue gpu_queue = sycl::queue(sycl::gpu_selector());
static sycl::queue cpu_queue = sycl::queue(sycl::host_selector());
#endif
 
enum dpcError_t {
    dpc_success,
    memcpy_failed,
    malloc_failed
};
 
enum dpcMemcpyKind {
    dpcMemcpyHostToDevice,
    dpcMemcpyDeviceToHost
};
 
template <typename T>
inline dpcError_t dpc_malloc(T **buffer, size_t size) {
    T *ptr = (T *) sycl::malloc_device(size, gpu_queue);
    if (ptr == nullptr) {
        return malloc_failed;
    }
    gpu_queue.memset(ptr, 0, size).wait_and_throw();
    (*buffer) = ptr;
    return dpc_success;
}
 
template <typename T>
inline dpcError_t dpc_memcpy(T *dst, T *ptr, size_t size, dpcMemcpyKind type) {    
    assert(type == dpcMemcpyHostToDevice);
    gpu_queue.memcpy((void *) dst, (void *) ptr, size).wait_and_throw();
    return dpc_success;
}
 
template <typename T>
inline dpcError_t dpc_free(T *ptr) {
    sycl::free((void *) ptr, gpu_queue);
    return dpc_success;
} 
 
static const char *dpc_get_error_string(dpcError_t e) {
    switch(e) {
        case dpc_success: return "DPC Success";
        case malloc_failed: return "Failed to allocate memory";
        case memcpy_failed: return "Failed to perform memcpy";
        default: return "Unknown error";
    }
}
 
#define gpuMalloc(p, s) dpc_malloc(p, s)
#define gpuMemcpy(d, p, s, k) dpc_memcpy(d, p, s, k)
#define gpuFree(p) dpc_free(p)
#define gpuMemcpyKind dpcMemcpyKind
#define gpuMemcpyHostToDevice dpcMemcpyHostToDevice
#define gpuMemcpyDeviceToHost dpcMemcpyDeviceToHost
#define gpuSuccess dpc_success
#define gpuGetErrorString(e) dpc_get_error_string(e)
 
template<typename...>
struct DPCBrick;
 
template<
    unsigned ... BDims,
    unsigned ... Folds>
struct DPCBrick<Dim<BDims...>, Dim<Folds...> > {
  typedef DPCBrick<Dim<BDims...>, Dim<Folds...> > mytype;    
  typedef BrickInfo<sizeof...(BDims)> myBrickInfo;        
 
  static constexpr unsigned VECLEN = cal_size<Folds...>::value;     
  static constexpr unsigned BRICKSIZE = cal_size<BDims...>::value;  
 
  myBrickInfo *bInfo;        
  size_t step;             
  bElem *dat;                
 
  FORCUDA
  inline _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
      typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type> operator[](unsigned b) {
    return _BrickAccessor<mytype, Dim<BDims...>, Dim<Folds...>,
        typename std::conditional<sizeof...(BDims) == sizeof...(Folds), bool, void>::type>(this, b, 0, 0, 0);
  }
 
  template<unsigned ... Offsets>
  FORCUDA
  inline bElem *neighbor(unsigned b) {
    unsigned off = cal_offs<sizeof...(BDims), Offsets...>::value;
    return &dat[bInfo->adj[b][off] * step];
  }
 
  DPCBrick(myBrickInfo *bInfo, bElem *bData, size_t bStep, unsigned offset) : bInfo(bInfo) {
    dat = bData + offset;
    step = bStep;
  }
};
 
#include "brick-gpu.h"
 
#endif //BRICK_BRICK_GPU_H