GopherProxy

	tset ghost nodes for div(tau) - sphere - GPU-based 3D discrete element method a…
	git clone git://src.adamsgaard.dk/sphere
	Log
	Files
	Refs
	LICENSE
	---
	commit 35a9742af2314b68fe8eb51056d0d996870d0bbd
	parent 9216f53f507b61c0f492f9903acfe0ab70058d37
	Author: Anders Damsgaard <[email protected]>
	Date: Thu, 5 Jun 2014 12:51:50 +0200

	set ghost nodes for div(tau)

	Diffstat:
	M src/debug.h \| 4 ++--
	M src/device.cu \| 22 ++++++++++++++++++++++
	M src/navierstokes.cuh \| 4 ++++

	3 files changed, 28 insertions(+), 2 deletions(-)
	---
	diff --git a/src/debug.h b/src/debug.h
	t@@ -40,10 +40,10 @@ const int conv_log_interval = 4;
	#define CHECK_NS_FINITE

	// Enable reporting of velocity prediction components to stdout
	-//#define REPORT_V_P_COMPONENTS
	+#define REPORT_V_P_COMPONENTS

	// Enable reporting of forcing function terms to stdout
	-//#define REPORT_FORCING_TERMS
	+#define REPORT_FORCING_TERMS

	// Choose solver model (see Zhou et al. 2010 "Discrete particle simulation of
	// particle-fluid flow: model formulations and their applicability", table. 1.
	diff --git a/src/device.cu b/src/device.cu
	t@@ -915,6 +915,23 @@ __host__ void DEM::startTime()
	cudaThreadSynchronize();
	checkForCudaErrorsIter("Post findFaceDivTau", iter);

	+ /*setUpperDivTau<<<dimGridFluidFace, dimBlockFluid>>>(
	+ dev_ns_p, ns.bc_bot, ns.bc_top);
	+ cudaThreadSynchronize();
	+ checkForCudaErrorsIter("Post setNSghostNodes(dev_ns_p)", iter)
	+ */
	+
	+ setNSghostNodesFace<Float>
	+ <<<dimGridFluidFace, dimBlockFluid>>>(
	+ dev_ns_div_tau_x,
	+ dev_ns_div_tau_y,
	+ dev_ns_div_tau_z,
	+ ns.bc_bot,
	+ ns.bc_top);
	+ cudaThreadSynchronize();
	+ checkForCudaErrorsIter("Post setNSghostNodes(dev_ns_div_tau)",
	+ iter);
	+
	setNSghostNodes<Float><<<dimGridFluid, dimBlockFluid>>>(
	dev_ns_p, ns.bc_bot, ns.bc_top);
	cudaThreadSynchronize();
	t@@ -1514,6 +1531,11 @@ __host__ void DEM::startTime()
	filetimeclock += time.dt;
	++iter;

	+ // Make sure all preceding tasks are complete
	+ if (cudaDeviceSynchronize() != cudaSuccess) {
	+ cerr << "Error during cudaDeviceSynchronize()" << endl;
	+ }
	+
	// Report time to console
	if (verbose == 1 && (iter % stdout_report == 0)) {

	diff --git a/src/navierstokes.cuh b/src/navierstokes.cuh
	t@@ -13,6 +13,9 @@
	#include "constants.cuh"
	#include "debug.h"

	+#define REPORT_V_P_COMPONENTS
	+#define REPORT_FORCING_TERMS
	+
	// Arithmetic mean of two numbers
	__inline__ __device__ Float amean(Float a, Float b) {
	return (a+b)*0.5;
	t@@ -2247,6 +2250,7 @@ __global__ void findPredNSvelocities(
	+ porosity_term
	+ advection_term;

	+ printf("\n\nhello, world\n\n");
	#ifdef REPORT_V_P_COMPONENTS
	// Report velocity components to stdout for debugging
	printf("[%d,%d,%d]\t"