 #endif
 
 	// Default setting for block sizes
-	size_t blockX = 32, blockY = 8, blockZ = 1;
+	size_t blockX = 64, blockY = 4, blockZ = 1;
 	getParam("blockX", blockX, argc, argv);
 	getParam("blockY", blockY, argc, argv);
 	getParam("blockZ", blockZ, argc, argv);


// Uncomment to use the live Kinect Camera
//#define KINECT
#include "aux.h"	// Helping functions for CUDA GPU Programming
#include <iostream>	// For standard IO on console
#include "constant.cuh"

    timer.start();

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

    // Setup texture reference for V
    vTexRef.addressMode[0] = cudaAddressModeClamp;
    vTexRef.addressMode[1] = cudaAddressModeClamp;
    vTexRef.filterMode = cudaFilterModeLinear;
    vTexRef.normalized = false;
    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaBindTexture2D(NULL, &vTexRef, dV, &desc, KINECT_SIZE_X, KINECT_SIZE_Y, KINECT_SIZE_X*sizeof(float));

    // Init block, grid, and shared memory size
    dim3 block = dim3(blockX, blockY, blockZ);
    dim3 grid = dim3((KINECT_SIZE_X+block.x-1)/block.x, (KINECT_SIZE_Y+block.y-1)/block.y, 1);
    size_t smBytes = (size_t)block.x*block.y*block.z*sizeof(float);

    // Iterate through main computation
    for(int n=0; n<N; n++)
    {
        theta *= decay;
        UpdateImageAndDualVariable<<<grid, block, smBytes>>>(dU, dP, KINECT_SIZE_X, KINECT_SIZE_Y, tau, theta); CUDA_CHECK;
        cudaDeviceSynchronize();
    }

    // Compute final output image
    theta *= decay;
    UpdateImageAndDualVariable<<<grid, block, smBytes>>>(dU, dP, KINECT_SIZE_X, KINECT_SIZE_Y, tau, theta); CUDA_CHECK;
    
    // Copy data back to CPU
    cudaMemcpy(fOutDepth, dU, (size_t)KINECT_SIZE_X*KINECT_SIZE_Y*sizeof(float), cudaMemcpyDeviceToHost); CUDA_CHECK;
    
    // Deallocate and unbind memory on the GPU
    cudaFree(dU); CUDA_CHECK;
    cudaFree(dV); CUDA_CHECK;
    cudaFree(dP); CUDA_CHECK;
    cudaUnbindTexture(vTexRef);

    // Display output image and save
    convert_layered_to_mat(mOutDepth, fOutDepth);
