moving helps to examples/util

examples/util/cuda_helpers.cuh

@@ -0,0 +1,25 @@
+#pragma once
+
+#define programCount 32
+#define programIndex (threadIdx.x & 31)
+#define taskIndex0 (blockIdx.x*4 + (threadIdx.x >> 5))
+#define taskCount0 (gridDim.x*4)
+#define taskIndex1 (blockIdx.y)
+#define taskCount1 (gridDim.y)
+#define taskIndex2 (blockIdx.z)
+#define taskCount2 (gridDim.z)
+#define taskIndex (taskIndex0 + taskCount0*(taskIndex1 + taskCount1*taskIndex2))
+#define taskCount (taskCount0*taskCount1*taskCount2)
+#define warpIdx (threadIdx.x >> 5)
+#define launch(ntx,nty,ntz,func) if (programIndex==0) func<<<dim3(((ntx)+4-1)/4,nty,ntz),128>>>
+#define sync cudaDeviceSynchronize()
+#define cif if
+__device__ __forceinline__ static double __shfl(double x, int lane)
+{
+  return __hiloint2double(
+      __shfl_xor(__double2hiint(x), lane),
+      __shfl_xor(__double2loint(x), lane));
+
+}
+#define shuffle(x,y) __shfl(x,y)
+#define broadcast(x,y) __shfl(x,y)

examples/util/ispc_malloc.cpp

@@ -0,0 +1,54 @@
+#include <cstdio>
+#include <cstdlib>
+#include <iostream>
+#include <cstring>
+#include "ispc_malloc.h"
+
+#ifdef _CUDA_
+
+void * operator new(size_t size) throw(std::bad_alloc)
+{
+  void *ptr;
+  ispc_malloc(&ptr, size);
+  return ptr;
+}
+void operator delete(void *ptr) throw()
+{
+  ispc_free(ptr);
+}
+
+#else
+
+void ispc_malloc(void **ptr, const size_t size)
+{
+  *ptr = malloc(size);
+}
+void ispc_free(void *ptr)
+{
+  free(ptr);
+}
+void ispc_memset(void *ptr, int value, size_t size)
+{
+  memset(ptr, value, size);
+}
+void ispcSetMallocHeapLimit(size_t value)
+{
+}
+void ispcSetStackLimit(size_t value)
+{
+}
+unsigned long long ispcGetMallocHeapLimit()
+{
+  return -1;
+}
+unsigned long long ispcGetStackLimit()
+{
+  return -1;
+}
+void * ispcMemcpy(void *dest,  void *src,  size_t num)
+{
+  memcpy(dest, src, num);
+  return dest;
+}
+
+#endif

examples/util/ispc_malloc.h

@@ -0,0 +1,10 @@
+#pragma once
+
+extern void ispc_malloc(void **ptr, const size_t size);
+extern void ispc_free(void *ptr);
+extern void ispc_memset(void *ptr, int value, size_t size);
+extern void ispcSetMallocHeapLimit(size_t value);
+extern void ispcSetStackLimit(size_t value);
+extern unsigned long long ispcGetMallocHeapLimit();
+extern unsigned long long ispcGetStackLimit();
+extern void * ispcMemcpy(void *dest,  void *src,  size_t num);

examples/util/nvcc_helpers.cu

@@ -0,0 +1,43 @@
+#ifndef _CUDA_
+#error "Something went wrong..."
+#endif
+
+void ispc_malloc(void **ptr, const size_t size)
+{
+  cudaMallocManaged(ptr, size);
+}
+void ispc_free(void *ptr)
+{
+  cudaFree(ptr);
+}
+void ispc_memset(void *ptr, int value, size_t size)
+{
+  cudaMemset(ptr, value, size);
+}
+void ispcSetMallocHeapLimit(size_t value)
+{
+  cudaDeviceSetLimit(cudaLimitMallocHeapSize,value);
+}
+void ispcSetStackLimit(size_t value)
+{
+  cudaDeviceSetLimit(cudaLimitStackSize,value);
+}
+unsigned long long ispcGetMallocHeapLimit()
+{
+  size_t value;
+  cudaDeviceGetLimit(&value, cudaLimitMallocHeapSize);
+  return value;
+}
+unsigned long long ispcGetStackLimit()
+{
+  size_t value;
+  cudaDeviceGetLimit(&value, cudaLimitStackSize);
+  return value;
+}
+void * ispcMemcpy(void *dest,  void *src,  size_t num)
+{
+  cudaMemcpy(dest, src, num, cudaMemcpyDefault);
+  return dest;
+}
+
+