diff options
Implement basic CUDA target
Currently only for the SSS streaming pattern.
CudaCodePrinter in `utility/printer.py` is required to add a 'f' suffix
to all single precision floating point literals. If this is not done
(when targeting single precision) most calculations happen in double
precision which destroys performance. (In OpenCL this is not necessary
as we can simply set the `-cl-single-precision-constant` flag. Sadly
such a flag doesn't seem to exist for nvcc.)
Diffstat (limited to 'boltzgen/kernel/template/update_sss_control_structure.cuda.mako')
-rw-r--r-- | boltzgen/kernel/template/update_sss_control_structure.cuda.mako | 21 |
1 files changed, 21 insertions, 0 deletions
diff --git a/boltzgen/kernel/template/update_sss_control_structure.cuda.mako b/boltzgen/kernel/template/update_sss_control_structure.cuda.mako new file mode 100644 index 0000000..2a98b94 --- /dev/null +++ b/boltzgen/kernel/template/update_sss_control_structure.cuda.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) +%> + +__global__ void init_sss_control_structure(${float_type}* f, ${float_type}** control) { +% for i, c_i in enumerate(descriptor.c): + control[${i}] = f + ${padding + layout.pop_offset(i, 2*padding)}; +% endfor +} + +__global__ void update_sss_control_structure(${float_type}** f) { +% for i, c_i in enumerate(descriptor.c): + ${float_type}* f_old_${i} = f[${i}]; +% endfor +% for i, c_i in enumerate(descriptor.c): + f[${i}] = f_old_${descriptor.c.index(-c_i)} + ${layout.neighbor_offset(-c_i)}; +% endfor +} |