The demotion happens because CUDA devices support double precision calculations at first with compute capability 1.3. NVCC knows the specifications and demotes every double to float for devices with CC < 1.3 just because the hardware cannot handle double precisions.
A good feature list could be found on wikipedia: CUDA
That you can’t see any doubles in this code doesn't mean that they are not there. Most commonly this error results from a missing f
postfix on a floating point constant. The compiler performance an implicit cast from all floats to double when one double is part of the expression. A floating point constant without the f
is a double value and the casting starts. However, for the less-operator a cast without constant expressions should not happen.
I can only speculate, but it seems to me that in your case a double precision value could be used within the thrust::sort implementation. Since you provide only a user function to a higher order function (functions that take functions as parameters).