aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-06-11 20:30:08 +0200
committerAdrian Kummerlaender2019-06-11 20:30:08 +0200
commit81b9c7d6541c3c377bfb2332e21ecd71028e5666 (patch)
tree58b5c9d91e56fdf02a88c4ce560d6ac52ad9715b
parent3660fbbd71c8579b60c1e062f9c1d288253c0d04 (diff)
downloadsymlbm_playground-81b9c7d6541c3c377bfb2332e21ecd71028e5666.tar
symlbm_playground-81b9c7d6541c3c377bfb2332e21ecd71028e5666.tar.gz
symlbm_playground-81b9c7d6541c3c377bfb2332e21ecd71028e5666.tar.bz2
symlbm_playground-81b9c7d6541c3c377bfb2332e21ecd71028e5666.tar.lz
symlbm_playground-81b9c7d6541c3c377bfb2332e21ecd71028e5666.tar.xz
symlbm_playground-81b9c7d6541c3c377bfb2332e21ecd71028e5666.tar.zst
symlbm_playground-81b9c7d6541c3c377bfb2332e21ecd71028e5666.zip
Templatize assignment loops
-rw-r--r--codegen_lbm.py47
1 files changed, 13 insertions, 34 deletions
diff --git a/codegen_lbm.py b/codegen_lbm.py
index a115ef1..c1ccd00 100644
--- a/codegen_lbm.py
+++ b/codegen_lbm.py
@@ -48,6 +48,8 @@ collide = [ Assignment(f_next[i], f_curr[i] + 1/tau * ( f_eq_i - f_curr[i] )) fo
collide_opt = cse(collide, optimizations='basic')
kernel = """
+__constant float tau = ${tau};
+
unsigned int indexOfDirection(int i, int j) {
return (i+1) + 3*(1-j);
}
@@ -79,17 +81,9 @@ __kernel void collide_and_stream(__global __write_only float* f_a,
return;
}
- const float f_curr_0 = f_i(f_b, cell.x+1, cell.y-1, -1, 1);
- const float f_curr_1 = f_i(f_b, cell.x , cell.y-1, 0, 1);
- const float f_curr_2 = f_i(f_b, cell.x-1, cell.y-1, 1, 1);
- const float f_curr_3 = f_i(f_b, cell.x+1, cell.y , -1, 0);
- const float f_curr_4 = f_i(f_b, cell.x , cell.y , 0, 0);
- const float f_curr_5 = f_i(f_b, cell.x-1, cell.y , 1, 0);
- const float f_curr_6 = f_i(f_b, cell.x+1, cell.y+1, -1,-1);
- const float f_curr_7 = f_i(f_b, cell.x , cell.y+1, 0,-1);
- const float f_curr_8 = f_i(f_b, cell.x-1, cell.y+1, 1,-1);
-
- const float tau = ${tau};
+% for i, c_i in enumerate(c):
+ const float f_curr_${i} = f_i(f_b, cell.x-(${c_i[0]}), cell.y-(${c_i[1]}), ${c_i[0]}, ${c_i[1]});
+% endfor
% for i, expr in enumerate(moments_helper):
const float ${expr[0]} = ${ccode(expr[1])};
@@ -112,15 +106,9 @@ __kernel void collide_and_stream(__global __write_only float* f_a,
const float ${ccode(expr)}
% endfor
- f_a[0*${nCells} + gid] = f_next_0;
- f_a[1*${nCells} + gid] = f_next_1;
- f_a[2*${nCells} + gid] = f_next_2;
- f_a[3*${nCells} + gid] = f_next_3;
- f_a[4*${nCells} + gid] = f_next_4;
- f_a[5*${nCells} + gid] = f_next_5;
- f_a[6*${nCells} + gid] = f_next_6;
- f_a[7*${nCells} + gid] = f_next_7;
- f_a[8*${nCells} + gid] = f_next_8;
+% for i in range(0,len(c)):
+ f_a[${i*nCells} + gid] = f_next_${i};
+% endfor
}
__kernel void collect_moments(__global __read_only float* f,
@@ -130,27 +118,17 @@ __kernel void collect_moments(__global __read_only float* f,
const uint2 cell = (uint2)(get_global_id(0), get_global_id(1));
- const float f_curr_0 = f[0*${nCells} + gid];
- const float f_curr_1 = f[1*${nCells} + gid];
- const float f_curr_2 = f[2*${nCells} + gid];
- const float f_curr_3 = f[3*${nCells} + gid];
- const float f_curr_4 = f[4*${nCells} + gid];
- const float f_curr_5 = f[5*${nCells} + gid];
- const float f_curr_6 = f[6*${nCells} + gid];
- const float f_curr_7 = f[7*${nCells} + gid];
- const float f_curr_8 = f[8*${nCells} + gid];
+% for i in range(0,len(c)):
+ const float f_curr_${i} = f[${i*nCells} + gid];
+% endfor
% for i, expr in enumerate(moments_helper):
const float ${expr[0]} = ${ccode(expr[1])};
% endfor
% for i, expr in enumerate(moments_assignment):
- const float ${ccode(expr)}
+ moments[${i*nCells} + gid] = ${ccode(expr.rhs)};
% endfor
-
- moments[0*${nCells} + gid] = rho;
- moments[1*${nCells} + gid] = u_0;
- moments[2*${nCells} + gid] = u_1;
}"""
@@ -230,6 +208,7 @@ class D2Q9_BGK_Lattice:
moments_assignment = moments_opt[1],
collide_helper = collide_opt[0],
collide_assignment = collide_opt[1],
+ c = c,
ccode = ccode
)).build()