diff options
Diffstat (limited to 'template')
| -rw-r--r-- | template/kernel.mako | 109 | ||||
| -rw-r--r-- | template/opengl.mako | 103 | ||||
| -rw-r--r-- | template/particles.mako | 30 | ||||
| -rw-r--r-- | template/sdf.cl.mako | 54 | ||||
| -rw-r--r-- | template/sdf.lib.glsl.mako | 75 | ||||
| -rw-r--r-- | template/streamline.mako | 57 |
6 files changed, 340 insertions, 88 deletions
diff --git a/template/kernel.mako b/template/kernel.mako index 529eb30..2e26124 100644 --- a/template/kernel.mako +++ b/template/kernel.mako @@ -1,3 +1,11 @@ +% if float_type == 'double': +#if defined(cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +#elif defined(cl_amd_fp64) +#pragma OPENCL EXTENSION cl_amd_fp64 : enable +#endif +% endif + <% def gid(): return { @@ -9,13 +17,13 @@ def pop_offset(i): return i * memory.volume %> -__kernel void equilibrilize(__global __write_only ${float_type}* f_next, - __global __write_only ${float_type}* f_prev) +__kernel void equilibrilize(__global ${float_type}* f_next, + __global ${float_type}* f_prev) { const unsigned int gid = ${gid()}; - __global __write_only ${float_type}* preshifted_f_next = f_next + gid; - __global __write_only ${float_type}* preshifted_f_prev = f_prev + gid; + __global ${float_type}* preshifted_f_next = f_next + gid; + __global ${float_type}* preshifted_f_prev = f_prev + gid; % if pop_eq_src == '': % for i, w_i in enumerate(descriptor.w): @@ -36,9 +44,9 @@ def neighbor_offset(c_i): %> -__kernel void collide_and_stream(__global __write_only ${float_type}* f_next, - __global __read_only ${float_type}* f_prev, - __global __read_only int* material, +__kernel void collide_and_stream(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + __global int* material, unsigned int time) { const unsigned int gid = ${gid()}; @@ -49,8 +57,8 @@ __kernel void collide_and_stream(__global __write_only ${float_type}* f_next, return; } - __global __write_only ${float_type}* preshifted_f_next = f_next + gid; - __global __read_only ${float_type}* preshifted_f_prev = f_prev + gid; + __global ${float_type}* preshifted_f_next = f_next + gid; + __global ${float_type}* preshifted_f_prev = f_prev + gid; % for i, c_i in enumerate(descriptor.c): const ${float_type} f_curr_${i} = preshifted_f_prev[${pop_offset(i) + neighbor_offset(-c_i)}]; @@ -79,12 +87,12 @@ __kernel void collide_and_stream(__global __write_only ${float_type}* f_next, % endfor } -__kernel void collect_moments(__global __read_only ${float_type}* f, - __global __write_only ${float_type}* moments) +__kernel void collect_moments(__global ${float_type}* f, + __global ${float_type}* moments) { const unsigned int gid = ${gid()}; - __global __read_only ${float_type}* preshifted_f = f + gid; + __global ${float_type}* preshifted_f = f + gid; % for i in range(0,descriptor.q): const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}]; @@ -97,79 +105,4 @@ __kernel void collect_moments(__global __read_only ${float_type}* f, % for i, expr in enumerate(moments_assignment): 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.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; - particle.w += min(min(particle.x, particle.y), particle.z) * aging; -% endif - } else { - particle.xyz = init_particles[pid].xyz; - particle.w = particle.w-1.0; - } - - particles[pid] = particle; -} +}
\ No newline at end of file diff --git a/template/opengl.mako b/template/opengl.mako new file mode 100644 index 0000000..daf5f66 --- /dev/null +++ b/template/opengl.mako @@ -0,0 +1,103 @@ +<% +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) + +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_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()}; + + __global ${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]; + } + + write_imagef(moments, ${moments_cell()}, data); +} + +__kernel void collect_gl_moments_to_texture(__global ${float_type}* f, +% 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 ${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 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 + + write_imagef(moments, ${moments_cell()}, data); +} diff --git a/template/particles.mako b/template/particles.mako new file mode 100644 index 0000000..0b9a67e --- /dev/null +++ b/template/particles.mako @@ -0,0 +1,30 @@ +__kernel void update_particles(__global float* moments, + __global int* material, + __global float4* particles, + __global 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 + + if (material[gid] == 1 && particle.w < 1.0) { + particle.x += moments[${1*memory.volume}+gid]; + particle.y += moments[${2*memory.volume}+gid]; +% if descriptor.d == 3: + particle.z += moments[${3*memory.volume}+gid]; +% endif + particle.w += min(particle.x, particle.y) * aging; + } else { + particle.xyz = init_particles[pid].xyz; + particle.w = particle.w-1.0; + } + + particles[pid] = particle; +} diff --git a/template/sdf.cl.mako b/template/sdf.cl.mako new file mode 100644 index 0000000..0a62f4b --- /dev/null +++ b/template/sdf.cl.mako @@ -0,0 +1,54 @@ +typedef float3 vec3; +typedef float2 vec2; + +float3 v3(float x, float y, float z) { + return (float3)(x,y,z); +} + +float2 v2(float x, float y) { + return (float2)(x,y); +} + +__constant float3 center = (float3)(${geometry.size_x/2.5}, ${geometry.size_y/2}, ${geometry.size_z/2}); + +<%include file="sdf.lib.glsl.mako"/> + +${sdf_src} + +__kernel void setup_channel_with_sdf_obstacle(__global int* material) { + const unsigned x = get_global_id(0); + const unsigned y = get_global_id(1); + const unsigned z = get_global_id(2); + + const unsigned gid = z*${geometry.size_x*geometry.size_y} + y*${geometry.size_x} + x; + + if (x == 0 || x == ${geometry.size_x-1} || + y == 0 || y == ${geometry.size_y-1} || + z == 0 || z == ${geometry.size_z-1}) { + material[gid] = 0; + return; + } + + if (x == 1) { + material[gid] = 3; + return; + } + + if (x == ${geometry.size_x-2}) { + material[gid] = 4; + return; + } + + if (y == 1 || y == ${geometry.size_y-2} || + z == 1 || z == ${geometry.size_z-2}) { + material[gid] = 2; + return; + } + + if (sdf((float3)(x,y,z)) < 0.0) { + material[gid] = 2; + return; + } + + material[gid] = 1; +} diff --git a/template/sdf.lib.glsl.mako b/template/sdf.lib.glsl.mako new file mode 100644 index 0000000..c58cc7b --- /dev/null +++ b/template/sdf.lib.glsl.mako @@ -0,0 +1,75 @@ +float sphere(vec3 v, float r) { + return length(v) - r; +} + +float torus(vec3 v, vec2 t) { + vec2 q = v2(length(v.xz)-t.x,v.y); + return length(q)-t.y; +} + +float box(vec3 v, vec3 b) { + vec3 q = fabs(v) - b; + return length(max(q,0.0)) + min(max(q.x,max(q.y,q.z)),0.0); +} + +float cylinder(vec3 v, float r, float h) { + return max(length(v.xy) - r, fabs(v.z) - 0.5*h); +} + +vec3 flip_xy(vec3 v) { + return v3(v.y, v.x, v.z); +} + +vec3 flip_yz(vec3 v) { + return v3(v.x, v.z, v.y); +} + +vec3 rotate_x(vec3 v, float r) { + v.yz = cos(r)*v.yz + sin(r)*v2(v.z, -v.y); + return v; +} + +vec3 rotate_y(vec3 v, float r) { + v.xz = cos(r)*v.xz + sin(r)*v2(v.z, -v.x); + return v; +} + +vec3 rotate_z(vec3 v, float r) { + v.xy = cos(r)*v.xy + sin(r)*v2(v.y, -v.x); + return v; +} + +vec3 translate(vec3 v, vec3 w) { + return v - w; +} + +float rounded(float a, float r) { + return a - r; +} + +float sadd(float a, float b, float k) { + float h = clamp(0.5 + 0.5 * (b - a) / k, 0.0, 1.0); + return mix(b, a, h) - k * h * (1 - h); +} + +float ssub(float b, float a, float k) { + float h = clamp(0.5 - 0.5*(b+a)/k, 0.0, 1.0); + return mix(b, -a, h) + k*h*(1.0-h); +} + +float sintersect(float a, float b, float k) { + float h = clamp(0.5 - 0.5*(b-a)/k, 0.0, 1.0); + return mix(b, a, h) + k*h*(1.0-h); +} + +float sub(float a, float b) { + return max(-b, a); +} + +float add(float a, float b) { + return min(a, b); +} + +float intersect(float a, float b) { + return max(a, b); +} diff --git a/template/streamline.mako b/template/streamline.mako new file mode 100644 index 0000000..adf5efd --- /dev/null +++ b/template/streamline.mako @@ -0,0 +1,57 @@ +<% +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) +%> + +__kernel void dillute(__global int* material, + __read_write image2d_t streamlines) +{ + const unsigned int gid = ${gid()}; + const int2 pos = (int2)(get_global_id(0), get_global_id(1)); + + float4 color = read_imagef(streamlines, pos); + + if (material[gid] == 1) { + color.xyz *= 0.975; + } else { + color.xyz = 0.2; + } + + write_imagef(streamlines, pos, color); +} + +float3 blueRedPalette(float x) { + return mix( + (float3)(0.0, 0.0, 1.0), + (float3)(1.0, 0.0, 0.0), + x + ); +} + +__kernel void draw_streamline(__global float* moments, + __global int* material, + __global float2* origins, + __read_write image2d_t streamlines) +{ + float2 particle = origins[get_global_id(0)]; + + for (int i = 0; i < ${2*memory.size_x}; ++i) { + const unsigned int gid = round(particle.y)*${memory.size_x} + round(particle.x); + if (material[gid] != 1) { + break; + } + + particle.x += 0.5 * moments[${1*memory.volume}+gid] / 0.01; + particle.y += 0.5 * moments[${2*memory.volume}+gid] / 0.01; + + const int2 pos = (int2)(round(particle.x), round(particle.y)); + + float4 color = read_imagef(streamlines, pos); + color.xyz += 0.05; + + write_imagef(streamlines, pos, color); + } +}
\ No newline at end of file |