aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2020-01-11 00:11:28 +0100
committerAdrian Kummerlaender2020-01-11 00:15:08 +0100
commitb5a24f31871d900342a3c47398cc75e22bad0b6f (patch)
tree5fa2889803d28f21536cbf964eca3850f7a66700
parentaa509dd4ebbb9d1d8ad6ebfe05111228fd9ae7c0 (diff)
downloadboltzgen-b5a24f31871d900342a3c47398cc75e22bad0b6f.tar
boltzgen-b5a24f31871d900342a3c47398cc75e22bad0b6f.tar.gz
boltzgen-b5a24f31871d900342a3c47398cc75e22bad0b6f.tar.bz2
boltzgen-b5a24f31871d900342a3c47398cc75e22bad0b6f.tar.lz
boltzgen-b5a24f31871d900342a3c47398cc75e22bad0b6f.tar.xz
boltzgen-b5a24f31871d900342a3c47398cc75e22bad0b6f.tar.zst
boltzgen-b5a24f31871d900342a3c47398cc75e22bad0b6f.zip
Implement SSS for OpenCL target
Sadly OpenCL kernels don't accept pointer-to-pointer arguments which complicates the control structure implementation. A workaround is to cast them into `uintptr_t` which is guaranteed to be large enough to fit any pointer on the device. Special care has to be taken to always perform the pointer shifts on actual floating point pointers and not on type-less pointers.
-rw-r--r--boltzgen/kernel/template/pattern/SSS.cl.mako121
-rw-r--r--boltzgen/kernel/template/update_sss_control_structure.cl.mako21
2 files changed, 142 insertions, 0 deletions
diff --git a/boltzgen/kernel/template/pattern/SSS.cl.mako b/boltzgen/kernel/template/pattern/SSS.cl.mako
new file mode 100644
index 0000000..ad5e854
--- /dev/null
+++ b/boltzgen/kernel/template/pattern/SSS.cl.mako
@@ -0,0 +1,121 @@
+<%def name="operator(name, params = None)">
+<%
+if layout.__class__.__name__ != 'SOA':
+ raise Exception('SSS pattern only works for the AOS memory layout')
+%>
+__kernel void ${name}(
+ __global uintptr_t* control
+% if 'cell_list_dispatch' in extras:
+ , __global unsigned* cells
+% else:
+ , std::size_t gid
+% endif
+% if params is not None:
+% for param_type, param_name in params:
+ , ${param_type} ${param_name}
+% endfor
+% endif
+) {
+% if 'cell_list_dispatch' in extras:
+ const unsigned int gid = cells[get_global_id(0)];
+% endif
+
+% for i, c_i in enumerate(descriptor.c):
+ __global ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')};
+% endfor
+
+% for i, c_i in enumerate(descriptor.c):
+ const ${float_type} f_curr_${i} = *preshifted_f_${i};
+% endfor
+
+ ${caller.body()}
+
+% for i, c_i in enumerate(descriptor.c):
+ *preshifted_f_${i} = f_next_${descriptor.c.index(-c_i)};
+% endfor
+}
+</%def>
+
+<%def name="operator_with_domain_dispatch(name, params = None)">
+__kernel void ${name}(
+ __global uintptr_t* control
+% if params is not None:
+% for param_type, param_name in params:
+ , ${param_type} ${param_name}
+% endfor
+% endif
+) {
+ const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)', 'get_global_id(2)')};
+
+% for i, c_i in enumerate(descriptor.c):
+ __global ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')};
+% endfor
+
+% for i, c_i in enumerate(descriptor.c):
+ const ${float_type} f_curr_${i} = *preshifted_f_${i};
+% endfor
+
+ ${caller.body()}
+
+% for i, c_i in enumerate(descriptor.c):
+ *preshifted_f_${i} = f_next_${descriptor.c.index(-c_i)};
+% endfor
+}
+</%def>
+
+<%def name="functor(name, params = None)">
+<%
+if layout.__class__.__name__ != 'SOA':
+ raise Exception('SSS pattern only works for the AOS memory layout')
+%>
+__kernel void ${name}(
+ __global const uintptr_t* control
+% if 'cell_list_dispatch' in extras:
+ , __global unsigned* cells
+% else:
+ , std::size_t gid
+% endif
+% if params is not None:
+% for param_type, param_name in params:
+ , ${param_type} ${param_name}
+% endfor
+% endif
+) {
+% if 'cell_list_dispatch' in extras:
+ const unsigned int gid = cells[get_global_id(0)];
+% endif
+
+% for i, c_i in enumerate(descriptor.c):
+ __global const ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')};
+% endfor
+
+% for i, c_i in enumerate(descriptor.c):
+ const ${float_type} f_curr_${i} = *preshifted_f_${descriptor.c.index(-c_i)};
+% endfor
+
+ ${caller.body()}
+}
+</%def>
+
+<%def name="functor_with_domain_dispatch(name, params = None)">
+__kernel void ${name}(
+ __global const uintptr_t* control
+% if params is not None:
+% for param_type, param_name in params:
+ , ${param_type} ${param_name}
+% endfor
+% endif
+) {
+ const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)', 'get_global_id(2)')};
+
+% for i, c_i in enumerate(descriptor.c):
+ __global const ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')};
+% endfor
+
+% for i, c_i in enumerate(descriptor.c):
+ const ${float_type} f_curr_${i} = *preshifted_f_${i};
+% endfor
+
+ ${caller.body()}
+}
+</%def>
diff --git a/boltzgen/kernel/template/update_sss_control_structure.cl.mako b/boltzgen/kernel/template/update_sss_control_structure.cl.mako
new file mode 100644
index 0000000..ec6a5d5
--- /dev/null
+++ b/boltzgen/kernel/template/update_sss_control_structure.cl.mako
@@ -0,0 +1,21 @@
+<%
+if streaming != 'SSS':
+ raise Exception('"update_sss_control_structure" function only makes sense for the SSS pattern')
+
+padding = (max(geometry.size_x,geometry.size_y,geometry.size_z)+1)**(descriptor.d-1)
+%>
+
+__kernel void init_sss_control_structure(__global ${float_type}* f, __global uintptr_t* control) {
+% for i, c_i in enumerate(descriptor.c):
+ control[${i}] = (uintptr_t)(f + ${padding + layout.pop_offset(i, 2*padding)});
+% endfor
+}
+
+__kernel void update_sss_control_structure(__global uintptr_t* control) {
+% for i, c_i in enumerate(descriptor.c):
+ __global ${float_type}* f_old_${i} = (__global ${float_type}*)(control[${i}]);
+% endfor
+% for i, c_i in enumerate(descriptor.c):
+ control[${i}] = (uintptr_t)(f_old_${descriptor.c.index(-c_i)} + ${layout.neighbor_offset(-c_i)});
+% endfor
+}