From c543a20179d495cb6d585c48637c24b248c29a07 Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Sat, 15 Jun 2019 14:18:32 +0200 Subject: Consistently name population buffers --- implosion.py | 10 +++++----- template/kernel.mako | 24 ++++++++++++------------ 2 files changed, 17 insertions(+), 17 deletions(-) diff --git a/implosion.py b/implosion.py index 514af46..90b1d55 100644 --- a/implosion.py +++ b/implosion.py @@ -4,8 +4,8 @@ import time import matplotlib import matplotlib.pyplot as plt matplotlib.use('AGG') -from lbm import Lattice, Geometry +from lbm import Lattice, Geometry import symbolic.D2Q9 as D2Q9 @@ -34,13 +34,13 @@ pop_eq = """ if ( sqrt(pow(get_global_id(0) - ${geometry.size_x//2}.f, 2.f) + pow(get_global_id(1) - ${geometry.size_y//2}.f, 2.f)) < ${geometry.size_x//10} ) { % for i, w_i in enumerate(descriptor.w): - preshifted_f_a[${i*geometry.volume}] = 1./24.f; - preshifted_f_b[${i*geometry.volume}] = 1./24.f; + preshifted_f_next[${i*geometry.volume}] = 1./24.f; + preshifted_f_prev[${i*geometry.volume}] = 1./24.f; % endfor } else { % for i, w_i in enumerate(descriptor.w): - preshifted_f_a[${i*geometry.volume}] = ${w_i}.f; - preshifted_f_b[${i*geometry.volume}] = ${w_i}.f; + preshifted_f_next[${i*geometry.volume}] = ${w_i}.f; + preshifted_f_prev[${i*geometry.volume}] = ${w_i}.f; % endfor }""" diff --git a/template/kernel.mako b/template/kernel.mako index f3de3da..7e930af 100644 --- a/template/kernel.mako +++ b/template/kernel.mako @@ -1,15 +1,15 @@ -__kernel void equilibrilize(__global __write_only float* f_a, - __global __write_only float* f_b) +__kernel void equilibrilize(__global __write_only float* f_next, + __global __write_only float* f_prev) { const unsigned int gid = get_global_id(1)*${geometry.size_x} + get_global_id(0); - __global __write_only float* preshifted_f_a = f_a + gid; - __global __write_only float* preshifted_f_b = f_b + gid; + __global __write_only float* preshifted_f_next = f_next + gid; + __global __write_only float* preshifted_f_prev = f_prev + gid; % if pop_eq_src == '': % for i, w_i in enumerate(descriptor.w): - preshifted_f_a[${i*geometry.volume}] = ${w_i}.f; - preshifted_f_b[${i*geometry.volume}] = ${w_i}.f; + preshifted_f_next[${i*geometry.volume}] = ${w_i}.f; + preshifted_f_prev[${i*geometry.volume}] = ${w_i}.f; % endfor % else: ${pop_eq_src} @@ -27,8 +27,8 @@ def neighbor_offset(c_i): return c_i[1]*geometry.size_x + c_i[0] %> -__kernel void collide_and_stream(__global __write_only float* f_a, - __global __read_only float* f_b, +__kernel void collide_and_stream(__global __write_only float* f_next, + __global __read_only float* f_prev, __global __read_only int* material) { const unsigned int gid = get_global_id(1)*${geometry.size_x} + get_global_id(0); @@ -39,11 +39,11 @@ __kernel void collide_and_stream(__global __write_only float* f_a, return; } - __global __write_only float* preshifted_f_a = f_a + gid; - __global __read_only float* preshifted_f_b = f_b + gid; + __global __write_only float* preshifted_f_next = f_next + gid; + __global __read_only float* preshifted_f_prev = f_prev + gid; % for i, c_i in enumerate(descriptor.c): - const float f_curr_${i} = preshifted_f_b[${direction_index(c_i)*geometry.volume + neighbor_offset(-c_i)}]; + const float f_curr_${i} = preshifted_f_prev[${direction_index(c_i)*geometry.volume + neighbor_offset(-c_i)}]; % endfor % for i, expr in enumerate(moments_subexpr): @@ -65,7 +65,7 @@ __kernel void collide_and_stream(__global __write_only float* f_a, % endfor % for i in range(0,descriptor.q): - preshifted_f_a[${i*geometry.volume}] = f_next_${i}; + preshifted_f_next[${i*geometry.volume}] = f_next_${i}; % endfor } -- cgit v1.2.3