萌新求助：kernel函数的一般参数是怎么传递的？使用常量内存代替参数没有提速？

[Reehac2016]

写了一个kernel函数，采用如下形式：
__global__ Func1（int * P1,float P2）
{
id=blockid.x*T_NUM＋threaded.x;
float t=id*P2;
........
}

在host端：
float a=3.14;
int b_num=128;
int t_num=128;
int * Data;
Cudamalloc( Data..........);

Func1<<<b_num,t_num,0>>>(Data,a);

程序顺利编译，运行了，但是速度没有达到老板的要求。
于是想办法：
哦，参数P2是在CPU端直接定义的值，然后直接作为参数传递给kernel函数。会有一次内存到显存的拷贝。该kernel函数被调用过1000次，也就是1000次hosttodevice的拷贝，耗时。
于是用常量内存代替该参数，并且在初始化时就对该常量赋值。之后1000次调用只读常量内存，不写。
__constant float pi=3.14;
__global__ Func1（int * P1）
{
id=blockid.x*T_NUM＋threaded.x;
float t=id*pi;
........
}

Func1<<<b_num,t_num,0>>>(Data);


编译，再运行。
速度没有丝毫提升，why？