From a3ac409a49803956986f4cab35a1f90f8d4b8e81 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Fri, 19 Jun 2020 22:01:57 +0200
Subject: Use lattice-internal moments buffer for particles

---
 template/opengl.mako    | 24 ++++++++++++++++--------
 template/particles.mako | 10 ++++------
 2 files changed, 20 insertions(+), 14 deletions(-)

(limited to 'template')

diff --git a/template/opengl.mako b/template/opengl.mako
index 47264e8..2029623 100644
--- a/template/opengl.mako
+++ b/template/opengl.mako
@@ -65,6 +65,7 @@ def neighbor_offset(c_i):
 %>
 
 __kernel void collect_gl_moments_to_texture(__global ${float_type}* f,
+                                            __global int* material,
 % if descriptor.d == 2:
                                             __write_only image2d_t moments)
 % elif descriptor.d == 3:
@@ -85,17 +86,24 @@ __kernel void collect_gl_moments_to_texture(__global ${float_type}* f,
 
     float4 data;
 
+    if (material[gid] == 1) {
 % if descriptor.d == 2:
-    data.x = ${ccode(moments_assignment[0].rhs)};
-    data.y = ${ccode(moments_assignment[1].rhs)};
-    data.z = ${ccode(moments_assignment[2].rhs)};
-    data.w = sqrt(data.y*data.y + data.z*data.z);
+      data.x = ${ccode(moments_assignment[0].rhs)};
+      data.y = ${ccode(moments_assignment[1].rhs)};
+      data.z = ${ccode(moments_assignment[2].rhs)};
+      data.w = sqrt(data.y*data.y + data.z*data.z);
 % elif descriptor.d == 3:
-    data.x = ${ccode(moments_assignment[0].rhs)};
-    data.y = ${ccode(moments_assignment[1].rhs)};
-    data.z = ${ccode(moments_assignment[2].rhs)};
-    data.w = ${ccode(moments_assignment[3].rhs)};
+      data.x = ${ccode(moments_assignment[0].rhs)};
+      data.y = ${ccode(moments_assignment[1].rhs)};
+      data.z = ${ccode(moments_assignment[2].rhs)};
+      data.w = ${ccode(moments_assignment[3].rhs)};
 % endif
+    } else {
+      data.x = 0.0;
+      data.y = 0.0;
+      data.z = 0.0;
+      data.w = -material[gid];
+    }
 
     write_imagef(moments, ${moments_cell()}, data);
 }
diff --git a/template/particles.mako b/template/particles.mako
index 48191d9..0b9a67e 100644
--- a/template/particles.mako
+++ b/template/particles.mako
@@ -1,4 +1,4 @@
-__kernel void update_particles(__global float4* moments,
+__kernel void update_particles(__global float*  moments,
                                __global int*    material,
                                __global float4* particles,
                                __global float4* init_particles,
@@ -14,13 +14,11 @@ __kernel void update_particles(__global float4* moments,
   const unsigned int gid = floor(particle.z)*${memory.size_x*memory.size_y} + floor(particle.y)*${memory.size_x} + floor(particle.x);
 % endif
 
-  const float4 moment = moments[gid];
-
   if (material[gid] == 1 && particle.w < 1.0) {
-    particle.x += moment.y;
-    particle.y += moment.z;
+    particle.x += moments[${1*memory.volume}+gid];
+    particle.y += moments[${2*memory.volume}+gid];
 % if descriptor.d == 3:
-    particle.z += moment.w;
+    particle.z += moments[${3*memory.volume}+gid];
 % endif
     particle.w += min(particle.x, particle.y) * aging;
   } else {
-- 
cgit v1.2.3