Commit edf4f2b8 authored by thomaset's avatar thomaset
Browse files

remove strange geany comments //~

parent 07baa963
......@@ -106,7 +106,8 @@ __device__ float3 _dpd_interaction( const int dpid, const float4 xdest, const fl
#ifdef PAdL
__device__ float3 _dpd_interaction_G_update( const int dpid, const float4 xdest, const float4 udest, const float4 xsrc, const float4 usrc, const int spid )
{
//~ const float myrandnr = Logistic::mean0var1( info.seed, xmin( spid, dpid ), xmax( spid, dpid ) );
// TODO: remove unused arguments of _dpd_interaction_G_update
//const float myrandnr = Logistic::mean0var1( info.seed, xmin( spid, dpid ), xmax( spid, dpid ) );
// check for particle types and compute the DPD force
float3 pos1 = make_float3(xdest.x, xdest.y, xdest.z), pos2 = make_float3(xsrc.x, xsrc.y, xsrc.z);
......@@ -275,8 +276,6 @@ __forceinline__ __device__ void core_ytang( const uint dststart, const uint psha
const uint sentry = xscale( spid, 2.f );
const float4 xdest = tex1Dfetch( texParticlesF4, dentry );
const float4 xsrc = tex1Dfetch( texParticlesF4, sentry );
//~ const float4 udest = tex1Dfetch( texParticlesF4, xadd( dentry, 1u ) );
//~ const float4 usrc = tex1Dfetch( texParticlesF4, xadd( sentry, 1u ) );
const float3 f = _dpd_interaction( dpid, xdest, xsrc, spid );
// the overhead of transposition acc back
......@@ -1467,13 +1466,13 @@ void _dpd_velocity_update_stage_0_2x2x2(
#ifdef LOWE
const float myrandnr2 = Logistic::mean0var1( info.seed, xmax( i, j ), xmin( i, j ) ); // non-symmetric, same seed can be used
#endif
//~ const float4 ri = tex1Dfetch( texParticlesH4, i);
//~ const float4 rj = tex1Dfetch( texParticlesH4, j);
//~ const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//~ const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//~ float *v_i = pp[i].v;
//~ float *v_j = pp[j].v;
// quickfix since tex1Dfetch produces bug, I don't know why...
//const float4 ri = tex1Dfetch( texParticlesH4, i);
//const float4 rj = tex1Dfetch( texParticlesH4, j);
//const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//float *v_i = pp[i].v;
//float *v_j = pp[j].v;
// TODO: quickfix since tex1Dfetch produces bug, I don't know why..., sometimes I get the ri==rj back... why
float *v_i = pp[i].v, *ri = pp[i].r;
float *v_j = pp[j].v, *rj = pp[j].r;
const float3 r_i = make_float3(ri[0], ri[1], ri[2]);
......@@ -1511,13 +1510,13 @@ printf("undefined velocity correction\n");
#ifdef LOWE
const float myrandnr2 = Logistic::mean0var1( info.seed, xmax( i, j ), xmin( i, j ) ); // non-symmetric, same seed can be used
#endif
//~ const float4 ri = tex1Dfetch( texParticlesH4, i);
//~ const float4 rj = tex1Dfetch( texParticlesH4, j);
//~ const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//~ const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//~ float *v_i = pp[i].v;
//~ float *v_j = pp[j].v;
// quickfix since tex1Dfetch produces bug, I don't know why...
//const float4 ri = tex1Dfetch( texParticlesH4, i);
//const float4 rj = tex1Dfetch( texParticlesH4, j);
//const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//float *v_i = pp[i].v;
//float *v_j = pp[j].v;
// TODO: quickfix since tex1Dfetch produces bug, I don't know why..., sometimes I get the ri==rj back... why.
float *v_i = pp[i].v, *ri = pp[i].r;
float *v_j = pp[j].v, *rj = pp[j].r;
const float3 r_i = make_float3(ri[0], ri[1], ri[2]);
......@@ -1612,13 +1611,13 @@ void _dpd_velocity_update_stage_1_7_2x2x2(
#ifdef LOWE
const float myrandnr2 = Logistic::mean0var1( info.seed, xmax( i, j ), xmin( i, j ) ); // non-symmetric, same seed can be used
#endif
//~ const float4 ri = tex1Dfetch( texParticlesH4, i);
//~ const float4 rj = tex1Dfetch( texParticlesH4, j);
//~ const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//~ const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//~ float *v_i = pp[i].v;
//~ float *v_j = pp[j].v;
// quickfix since tex1Dfetch produces bug, I don't know why...
//const float4 ri = tex1Dfetch( texParticlesH4, i);
//const float4 rj = tex1Dfetch( texParticlesH4, j);
//const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//float *v_i = pp[i].v;
//float *v_j = pp[j].v;
// TODO: quickfix since tex1Dfetch produces bug, I don't know why..., sometimes I get the ri==rj back... why
float *v_i = pp[i].v, *ri = pp[i].r;
float *v_j = pp[j].v, *rj = pp[j].r;
const float3 r_i = make_float3(ri[0], ri[1], ri[2]);
......@@ -1862,13 +1861,13 @@ _dpd_velocity_update_stage_0(
#ifdef LOWE
const float myrandnr2 = Logistic::mean0var1( info.seed, xmax( i, j ), xmin( i, j ) ); // non-symmetric, same seed can be used
#endif
//~ const float4 ri = tex1Dfetch( texParticlesH4, i);
//~ const float4 rj = tex1Dfetch( texParticlesH4, j);
//~ const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//~ const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//~ float *v_i = pp[i].v;
//~ float *v_j = pp[j].v;
// quickfix since tex1Dfetch produces bug, I don't know why...
//const float4 ri = tex1Dfetch( texParticlesH4, i);
//const float4 rj = tex1Dfetch( texParticlesH4, j);
//const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//float *v_i = pp[i].v;
//float *v_j = pp[j].v;
// TODO: quickfix since tex1Dfetch produces bug, I don't know why..., sometimes I get the ri==rj back... why
float *v_i = pp[i].v, *ri = pp[i].r;
float *v_j = pp[j].v, *rj = pp[j].r;
const float3 r_i = make_float3(ri[0], ri[1], ri[2]);
......@@ -1962,10 +1961,10 @@ _dpd_velocity_update_stage_1_26(
for(int l = idx; l < idx + i_per_thread; ++l){
if(l < start_count_smaller.y){
const uint i = start_count_smaller.x + l;
//~ const float4 ri = tex1Dfetch( texParticlesH4, i);
//~ const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//~ float *v_i = pp[i].v;
// quickfix since tex1Dfetch produces bug, I don't know why...
//const float4 ri = tex1Dfetch( texParticlesH4, i);
//const float3 r_i = make_float3(ri.x, ri.y, ri.z);
//float *v_i = pp[i].v;
// TODO: quickfix since tex1Dfetch produces bug, I don't know why..., sometimes I get the ri==rj back... why
float *v_i = pp[i].v, *ri = pp[i].r;
const float3 r_i = make_float3(ri[0], ri[1], ri[2]);
for(uint interaction = 0; interaction < start_count_larger.y; ++interaction){
......@@ -1974,10 +1973,10 @@ _dpd_velocity_update_stage_1_26(
#ifdef LOWE
const float myrandnr2 = Logistic::mean0var1( info.seed, xmax( i, j ), xmin( i, j ) ); // non-symmetric, same seed can be used
#endif
//~ const float4 rj = tex1Dfetch( texParticlesH4, j);
//~ const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//~ float *v_j = pp[j].v;
// quickfix since tex1Dfetch produces bug, I don't know why...
//const float4 rj = tex1Dfetch( texParticlesH4, j);
//const float3 r_j = make_float3(rj.x, rj.y, rj.z);
//float *v_j = pp[j].v;
// TODO: quickfix since tex1Dfetch produces bug, I don't know why..., sometimes I get the ri==rj back... why
float *v_j = pp[j].v, *rj = pp[j].r;
const float3 r_j = make_float3(rj[0], rj[1], rj[2]);
// synchronize threads to have the necessary serial velocity update
......@@ -2065,13 +2064,13 @@ _dpd_velocity_update_stage_0_26(
#ifdef LOWE
const float myrandnr2 = Logistic::mean0var1( info.seed, xmax( i, j ), xmin( i, j ) ); // non-symmetric, same seed can be used
#endif
//~ const float4 ri = tex1Dfetch( texParticlesH4, i); // moving in 1st loop is no improvement
//~ const float4 rj = tex1Dfetch( texParticlesH4, j);
//~ const float r_i[3] = {ri.x, ri.y, ri.z};
//~ const float r_j[3] = {rj.x, rj.y, rj.z};
//~ float *v_i = pp[i].v;
//~ float *v_j = pp[j].v;
// quickfix since tex1Dfetch produces bug, I don't know why...
//const float4 ri = tex1Dfetch( texParticlesH4, i); // moving in 1st loop is no improvement
//const float4 rj = tex1Dfetch( texParticlesH4, j);
//const float r_i[3] = {ri.x, ri.y, ri.z};
//const float r_j[3] = {rj.x, rj.y, rj.z};
//float *v_i = pp[i].v;
//float *v_j = pp[j].v;
// TODO: quickfix since tex1Dfetch produces bug, I don't know why..., sometimes I get the ri==rj back... why
float *v_i = pp[i].v, *r_i = pp[i].r;
float *v_j = pp[j].v, *r_j = pp[j].r;
#ifdef SSA
......@@ -2791,7 +2790,7 @@ void pairwise_velocity_update_cuda_nohost(
c.invdomainsize = make_float3( 1 / XL, 1 / YL, 1 / ZL );
c.domainstart = make_float3( -XL * 0.5, -YL * 0.5, -ZL * 0.5 );
c.invrc = 1.f / rc;
//~ c.axayaz = axayaz;
//c.axayaz = axayaz;
c.seed = seed;
if (!is_mps_enabled)
......@@ -2803,9 +2802,9 @@ void pairwise_velocity_update_cuda_nohost(
static int cetriolo = 0;
cetriolo++;
//~ int np32 = np;
//~ if( np32 % 32 ) np32 += 32 - np32 % 32;
//~ CC( cudaMemsetAsync( axayaz, 0, sizeof( float )* np32 * 3) );
//int np32 = np;
//if( np32 % 32 ) np32 += 32 - np32 % 32;
//CC( cudaMemsetAsync( axayaz, 0, sizeof( float )* np32 * 3) );
#if !defined(BLOCK_PARALLELIZATION) && !defined(SINGLE_THREAD_PER_CELL)
// Own parallelization scheme, more threads working than Larentzos, but also more stages
......@@ -2977,9 +2976,9 @@ void pairwise_velocity_update_cuda_nohost(
}
#endif
}
//~ #ifdef TRANSPOSED_ATOMICS
//~ transpose_acc <<< 28, 1024, 0>>>( np );
//~ #endif
//#ifdef TRANSPOSED_ATOMICS
//transpose_acc <<< 28, 1024, 0>>>( np );
//#endif
} else {
fprintf( stderr, "Incompatible grid config in velocity_update_ssa_cuda_nohost \n" );
}
......@@ -3102,7 +3101,6 @@ void G_update_dpd_cuda_nohost( const float * const xyzuvw, const float4 * const
CC( cudaMemsetAsync( G, 0, sizeof( float )* np32 * 3) );
if( c.ncells.x % MYCPBX == 0 && c.ncells.y % MYCPBY == 0 && c.ncells.z % MYCPBZ == 0 ) {
//~ _dpd_forces_symm_merged <<< dim3( c.ncells.x / MYCPBX, c.ncells.y / MYCPBY, c.ncells.z / MYCPBZ ), dim3( 32, MYWPB ), 0>>> ();
G_update_dpd_symm_merged <<< dim3( c.ncells.x / MYCPBX, c.ncells.y / MYCPBY, c.ncells.z / MYCPBZ ), dim3( 32, MYWPB ), 0>>> ();
#ifdef TRANSPOSED_ATOMICS
transpose_acc <<< 28, 1024, 0>>>( np );
......
......@@ -7,7 +7,7 @@
#define numberdensity 10
#define kBT 0.0444302
//~ #define dt 5e-4 /* timestep */
//#define dt 5e-4 /* timestep */
#define rbc_mass 0.5 /* ratio of RBC particle mass to solvent particle mass */
#define gamma_dot 0.5f /* shear rate */
#define hydrostatic_a 0.0f /* flow acceleration for Poiseuille setup */
......
......@@ -7,7 +7,7 @@
#define numberdensity 10
#define kBT 0.0444302
//~ #define dt 5e-4 /* timestep */
//#define dt 5e-4 /* timestep */
#define rbc_mass 0.5 /* ratio of RBC particle mass to solvent particle mass */
#define gamma_dot 0.0 /* shear rate */
#define hydrostatic_a 0.0333333 /* flow acceleration for Poiseuille setup */
......
......@@ -34,11 +34,6 @@ __device__ float3 compute_dpd_force_traced_stage1(int type1, int type2,
const float yr = _yr * invrij;
const float zr = _zr * invrij;
//~ const float rdotv =
//~ xr * (vel1.x - vel2.x) +
//~ yr * (vel1.y - vel2.y) +
//~ zr * (vel1.z - vel2.z);
const float gammadpd_pair = 0.5 * (gammadpd_device[type1] + gammadpd_device[type2]);
const float sigmaf_pair = sqrtf(2*gammadpd_pair*kBT / dt);
float strength = (sigmaf_pair * myrandnr) * wr;
......
......@@ -110,14 +110,14 @@ void forces_fsi(SolventWrap *w_s, std::vector<ParticlesWrap> *w_r) {
// update all forces
// TODO: check if working for fsi, rbc, ...
void forces_stage1(bool wall_created) {
//~ SolventWrap w_s(s_pp, s_n, s_ff, cells->start, cells->count);
//~ std::vector<ParticlesWrap> w_r;
//~ if (rbcs0) w_r.push_back(ParticlesWrap(r_pp, r_n, r_ff));
//SolventWrap w_s(s_pp, s_n, s_ff, cells->start, cells->count);
//std::vector<ParticlesWrap> w_r;
//if (rbcs0) w_r.push_back(ParticlesWrap(r_pp, r_n, r_ff));
clear_forces(s_ff, s_n);
clear_forces(s_ff_diss, s_n);
//~ if (rbcs0) clear_forces(r_ff, r_n);
//if (rbcs0) clear_forces(r_ff, r_n);
forces_dpd_stage1();
if (wall_created){
......@@ -125,47 +125,47 @@ void forces_stage1(bool wall_created) {
forces_wall_stage2();
}
//~ if (rbcs0) {
//~ forces_fsi(&w_s, &w_r);
//if (rbcs0) {
//forces_fsi(&w_s, &w_r);
//~ rex::bind_solutes(w_r);
//~ rex::pack_p();
//~ rex::post_p();
//~ rex::recv_p();
//rex::bind_solutes(w_r);
//rex::pack_p();
//rex::post_p();
//rex::recv_p();
//~ rex::halo(); /* fsi::halo(); */
//rex::halo(); /* fsi::halo(); */
//~ rex::post_f();
//~ rex::recv_f();
//~ }
//rex::post_f();
//rex::recv_f();
//}
}
// update only dissipative forces
void forces_stage2(bool wall_created) {
// SolventWrap w_s(s_pp, s_n, s_ff, cells->start, cells->count);
// std::vector<ParticlesWrap> w_r;
// if (rbcs0) w_r.push_back(ParticlesWrap(r_pp, r_n, r_ff));
//SolventWrap w_s(s_pp, s_n, s_ff, cells->start, cells->count);
//std::vector<ParticlesWrap> w_r;
//if (rbcs0) w_r.push_back(ParticlesWrap(r_pp, r_n, r_ff));
clear_forces(s_ff_diss, s_n);
// if (rbcs0) clear_forces(r_ff, r_n);
//if (rbcs0) clear_forces(r_ff, r_n);
forces_dpd_stage2();
// is also some kind of dissipative force!
//is also some kind of dissipative force!
if (wall_created) forces_wall_stage2();
// if (rbcs0) {
// forces_fsi(&w_s, &w_r);
//if (rbcs0) {
//forces_fsi(&w_s, &w_r);
// rex::bind_solutes(w_r);
// rex::pack_p();
// rex::post_p();
// rex::recv_p();
//rex::bind_solutes(w_r);
//rex::pack_p();
//rex::post_p();
//rex::recv_p();
// rex::halo(); /* fsi::halo(); */
//rex::halo(); /* fsi::halo(); */
// rex::post_f();
// rex::recv_f();
// }
//rex::post_f();
//rex::recv_f();
//}
}
#elif defined(SSA) || defined(LOWE) || defined(PETERS) || defined(PAdL)
void forces(bool wall_created) {
......@@ -427,8 +427,8 @@ void run0(float driving_force, bool wall_created, int it) {
body_force(driving_force);
// only velocity update --> no redistribution
k_sim::update_stage2<<<k_cnf(s_n)>>> (false, s_pp, s_ff, s_n);
// TODO: is this the correct order for these functions?
if (rbcs0) update_r(); // not implemented VV
if (rbcs0) update_r();
if (rbcs0) bounce_solid();
dumps_diags(it);
}
......@@ -446,22 +446,18 @@ void run0(float driving_force, bool wall_created, int it) {
k_sim::update_stage2<<<k_cnf(s_n)>>> (false, s_pp, s_ff, s_ff_diss, s_n);
forces_stage2(wall_created);
// TODO: is this the correct order for these functions?
if (rbcs0) update_r();
if (rbcs0) bounce_solid();
dumps_diags(it);
}
#elif defined SSA
void run0(float driving_force, bool wall_created, int it) {
//~ if(it < 2)
//~ cudaProfilerStart();
// update zipped arrays for velocity update
update_helper_arrays();
// do pairwise velocity update
DPD::pairwise_velocity_update(s_pp, s_zip0, s_zip1,
s_n, cells->start, cells->count);
//~ if(it < 2)
//~ cudaProfilerStop();
// update velocity with conservative force (stage 1)
k_sim::update_stage1<<<k_cnf(s_n)>>> (false, s_pp, s_ff, s_n);
// bounceback of solvent after the update of the particle positions
......@@ -473,7 +469,7 @@ void run0(float driving_force, bool wall_created, int it) {
body_force(driving_force);
// update velocities with conservative force (stage 2)
k_sim::update_stage2<<<k_cnf(s_n)>>> (false, s_pp, s_ff, s_n);
// TODO: is this the correct order for these functions?
if (rbcs0) update_r();
if (rbcs0) bounce_solid();
dumps_diags(it);
......@@ -491,16 +487,13 @@ void run0(float driving_force, bool wall_created, int it) {
body_force(driving_force);
// update velocities with conservative force (stage 2)
k_sim::update_stage2<<<k_cnf(s_n)>>> (false, s_pp, s_ff, s_n);
//~ if(it < 2)
//~ cudaProfilerStart();
// update zipped arrays for velocity update
update_helper_arrays();
// do pairwise velocity update
DPD::pairwise_velocity_update(s_pp, s_zip0, s_zip1,
s_n, cells->start, cells->count);
//~ if(it < 2)
//~ cudaProfilerStop();
// TODO: is this the correct order for these functions?
if (rbcs0) update_r();
if (rbcs0) bounce_solid();
dumps_diags(it);
......@@ -519,8 +512,6 @@ void run0(float driving_force, bool wall_created, int it) {
body_force(driving_force);
// update velocities with conservative force (stage 2)
k_sim::update_stage2<<<k_cnf(s_n)>>> (false, s_pp, s_ff, s_n);
//~ if(it < 2)
//~ cudaProfilerStart();
// update zipped arrays for velocity update
update_helper_arrays();
// do 1st pairwise velocity update
......@@ -540,12 +531,11 @@ void run0(float driving_force, bool wall_created, int it) {
// do 2nd pairwise veloctity update
DPD::pairwise_velocity_update(s_pp, s_zip0, s_zip1,
s_n, cells->start, cells->count, xi);
//~ if(it < 2)
//~ cudaProfilerStop();
// update position and velocity
// --> merged in start of run0
// TODO: is this the correct order for these functions?
// TODO: should maybe moved up after position update?
if (rbcs0) update_r();
if (rbcs0) bounce_solid();
dumps_diags(it);
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment