Skip to content

Commit

Permalink
Merge branch 'master' of github.com:mosaic-group/openfpm
Browse files Browse the repository at this point in the history
  • Loading branch information
Serhii Yaskovets committed Jun 26, 2024
2 parents 53fdca7 + 859f474 commit 8b3531c
Show file tree
Hide file tree
Showing 5 changed files with 17 additions and 20 deletions.
5 changes: 2 additions & 3 deletions example/Vector/7_SPH_dlb/main.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -303,7 +303,7 @@ inline double Wab(double r)
if (r < 1.0)
return (1.0 - 3.0/2.0*r*r + 3.0/4.0*r*r*r)*a2;
else if (r < 2.0)
return (1.0/4.0*(2.0 - r*r)*(2.0 - r*r)*(2.0 - r*r))*a2;
return (1.0/4.0*(2.0 - r)*(2.0 - r)*(2.0 - r))*a2;
else
return 0.0;
}
Expand Down Expand Up @@ -1031,7 +1031,7 @@ int main(int argc, char* argv[])

// extended boundary around the domain, and the processor domain
Ghost<3,double> g(2*H);

//! \cond [Initialization and parameters] \endcond

/*!
Expand Down Expand Up @@ -1493,4 +1493,3 @@ int main(int argc, char* argv[])
*
*/
}

14 changes: 7 additions & 7 deletions example/Vector/7_SPH_dlb_gpu/main.cu
Original file line number Diff line number Diff line change
Expand Up @@ -45,12 +45,12 @@
*
* my_kernel<<<wthr,thr>>>(arguments ... )
*
* Where wthr is the number of workgroups and thr is the number of threads in a workgroup and arguments... are the arguments to pass to the kernel.
* Where wthr is the number of workgroups and thr is the number of threads in a workgroup and arguments... are the arguments to pass to the kernel.
* Equivalently we can launch a kernel with the macro CUDA_LAUNCH_DIM3(my_kernel,wthr,thr,arguments...) or CUDA_LAUNCH(my_kernel,ite,arguments) where
* ite has been taken using getDomainIteratorGPU. There are several advantage on using CUDA_LAUNCH. The first advantage in using the macro is enabling SE_CLASS1
* all kernel launch become synchronous and an error check is performed before continue to the next kernel making debugging easier. Another feature is the possibility
* to run CUDA code on CPU without a GPU. compiling with "CUDA_ON_CPU=1 make" (Note openfpm must be compiled with GPU support (-g) or with CUDA_ON_CPU support
* (-c "... --enable_cuda_on_cpu"). You can compile this example on CPU. You do not have to change a single line of code for this example. (Check the video to see this
* to run CUDA code on CPU without a GPU. compiling with "CUDA_ON_CPU=1 make" (Note openfpm must be compiled with GPU support (-g) or with CUDA_ON_CPU support
* (-c "... --enable_cuda_on_cpu"). You can compile this example on CPU. You do not have to change a single line of code for this example. (Check the video to see this
* feature in action). All the openfpm GPU example and CUDA example can run on CPU if they use CUDA_LAUNCH as macro. We are planning to support
* AMD GPUs as well using this system.
*
Expand Down Expand Up @@ -223,7 +223,7 @@ inline __device__ __host__ real_number Wab(real_number r)
if (r < 1.0)
return (1.0 - 3.0/2.0*r*r + 3.0/4.0*r*r*r)*a2;
else if (r < 2.0)
return (1.0/4.0*(2.0 - r*r)*(2.0 - r*r)*(2.0 - r*r))*a2;
return (1.0/4.0*(2.0 - r)*(2.0 - r)*(2.0 - r))*a2;
else
return 0.0;
}
Expand Down Expand Up @@ -839,7 +839,7 @@ int main(int argc, char* argv[])

++obstacle_box;
}

vd.map();

// Now that we fill the vector with particles
Expand All @@ -848,7 +848,7 @@ int main(int argc, char* argv[])
vd.addComputationCosts(md);
vd.getDecomposition().decompose();
vd.map();

///////////////////////////

// Ok the initialization is done on CPU on GPU we are doing the main loop, so first we offload all properties on GPU
Expand Down Expand Up @@ -958,7 +958,7 @@ int main(int argc, char* argv[])

openfpm_finalize();
}

