aboutsummaryrefslogtreecommitdiff
path: root/template
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-09-01 23:06:10 +0200
committerAdrian Kummerlaender2019-09-01 23:06:10 +0200
commita89d3eba9aec11a4f55817fcdb189fcd6f26b574 (patch)
tree929f411a20564c9e267116f3e6a8966fe6c3a781 /template
parent0a64777ddc063d9b3f4abd68383e4ef998977bb7 (diff)
downloadsymlbm_playground-a89d3eba9aec11a4f55817fcdb189fcd6f26b574.tar
symlbm_playground-a89d3eba9aec11a4f55817fcdb189fcd6f26b574.tar.gz
symlbm_playground-a89d3eba9aec11a4f55817fcdb189fcd6f26b574.tar.bz2
symlbm_playground-a89d3eba9aec11a4f55817fcdb189fcd6f26b574.tar.lz
symlbm_playground-a89d3eba9aec11a4f55817fcdb189fcd6f26b574.tar.xz
symlbm_playground-a89d3eba9aec11a4f55817fcdb189fcd6f26b574.tar.zst
symlbm_playground-a89d3eba9aec11a4f55817fcdb189fcd6f26b574.zip
Prototype "ink" particles visualization
Diffstat (limited to 'template')
-rw-r--r--template/kernel.mako41
1 files changed, 33 insertions, 8 deletions
diff --git a/template/kernel.mako b/template/kernel.mako
index 00a4345..f6bac7f 100644
--- a/template/kernel.mako
+++ b/template/kernel.mako
@@ -100,6 +100,7 @@ __kernel void collect_moments(__global __read_only ${float_type}* f,
}
__kernel void collect_gl_moments(__global __read_only ${float_type}* f,
+ __global __read_only int* material,
__global __write_only float4* moments)
{
const unsigned int gid = ${gid()};
@@ -116,17 +117,41 @@ __kernel void collect_gl_moments(__global __read_only ${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];
+ }
moments[gid] = data;
}
+
+__kernel void update_particles(__global __read_only float4* moments,
+ __global __write_only float4* particles)
+{
+ const unsigned int pid = get_global_id(0);
+
+ float4 particle = particles[pid];
+
+ const unsigned int gid = floor(particle.y)*${memory.size_x} + floor(particle.x);
+
+ float4 moment = moments[gid];
+
+ particle.x += moment.y;
+ particle.y += moment.z;
+
+ particles[pid] = particle;
+}