Annotate CUDA kernels with __launch_bounds__ and launder the (statically-known) block size down to the launch site.

Fixes #14