Workaround to free device memory

* ConstHybridVector provices the method free to manually free device memory
* A better solution would be a reference counting implementation,
  but there is no out of the box solution like boost::shared_ptr
* Leaves ComputationalRadiationPhysics#62 still open (smart pointers necessary)
* closes: ComputationalRadiationPhysics#18

include/ConstHybridVector.hpp

@@ -0,0 +1,62 @@
+#pragma once
+
+// STL
+#include <vector> /* std::vector */
+
+#include <cuda_utils.hpp> /* copyToDevice */
+
+
+/**
+ * @brief Vector on host and array on device
+ * with transparent access. Access is
+ * triggered based on a compiler macro.
+ *
+ **/
+template<class T>
+class ConstHybridVector {
+public:
+
+ ConstHybridVector(std::vector<T> &srcV) :
+ hostV(srcV),
+ deviceV(copyToDevice(srcV)){
+
+ }
+
+ __forceinline__ __host__ __device__ T at(int i) const{
+#ifdef __CUDA_ARCH__
+ return deviceV[i];
+#else
+ return hostV.at(i);
+#endif
+ }
+
+ __forceinline__ __host__ __device__ const T operator[] (int i) const {
+#ifdef __CUDA_ARCH__
+ return deviceV[i];
+#else
+ return hostV[i];
+#endif
+ }
+
+ __host__ operator T*(){
+ return &(hostV.at(0));
+ }
+
+ __host__ T* toArray() const{
+ std::vector<T> copyV = hostV;
+ return &(copyV.at(0));
+ }
+
+ __host__ std::vector<T> toVector() const{
+ return hostV;
+ }
+
+ __host__ void free() {
+ cudaFree(deviceV);
+ }
+
+private:
+ T *deviceV;
+ std::vector<T> hostV;
+
+};

include/calc_phi_ase_graybat.hpp

@@ -45,5 +45,5 @@
 
 float calcPhiAseGrayBat ( const ExperimentParameters &experiment,
  const ComputeParameters &compute,
- const Mesh& mesh,
+ Mesh& mesh,
  Result &result );

include/calc_phi_ase_mpi.hpp

@@ -44,6 +44,6 @@
  */
 float calcPhiAseMPI (const ExperimentParameters &experiment,
  const ComputeParameters &compute,
- const Mesh& mesh,
+ Mesh& mesh,
  Result &result);
 

include/cuda_utils.hpp

@@ -88,67 +88,7 @@ T copyFromDevice(const T* deviceV){
  return a;
 }
 
-/**
- * @brief Vector on host and array on device
- * with transparent access. Access is
- * triggered based on a compiler macro.
- *
- **/
-template<class T>
-class constHybridVector {
-public:
-
- constHybridVector(std::vector<T> &srcV) :
- hostV(srcV),
- deviceV(copyToDevice(srcV)){
-
- }
-
- constHybridVector(){
-
- }
-
- __forceinline__ __host__ __device__ T at(int i) const{
-#ifdef __CUDA_ARCH__
- return deviceV[i];
-#else
- return hostV.at(i);
-#endif
- }
-
- __forceinline__ __host__ __device__ const T operator[] (int i) const {
-#ifdef __CUDA_ARCH__
- return deviceV[i];
-#else
- return hostV[i];
-#endif
- }
-
- // Assignment operator not needed because vector should be constant
- /* __host__ constHybridVector& operator= (const std::vector<T> &otherV){ */
- /* deviceV = copyToDevice(otherV); */
- /* hostV = otherV; */
- /* return *this; */
- /* } */
-
- __host__ operator T*(){
- return &(hostV.at(0));
- }
-
- __host__ T* toArray() const{
- std::vector<T> copyV = hostV;
- return &(copyV.at(0));
- }
-
- __host__ std::vector<T> toVector() const{
- return hostV;
- }
-
-private:
- T *deviceV;
- std::vector<T> hostV;
 
-};
 
 
 /** 

include/mesh.hpp

@@ -32,8 +32,8 @@
 #include <string>
 #include <curand_kernel.h> /* curand_uniform */
 
-#include <cuda_utils.hpp>
 #include <geometry.hpp>
+#include <ConstHybridVector.hpp>
 
 #define REFLECTION_SMALL 1E-3
 #define SMALL 1E-5
