Error with Doubles when using 1.1 Double is not supported. Demoting to float

Good afternoon,

I’m producing a ray-tracing application using CUDA as part of a university project. The target platform is Nvidia 9900 GPUs which I believe have the 1.1 architecture.
Having done a bit of googling I understand that there are some issues with cuda pre-1.3 and double’s so I removed them from my code. However, when I compile I get the following error:

ptxas ./nvccoutput/traverser.ptx, line 136; warning : Double is not supported. Demoting to float

The file has no ‘double’ variables or functions and does not call any. My instinct tells me it something to do with the autogenerated code as this seems to include lots of doubles. Other fixes have suggested using the nvcc flag -arch sm_13 but my target hardware is sm_11 so this isnt a fix.

Please can you tell me how I can solve this problem? :thanks:

// Traverser.c
#include “constants.h”
#include “world.cuh”
#include “ray.cuh”
#include “cuda_mathlib.cuh”
#include “camera.cuh”
#include “voxel_translation.cuh”
#include “traverser.cuh”

device void TRAVERSER_GetVec(int i, int j, Camera camera, vec3_t topleft,
vec3_t output) {
vec3_t w, h;
CUDAMATH_VectorScale(camera->wdir, (float) SCREENW
(float) i/(float) PIXELW, w);
CUDAMATH_VectorScale(camera->hdir, (float) SCREENH*(float) j/(float) PIXELH, h);
CUDAMATH_VectorAdd(w, h, w);
CUDAMATH_VectorAdd(w, topleft, output);

device void TRAVERSER_ConstructRay(int i, int j, Camera *camera,
vec3_t topleftloc, Ray *output) {
vec3_t dir, loc;
TRAVERSER_GetVec(i, j, camera, topleftloc, loc);
CUDAMATH_VectorSubtract(loc, camera->loc, dir);
CUDAMATH_VectorNormalize(dir, dir);
InitRay(loc, dir, output);

global void TRAVERSER_RayTrace(VOXEL* dev_world, Camera camera, vec3_t
topleftloc, int
dev_screen) {
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
Ray ray;
int ispresent = 0;
TRAVERSER_ConstructRay(x, y, camera, topleftloc, &ray);
while(IsValidVoxel(ray.X, ray.Y, ray.Z) && ispresent == 0) {
if (ObjectPresent(ray.X, ray.Y, ray.Z, dev_world)) {
ispresent = 1;
if (ispresent == 0) {
dev_screen[x + yblockDim.xgridDim.x] = 0;
} else {
dev_screen[x + yblockDim.xgridDim.x] = 1;

// From (–keep) output code
grep ‘double’ traverser.*
traverser.cpp1.ii: We put this #include_next outside the double inclusion check because
traverser.cpp1.ii:struct double1
traverser.cpp1.ii: double x;
traverser.cpp1.ii:struct attribute((aligned(16))) double2
traverser.cpp1.ii: double x, y;
traverser.cpp1.ii:struct double3
traverser.cpp1.ii: double x, y, z;
traverser.cpp1.ii:struct attribute((aligned(16))) double4
traverser.cpp1.ii: double x, y, z, w;
traverser.cpp1.ii:typedef struct double1 double1;
traverser.cpp1.ii:typedef struct double2 double2;
traverser.cpp1.ii:typedef struct double3 double3;
traverser.cpp1.ii:typedef struct double4 double4;…

Other details:

  • Ubuntu 10.10 64bit
  • nvcc and toolkit from 1st December 2010
  • Test platform: 9800 GTX+
  • Target platform: 9900 (i believe both use cuda 1.1)
  • Compile options:
    nvcc -arch sm_11 --keep -keep-dir ./nvccoutput

only contains a kernel though

Double arithmetic can be implicitly introduced by using floating point constants (which are double by default) or intrinsic functions. Both can be turned into explicit single arithmetic by appending an [font=“Courier New”]f[/font].
E.g., [font=“Courier New”]1.2[/font] → [font=“Courier New”]1.2f[/font], [font=“Courier New”]sin(x)[/font] → [font=“Courier New”]sinf(x)[/font].

Cheers thankyou that worked :)

got another error but its completely different (‘unspecified launch failure’).