Commit 12c37579 authored by incardon's avatar incardon

PPM Core integration with thrust + boost::array, compile, to debug

parent a8f1cdfa
......@@ -12,7 +12,6 @@ m4_ifdef([MYSQL_FOUND],,[m4_include([m4/ax_lib_mysql.m4])])
m4_ifdef([AX_CHECK_COMPILER_FLAGS],,[m4_include([m4/ax_check_compiler_flags.m4])])
m4_ifdef([ACX_PTHREAD],,[m4_include([m4/acx_pthread.m4])])
m4_ifdef([AX_CHECK_CL],,[m4_include([m4/ax_opencl.m4])])
m4_ifdef([AC_PROG_SED],,[m4_include([m4/programs.m4])])
m4_ifdef([AX_BOOST_BASE],,[m4_include([m4/ax_boost_base.m4])])
m4_ifdef([AX_BOOST_PROGRAM_OPTIONS],,[m4_include([m4/ax_boost_program_options.m4])])
m4_ifdef([AX_BOOST_THREAD],,[m4_include([m4/ax_boost_thread.m4])])
......@@ -20,8 +19,10 @@ m4_ifdef([ACX_MPI],,[m4_include([m4/acx_mpi.m4])])
m4_ifdef([AX_OPENMP],,[m4_include([m4/ax_openmp.m4])])
m4_ifdef([AX_GCC_X86_CPUID],,[m4_include([m4/ax_gcc_x86_cpuid.m4])])
m4_ifdef([AX_GCC_ARCHFLAG],,[m4_include([m4/ax_gcc_archflag.m4])])
m4_ifdef([AX_CUDA],,[m4_include([m4/ax_cuda.m4])])
CXXFLAGS+=" --std=c++11 "
NVCCFLAGS=" "
# Checks for programs.
AC_PROG_CXX
......@@ -53,10 +54,40 @@ if test x"$debuger" = x"yes"; then
AC_DEFINE([DEBUG_MODE],[],[Debug])
AC_DEFINE([DEBUG],[],[Debug])
CXXFLAGS="$CXXFLAGS -g3 -Wall -O0 "
NVCCFLAGS+="$NVCCFLAGS -g -O0 "
else
CXXFLAGS="$CXXFLAGS -Wall -O3 -g3 "
NVCCFLAGS+="$NVCCFLAGS -O3 "
fi
####### Checking for GPU support
AX_CUDA
AC_MSG_CHECKING(whether to build with GPU support)
gpu_support=no
AC_ARG_ENABLE(gpu,
AC_HELP_STRING(
[--enable-gpu],
[enable gpu support]
),
gpu_support="$enableval"
)
AC_MSG_RESULT($gpu_support)
if test x"$gpu_support" = x"yes"; then
AC_DEFINE([GPU],[],[GPU support])
fi
# Set this conditional if cuda is wanted
AM_CONDITIONAL(BUILDCUDA, test ! x$NVCC = x"no")
###########################
no_avx=no
no_sse42=no
no_sse41=no
......@@ -73,6 +104,7 @@ AX_CHECK_COMPILER_FLAGS([-msse],[CXXFLAGS="$CXXFLAGS -msse"],[no_sse=yes])
AX_CHECK_COMPILER_FLAGS([-mmmx],[CXXFLAGS="$CXXFLAGS -mmmx"],[no_mmx=yes])
AX_CHECK_COMPILER_FLAGS([-Wno-unused-but-set-variable],[CXXFLAGS="$CXXFLAGS -Wno-unused-but-set-variable"],[])
AC_SUBST(NVCCFLAGS)
# Checks for typedefs, structures, and compiler characteristics.
......@@ -136,5 +168,11 @@ if [ test x"$no_64" = x"no" ]; then
else
echo "* 64 bit: no *"
fi
if [ test x"$gpu_support" = x"no" ]; then
echo "* gpu: no *"
else
echo "* gpu: yes *"
fi
echo "* *"
echo "***********************************"
......@@ -7,6 +7,8 @@
#include <boost/fusion/include/vector.hpp>
#include <boost/fusion/container/vector/vector_fwd.hpp>
#include <boost/fusion/include/vector_fwd.hpp>
#include "boost/multi_array.hpp"
#include "base_type.hpp"
template<typename T> class Point_orig
{
......@@ -32,7 +34,7 @@ template<typename T> class Point
{
public:
typedef boost::fusion::vector<T,T,T,T,T[3],T[3][3]> type;
typedef boost::fusion::vector<T,T,T,T,boost::multi_array_ref<T,1>,boost::multi_array_ref<T,2> > type;
type data;
......
#ifndef BASE_TYPE_HPP
#define BASE_TYPE_HPP
#include <boost/multi_array.hpp>
template<class T>
struct base_type
{
typedef T type;
};
template<class T>
struct base_type<boost::multi_array_ref<T,1>>
{
typedef T type;
};
template<class T>
struct base_type<boost::multi_array_ref<T,2>>
{
typedef T type;
};
template<class T>
struct base_type<boost::multi_array_ref<T,3>>
{
typedef T type;
};
template<class T>
struct base_type<boost::multi_array_ref<T,4>>
{
typedef T type;
};
template<class T>
struct base_type<boost::multi_array_ref<T,5>>
{
typedef T type;
};
#endif
#include <iostream>
#include <boost/shared_array.hpp>
#include <vector>
#include <memory.hpp>
#include "memory_cpu.hpp"
#include "memory_gpu.hpp"
#include <initializer_list>
#include <array>
......
#include "map.hpp"
#include "memory.hpp"
#include "memory_cpu.hpp"
#include "memory_gpu.hpp"
#include "Particle.hpp"
#include <boost/mpl/int.hpp>
#include <typeinfo>
......@@ -8,8 +9,9 @@
#include <test_3.hpp>
#include <test_4.hpp>
#include <test_5.hpp>
#include "memory_gpu_thrust.hpp"
float Wi(float dist_x)
/*float Wi(float dist_x)
{
if (dist_x <= 1)
return 1.5f*dist_x * dist_x * dist_x - 2.5f * dist_x * dist_x + 1.0f;
......@@ -17,7 +19,7 @@ float Wi(float dist_x)
return -0.5f*dist_x * dist_x * dist_x * dist_x + 2.5f * dist_x * dist_x - 4.0f * dist_x + 2.0f;
else
return 0.0;
}
}*/
int main()
{
......@@ -29,23 +31,25 @@ int main()
sz.push_back(GS_SIZE);
sz.push_back(GS_SIZE);
layout_cpu< grid<Point<float>>, memory_cpu<memory_cpu_type<Point<float>>::type> > c3(sz);
layout_gpu< grid<Point<float>>, memory_gpu<memory_gpu_type<Point<float>>::type> > c3(sz);
// cpu test
test1();
// test1();
// layout_cpu< Particles<Point<float>, memory_cpu<float> >, particle_key > p1;
test2(c3);
// test2(c3);
test3(c3);
// test3(c3);
test4(c3);
// test4(c3);
// test5(c3);
memory_gpu_thrust<float> mgt;
test5(c3);
// k.set(2,2,2);
// c3.get(k).x = 6.0;
......
......@@ -9,7 +9,8 @@
#include "grid.hpp"
#include "Particle.hpp"
#include "Point.hpp"
#include "memory.hpp"
#include "memory_cpu.hpp"
#include "memory_gpu.hpp"
#ifndef MAP_HPP_
#define MAP_HPP_
......@@ -145,6 +146,12 @@ map_cpu<grid<grid<Point>>>
}
}*/
template<unsigned int p>
struct Point_type_prop
{
typedef typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type type;
};
template<typename T, typename Mem>
class layout_gpu<grid<Point<T>>,Mem>
{
......@@ -165,37 +172,38 @@ public:
boost::fusion::at_c<5>(data).allocate(g1.size());
}
template <unsigned int p>inline typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type & get(grid_key<p> & v1)
template <unsigned int p>inline typename mem_reference<typename Point_type_prop<p>::type>::type get(grid_key<p> & v1)
{
return boost::fusion::at_c<p>(data).get(g1.LinId(v1.getId()));
}
template <unsigned int p>inline typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type & get(grid_key_1<p> & v1)
template <unsigned int p>inline typename mem_reference<typename Point_type_prop<p>::type >::type get(grid_key_1<p> & v1)
{
return boost::fusion::at_c<p>(data).get(g1.LinId(v1.k[0]));
}
template <unsigned int p>inline typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type & get(grid_key_2<p> & v1)
template <unsigned int p>inline typename mem_reference<typename Point_type_prop<p>::type >::type get(grid_key_2<p> & v1)
{
return boost::fusion::at_c<p>(data).get(g1.LinId(v1.k[1],v1.k[0]));
}
template <unsigned int p>inline typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type & get(grid_key_3<p> & v1)
template <unsigned int p>inline typename mem_reference<typename Point_type_prop<p>::type >::type get(grid_key_3<p> & v1)
{
return boost::fusion::at_c<p>(data).get(g1.LinId(v1.k[2],v1.k[1],v1.k[0]));
}
template <unsigned int p>inline typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type & get(grid_key_4<p> & v1)
template <unsigned int p>inline typename mem_reference<typename Point_type_prop<p>::type >::type get(grid_key_4<p> & v1)
{
return boost::fusion::at_c<p>(data).get(g1.LinId(v1.k[3],v1.k[2],v1.k[1],v1.k[0]));
}
template <unsigned int p, unsigned int dim>inline typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type & get(grid_key_d<dim,p> & v1)
template <unsigned int p, unsigned int dim>inline typename mem_reference<typename Point_type_prop<p>::type >::type get(grid_key_d<dim,p> & v1)
{
return boost::fusion::at_c<p>(data).get(g1.LinId(v1));
}
template <unsigned int p, unsigned int dim>inline typename boost::fusion::result_of::at<Point<float>::type,boost::mpl::int_<p> >::type & get(grid_key_dx<dim> & v1)
template <unsigned int p, unsigned int dim>inline typename mem_reference<typename Point_type_prop<p>::type >::type get(grid_key_dx<dim> & v1)
{
return boost::fusion::at_c<p>(data).get(g1.LinId(v1));
}
......
#ifndef MEMORY_CPU_HPP_
#define MEMORY_CPU_HPP_
#include <boost/shared_array.hpp>
#include <boost/type_traits/remove_reference.hpp>
#include <boost/fusion/sequence/intrinsic/at_c.hpp>
#include <boost/fusion/include/at_c.hpp>
#include <boost/fusion/container/vector.hpp>
#include <boost/fusion/include/vector.hpp>
#include <boost/fusion/container/vector/vector_fwd.hpp>
#include <boost/fusion/include/vector_fwd.hpp>
#include <base_type.hpp>
typedef long int mem_id;
template<typename T>
class memory_cpu
{
boost::shared_array<T> mem;
size_t sz;
public:
T & get(mem_id id)
{
return mem.get()[id];
}
inline void set(mem_id id, T m)
{
mem.get()[id] = m;
}
inline bool allocate(size_t sz)
{
mem = boost::shared_array<T>(new T[sz]);
return true;
}
inline void destroy()
{
mem.reset();
}
inline bool copy(memory_cpu m)
{
memcpy(mem.get(),m.mem.get(),m.size());
return true;
}
inline size_t size()
{
return sz;
}
};
////////////////////////////// compile based specialization
#include "Point.hpp"
template<typename T>
struct memory_cpu_type
{
typedef typename T::type type;
};
template<typename T>
struct memory_cpu_type<Point<T>>
{
typedef typename Point<T>::type type;
};
////////////////// Compile based specialization
struct memory_null
{
};
#endif
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <algorithm>
#include <cstdlib>
#include "memory_cpu_thrust.hpp"
typedef long int mem_id;
#define STC_D(dv) (static_cast<thrust::device_vector<T> *>(dv))
#define STC_H(dv) (static_cast<thrust::host_vector<T> *>(dv))
template<class T> void memory_cpu_thrust<T>::load(memory_gpu_thrust<T> & d_vec)
{
// transfer data to the device
thrust::copy(STC_H(dv)->begin(), STC_H(dv)->end(), STC_D(d_vec.dv)->begin());
}
template<class T> void memory_cpu_thrust<T>::sort()
{
// sort data into device
thrust::sort(STC_H(dv)->begin(), STC_H(dv)->end());
}
template<class T> void memory_gpu_thrust<T>::resize(size_t sz)
{
// sort data into device
return STC_H(dv)->resize(sz);
}
template<class T> T memory_cpu_thrust<T>::reduce()
{
// sort data into device
return thrust::reduce(STC_H(dv)->begin(), STC_H(dv)->end());
}
template<class T> void * memory_cpu_thrust<T>::getThrustObj()
{
return static_cast<void *>(STC_H(dv));
}
// Compiled Forced specialization
template class memory_cpu_thrust<int>;
template class memory_cpu_thrust<float>;
template class memory_cpu_thrust<double>;
#ifndef MEMORY_CPU_THRUST_HPP
#define MEMORY_CPU_THRUST_HPP
#include <algorithm>
#include <cstdlib>
template<typename T>
class memory_cpu_thrust;
#include "memory_gpu_thrust.hpp"
typedef long int mem_id;
template<typename T>
class memory_cpu_thrust
{
void * dv;
public:
void load(memory_gpu_thrust<T> & d_vec);
void load(boost::shared_array<T> & mem);
void sort();
void resize(size_t sz);
T reduce();
void * getThrustObj();
// void foreach();
// void test();
};
#endif
#ifndef MEMORY_HPP_
#define MEMORY_HPP_
#ifndef MEMORY_GPU_HPP_
#define MEMORY_GPU_HPP_
#include "config.h"
#include <algorithm>
#include <cstdlib>
#include <boost/shared_array.hpp>
#include <boost/type_traits/remove_reference.hpp>
#include <boost/fusion/sequence/intrinsic/at_c.hpp>
#include <boost/fusion/include/at_c.hpp>
#include <boost/fusion/container/vector.hpp>
#include <boost/fusion/include/vector.hpp>
#include <boost/fusion/container/vector/vector_fwd.hpp>
#include <boost/fusion/include/vector_fwd.hpp>
#include "base_type.hpp"
typedef long int mem_id;
#ifndef NVCC
template<typename T>
class memory_cpu
struct mem_reference<T>
{
boost::shared_array<T> mem;
size_t sz;
public:
T & get(mem_id id)
{
return mem.get()[id];
}
inline void set(mem_id id, T m)
{
mem.get()[id] = m;
}
inline bool allocate(size_t sz)
{
mem = boost::shared_array<T>(new T[sz]);
return true;
}
inline void destroy()
{
mem.reset();
}
inline bool copy(memory_cpu m)
{
memcpy(mem.get(),m.mem.get(),m.size());
return true;
}
inline size_t size()
{
return sz;
}
typedef T& type;
};
template<typename T>
class memory_gpu
{
boost::shared_array<T> mem;
// Thrust memory management
size_t sz;
public:
......@@ -97,6 +62,158 @@ class memory_gpu
}
};
#endif
#ifdef NVCC
#include <memory_gpu_thrust.hpp>
#include <boost/type_traits/remove_reference.hpp>
template<typename T>
struct mem_reference
{
typedef typename boost::remove_reference<T>::type type;
};
template<> struct mem_reference<float&>
{
typedef float& type;
};
template<> struct mem_reference<int&>
{
typedef int& type;
};
template<> struct mem_reference<double&>
{
typedef double& type;
};
template<> struct mem_reference<float>
{
typedef float& type;
};
template<> struct mem_reference<int>
{
typedef int& type;
};
template<> struct mem_reference<double>
{
typedef double& type;
};
template<typename T>
class memory_gpu
{
memory_gpu_thrust< typename base_type<T>::type > mem;
T * mem_ptr;
size_t sz;
public:
typedef T type;
typename mem_reference<T>::type get(mem_id id)
{
return mem_ptr[id];
}
inline void set(mem_id id, T m)
{
mem_ptr[id] = m;
}
inline bool allocate(size_t sz)
{
mem.resize(sz);
return true;
}
inline void destroy()
{
}
inline bool copy(memory_gpu m)
{
memory_gpu_thrust<T>::copy(mem.dv,m.mem.dv);
return true;
}
inline size_t size()
{
return mem->dv.size();
}
};
template<typename T, unsigned int i, unsigned int j>
class memory_gpu_array2D
{
memory_gpu_thrust<T> mem;
boost::multi_array_ref<T,3> ma;
size_t sz;
public:
typedef T type;
typename boost::multi_array_ref<T,2> get(mem_id id)
{
return ma[id];
}
inline void set(mem_id id, T m)
{
mem.dv[id] = m;
}
inline bool allocate(size_t sz)
{
mem.dv.resize(sz);
// shape on ma
typename boost::multi_array_ref<T,3>::size_type xDim(i), yDim(j), zDim(sz);
typename boost::multi_array<T,3>::size_type ordering[] = {2,1,0};
ma = boost::multi_array_ref<T,3>(mem.getPointer(),boost::extents[xDim][yDim][zDim],ordering);
return true;
}
inline void destroy()
{
}
inline bool copy(memory_gpu_array2D m)
{
memory_gpu_thrust<T>::copy(mem.dv,m.mem.dv);
return true;
}