First commit of CUDA FVM Vlasov solver.

Makefile

@@ -0,0 +1,107 @@
+include Makefile.inc
+
+#LIB = -L/usr/local/cuda/lib64 -lcudart -lsilo
+
+#INC_CUDA = -I/usr/local/cuda/include
+
+# Define dependencies of each object file
+DEPS_COMMON = common.h definitions.h logger.h
+DEPS_CELL_SPATIAL = cell_spatial.h cell_spatial.cpp
+DEPS_CELLSYNC = cell_spatial.h cellsync.cpp
+DEPS_CUDA_ACC = cuda_common.cu cudalaunch.h devicegrid.h grid.h parameters.h project.cu cuda_acc.cu
+DEPS_CUDA_TRANS = cuda_common.cu cudalaunch.h devicegrid.h grid.h parameters.h project.cu cuda_trans.cu
+DEPS_CUDAFUNCS = cudafuncs.h cudafuncs.cpp
+DEPS_GPU_DEVICE_GRID = cell_spatial.h parameters.h devicegrid.h gpudevicegrid.cpp
+DEPS_GRID = cell_spatial.h cudafuncs.h parameters.h grid.h grid.cpp
+DEPS_GRIDBUILDER = cell_spatial.h devicegrid.h grid.h parameters.h project.h gridbuilder.cpp
+DEPS_LOGGER = logger.h logger.cpp
+DEPS_MAIN = cudafuncs.h parameters.h grid.h devicegrid.h writevars.h main.cpp
+DEPS_PARAMETERS = parameters.h parameters.cpp
+DEPS_PROJECT = project.h project.cpp
+DEPS_SILOWRITER = cell_spatial.h silowriter.h silowriter.cpp
+DEPS_WRITEVARS = cudafuncs.h grid.h devicegrid.h silowriter.h writevars.h writevars.cpp
+
+DEPS_CELL_SPATIAL += $(DEPS_COMMON)
+DEPS_CELLSYNC += $(DEPS_COMMON)
+DEPS_CUDA_ACC += $(DEPS_COMMON)
+DEPS_CUDA_TRANS += $(DEPS_COMMON)
+DEPS_CUDAFUNCS += $(DEPS_COMMON)
+DEPS_GPU_DEVICE_GRID += $(DEPS_COMMON)
+DEPS_GRID += $(DEPS_COMMON)
+DEPS_GRIDBUILDER += $(DEPS_COMMON)
+DEPS_LOGGER += $(DEPS_COMMON)
+DEPS_MAIN += $(DEPS_COMMON)
+DEPS_PARAMETERS += $(DEPS_COMMON)
+DEPS_PROJECT += $(DEPS_COMMON)
+DEPS_SILOWRITER += $(DEPS_COMMON)
+DEPS_WRITEVARS += ${DEPS_COMMON}
+
+HDRS = cell_spatial.h common.h cudafuncs.h cudalaunch.h definitions.h devicegrid.h grid.h \
+	 logger.h parameters.h project.h silowriter.h writevars.h
+
+SRC = cell_spatial.cpp cellsync.cpp cuda_acc.cu cuda_common.cu cuda_trans.cu \
+	cudafuncs.cpp gpudevicegrid.cpp grid.cpp gridbuilder.cpp \
+	logger.cpp main.cpp parameters.cpp silowriter.cpp writevars.cpp
+
+OBJS = cell_spatial.o cellsync.o cuda_acc.o cuda_trans.o cudafuncs.o grid.o\
+	gridbuilder.o logger.o main.o parameters.o gpudevicegrid.o project.o\
+	silowriter.o writevars.o
+
+clean:
+	make clean -C projects
+	rm -rf *.o *.ptx *.tar* *.txt *.silo project.cu project.cpp main *~
+
+# Rules for making each object file needed by the executable
+cell_spatial.o: $(DEPS_CELL_SPATIAL)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c cell_spatial.cpp
+
+cellsync.o: $(DEPS_CELLSYNC)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c cellsync.cpp $(INC_CUDA)
+
+cuda_acc.o: $(DEPS_CUDA_ACC)
+	$(NVCC) $(NVCCFLAGS) $(FLAGS) -c cuda_acc.cu
+
+cuda_trans.o: $(DEPS_CUDA_TRANS)
+	$(NVCC) $(NVCCFLAGS) $(FLAGS) -c cuda_trans.cu
+
+cudafuncs.o: $(DEPS_CUDAFUNCS)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c cudafuncs.cpp $(INC_CUDA)
+
+gpudevicegrid.o: $(DEPS_GPU_DEVICE_GRID)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c gpudevicegrid.cpp $(INC_CUDA)
+
+grid.o: $(DEPS_GRID)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c grid.cpp
+
+gridbuilder.o: $(DEPS_GRIDBUILDER)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c gridbuilder.cpp
+
+logger.o: $(DEPS_LOGGER)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c logger.cpp
+
+main.o: $(DEPS_MAIN)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c main.cpp
+
+parameters.o: $(DEPS_PARAMETERS)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c parameters.cpp
+
+project.o: $(DEPS_PROJECT)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c project.cpp
+
+projinstall:
+	make project -C projects
+
+silowriter.o: $(DEPS_SILOWRITER)
+	$(CMP) $(CXXFLAGS) $(FLAGS) -c silowriter.cpp ${INC_SILO}
+
+writevars.o: ${DEPS_WRITEVARS}
+	${CMP} ${CXXFLAGS} ${FLAGS} -c writevars.cpp ${INC_SILO}
+
+# Make a tar file containing the source code
+dist:
+	tar -cf cudaFVM.tar $(HDRS) $(SRC) Doxyfile Makefile Makefile.inc projects
+	gzip cudaFVM.tar
+
+# Make executable
+main: projinstall $(OBJS)
+	$(LINK) -o main $(OBJS) $(LIB)

