double precision in Kernel how to use double precision in Kernel

[Quoc_Vinh]

Hi every body.

I am using Geforce 8800GT. I know that the compute capacity of this card is 1.1 and it does not support double precision.

I was trying to find how to use double precision with this card in this forum, but I could not get an expected answer.

if any one know how to use double precision in CUDA2.0 with geforce8800GT, please help me.

for example:

[codebox]global void Kernel(double *globalData, double value)

{

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

globalData[tx] += value;

}[/codebox]

Thank!

[Ailleur]

As you have stated yourself, it does not support double precision. Does not!
You wont find an answer on how to do it, as it does not support it!

The only thing you can do is use double single precision, and afaik, youre on your own to do that.

[e.ping]

On cards with compute capability less than 1.3, the CUDA compiler automatically demotes doubles to floats for you. Which is why the programming guide (and the NVIDIA folks in technical presentations) always emphasize that you must use single precision literals (0.0f instead of 0.0) if you want portable single precision code: Otherwise, you implicitly rely on the demoted types, and your single precision code on DP-HW will suddenly run much slower, because it actually uses doubles for computation (and stores results again in floats, obviously).

Someone posted a neat emulated precision template header recently, you might want to search the forums for Bailey or dsfun90. This gives you an s46e8 type, which is about halfway between a float and a double.

[Quoc_Vinh]

On cards with compute capability less than 1.3, the CUDA compiler automatically demotes doubles to floats for you. Which is why the programming guide (and the NVIDIA folks in technical presentations) always emphasize that you must use single precision literals (0.0f instead of 0.0) if you want portable single precision code: Otherwise, you implicitly rely on the demoted types, and your single precision code on DP-HW will suddenly run much slower, because it actually uses doubles for computation (and stores results again in floats, obviously).

Someone posted a neat emulated precision template header recently, you might want to search the forums for Bailey or dsfun90. This gives you an s46e8 type, which is about halfway between a float and a double.

Thank you.

I made reference to this page http://forums.nvidia.com/index.php?showtopic=73067.

[tmurray]

Just so you know, double* does not behave properly on pre-1.3 hardware. It reads four bytes instead of eight and doesn’t do a double → float cast, so it just doesn’t work at all.