ptxas "double is not supported" warning when using thrust::sort on a struct array

Question

I'm trying to sort an array of structs on my GPU with thrust::sort. However, when I compile with nvcc, I get this warning:

ptxas /tmp/tmpxft_00005186_00000000-5_antsim.ptx, line 1520; warning : Double is not supported. Demoting to float

I've isolated the problem to my call to thrust::sort, here:

thrust::sort(thrustAnts, thrustAnts + NUM_ANTS, antSortByX());

thrustAnts is an array of Ant structs located on the GPU, while antSortByX is a functor as defined below:

typedef struct {
   float posX;
   float posY;
   float direction;
   float speed;
   u_char life;
   u_char carrying;
   curandState rngState;
} Ant;

struct antSortByX {
   __host__ __device__ bool operator()(Ant &antOne, Ant &antTwo) {
      return antOne.posX < antTwo.posX;
   }
};

It seems to me as though there aren't any doubles in this, though I'm suspicious the less-than operator in my functor evaluates those floats as doubles. I can solve this problem by compiling with -arch sm_13, but I'm curious as to why this is complaining at me in the first place.

Answer 1

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).