cupy questions on parameter passing to RawKernel

I'm having a question with the parameter passing of the RawKernel. If the parameter list contains double parameter, the running speed dropped dramatically. below is my test code

import cupy as cp
import numpy as np
import time

kernelStr = """
extern "C"
{
    __global__ void doubleAdd(
        float* output,
        int width, 
        int height,
        double inc
        )
    {
        unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
        unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
        output[y * width + x] += inc;
    }
    
    __global__ void floatAdd(
        float* output,
        int width, int height,
        float inc
        )
    {
        unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
        unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
        output[y * width + x] += inc;
    }
}
"""

doubleAdd = cp.RawKernel(kernelStr, "doubleAdd")
floatAdd = cp.RawKernel(kernelStr, "floatAdd")


if __name__ == '__main__':

    width = 4096
    height = 4096
    block = (32, 32)
    grid = ((width + block[0] - 1) // block[0], (width + block[1] - 1) // block[1])


    x = cp.zeros((1, 1))

    s = time.time()
    for i in range(10):
        doubleAdd(grid, block, (cp.zeros(shape=(height, width), dtype=cp.float32), int(width), int(height), float(1.0)))
        cp.cuda.Stream.null.synchronize()
    print("doubleAdd time used: ", time.time() - s, "s")

    s = time.time()
    for i in range(10):
        floatAdd(grid, block, (cp.zeros(shape=(height, width), dtype=cp.float32), int(width), int(height), np.float32(1.0)))
        cp.cuda.Stream.null.synchronize()
    print("floatAdd time used: ", time.time() - s, "s")

and here is the result:

doubleAdd time used: 0.3446180820465088 s floatAdd time used: 0.008989095687866211 s

My GPU is TITAN X

Thanks!