From 536b25e2c2b742c17d023d2b3386fed4dc60a339 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Sun, 27 Oct 2019 22:22:24 +0100
Subject: Optionally generate cell-list-based OpenCL dispatch functions

Requires different function naming as OpenCL 1.2 doesn't support overloads.

The OpenCL kernel code generated using this commit was successfully tested
on an actual GPU. Time to set up some automatic validation.
---
 boltzgen/kernel/template/collide_and_stream.cl.mako | 15 ++++++++++++---
 1 file changed, 12 insertions(+), 3 deletions(-)

(limited to 'boltzgen/kernel/template/collide_and_stream.cl.mako')

diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako
index 28cfa57..a8fe532 100644
--- a/boltzgen/kernel/template/collide_and_stream.cl.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cl.mako
@@ -1,6 +1,6 @@
-__kernel void collide_and_stream(__global ${float_type}* f_next,
-                                 __global ${float_type}* f_prev,
-                                 unsigned int gid)
+__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;
@@ -29,3 +29,12 @@ __kernel void collide_and_stream(__global ${float_type}* f_next,
     preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
 % endfor
 }
+
+% if 'cell_list_dispatch' in extras:
+__kernel void collide_and_stream_cells(__global ${float_type}* f_next,
+                                       __global ${float_type}* f_prev,
+                                       __global unsigned int*  cells)
+{
+    collide_and_stream_gid(f_next, f_prev, cells[get_global_id(0)]);
+}
+% endif
-- 
cgit v1.2.3