Makefile.inc

@@ -0,0 +1,42 @@
+# Compile directory: 
+INSTALL = /home/sandroos/codes/cuda/FVM2
+
+# Path to CUDA programming examples:
+CUDASDK = /home/sandroos/NVIDIA_GPU_Computing_SDK
+# Path to CUDA:
+CUDAPTH = /usr/local/cuda
+# Path to SILO:
+SILOPTH = 
+
+# Set the device:
+DEVICE = DEVICE_DEFAULT
+# Set the simulated project:
+PROJ = harm1D
+
+# C++ compiler:
+CMP = /usr/bin/g++-4.3
+# CUDA compiler:
+NVCC = nvcc
+# Linker:
+LINK = /usr/bin/g++-4.3
+# Flags for C++ and CUDA compilers:
+FLAGS = -O3
+
+# -------------------------------------------
+# Automatically set parameters, DO NOT CHANGE
+# -------------------------------------------
+
+INC_CUDA = -I${CUDAPTH}/include
+INC_SDK = -I${CUDASDK}/C/common/inc
+INC_SILO = -I${SILOPTH}/include
+
+LIB_CUDA = -L${CUDAPTH}/lib64 -lcudart
+LIB_CUTIL = -L${CUDASDK}/C/lib -lcutil_x86_64
+LIB_SILO = -L${SILOPTH}/lib -lsilo
+
+LIB = ${LIB_CUDA}
+LIB += ${LIB_SILO}
+
+CXXFLAGS =
+NVCCFLAGS = --ptxas-options=-v -ccbin $(CMP) -D${DEVICE}
+

cell_spatial.cpp

@@ -0,0 +1,73 @@
+#include <cstdlib>
+#include <iostream>
+#include <cmath>
+#include <limits>
+
+#include "cell_spatial.h"
+
+using namespace std;
+
+SpatialCell::SpatialCell() {
+   gpuIndex = numeric_limits<uint>::max();
+
+   cellType = Cell::UNINITIALIZED;
+   cellIndex = numeric_limits<uint>::max();
+   velBlockIndex = numeric_limits<uint>::max();
+   i_ind = numeric_limits<uint>::max();
+   j_ind = numeric_limits<uint>::max();
+   k_ind = numeric_limits<uint>::max();
+
+   cpu_nbrsSpa = NULL;
+   cpu_cellParams = NULL;
+
+   cpu_nbrsVel = NULL;
+   cpu_avgs = NULL;
+   cpu_blockParams = NULL;
+
+   gpu_cellParams = NULL;
+
+   gpu_nbrsVel = NULL;
+   gpu_avgs = NULL;
+   gpu_blockParams = NULL;
+}
+
+SpatialCell::SpatialCell(const SpatialCell& s) {
+   gpuIndex = s.gpuIndex;
+
+   cellType = s.cellType;
+   cellIndex = s.cellIndex;
+   velBlockIndex = s.velBlockIndex;
+   i_ind = s.i_ind;
+   j_ind = s.j_ind;
+   k_ind = s.k_ind;
+
+   cpu_nbrsSpa = s.cpu_nbrsSpa;
+   cpu_cellParams = s.cpu_cellParams;
+
+   cpu_nbrsVel = s.cpu_nbrsVel;
+   cpu_avgs = s.cpu_avgs;
+   cpu_blockParams = s.cpu_blockParams;
+
+   gpu_cellParams = s.gpu_cellParams;
+
+   gpu_nbrsVel = s.gpu_nbrsVel;
+   gpu_avgs = s.gpu_avgs;
+   gpu_blockParams = s.gpu_blockParams;
+}
+
+SpatialCell::~SpatialCell() {
+   cpu_nbrsSpa = NULL;
+   cpu_cellParams = NULL;
+
+   cpu_nbrsVel = NULL;
+   cpu_avgs = NULL;
+   cpu_blockParams = NULL;
+
+   gpu_cellParams = NULL;
+
+   gpu_nbrsVel = NULL;
+   gpu_avgs = NULL;
+   gpu_blockParams = NULL;
+}
+
+

