diff options
Diffstat (limited to 'template/opengl.mako')
| -rw-r--r-- | template/opengl.mako | 28 |
1 files changed, 15 insertions, 13 deletions
diff --git a/template/opengl.mako b/template/opengl.mako index 47264e8..daf5f66 100644 --- a/template/opengl.mako +++ b/template/opengl.mako @@ -13,11 +13,22 @@ def moments_cell(): 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) + +def neighbor_offset(c_i): + return { + 2: lambda: c_i[1]*memory.size_x + c_i[0], + 3: lambda: c_i[2]*memory.size_x*memory.size_y + c_i[1]*memory.size_x + c_i[0] + }.get(descriptor.d)() + %> -__kernel void collect_gl_moments(__global ${float_type}* f, - __global int* material, - __global float4* moments) +__kernel void collect_gl_moments_and_materials_to_texture(__global ${float_type}* f, + __global 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()}; @@ -52,18 +63,9 @@ __kernel void collect_gl_moments(__global ${float_type}* f, data.w = -material[gid]; } - moments[gid] = data; + write_imagef(moments, ${moments_cell()}, data); } -<% -def neighbor_offset(c_i): - return { - 2: lambda: c_i[1]*memory.size_x + c_i[0], - 3: lambda: c_i[2]*memory.size_x*memory.size_y + c_i[1]*memory.size_x + c_i[0] - }.get(descriptor.d)() - -%> - __kernel void collect_gl_moments_to_texture(__global ${float_type}* f, % if descriptor.d == 2: __write_only image2d_t moments) |