Hi,
It is exciting to know that NVIDIA GPUs will finally have support for double precision. Since its a new feature, I was wondering what it means for performance of raw compute or memory operations. It would be great if someone could give numbers or explain what to theoretically expect.
I performed a quick experiment to see the effect of using doubles in clearly compute-bound and clearly memory-bound situations. I notice that in the first case, kernel execution time increases to around 7.6 times when all data (both host and device) is changed to doubles. In the memory bound case, the performance drop is a factor of around 3.3. I guess it would be better for real-life kernels.
I’m not sure what to infer from these numbers, could someone please shed some light at what would have been expected, and why? Here are the kernels I used:-
- Compute-bound: This kernel performs a series of dependent MADs, reading an input from one array and eventually writing to another.
__global__ void
gpu_cb(float myseed, float *myI, float *myO, unsigned int length)
{
float x, acc;
int idx=blockIdx.x*blockDim.x + threadIdx.x;
x=*(myI+idx);
acc=x*myseed;
#pragma unroll
for(int i=0;i<256;i++)
{
acc+=(acc*x);
}
*(myO+idx)=acc;
}
- Memory bound: This kernel performs a series of dependent (coalesced) memory operations. I can see from the ptx and the profiler that the equivalent number of memory operations do occur (compiler does not optimize any out). For this case the speed of the ‘add’ operation after every access might have interfered with performance measurement, especially because accesses are coalesced. What is a better way to design such a micro-benchmark?
__global__ void
gpu_mb(float myseed, float *myI, float *myO, unsigned int length)
{
float x, acc;
int idx=blockIdx.x*blockDim.x + threadIdx.x;
float *add;
add=myI+idx;
x=*add;
#pragma unroll
for(int i=0;i<256;i++)
{
add += 16;
x+=*add;
}
acc=x*myseed;
*(myO+idx)=acc;
}
Thanks for reading and I apologize for errors,
Anjul