aboutsummaryrefslogtreecommitdiff
path: root/template
diff options
context:
space:
mode:
Diffstat (limited to 'template')
-rw-r--r--template/kernel.mako79
-rw-r--r--template/opengl.mako100
-rw-r--r--template/particles.mako38
3 files changed, 138 insertions, 79 deletions
diff --git a/template/kernel.mako b/template/kernel.mako
index a57c2ec..025baf5 100644
--- a/template/kernel.mako
+++ b/template/kernel.mako
@@ -98,82 +98,3 @@ __kernel void collect_moments(__global __read_only ${float_type}* f,
moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)};
% endfor
}
-
-__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()};
-
- __global __read_only ${float_type}* preshifted_f = f + gid;
-
-% for i in range(0,descriptor.q):
- const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
- 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);
-% 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)};
-% 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 __read_only int* material,
- __global __write_only float4* particles,
- __global __read_only float4* init_particles,
- float aging)
-{
- const unsigned int pid = get_global_id(0);
-
- float4 particle = particles[pid];
-
-% if descriptor.d == 2:
- const unsigned int gid = floor(particle.y)*${memory.size_x} + floor(particle.x);
-% elif descriptor.d == 3:
- 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.x += moment.y;
- particle.y += moment.z;
-% if descriptor.d == 2:
- particle.w += min(particle.x, particle.y) * aging;
-% elif descriptor.d == 3:
- particle.z += moment.w;
- float dy = (particle.y-${geometry.size_y/2.0});
- float dz = (particle.z-${geometry.size_z/2.0});
- dy *= dy;
- dz *= dz;
- particle.w = 10.0*sqrt(moment.y*moment.y+moment.z*moment.z+moment.w*moment.w);
-% endif
- } else {
- particle.xyz = init_particles[pid].xyz;
- particle.w = 0.0;
- }
-
- particles[pid] = particle;
-}
diff --git a/template/opengl.mako b/template/opengl.mako
new file mode 100644
index 0000000..6b09f69
--- /dev/null
+++ b/template/opengl.mako
@@ -0,0 +1,100 @@
+<%
+def gid():
+ return {
+ 2: 'get_global_id(1)*%d + get_global_id(0)' % memory.size_x,
+ 3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (memory.size_x*memory.size_y, memory.size_x)
+ }.get(descriptor.d)
+
+def pop_offset(i):
+ return i * memory.volume
+
+def moments_cell():
+ return {
+ 2: '(int2)(get_global_id(0), get_global_id(1))',
+ 3: '(int4)(get_global_id(0), get_global_id(1), get_global_id(2), 0)'
+ }.get(descriptor.d)
+%>
+
+__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()};
+
+ __global __read_only ${float_type}* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+ const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+ 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);
+% 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)};
+% endif
+ } else {
+ data.x = 0.0;
+ data.y = 0.0;
+ data.z = 0.0;
+ data.w = -material[gid];
+ }
+
+ moments[gid] = data;
+}
+
+__kernel void collect_gl_moments_to_texture(__global __read_only ${float_type}* f,
+ __global __read_only int* material,
+% if descriptor.d == 2:
+ __write_only image2d_t moments)
+% elif descriptor.d == 3:
+ __write_only image3d_t moments)
+% endif
+{
+ const unsigned int gid = ${gid()};
+
+ __global __read_only ${float_type}* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+ const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+ float4 data;
+
+ if (material[gid] != 5) {
+% 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);
+% 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)};
+% endif
+ } else {
+ data.x = 0.0;
+ data.y = 0.0;
+ data.z = 0.0;
+ data.w = 1.0;
+ }
+
+ write_imagef(moments, ${moments_cell()}, data);
+}
diff --git a/template/particles.mako b/template/particles.mako
new file mode 100644
index 0000000..6235307
--- /dev/null
+++ b/template/particles.mako
@@ -0,0 +1,38 @@
+__kernel void update_particles(__global __read_only float4* moments,
+ __global __read_only int* material,
+ __global __write_only float4* particles,
+ __global __read_only float4* init_particles,
+ float aging)
+{
+ const unsigned int pid = get_global_id(0);
+
+ float4 particle = particles[pid];
+
+% if descriptor.d == 2:
+ const unsigned int gid = floor(particle.y)*${memory.size_x} + floor(particle.x);
+% elif descriptor.d == 3:
+ 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;
+% if descriptor.d == 2:
+ particle.w += min(particle.x, particle.y) * aging;
+% elif descriptor.d == 3:
+ particle.z += moment.w;
+ float dy = (particle.y-${geometry.size_y/2.0});
+ float dz = (particle.z-${geometry.size_z/2.0});
+ dy *= dy;
+ dz *= dz;
+ particle.w = 10.0*sqrt(moment.y*moment.y+moment.z*moment.z+moment.w*moment.w);
+% endif
+ } else {
+ particle.xyz = init_particles[pid].xyz;
+ particle.w = particle.w-1.0;
+ }
+
+ particles[pid] = particle;
+}