CUDA上add.f32指令的执行周期到底是多少？

这两天，受网友cuda2010的提醒，发现了一个令人迷惑的问题，即CUDA上add.f32指令的执行周期到底是多少．
从cuda手册以及各类分析文章上说，cuda上的单精度浮点加法指令的执行周期是4周期．

 

即如果程序都是串行执行单精度浮点加法计算的话,性能应该是频率/4fps.

 

为验证上面的认识是否正确，尝试运行了一个简单的程序，发现结果并非如此.下面是本次试验的情况．

 

试验平台：
GPU NV 130M 4SM 1.5GHz(sp频率)
OS visita 32
CUDA SDK 3.0 + VS2008

 

为了保证是串行执行，在启动kernel时使用的参数是<<<1,1>>>.如下(重复执行10次)：
for(int ii=0; ii<10; ii++) test1<<<1,1>>>(dres);

 

而kernel程序test1如下:(修改自cuda2010的程序)
#define K    1125000
#define FLOP    2
#define TYPE    float
__global__ void test1(TYPE *res) {
    int inx=blockDim.x*blockIdx.x+threadIdx.x;
    TYPE s=(TYPE)1e-8f, d=(TYPE)0.0f;
#pragma unroll 3750
    for(int i=0; i<K; i++) {
        s+=d;
        d+=s;
    }
    res[inx]=s+d;
}
编译成ptx后指令序列如下：
.entry _Z5test1Pf (
 .param .u32 __cudaparm__Z5test1Pf_res)
{
.reg .u16 %rh<4>;
.reg .u32 %r<10>;
.reg .f32 %f<5>;
.reg .pred %p<3>;
.loc 2 177 0
$LBB1__Z5test1Pf:
.loc 2 24 0
mov.s32  %r1, 0;
mov.f32  %f1, 0f00000000;      // 0
mov.f32  %f2, 0f322bcc77;      // 1e-008
$Lt_4_1794:
//<loop> Loop body line 24, nesting depth: 1, iterations: 300
.loc 2 182 0
add.f32  %f2, %f1, %f2;
.loc 2 183 0
add.f32  %f1, %f1, %f2;
.loc 2 182 0
add.f32  %f2, %f1, %f2;
.loc 2 183 0
．．．．．．．
．．．．．．．
.loc 2 182 0
add.f32  %f2, %f1, %f2;
.loc 2 183 0
add.f32  %f1, %f1, %f2;
add.s32  %r1, %r1, 3750;
mov.u32  %r2, 1125000;
setp.ne.s32  %p1, %r1, %r2;
@%p1 bra  $Lt_4_1794;
.loc 2 185 0
add.f32  %f3, %f1, %f2;
ld.param.u32  %r3, [__cudaparm__Z5test1Pf_res];
cvt.u32.u16  %r4, %tid.x;
mov.u16  %rh1, %ctaid.x;
mov.u16  %rh2, %ntid.x;
mul.wide.u16  %r5, %rh1, %rh2;
add.u32  %r6, %r4, %r5;
mul.lo.u32  %r7, %r6, 4;
add.u32  %r8, %r3, %r7;
st.global.f32  [%r8+0], %f3;
.loc 2 186 0
exit;
$LDWend__Z5test1Pf:
} // _Z5test1Pf
test1程序的主要指令序列是7500条add.f32指令，其它指令所占比例很小，可以基本忽略．

 

性能计算公式是：
(1E-9*K*10*FLOP)/(tend-tbeg)
tend和tbeg是用QueryPerformanceCounter获取的时间(秒).

 

下面是运行结果：
时间：0.30173842s
实测性能：0.0745679Gflops

频率/实测性能=1.5/0.0745679=20.116>4!

 

多出来的周期GPU在干什么哪？？？迷惑！！！