 
 // Uncomment to use the live Kinect Camera
 //#define KINECT
-#define IMG_WIDTH 640
-#define IMG_HEIGHT 480
 
 __global__ void DiamondDotProduct(float *p, float *d, int w, int h)
 {

    // Display images
	showImage("Input Depth Image", mInDepth, 40, 100);
    convert_layered_to_mat(mOutDepth, fOutDepth);
	showImage("Output Depth Image", mOutDepth, 100+IMG_WIDTH, 100);

    // Save images
    cv::imwrite("image_input.png",mInDepth*255.f);  // "imwrite" assumes channel range [0,255]


    // Allocate memory on the GPU and copy data
    float *dU, *dV, *dP, *dD, *dDiaD;
    cudaMalloc(&dU, (size_t)width*height*sizeof(float)); CUDA_CHECK;
    cudaMalloc(&dV, (size_t)width*height*sizeof(float)); CUDA_CHECK;
    cudaMalloc(&dD, (size_t)width*height*sizeof(float)); CUDA_CHECK;
    cudaMalloc(&dP, (size_t)3*width*height*sizeof(float)); CUDA_CHECK;
    cudaMalloc(&dDiaD, (size_t)3*width*height*sizeof(float)); CUDA_CHECK;
    cudaMemcpy(dU, fInDepth, (size_t)width*height*sizeof(float), cudaMemcpyHostToDevice); CUDA_CHECK;
    cudaMemcpy(dV, dU, (size_t)width*height*sizeof(float), cudaMemcpyDeviceToDevice); CUDA_CHECK;
    cudaMemset(dD, 0, (size_t)width*height*sizeof(float)); CUDA_CHECK;
    cudaMemset(dP, 0, (size_t)3*width*height*sizeof(float)); CUDA_CHECK;
    cudaMemset(dDiaD, 0, (size_t)3*width*height*sizeof(float)); CUDA_CHECK;

    // Init block and grid sizes
    dim3 block = dim3(blockX, blockY, blockZ);
    dim3 grid = dim3((width+block.x-1)/block.x, (height+block.y-1)/block.y, 1);

    for(int n=0; n<N; n++)
    {
        DiamondDotProduct<<<grid, block>>>(dP, dD, width, height);
        cudaDeviceSynchronize();

        theta *= decay;
        ComputeU<<<grid, block>>>(dV, dD, dU, width, height, theta);
        cudaDeviceSynchronize();
        
        DiamondOperator<<<grid, block>>>(dU, dDiaD, width, height);
        cudaDeviceSynchronize();

        ComputeDualVariable<<<grid, block>>>(dDiaD, dP, width, height, theta, tau);
        cudaDeviceSynchronize();
    }

    DiamondDotProduct<<<grid, block>>>(dP, dD, width, height);
    cudaDeviceSynchronize();

    //theta *= decay;
    ComputeU<<<grid, block>>>(dV, dD, dU, width, height, theta);
    
    // Copy data back to CPU
    cudaMemcpy(fOutDepth, dU, (size_t)width*height*sizeof(float), cudaMemcpyDeviceToHost); CUDA_CHECK;
    
    // Deallocate memory on the GPU
    cudaFree(dU); CUDA_CHECK;
