From f5ee688a95a3ad87cd5a7a996e657bf871f52384 Mon Sep 17 00:00:00 2001
From: Jean-Matthieu Etancelin <jean-matthieu.etancelin@univ-reims.fr>
Date: Mon, 16 Mar 2015 10:11:17 +0100
Subject: [PATCH] Fix multiresolution filter. Fix gravity term in baroclinic
operator.
---
.../hysop/gpu/cl_src/kernels/fine_to_coarse_filter.cl | 4 ++++
HySoP/hysop/gpu/gpu_multiresolution_filter.py | 10 ++++++++--
HySoP/hysop/operator/discrete/baroclinic.py | 6 +++++-
3 files changed, 17 insertions(+), 3 deletions(-)
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 9f0436f7b..2851c6038 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 091a53490..f388f61c8 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 2dabb1879..b02527fb2 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)
--
GitLab