diff options
Diffstat (limited to 'template/kernel.mako')
| -rw-r--r-- | template/kernel.mako | 41 |
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; +} |