GopherProxy

	tBugfixes - sphere - GPU-based 3D discrete element method algorithm with option…
	git clone git://src.adamsgaard.dk/sphere
	Log
	Files
	Refs
	LICENSE
	---
	commit ff0ef120a58910bd4dd4e366c584a7d4b451c9a7
	parent 5bfd34d61e71717237371f1b319a5ca90987abaa
	Author: Anders Damsgaard Christensen <[email protected]>
	Date: Fri, 26 Oct 2012 23:13:27 +0200

	Bugfixes

	Diffstat:
	M src/device.cu \| 21 +++++++++++----------
	M src/sphere.h \| 8 ++++----

	2 files changed, 15 insertions(+), 14 deletions(-)
	---
	diff --git a/src/device.cu b/src/device.cu
	t@@ -11,6 +11,7 @@
	#include "datatypes.h"
	#include "utility.cuh"
	#include "constants.cuh"
	+#include "debug.h"

	#include "sorting.cuh"
	#include "contactmodels.cuh"
	t@@ -337,10 +338,10 @@ __host__ void DEM::transferToGlobalDeviceMemory()
	cout << " Transfering data to the device: ";

	// Copy structure data from host to global device memory
	- cudaMemcpy(k, dev_k, sizeof(k), cudaMemcpyDeviceToHost);
	- cudaMemcpy(e, dev_e, sizeof(e), cudaMemcpyDeviceToHost);
	- cudaMemcpy(time, dev_time, sizeof(time), cudaMemcpyDeviceToHost);
	- cudaMemcpy(walls, dev_walls, sizeof(walls), cudaMemcpyDeviceToHost);
	+ cudaMemcpy(&k, dev_k, sizeof(k), cudaMemcpyDeviceToHost);
	+ cudaMemcpy(&e, dev_e, sizeof(e), cudaMemcpyDeviceToHost);
	+ cudaMemcpy(&time, dev_time, sizeof(time), cudaMemcpyDeviceToHost);
	+ cudaMemcpy(&walls, dev_walls, sizeof(walls), cudaMemcpyDeviceToHost);

	checkForCudaErrors("End of transferFromGlobalDeviceMemory");
	if (verbose == 1)
	t@@ -361,8 +362,8 @@ __host__ void DEM::startTime()


	// Particle memory size
	- unsigned int memSizeF = sizeof(Float) * np;
	- unsigned int memSizeF4 = sizeof(Float4) * np;
	+ //unsigned int memSizeF = sizeof(Float) * np;
	+ //unsigned int memSizeF4 = sizeof(Float4) * np;

	// Allocate device memory for particle variables,
	// tied to previously declared pointers in structures
	t@@ -384,7 +385,7 @@ __host__ void DEM::startTime()
	// GPU workload configuration
	unsigned int threadsPerBlock = 256;
	// Create enough blocks to accomodate the particles
	- unsigned int blocksPerGrid = iDivUp(p.np, threadsPerBlock);
	+ unsigned int blocksPerGrid = iDivUp(np, threadsPerBlock);
	dim3 dimGrid(blocksPerGrid, 1, 1); // Blocks arranged in 1D grid
	dim3 dimBlock(threadsPerBlock, 1, 1); // Threads arranged in 1D block
	// Shared memory per block
	t@@ -749,8 +750,8 @@ __host__ void DEM::startTime()
	freeGlobalDeviceMemory();

	// Contact info arrays
	- delete[] host_contacts;
	- delete[] host_distmod;
	- delete[] host_delta_t;
	+ delete[] k.contacts;
	+ delete[] k.distmod;
	+ delete[] k.delta_t;

	} /* EOF */
	diff --git a/src/sphere.h b/src/sphere.h
	t@@ -21,11 +21,11 @@ class DEM {

	// Structure containing individual particle kinematics
	Kinematics k; // host
	- Kinematics dev_k; // device
	+ Kinematics *dev_k; // device

	// Structure containing energy values
	Energies e; // host
	- Energies dev_e; // device
	+ Energies *dev_e; // device

	// Structure of global parameters
	Params params; // host
	t@@ -34,7 +34,7 @@ class DEM {
	Grid grid; // host

	// Structure containing sorting arrays
	- Sorting dev_sort; // device
	+ Sorting *dev_sort; // device

	// Structure of temporal parameters
	Time time; // host
	t@@ -42,7 +42,7 @@ class DEM {

	// Structure of wall parameters
	Walls walls; // host
	- Walls dev_walls; // device
	+ Walls *dev_walls; // device

	// GPU initialization, must be called before startTime()
	void initializeGPU(void);