aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-29 10:56:41 +0100
committerAdrian Kummerlaender2019-10-29 10:56:41 +0100
commitdbd9a340a7809a770d52d10154712278431acdc3 (patch)
treeda5bd5d706ebd0607b55062d87ac6f6ebb5e3bf6 /boltzgen/kernel/template
parent15c0cba693646269c04f245ca52f405ddfdb4a07 (diff)
downloadboltzgen-dbd9a340a7809a770d52d10154712278431acdc3.tar
boltzgen-dbd9a340a7809a770d52d10154712278431acdc3.tar.gz
boltzgen-dbd9a340a7809a770d52d10154712278431acdc3.tar.bz2
boltzgen-dbd9a340a7809a770d52d10154712278431acdc3.tar.lz
boltzgen-dbd9a340a7809a770d52d10154712278431acdc3.tar.xz
boltzgen-dbd9a340a7809a770d52d10154712278431acdc3.tar.zst
boltzgen-dbd9a340a7809a770d52d10154712278431acdc3.zip
Unify AOS, SOA specific cell preshift between targets
SOA and AOS should not be target specific, neighbor offset calculation / bijection between gid and cell coordinates should be customizable.
Diffstat (limited to 'boltzgen/kernel/template')
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cl.mako4
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cpp.mako4
-rw-r--r--boltzgen/kernel/template/collect_moments.cl.mako2
-rw-r--r--boltzgen/kernel/template/collect_moments.cpp.mako2
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cl.mako4
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cpp.mako4
-rw-r--r--boltzgen/kernel/template/equilibrilize.cl.mako4
-rw-r--r--boltzgen/kernel/template/equilibrilize.cpp.mako4
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cl.mako4
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cpp.mako4
10 files changed, 18 insertions, 18 deletions
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cl.mako b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
index 0762a09..cb59cbd 100644
--- a/boltzgen/kernel/template/bounce_back_boundary.cl.mako
+++ b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
@@ -2,8 +2,8 @@ __kernel void bounce_back_boundary_gid(__global ${float_type}* f_next,
__global ${float_type}* f_prev,
unsigned int gid)
{
- __global ${float_type}* preshifted_f_next = f_next + gid;
- __global ${float_type}* preshifted_f_prev = f_prev + gid;
+ __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, c_i in enumerate(descriptor.c):
const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
index 06dd718..c7abd2a 100644
--- a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
+++ b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
@@ -2,8 +2,8 @@ void bounce_back_boundary( ${float_type}* f_next,
const ${float_type}* f_prev,
std::size_t gid)
{
- ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
- const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+ ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ const ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, c_i in enumerate(descriptor.c):
const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako
index 0ab42d1..06b92c9 100644
--- a/boltzgen/kernel/template/collect_moments.cl.mako
+++ b/boltzgen/kernel/template/collect_moments.cl.mako
@@ -2,7 +2,7 @@ __kernel void collect_moments_gid(__global ${float_type}* f,
__global ${float_type}* moments,
unsigned int gid)
{
- __global ${float_type}* preshifted_f = f + gid;
+ __global ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')};
% for i in range(0,descriptor.q):
const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
diff --git a/boltzgen/kernel/template/collect_moments.cpp.mako b/boltzgen/kernel/template/collect_moments.cpp.mako
index 8c37db2..d72ff5d 100644
--- a/boltzgen/kernel/template/collect_moments.cpp.mako
+++ b/boltzgen/kernel/template/collect_moments.cpp.mako
@@ -3,7 +3,7 @@ void collect_moments(const ${float_type}* f,
${float_type}& rho,
${float_type} u[${descriptor.d}])
{
- const ${float_type}* preshifted_f = f + gid*${layout.gid_offset()};
+ const ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')};
% for i in range(0,descriptor.q):
const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako
index a8fe532..c586884 100644
--- a/boltzgen/kernel/template/collide_and_stream.cl.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cl.mako
@@ -2,8 +2,8 @@ __kernel void collide_and_stream_gid(__global ${float_type}* f_next,
__global ${float_type}* f_prev,
unsigned int gid)
{
- __global ${float_type}* preshifted_f_next = f_next + gid;
- __global ${float_type}* preshifted_f_prev = f_prev + gid;
+ __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, c_i in enumerate(descriptor.c):
const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
diff --git a/boltzgen/kernel/template/collide_and_stream.cpp.mako b/boltzgen/kernel/template/collide_and_stream.cpp.mako
index 71c62b7..20fc529 100644
--- a/boltzgen/kernel/template/collide_and_stream.cpp.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cpp.mako
@@ -2,8 +2,8 @@ void collide_and_stream( ${float_type}* f_next,
const ${float_type}* f_prev,
std::size_t gid)
{
- ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
- const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+ ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ const ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, c_i in enumerate(descriptor.c):
const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako
index 4b9b984..4ee8d41 100644
--- a/boltzgen/kernel/template/equilibrilize.cl.mako
+++ b/boltzgen/kernel/template/equilibrilize.cl.mako
@@ -2,8 +2,8 @@ __kernel void equilibrilize_gid(__global ${float_type}* f_next,
__global ${float_type}* f_prev,
unsigned int gid)
{
- __global ${float_type}* preshifted_f_next = f_next + gid;
- __global ${float_type}* preshifted_f_prev = f_prev + gid;
+ __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, w_i in enumerate(descriptor.w):
preshifted_f_next[${layout.pop_offset(i)}] = ${w_i}.f;
diff --git a/boltzgen/kernel/template/equilibrilize.cpp.mako b/boltzgen/kernel/template/equilibrilize.cpp.mako
index 9f082a1..3b95a31 100644
--- a/boltzgen/kernel/template/equilibrilize.cpp.mako
+++ b/boltzgen/kernel/template/equilibrilize.cpp.mako
@@ -2,8 +2,8 @@ void equilibrilize(${float_type}* f_next,
${float_type}* f_prev,
std::size_t gid)
{
- ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
- ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+ ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, w_i in enumerate(descriptor.w):
preshifted_f_next[${layout.pop_offset(i)}] = ${w_i.evalf()};
diff --git a/boltzgen/kernel/template/momenta_boundary.cl.mako b/boltzgen/kernel/template/momenta_boundary.cl.mako
index e4a8ff3..f5edb61 100644
--- a/boltzgen/kernel/template/momenta_boundary.cl.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cl.mako
@@ -4,8 +4,8 @@ __kernel void ${name}_momenta_boundary_gid(
__global ${float_type}* f_prev,
unsigned int gid, ${param})
{
- __global ${float_type}* preshifted_f_next = f_next + gid;
- __global ${float_type}* preshifted_f_prev = f_prev + gid;
+ __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, c_i in enumerate(descriptor.c):
const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako
index 5475241..3f51087 100644
--- a/boltzgen/kernel/template/momenta_boundary.cpp.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako
@@ -4,8 +4,8 @@ void ${name}_momenta_boundary(
const ${float_type}* f_prev,
std::size_t gid, ${param})
{
- ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
- const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+ ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
+ const ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
% for i, c_i in enumerate(descriptor.c):
const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];