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