Working on KickDriftKernel #93

luchete80 · Jun 10, 2022 · 48dc553 · 48dc553
1 parent f702a8f
commit 48dc553
Show file tree

Hide file tree

Showing 3 changed files with 23 additions and 12 deletions.
diff --git a/src/cuda/Domain_d.cuh b/src/cuda/Domain_d.cuh
@@ -290,7 +290,8 @@ __device__                      inline void CalcDensInc(
 __device__ inline void UpdateDensity(const double &dt);
 
   __device__ inline void UpdateVel(const double &dt);
-
+  __device__ inline void UpdatePos(const double &dt);
+
 	__device__ /*__forceinline__*/inline void CalcRateTensors(const uint *particlenbcount,
 																													const uint *neighborWriteOffsets,
 																													const uint *neighbors);

diff --git a/src/cuda/KickDriftSolver.cu b/src/cuda/KickDriftSolver.cu
@@ -193,11 +193,8 @@ void Domain_d::MechKickDriftSolve(const double &tf, const double &dt_out){
       CudaHelper::GetPointer(nsearch.deviceData->d_Neighbors)		
 		);
     cudaDeviceSynchronize(); //REQUIRED!!!!
-    UpdateDensityKernel<<<blocksPerGrid,threadsPerBlock >>>(this,
-      CudaHelper::GetPointer(nsearch.deviceData->d_NeighborCounts),
-      CudaHelper::GetPointer(nsearch.deviceData->d_NeighborWriteOffsets),
-      CudaHelper::GetPointer(nsearch.deviceData->d_Neighbors)		
-		);
+
+    UpdateDensityKernel<<<blocksPerGrid,threadsPerBlock >>>(this,deltat);		
     cudaDeviceSynchronize(); //REQUIRED!!!!
 
     if (contact){
@@ -225,12 +222,13 @@ void Domain_d::MechKickDriftSolve(const double &tf, const double &dt_out){
 		}
 
 		//Move particle and then calculate streses and strains ()
-		MoveKernelExt<<<blocksPerGrid,threadsPerBlock >>> (v, va,vb,
-														rho, rhoa, rhob, drho,
-														x, a,
-														u, /*Mat3_t I, */deltat,
-														isfirst_step, particle_count);	
-		cudaDeviceSynchronize(); //REQUIRED!!!!
+		// MoveKernelExt<<<blocksPerGrid,threadsPerBlock >>> (v, va,vb,
+														// rho, rhoa, rhob, drho,
+														// x, a,
+														// u, /*Mat3_t I, */deltat,
+														// isfirst_step, particle_count);	
+		UpdatePosKernel<<<blocksPerGrid,threadsPerBlock >>>(this,deltat);
+    cudaDeviceSynchronize(); //REQUIRED!!!!
 
 		//If kernel is the external, calculate pressure
 		//Calculate pressure!

diff --git a/src/cuda/Mechanical.cu b/src/cuda/Mechanical.cu
@@ -150,6 +150,18 @@ __global__ void UpdateVelKernel(Domain_d *dom, const double &dt) {
   dom->UpdateVel(dt);
 }
 
+__device__ inline void Domain_d::UpdatePos(const double &dt){
+	int i = threadIdx.x + blockDim.x*blockIdx.x;
+
+	if ( i < particle_count ) {
+    	x[i] += dt*v[i];   
+  }
+}
+
+__global__ void UpdatePosKernel(Domain_d *dom, const double &dt) {
+  dom->UpdatePos(dt);
+}
+
 __global__ void UpdateDensityKernel(Domain_d *dom, const double &dt) {
   dom->UpdateDensity(dt);
 }