#else

int main(int argc, char* argv[])
Expand Down
4 changes: 2 additions & 2 deletions example/Vector/7_SPH_dlb_gpu_more_opt/main.cu
Original file line number Diff line number Diff line change
Expand Up @@ -225,7 +225,7 @@ inline __device__ __host__ real_number Wab(real_number r)
if (r < 1.0)
return (1.0 - 3.0/2.0*r*r + 3.0/4.0*r*r*r)*a2;
else if (r < 2.0)
return (1.0/4.0*(2.0 - r*r)*(2.0 - r*r)*(2.0 - r*r))*a2;
return (1.0/4.0*(2.0 - r)*(2.0 - r)*(2.0 - r))*a2;
else
return 0.0;
}
Expand Down Expand Up @@ -1101,7 +1101,7 @@ int main(int argc, char* argv[])

openfpm_finalize();
}

#else

int main(int argc, char* argv[])
Expand Down
5 changes: 2 additions & 3 deletions example/Vector/7_SPH_dlb_gpu_opt/main.cu
Original file line number Diff line number Diff line change
Expand Up @@ -215,7 +215,7 @@ inline __device__ __host__ real_number Wab(real_number r)
if (r < 1.0)
return (1.0 - 3.0/2.0*r*r + 3.0/4.0*r*r*r)*a2;
else if (r < 2.0)
return (1.0/4.0*(2.0 - r*r)*(2.0 - r*r)*(2.0 - r*r))*a2;
return (1.0/4.0*(2.0 - r)*(2.0 - r)*(2.0 - r))*a2;
else
return 0.0;
}
Expand Down Expand Up @@ -681,7 +681,6 @@ inline void sensor_pressure(Vector & vd,
CUDA_LAUNCH_DIM3(sensor_pressure_gpu,1,1,vd.toKernel(),NN.toKernel(),probe,(real_number *)press_tmp_.toKernel());

vd.template restoreOrder<>(NN);

// move calculated pressure on
press_tmp_.deviceToHost();
press_tmp = *(real_number *)press_tmp_.getPointer();
Expand Down Expand Up @@ -998,7 +997,7 @@ int main(int argc, char* argv[])

openfpm_finalize();
}

#else

int main(int argc, char* argv[])
Expand Down
9 changes: 4 additions & 5 deletions example/Vector/7_SPH_dlb_opt/main.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -383,11 +383,11 @@ inline double Wab(double r)
r /= H;

if (r < 1.0)
{return (1.0 - 3.0/2.0*r*r + 3.0/4.0*r*r*r)*a2;}
return (1.0 - 3.0/2.0*r*r + 3.0/4.0*r*r*r)*a2;
else if (r < 2.0)
{return (1.0/4.0*(2.0 - r*r)*(2.0 - r*r)*(2.0 - r*r))*a2;}
return (1.0/4.0*(2.0 - r)*(2.0 - r)*(2.0 - r))*a2;
else
{return 0.0;}
return 0.0;
}

const double c1 = -3.0/M_PI/H/H/H/H;
Expand Down Expand Up @@ -881,7 +881,7 @@ int main(int argc, char* argv[])
// extended boundary around the domain, and the processor domain
// by the support of the cubic kernel
Ghost<3,double> g(r_gskin);

/*! \cond [skin_calc] \endcond */

// Eliminating the lower part of the ghost
Expand Down Expand Up @@ -1141,4 +1141,3 @@ int main(int argc, char* argv[])

openfpm_finalize();
}

0 comments on commit 8b3531c

Please sign in to comment.