From a90f1faedac705242a674ba94ea7bc5438cab078 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Sun, 2 Feb 2020 22:43:41 +0100
Subject: Implement basic multi-cuboid communication for CUDA target

---
 boltzgen/kernel/template/communicate.cuda.mako | 49 ++++++++++++++++++++++++++
 1 file changed, 49 insertions(+)
 create mode 100644 boltzgen/kernel/template/communicate.cuda.mako

(limited to 'boltzgen/kernel/template')

diff --git a/boltzgen/kernel/template/communicate.cuda.mako b/boltzgen/kernel/template/communicate.cuda.mako
new file mode 100644
index 0000000..ebc2a21
--- /dev/null
+++ b/boltzgen/kernel/template/communicate.cuda.mako
@@ -0,0 +1,49 @@
+<%
+import numpy
+
+directions = list(filter(lambda c: numpy.sum(numpy.abs(c)) == 1, descriptor.c))
+
+def name(direction):
+    names = [ ('west','','east'), ('south','','north'), ('down','','up') ]
+    return ''.join([ names[i][v+1] for i, v in enumerate(direction) ])
+
+def neighbors(direction):
+    j = list(direction).index(next(x for x in direction if x != 0))
+    return [ (i,c_i) for i, c_i in enumerate(descriptor.c) if c_i[j] == direction[j] ]
+%>
+
+% for direction in directions:
+__global__ void communicate_${name(direction)}(
+      ${float_type}** f_a
+    , ${float_type}** f_b
+    , std::size_t* cells_a
+    , std::size_t* cells_b
+    , std::size_t  cell_count
+) {
+    const std::size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+    if (!(index < cell_count)) {
+        return;
+    }
+    const std::size_t gid_a = cells_a[index];
+    const std::size_t gid_b = cells_b[index];
+
+% for i, c_i in neighbors(direction):
+    ${float_type}* preshifted_f_a_${i} = f_a[${i}] + ${layout.cell_preshift('gid_a')};
+    ${float_type}* preshifted_f_b_${i} = f_b[${i}] + ${layout.cell_preshift('gid_b')};
+% endfor
+
+% for i, c_i in neighbors(direction):
+    *preshifted_f_b_${i} = *preshifted_f_a_${i};
+% endfor
+
+% for i, c_i in neighbors(-direction):
+    ${float_type}* preshifted_f_a_${i} = f_a[${i}] + ${layout.cell_preshift('gid_a')};
+    ${float_type}* preshifted_f_b_${i} = f_b[${i}] + ${layout.cell_preshift('gid_b')};
+% endfor
+
+% for i, c_i in neighbors(-direction):
+    *preshifted_f_a_${i} = *preshifted_f_b_${i};
+% endfor
+}
+
+% endfor
-- 
cgit v1.2.3