cell_spatial.h

@@ -0,0 +1,49 @@
+#ifndef CELL_SPATIAL
+#define CELL_SPATIAL
+
+#include <utility>
+#include <vector>
+
+#include "definitions.h"
+
+namespace Cell {
+   enum CellType {INNER,BOUNDARY,GHOST,UNINITIALIZED};
+   enum Array {Blocks,BlockParams,NbrsVel,CellParams,Fx,Fy,Fz};
+   enum Dir {CpuToDev,DevToCpu};
+};
+
+struct SpatialCell {
+   uint gpuIndex;
+   uint N_blocks;
+
+   Cell::CellType cellType; 
+   uint cellIndex; 
+   uint velBlockIndex; 
+   uint i_ind; /**< The index of the spatial cell in x-coordinate.*/
+   uint j_ind; /**< The index of the spatial cell in y-coordinate.*/
+   uint k_ind; /**< The index of the spatial cell in z-coordinate.*/
+
+   // Pointers to arrays containing spatial cell parameters in CPU memory
+   real* cpu_cellParams;  /**< Pointer to physical cell parameters in cpu memory.*/
+   uint* cpu_nbrsSpa;     /**< Pointer to spatial  neighbout list in CPU memory.*/
+
+   // Pointers to arrays containing velocity block parameters in CPU memory
+   uint* cpu_nbrsVel;     /**< Pointer to velocity neighbour list in CPU memory.*/
+   real* cpu_avgs;        /**< Pointer to velocity block array in CPU memory.*/
+   real* cpu_blockParams; /**< Pointer to velocity block parameter array in CPU memory.*/
+
+   // Pointers to arrays containing velocity block parameters in device memory
+   real* gpu_cellParams;  /**< Pointer to spatial cell parameter array in GPU memory.*/
+   real* gpu_avgs;        /**< Pointer to velocity block array in GPU memory.*/
+   real* gpu_blockParams; /**< Pointer to velocity block parameter array in GPU memory.*/
+   uint* gpu_nbrsVel;     /**< Pointer to velocity neighbout list in GPU memory.*/
+
+   SpatialCell();
+   SpatialCell(const SpatialCell& s);
+   ~SpatialCell();
+
+   bool devSync(const Cell::Array& array,const Cell::Dir direction);
+};
+
+#endif
+

cellsync.cpp

@@ -0,0 +1,68 @@
+#include <cstdlib>
+#include <iostream>
+#include <cuda_runtime_api.h>
+
+#include "common.h"
+#include "cell_spatial.h"
+#include "logger.h"
+
+using namespace std;
+
+extern Logger logger;
+
+bool SpatialCell::devSync(const Cell::Array& array,const Cell::Dir direction) {
+   // Set pointers to cpu and gpu arrays which are to be synced, 
+   // and calculate the amount of bytes to transfer:
+   void* cpuArray;
+   void* gpuArray;
+   size_t bytes;
+   switch (array) {
+    case Cell::Blocks:
+      cpuArray = cpu_avgs;
+      gpuArray = gpu_avgs;
+      bytes = N_blocks*SIZE_VELBLOCK*sizeof(real);
+      break;
+    case Cell::BlockParams:
+      cpuArray = cpu_blockParams;
+      gpuArray = gpu_blockParams;
+      bytes = N_blocks*SIZE_BLOCKPARAMS*sizeof(real);
+      break;
+    case Cell::NbrsVel:
+      cpuArray = cpu_nbrsVel;
+      gpuArray = gpu_nbrsVel;
+      bytes = N_blocks*SIZE_NBRS_VEL*sizeof(uint);
+      break;
+    case Cell::CellParams:
+      cpuArray = cpu_cellParams;
+      gpuArray = gpu_cellParams;
+      bytes = SIZE_CELLPARAMS*sizeof(real);
+      break;
+   }
+   // Set destination and source arrays:
+   void* src;
+   void* dst;
+   cudaMemcpyKind copyDirection;
+   switch (direction) {
+    case Cell::CpuToDev:
+      src = cpuArray;
+      dst = gpuArray;
+      copyDirection = cudaMemcpyHostToDevice;
+      break;
+    case Cell::DevToCpu:
+      src = gpuArray;
+      dst = cpuArray;
+      copyDirection = cudaMemcpyDeviceToHost;
+      break;
+   }
+   // Perform copy and return result:
+   //cerr << "\t Copying from " << src << " to " << dst << endl;
+   cudaError_t result = cudaMemcpy(dst,src,bytes,copyDirection);   
+   if (result == cudaSuccess) return true;
+
+   logger << "(CELLSYNC): Error occurred while copying data, error = '" << cudaGetErrorString(result) << "'" << endl;
+   return false;
+}
+
+
+
+