diff --git a/HySoP/hysop/gpu/cl_src/kernels/fine_to_coarse_filter.cl b/HySoP/hysop/gpu/cl_src/kernels/fine_to_coarse_filter.cl index 9f0436f7b562335cd172c089551b8d65d07cb5cb..2851c60383e2d634a729ab5cbeb0360e5c5da5a9 100644 --- a/HySoP/hysop/gpu/cl_src/kernels/fine_to_coarse_filter.cl +++ b/HySoP/hysop/gpu/cl_src/kernels/fine_to_coarse_filter.cl @@ -1,3 +1,7 @@ +__kernel void initialize_output(__global float* scal_out) { + scal_out[get_global_id(0) + get_global_id(1)*NB_OUT_X + get_global_id(2)*NB_OUT_X*NB_OUT_Y] = 0.0; +} + __kernel void coarse_to_fine_filter(__global const float* scal_in, __global float* scal_out, float scale_factor, diff --git a/HySoP/hysop/gpu/gpu_multiresolution_filter.py b/HySoP/hysop/gpu/gpu_multiresolution_filter.py index 091a53490d81fd65c950cd6ea9c9567501dd00f1..f388f61c8cd46b73f84d99a3eb6bd83ce71eab91 100644 --- a/HySoP/hysop/gpu/gpu_multiresolution_filter.py +++ b/HySoP/hysop/gpu/gpu_multiresolution_filter.py @@ -124,9 +124,14 @@ class GPUFilterFineToCoarse(FilterFineToCoarse, GPUOperator): prg = self.cl_env.build_src(src, build_options, vec) self.fine_to_coarse = KernelLauncher( prg.coarse_to_fine_filter, self.cl_env.queue, gwi, lwi) + self.initialize = KernelLauncher( + prg.initialize_output, self.cl_env.queue, + self.field_out[0].data[0].shape, None) def apply(self, simulation=None): #evts = [] + self.field_out[0].events.append( + self.initialize(self.field_out[0].gpu_data[0])) for iy in xrange(len(self._rmsh.weights)): for iz in xrange(len(self._rmsh.weights)): evt = self.fine_to_coarse(self.field_in[0].gpu_data[0], @@ -307,5 +312,6 @@ class GPUFilterFineToCoarse(FilterFineToCoarse, GPUOperator): # set_gh_zr.wait() def get_profiling_info(self): - for p in self.fine_to_coarse.profile: - self.profiler += p + for k in (self.fine_to_coarse, self.initialize): + for p in k.profile: + self.profiler += p diff --git a/HySoP/hysop/operator/discrete/baroclinic.py b/HySoP/hysop/operator/discrete/baroclinic.py index 2dabb187900bc38c4a6d705b5f3ec12990f783f1..b02527fb21f4b62ea9f2947b8490e587f42638f8 100644 --- a/HySoP/hysop/operator/discrete/baroclinic.py +++ b/HySoP/hysop/operator/discrete/baroclinic.py @@ -64,6 +64,9 @@ class Baroclinic(DiscreteOperator): indices=self.velocity.topology.mesh.iCompute, method=self.method[SpaceDiscretisation]) + # Gravity vector + self._gravity = npw.asrealarray([0., 0., -9.81]) + # Time stem of the previous iteration self._old_dt = None @@ -139,7 +142,8 @@ class Baroclinic(DiscreteOperator): self._result[d][iCompute] -= self._tempGrad[d][iCompute] # gravity term : result = result - g - ### Seems not working: self._result[2][iCompute] -= 1. + for d in xrange(self.velocity.dimension): + self._result[2][iCompute] -= self._gravity[d] # baroclinicTerm = -(gradRho/rho) x (gradP/rho) self._tempGrad = self._gradOp(self.density[0:1], self._tempGrad)