Unexpected behavior using doubles in a kernel

Hi,

the Programming Guide 1.0 mentions on page 42 that double types will be demoted to float by default, if there is no native double support.

I have tried this in my kernel using a typedef, but it didn’t work as expected.

The typedef is declared in the header “config.h”

typedef double ColorType;

typedef unsigned char ImgType;

The kernel gets a ColorType passed by parameter

#include "config.h"

__global__ void simpleKernel(ImgType* image, int width,ColorType a)

{

	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;

	unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

	image[(x+y*width)*3] = (ImgType)((a+a)*x);

	image[(x+y*width)*3+1] = (ImgType)((a*2.0)*y);

	image[(x+y*width)*3+2] =  0;

}

and will be called like this:

ColorType a = 0.1;

dim3 block(8, 8, 1);

dim3 grid(ImageWidth/8, ImageHeight/8, 1);

simpleKernel<<<grid,block>>>(imageDevice,ImageWidth,a);

I get different results when switching the ColorType typedef from double to float. Is this a bug?

I’m using CUDA 1.0 on Ubuntu Linux 7.04 with a 8600GT 256MB card.

I’ve attached my code, it would be nice if anybody could test it.

Thanks.
doubleTest.tar.gz (2.16 KB)

It seems to me that double parameters will be converted to int and not to float.

If i add a ColorType variable directly to the kernel, float and double will produce the same results.

__global__ void simpleKernel(ImgType* image, int width)

{

	ColorType a = 0.1;

	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;

	unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

	image[(x+y*width)*3] = (ImgType)((a+a)*x);

	image[(x+y*width)*3+1] = (ImgType)((a*2.0)*y);

	image[(x+y*width)*3+2] =  0;

}