@@ -109,9 +109,6 @@
 class Mesh {
  public:
 
- Mesh(){}
-
- //Mesh();
  Mesh(// Constants
  double claddingAbsorption,
  float surfaceTotal,
@@ -141,23 +138,23 @@ class Mesh {
  std::vector<unsigned> triangleNormalPoint);
 
 
- constHybridVector<double> points;
- constHybridVector<double> betaVolume;
- constHybridVector<double> normalVec;
- constHybridVector<double> centers;
- constHybridVector<float> triangleSurfaces;
- constHybridVector<int> forbiddenEdge;
- constHybridVector<double> betaCells;
- constHybridVector<unsigned> claddingCellTypes;
+ ConstHybridVector<double> points;
+ ConstHybridVector<double> betaVolume;
+ ConstHybridVector<double> normalVec;
+ ConstHybridVector<double> centers;
+ ConstHybridVector<float> triangleSurfaces;
+ ConstHybridVector<int> forbiddenEdge;
+ ConstHybridVector<double> betaCells;
+ ConstHybridVector<unsigned> claddingCellTypes;
 
- constHybridVector<float> refractiveIndices; 
- constHybridVector<float> reflectivities; //based on triangleIndex, with offset from bottom/top
- constHybridVector<float> totalReflectionAngles;
+ ConstHybridVector<float> refractiveIndices; 
+ ConstHybridVector<float> reflectivities; //based on triangleIndex, with offset from bottom/top
+ ConstHybridVector<float> totalReflectionAngles;
 
  // Indexstructs
- constHybridVector<unsigned> trianglePointIndices;
- constHybridVector<int> triangleNeighbors;
- constHybridVector<unsigned> triangleNormalPoint;
+ ConstHybridVector<unsigned> trianglePointIndices;
+ ConstHybridVector<int> triangleNeighbors;
+ ConstHybridVector<unsigned> triangleNormalPoint;
 
  // Constants
  double claddingAbsorption;
@@ -174,6 +171,8 @@ class Mesh {
 
  ~Mesh();
 
+ void free();
+
  __device__ int getNeighbor(unsigned triangle, int edge) const;
  __device__ Point genRndPoint(unsigned triangle, unsigned level, curandStateMtgp32 *globalState) const;
  __device__ double getBetaVolume(unsigned triangle, unsigned level) const;

src/calc_phi_ase_graybat.cc

@@ -157,7 +157,7 @@ void processSamples(const Vertex slave,
 
 float calcPhiAseGrayBat ( const ExperimentParameters &experiment,
  const ComputeParameters &compute,
- const Mesh& mesh,
+ Mesh& mesh,
  Result &result ){
 
  /***************************************************************************
@@ -202,6 +202,7 @@ float calcPhiAseGrayBat ( const ExperimentParameters &experiment,
  if(vertex != master){
  processSamples(vertex, master, cage, experiment, compute, mesh, result);
  cage.~Cage();
+ mesh.free();
  exit(0);
 
  } 

src/calc_phi_ase_mpi.cc

@@ -172,7 +172,7 @@ void mpiCompute( const ExperimentParameters &experiment,
 
 float calcPhiAseMPI ( const ExperimentParameters &experiment,
  const ComputeParameters &compute,
- const Mesh& mesh,
+ Mesh& mesh,
  Result &result ){
 
  // Init MPI
@@ -203,6 +203,7 @@ float calcPhiAseMPI ( const ExperimentParameters &experiment,
 
  cudaDeviceReset(); 
  MPI_Finalize();
+ mesh.free();
  break;
 
 

src/main.cc

@@ -289,7 +289,13 @@ int main(int argc, char **argv){
  dout(V_STAT) << std::endl;
 
  }
-
+
+ // Cleanup device memory
+ // TODO: replace by smart pointer for device memory
+ for(Mesh &mesh : meshs){
+ mesh.free();
+ }
+
  return 0;
 
 }

src/mesh.cu

@@ -107,7 +107,27 @@ Mesh::Mesh(// Constants
  triangleNormalPoint(triangleNormalPoint){
 
 }
+
 Mesh::~Mesh() {
+
+}
+
+__host__ void Mesh::free(){
+ points.free();
+ betaVolume.free();
+ normalVec.free();
+ centers.free();
+ triangleSurfaces.free();
+ forbiddenEdge.free();
+ betaCells.free();
+ claddingCellTypes.free();
+ refractiveIndices.free(); 
+ reflectivities.free(); 
+ totalReflectionAngles.free();
+ trianglePointIndices.free();
+ triangleNeighbors.free();
+ triangleNormalPoint.free();
+
 }
 
 /**