1 #define N (1<<28)
2 float a[N], b[N], c[N], d[N];
3 
4 int main(void) {
5     for(unsigned int i=0; i<N; ++i)
6         d[i] = a[i]*b[i]+c[i];
7 
8     return 0;
9 }

 1 #include <cuda.h>
 2 
 3 #define N (1<<28)
 4 #define BLOCK_SIZE 128
 5 
 6 __device__ float a[N], b[N], c[N], d[N];
 7 
 8 __global__ void ma4(void) {
 9     unsigned int tid = blockIdx.x*blockDim.x + threadIdx.x;
10     d[tid] = a[tid]*b[tid]+c[tid];
11 }
12 
13 int main(void) {
14     ma4<<<N/BLOCK_SIZE, BLOCK_SIZE>>>();
15     cudaDeviceSynchronize();
16     return 0;
17 }

1 $ nvcc fma.cu -O3 -arch=sm_20 --ptxas-options=-v -o fmacuda
2 ptxas info    : Compiling entry function '_Z3fmav' for 'sm_20'
3 ptxas info    : Function properties for _Z3fmav
4     0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
5 ptxas info    : Used 13 registers, 32 bytes cmem[0], 32 bytes cmem[14]
6 $ ./fmacuda
7 $