Skip to content
Snippets Groups Projects
Commit d8f49005 authored by Geray Karademir's avatar Geray Karademir
Browse files

Removing outdated file

parent 1359ca4a
Branches
No related tags found
No related merge requests found
Hide whitespace changes
Inline Side-by-side
Showing
with 0 additions and 241 deletions
OpenACC/density.hpp deleted 100644 → 0
+ 0
− 241
View file @ 1359ca4a
#include <openacc.h>
#if defined(LT_METAL_COOLING_WAL)
#define set_metallicities_gpu(i, Metallicities, dens_to_phys_fact) \
{ \
double nonHydmass; \
double mass; \
int pos, jjj, kkk; \
dens_to_phys_fact *= All.UnitDensity_in_cgs * All.HubbleParam * All.HubbleParam; \
mass = P[i].Mass; \
for (pos = 0, jjj = 0; jjj < WalCool_n_Ab; jjj++, pos++) \
if (SpeciesIdx[jjj] >= 0) \
{ \
if (jjj == myHyd) \
{ \
for (nonHydmass = 0, kkk = 0; kkk < LT_NMetP; kkk++) \
{ \
nonHydmass += SphP[i].Metals[kkk]; \
} \
Metallicities[SpeciesPos[jjj]] = \
(double)(mass - nonHydmass) / mass * SphP[i].Density * dens_to_phys_fact; \
} \
else \
{ \
Metallicities[SpeciesPos[jjj]] = (double)SphP[i].Metals[SpeciesIdx[jjj]] / mass; \
} \
if (jjj == myHel && UseHeNumberRatio) \
{ \
double d = dens_to_phys_fact * SphP[i].Density; \
Metallicities[SpeciesPos[jjj]] = d; \
} \
} \
Metallicities[HydPos] /= PROTONMASS; \
if (UseHeNumberRatio) \
{ \
Metallicities[HelPos] /= (4 * PROTONMASS * Metallicities[HydPos]); \
} \
Metallicities; \
}
#endif
inline int openacc_density_iteration_init(int &should_drift, int &DoPrimaryFirstIterationOnGPU, int &DoPrimaryOnGPU,
int *&DensityActiveParticleList, struct densdata_out *&LiteDOut)
{
DoPrimaryFirstIterationOnGPU = acc_all.go_gpu;
should_drift = !acc_all.go_gpu;
DensityActiveParticleList = (int *)mymalloc("DAPL", NActivePart * sizeof(int));
int NDensityActivePart = 0;
DoPrimaryOnGPU = DoPrimaryFirstIterationOnGPU;
for (int il = 0; il < NActivePart; il++)
{
int i = ActiveParticleList[il];
if (density_isactive(i))
{
DensityActiveParticleList[NDensityActivePart] = i;
NDensityActivePart++;
}
}
if (DoPrimaryFirstIterationOnGPU)
{
if (DoPrimaryOnGPU)
DoPrimaryOnGPU = NDensityActivePart > ACC_ACTIVEPART_CPU;
}
#pragma omp parallel for
for (int il = 0; il < NDensityActivePart; il++)
{
int i = DensityActiveParticleList[il];
struct densdata_out out;
memset(&out, 0, sizeof(struct densdata_out));
out2particle_density(&out, i, 0);
}
if (DoPrimaryOnGPU)
{
LiteDOut = (struct densdata_out *)mymalloc("LiteDO", NDensityActivePart * sizeof(struct densdata_out));
// let's zero all output-data.
}
return NDensityActivePart;
}
#pragma acc routine
int density_evaluate_local(int active_target, int p_target, int mode, int *ngblist, int numngb,
struct densdata_in *DensDataGet, struct densdata_out *DensDataResult,
struct gravity_walkdata *GP, struct densdata_in *P, int i_node, int node);
#pragma acc routine
int ngb_treefind_variable_threads_local(MyLongDouble *searchcenter, MyFloat hsml, int mode, int *startnode,
struct gravity_walkdata *P, int *ngblist, int maxnumngb, int i_node);
void openacc_density_first_phase(int DoPrimaryOnGPU, int NDensityActivePart, int *DensityActiveParticleList,
struct densdata_out *LiteDOut, struct densdata_in *LiteDIn)
{
if (DoPrimaryOnGPU && NDensityActivePart > 0)
{
#ifndef ACC_NOACC
int MiniNGBList[ACC_GPU_NGB];
#pragma acc parallel loop gang vector present(WalkLiteP [0:NumPart], LiteDOut [0:NDensityActivePart]) private( \
MiniNGBList) async(1) if (DoPrimaryOnGPU)
#else
#pragma omp parallel for if (DoPrimaryOnGPU)
#endif
for (int i = 0; i < NDensityActivePart; i++)
{
#ifdef ACC_NOACC
int MiniNGBList[ACC_GPU_NGB];
#endif
int p_target = DensityActiveParticleList[i];
int startnode = All.MaxPart;
int endnode = startnode;
int i_node = 0;
while (endnode >= 0)
{
startnode = endnode;
int numngb = ngb_treefind_variable_threads_local(WalkLiteP[p_target].Pos, WalkLiteP[p_target].Hsml, 2,
&endnode, WalkLiteP, MiniNGBList, ACC_GPU_NGB, i_node);
int ret = density_evaluate_local(i, p_target, 2, MiniNGBList, numngb, NULL, LiteDOut, WalkLiteP,
LiteDIn, i_node, startnode);
// int ret = 0;
i_node++;
}
// mode is 2 because the computation is very ismilar to the second phae (mode==1) where you have
// datain/dataout structures and the local P array
// break;
} // end parallel for
}
}
inline void openacc_density_secondary_phase(struct densdata_in *DensDataGet, struct densdata_out *DensDataResult,
struct densdata_in *LiteDIn, int DoPrimaryOnGPU, int should_drift)
{
int DoSecondaryOnGPU =
DoPrimaryOnGPU && DoPrimaryOnGPU && (Nimport > ACC_ACTIVEPART_II_CPU); // && acc_async_test(1);
if (!DoSecondaryOnGPU)
{
#pragma omp parallel
{
int mainthreadid = ThisThread;
ar_green_evaluate_secondary(DensDataGet, density_evaluate, Module_Density, TreeWalkType_GreenTreeWalk,
NeighbourSearchCriteria_Density, &ThisThread, should_drift);
}
}
else if (Nimport > 0)
{
int MiniNGBListII[ACC_GPU_NGB];
#pragma acc wait(1)
#pragma acc parallel loop gang vector copyin(DensDataGet [0:Nimport], DataNodeListGet [0:Nimport]) \
present(WalkLiteP [0:NumPart]) present(LiteDIn [0:NumPart]) \
copyout(DensDataResult [0:Nimport]) private(MiniNGBListII)
for (int i = 0; i < Nimport; i++)
{
int dummy = 55;
int node_list_i = 0;
int i_node = 0;
for (int node_list_i = 0; node_list_i < NODELISTLENGTH; node_list_i++)
{
if (DataNodeListGet[i].NodeList[node_list_i] == -1)
break;
int startnode = DataNodeListGet[i].NodeList[node_list_i];
int endnode = startnode;
int first_iter = 1;
while (endnode >= 0)
{
startnode = endnode;
int numngb = ngb_treefind_variable_threads_local(
DensDataGet[i].Pos, DensDataGet[i].Hsml, 1, &endnode, WalkLiteP,
/* WalkLiteSPH,*/ MiniNGBListII, ACC_GPU_NGB, first_iter);
int ret = density_evaluate_local(i, i, 1, MiniNGBListII, numngb, DensDataGet, DensDataResult,
WalkLiteP, LiteDIn, i_node, startnode);
i_node++;
first_iter = 0;
} // end while
} // end for nodes
} // end parallel for
VERBOSE(3, " DONE SECOND FASE 1");
}
}
inline struct densdata_in *openacc_density_init(int &should_drift, int &DoPrimaryFirstIterationOnGPU)
{
TIMER_PUSH(8, "ACC-density_init");
if (acc_all.go_gpu == -1)
{
// first iteration starts with a density smoothing length finder:
update_walk_lite_p();
acc_all.go_gpu = 1;
// next time we will run or not on GPU depending on what the gravity loop decided.
}
VERBOSE(3, " acc_all.inside_data_region %d", acc_all.inside_data_region);
if (!acc_all.inside_data_region)
{
update_openacc_structures();
}
VERBOSE(3, " acc_all.inside_data_region %d 0", acc_all.inside_data_region);
DoPrimaryFirstIterationOnGPU = acc_all.go_gpu;
should_drift = !acc_all.go_gpu;
struct densdata_in *LiteDIn = (densdata_in *)mymalloc("LiteDI", NumPart * sizeof(struct densdata_in));
VERBOSE(3, " acc_all.inside_data_region %d 1", acc_all.inside_data_region);
if (DoPrimaryFirstIterationOnGPU)
{
#pragma omp parallel for
for (int i = 0; i < NumPart; i++)
{
WalkLiteP[i].Hsml = P[i].Hsml;
particle2in_density(&LiteDIn[i], i);
}
}
VERBOSE(3, " acc_all.inside_data_region %d 2", acc_all.inside_data_region);
TIMER_POP("ACC-density_init");
return LiteDIn;
}
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment