     cudaFree(fp_dev);
     cudaFree(gm_dev);
     cudaFree(gp_dev);
+
+    cudaEventDestroy(start);
+    cudaEventDestroy(stop);
 #endif
 
     FCLAW_ASSERT(ierror == 0);

    dim3 dimGrid(mx, my);
    dim3 dimBlock(1, 1);
    int x_stride = mx + 2 * mbc;
    fclaw2d_timer_start (&glob->timers[FCLAW2D_TIMER_CUDA_KERNEL1]);
    cudaclaw5_update_q_cuda<<<dimGrid, dimBlock>>>(x_stride, mbc, dtdx, dtdy,
                                                   qold_dev, fm_dev, fp_dev,
                                                   gm_dev, gp_dev);
    fclaw2d_timer_stop (&glob->timers[FCLAW2D_TIMER_CUDA_KERNEL1]);
    cudaError_t code = cudaPeekAtLastError();
    if(code!=cudaSuccess){
        printf("ERROR: %s\n",cudaGetErrorString(code));

    double* gm_dev;
    double* gp_dev;

    cudaEventRecord(start);
    cudaMalloc((void**)&qold_dev, size * sizeof(double));
    cudaMalloc((void**)&fm_dev, size * sizeof(double));
    cudaMalloc((void**)&fp_dev, size * sizeof(double));
    cudaMalloc((void**)&gm_dev, size * sizeof(double));
    cudaMalloc((void**)&gp_dev, size * sizeof(double));
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&milliseconds, start, stop);
    glob->timers[FCLAW2D_TIMER_CUDA_ALLOCATE].cumulative += milliseconds*1e-3;

    cudaMemcpy(qold_dev, qold, size * sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(fm_dev, fm, size * sizeof(double), cudaMemcpyHostToDevice);
