Unexpected behavior using doubles in a kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

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)

2

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;

}