AccelerateHS/accelerate-cuda — pull request #5 by bmsherman (2013-11-07 13:48:34+00:00)
Title: Fixed issue with warp synchronization in foldAll kernel.

This change fixes the issue documented in https://github.com/AccelerateHS/accelerate/issues/123.

The issue was that accelerate-cuda was generating code that relied on __syncthreads() statements made inside conditionals. These conditionals sometimes evaluated differently for different threads in the same warp. In this case, __syncthreads() was not guaranteed to work.

I'm not a CUDA/C expert, so I can't say that this code is the cleanest or most efficient way to accomplish the desired change. But I have confirmed that it performs correctly on the test program described in https://github.com/AccelerateHS/accelerate/issues/123 with my computer setup.